CUDA - 子网格将在哪个流中？

Posted 2023-03-23

【中文标题】CUDA - 子网格将在哪个流中？

【英文标题】：CUDA - Which stream will child grid be in?

【发布时间】：2022-01-14 03:03:33

【问题描述】：

如果我使用动态并行，子网格将在哪个流中运行？

例如，我有一个叫A的内核，另一个叫B的内核。B是由A启动的。

如果内核 A 在 stream_A 中运行，并且如果我没有为内核 B 指定流 ID，内核将在哪个流中运行？它是默认流，还是会继承 A 正在运行的流？

【参考方案1】：

根据我在文档中看到的内容以及我自己的实验，我观察到通常主机上的流和设备上的流之间没有（排序）关系。

我们可以考虑两种情况：

第一种情况，对于创建的流，我会说这是覆盖explicitly in the documentation。

在第二种情况下，虽然文档提到了设备 NULL 流，但这可能有点不清楚。我们可以编写一个简单的测试用例来为我们解决这个问题：

$ cat cdp.cu
#include <cstdio>
#include <cassert>
#include <unistd.h>
const unsigned long long dt = 10000000000ULL;
__device__ void delay()
        unsigned long long start = clock64();
        while (start+dt > clock64());

__global__ void child()
        delay();

__global__ void parent()
        child<<<1,1>>>();
        cudaError_t err = cudaGetLastError();
        assert(err == cudaSuccess);
        err = cudaDeviceSynchronize();
        assert(err == cudaSuccess);

__global__ void pk()
        printf("hello\n");

int main()
        cudaStream_t s1, s2;
        cudaStreamCreate(&s1);
        cudaStreamCreate(&s2);
        parent<<<1,1, 0, s1>>>();
        sleep(1);
        pk<<<1,1,0,s2>>>();
        cudaDeviceSynchronize();

$ nvcc -o cdp cdp.cu -rdc=true -lcudadevrt -lineinfo
$ time ./cdp                 
hello

real    0m7.276s
user    0m4.193s
sys     0m2.061s
$

（Ubuntu 18.04、CUDA 11.4、V100）

当我运行上述代码时，我观察到在启动 cdp 可执行文件后，控制台保持空闲大约 1 秒。之后控制台打印“hello”，然后控制台保持空闲 6-7 秒，然后应用程序退出。 （请不要期望在 WDDM GPU 中出现相同的行为。）

如果设备上的 NULL 流与主机上的 NULL 流“相同”或“继承自”，那么我们预计“hello”打印输出直到应用程序退出前才会出现。如果涉及主机 NULL 流语义，则在 child 内核完成之前，不允许运行 pk 内核。因此我们必须得出结论，设备 NULL 流与主机 NULL 流不同。

我们可以使用一点逻辑来说服自己，这与继承用于parent 内核启动的主机（创建）流也不完全相同。我们可以在文档中读到，在同一个线程块中使用 NULL 流启动子内核将具有 NULL 流行为。但是，如果设备 NULL 流与主机创建的流完全相同/继承自主机创建的流，则这不可能是真的。根据文档，如果我有两个子内核从同一个线程块启动，一个进入 NULL 流，另一个进入设备创建的流，那么我们不希望它们能够同时运行——这就是空流。但是，如果 NULL 流只是继承的主机流，即创建的流，则没有理由期望它们不能同时运行。

所以我们得出结论，设备 NULL 流不是主机 NULL 流，设备 NULL 流也不是主机创建的流。这些陈述在我看来与文档一致。

如果您想在文档中得到澄清，通常的建议是file a bug。

与其担心 NULL 流行为的特殊性，我在教授 CUDA 时给出的建议是，如果您担心复杂的并发场景，请使用已创建的流来完成所有工作。留下 NULL 流。您想做的任何事情都可以完全使用创建的流来完成。当然，随心所欲。