由于占用率低而导致 GPU 利用率不足是什么意思?

2024-01-18

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

  1. NumbaPerformanceWarning:网格大小 (27)

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

有谁知道这两个警告到底意味着什么?那我应该如何改进我的代码呢?


NumbaPerformanceWarning:网格大小 (27)

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 部分这个课程 https://www.olcf.ornl.gov/cuda-training-series/.

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

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

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

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

  1. 创建一个设备数组,用于存储a在设备内存中,我们称之为d_a
  2. 复制数据来自a to 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-4again为推出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 https://nyu-cds.github.io/python-numba/05-cuda/,或者只是 numba cuda 文档,将会很有用。

本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)

由于占用率低而导致 GPU 利用率不足是什么意思? 的相关文章

随机推荐