如何在内核启动后让 CUDA 返回控制权？

Question

这可能是一个愚蠢的问题，但是有没有办法从内核异步返回？例如，我有这个内核，它执行第一次流压缩并输出给用户，但在此之前必须执行第二次流压缩以更新其内部结构。

有没有办法在第一次流压缩完成后将控制权返回给用户，同时 GPU 在后台继续进行第二次流压缩？当然，第二次流压缩仅适用于共享内存和全局内存，但用户不应检索任何内容。

我无法使用推力。

Answer 1

GPU 内核本身并不从“用户”（即带有 GPU 的系统上的 CPU 线程）获取控制权。

但是，对于 CUDA 的运行时，调用 GPU 内核的默认方式是让线程等待，直到内核执行结束：

my_kernel<<<my_grid_dims,my_block_dims,dynamic_shared_memory_size>>>(args,go,here);

但您也可以使用streams。这些是硬件支持的执行队列，您可以在其中异步地排队工作（内存复制、内核执行等），就像您所要求的那样。

在这种情况下，您的启动可能如下所示：

cudaStream_t my_stream;
cudaError_t result = cudaStreamCreateWithFlags(&my_stream, cudaStreamNonBlocking);  
if (result != cudaSuccess) { /* error handling */ }

my_kernel<<<my_grid_dims,my_block_dims,dynamic_shared_memory_size,my_stream>>>(args,go,here);

有很多关于使用流的资源；首先尝试这篇博客文章。 CUDA 编程指南有很大一部分是关于 异步执行 .

流和各种库

Thrust 使用

thrust::future

和其他结构提供异步功能已经有一段时间了。请参阅此处。

我自己的 Modern-C++ CUDA API 包装器 使使用流变得更加容易，使您无需始终检查错误并记住在流超出范围之前销毁流并释放内存。使使用流变得更容易。请参阅此示例；语法看起来像这样：

auto stream = device.create_stream(cuda::stream::async);
stream.enqueue.copy(d_a.get(), a.get(), nbytes);
stream.enqueue.kernel_launch(my_kernel, launch_config, d_a.get(), more, args);

（错误会引发异常）