错误的CUDA动态并行代码结果

Question

我最近遇到了Uncorrectable ECC error所示的问题。简而言之，我不时收到一个不可纠正的ECC错误，我的动态并行代码会产生不正确的结果。不可纠正的ECC错误的最可能的假设是驱动程序堆栈已损坏，这也已被另一用户的经验间接证实（参见上面的帖子）。现在我想面对第二个问题，即算法问题。为此，我正在处理下面报告的重现器，由于生成不正确结果的原始代码使用动态并行性，因此也使用此CUDA特性。

我没有看到这个代码的任何明显的问题。我认为关于子内核启动的同步应该是正确的：第一个__syncthreads()不应该是必需的，并且cudaDeviceSynchronize()应该确保子内核的所有内存写入在printf之前完成。

我的问题是：这段代码是错误的还是错误的结果是由于非编程问题？

我的配置：CUDA 5.0，Windows 7,4 GPU系统，配备Kepler K20c，驱动程序327.23。

#include <stdio.h> 
#include <conio.h> 

#define K 6 
#define BLOCK_SIZE 256 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) { getch(); exit(code); } 
    } 
} 

int iDivUp(int a, int b) { return ((a % b) != 0) ? (a/b + 1) : (a/b); } 

__global__ void child_kernel(double* P1) 
{ 
    int m = threadIdx.x; 

    P1[m] = (double)m; 
} 

__global__ void parent_kernel(double* __restrict__ x, int M) 
{ 
    int i = threadIdx.x + blockDim.x * blockIdx.x; 

    if(i<M) { 

     double* P1 = new double[13]; 

     dim3 dimBlock(2*K+1,1); dim3 dimGrid(1,1); 

     __syncthreads(); 
     child_kernel<<<dimGrid,dimBlock>>>(P1); 
     cudaDeviceSynchronize(); 

     for(int m=0; m<2*K+1; m++) printf("%f %f\n",P1[m],(double)m); 

    } 
} 

int main() { 

    const int M = 19000; 

//gpuErrchk(cudaSetDevice(0)); 

    double* x = (double*)malloc(M*sizeof(double)); 
    for (int i=0; i<M; i++) x[i] = (double)i; 

    double* d_x; gpuErrchk(cudaMalloc((void**)&d_x,M*sizeof(double))); 

    gpuErrchk(cudaMemcpy(d_x,x,M*sizeof(double),cudaMemcpyHostToDevice)); 

    dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(iDivUp(M,BLOCK_SIZE)); 
    parent_kernel<<<dimGrid,dimBlock>>>(d_x,M); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    getch(); 

    return 0; 
} 

Answer 1

我敢肯定，你超出launch pending limit。几乎不可能用你的代码原样告诉，但是我修改了它并在子内核启动时添加了错误检查。

当我这样做，我得到启动错误，打印输出!。跳绳推出错误的情况下，我的所有的P1[m]与m通行证（我没有得到任何*打印在所有。）

#include <stdio.h> 

#define K 6 
#define BLOCK_SIZE 256 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) { exit(code); } 
    } 
} 

int iDivUp(int a, int b) { return ((a % b) != 0) ? (a/b + 1) : (a/b); } 

__global__ void child_kernel(unsigned long long* P1) 
{ 
    int m = threadIdx.x; 

    P1[m] = (unsigned long long)m; 
} 

__global__ void parent_kernel(double* __restrict__ x, int M) 
{ 
    int i = threadIdx.x + blockDim.x * blockIdx.x; 

    if(i<M) { 

     unsigned long long* P1 = new unsigned long long[13]; 

     dim3 dimBlock(2*K+1,1); dim3 dimGrid(1,1); 

     __syncthreads(); 
     child_kernel<<<dimGrid,dimBlock>>>(P1); 
     cudaDeviceSynchronize(); 
     cudaError_t err = cudaGetLastError(); 
     if (err != cudaSuccess) printf("!"); 
     else for(unsigned long long m=0; m<dimBlock.x; m++) if (P1[m] != m) printf("*"); 

    } 
} 

int main() { 

    const int M = 19000; 

//gpuErrchk(cudaSetDevice(0)); 

    double* x = (double*)malloc(M*sizeof(double)); 
    for (int i=0; i<M; i++) x[i] = (double)i; 

    double* d_x; gpuErrchk(cudaMalloc((void**)&d_x,M*sizeof(double))); 

    gpuErrchk(cudaMemcpy(d_x,x,M*sizeof(double),cudaMemcpyHostToDevice)); 

    dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(iDivUp(M,BLOCK_SIZE)); 
    parent_kernel<<<dimGrid,dimBlock>>>(d_x,M); 
    gpuErrchk(cudaPeekAtLastError()); 
    gpuErrchk(cudaDeviceSynchronize()); 

    return 0; 
} 

随意添加err变量的进一步解码父内核内核检查的说服自己，你超出了启动待定限制。作为另一种测试，您可以在您的主机代码中将M设置为2048而不是19000，并且所有!打印输出都将消失。 （启动挂起限制默认== 2048）

正如我在评论中所述，我认为不可纠正的ECC错误是一个单独的问题，我建议尝试我在评论中链接的驱动程序321.01。