为什么这个内核在 GK210 上没有达到峰值 IPC？

Question

我认为尝试编写一个达到峰值 IPC 的 CUDA 内核对我来说是有教育意义的，所以我想出了这个内核（为简洁起见省略了主机代码，但可用 here）

#define WORK_PER_THREAD 4

__global__ void saxpy_parallel(int n, float a, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    i *= WORK_PER_THREAD;
    
    if (i < n)
    {
        #pragma unroll
        for(int j=0; j<WORK_PER_THREAD; j++)
            y[i+j] = a * x[i+j] + y[i+j];
    }
}

我运行这个内核在 GK210 上，有 n=32*1000000 个元素，并期望看到一个接近 4 的 IPC，但最终得到了一个糟糕的 0.186

ubuntu@ip-172-31-60-181:~/ipc_example$ nvcc saxpy.cu
ubuntu@ip-172-31-60-181:~/ipc_example$ sudo nvprof --metrics achieved_occupancy --metrics ipc ./a.out

==5828== NVPROF is profiling process 5828, command: ./a.out
==5828== Warning: Auto boost enabled on device 0. Profiling results may be inconsistent.
==5828== Profiling application: ./a.out
==5828== Profiling result:
==5828== Metric result:
Invocations                               Metric Name                        Metric Description         Min         Max         Avg
Device "Tesla K80 (0)"
    Kernel: saxpy_parallel(int, float, float*, float*)
          1                        achieved_occupancy                        Achieved Occupancy    0.879410    0.879410    0.879410
          1                                       ipc                              Executed IPC    0.186352    0.186352    0.186352

当我设置 WORK_PER_THREAD=16 时，我更加困惑，导致启动的线程更少，但是 16，而不是 4，每个执行的独立指令，IPC 下降到 0.01

我的两个问题是：

GK210 的 IPC 峰值是多少？我认为每个周期 8 = 4 个 warp 调度程序 * 2 个指令调度，但我想确定一下。
为什么这个内核实现了这么低的IPC却实现了很高的占用率，为什么IPC随着WORK_PER_THREAD的增加而降低，如何提高这个内核的IPC？

Answer 1

What is the peak IPC I can expect on a GK210?

每个SM的峰值IPC等于SM中warp调度器的数量乘以每个warp调度器的发布率。此信息可在特定 GPU 的白皮书中找到。 GK210白皮书是here。从该文档（例如 p8 上的 SM 图）我们看到每个 SM 都有 4 个能够进行双重发布的 warp 调度器。因此，理论上可实现的 IPC 峰值是每个 SM 每个时钟 8 条指令。（但实际上，即使是精心设计的代码，您也不太可能看到高于 6 或 7 的值）。

Why does this kernel achieve such low IPC while achieved occupancy is high, why does IPC decrease as WORK_PER_THREAD increases, and how can I improve the IPC of this kernel?

您的内核几乎在每个操作中都需要全局事务。全局加载甚至 L2 缓存加载都有延迟。当你所做的一切都依赖于这些时，就无法避免延迟，因此你的 warp 经常停滞不前。 GK210 上每个 SM 的可观察 IPC 峰值在 6 附近，但连续加载和存储操作不会达到这一点。您的内核为每个 multiply/add 执行 2 次加载和一次存储（总共移动 12 个字节）。你将无法改进它。（你的内核占用率高，因为 SM 加载了 warp，但 IPC 低，因为这些 warp 经常停滞，无法发出指令，等待加载操作的延迟到期。）你需要找到其他有用的工作要做。

那可能是什么？好吧，如果您执行矩阵乘法运算，它具有大量数据重用且每个数学运算的字节数相对较少，您可能会看到更好的测量结果。

你的代码呢？有时你需要做的工作就是这样。我们称之为内存绑定代码。对于这样的内核，用于判断“好”的品质因数不是 IPC，而是实现的带宽。如果您的内核需要加载和存储特定数量的字节来执行其工作，那么如果我们将内核持续时间与内存事务进行比较，我们可以获得良好的衡量标准。换句话说，对于纯内存绑定代码（即您的内核），我们将通过测量加载和存储的总字节数来判断优劣（分析器对此有度量，或者对于简单代码，您可以通过检查直接计算它），并将其除以内核持续时间。这给出了实现的带宽。然后，我们将其与基于代理测量的可实现带宽进行比较。一个可能的代理测量工具是 bandwidthTest CUDA 示例代码。

当这两个带宽的比率接近 1.0 时，考虑到它试图完成的内存限制工作，您的内核表现“良好”。

为什么这个内核在 GK210 上没有达到峰值 IPC？

Why does this kernel not achieve peak IPC on a GK210?

cuda