跨步访问的银行冲突

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），假设bank数为32：

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 for some k and d = gcd(s, 32)

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

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

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