共享内存何时真正需要填充？

Question

我对NVidia的两份文件感到困惑。CUDA最佳实践指出，共享内存是以银行方式组织的，在32位模式下，每4个字节就是一个银行（这是我的理解）。然而使用CUDA进行并行前缀求和（扫描）详细介绍了由于银行冲突而需要向扫描算法添加填充的细节。

问题在于，对于我所提供的此算法的基本类型为float，其大小为4字节。因此，每个float是一个银行，不存在银行冲突。

因此，我的理解是否正确--即，如果您处理的是4*N字节类型，则无需担心银行冲突，因为根据定义不会存在？如果不是，请告诉我应该如何理解它（何时使用填充）？

- greenoldman

Answer 1

您可能对NVIDIA CUDA网络研讨会页面上的此网络研讨会感兴趣。从此网络研讨会的幻灯片35-45中还介绍了共享内存，包括银行。

一般来说，当两个不同的线程尝试从共享内存中访问（来自同一内核指令）地址的低4位（pre-cc2.0设备）或5位（cc2.0及更高版本设备）相同时，就会发生共享内存银行冲突。当发生银行冲突时，共享内存系统会序列化访问相同银行中的位置，从而降低性能。填充尝试避免某些访问模式的这种情况。请注意，对于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中的线程之间访问共享内存，可能会遇到银行冲突：

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

让我们仔细看一下上面的双向银行冲突。由于我们将threadIdx.x乘以2，因此线程0访问位于银行0中的位置0，但线程16访问位于银行0中的位置32，从而创建了一个银行冲突。对于上面的32路示例，所有地址都对应于银行0。因此，必须发生32次共享内存事务才能满足此请求，因为它们都是串行化的。

因此，回答这个问题，如果我知道我的访问模式会像这样：

my = A[threadIdx.x*32]; 

我可能需要填充我的数据存储，使得A[32]，A[64]，A[96]等位置成为虚拟/填充位置。这样，我可以像这样获取相同的数据：

my = A[threadIdx.x*33]; 

并且无需发生银行冲突即可获取我的数据。

希望这能帮到你。

- Robert Crovella

---

谢谢你，再次感谢 :-) 你拯救了我的理解！ - greenoldman

您IP地址为143.198.54.68，由于运营成本限制，当前对于免费用户的使用频率限制为每个IP每72小时10次对话，如需解除限制，请点击左下角设置图标按钮（手机用户先点击左上角菜单按钮）。 - user1197918

抱歉，我觉得你的问题不太清晰。 填充元素的计算并不是 Kepler 独有的。填充必须根据预期的访问模式进行组织。 我不知道共享内存配置和双精度支持之间有什么联系。 CUDA 7 支持的所有 CUDA GPU 都支持双精度。 Kepler GPU 支持 8 字节银行模式。也许这就是你所指的。 - Robert Crovella

你们应该在文档中更明确地说明。我仔细阅读了指南，但在这种情况下为什么会有32个冲突，直到我偶然发现这个帖子，我才知道原因。 - Íhor Mé

Answer 2

你的理解是错误的。当来自同一warp的线程正在访问驻留在同一个bank中的不同值时，就会发生Bank冲突。

从CUDA C编程指南中得知： 为了实现高带宽，共享内存被划分为大小相等的内存模块，称为Banks，可以同时访问。任何由n个地址组成且落在n个独立内存bank中的内存读写请求都可以同时完成，从而产生的总带宽是单个模块带宽的n倍。 但是，如果一个内存请求的两个地址落入同一个内存Bank中，则会发生Bank冲突，需要序列化访问。硬件将具有Bank冲突的内存请求分割为尽可能多的单独无冲突请求，通过将吞吐量减少与多个内存请求数量相等的因子来进行，如果单独的内存请求数量为n，则初始内存请求则导致n路Bank冲突。

填充用于避免Bank冲突。当您知道共享内存访问模式时，可以确定如何填充共享内存数组以避免Bank冲突。 例如，如果假设您有__shared__float x [32] [32] ;并且每个线程的线程索引tid都像这样访问x：somevariable = x [tid] [0];这将导致32路Bank冲突，因为所有线程都从同一Bank中访问不同的值。为了避免冲突，您可以在第一个维度上填充一个更多的元素:__shared__float x [32] [33];。这将完全消除Bank冲突，因为现在每行都有一个与上一行偏移一个bank位置的银行位置。

- RoBiK

---

谢谢你的加一，我接受了Robert的答案，因为他提供了示例。 - greenoldman