CUDA中有多少个网格？

Question

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

一个GPU设备上是否可以存在两个网格？或者说，一个GPU设备只能有一个网格吗？

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

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

- Cholgyun Ri

Answer 1

如果按照上述方式发出两个内核，则它们将被串行化（按顺序运行）。这是因为没有其他代码（例如切换流），这两个内核将被发到同一个cuda流中。所有发到同一流中的cuda调用都会按顺序执行，即使您认为应该看到其他情况，因为您正在使用cudaMemcpyAsync或类似方法。

当然可以让多个内核相对异步地运行（因此可能是并发的），但需要使用cuda流API才能实现这一点。

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

此外，回答你的第一个问题，来自CUDA C编程指南3.2.5.2节，“设备可以同时执行的最大内核启动次数因设备而异，但某些设备可能高达128次”。参考一下，“网格”是与单个内核启动相关联的整个线程数组。

- Robert Crovella

---

您还需要回答最初的问题：“GPU 上可能有多少[并发]网格”- CC <2.0 时为1个，2.0 <= CC <= 3.0 时为16个，CC = 3.5 时为32个。 - harrism

@harrism 这可能与主题无关，但在这种情况下Hyper-Q提供了什么？ - lashgar

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

“一个设备可以同时执行的最大内核启动次数为16”是否与设备有关？即它是否因GPU而异？ - AAEM

是的，它在编程指南第3.2.5.2节中有详细介绍，可以查看这里。 - Robert Crovella

Answer 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，即使您将内核分布到多个流中，您的工作有时也可能被序列化。

- solvingPuzzles