1

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

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

4

1 回答 1

3

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

__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, 4 个线程将访问每个 bank,对于 z,8 个线程将访问每个 bank。

于 2012-07-01T11:19:47.877 回答