Получение неправильного вывода с CUDA при использовании более одного блока

Question

Я пытаюсь разработать реализацию БПФ в CUDA с помощью visual studio 2010, до сих пор я работал на нем до 1024 точек внутри одного блока. Проблема в том, что всякий раз, когда я использую более одного блока, результаты для блока 1 будут в порядке, а остальные вернут неправильное значение (они не кажутся случайными, они не меняются в нескольких прогонах). Вот мое ядро ​​

__device__ void FFT(int idxS,int bfsize, Complex* data1, Complex* data0, int k, int N){ 
     Complex alpha; 
     if((idxS % bfsize) < (bfsize/2)){ 
      data1[idxS] = ComplexAdd(data0[idxS],data0[idxS+bfsize/2]); 
     } 
     else 
     { 
      float angle = -PI*2*((idxS*(1<<k)%(bfsize/2)))/N; 
      alpha.x = cos(angle); 
      alpha.y= sin(angle); 
      Complex v0; 
      v0 = ComplexAdd(data0[idxS-bfsize/2] ,ComplexScale(data0[idxS],-1)); 
      data1[idxS] = ComplexMul(v0, alpha); 
     } 
     } 

__device__ void Ordenador(int r, int idxS ,Complex* data1, Complex* data0){ 
    int p = 0; 
    for(int k = 0;k < r;k++) 
     { 
      if(idxS & (1<<k)) 
      p+=1<<(r - k - 1); 
     } 
    data1[idxS] = data0[p]; 
    __syncthreads(); 
} 


__global__ void GPU_FFT(int N, int r, Complex* data0, Complex* data1, int k) { 
    int idxS = threadIdx.x+ blockIdx.x * blockDim.x; 
     __syncthreads; 
     int bfsize = 1<<(r - k); 
     FFT(idxS, bfsize, data1, data0, k, N); 
     data0[idxS] = data1[idxS]; 
    } 
int prepFFT(float *Entrada, Complex* saida, int N){ 
    if(ceilf(log2((float)N)) == log2((float)N)){ 
     for (int i=0; i<N; i++){ 
      saida[i].x = Entrada[i]; 
      saida[i].y = 0; 
     } 
     Complex *d_saida; 
     int m = (int)log2((float)N); 
     Complex *data1 = new Complex[N]; 
     Complex *data1_d; 
     if (N<1024){ 
     HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     const dim3 numThreads (N,1,1); 
     const dim3 numBlocks(1,1,1); 
      for(int k = 0 ;k < m ; k++) 
    { 
     GPU_FFT<<<numBlocks,numThreads, N*2>>>(N, m, d_saida, data1_d, k); 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
    } 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
     HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost)); 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
     } 
     else{ 
     HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     const dim3 numThreads (1024,1,1); 
     const dim3 numBlocks(N/1024 +1,1,1); 
      for(int k = 0;k < m;k++) 
    { 
     GPU_FFT<<<numBlocks,numThreads, N*2>>>(N, m, d_saida, data1_d, k); 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
    } 
     HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost)); 
     HANDLE_ERROR (cudaDeviceSynchronize());  
     cudaFree(data1_d); 
     cudaFree(d_saida); 
     delete data1; 

     } 
     return 1; 
    } 
    else 
     return 0; 
} 

Я попытался использовать общую память, однако она вернет все 0s, и я понял, что CUDA не копирует из глобального в общий (NSight скажет мне, что значение для этой позиции памяти было ????). Этот код должен быть просто доказательством концепции на данный момент, не нужно оптимизировать, просто верните правильные значения. Если вам нужен полный код, я его предоставил. Я искал решение для этого уже более месяца, это мой отчаянный звонок.

Спасибо, Джон

------- -------- Update

Я изменил код для отладки запуска 2 потока в каждом из 2-х блоков.

int prepFFT(float *Entrada, Complex* saida, int N){ 
    if(ceilf(log2((float)N)) == log2((float)N)){ 
     for (int i=0; i<N; i++){ 
      saida[i].x = Entrada[i]; 
      saida[i].y = 0; 
     } 
     Complex *d_saida; 
     int m = (int)log2((float)N); 

     Complex *data1 = new Complex[N]; 
     Complex *data1_d; 

     if (N<1024){ 
     HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     const dim3 numThreads (2,1,1); 
     const dim3 numBlocks(2,1,1); 
      for(int k = 0 ;k < m ; k++) 
    { 
     GPU_FFT<<<numBlocks,numThreads, N*2>>>(N, m, d_saida, data1_d, k); 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
    } 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
     HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost)); 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
     } 
     else{ 
     HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N)); 
     HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice)); 
     const dim3 numThreads (1024,1,1); 
     const dim3 numBlocks(N/1024 +1,1,1); 
      for(int k = 0;k < m;k++) 
    { 
     GPU_FFT<<<numBlocks,numThreads, N*2>>>(N, m, d_saida, data1_d, k); 
     HANDLE_ERROR (cudaDeviceSynchronize()); 
    } 
     HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost)); 
     HANDLE_ERROR (cudaDeviceSynchronize());  
     cudaFree(data1_d); 
     cudaFree(d_saida); 
     delete data1; 

     } 
     return 1; 
    } 
    else 
     return 0; 

} 

--------------------- Редактировать 2 -------------------- -

Что действительно странно, так это то, что при использовании memcheck (в любом режиме) программа возвращает правильные результаты.

---- Final Edit ---------------

Я обнаружил, что проблема была в этом кусочке кода

FFT(idxS, bfsize, data1, data0, k, N); 
data0[idxS] = data1[idxS]; 

Я нашел что разделение последней строки в новой функции и вызов ее с процессором дали правильные результаты для меня. Спасибо за помощь! С наилучшими пожеланиями!

Answer 1

Прежде всего, вы должны проверять ваши основные функции ядра __global__ void GPU_FFT для выпуска

Просто измените его следующим образом:

__global__ void GPU_FFT(int N, int r, Complex* data0, Complex* data1, int k) { 
    int idxS = threadIdx.x+ blockIdx.x * blockDim.x; 
     int bfsize = 1<<(r - k); 
     //FFT(idxS, bfsize, data1, data0, k, N); 
     //data0[idxS] = data1[idxS]; 
     if (idxS <= N) data0[idxS] = idxS; 
    } 

Что происходит во втором блоке в настоящее время?

Если это нормально раскомментируйте //FFT(idxS, bfsize, data1, data0, k, N);

и изменить последнюю строку:

if (idxS <= N) data0[idxS] = data1[idxS];

Что происходит сейчас? Еще старая ошибка?

p.s. и вам не нужно __syncthreads; сразу после получения ваших индексов потоков

upd.

if((idxS % bfsize) < (bfsize/2)){ 
__syncthreads; 
...}