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数组是全局数组，您可以正确地为其指定一个数组。