cuda共享內存覆蓋？

Question

我試圖按照此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

我看到兩個問題在你的代碼

問題1：EXTERN共享內存

哎....我恨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]不同。但是，您按照blockIdx.x*blockDim.x + threadIdx.x的索引編寫索引，就好像臨時數組在塊之間共享一樣。

相反，您應該最有可能將您的臨時數組索引爲threadIdx.x。

當然，idx數組是全局數組，您可以正確地爲其指定一個數組。