cuda - 什么时候真正需要填充共享内存？

Question

我对来自 NVidia 的 2 个文件感到困惑。 “CUDA 最佳实践”描述了共享内存是按组组织的，通常在 32 位模式下，每 4 个字节是一个组(这就是我的理解)。然而 Parallel Prefix Sum (Scan) with CUDA详细介绍了由于组冲突而应如何将填充添加到扫描算法中。

对我来说问题是，这个算法的基本类型是浮点数，大小是 4 个字节。因此，每个 float 都是一家银行，不存在银行冲突。

那么我的理解是否正确——即，如果您使用 4*N 字节类型，则不必担心银行冲突，因为根据定义，不会有任何冲突？如果不是，我应该如何理解(何时使用填充)？

Answer 1

您可能对 this webinar 感兴趣来自 NVIDIA CUDA webinar page来自 this webinar 的幻灯片 35-45 中也描述了包括存储体在内的共享内存。 .

通常，当两个不同的线程尝试访问(从同一内核指令)共享内存中的低 4 位(cc2.0 之前的设备)或 5 位(cc2.0 和更新的设备)的地址相同。当确实发生 bank 冲突时，共享内存系统会序列化对同一 bank 中的位置的访问，从而降低性能。对于某些访问模式，填充尝试避免这种情况。请注意，对于 cc2.0 和更新版本，如果所有位都相同(即相同的位置)，这不会导致库冲突。

从图片上看，我们可以这样看:

__shared__ int A[2048];
int my;
my = A[0]; // A[0] is in bank 0
my = A[1]; // A[1] is in bank 1
my = A[2]; // A[2] is in bank 2
...
my = A[31]; // A[31] is in bank 31 (cc2.0 or newer device)
my = A[32]; // A[32] is in bank 0
my = A[33]; // A[33] is in bank 1

现在，如果我们在 warp 中跨线程访问共享内存，我们可能会遇到 bank 冲突:

my = A[threadIdx.x];    // no bank conflicts or serialization - handled in one trans.
my = A[threadIdx.x*2];  // 2-way bank conflicts - will cause 2 level serialization
my = A[threadIdx.x*32]; // 32-way bank conflicts - will cause 32 level serialization

让我们仔细看看上面的 2-way bank 冲突。由于我们正在乘以 threadIdx.x通过 2，线程 0 访问存储区 0 中的位置 0，但线程 16 访问也在存储区 0 中的位置 32，从而产生存储区冲突。对于上面的 32 路示例，所有地址都对应于 bank 0。因此，必须发生 32 个到共享内存的事务才能满足此请求，因为它们都是序列化的。

所以回答这个问题，如果我知道我的访问模式会是这样的，例如:

my = A[threadIdx.x*32]; 

然后我可能想要填充我的数据存储，以便 A[32]是一个虚拟/焊盘位置，正如 A[64] , A[96]等等。
然后我可以像这样获取相同的数据:

my = A[threadIdx.x*33]; 

并在没有银行冲突的情况下获取我的数据。

希望这可以帮助。