Python Numba | 多流和共享内存CUDA优化技术介绍和代码示例

最新推荐文章于 2023-05-26 21:39:09 发布

皮皮鲁同学

最新推荐文章于 2023-05-26 21:39:09 发布

阅读量4.9k

点赞数 13

分类专栏： GPU Python 计算加速文章标签： GPU Python Numba CUDA 计算加速

本文链接：https://blog.csdn.net/qq_42596142/article/details/103183209

版权

公众号

本文为英伟达GPU计算加速系列的第三篇，前两篇文章为：

AI时代人人都应该了解的GPU知识：主要介绍了CPU与GPU的区别、GPU架构、CUDA软件栈简介。
超详细Python Cuda零基础入门教程：主要介绍了CUDA核函数，Thread、Block和Grid概念，内存分配，并使用Python Numba进行简单的并行计算。

阅读完前两篇文章后，相信读者应该能够将一些简单的CPU代码修改成GPU并行代码，但是对计算密集型任务，仅仅使用前文的方法还是远远不够的，GPU的并行计算能力未能充分利用。本文将主要介绍一些常用性能优化的进阶技术，这部分对编程技能和硬件知识都有更高的要求，建议读者先阅读本系列的前两篇文章，甚至阅读英伟达官方的编程手册，熟悉CUDA编程的底层知识。当然，将这些优化技巧应用之后，程序将获得更大的加速比，这对于需要跑数小时甚至数天的程序来说，收益非常之大。

本文仍然使用Python版的Numba库调用CUDA，有更复杂需求的朋友可以直接使用C/C++调用CUDA，并阅读英伟达的官方文档。C/C++对数据的控制更细致，是英伟达官方推荐的编程语言，所能提供的编程接口更全面。

CUDA C Programming Guide ：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

下一篇文章将提供实战案例，包括金融领域期权定价的GPU实现。

CUDA优化方向

我之前的文章中提到，CPU + GPU 是一种异构计算的组合，各有独立的内存，GPU的优势是更多的计算核心。该架构在并行计算上有很大优势，但是数据需要从主机和设备间相互拷贝，会造成一定的延迟。因此，要从下面两个方面来优化GPU程序：

充分利用GPU的多核心，最大化并行执行度
优化内存使用，最大化数据吞吐量，减少不必要的数据拷贝

哪个方向有更大收益，最终还是要看具体的计算场景。英伟达提供了非常强大的性能分析器nvprof和可视化版nvvp，使用性能分析器能监控到当前程序的瓶颈。据我了解，分析器只支持C/C++编译后的可执行文件，Python Numba目前应该不支持。

Profiler User’s Guide : https://docs.nvidia.com/cuda/profiler-users-guide/index.html

并行计算优化

网格跨度

在上一篇文章中，我曾提到，CUDA的执行配置：[gridDim, blockDim]中的blockDim最大只能是1024，但是并没提到gridDim的最大限制。英伟达给出的官方回复是gridDim最大为一个32位整数的最大值，也就是2,147,483,648，大约二十亿。这个数字已经非常大了，足以应付绝大多数的计算，但是如果对并行计算的维度有更高需求呢？网格跨度有更好的并行计算效率。

并行计算数大于线程数

这里仍然以[2, 4]的执行配置为例，该执行配置中整个grid只能并行启动8个线程，假如我们要并行计算的数据是32，会发现后面8号至31号数据共计24个数据无法被计算。

网格跨度

我们可以在0号线程中，处理第0、8、16、24号数据，就能解决数据远大于执行配置中的线程总数的问题，用程序表示，就是在核函数里再写个for循环。以打印为例，代码如下：

from numba import cuda

@cuda.jit
def gpu_print(N):
    idxWithinGrid = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x 
    gridStride = cuda.gridDim.x * cuda.blockDim.x
    # 从 idxWithinGrid 开始
    # 每次以整个网格线程总数为跨步数
    for i in range(idxWithinGrid, N, gridStride):
        print(i)

def main():
    gpu_print[2, 4](32)
    cuda.synchronize()

if __name__ == "__main__":
    main()

注意，跨步大小为网格中线程总数，用gridDim.x * blockDim.x来计算。for循环的step是网格中线程总数，这也是为什么将这种方式称为网格跨步。如果网格总线程数为1024，那么0号线程将计算第0、1024、2048…号的数据。这里我们也不用再明确使用if (idx < N)来判断是否越界，因为for循环也有这个判断。

使用网格跨步的优势主要有：

扩展性：可以解决数据量比线程数大的问题
线程复用：CUDA线程启动和销毁都有开销，主要是线程内存空间初始化的开销；不使用网格跨步，CUDA需要启动大于计算数的线程，每个线程内只做一件事情，做完就要被销毁；使用网格跨步，线程内有for循环，每个线程可以干更多事情，所有线程的启动销毁开销更少。
方便调试：我们可以把核函数的执行配置写为[1, 1]，如下所示，那么核函数的跨步大小就成为了1，核函数里的for循环与CPU函数中顺序执行的for循环的逻辑一样，非常方便验证CUDA并行计算与原来的CPU函数计算逻辑是否一致。

kernel_function[1,1](...)

多流

之前我们讨论的并行，都是线程级别的，即CUDA开启多个线程，并行执行核函数内的代码。GPU最多就上千个核心，同一时间只能并行执行上千个任务。当我们处理千万级别的数据，整个大任务无法被GPU一次执行，所有的计算任务需要放在一个队列中，排队顺序执行。CUDA将放入队列顺序执行的一系列操作称为流（Stream）。

由于异构计算的硬件特性，CUDA中以下操作是相互独立的，通过编程，是可以操作他们并发地执行的：

主机端上的计算
设备端的计算（核函数）
数据从主机和设备间相互拷贝
数据从设备内拷贝或转移
数据从多个GPU设备间拷贝或转移

数据拷贝和计算的重叠

针对这种互相独立的硬件架构，CUDA使用多流作为一种高并发的方案：把一个大任务中的上述几部分拆分开，放到多个流中，每次只对一部分数据进行拷贝、计算和回写，并把这个流程做成流水线。因为数据拷贝不占用计算资源，计算不占用数据拷贝的总线（Bus）资源，因此计算和数据拷贝完全可以并发执行。如图所示，将数据拷贝和函数计算重叠起来的，形成流水线，能获得非常大的性能提升。实际上，流水线作业的思想被广泛应用于CPU和GPU等计算机芯片设计上，以加速程序。

默认流与多流

以向量加法为例，上图中第一行的Stream 0部分是我们之前的逻辑，没有使用多流技术，程序的三大步骤是顺序执行的：先从主机拷贝初始化数据到设备（Host To Device）；在设备上执行核函数（Kernel）；将计算结果从设备拷贝回主机（Device To Host）。当数据量很大时，每个步骤的耗时很长，后面的步骤必须等前面执行完毕才能继续，整体的耗时相当长。以2000万维的向量加法为例，向量大约有几十M大小，将整个向量在主机和设备间拷贝将占用占用上百毫秒的时间，有可能远比核函数计算的时间多得多。将程序改为多流后，每次只计算一小部分，流水线并发执行，会得到非常大的性能提升。

默认情况下，CUDA使用0号流，又称默认流。不使用多流时，所有任务都在默认流中顺序执行，效率较低。在使用多流之前，必须先了解多流的一些规则：

给定流内的所有操作会按序执行。
非默认流之间的不同操作，无法保证其执行顺序。
所有非默认流执行完后，才能执行默认流；默认流执行完后，才能执行其他非默认流。

参照上图，可将这三个规则解释为：

非默认流1中，根据进流的先后顺序，核函数1和2是顺序执行的。
无法保证核函数2与核函数4的执行先后顺序，因为他们在不同的流中。他们执行的开始时间依赖于该流中前一个操作结束时间，例如核函数2的开始依赖于核函数1的结束，与核函数3、4完全不相关。
默认流有阻塞的作用。如图中红线所示，如果调用默认流，那么默认流会等非默认流都执行完才能执行；同样，默认流执行完，才能再次执行其他非默认流。

可见，某个流内的操作是顺序的，非默认流之间是异步的，默认流有阻塞作用。

如果想使用多流时，必须先定义流：

stream = numba.cuda.stream()

CUDA的数据拷贝以及核函数都有专门的stream参数来接收流，以告知该操作放入哪个流中执行：

numba.cuda.to_device(obj, stream=0, copy=True, to=None)
numba.cuda.copy_to_host(self, ary=None, stream=0)

核函数调用的地方除了要写清执行配置，还要加一项stream参数：

kernel[blocks_per_grid, threads_per_block, stream=0]

根据这些函数定义也可以知道，不指定stream参数时，这些函数都使用默认的0号流。

对于程序员来说，需要将数据和计算做拆分，分别放入不同的流里，构成一个流水线操作。

将之前的向量加法的例子改为多流处理，完整的代码为：

from numba import cuda
import numpy as np
import math
from time import time

@cuda.jit
def gpu_add(a, b, result, n):
    idx = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
    if idx < n :
        result[idx] =