CUDA中有多少个网格

Posted 2023-04-15

【中文标题】CUDA中有多少个网格

【英文标题】：How many grids in CUDA

【发布时间】：2012-09-29 20:03:18

【问题描述】：

GPU 中可能有多少个 CUDA 网格？

两个网格可以同时存在于 GPU 中吗？还是一台 GPU 设备只有一个网格？

Kernel1<<gridDim, blockDim>>(dst1, param1);
Kernel1<<gridDim, blockDim>>(dst2, param2);

上面的两个内核是同时运行还是顺序运行？

【参考方案1】：

如果两个内核如上所示发出，它们将被序列化（它们将按顺序运行）。这是因为没有任何其他代码（即切换流），两个内核将被发布到同一个 cuda 流。对同一流发出的所有 cuda 调用都按顺序执行，即使您认为您应该看到其他情况，因为您使用的是 cudaMemcpyAsync 或类似的东西。

当然可以让多个内核彼此异步运行（因此可能同时运行），但必须使用 cuda 流 API 来完成此操作。

您可能想查看CUDA C Programmers Guide 中的第 3.2.5 节“异步并发执行”以了解有关流和并发内核执行的更多信息。此外，nvidia CUDA SDK 中有许多示例，例如simple streams，将说明这些概念。 concurrent kernels 示例展示了如何同时运行多个内核（使用流）。请注意，同时运行内核需要计算能力 2.0 或“更高”的硬件。

另外，回答您的第一个问题，来自 CUDA C 编程指南的section 3.2.5.2，“设备可以同时执行的内核启动的最大数量因设备而异，但对于某些设备可能高达 128”

作为参考，“网格”是与单个内核启动相关联的整个线程数组。

【讨论】：

您还应该回答最初的问题：“GPU 上可能有多少 [并发] 网格”——CC

@harrism 可能完全不相关，但是 Hyper-Q 在这种情况下提供了什么？

当内核由 CPU 上的独立进程启动时，Hyper-Q 提供了更有效地使用 GPU 的可能性。 cuda 流 API 通常用于管理从单个进程启动的重叠内核。

@RobertCrovella "设备可以同时执行的最大内核启动次数为 16" 与设备有关吗？即它与一个 GPU 不同？

是的，它在编程指南第 3.2.5.2 节中有介绍，请参阅here

【参考方案2】：

为了详细说明 Robert 的回答，这里有一个示例，说明如何使用流来使 Kernel1 的两个实例同时运行：

cudaStream_t stream1; cudaStreamCreate(&stream1);
cudaStream_t stream2; cudaStreamCreate(&stream2);

Kernel1<<gridDim, blockDim, 0, stream1>>(dst1, param1);
Kernel1<<gridDim, blockDim, 0, stream2>>(dst2, param2);

关于流并发执行的更多注意事项：

如果我们在没有指定流Kernel1<<<g, b>>>() 的情况下启动内核，然后使用特定流Kernel2<<<g, b, 0, stream>>>() 启动内核，那么Kernel2 将等待Kernel1 完成。 在没有流 (Kernel1<<<g, b>>>()) 的情况下启动内核时，Nvidia 将其称为“使用 NULL 流”。 如果您使用 cudaEvents，即使您将内核分发到多个流中，您的工作有时也会被序列化。