CUDAおそらく '不十分な'共有メモリを使用した不正なメモリアクセス

Question

私は基本的な削減によってベクトル蓄積を行うことができる単純なCUDAカーネルを持っています。私はそれを複数のブロックに分割してより大きなデータを扱うことができるように拡大しています。しかし、私は、適切な量の共有メモリをカーネルで使用するために割り当てることは、不正なメモリアクセスでは失敗しています。この制限を増やすと消えてしまいますが、その理由を知りたいのです。

コアカーネル： は、ここで私が話していたコードである

__global__ static 
    void vec_add(int *buffer, 
       int numElem, // The actual number of elements 
       int numIntermediates) // The next power of two of numElem 
    { 
     extern __shared__ unsigned int interim[]; 

     int index = blockDim.x * blockIdx.x + threadIdx.x; 

     // Copy global intermediate values into shared memory. 
     interim[threadIdx.x] = 
      (index < numElem) ? buffer[index] : 0; 

     __syncthreads(); 

     // numIntermediates2 *must* be a power of two! 
     for (unsigned int s = numIntermediates/2; s > 0; s >>= 1) { 
      if (threadIdx.x < s) { 
       interim[threadIdx.x] += interim[threadIdx.x + s]; 
      } 
      __syncthreads(); 
     } 

     if (threadIdx.x == 0) { 
      buffer[blockIdx.x] = interim[0]; 
     } 
    } 

そして、これは、発信者である：

void accumulate (int* buffer, int numElem) 
{ 
    unsigned int numReductionThreads = 
     nextPowerOfTwo(numElem); // A routine to return the next higher power of 2. 

    const unsigned int maxThreadsPerBlock = 1024; // deviceProp.maxThreadsPerBlock 

    unsigned int numThreadsPerBlock, numReductionBlocks, reductionBlockSharedDataSize; 

    while (numReductionThreads > 1) { 

     numThreadsPerBlock = numReductionThreads < maxThreadsPerBlock ?   
      numReductionThreads : maxThreadsPerBlock; 

     numReductionBlocks = (numReductionThreads + numThreadsPerBlock - 1)/numThreadsPerBlock; 

     reductionBlockSharedDataSize = numThreadsPerBlock * sizeof(unsigned int); 

     vec_add <<< numReductionBlocks, numThreadsPerBlock, reductionBlockSharedDataSize >>> 
      (buffer, numElem, numReductionThreads); 

     numReductionThreads = nextPowerOfTwo(numReductionBlocks); 
    } 

} 

私は1152個の要素のサンプルセットでこのコードを試してみました私のGPUは次の設定で構成されています： タイプ：Quadro 600 MaxThreadsPerBlock：1024 MaxSharedMemory：48KB

OUTPUT：

Loop 1: numElem = 1152, numReductionThreads = 2048, numReductionBlocks = 2, numThreadsPerBlock = 1024, reductionBlockSharedDataSize = 4096 
Loop 2: numElem = 1152, numReductionThreads = 2, numReductionBlocks = 1, numThreadsPerBlock = 2, reductionBlockSharedDataSize = 8 
CUDA Error 77: an illegal memory access was encountered 

私の「暫定」共有メモリが不正なメモリアクセスを引き起こしたことを疑っては、私は、任意に次の行に2回で共有メモリを増加：

reductionBlockSharedDataSize = 2 * numThreadsPerBlock * sizeof(unsigned int); 

そして、私のカーネルが起動しましたうまく動作します！

私が理解していないのは、なぜこの問題を解決するためにこの余分な共有メモリを提供しなければならないのかということです。

このマジックナンバーを確認するためのさらなる実験として、6912ポイントのデータセットを使ってコードを実行しました。今回は、2倍または4倍でさえ私を助けませんでした。

Loop 1: numElem = 6912, numReductionThreads = 8192, numReductionBlocks = 8, numThreadsPerBlock = 1024, reductionBlockSharedDataSize = 16384 

Loop 2: numElem = 6912, numReductionThreads = 8, numReductionBlocks = 1, numThreadsPerBlock = 8, reductionBlockSharedDataSize = 128 
CUDA Error 77: an illegal memory access was encountered 

しかし、私は共有メモリのサイズを8倍に増やしても問題はなくなりました。

私は間もなく48KBの共有メモリ制限を使い果たしてしまいますので、大規模なデータセットでもこのスケーリングファクタを任意に選ぶことはできません。だから私は自分の問題を修正する正当な方法を知りたい。

Answer 1

インデックス外のアクセスを指摘してくれた@havogtに感謝します。 問題は、vec_addメソッドのnumIntermediatesとして誤った引数を使用していたことでした。その意図は、カーネルが常にスレッド数と同じ数のデータポイントで動作することで、これは常に1024であったはずです。 私は、引数としてnumThreadsPerBlockを使用してそれを修正：

vec_add <<< numReductionBlocks, numThreadsPerBlock, reductionBlockSharedDataSize >>> 
     (buffer, numElem, numThreadsPerBlock);