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

我对 NVidia 的 2 个文档感到困惑。CUDA 最佳实践描述了共享内存是按bank组织的，一般来说在32位模式下每4个字节就是一个bank（我就是这么理解的）。然而使用 CUDA 并行前缀和（扫描）详细介绍了由于库冲突而应如何将填充添加到扫描算法中。

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

我的理解是否正确——也就是说，如果你继续工作4*N-byte 类型，您不必担心银行冲突，因为根据定义，不会有银行冲突？如果不是，我应该如何理解它（何时使用填充）？

您可能感兴趣这次网络研讨会来自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

让我们仔细看看上面的双向银行冲突。既然我们在乘法threadIdx.x到 2，线程 0 访问存储体 0 中的位置 0，但线程 16 访问位置 32，即also在银行 0 中，从而造成银行冲突。对于上面的 32 路示例，所有地址都对应于存储体 0。因此，必须发生对共享内存的 32 个事务才能满足此请求，因为它们都是序列化的。

所以要回答这个问题，if我知道我的访问模式将是这样的，例如：

my = A[threadIdx.x*32]; 

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

my = A[threadIdx.x*33]; 

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

希望这可以帮助。