使用多个块时使用 CUDA 得到错误的输出答案

【问题标题】：getting the wrong output with CUDA when using more than one block使用多个块时使用 CUDA 得到错误的输出
【发布时间】：2013-06-10 03:44:02
【问题描述】：

我正在尝试使用 Visual Studio 2010 在 CUDA 中开发 FFT 的实现，到目前为止，我已经让它在一个块内的 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;
}

我尝试过使用共享内存，但是它会返回全 0，并且我认为 CUDA 没有从全局复制到共享（NSight 会告诉我该内存位置的值是????）。这段代码现在应该只是一个概念证明，不需要优化，只需返回正确的值。如果你们需要完整的代码，我会提供。一个多月以来，我一直在寻找解决方案，这是我绝望的呼唤。

谢谢，约翰

------- 更新--------

出于调试目的，我更改了代码，在 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（在任何模式下）时，程序会返回正确的结果。

----最终编辑---------------

我发现问题出在这段代码中

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

我发现在新函数中分离最后一行并用 CPU 调用它对我来说产生了正确的结果。感谢您的帮助！！最好的问候！

【问题讨论】：

我想您这样做是出于学习目的，但如果不是，请指出有一个 CUDA 库 (cufft) 可以为您执行 FFT。跨度>
是的，我这样做是为了学习，稍后我将使用 cufft 进行比较。感谢您的提醒。
您尝试在失败的情况下使用cuda-memcheck 运行您的代码，看看它是否报告任何访问错误。
也许您应该使用块 2 中的线程 317 作为示例来完成索引，看看您是否可以发现全局内存中数组中的索引问题。
"当我使用 memcheck 时，它会返回正确的结果，但是如果我在没有 memcheck 的情况下运行，我会得到奇怪的结果"。现在这听起来像是一个竞争条件。使用 memcheck 运行可能会影响块的执行顺序。您的代码中可能存在块顺序依赖性。尝试使用 racecheck 工具运行 memcheck。运行cuda-memcheck --help 了解如何操作。

标签： cuda gpu fft

【解决方案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];

现在会发生什么？还是老错误？

附言检索线程索引后，您不需要__syncthreads;

更新。

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

【讨论】：

data0 是一个复杂的，所以我更改为以下内容：global void GPU_FFT(int N, int r, Complex* data0, Complex* data1, int k) { int idxS = threadIdx.x+ blockIdx.x * blockDim.x; if (idxS
第一部分没问题，我无法理解您的最后更改，因为 [idxS] 返回错误
好的，这样做了，我仍然得到旧的错误。真正奇怪的是，如果我使用 memcheck 运行它，我会得到正确的结果（尝试了每个选项，没有错误）
好的，我们确实本地化了问题。尝试简化您的设备 FFT 功能。但首先你可以尝试把 __syncthreads;就在 if((idxS % bfsize)
好的，这是我到目前为止所得到的：没有任何 __syncthreads 我得到了很多错误的输出，其中一个在 if 中少一点，一个在完成 if else 语句之后我从程序中得到了最大的收益，一半的输出是正确的。我将尝试简化函数，我正在考虑将 bfsize 计算发送到 CPU，将其作为参数传递并将 if else 语句也传递给 CPU，这样我一次可以调用一半的代码。