CUDA中的共享内存库冲突:内存如何与库对齐

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个,费米),但是连续的32位字存在于连续的存储体中,即地址空间在存储体之间交错.这意味着,如果你总是有一个可以分解为的数组索引x + threadIdx.x(其中x不依赖于threadIdx.x,或者至少在16个线程的组中是不变的),你就不会有银行冲突.

当您沿阵列进一步访问矩阵时,仍然可以在一个连续的块中访问它们,因此您不会发生银行冲突.只有当您开始访问非相邻元素时才会发生银行冲突.

SDK中的简化示例通过从简单的实现构建到优化的实现来很好地说明了银行冲突,可能值得一看.