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

Question

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

#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 的糟糕 IPC

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

我的两个问题是：

1. 我可以在 GK210 上期望的峰值 IPC 是多少？我认为它是 8 = 4 个 warp 调度程序 * 每个周期 2 个指令调度，但我想确定一下。
2. 为什么这个内核在实现占用率高的情况下实现了如此低的IPC，为什么IPC随着WORK_PER_THREAD的增加而减少，我该如何提高这个内核的IPC？

Answer 1

我可以在 GK210 上期望的峰值 IPC 是多少？

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

为什么这个内核在实现占用率高的情况下实现了如此低的IPC，为什么IPC随着WORK_PER_THREAD的增加而减少，我该如何提高这个内核的IPC？

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

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

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

当这两个带宽的比率接近 1.0 时，您的内核运行“良好”，考虑到它正在尝试执行的内存绑定工作。