cuda общая память перезаписывается?

Question

Я пытаюсь написать параллельный префикс сканирования на CUDA следуя tutorial -

Я пытаюсь рабочую неэффективные «двойной буферизации один», как объяснено в руководстве.

Это то, что у меня есть:

// 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: ехЬегп разделяемой памяти

Agh .... Я ненавижу extern __shared__ память. Проблема в том, что компилятор не знает, насколько велики массивы. В результате они оба указывают на одну и ту же память! Итак, в вашем случае: temp_in[5] и temp_out[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, как если бы массивы temp были разделены между блоками.

Вместо этого, скорее всего, вы должны указывать свои временные массивы только на threadIdx.x.

Конечно, массив idx является глобальным, и вы правильно его индексируете.