我怎样才能确保编译器从全局内存中并行化我的加载？

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 线程将开始加载第二次迭代的数据，而其余线程仍在计算结果第一次迭代。

我在考虑分配多少线程用于加载和计算，我推断我只需要一个 warp 来加载，而不管缓冲区大小，因为内部 for 循环由来自全局的独立加载组成记忆：它们可以同时飞行。这是一个有效的推理线吗？

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

buff 或 globalMem 上的附加限定符（const、__restrict__ 等）是否有助于确保编译器执行我想要的操作？

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

Answer 1

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

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

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

您也可以使用部分展开。用 #pragma unroll 2 注释加载循环将允许编译器发出两次加载，然后每两次循环迭代两次存储，从而实现与加倍 nLoadThreads 类似的效果。用更高的数字替换 2 是可能的，但你会在某个时候达到最大的飞行事务数（使用 float2 或 float4 移动以传输更多具有相同事务数的数据）。此外，很难预测编译器是否更喜欢对指令重新排序，而不是更复杂的代码成本，以完成最终的、可能是部分的、通过展开循环的行程。

所以建议是：

使用尽可能多的加载线程。
通过对循环迭代次数进行模板化并针对所有可能的循环次数（或最常见的次数，使用通用回退）对其进行实例化，或使用部分循环展开来展开加载循环。
如果数据适当对齐，将其移动为 float2 或 float4 以移动具有相同事务数的更多数据。

我怎样才能确保编译器从全局内存中并行化我的加载？

How can I make sure the compiler parallelizes my loads from global memory?

cuda

gpu