CUDA在多个数据上合并了一个扭曲

CUDA coalesced one warp on multiple data

本文关键字：一个 合并 数据 CUDA      更新时间：2023-10-16

我有一个关于联合cuda访问的基本问题。

例如，我有一个由32个元素和32个线程组成的数组，每个线程访问一个元素。

__global__ void co_acc ( int A[32], int B[32] ) {
    int inx = threadIdx.x + (gridDim.x * blockDim.x);
    B[inx] = A[inx]
}

现在，我想知道的是：如果我有32个线程，但有一个64个元素的数组，那么每个线程必须复制2个元素。为了保持联合访问，我应该转换我拥有的线程数的数组访问索引。例如：ID为0的线程将访问A[0]和A[0+32]。我的假设是对的吗？

__global__ void co_acc ( int A[64], int B[64] ) {
    int inx = threadIdx.x + (gridDim.x * blockDim.x);
    int actions = 64/blockDim.x;
    for ( int i = 0; i < actions; ++i )
        B[inx+(i*blockDim.x)] = A[inx+(i*blockDim.x)]
}

为了保持联合访问，我应该将数组访问的索引移动我所拥有的线程数。例如：ID为0的线程将访问A[0]和A[0+32]。我的假设是对的吗？

是的，这是一个正确的方法。

严格来说，这不是应该，而是可以：只要warp请求中的所有线程都位于同一（对齐）128字节行内，任何内存访问都将合并。这意味着你可以排列线程索引，你的访问仍然会被合并（但既然你可以做简单的事情，为什么要做复杂的事情呢）。

另一种解决方案是让每个线程加载一个int2:

__global__ void co_acc ( int A[64], int B[64] ) {
    int inx = threadIdx.x + (gridDim.x * blockDim.x);
    reinterpret_cast<int2*>(B)[inx] = reinterpret_cast<int2*>(A)[inx];
}

这是（在我看来）更简单、更清晰的代码，可能会提供更好的性能，因为这可能会减少编译器发出的指令数量和内存请求之间的延迟（免责声明：我没有尝试过）。

注意：正如Robert Crovella在评论中提到的那样，如果你真的在使用32个线程的线程块，那么你很可能严重不足了GPU的容量。

相关文章：