Why can't i launch concurrent kernels with the same pointer?

When launching concurrent kernels with the same data pointer, i get an "error reading location" exception.

I can easily fix this problem with cudaDeviceSynchronize(), but i want to understand what goes wrong.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

struct Small {
    int a = 0;
};

struct Big {
    Small* small;
};

__global__ void kernelA(Small* small) {
    // Do nothing
}
__global__ void kernelB(Small* small) {
    // Do nothing
}

void myFunc(Big* big) {
    kernelA << <10, 10 >> > (big->small);
    //cudaDeviceSynchronize();
    kernelB << <10, 10 >> > (big->small);
}


int main()
{
    Big* big;
    cudaMallocManaged(&big, sizeof(Big));
    cudaMallocManaged(&big->small, sizeof(Small));

    myFunc(big);

    return 0;
}

If i uncomment "cudaDeviceSynchronize();" it works, but i don't understand why it wont work either way.

I have read this issue which is similar, but not quite: https://stackoverflow.com/questions/30706945/cuda-kernel-concurrency-with-common-shared-argument

So what is going wrong, why can't two kernels have the same pointer?

Additional information

I am using Visual Studio to compile compute_52,sm_52 for windows

You appear to be on windows, which is a pre-pascal regime for managed memory, regardless of what GPU you are running on.

In this modality, as soon as you launch a kernel, all managed allocations become inaccessible to host code. So big is a location in host memory, and therefore retrieving the pointer value of big is always legal in host code. But what big points to in your case is managed memory. To get that value, it's necessary to dereference big and retrieve a value offset from that pointer. That dereference operation results in accessing a managed allocation, which is illegal in host code (in a pre-pascal regime, after a kernel launch, before any cudaDeviceSynchronize()). The kernel launch process involves host code retrieval of the kernel arguments.

Therefore you are going to hit a seg fault.

And this doesn't have much to do with the "same pointer". It has to do with managed allocations, that must be accessed from host code in pre-pascal regime.

Some fix ideas:

• as you already discovered, put a cudaDeviceSynchronize() in between, which restores host-code accessibility to the managed allocation.

• switch to a demand-paged managed environment (linux, on Pascal or newer GPU).

• capture the pointer in question in a host variable, so that dereferencing of a managed pointer is not required for host code access at the point of the second kernel launch:

  Small *temp = big->small;
  kernelA << <10, 10 >> > (temp);  // or (big->small); here
  kernelB << <10, 10 >> > (temp);  // must use temp here
  

As an aside, not having a cudaDeviceSynchronize() or other synchronizing function (like cudaMemcpy()) after the last kernel launch means that the process may get torn down before your kernels actually execute, but I assume you're just showing a minimal example here. Still I would advise the inclusion of such.