如何让 CUDA 流等待尚未安排的工作？ （即类似用户事件的模式）

Question

我想在 CUDA 流上做一些工作，比如内核

K

，这取决于之前需要在 CPU 上完成的工作。我在调度时并不知道CPU工作的具体细节

K

；我只是希望

K

在指示一切准备就绪之前不要开始。

现在，如果我确切地知道要完成哪些 CPU 工作，例如

K

可以在某些功能

foo()

得出结论后启动，我可以执行以下操作：

• 在流 SideStream 上对

  foo()

  的调用进行排队
• 在 SideStream 上将事件入队

  E1

• 在 MainStream 上排队等待事件

  E1

在主流上
• K

  入队

但是 - 我的 CUDA 调度代码无法访问这样的

foo()

？我想允许代码中的其他任意位置在 E1 良好且准备就绪时触发 E1，并在 MainStream 上触发 K。 ...但我不能这样做，因为在 CUDA 中，您只能等待已经排队（已经“记录”）的事件。

这似乎是 OpenCL 提供比 CUDA 更丰富的 API 的少数几个领域之一：“用户事件”。它们可以被等待，并且它们的执行完成状态可以由用户设置。参见：

• https://registry.khronos.org/OpenCL/sdk/3.0/docs/man/html/clCreateUserEvent.html
• https://registry.khronos.org/OpenCL/sdk/3.0/docs/man/html/clSetUserEventStatus.html

但是，CUDA 本身当然能够提供类似的功能，即使只是为了实现 OpenCL API 调用。那么，使用 CUDA 实现这种效果的惯用方法是什么？

Answer 1

可以在 K 之前启动一个内核，它只是等待主机设置一个标志。对于较新的 GPU，

cuda::latch

 可能更高效，因为它似乎在旋转时使用 nanosleep 函数

#include <cstdio>
#include <chrono>
#include <thread>

#include <cuda/latch>

__global__ 
void kernel(){
    printf("kernel\n");
}

__global__ 
void waitKernel(volatile int* flag){
    while(*flag != 1);
}

__global__ 
void waitKernelLatch(cuda::latch<cuda::thread_scope_system>* latchPtr){
    latchPtr->wait();
}

int main(){
    int* waitFlag;
    cudaMallocHost(&waitFlag, sizeof(int));

    cuda::latch<cuda::thread_scope_system>* latchPtr;
    cudaMallocHost(&latchPtr, sizeof(cuda::latch<cuda::thread_scope_system>));
    

    printf("wait using flag\n");
    *waitFlag = 0;
    waitKernel<<<1,1>>>(waitFlag);
    kernel<<<1,1>>>();

    printf("do some cpu stuff\n");
    std::this_thread::sleep_for(std::chrono::seconds(3));

    *waitFlag = 1;
    cudaDeviceSynchronize();



    printf("wait using latch\n");
    new (latchPtr) cuda::latch<cuda::thread_scope_system>(1);
    waitKernelLatch<<<1,1>>>(latchPtr);
    kernel<<<1,1>>>();

    printf("do some cpu stuff\n");
    std::this_thread::sleep_for(std::chrono::seconds(3));

    latchPtr->count_down();
    cudaDeviceSynchronize();


    cudaFreeHost(waitFlag);
}

Answer 2

人们可以使用 CUDA 的“流排序内存操作”功能，完全避免主机函数调度：

指定一个 32 位值（最好对齐良好）
• v

   用于向 MainStream 发送信号。

在
• cuMemHostRegister() 上调用 v

  ，以获取其设备地址（可能与其主机地址相同）。

在 MainStream 上的
• v

   上排队等待（使用 

  cuStreamWaitValue32()）
入队
• K
将
• v

   传递给任何安排额外 CPU 工作的代码。

当 CPU 工作完成时，确保将
• 1

   写入 

  v

  。

Answer 3

这是一个可能的想法 - 基于@AbatorAbetor 的评论，尽管我不知道这是否是人们在实践中使用的。

编写一个函数
• foo()

  ，它将条件变量作为参数并等待该变量。例如，您可以使用 

  std::condition_variable。
定义条件变量。

现在按照您的问题进行操作 - 因为您恰好拥有您缺少的功能：

在流 SideStream 上对
• foo()

   的调用进行排队

在 SideStream 上将事件入队
• E1
在 MainStream 上排队等待事件
• E1在主流上
K
• 入队
但是您还没有完全完成：您的调度程序现在向前/向外传递条件变量（同时保持其活动状态！），以便最终您提到的“CPU 工作”具有对它的引用。完成后，它需要做的就是对条件变量进行通知操作：这将唤醒

foo()

，然后立即触发

E

，然后

K

。警告：我假设让 CUDA 回调函数块像这样不会干扰其他 CUDA 运行时/驱动程序工作。