CUDA 共享内存中的扁平化与多维数组

Question

在 cuda 共享内存中使用扁平化数组比使用多维数组有任何性能优势吗？

我知道当涉及到主机内存时，展平数组比多维数组具有某些优势 space，但我不确定这是否适用于 gpu 共享内存中的数组。这主要是因为我看到很多代码示例确实在共享内存中使用了多维，例如CUDA Best Practices Guide中的矩阵乘法示例。

Answer 1

多维数组最终被展平以映射到 flat/linear 共享内存地址 space。因此，使用其中一种并没有性能优势。

通过在扁平化上使用多维共享数组，您可以避免手动索引计算的负担。对于多维数组，索引计算会自动添加到幕后最终编译代码中，方便源码阅读。

与多维表示形式相比，使用平面表示形式的一个优势可能是您可以更轻松地推理和发现共享内存访问中的存储体冲突。考虑这个共享内存缓冲区：

__shared__ int A[ 64 ];

如果您像这样访问缓冲区：

int laneID = threadIdx.x & 31;
int ret    = A[ laneID * 2 ];

与以下示例相比，可能更容易区分线程之间的银行冲突：

__shared__ int B[ 32 ][ 2 ];

你在哪里访问它：

int laneID = threadIdx.x & 31;
int ret    = B[ laneID ][ 0 ];

请注意，以上两种情况在功能上是等价的。

Flattened vs multi-dimensional arrays in shared memory for CUDA