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 编译会发生什么？ “每线程默认流”的文本说：

每个线程的默认流是线程和CUcontext 本地的隐式流，并且不与其他流同步（就像显式创建的流一样）。每个线程的默认流不是非阻塞流，如果在程序中同时使用两者，它将与旧的默认流同步。

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

---

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

Answer 1

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

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

但是，如果我使用 CUDA_API_PER_THREAD_DEFAULT_STREAM 编译会发生什么？

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

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