c - 如何确保CUDA中3D共享数据访问不发生bank冲突

Question

我正在使用 CUDA 对几个相同大小的大型三维数据集(每个数据集由 float 组成)执行一些操作。

示例如下:

out[i+j+k]=in_A[i+j+k]*out[i+j+k]-in_B[i+j+k]*(in_C[i+j+k] +1]-in_C[i+j+k]);

其中(numCols、numDepth 指 3D 集的 y 和 z 维度(例如 out、in_A、in_C 等)并且:

int tx=blockIdx.x*blockDim.x + threadIdx.x; int i=tx*numCols*numDepth;

int ty=blockIdx.y*blockDim.y + threadIdx.y; int j=ty*numDepth

int tz=blockIdx.z*blockDim.z + threadIdx.z; int k=tz;

我已将内核设置为在 (11,14,4) block 上运行，每个 block 中有 (8,8,8) 个线程。通过这种方式设置，每个线程对应于每个数据集中的一个元素。 为了与我设置内核的方式保持一致，我使用 3D 共享内存来减少 in_C 的冗余全局读取:

(8x8x9 而不是 8x8x8，这样边缘 in_C[i+j+k+1] 也可以加载)

__shared__ float s_inC[8][8][9];

还有其他 Stack Exchange 帖子 ( ex link ) 和 CUDA 文档，它们处理 2D 共享内存并描述了可以采取哪些措施来确保不存在存储体冲突，例如将列维度填充 1 并使用 threadIdx 访问共享数组.y 然后是 threadIdx.x，但我找不到一个描述使用 3D 情况时会发生什么的情况。

我想，同样的规则也适用于 2D 情况和 3D 情况，只需考虑在 2D 方案中应用 Z 次即可。

因此，通过这种想法，可以通过以下方式访问 s_inC:

s_inC[threadIdx.z][threadIdx.y][threadIdx.x]=in_C[i+j+k];

将阻止半扭曲中的线程同时访问同一存储体，并且共享内存应声明为:

__shared__ float s_inC[8][8+1][9];

(省略同步、边界检查、包含_C[i+j+k+1] 中的极端情况等)。

前两个假设是否正确并可以防止银行冲突？

我使用的是 Fermi 硬件，因此有 32 个 32 位共享内存库

Answer 1

我认为您关于银行冲突预防的结论值得怀疑。

假设8x8x8线程阻塞，那么访问就像

__shared__ int shData[8][8][8];
...
shData[threadIdx.z][threadIdx.y][threadIdx.x] = ...

将无银行冲突。

与此相反，使用 8x8x8 线程 block ，然后进行类似的访问

__shared__ int shData[8][9][9];
...
shData[threadIdx.z][threadIdx.y][threadIdx.x] = ...

会产生银行冲突。

下图对此进行了说明，其中黄色单元表示来自同一经纱的 yarn 。该图报告了每个 32 位存储体的线程以元组 (threadIdx.x, threadIdy.y, threadIdz.z) 的形式访问它。红色单元格是您正在使用的填充单元格，任何线程都不会访问它们。