How do I know that cudaMemcpyAsync is done reading host memory?

Question

For example... Here's what I see in NVIDIA's docs:

cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 0);
kernel<<<grid, block>>>(a_d);
cpuFunction();

Let's say this is wrapped in a function...

void consume() {
  cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 0);
  kernel<<<grid, block>>>(a_d);
}

What if I also have a function

void produce() {
  // do stuff
  a_h[0] = 1;
  a_h[1] = 3;
  a_h[2] = 5;
  //...
}

If I call:

produce();
consume();
produce(); // problem??

The second produce() will start to change the memory on the host at a_h

How do I know that CUDA isn't still reading the host memory during the asynchronous memory copy routine?

How can I safely write to the host a_h memory without disrupting that asynchronous mem copy?

EDIT---

I know I can call cudaDeviceSynchronize() or cudaStreamSynchronize() but that will also wait for kernel to complete. I would prefer to not wait until kernel is done.

I want to start writing to host a_h as soon as possible, while not waiting for kernel to finish.

Answer 1

If you use a stream for the cudaMemcpyAsync call, you can insert an event into the stream after the asynchronous transfer and then use cudaEventSynchronize to synchronize on that event. This guarantees that the copy has finished, but doesn't rely on the device being idle or the stream being empty.