GPU 利用率低是什么意思？

Question

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

1. NumbaPerformanceWarning：网格大小 (27) < 2 * SM count (216) will likely result in GPU under utilization due to low occupancy.

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

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

Answer 1

NumbaPerformanceWarning：网格大小 (27) < 2 * SM count (216) will likely result in GPU under utilization due to low occupancy.

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

解决方案是在内核启动时公开足够的并行性（足够的线程）以完全“占用”或“利用”GPU。 216 个线程块（每个线程块有 1024 个线程）对于 A100 来说足够了。少了可能就不行了。

为了进一步理解这里，我推荐本课程的前 4 部分。

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

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

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

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

1. 创建一个设备数组，用于在设备内存中存储

   a

   ，我们称之为

   d_a

2. 将数据从

   a

   复制到

   d_a

   （主机->设备）
3. 启动你的内核，内核实际上正在使用

   d_a

4. 内核完成后，将

   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()

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

为了获得更多理解，可能任何在线教程（例如这个）或 numba cuda 文档都会有用。

Answer 2

关于第一条消息，如果您仍然想以小网格尺寸运行它，只需在终端中将环境变量

NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS

设置为0即可：

$ conda env config vars set NUMBA_CUDA_LOW_OCCUPANCY_WARNINGS=0

要查看更多详细信息，请查看 Numba - 参考手册 - 环境变量