銀行衝突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; 
}

來源

2015-02-06 Bart

爲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第二行，等等

來源

2015-02-09 09:23:08 Bart

銀行衝突CUDA共享內存？

回答

相關問題