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

【问题标题】：Device sync during async memcpy in CUDACUDA 中异步 memcpy 期间的设备同步
【发布时间】：2011-10-25 06:43:47
【问题描述】：

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

【问题讨论】：

【解决方案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 线程块完全独立于其他执行线程执行。

【讨论】：

好的，这意味着没有办法实现我想要的。这很糟糕，因为如果我有一个内核，其中第一部分可以在不访问 memcpy 区域的情况下执行，而第二部分需要它，则在执行异步复制之前我无法启动内核。这迫使我在 CPU 上做第一部分。
是否可以将问题划分为依赖和独立的部分？如果是这样，则可以将计算分成两个内核启动——一个依赖于异步传输，一个不依赖于异步传输。
这也是一种可能。我没有考虑任何实际问题，所以我无法回答你的问题。我只是在学习，我得到了这个问题。