分析我的CUDA內核的內存訪問合併

Question

我想通過BS_x * BS_Y線程將內容移動到共享內存來讀取（BS_X + 1）*（BS_Y + 1）全局內存位置，並開發了以下代碼。

int i  = threadIdx.x; 
int j  = threadIdx.y; 
int idx  = blockIdx.x*BLOCK_SIZE_X + threadIdx.x; 
int idy  = blockIdx.y*BLOCK_SIZE_Y + threadIdx.y; 

int index1 = j*BLOCK_SIZE_Y+i; 

int i1  = (index1)%(BLOCK_SIZE_X+1); 
int j1  = (index1)/(BLOCK_SIZE_Y+1); 

int i2  = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)%(BLOCK_SIZE_X+1); 
int j2  = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)/(BLOCK_SIZE_Y+1); 

__shared__ double Ezx_h_shared_ext[BLOCK_SIZE_X+1][BLOCK_SIZE_Y+1];  

Ezx_h_shared_ext[i1][j1]=Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1)]; 

if ((i2<(BLOCK_SIZE_X+1))&&(j2<(BLOCK_SIZE_Y+1))) 
Ezx_h_shared_ext[i2][j2]=Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j2)*xdim+(blockIdx.x*BLOCK_SIZE_X+i2)]; 

在我的理解中，合併是順序處理的連續內存讀取的並行等價物。現在如何檢測全局內存訪問是否合併？我指出從（i1，j1）到（i2，j2）有一個索引跳轉。 在此先感謝。

Answer 1

我已經使用手寫合併分析器評估了代碼的內存訪問。評估顯示代碼較少利用合併。下面是你可能會發現有用的凝聚分析：

#include <stdio.h> 
#include <malloc.h> 

typedef struct dim3_t{ 
    int x; 
    int y; 
} dim3; 


// KERNEL LAUNCH PARAMETERS 
#define GRIDDIMX 4 
#define GRIDDIMY 4 
#define BLOCKDIMX 16 
#define BLOCKDIMY 16 


// ARCHITECTURE DEPENDENT 
// number of threads aggregated for coalescing 
#define COALESCINGWIDTH 32 
// number of bytes in one coalesced transaction 
#define CACHEBLOCKSIZE 128 
#define CACHE_BLOCK_ADDR(addr,size) (addr*size)&(~(CACHEBLOCKSIZE-1)) 


int main(){ 
    // fixed dim3 variables 
    // grid and block size 
    dim3 blockDim,gridDim; 
    blockDim.x=BLOCKDIMX; 
    blockDim.y=BLOCKDIMY; 
    gridDim.x=GRIDDIMX; 
    gridDim.y=GRIDDIMY; 

    // counters 
    int unq_accesses=0; 
    int *unq_addr=(int*)malloc(sizeof(int)*COALESCINGWIDTH); 
    int total_unq_accesses=0; 

    // iter over total number of threads 
    // and count the number of memory requests (the coalesced requests) 
    int I, II, III; 
    for(I=0; I<GRIDDIMX*GRIDDIMY; I++){ 
     dim3 blockIdx; 
     blockIdx.x = I%GRIDDIMX; 
     blockIdx.y = I/GRIDDIMX; 
     for(II=0; II<BLOCKDIMX*BLOCKDIMY; II++){ 
      if(II%COALESCINGWIDTH==0){ 
       // new coalescing bunch 
       total_unq_accesses+=unq_accesses; 
       unq_accesses=0; 
      } 
      dim3 threadIdx; 
      threadIdx.x=II%BLOCKDIMX; 
      threadIdx.y=II/BLOCKDIMX; 

      //////////////////////////////////////////////////////// 
      // Change this section to evaluate different accesses // 
      //////////////////////////////////////////////////////// 
      // do your indexing here 
      #define BLOCK_SIZE_X BLOCKDIMX 
      #define BLOCK_SIZE_Y BLOCKDIMY 
      #define xdim 32 
      int i  = threadIdx.x; 
      int j  = threadIdx.y; 
      int idx  = blockIdx.x*BLOCK_SIZE_X + threadIdx.x; 
      int idy  = blockIdx.y*BLOCK_SIZE_Y + threadIdx.y; 

      int index1 = j*BLOCK_SIZE_Y+i; 

      int i1  = (index1)%(BLOCK_SIZE_X+1); 
      int j1  = (index1)/(BLOCK_SIZE_Y+1); 

      int i2  = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)%(BLOCK_SIZE_X+1); 
      int j2  = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)/(BLOCK_SIZE_Y+1); 
      // calculate the accessed location and offset here 
      // change the line "Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1)];" to 
      int addr = (blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1); 
      int size = sizeof(double); 
      ////////////////////////// 
      // End of modifications // 
      ////////////////////////// 

      printf("tid (%d,%d) from blockid (%d,%d) accessing to block %d\n",threadIdx.x,threadIdx.y,blockIdx.x,blockIdx.y,CACHE_BLOCK_ADDR(addr,size)); 
      // check whether it can be merged with existing requests or not 
      short merged=0; 
      for(III=0; III<unq_accesses; III++){ 
       if(CACHE_BLOCK_ADDR(addr,size)==CACHE_BLOCK_ADDR(unq_addr[III],size)){ 
        merged=1; 
        break; 
       } 
      } 
      if(!merged){ 
       // new cache block accessed over this coalescing width 
       unq_addr[unq_accesses]=CACHE_BLOCK_ADDR(addr,size); 
       unq_accesses++; 
      } 
     } 
    } 
    printf("%d threads make %d memory transactions\n",GRIDDIMX*GRIDDIMY*BLOCKDIMX*BLOCKDIMY, total_unq_accesses); 
} 

的代碼將網格的每個線程運行，計算合併請求，度量的存儲器存取合併的數量。

要使用分析器，請將代碼的索引計算部分粘貼到指定區域，然後將內存訪問（數組）分解爲'地址'和'大小'。我已經做到了這一點爲您的代碼，其中indexings是：

int i  = threadIdx.x; 
int j  = threadIdx.y; 
int idx  = blockIdx.x*BLOCK_SIZE_X + threadIdx.x; 
int idy  = blockIdx.y*BLOCK_SIZE_Y + threadIdx.y; 

int index1 = j*BLOCK_SIZE_Y+i; 

int i1  = (index1)%(BLOCK_SIZE_X+1); 
int j1  = (index1)/(BLOCK_SIZE_Y+1); 

int i2  = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)%(BLOCK_SIZE_X+1); 
int j2  = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)/(BLOCK_SIZE_Y+1); 

和內存訪問：

Ezx_h_shared_ext[i1][j1]=Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1)]; 

的分析報告4096個線程訪問4064個高速緩存塊。運行您的實際網格和塊大小的代碼並分析合併行爲。

Answer 2

visual profiler是檢查您的工作的好工具。當你有一段代碼在功能上正確後，再從視覺分析器中運行它。例如，在Linux上，假設你有一個X會話，只需從終端窗口運行nvvp即可。然後您將得到一個嚮導，它將提示您輸入應用程序以及任何命令行參數。

分析器然後會對您的應用程序進行基本運行以收集統計信息。您還可以選擇更高級的統計信息收集（需要附加運行），其中之一將是內存利用率統計信息。它會報告內存利用率的峯值百分比，並將標記警告，認爲它是低利用率，值得您關注。

如果您的應用程序數量超過50％，您的應用程序可能會以您期望的方式運行。如果您的電話號碼較低，則可能錯過了一些融合的細節。它將分別報告內存讀取和內存寫入的統計信息。要獲得100％或接近它，您還需要確保從warp合併的讀取和寫入在128字節邊界上對齊。

這些情況下的常見錯誤是使用基於threadIdx.y的變量作爲最快速變化的索引。在我看來你並沒有犯這個錯誤。例如這是一個常見的錯誤，因爲這通常是我們在C中考慮它的方式。但是線程首先在X軸上組合在一起，因此我們想要使用shared[threadIdx.y][threadIdx.x]或類似的東西。如果您確實犯了這個錯誤，那麼您的代碼仍然可以在功能上正確，但是您將在分析器中獲得較低的百分比利用率數字，例如12％甚至3％。如上所述，爲了獲得50％以上並接近100％，您需要確保不僅所有線程請求都相鄰，而且它們在128B邊界上對齊。由於L1/L2緩存，這些並不是硬性規則，而是指導原則。緩存可以在一定程度上緩解一些錯誤。

Answer 3

隨着GPU的發展，獲取合併訪問的要求變得不那麼嚴格。您對早期GPU架構的合併訪問描述比較新近的描述更準確。特別是Fermi（計算能力2.0）顯着放寬了需求。在費米以後，連續訪問內存位置並不重要。相反，焦點轉移到儘可能少的內存交易來訪問內存。在Fermi上，全局內存事務是128字節寬。所以，當一個warp中的32個線程碰到一條執行加載或存儲的指令時，將會安排128個字節的transaction來處理warp中的所有線程。然後性能取決於需要多少事務。如果所有線程都訪問128字節區域內的值，則需要一個事務處理。如果所有線程訪問不同的128字節區域的值，則需要32個事務。這將是服務經線中單個指令請求的最壞情況。

您可以使用CUDA分析器之一來確定服務請求需要多少事務的平均值。該數字應儘可能接近1。數字越高意味着您應該看看是否有機會優化內核中的內存訪問。