如何确保与 CUDA 中的 3D 共享数据访问没有银行冲突

Posted 2023-03-23

【中文标题】如何确保与 CUDA 中的 3D 共享数据访问没有银行冲突

【英文标题】：How to ensure no bank conflict with 3D shared data access in CUDA

【发布时间】：2014-04-05 13:42:01

【问题描述】：

我正在使用 CUDA 对几个相同大小的大型三维数据集进行一些操作，每个数据集都由浮点数组成。

下面的例子：

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) 块上运行，每个块中有 (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 案例，只是在应用 Z 次的 2D 方案中考虑它。

所以通过这种想法，通过以下方式访问s_inC：

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

会阻止半扭曲中的线程同时访问同一个bank，共享内存应该声明为：

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

（省略同步、边界检查、包含非常边缘的情况 in_C[i+j+k+1] 等）。

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

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

【参考方案1】：

我认为你关于银行冲突预防的结论是有问题的。

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

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

不会产生银行冲突。

与此相反，使用8x8x8 线程块，然后访问类似

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

会产生银行冲突。

下图说明了这一点，其中黄色单元格表示来自同一经线的线。该图报告，对于每个32 位库，访问它的线程作为元组(threadIdx.x, threadIdy.y, threadIdz.z)。红色单元格是您正在使用的填充单元格，任何线程都无法访问它们。