同步多个 Cuda 流

Question

对于我目前正在开发的应用程序，我希望有一个长内核（即相对于其他内核需要很长时间才能完成的内核）与一系列同时运行的多个较短内核同时执行.然而，更复杂的是，四个较短的内核在完成后都需要同步，以便执行另一个短内核来收集和处理其他短内核输出的数据。

以下是我想到的示意图，带有编号的绿色条代表不同的内核：

为了实现这一点，我编写了如下代码：

// definitions of kernels 1-6

class Calc
{
    Calc()
    {
        // ...
        cudaStream_t stream[5];
        for(int i=0; i<5; i++) cudaStreamCreate(&stream[i]);
        // ...
    }

    ~Calc()
    {
        // ...
        for(int i=0; i<5; i++) cudaStreamDestroy(stream[i]);
        // ...
    }

    void compute()
    {
        kernel1<<<32, 32, 0, stream[0]>>>(...);
        for(int i=0; i<20; i++) // this 20 is a constant throughout the program
        {
            kernel2<<<1, 32, 0, stream[1]>>>(...);
            kernel3<<<1, 32, 0, stream[2]>>>(...);
            kernel4<<<1, 32, 0, stream[3]>>>(...);
            kernel5<<<1, 32, 0, stream[4]>>>(...);
            // ?? synchronisation ??
            kernel6<<<1, 32, 0, stream[1]>>>(...);
        }
    }
}

int main()
{
    // preparation

    Calc C;

    // run compute-heavy function as many times as needed
    for(int i=0; i<100; i++)
    {
        C.compute();
    }

    // ...

    return 0;
}

注意：块、线程和共享内存的数量只是任意数字。

现在，我将如何在每次迭代时正确同步内核 2-5？一方面，我不知道哪个内核需要最长时间才能完成，因为这可能取决于用户输入。此外，我尝试过使用cudaDeviceSynchronize() 和cudaStreamSynchronize()，但它们的总执行时间增加了三倍多。

Cuda 事件可能是要走的路吗？如果是这样，我应该如何应用它们？如果不是，那么正确的方法是什么？

非常感谢。

Answer 1

需要先制作两个cmets。

1. 启动小内核（一个块）通常不是从 GPU 获得良好性能的方法。同样，每个块 (32) 具有少量线程的内核通常会施加占用限制，这将阻止 GPU 的全部性能。启动多个并发内核并不能减轻第二个考虑。我不会在这里再花时间，因为你说这些数字是任意的（但请参阅下面的下一条评论）。

2. 目睹实际的内核并发性很难。我们需要执行时间相对较长但对 GPU 资源需求相对较低的内核。 <<<32,32>>> 的内核可能填充您正在运行的 GPU，从而阻止来自并发内核的任何块运行。

您的问题似乎归结为“我如何防止kernel6 开始直到kernel2-5 完成。

可以为此使用事件。基本上，您将在 kernel2-5 启动后将record an event 放入每个流中，然后在@987654327 启动之前为 4 个事件中的每一个调用一个 cudaStreamWaitEvent @。

像这样：

        kernel2<<<1, 32, 0, stream[1]>>>(...);
        cudaEventRecord(event1, stream[1]);
        kernel3<<<1, 32, 0, stream[2]>>>(...);
        cudaEventRecord(event2, stream[2]);
        kernel4<<<1, 32, 0, stream[3]>>>(...);
        cudaEventRecord(event3, stream[3]);
        kernel5<<<1, 32, 0, stream[4]>>>(...);
        cudaEventRecord(event4, stream[4]);
        // ?? synchronisation ??
        cudaStreamWaitEvent(stream[1], event1);
        cudaStreamWaitEvent(stream[1], event2);
        cudaStreamWaitEvent(stream[1], event3);
        cudaStreamWaitEvent(stream[1], event4);
        kernel6<<<1, 32, 0, stream[1]>>>(...);

请注意，以上所有调用都是异步。它们都不会花费超过几微秒的时间来处理，并且它们都不会阻止 CPU 线程继续运行，这与您使用 cudaDeviceSynchronize() 或 cudaStreamSynchronize() 不同，它们通常会阻塞 CPU线程。

因此，您可能希望在循环执行上述序列（例如cudaStreamSynchronize(stream[1]);）之后进行某种同步，否则所有这些的异步性质将变得难以弄清楚（另外，基于在您的示意图上，您似乎不希望迭代 i+1 的 kernel2-5 开始，直到迭代 i 的 kernel6 完成？）请注意，我已经省略了事件创建以及其他样板，我' m 假设您可以弄清楚或参考任何使用事件的示例代码，或参考文档。

即使您实现了所有这些基础架构，您是否能够见证（或不见证）实际内核并发将取决于您的内核本身，不是我在此答案中建议的任何内容。因此，如果您回来说“我这样做了，但我的内核没有同时运行”，这实际上与您在这里提出的问题不同，我建议您首先参考我上面的评论 #2。