银行冲突CUDA共享内存？

问题描述：

我在CUDA内核中遇到（我认为是）共享内存bank冲突。代码本身相当复杂，但我在下面的简单示例中转载了它。银行冲突CUDA共享内存？

在这种情况下，使用可能在右侧填充的共享存储器阵列（变量ng）将其简化为来自全局 - >共享 - >全局存储器的大小为16x16的二维数组的简单副本，。

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

EDIT如下面格雷格史密斯提到，开普勒有8个字节的32家银行宽（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

没有填充，半翘曲写（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银行布局是依赖于银行宽度可配置为4字节或8字节。 – 2015-02-07 01:33:12

这是否意味着在8字节模式下双倍存储'D01..D31'存储在不同的存储区中，并且'D01'和'D32'共享一个银行？我似乎无法找到任何详细的信息。 – Bart 2015-02-07 10:27:23

似乎是这样; http://gpgpu.org/wp/wp-content/uploads/2013/09/08-opti-smem-instr.pdf。我将其添加到我的帖子中 – Bart 2015-02-07 11:12:51

答

事实证明，我没有正确理解开普勒架构。正如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），写入/从例如共享存储器读出到/前两行（threadIdx.y=0,threadIdx.y=1）在一个warp内处理。这意味着对于threadIdx.y=0值D00..D15存储在B0..B15中，则存在两个双打的填充，之后在相同的warp值D18..D33内存储B18..B31+B00..B01，这导致B00-B01上的银行冲突。如果没有填充（ng=0），则第一行将被写入D00..D15的B00..B15,D16..D31的第二行B16..B31，因此不会发生银行冲突。

对于blockDim.x>=32的线程块应该不会发生问题。例如，对于itot=32，blockDim.x=32，ng=2，第一行被存储在银行B00..B31，然后两个小区的填充，在B02..B31+B00..B01第二行，等等

银行冲突CUDA共享内存？

相关推荐