由于占用率低而说GPU处于利用率低下是啥意思？

Posted 2023-03-23

【中文标题】由于占用率低而说GPU处于利用率低下是啥意思？

【英文标题】：What does it mean by say GPU under ultilization due to low occupancy?由于占用率低而说GPU处于利用率低下是什么意思？

【发布时间】：2022-01-14 06:49:04

【问题描述】：

我正在使用 NUMBA 和 cupy 执行 GPU 编码。现在我已将代码从 V100 NVIDIA 卡切换到 A100，但随后收到以下警告：

NumbaPerformanceWarning：网格大小 (27)

NumbaPerformanceWarning：CUDA 内核中使用的主机数组将产生与设备之间的复制开销。

有人知道这两个警告的真正含义吗？那我应该如何改进我的代码呢？

【参考方案1】：

NumbaPerformanceWarning：网格大小 (27)

GPU 被细分为 SM。每个 SM 可以容纳一个线程块的补充（这就像说它可以容纳一个线程的补充）。为了“充分利用”GPU，您可能希望每个 SM 都“满”，这大致意味着每个 SM 有足够的线程块来填充其线程补充。 A100 GPU 有 108 个 SM。如果您的内核在内核启动（即网格）中的线程块少于 108 个，那么您的内核将无法充分利用 GPU。一些 SM 将是空的。一个线程块不能同时驻留在 2 个或更多 SM 上。甚至 108 个（每个 SM 一个）可能还不够。一个 A100 SM 可以容纳 2048 个线程，这至少是两个线程块，每个线程块有 1024 个线程。内核启动中小于 2*108 的线程块可能无法充分利用 GPU。当您没有充分利用 GPU 时，您的性能可能不会尽可能好。

解决方案是在您的内核启动中公开足够的并行性（足够的线程）以完全“占用”或“利用”GPU。 1024 个线程的 216 个线程块对于 A100 来说已经足够了。少一点就不一定了。

这里为了进一步理解，推荐this course的前4段。

NumbaPerformanceWarning：CUDA 内核中使用的主机数组将产生与设备之间的复制开销。

关于启动 numba 内核的一个很酷的事情是我可以向它传递一个主机数据数组：

a = numpy.ones(32, dtype=numpy.int64)
my_kernel[blocks, threads](a)

numba 会“做正确的事”。在上面的例子中，它将：

创建一个设备数组，用于在设备内存中存储a，我们称之为d_a 将数据从a复制到d_a（主机->设备） 启动你的内核，内核实际使用d_a 内核完成后，将d_a的内容复制回a（设备->主机）

这一切都非常方便。但是如果我正在做这样的事情呢：

a = numpy.ones(32, dtype=numpy.int64)
my_kernel1[blocks, threads](a)
my_kernel2[blocks, threads](a)

numba 将执行上述步骤 1-4 以启动 my_kernel1，然后再次执行步骤 1-4 以启动 my_kernel2。在大多数情况下，这可能不是您作为 numba cuda 程序员想要的。

这种情况下的解决方案是“控制”数据移动：

a = numpy.ones(32, dtype=numpy.int64)
d_a = numba.cuda.to_device(a)
my_kernel1[blocks, threads](d_a)
my_kernel2[blocks, threads](d_a)
a = d_a.to_host()

这消除了不必要的复制，并且在许多情况下通常会使您的程序运行得更快。 （对于涉及单个内核启动的简单示例，可能没有区别。）

为了进一步理解，可能任何在线教程（例如 this one）或只是 numba cuda 文档都会很有用。