我对来自 NVidia 的 2 份文件感到困惑。“CUDA 最佳实践”描述了共享内存是按银行组织的，通常在 32 位模式下，每 4 个字节就是一个银行（这就是我的理解）。然而，使用 CUDA 的并行前缀和（扫描）详细说明了由于存储库冲突，应如何将填充添加到扫描算法中。
对我来说问题是，这个算法的基本类型是浮点数，它的大小是 4 个字节。因此每个浮动都是一个银行，没有银行冲突。
那么我的理解是否正确——即，如果您使用 4*N 字节类型，您不必担心银行冲突，因为根据定义不会有冲突？如果不是，我应该如何理解它（何时使用填充）？
您可能对此网络研讨会感兴趣，来自NVIDIA CUDA 网络研讨会页面 共享内存（包括存储库）也在本网络研讨会的幻灯片 35-45 中进行了描述。
通常，当两个不同的线程试图（从同一内核指令）访问共享内存中的低 4 位（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
让我们仔细看看上面的 2-way bank 冲突。由于我们乘以2，线程 0 访问了 bank 0 中的位置 0，但线程 16 访问了也在threadIdx.xbank 0 中 的位置 32 ，从而产生了 bank 冲突。对于上面的 32 路示例，所有地址都对应于银行 0。因此，必须发生 32 次到共享内存的事务才能满足此请求，因为它们都是序列化的。
所以回答这个问题，如果我知道我的访问模式会是这样的，例如：
my = A[threadIdx.x*32];
然后我可能想要填充我的数据存储，以便它A[32]是一个虚拟/填充位置，A[64]等等A[96] 。然后我可以像这样获取相同的数据：
my = A[threadIdx.x*33];
并在没有银行冲突的情况下获取我的数据。
希望这可以帮助。
你的理解是错误的。当来自同一个 warp 的线程访问驻留在同一个 bank 中的不同值时，就会发生 Bank 冲突。
来自 CUDA C 编程指南：
为了实现高带宽，共享内存被分成大小相等的内存模块，称为银行，可以同时访问。因此，可以同时处理由落在 n 个不同存储器组中的 n 个地址构成的任何存储器读取或写入请求，从而产生的总带宽是单个模块带宽的 n 倍。
但是，如果一个内存请求的两个地址落在同一个内存 bank 中，就会发生 bank 冲突，访问必须串行化。硬件根据需要将具有存储体冲突的内存请求拆分为多个单独的无冲突请求，从而将吞吐量降低等于单独内存请求数量的因子。如果单独的内存请求的数量为 n，则称初始内存请求会导致 n-way bank 冲突。
填充用于避免银行冲突。当您知道您的共享内存访问模式时，您可以确定如何填充您的共享内存阵列以避免银行冲突。
例如，假设您有__shared__ float x[32][32];并且每个线程索引为 tid 的线程都像这样访问 x somevariable = x[tid][0];。这将导致 32-way bank 冲突，因为所有线程都从同一个 bank 访问不同的值。
为避免冲突，您在第一个维度中用一个元素填充数组：__shared__ float x[32][33];。这将完全消除存储库冲突，因为现在每一行都有一个存储库位置，相对于前一行偏移一个。