cuda - 跨步访问的银行冲突

Question

我正在阅读 CUDA_C_Programming_Guide，在共享内存主题中，我遇到了一个示例：设备计算能力：1.0，共享内存中有 16 个库

extern __shared__ float shared[]; 
float data = shared[BaseIndex + s * tid];

在他们得出的解释中，“s”必须是奇数，谁能帮我理解s偶数时会发生什么，s奇数时会发生什么？

Answer 1

奇数s的结论不容易直接看出，但是如果你尝试在bank冲突发生时推导（两个线程tid和tid'访问同一个bank），假设32是bank数：

s*tid == s*tid' (mod 32)

s*tid == s*(tid + n) (mod 32) 其中 tid' = tid + n

s*tid == s*tid + s*n (mod 32)

s*n == 0 (mod 32)

n = (32/d)*k 对于一些 k 和 d = gcd(s, 32)

所以当32小于等于32/d时不会发生bank冲突

并且由于 d = gcd(s, 2^5)，s 必须是奇数。

关于您在评论中的问题，我没有完全理解您不理解的内容，但简单的解释是：如果两个线程尝试访问同一个银行（这意味着访问同一行中的两个单词），则访问将被序列化。