GPU加速02:超详细Python Cuda零基础入门教程，没有显卡也能学！

我的网站上代码和公式显示效果更好，并不断更正迭代，欢迎访问： https:// lulaoshi.info/gpu/pytho n-cuda/cuda-intro.html 。

我和滴滴云有一些合作，没有GPU的朋友可以前往滴滴云上购买GPU/vGPU/机器学习产品，记得输入AI大师码：1936，可享受9折优惠。GPU产品分时计费，比自己购买硬件更划算，请前往滴滴云官网 http://www. didiyun.com 购买。

Python是当前最流行的编程语言，被广泛应用在深度学习、金融建模、科学和工程计算上。作为一门解释型语言，它运行速度慢也常常被用户诟病。著名Python发行商Anaconda公司开发的Numba库为程序员提供了Python版CPU和GPU编程工具，速度比原生Python快数十倍甚至更多。使用Numba进行GPU编程，你可以享受：

1. Python简单易用的语法；
2. 极快的开发速度；
3. 成倍的硬件加速。

为了既保证Python语言的易用性和开发速度，又达到并行加速的目的，本系列主要从Python的角度给大家分享GPU编程方法。关于Numba的入门可以参考 我的另一篇文章 。 更加令人兴奋的是，Numba提供了一个GPU模拟器，即使你手头暂时没有GPU机器，也可以先使用这个模拟器来学习GPU编程！

本系列为英伟达GPU入门介绍的第二篇，主要介绍CUDA编程的基本流程和核心概念，并使用Python Numba编写GPU并行程序。为了更好地理解GPU的硬件架构，建议读者先阅读我的第一篇文章。

1. GPU硬件知识和基础概念：包括CPU与GPU的区别、GPU架构、CUDA软件栈简介。
2. GPU编程入门：主要介绍CUDA核函数，Thread、Block和Grid概念，并使用Python Numba进行简单的并行计算。
3. GPU编程进阶：主要介绍一些优化方法。
4. GPU编程实践：使用Python Numba解决复杂问题。

初识GPU编程

兵马未动，粮草先行。在开始GPU编程前，需要明确一些概念，并准备好相关工具。

CUDA是英伟达提供给开发者的一个GPU编程框架，程序员可以使用这个框架轻松地编写并行程序。本系列第一篇文章提到，CPU和主存被称为 主机（Host） ，GPU和显存（显卡内存）被称为 设备（Device） ，CPU无法直接读取显存数据，GPU无法直接读取主存数据，主机与设备必须通过总线（Bus）相互通信。

在进行GPU编程前，需要先确认是否安装了CUDA工具箱，可以使用 echo $CUDA_HOME 检查CUDA环境变量，返回值不为空说明已经安装好CUDA。也可以直接用Anaconda里的 conda 命令安装CUDA：

$ conda install cudatoolkit

然后可以使用 nvidia-smi 命令查看显卡情况，比如这台机器上几张显卡，CUDA版本，显卡上运行的进程等。我这里是一台32GB显存版的Telsa V100机器。

安装Numba库：

$ conda install numba

检查一下CUDA和Numba是否安装成功：

from numba import cuda
print(cuda.gpus)

如果上述步骤没有问题，可以得到结果： <Managed Device 0>... 。如果机器上没有GPU或没安装好上述包，会有报错。CUDA程序执行时会独霸一张卡，如果你的机器上有多张GPU卡，CUDA默认会选用0号卡。如果你与其他人共用这台机器，最好协商好谁在用哪张卡。一般使用 CUDA_VISIBLE_DEVICES 这个环境变量来选择某张卡。如选择5号GPU卡运行你的程序。

CUDA_VISIBLE_DEVICES='5' python example.py

如果手头暂时没有GPU设备，Numba提供了一个模拟器，供用户学习和调试，只需要在命令行里添加一个环境变量。

Mac/Linux:

export NUMBA_ENABLE_CUDASIM=1

Windows:

SET NUMBA_ENABLE_CUDASIM=1

需要注意的是，模拟器只是一个调试的工具，在模拟器中使用Numba并不能加速程序，有可能速度更慢，而且在模拟器能够运行的程序，并不能保证一定能在真正的GPU上运行，最终还是要以GPU为准。

有了以上的准备工作，我们就可以开始我们的GPU编程之旅了！

GPU程序与CPU程序的区别

一个传统的CPU程序的执行顺序如下图所示：

CPU程序是顺序执行的，一般需要：

1. 初始化。
   
2. CPU计算。
   
3. 得到计算结果。
   

在CUDA编程中，CPU和主存被称为主机（Host），GPU被称为设备（Device）。

当引入GPU后，计算流程变为：

1. 初始化，并将必要的数据拷贝到GPU设备的显存上。
2. CPU调用GPU函数，启动GPU多个核心同时进行计算。
3. CPU与GPU异步计算。
4. 将GPU计算结果拷贝回主机端，得到计算结果。

一个名为 gpu_print.py 的GPU程序如下所示：

from numba import cuda
def cpu_print():
    print("print by cpu.")
@cuda.jit
def gpu_print():
    # GPU核函数
    print("print by gpu.")
def main():
    gpu_print[1, 2]()
    cuda.synchronize()
    cpu_print()
if __name__ == "__main__":
    main()

使用 CUDA_VISIBLE_DEVICES='0' python gpu_print.py 执行这段代码，得到的结果为：

print by gpu.
print by gpu.
print by cpu.

与传统的Python CPU代码不同的是：

• 使用 from numba import cuda 引入 cuda 库
• 在GPU函数上添加 @cuda.jit 装饰符，表示该函数是一个在GPU设备上运行的函数，GPU函数又被称为 核函数 。
• 主函数调用GPU核函数时，需要添加如 [1, 2] 这样的 执行配置 ，这个配置是在告知GPU以多大的并行粒度同时进行计算。 gpu_print[1, 2]() 表示同时开启2个线程并行地执行 gpu_print 函数，函数将被并行地执行2次。下文会深入探讨如何设置执行配置。
• GPU核函数的启动方式是 异步 的：启动GPU函数后，CPU不会等待GPU函数执行完毕才执行下一行代码。必要时，需要调用 cuda.synchronize() ，告知CPU等待GPU执行完核函数后，再进行CPU端后续计算。这个过程被称为 同步 ，也就是GPU执行流程图中的红线部分。如果不调用 cuda.synchronize() 函数，执行结果也将改变，"print by cpu.将先被打印。虽然GPU函数在前，但是程序并没有等待GPU函数执行完，而是继续执行后面的 cpu_print 函数，由于CPU调用GPU有一定的延迟，反而后面的 cpu_print 先被执行，因此 cpu_print 的结果先被打印了出来。

Thread层次结构

前面的程序中，核函数被GPU并行地执行了2次。在进行GPU并行编程时需要定义执行配置来告知以怎样的方式去并行计算，比如上面打印的例子中，是并行地执行2次，还是8次，还是并行地执行20万次，或者2000万次。2000万的数字太大，远远多于GPU的核心数，如何将2000万次计算合理分配到所有GPU核心上。解决这些问题就需要弄明白CUDA的Thread层次结构。

CUDA将核函数所定义的运算称为 线程（Thread） ，多个线程组成一个 块（Block） ，多个块组成 网格（Grid） 。这样一个grid可以定义成千上万个线程，也就解决了并行执行上万次操作的问题。例如，把前面的程序改为并行执行8次：可以用2个block，每个block中有4个thread。原来的代码可以改为 gpu_print[2, 4]() ，其中方括号中第一个数字表示整个grid有多少个block，方括号中第二个数字表示一个block有多少个thread。

实际上，线程（thread）是一个编程上的软件概念。从硬件来看，thread运行在一个CUDA核心上，多个thread组成的block运行在Streaming Multiprocessor（SM的概念详见本系列第一篇文章），多个block组成的grid运行在一个GPU显卡上。

CUDA提供了一系列内置变量，以记录thread和block的大小及索引下标。以 [2, 4] 这样的配置为例： blockDim.x 变量表示block的大小是4，即每个block有4个thread， threadIdx.x 变量是一个从0到 blockDim.x - 1 （4-1=3）的索引下标，记录这是第几个thread； gridDim.x 变量表示grid的大小是2，即每个grid有2个block， blockIdx.x 变量是一个从0到 gridDim.x - 1 （2-1=1）的索引下标，记录这是第几个block。

某个thread在整个grid中的位置编号为： threadIdx.x + blockIdx.x * blockDim.x 。

利用上述变量，我们可以把之前的代码丰富一下：

from numba import cuda
def cpu_print(N):
    for i in range(0, N):
        print(i)
@cuda.jit
def gpu_print(N):
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x 
    if (idx < N):
        print(idx)
def main():
    print("gpu print:")
    gpu_print[2, 4](8)
    cuda.synchronize()
    print("cpu print:")
    cpu_print(8)
if __name__ == "__main__":
    main()

执行结果为：

gpu print:
cpu print:
7

这里的GPU函数在每个CUDA thread中打印了当前thread的编号，起到了CPU函数 for 循环同样的作用。因为 for 循环中的计算内容互相不依赖，也就是说，某次循环只是专心做自己的事情，循环第i次不影响循环第j次的计算，所以这样互相不依赖的 for 循环非常适合放到CUDA thread里做并行计算。在实际使用中，我们一般将CPU代码中互相不依赖的的 for 循环适当替换成CUDA代码。

这份代码打印了8个数字，核函数有一个参数 N ， N = 8 ，假如我们只想打印5个数字呢？当前的执行配置共2 * 4 = 8个线程，线程数8与要执行的次数5不匹配，不过我们已经在代码里写好了 if (idx < N) 的判断语句，判断会帮我们过滤不需要的计算。我们只需要把 N = 5 传递给 gpu_print 函数中就好，CUDA仍然会启动8个thread，但是大于等于 N 的thread不进行计算。 注意，当线程数与计算次数不一致时，一定要使用这样的判断语句，以保证某个线程的计算不会影响其他线程的数据。

Block大小设置

不同的执行配置会影响GPU程序的速度，一般需要多次调试才能找到较好的执行配置，在实际编程中，执行配置 [gridDim, blockDim] 应参考下面的方法：

• block运行在SM上，不同硬件架构（Turing、Volta、Pascal...）的CUDA核心数不同，一般需要根据当前硬件来设置block的大小 blockDim （执行配置中第二个参数）。一个block中的thread数最好是32、128、256的倍数。==注意，限于当前硬件的设计，block大小不能超过1024。==
• grid的大小 gridDim （执行配置中第一个参数），即一个grid中block的个数可以由总次数 N 除以 blockDim ，并向上取整。

例如，我们想并行启动1000个thread，可以将blockDim设置为128， 1000 ÷ 128 = 7.8 ，向上取整为8。使用时，执行配置可以写成 gpuWork[8, 128]() ，CUDA共启动 8 * 128 = 1024 个thread，实际计算时只使用前1000个thread，多余的24个thread不进行计算。

注意，这几个变量比较容易混淆，再次明确一下： blockDim 是block中thread的个数，一个block中的 threadIdx 最大不超过 blockDim ； gridDim 是grid中block的个数，一个grid中的 blockIdx 最大不超过 gridDim 。

以上讨论中，block和grid大小均是一维，实际编程使用的执行配置常常更复杂，block和grid的大小可以设置为二维甚至三维，如下图所示。这部分内容将在下篇文章中讨论。

内存分配

前文提到，GPU计算时直接从显存中读取数据，因此每当计算时要将数据从主存拷贝到显存上，用CUDA的术语来说就是要把数据从主机端拷贝到设备端。CUDA强大之处在于它能自动将数据从主机和设备间相互拷贝，不需要程序员在代码中写明。这种方法对编程者来说非常方便，不必对原有的CPU代码做大量改动。

我们以一个向量加法为例，编写一个向量加法的核函数如下：

@cuda.jit
def gpu_add(a, b, result, n):
    # a, b为输入向量，result为输出向量
    # 所有向量都是n维
    # 得到当前thread的索引
    idx = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
    if idx < n:
        result[idx] = a[idx] + b[idx]

初始化两个2千万维的向量，作为参数传递给核函数：

n = 20000000
x = np.arange(n).astype(np.int32)
y = 2 * x
gpu_result = np.zeros(n)
# CUDA执行配置
threads_per_block = 1024
blocks_per_grid = math.ceil(n / threads_per_block)
gpu_add[blocks_per_grid, threads_per_block](x, y, gpu_result, n)

把上述代码整合起来，与CPU代码做对比，并验证结果正确性：

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] = a[idx] + b[idx]
def main():
    n = 20000000
    x = np.arange(n).astype(np.int32)
    y = 2 * x
    gpu_result = np.zeros(n)
    cpu_result = np.zeros(n)
    threads_per_block = 1024
    blocks_per_grid = math.ceil(n / threads_per_block)
    start = time()
    gpu_add[blocks_per_grid, threads_per_block](x, y, gpu_result, n)
    cuda.synchronize()
    print("gpu vector add time " + str(time() - start))
    start = time()
    cpu_result = np.add(x, y)
    print("cpu vector add time " + str(time() - start))
    if (np.array_equal(cpu_result, gpu_result)):
        print("result correct")
if __name__ == "__main__":
    main()

运行结果，GPU代码竟然比CPU代码慢10+倍！

gpu vector add time 13.589356184005737
cpu vector add time 1.2823548316955566
result correct

说好的GPU比CPU快几十倍上百倍呢？这里GPU比CPU慢很多原因主要在于：

1. 向量加法的这个计算比较简单，CPU的numpy已经优化到了极致，无法突出GPU的优势，我们要解决实际问题往往比这个复杂得多，当解决复杂问题时，优化后的GPU代码将远快于CPU代码。
2. 这份代码使用CUDA默认的统一内存管理机制，没有对数据的拷贝做优化。CUDA的统一内存系统是当GPU运行到某块数据发现不在设备端时，再去主机端中将数据拷贝过来，当执行完核函数后，又将所有的内存拷贝回主存。在上面的代码中，输入的两个向量是只读的，没必要再拷贝回主存。
3. 这份代码没有做流水线优化。CUDA并非同时计算2千万个数据，一般分批流水线工作：一边对2000万中的某批数据进行计算，一边将下一批数据从主存拷贝过来。计算占用的是CUDA核心，数据拷贝占用的是总线，所需资源不同，互相不存在竞争关系。这种机制被称为流水线。这部分内容将在下篇文章中讨论。

原因2中本该程序员动脑思考的问题交给了CUDA解决，增加了时间开销，所以CUDA非常方便的统一内存模型缺点是计算速度慢。针对原因2，我们可以继续优化这个程序，告知GPU哪些数据需要拷贝到设备，哪些需要拷贝回主机。

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] = a[idx] + b[idx]
def main():
    n = 20000000
    x = np.arange(n).astype(np.int32)
    y = 2 * x
    # 拷贝数据到设备端
    x_device = cuda.to_device(x)
    y_device = cuda.to_device(y)
    # 在显卡设备上初始化一块用于存放GPU计算结果的空间
    gpu_result = cuda.device_array(n)
    cpu_result = np.empty(n)
    threads_per_block = 1024
    blocks_per_grid = math.ceil(n / threads_per_block)
    start = time()
    gpu_add[blocks_per_grid, threads_per_block](x_device, y_device, gpu_result, n)
    cuda.synchronize()
    print("gpu vector add time " + str(time() - start))
    start = time()
    cpu_result = np.add(x, y)
    print("cpu vector add time " + str(time() - start))
    if (np.array_equal(cpu_result, gpu_result.copy_to_host())):
        print("result correct!")