银行冲突CUDA共享内存？

Question

我运行陷入（我认为是）CUDA 内核中的共享内存库冲突。代码本身相当复杂，但我在下面附带的简单示例中复制了它。

在这种情况下，它被简化为从全局 -> 共享 -> 全局内存的简单副本，大小为 16x16 的二维数组，使用可能在右侧填充的共享内存数组（变量 ng).

如果我用 ng=0 编译代码并用 NVVP 检查 共享内存访问模式 ，它告诉我有 "no issues"。例如ng=2 我在标有 "NVVP warning" 的行得到 "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 个 bank (B1, B2, ..) 的 4 个字节，double (D01, D02, ..) 存储为：

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

没有填充，half warp 写入（as[ijs] = in[ij]）到共享内存 D01 .. D15，D16 .. D31，等等。使用填充（大小为 2），前半个 warp 写入到D01 .. D15，填充到 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;
}

Answer 1

原来是我没有正确理解Keppler架构。正如 Greg Smith 在上面的评论之一中指出的那样，Keppler 可以配置为具有 32 个 8 字节的共享内存组。在这种情况下，使用 cudaDeviceSetSharedMemConfig( cudaSharedMemBankSizeEightByte )，共享内存布局如下所示：

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

现在，对于我的简单示例（使用 itot=16），例如 writing/reading to/from 上的共享内存前两行（threadIdx.y=0、threadIdx.y=1）在一个 warp 中处理。这意味着对于 threadIdx.y=0 值 D00..D15 存储在 B0..B15 中，然后有两个双精度的填充，之后在相同的扭曲值 D18..D33 中存储在 B18..B31+B00..B01，这会导致 B00-B01 上的银行冲突。如果没有填充 (ng=0)，第一行将写入 B00..B15 中的 D00..D15，第二行将写入 D16..D31 中的 B16..B31，因此不会发生内存冲突。

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

银行冲突CUDA共享内存？

Bank conflict CUDA shared memory?

cuda

shared-memory

bank-conflict