cuda - 如何确保编译器并行化全局内存中的加载？

Question

我编写了一个 CUDA 内核，如下所示:

int tIdx = threadIdx.x; // Assume a 1-D thread block and a 1-D grid
int buffNo = 0;
for (int offset=buffSz*blockIdx.x; offset<totalCount; offset+=buffSz*gridDim.x) {
    // Select which "page" we're using on this iteration
    float *buff = &sharedMem[buffNo*buffSz];
    // Load data from global memory
    if (tIdx < nLoadThreads) {
        for (int ii=tIdx; ii<buffSz; ii+=nLoadThreads)
            buff[ii] = globalMem[ii+offset];
    }
    // Wait for shared memory
    __syncthreads();
    // Perform computation
    if (tIdx >= nLoadThreads) {
        // Perform some computation on the contents of buff[]
    }
    // Switch pages
    buffNo ^= 0x01;
}

请注意，只有一个 __syncthreads()在循环中，所以第一个 nLoadThreads线程将开始加载第二次迭代的数据，而其余线程仍在计算第一次迭代的结果。

我正在考虑分配多少线程用于加载与计算，并且我推断我只需要一个用于加载的线程，无论缓冲区大小如何，因为内部 for 循环由来自全局内存的独立加载组成:他们可以同时飞行。这是一个有效的推理吗？

然而，当我尝试这个时，我发现(1)增加负载扭曲的数量可以显着提高性能，(2)nvvp中的反汇编显示buff[ii] = globalMem[ii+offset]被编译为从全局内存加载，然后是 2 条指令，然后存储到共享内存，这表明编译器在这里没有应用指令级并行性。

是否会在 const 上添加额外的限定符( __restrict__ 、 buff 等)或globalMem帮助确保编译器执行我想要的操作？

我怀疑这个问题与buffSz有关。在编译时未知(实际数据是二维的，适当的缓冲区大小取决于矩阵维度)。为了完成我想要的操作，编译器需要为运行中的每个 LD 操作分配一个单独的寄存器，对吧？如果我手动展开循环，编译器会重新排序指令，以便在相应的 ST 需要访问该寄存器之前有一些 LD 正在运行。我尝试了#pragma unroll但编译器仅展开循环而没有重新排序指令，因此这没有帮助。我还能做什么？

Answer 1

编译器没有机会将共享内存的存储重新排序，以远离全局内存的加载，因为紧随其后的是 __syncthreads() 屏障。 由于所有线程无论如何都必须在屏障处等待，因此使用更多线程进行加载会更快。这意味着更多的全局内存事务可以随时进行，并且每个加载线程必须减少全局内存延迟。

到目前为止，所有 CUDA 设备都不支持乱序执行，因此加载循环每次循环迭代都会产生一次全局内存延迟，除非编译器可以在存储之前展开它并重新排序加载。

为了允许完全展开，需要在编译时知道循环迭代的次数。您可以使用 talonmies 的循环行程模板建议来实现此目的。

您还可以使用部分展开。使用#pragma unroll 2注释加载循环将允许编译器发出两次加载，然后每两次循环迭代两次存储，从而达到与加倍nLoadThreads类似的效果。用更大的数字替换 2 是可能的，但您会在某个时刻达到飞行中的最大交易数量(使用 float2 或 float4 移动以相同数量的交易传输更多数据)。此外，很难预测编译器是否会更喜欢重新排序指令，而不是最终(可能是部分)展开循环的更复杂代码的成本。

所以建议是:

1. 使用尽可能多的加载线程。
2. 通过模板化循环迭代次数并针对所有可能的循环次数(或最常见的循环行程，带有通用后备)实例化它，或使用部分循环展开来展开加载循环。
3. 如果数据已适当对齐，请将其移动为 float2 或 float4，以便在事务数量相同的情况下移动更多数据。