如何在 CUDA 中同时执行子内核 256 次答案

【问题标题】：How to execute a child kernel 256 times concurrently in CUDA如何在 CUDA 中同时执行子内核 256 次
【发布时间】：2019-08-30 18:54:59
【问题描述】：

我是 CUDA 编程的新手，但我需要在复杂的项目中使用它。我真的需要一些帮助。

我的问题是如果我想同时执行一个子内核 256 次，我可以用动态并行做什么？

我读到了 NVIDIA blog，上面写着：

默认情况下，在线程块中启动的网格会被执行顺序：下一个网格仅在前一个网格之后开始执行一个已经完成。即使网格是由不同的启动的，也会发生这种情况块内的线程。

所以，我的想法是为父内核设置块大小（1,1）和网格大小（256,1），我可以在不同的块中使用 256 个线程同时启动子内核。会不会非常低效？有什么更好的解决方案？

【问题讨论】：

不太清楚，为什么要同时启动它们。正如您所描述的，只需启动一个具有 256 个线程的块就足够了……
在一个块中启动单个线程效率非常低，应尽可能避免。如果可能的话，你应该以 warp 大小的倍数（当前硬件为 32）启动。
@DanielBauer 抱歉，这是我第一次在这个论坛上提问。让我更具体一点。我有一个大小为 (512,1024) 的数据矩阵。我需要对每一行对执行一次操作，并且我希望此操作同时执行 256 次以节省时间。该操作包括计算FFT（或卷积），获取FFT结果中最大值的索引，并将数据移动到其原始位置。（实际上，这个过程在逆合成孔径雷达成像中称为距离对齐）。
尝试制作一个同时使用尽可能多的资源的内核，然后在一个大网格中运行这个内核。 CUDA 应该自动同时处理尽可能多的块。通常，如果您只重复运行一项任务，则不必对并发性做任何事情。如果您要计算不同的任务，其中一个不清楚是否需要来自另一个的数据，您需要更多地考虑这些东西

标签： cuda gpu

【解决方案1】：

那句话继续

然而，通常需要更多的并发性；与主机端内核启动一样，我们可以使用 CUDA 流来实现这一点。在设备上创建的所有流都是非阻塞的；也就是说，它们不支持与 default NULL 流的隐式同步。因此，以下是在设备代码中创建流的唯一方法。

cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);

然后为每个 CUDA 线程使用不同的（设备端）流应该使它们独立运行，而不是默认运行。

此外，您可以使用父块中父线程之间的某种缩减算法将多个启动合并为一个大型启动。增加子内核的线程总数及其从线程 id 到问题空间的映射。这应该可以克服小内核的性能问题以及硬件支持的每个设备的最大并发内核执行数（4 到 128，具体取决于 Cuda 计算能力）。

【讨论】：

当前最大网格数在4-128之间。这因计算能力而异。请参阅docs.nvidia.com/cuda/cuda-c-programming-guide/…“每个设备的最大驻留网格数”。 32是从GK110开始的最大通讯通道数。