memory - 二维 block 的 CUDA 合并访问

Question

对于 1D 情况，我非常了解 CUDA 中全局内存的整个合并访问要求。

但是，我对二维情况有点困惑(即我们有一个由 2D block 组成的 2D 网格)。

假设我有一个向量 in_vector 并且在我的内核中我想以合并的方式访问它。就像这样:

__global__ void my_kernel(float* out_matrix, float* in_vector, int size)
{
   int i = blockIdx.x * blockDim.x + threadIdx.x;
   int j = blockIdx.y * blockDim.y + threadIdx.y;
   // ...
   float vx = in_vector[i]; // This is good. Here we have coalesced access
   float vy = in_vector[j]; // Not sure about this. All threads in my warp access the same global address. (See explanation)
   // ...
   // Do some calculations... Obtain result
}

根据我对这种 2D 情况的理解， block 内的线程以列为主的方式“排列”。例如:假设 (threadIdx.x, threadIdx.y) 表示法:

• 第一个扭曲将为:(0, 0), (1, 0), (2, 0), ..., (31, 0),
• 第二个扭曲将为:(0, 1), (1, 1), (2, 1), ..., (31, 1),
• 等等...

在这种情况下，调用 in_vector[i] 为我们提供了合并访问，因为同一 warp 中的每个连续线程都将访问连续的地址。然而，调用 in_vector[j] 似乎是一个坏主意，因为每个连续的线程都将访问全局内存中的相同地址(例如，warp 0 中的所有线程都将访问 in_vector[0]，这将为我们提供32种不同的全局内存请求)

我理解正确吗？如果是这样，我如何使用in_vector[j]对全局内存进行合并访问？

Answer 1

您在问题中显示的内容仅适用于某些 block 大小。您的“合并”访问权限:

int i = blockIdx.x * blockDim.x + threadIdx.x;
float vx = in_vector[i];

仅当 blockDim.x 大于或等于 32 时，才会导致从全局内存中合并访问 in_vector。即使在合并的情况下，一个线程中的每个线程共享相同 threadIdx.x 值的 block 从全局内存中读取相同的单词，这似乎违反直觉且浪费。

确保每个线程读取唯一且合并的正确方法是计算 block 内的线程数和网格内的偏移量，可能类似于:

int tid = threadIdx.x + blockDim.x * threadIdx.y; // must use column major order
int bid = blockIdx.x + gridDim.x * blockDim.y; // can either use column or row major
int offset = (blockDim.x * blockDim.y) * bid; // block id * threads per block
float vx = in_vector[tid + offset];

如果您的目的确实不是要读取每个线程的唯一值，那么您可以节省大量内存带宽并使用共享内存实现合并，如下所示:

__shared__ float vx[32], vy[32]; 

int tid = threadIdx.x + blockDim.x * threadIdx.y;

if (tid < 32) {
    vx[tid] = in_vector[blockIdx.x * blockDim.x + tid];
    vy[tid] = in_vector[blockIdx.y * blockDim.y + tid];
}
__syncthread();

您将得到一个将唯一值读取到共享内存中的单个扭曲。然后，其他线程可以从共享内存中读取值，而无需任何进一步的全局内存访问。请注意，在上面的示例中，我遵循了代码的约定，即使以这种方式读取 in_vector 两次不一定有多大意义。