如何暂停 CUDA 流然后恢复它？答案

【问题标题】：How can I pause a CUDA stream and then resume it?如何暂停 CUDA 流然后恢复它？
【发布时间】：2022-01-30 20:28:37
【问题描述】：

假设我们有两个 CUDA 流同时在 GPU 上运行两个 CUDA 内核。如何暂停使用我放入主机代码的指令运行的 CUDA 内核并使用主机代码中的指令恢复它？我不知道在这种情况下如何编写示例代码，例如继续这个问题。

我的问题是，CUDA 中是否有一条指令可以暂停在 CUDA 流中运行的 CUDA 内核，然后恢复它？

【问题讨论】：

AFAIK，这不可能“暂停”正在运行的 CUDA 内核。所有基于流的函数都可用here。
CUDA 内核是运行一次并且您真的想要中断它，还是多次运行并且您想要调用之间的中断？
suspendDevice 函数可能是一种方法吗？ docs.nvidia.com/cuda/debugger-api/group__EXEC.html#group__EXEC 否则将内核重写为几个较小的内核或检查内核内部的标志，何时存储状态并终止。还请告诉我们，您为什么要暂停？为第一个内核或例如获得更多性能只是等待更多数据。在第二种情况下，内核可以继续运行，但一段时间内什么也不做。
@Sebastian，CUDA 内核运行一次。我有两个 CUDA 流，每个流都运行一个内核，这两个内核同时运行，我希望这两个 CUDA 流中的一个在出现条件时暂停运行其内核，并在需要时继续运行。
您能否说一下，是否出于性能原因暂停（例如，其他内核是否应该运行得更快）？您的内核是否有很多块（比 GPU 上的多处理器多得多）？线程的状态有多复杂，可以手动保存吗？

标签： cuda cuda-streams

【解决方案1】：

您可以使用带有参数的动态并行性与主机进行信号通信。然后启动一个只有 1 个 cuda 线程的父内核，让它连续启动子内核，直到工作完成或收到信号。如果子内核没有完全占用 GPU，就会失去性能。

__global__ void parent(int * atomicSignalPause, int * atomicSignalExit, Parameters * prm)
{
       int progress = 0;
       while(checkSignalExit(atomicSignalExit) && progress<100)
       {
           while(checkSignalPause(atomicSignalPause))
           {
                 child<<<X,Y>>>(prm,progress++);
                 cudaDeviceSynchronize();
           }
       }
}

没有暂停流的命令。对于多个 GPU，您应该使用统一的内存分配进行通信（GPU 之间）。

为了克服 gpu 利用率问题，您可以为子内核发明一个任务队列。它推送工作 N 次（大致足以保持 GPU 在功率/计算方面的效率），然后对于每个完成的子内核，它会在父内核中增加一个专用计数器并推送一个新工作，直到所有工作完成（同时尝试保持并发N 处的内核）。

可能是这样的：

// producer kernel
// N: number of works that make gpu fully utilized
while(hasWork)
{
     // concurrency is a global parameter
     while(checkConcurrencyAtomic(concurrency)<N)
     {
         incrementConcurrencyAtomic(concurrency);

         // a "consumer" parent kernel will get items from queue
         // it will decrement concurrency when a work is done             
         bool success = myQueue.tryPush(work, concurrency);
         if(success)
         {
            // update status of whole work or signal the host
         }
     }

     // synchronization once per ~N work
     cudaDeviceSynchronize();
     ... then check for pause signals and other tasks
}

如果总工作时间超过几秒钟，这些原子值更新不应该是性能问题，但如果您有太多子内核要启动，那么您可以启动更多生产者/消费者（父）cuda-threads。

【讨论】：