3

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

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

4

1 回答 1

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。

于 2022-01-30T11:52:15.157 回答