2.x 设备中的存储库冲突答案

【问题标题】：Bank conflicts in 2.x devices2.x 设备中的存储库冲突
【发布时间】：2012-07-01 19:21:03
【问题描述】：

什么是带有 2.x 设备的设备中的存储库冲突？据我了解 CUDA C 编程指南，在 2.x 设备中，如果两个线程访问同一个共享内存库中的同一个 32 位字，则不会导致库冲突。相反，这个词被广播。当两个线程在同一个共享内存bank中写入相同的32位字时，只有一个线程成功。

由于片上内存为 64 KB（共享内存为 48 KB，L1 为 16 KB，反之亦然），并且它分为 32 个存储区，我假设每个存储区包含 2 KB。所以我认为如果两个线程访问同一个共享内存银行中的两个不同的 32 位字，就会出现银行冲突。这是正确的吗？

【问题讨论】：

没错。我有点想用 -1 来回答你的问题，因为在 CUDA C 编程指南中对银行冲突进行了如此详尽的解释。

标签： cuda gpu shared-memory bank-conflict

【解决方案1】：

您的描述是正确的。有许多访问模式会产生银行冲突，但这里有一个简单且常见的示例：跨步访问。

__shared__ int smem[512];

int tid = threadIdx.x;

x = smem[tid * 2]; // 2-way bank conflicts
y = smem[tid * 4]; // 4-way bank conflicts
z = smem[tid * 8]; // 8-way bank conflicts
// etc.

Bank ID = index % 32，因此如果查看 x、y 和 z 访问中的地址模式，您可以看到在 32 个线程的每个 warp 中，对于 x，2 个线程将访问每个 bank，对于 y，每个 bank 有 4 个线程访问，对于 z，每个 bank 有 8 个线程访问。

【讨论】：