Wrong results of a CUDA dynamic parallelism code

Question

I recently bumped in the problem illustrated at Uncorrectable ECC error. Shortly speaking, from time to time I receive an Uncorrectable ECC error and my dynamic parallelism code generates uncorrect results. The most probable hypothesis of the uncorrectable ECC error is a corrupted driver stack, which has also been indirectly confirmed by the experience of another user (see the above post). I would now like to face the second issue, i.e., the algorithmic one. To this end, I'm dealing with the reproducer reported below which, since the original code generating uncorrect results uses dynamic parallelism, uses this CUDA feature too.

I do not see any evindent issue with this code. I think that the synchronization regarding the child kernel launch should be ok: the first __syncthreads() should not be necessary and the cudaDeviceSynchronize() should ensure that all the memory writes of the child kernel are accomplished before the printf.

My question is: is this code wrong or the wrong results are due to a non-programming issue?

My configuration: CUDA 5.0, Windows 7, 4-GPU system equipped with Kepler K20c, driver 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

I'm pretty sure you're exceeding the launch pending limit. It's nearly impossible to tell with your code as-is, but I've modified it and added error checking on the child kernel launch.

When I do that, I get launch errors, signified by a printout of !. Skipping the launch error cases, all of my in-kernel checking of P1[m] vs. m passes (I get no * printout at all.)

#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;
}

Feel free to add further decoding of the err variable in the parent kernel to convince yourself that you are exceeding the launch pending limit. As another test, you can set M to 2048 instead of 19000 in your host code, and all the ! printouts go away. (launch pending limit default == 2048)

As I've stated in the comments, I think the uncorrectable ECC error is a separate issue, and I suggest trying the driver 321.01 that I linked in the comments.