CUDA 2D陣列nvidia

Question

我對於2D陣列使用cudaMallocPitch和cudaMemcpy2D。即使無法正確輸出，我也不確定是否編碼正確。任何人都可以幫忙嗎？任何一個可以調試我的錯誤？提前致謝。

#include<stdio.h> 
#include<cuda.h> 
#define siz 4*sizeof(int) 
__global__ void addmatrix(int *m1,int *m2,size_t pitch) 
{ 
    int r=threadIdx.x; 
    int *r1=m1+r*pitch; 
    int *r2=m2+r*pitch; 
    int c; 
    for(c=1;c<=4;c++) 
    { 
     r1[c]+=r2[c]; 
    } 
} 
int main() 
{ 
    int i,j; 
    int **m1_c,**m2_c; 
    int *m1_d,*m2_d; 
    size_t pitch; 
    cudaError_t err; 
    m1_c=(int **)malloc(4*sizeof(int *)); 
    for(i=1;i<=4;i++) 
    { 
     m1_c[i]=(int *)malloc(siz); 
    } 
    m2_c=(int **)malloc(4*sizeof(int *)); 
    for(i=1;i<=4;i++) 
    { 
     m2_c[i]=(int *)malloc(siz); 
    } 
    for(i=1;i<=4;i++) 
    { 
     for(j=1;j<=4;j++) 
     { 
      m1_c[i][j]=rand()%10; 
      m2_c[i][j]=rand()%10; 
     } 
    } 
    for(i=1;i<=4;i++) 
    { 
     for(j=1;j<=4;j++) 
     { 
      printf("%d\t",m1_c[i][j]); 
     } 
     printf("\n"); 
    } 
    printf("\n\n"); 
    for(i=1;i<=4;i++) 
    { 
     for(j=1;j<=4;j++) 
     { 
      printf("%d\t",m2_c[i][j]); 
     } 
     printf("\n"); 
    } 
    err=cudaMallocPitch((void **)&m1_d,&pitch,siz,siz); 
    err=cudaMallocPitch((void **)&m2_d,&pitch,siz,siz); 
    err=cudaMemcpy2D(m1_d,pitch,m1_c,siz,siz,4,cudaMemcpyHostToDevice); 
    err=cudaMemcpy2D(m2_d,pitch,m2_c,siz,siz,4,cudaMemcpyHostToDevice); 
    dim3 grid(1); 
    dim3 block(16); 
    addmatrix<<<grid,block>>>(m1_d,m2_d,siz); 
    cudaMemcpy2D(m1_c,siz,m1_d,pitch,siz,4,cudaMemcpyDeviceToHost); 

    for(i=1;i<=4;i++) 
    { 
     for(j=1;j<=4;j++) 
     { 
      printf("%d\t",m1_c[i][j]); 
     } 
     printf("\n"); 
    } 
    err=cudaFree(m1_d); 
    err=cudaFree(m2_d); 
    err=cudaDeviceReset();  
} 

Answer 1

所以這個代碼有幾個問題。在沒有特定的順序：

1. 您通過您在不同的陣列索引從1到4，但這不是。C索引正確從零開始，去到一個小於尺寸。這與CUDA無關。
2. cudaMemcpy2D需要兩個指針（src和dst），它們都是指向內存中線性數組的指針。我意識到這是令人困惑的，因爲2D出現在整個描述中，但兩個指針參數基本上都是相同類型（一個指向內存的指針），並且你傳遞了2種不同類型的指針（一個是指向內存的指針，是指向內存的指針）。從cudaMemcpy2D的定義中清楚，你的用法是不正確的。有很多關於如何使用cudaMemcpy2D的示例回答問題，我建議您搜索並查看其中的一些。請注意，解決此問題可能會導致您從根本上重新考慮如何將數據存儲在主機矩陣中。關於處理多維矩陣有很多問題，如this one - 如果可能的話，您應該將它們弄平。請注意，在您當前的代碼中，使用cudaMemcpy2D這個錯誤會導致您的主機矩陣上的指針數組被破壞，當您嘗試打印結果時會導致seg錯誤。
3. 傳遞給cudaMallocPitch的參數不太正確。對於width和height參數，您將通過siz這是矩陣尺寸的字節。但是您應該只傳遞width參數的字節維數。對於height參數，您應該傳遞行數，即4。對cudaMemcpy2D的調用也有類似的要求，但你已經明白了。
4. 現在我們來看看你的內核。在調用中，您將啓動一個由16個線程組成的網格。由於你的矩陣有16個元素，這似乎是明智的。這意味着線程策略其中每個線程將負責結果的單個元素。但是看看你的內核代碼，你有每個線程計算整行的結果，即4個元素。有兩種方法可以解決這個問題：可以將網格減少爲4個線程而不是16個線程（可能更簡單一些，或許可以從代碼修改的角度考慮），也可以重新編寫內核（消除for循環）每個線程計算一個單個的輸出元素（這可能會並行地完成更多的工作）。
5. 此外，在您的內核中，您正在基於指針算術的索引中使用pitch參數。但請記住，音高在字節中，並且對於指針算術索引，編譯器希望參數位於元素--它根據數據類型爲您進行字節轉換。再次，這是一個真正的C問題，而不是CUDA特有的。你可以通過在內核中使用pitch的地方使用(pitch/sizeof(int))來解決這個問題。
6. 你正在向內核傳遞siz。你應該通過pitch作爲音高參數。 siz實際上是主機數據存儲上的「間距」，但是pitch是設備上存儲的間距。內核在設備存儲上運行，所以它需要正確的音調。
7. 作爲建議，請在所有cuda API調用和內核調用上執行cuda error checking。

下面是一些代碼，解決所有上述問題，在一個時尚或其他：

#include<stdio.h> 
#define siz (4*sizeof(int)) 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

__global__ void addmatrix(int *m1,int *m2,size_t pitch) 
{ 
    int r=threadIdx.x; 
    int *r1=m1+r*(pitch/sizeof(int)); 
    int *r2=m2+r*(pitch/sizeof(int)); 
    int c; 
    for(c=0;c<4;c++) 
    { 
     r1[c]+=r2[c]; 
    } 
} 
int main() 
{ 
    int i,j; 
    int *m1_c,*m2_c; 
    int *m1_d,*m2_d; 
    size_t pitch; 
    cudaError_t err; 
    m1_c=(int *)malloc(16*sizeof(int)); 
    m2_c=(int *)malloc(16*sizeof(int)); 
    for(i=0;i<4;i++) 
    { 
     for(j=0;j<4;j++) 
     { 
      m1_c[(i*4)+j]=rand()%10; 
      m2_c[(i*4)+j]=rand()%10; 
     } 
    } 
    for(i=0;i<4;i++) 
    { 
     for(j=0;j<4;j++) 
     { 
      printf("%d\t",m1_c[(i*4)+j]); 
     } 
     printf("\n"); 
    } 
    printf("\n\n"); 
    for(i=0;i<4;i++) 
    { 
     for(j=0;j<4;j++) 
     { 
      printf("%d\t",m2_c[(i*4)+j]); 
     } 
     printf("\n"); 
    } 
    err=cudaMallocPitch((void **)&m1_d,&pitch,siz,4); 
    cudaCheckErrors("cm1"); 
    err=cudaMallocPitch((void **)&m2_d,&pitch,siz,4); 
    cudaCheckErrors("cm2"); 
    err=cudaMemcpy2D(m1_d,pitch,m1_c,siz,siz,4,cudaMemcpyHostToDevice); 
    cudaCheckErrors("cm3"); 
    err=cudaMemcpy2D(m2_d,pitch,m2_c,siz,siz,4,cudaMemcpyHostToDevice); 
    cudaCheckErrors("cm4"); 
    dim3 grid(1); 
    dim3 block(4); 
    addmatrix<<<grid,block>>>(m1_d,m2_d,pitch); 
    cudaMemcpy2D(m1_c,siz,m1_d,pitch,siz,4,cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cm5"); 

    for(i=0;i<4;i++) 
    { 
     for(j=0;j<4;j++) 
     { 
      printf("%d\t",m1_c[(i*4)+j]); 
     } 
     printf("\n"); 
    } 
    err=cudaFree(m1_d); 
    err=cudaFree(m2_d); 
    err=cudaDeviceReset(); 
}