银行冲突 CUDA 共享内存？

Question

我在CUDA内核中遇到了（我认为是）共享内存冲突的问题。代码本身相当复杂，但我在下面附上了一个简单的示例来重现它。

在这种情况下，它被简化为从全局 -> 共享 -> 全局内存的简单复制，使用可能在右侧填充的共享内存数组（变量ng），大小为16x16的2D数组。

如果我使用ng = 0编译代码，并使用NVVP检查共享内存访问模式，它告诉我“没有问题”。使用例如ng = 2时，在标有“NVVP警告”的行处得到“Shared Store Transactions/Access = 2, Ideal Transactions/Acces = 1”。我不明白为什么（或更具体地说：为什么填充会导致警告）。

编辑：如Greg Smith在下面所提到的，在Kepler上有32个8字节宽的银行(http://gpgpu.org/wp/wp-content/uploads/2013/09/08-opti-smem-instr.pdf，第18页)。但我不知道这会如何改变问题。

如果我理解正确，使用32个4字节的银行(B1, B2, ..)，双精度浮点数(D01, D02, ..)被存储为：

B1   B2   B3   B4   B5    ..   B31
----------------------------------
D01       D02       D03   ..   D15
D16       D17       D18   ..   D31
D32       D33       D34   ..   D47

没有填充的情况下，半个warp会将(as[ijs] = in[ij])写入共享内存D01 .. D15，D16 .. D31等。使用大小为2的填充后，第一个半warp写入D01 .. D15，第二个半warp在填充后写入D18 .. D33，这仍然不应该导致银行冲突吗？

有任何想法可能出了什么问题吗？

简化示例（已测试cuda 6.5.14）：

// Compiled with nvcc -O3 -arch=sm_35 -lineinfo 

__global__ void copy(double * const __restrict__ out, const double * const __restrict__ in, const int ni, const int nj, const int ng)

{
    extern __shared__ double as[];
    const int ij=threadIdx.x + threadIdx.y*blockDim.x;
    const int ijs=threadIdx.x + threadIdx.y*(blockDim.x+ng);

    as[ijs] = in[ij]; // NVVP warning
    __syncthreads();
    out[ij] = as[ijs]; // NVVP warning
}

int main()
{
    const int itot = 16;
    const int jtot = 16;
    const int ng = 2;
    const int ncells = itot * jtot;

    double *in  = new double[ncells];
    double *out = new double[ncells];
    double *tmp = new double[ncells];
    for(int n=0; n<ncells; ++n)
        in[n]  = 0.001 * (std::rand() % 1000) - 0.5;

    double *ind, *outd;
    cudaMalloc((void **)&ind,  ncells*sizeof(double));
    cudaMalloc((void **)&outd, ncells*sizeof(double));
    cudaMemcpy(ind,  in,  ncells*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(outd, out, ncells*sizeof(double), cudaMemcpyHostToDevice);

    dim3 gridGPU (1, 1 , 1);
    dim3 blockGPU(16, 16, 1);

    copy<<<gridGPU, blockGPU, (itot+ng)*jtot*sizeof(double)>>>(outd, ind, itot, jtot, ng);

    cudaMemcpy(tmp, outd, ncells*sizeof(double), cudaMemcpyDeviceToHost);

    return 0;
}

- Bart

---

gk110的银行布局取决于可配置为4字节或8字节的银行宽度。 - Greg Smith

这是否意味着在8字节模式下，双精度数 D01..D31 存储在不同的存储区块中，而 D01 和 D32 共享一个存储区块？我似乎找不到任何详细信息。 - Bart

看起来确实是这样的；http://gpgpu.org/wp/wp-content/uploads/2013/09/08-opti-smem-instr.pdf。我已将其添加到我的帖子中。 - Bart

Answer 1

原来我对Keppler架构的理解有误。就像Greg Smith在上面的评论中指出的那样，Keppler可以配置为具有32个8字节的共享内存Bank。在这种情况下，使用cudaDeviceSetSharedMemConfig（cudaSharedMemBankSizeEightByte），共享内存布局如下：

bank:  B0   B1   B2   B3   B4    ..   B31
       ----------------------------------
index: D00  D01  D02  D03  D04   ..   D31
       D32  D33  D34  D35  D36   ..   D63   

现在，以我的简单示例（使用itot=16）为例，在共享内存上的写入/读取，例如前两行（threadIdx.y=0，threadIdx.y=1），由一个warp处理。这意味着对于threadIdx.y=0的值，D00..D15存储在B0..B15中，然后有两个double的填充，之后在同一warp中，D18..D33的值存储在B18..B31+B00..B01中，这会导致B00-B01上的bank冲突。没有填充（ng=0），第一行被写入B00..B15中的D00..D15，第二行被写入B16..B31中的D16..D31，因此不会发生bank冲突。

对于一个blockDim.x>=32的线程块，问题不应该出现。例如，对于itot=32，blockDim.x=32，ng=2，第一行存储在bank B00..B31中，然后是两个单元格的填充，第二行在B02..B31+B00..B01中，以此类推。

- Bart