使用CUDA共享內存來改善全球訪問模式

Question

我有以下內核得到一堆向量幅度：

__global__ void norm_v1(double *in, double *out, int n) 
{ 
    const uint i = blockIdx.x * blockDim.x + threadIdx.x; 

    if (i < n) 
    { 
     double x = in[3*i], y = in[3*i+1], z = in[3*i+2]; 
     out[i] = sqrt(x*x + y*y + z*z); 
    } 
} 

但是由於in包裝作爲[x0,y0,z0,...,xn,yn,zn]它表明與探查表現不佳32％的全球負載效率。將數據重新包裝爲[x0, x1, ..., xn, y0, y1, ..., yn, z0, z1, ..., zn]可以大大提高工作效率（其中x,y和z的偏移量會相應改變）。運行時間縮短，效率高達100％。

但是，這種包裝對我的應用來說根本不實用。因此，我希望調查共享內存的使用。我的想法是，塊中的每個線程都可以從全局內存中複製三個值（分開爲blockDim.x） - 從而產生合併訪問。在最大blockDim.x = 256假設我想出了：

#define BLOCKDIM 256 

__global__ void norm_v2(double *in, double *out, int n) 
{ 
    __shared__ double invec[3*BLOCKDIM]; 

    const uint i = blockIdx.x * blockDim.x + threadIdx.x; 

    invec[0*BLOCKDIM + threadIdx.x] = in[0*BLOCKDIM+i]; 
    invec[1*BLOCKDIM + threadIdx.x] = in[1*BLOCKDIM+i]; 
    invec[2*BLOCKDIM + threadIdx.x] = in[2*BLOCKDIM+i]; 
    __syncthreads(); 

    if (i < n) 
    { 
     double x = invec[3*threadIdx.x]; 
     double y = invec[3*threadIdx.x+1]; 
     double z = invec[3*threadIdx.x+2]; 

     out[i] = sqrt(x*x + y*y + z*z); 
    } 
} 

然而，這顯然是不足的時候n % blockDim.x != 0，需要預先知道的最大blockDim，當與n = 1024測試out[i > 255]產生不正確的結果。我應該如何最好地解決這個問題？

Answer 1

我認爲這是可以解決的問題out[i > 255]：

__shared__ double shIn[3*BLOCKDIM]; 

const uint blockStart = blockIdx.x * blockDim.x; 

invec[0*blockDim.x+threadIdx.x] = in[ blockStart*3 + 0*blockDim.x + threadIdx.x]; 
invec[1*blockDim.x+threadIdx.x] = in[ blockStart*3 + 1*blockDim.x + threadIdx.x]; 
invec[2*blockDim.x+threadIdx.x] = in[ blockStart*3 + 2*blockDim.x + threadIdx.x]; 
__syncthreads(); 

double x = shIn[3*threadIdx.x]; 
double y = shIn[3*threadIdx.x+1]; 
double z = shIn[3*threadIdx.x+2]; 

out[blockStart+threadIdx.x] = sqrt(x*x + y*y + z*z); 

至於n % blockDim.x != 0我建議填充所述輸入/輸出陣列用0至符合要求。

如果你不喜歡的BLOCKDIM宏 - 探索使用extern __shared__ shArr[]，然後使第三參數到內核配置：

norm_v2<<<gridSize,blockSize,dynShMem>>>(...) 

的dynShMem是動態共享存儲器使用量（以字節爲單位）。這是額外的共享內存池，其大小在運行時指定，其中最初將分配所有extern __shared__變量。

---

您使用的是什麼GPU？費米或開普勒可能幫助您的原代碼與他們的L1緩存。

---

如果你不想觸摸板您in陣列，或者你最終做類似的把戲在其他地方，你可能要考慮實施的設備側memcopy，這樣的事情：

template <typename T> 
void memCopy(T* destination, T* source, size_t numElements) { 
    //assuming sizeof(T) is a multiple of sizeof(int) 
    //assuming one-dimentional kernel (only threadIdx.x and blockDim.x matters) 
    size_t totalSize = numElements*sizeof(T)/sizeof(int); 
    int* intDest = (int*)destination; 
    int* intSrc = (int*)source; 
    for (size_t i = threadIdx.x; i < totalSize; i += blockDim.x) { 
     intDest[i] = intSrc[i]; 
    } 
    __syncthreads(); 
} 

它基本上將任何數組視爲int-s的數組，並將數據從一個位置複製到另一個位置。如果僅使用64位類型，則可能需要將底層的int類型替換爲double -s或long long int。

然後你就可以替換複製行：

memCopy(invec, in+blockStart*3, min(blockDim.x, n-blockStart));