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 会“做正确的事”。在上面的例子中它将:
- 创建一个设备数组,用于存储
a
在设备内存中,我们称之为d_a
- 复制数据来自
a
to 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-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 文档,将会很有用。