共享内存库冲突、多播和广播性能 Cuda

Question

我了解，对于 Kepler 设备 (cc 3.0) 及更高版本，共享内存冲突仅在来自同一 warp 的线程访问同一 bank 中的不同单词时才会发生。如果所有线程访问同一个单词（广播）或某些线程访问同一个单词（多播），则不会发生冲突。

在以下代码中：

__shared__ float3 nodeCoefficient[sideX * sideY * sideZ];

...

for (unsigned int zIdx = 0; zIdx < 4; zIdx++) {
    for (unsigned int yIdx = 0; yIdx < 4; yIdx++) {
        for (unsigned int xIdx = 0; xIdx < 4; ++xIdx) {
            int indexXYZ = ((threadidx.z/5.5 + zIdx) * sideY + (threadidx.y/5.5+ yIdx)) * sideX + (threadidx.x/5.5 + xIdx);
            displace += nodeCoefficient[indexXYZ] * (bValues[xIdx].x * bValues[yIdx].y);
        }
    }
}

共享内存访问中存在多播

现在，如果我们将 indexXYZ 更改为：

indexXYZ = (( zIdx) * sideY + ( yIdx)) * sideX + ( xIdx);

我们已经广播了。

最后如果我们把 indexXYZ 改成：

int indexXYZ = ((threadidx.z + zIdx) * sideY + (threadidx.y+ yIdx)) * sideX + (threadidx.x + xIdx);

我们有一个线性访问模式。

上述包括故意银行冲突版本的性能比较如下在gtx750m上：

1.组播：18 毫秒 2.广播：9毫秒 3.线性：5.5ms 4.银行冲突：90ms

我希望无银行冲突的代码会表现得类似。为什么广播、多播和线性访问之间存在差异？

干杯， T

（问题随后被编辑，因为原始版本被标记为过于宽泛）

Answer 1

如果我们有这样的共享内存定义：

__shared__ int sdata[BLOCK_SIZE*2];

以下代码行将生成代表here（中间列）的银行访问模式，没有银行冲突：

int a = sdata[2*(threadIdx.x%16)];

以下代码行将生成相同的银行访问模式（尽管位置不同），但存在 2 路银行冲突：

int a = sdata[2*threadIdx.x];