c++ - CUDA 在多个数据上合并一个 warp

Question

我有一个关于 union 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 个元素。为了保持 union 访问，我应该转移通过我拥有的线程数访问数组的索引。例如: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)]
}

Answer 1

To keep a coalesced access, I should shift the index for the array access by the number of threads I have. eg: Thread with ID 0 will access A[0] and A[0+32]. Am I right with this assumption?

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

严格来说，这不是应该，而是可以:任何内存访问都将被合并，只要 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];
}

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

Note: as Robert Crovella has mentioned in his comment, if you really are using thread blocks of 32 threads, then you are likely seriously underusing the capacity of your GPU.