银行冲突CUDA共享内存？答案

【问题标题】：Bank conflict CUDA shared memory?银行冲突CUDA共享内存？
【发布时间】：2015-04-07 03:04:27
【问题描述】：

我在 CUDA 内核中遇到了（我相信是）共享内存库冲突。代码本身相当复杂，但我在下面附加的简单示例中复制了它。

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

如果我使用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 个库 (B1, B2, ..) 的 4 个字节，双精度 (D01, D02, ..) 存储为：

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

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

【问题讨论】：

gk110 的 bank 布局取决于 bank 宽度，可配置为 4bytes 或 8bytes。
这是否意味着在 8 字节模式下双精度数 D01..D31 存储在不同的 bank 中，而 D01 和 D32 共享一个 bank？我似乎找不到任何详细信息。
好像是这样的； gpgpu.org/wp/wp-content/uploads/2013/09/08-opti-smem-instr.pdf。我将其添加到我的帖子中

标签： cuda shared-memory bank-conflict

【解决方案1】：

原来我没有正确理解 Keppler 架构。正如 Greg Smith 在上述 cmets 之一中所指出的，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），例如在共享内存上的写入/读取前两行（threadIdx.y=0、threadIdx.y=1）在一个扭曲中处理。这意味着对于threadIdx.y=0 值D00..D15 存储在B0..B15 中，然后有两个双精度值的填充，之后在相同的warp 值D18..D33 中存储在B18..B31+B00..B01 中，这会导致银行冲突B00-B01。如果没有填充 (ng=0)，第一行将写入B00..B15 中的D00..D15，第二行写入B16..B31 中的D16..D31，因此不会发生银行冲突。

对于blockDim.x>=32 的线程块不应该出现问题。比如itot=32、blockDim.x=32、ng=2，第一行存入bankB00..B31，然后两个cell padding，第二行存B02..B31+B00..B01等。

【讨论】：