CUDA 中的异步 memcpy 期间的设备同步

Question

假设我想在 CUDA 中执行异步 memcpy 主机到设备，然后立即运行内核。如何在内核中测试异步传输是否已完成？

Answer 1

使用 CUDA“流”对异步复制和内核启动进行排序可确保内核在异步传输完成后执行。下面的代码示例演示了：

#include <stdio.h>

__global__ void kernel(const int *ptr)
{
  printf("Hello, %d\n", *ptr);
}

int main()
{
  int *h_ptr = 0;

  // allocate pinned host memory with cudaMallocHost
  // pinned memory is required for asynchronous copy
  cudaMallocHost(&h_ptr, sizeof(int));

  // look for thirteen in the output
  *h_ptr = 13;

  // allocate device memory
  int *d_ptr = 0;
  cudaMalloc(&d_ptr, sizeof(int));

  // create a stream
  cudaStream_t stream;
  cudaStreamCreate(&stream);

  // sequence the asynchronous copy on our stream
  cudaMemcpyAsync(d_ptr, h_ptr, sizeof(int), cudaMemcpyHostToDevice, stream);

  // sequence the kernel on our stream after the copy
  // the kernel will execute after the copy has completed
  kernel<<<1,1,0,stream>>>(d_ptr);

  // clean up after ourselves
  cudaStreamDestroy(stream);
  cudaFree(d_ptr);
  cudaFreeHost(h_ptr);
}

输出：

$ nvcc -arch=sm_20 async.cu -run
Hello, 13

我不相信有任何受支持的方法可以从内核内部测试是否满足某些异步条件（例如异步传输的完成）。假定 CUDA 线程块完全独立于其他执行线程执行。