熱方程式の行列 - 不正なアドレスエラー

Question

公式ガイドのshared memory exampleを参照してthis questionに続いて、私はちょうど私が

を作ったこの悪い描かれた画像のようである熱方程式の行列を構築しようとしていますここで

は、最小限の例では、コードは大きなN（32より大きい）を合わせて配置されている

#define N 32 
#define BLOCK_SIZE 16 
#define NUM_BLOCKS ((N + BLOCK_SIZE - 1)/ BLOCK_SIZE) 

__global__ void heat_matrix(int* A) 
{ 
    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    __shared__ int temp_sm_A[N*N]; 
    int* temp_A = &temp_sm_A[0]; memset(temp_A, 0, N*N*sizeof(int)); 

    if (tid < N) //(*) 
    { 
     #pragma unroll 
     for (unsigned int m = 0; m < NUM_BLOCKS; ++m) 
     {   
      #pragma unroll 
      for (unsigned int e = 0; e < BLOCK_SIZE ; ++e) 
      { 
       if ((tid == 0 && e == 0) || (tid == (N-1) && e == (BLOCK_SIZE-1))) 
       { 
        temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2; 
        temp_A[tid + (e + BLOCK_SIZE * m) * N + (tid==0 ? 1 : -1)] = 1; 
       } 
       if (tid == e) 
       { 
        temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1; 
        //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N -1)); 
        temp_A[tid + (e + BLOCK_SIZE * m) * N] = -2; 
        //printf("temp_A[%d] = -2;\n", (tid + (e + BLOCK_SIZE * m) * N)); 
        temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1; 
        //printf("temp_A[%d] = 1;\n", (tid + (e + BLOCK_SIZE * m) * N +1)); 
       } 
      } 
     } 
     __syncthreads(); //(**) 
     memcpy(A, temp_A, N*N*sizeof(int)); 
    } 
} 
int main(){ 
    int* h_A = (int*)malloc(N*N*sizeof(int)); memset(h_A, 0, N*N*sizeof(int)); 
    int* d_A; 
    checkCudaErrors(cudaMalloc((void**)&d_A, N*N*sizeof(int))); 
    checkCudaErrors(cudaMemcpy(d_A, h_A, N*N*sizeof(int), cudaMemcpyHostToDevice)); 
    dim3 dim_grid((N/2 + BLOCK_SIZE -1)/ BLOCK_SIZE); 
    dim3 dim_block(BLOCK_SIZE); 

    heat_matrix <<< dim_grid, dim_block >>> (d_A); 
    checkCudaErrors(cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)); 
... 
} 

、私がこれまで行ってきたものです。私はブロック分割を利用しました。 nvcc利回り

CUDA error at matrix.cu:102 code=77(cudaErrorIllegalAddress) "cudaMemcpy(h_A, d_A, N*N*sizeof(int), cudaMemcpyDeviceToHost)" 

そしてcuda-memcheckを実行するとき、私はコードで間違っていたところ、私が見ることができない唯一のエラー（実際にそこに別のがあるが、それはcudasuccess=checkCudaErrors(cudaDeviceReset()); ...から来ている）

========= CUDA-MEMCHECK 
========= Invalid __shared__ write of size 4 
=========  at 0x00000cd0 in heat_matrix(int*) 
=========  by thread (0,0,0) in block (0,0,0) 
=========  Address 0xfffffffc is out of bounds 
... 

を提供します。最初のブロックのスレッド0はどのように不正なアクセスを引き起こすことができますか？それに対処する特定のifケースもあり、エラーが発生したコード行は報告されていません。

また、私のコードには、ifのすべてを処理するよりも効率的な方法がありますか？確かに存在しますが、ケースを2番目のforに分割するより良い並行表現は見つかりませんでした。

---

私にとっては、(*)は不要です。代わりに(**)が必要です。他のGPU関数呼び出しに従いたいと思っています。私は正しい？

Answer 1

1. チェックアウトこの行：最初の反復の間ゼロに等しいtidとスレッドの

     temp_A[tid + (e + BLOCK_SIZE * m) * N - 1] = 1; 
   

   をtid + (e + BLOCK_SIZE * m) * N - 1 -1のインデックスを評価します。これはまさにcuda-memcheckの出力が（アドレスがアンダーフローのために折り返されていると）不平を言っているものです。

2. 同様のアウトオブバウンドアクセスライン

     temp_A[tid + (e + BLOCK_SIZE * m) * N + 1] = 1; 
   

   tid

   、e、及びm全てについて後で発生するが、その最大値をとります。

3. 複数のスレッドが同じメモリ位置に書き込みます。各スレッドは、内部ループ反復ごとに1つの配列要素に書き込む必要があります。隣接する要素は既にスレッドでカバーされているため、書き込む必要はありません。

4. memset()の初期化とメインループ内のストアとの競合状態があります。 memset()の後ろにsyncthreads()を置きます。

5. memset()とmemcpy()を呼び出すと、各スレッドはフルセット/コピーを行い、操作は1回ではなくN回になります。
   これを処理する一般的な方法は、明示的に操作を書き出し、ブロックのスレッド間で作業を分割することです。
   

6. 共有メモリに最初にマトリックスを作成し、それを後でグローバルメモリにコピーすると、メリットはありません。グローバルメモリ内のAに直接書き込むと、memset(),memcpy()およびsyncthreads()の必要はありません。

   ブロックサイズが16スレッドの場合、スレッドブロックは32スレッド（ワープ）単位で割り当てられるため、リソースの半分は使用されません。

Thread Hierarchyについては、「CUDA Cプログラミングガイド」のセクションをもう一度読んでみてください。

Answer 2

カーネルでは、temp_Aは共有メモリ配列の先頭へのポインタです。考察：

N = 32;

BLOCK_SIZE = 16;

m（0,1）;

E（0、BLOCK_SIZE）

が容易1024要素長配列の境界の外に出ることができるようtemp_A[tid + (e + BLOCK_SIZE * m) * N]アクセス。