调用 cudaMemcpyAsync 时何时可以安全地重用 CPU 缓冲区？

Question

我的项目将有多个线程，每个线程在不同的 cudaStreams 上发布内核执行。其他一些线程将消耗将存储在队列中的结果这里的一些伪代码：

while(true) {
  cudaMemcpyAsync(d_mem, h_mem, some_stream) 
  kernel_launch(some_stream)
  cudaMemcpyAsync(h_queue_results[i++], d_result, some_stream)
}

在第一个 cudaMemcpyAsync 返回后重用 h_mem 是否安全？还是应该使用 N 个主机缓冲区来发出 gpu 计算？

如何知道h_mem什么时候可以重用？我应该使用 cudaevents 进行一些同步吗？

顺便说一句。 h_mem 是主机固定的。如果它是可分页的，我可以直接重用 h_mem 吗？从我在这里读到的内容看来，我可以在 memcpyasync 返回后立即重用，对吗？

异步

对于从可分页主机内存到设备内存的传输，主机内存 立即复制到暂存缓冲区（无设备同步 被执行）。一旦可分页缓冲区具有，该函数将返回 已复制到暂存内存。 DMA传输到最终 目的地可能尚未完成。用于固定主机之间的传输 内存和设备内存，功能完全异步。为了 从设备内存传输到可分页主机内存，函数 只有在复制完成后才会返回。对于所有其他转移， 该函数是完全异步的。如果可分页内存必须首先 暂存到固定内存，这将使用异步处理 工作线程。对于从任何主机内存到任何主机内存的传输， 该函数与主机完全同步。

MemcpyAsynchronousBehavior

谢谢！

Answer 1

为了获得复制/计算重叠，您必须使用固定内存。原因包含在您摘录的段落中。大概你的多流方法的全部原因是复制/计算重叠，所以我认为正确的答案不是切换到使用可分页内存缓冲区。

关于您的问题，假设 h_mem 仅用作您在此处显示的伪代码的源缓冲区（即其中的数据仅参与该 cudaMemcpyAsync 调用），则 h_mem 缓冲区为一旦该流中的 next cuda 操作开始，就不再需要了。因此，如果您的kernel_launch 是实际的kernel<<<...>>>(...)，那么一旦kernel 开始，您就可以确信之前的cudaMemcpyAsync 是完整的。

您可以将 cudaEvents 与 cudaEventSynchronize() 或 cudaStreamWaitEvent() 一起使用，或者您可以直接在流中使用 cudaStreamSynchronize()。例如，如果您在显示的流伪代码中的某处有一个cudaStreamSynchronize() 调用，并且它在cudaMemcpyAsync 调用之后，那么cudaStreamSynchronize() 调用之后的任何代码都保证在cudaMemcpyAsync() 调用之后执行做完了。我引用的所有调用都记录在 usual place 中。