optimization - CUDA中的共享存储库冲突: How memory is aligned to banks

Question

据我所知，共享内存被划分为存储体，并且多个线程对同一存储体内单个数据元素的访问将导致冲突(或广播)。

目前，我分配了一个相当大的数组，该数组在概念上表示两个矩阵的几对:

__shared__ float A[34*N]

其中 N是对的数目，一对的前16个浮点数是一个矩阵，随后的18个浮点数是第二个矩阵。

事实是，访问第一个矩阵没有冲突，但是访问第二个矩阵有冲突。这些冲突是不可避免的，但是，我的想法是，由于第二个矩阵为18，因此所有将来的矩阵都将与库对齐，因此发生的冲突将超过必要的数量。

是真的，如果可以的话，我该如何避免呢？

每次分配共享内存时，它都从新的存储区开始吗？所以我可能会做

__shared__ Apair1[34]
__shared__ Apair2[34]
...

有任何想法吗？

谢谢

Answer 1

如果您的矩阵对是连续存储的，并且通过线程索引线性访问元素，则不会有共享内存库冲突。

换句话说，如果您有:

A[0]  <- mat1 element1
A[1]  <- mat1 element2
A[2]  <- mat1 element3
A[15] <- mat1 element16
A[16] <- mat2 element1
A[17] <- mat2 element2
A[33] <- mat2 element18

您可以使用以下命令访问此文件:

float element;
element = A[pairindex * 34 + matindex * 16 + threadIdx.x];

然后，相邻线​​程正在访问矩阵中的相邻元素，并且您没有冲突。

在回应您的评论(如下)时，您似乎在理解上是错误的。确实有16个存储体(在当前的世代中为32个，下一代为Fermi)，但是连续的32位字驻留在连续的存储体中，即地址空间在存储体中交错。这意味着，只要您始终具有可以分解为 x + threadIdx.x的数组索引(其中 x不依赖于threadIdx.x，或者至少在16个线程的组中是恒定的)，就不会出现库冲突。

当您沿阵列进一步访问矩阵时，您仍将以连续的块形式访问矩阵，因此不会出现存储体冲突。只有当您开始访问不相邻的元素时，才会出现银行冲突。

SDK中的减少样本通过从单纯的实现到优化的实现构建，很好地说明了银行冲突，可能值得一看。