uda共有メモリ上書き？

Question

は、私は、チュートリアルで説明したように、作業効率の悪い「ダブルバッファ1」をしようとしています。このtutorial -

に従うことによって、CUDAの並列プレフィックススキャンを記述しようとしています。

これは私が持っているものです。

// double buffered naive. 

// d = number of iterations, N - size, and input. 
__global__ void prefixsum(int* in, int d, int N) 
{ 

     //get the block index 
     int idx = blockIdx.x*blockDim.x + threadIdx.x; 

     // allocate shared memory 
     extern __shared__ int temp_in[], temp_out[]; 

     // copy data to it. 
     temp_in[idx] = in[idx]; 
     temp_out[idx] = 0; 

     // block until all threads copy 

     __syncthreads(); 

     int i = 1; 
     for (i; i<=d; i++) 
     { 
       if (idx < N+1 && idx >= (int)pow(2.0f,(float)i-1)) 
       { 
         // copy new result to temp_out 
         temp_out[idx] += temp_in[idx - (int)pow(2.0f,(float)i-1)] + temp_in[idx]; 
       } 
       else 
       { 
         // if the element is to remain unchanged, copy the same thing 
         temp_out[idx] = temp_in[idx]; 
       } 
       // block until all theads do this 
       __syncthreads(); 
       // copy the result to temp_in for next iteration 
       temp_in[idx] = temp_out[idx]; 
       // wait for all threads to do so 
       __syncthreads(); 
     } 

     //finally copy everything back to global memory 
     in[idx] = temp_in[idx]; 
} 

あなたはこれで間違っているものを指摘することはできますか？私は何が起こるべきかについてのコメントを書いている。

これは、カーネルの呼び出しである - これはグリッドとブロックの割り当てである

prefixsum<<<dimGrid,dimBlock>>>(d_arr, log(SIZE)/log(2), N); 

：

dim3 dimGrid(numBlocks); 
dim3 dimBlock(numThreadsPerBlock); 

問題は、私はより多くのだ任意の入力に対して正しい出力を得ることはありませんということです8要素長い。

Answer 1

私はあなたのコード内の2つの問題

問題1参照ます。extern共有メモリ

AGHを....私はextern __shared__メモリを憎みます。問題は、コンパイラが配列の大きさを知らないことです。結果として、彼らは両方とも同じメモリを指しています！ したがって、あなたの場合：temp_in[5]と​​は、共有メモリ内の同じ単語を参照しています。

あなたが本当にextern __shared__メモリをしたい場合は、手動で、たとえば次のようなもの二番目の配列、オフセットすることができます

size_t size = .... //the size of your array 
extern __shared__ int memory[]; 
int* temp_in=memory; 
int* temp_out=memory+size; 

問題2：共有配列のインデックス

共有メモリがためにプライベートです各ブロック。つまり、あるブロック内のtemp[0]は、別のブロック内のtemp[0]と異なる場合があります。ただし、temp配列がブロック間で共有されているかのように、インデックスをblockIdx.x*blockDim.x + threadIdx.xでインデックスします。

代わりに、おそらくthreadIdx.xでテンポラリ配列のインデックスを作成する必要があります。

もちろん、idxアレイはグローバルであり、正しく索引付けします。