同步多个 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()and cudaStreamSynchronize()，但它们使总执行时间增加了两倍多。

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

非常感谢。

Answer 1

首先需要提出两个意见。

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

2. 见证实际的内核并发性很困难。我们需要执行时间相对较长但对 GPU 资源需求相对较低的内核。内核<<<32,32>>>可能会填满您正在运行的 GPU，从而阻止并发内核中的块运行。

您的问题似乎可以归结为“我如何防止kernel6在kernel2-5完成之前开始”。

为此可以使用事件。基本上，您将在 kernel2-5 启动后将一个事件记录到每个流中，cudaStreamWaitEvent并在.kernel6

就像这样：

        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()or不同cudaStreamSynchronize()，后者通常会阻止 CPU 线程。

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

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