Array being overwritten in Dynamic Parallelism


#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <curand_kernel.h>

// nvcc -rdc=true -lcudadevrt -o at at.cu

__global__ void childKernel(int *arr, int size, int depth)
{
    if (depth == 10) return;

    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    for(int i = 0; i < size; i++)
    {
        if (arr[i] != i)
        {
            printf("tid:%d\t Error at index %d: read %d || depth: %d\n", tid, i, arr[i], depth);
        }
    }

    curandState state;
    curand_init(depth, tid, size, &state);

    int size1 = curand(&state) % 100;
    int *arr1 = (int *)malloc(size1 * sizeof(int));

    for (int i = 0; i < size; i++)
    {
        arr1[i] = i;
    }

    childKernel<<<1, 32>>>(arr1, size1, depth + 1);


    
}

__global__ void generateRandomInts(int* array, int size, unsigned long long seed) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    curandState state;
    curand_init(seed, tid, 0, &state);

    for (int i = tid; i < size; i += blockDim.x * gridDim.x) {
        array[i] = curand(&state); // Generate random integer
    }
}

__global__ void parentKernel()
{
    curandState state;
    curand_init(0, 0, 0, &state);

    int size = curand(&state) % 100000;
    int *arr = (int *)malloc(size * sizeof(int));

    for (int i = 0; i < size; i++)
    {
        arr[i] = i;
    }

    childKernel<<<1, 32>>>(arr, size, 0);


}

int main()
{

    // each kernel creates a sorted array with variable size
    // then launches a child kernel to read through it and then does the same thing again
    parentKernel<<<1,32>>>();


    return 0;
}

===


> nvcc -rdc=true -lcudadevrt -o at at.cu
> .\at.exe

Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 0: read 4928 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 1: read 4929 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 2: read 4930 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 3: read 4931 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 4: read 4932 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 5: read 4933 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 6: read 4934 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 7: read 4935 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 8: read 4936 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 9: read 4937 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 10: read 4938 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1
Error at index 11: read 4939 || depth: 1

My code has a parent kernel create a sorted array and has the child kernels read and v

My code generates the following results and I am not sure why the array of one thread is being overwritten by another. Is there a name to describe this problem and also how do I go around fixing this issue?

The first step would be to use proper CUDA error checking, throughout (that is: including in device code. See here).

  1. It looks like you are launching 32 threads in the parent kernel. each of those threads launches a child kernel with 32 threads, and then each child kernel launches another child kernel with 32 threads to a depth of 10. At depth 1 you have 1024 kernel launches, at depth 2 you have 32k kernel launches, and so on. You cannot launch that many child kernels, and this is documented in the dynamic parallelism section of the CUDA programming guide.

  2. In kernel new and malloc have various limits. With such a huge number of kernel launches and threads, you are going to have a huge number of malloc requests. Eventually (i.e. very soon) you will run out of heap space. This is also documented in the CUDA programming guide. The canonical method to test for problems here is to check if the returned pointer is NULL.