c++ - CUDA 默认流和 CUDA_API_PER_THREAD_DEFAULT_STREAM

Question

The documentation here试图解释如何处理默认流。

给定这样的代码(忽略分配错误):

char *ptr;
char source[1000000];
cudaMalloc((void**)&ptr, 1000000);
cudaMemcpyAsync(ptr, source, 1000000, cudaMemcpyHostToDevice);
myKernel<<<1000, 1000>>>(ptr);

是否存在 myKernel 在 cudaMemcpyAsync 完成复制之前启动的风险？我认为“否”，因为这是文档中所述的“旧版默认流”。

但是，如果我使用 CUDA_API_PER_THREAD_DEFAULT_STREAM 进行编译，会发生什么情况？ “每线程默认流”的文本说:

The per-thread default stream is an implicit stream local to both the thread and the CUcontext, and which does not synchronize with other streams (just like explcitly created streams). The per-thread default stream is not a non-blocking stream and will synchronize with the legacy default stream if both are used in a program.

我认为这也可以，因为 cudaMemcpyAsync 和 myKernel 都有效地使用了 CU_STREAM_PER_THREAD；我说得对吗？

---

_{我问的原因是我在内核中有一个非常奇怪的间歇性 CUDA 错误 77，我只能用 cudaMemcpyAsync 在调用 myKernel 之前未完成来解释>，这意味着我不理解文档。不过，真正的代码过于复杂且过于专有，无法制作 MCVE。}

Answer 1

Is there a risk that myKernel will start before cudaMemcpyAsync finishes copying? I think "No" because this is a "Legacy default stream" as described in the documentation.

不，这不可能发生，因为正如您所注意到的，旧版默认流(流 0)在所有情况下都是阻塞的。

However, if I compile with CUDA_API_PER_THREAD_DEFAULT_STREAM what happens?

几乎没有任何变化。每线程默认流不阻塞，因此其他流和使用其默认流的其他线程可以在上下文中并发运行。然而，这两个操作仍然在同一个流中并且相对于彼此是顺序的。如果 source 是一个不可分页的内存分配，这两个操作之间可能发生重叠的唯一方法是允许传输和内核执行之间发生重叠。否则，由于流的顺序属性和主机源内存施加的限制，它们将按顺序运行。

如果您怀疑操作的意外重叠确实存在问题，您应该能够通过分析来确认这一点。