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

我正在使用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情况下发生的情况的描述。

我可以想象，仅在Z次应用2D方案中考虑一下，相同的规则就适用于2D情况就适用于3D情况。

因此，通过这种思考，可以通过以下方式访问s_inC：

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

会阻止一半线程束中的线程同时访问同一存储体，并且共享内存应声明为：

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

（省去同步，边界检查，在in_C [i + j + k + 1]中包含最边缘情况，等等）。

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

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

最佳答案

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

假设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)访问它的线程。红色单元格是您正在使用的填充单元格，任何线程都无法访问它们。