Python与GPU
Python
作为解释型语言,.py
文件一般是没法直接用GPU
加速的,关于Python
与GPU
的结合点,以及GPU
、CPU
、CUDA
、多核
、并行
、机器码
…等底层实现
Numba加速Python:在CPU下
Numba
加速Python
,可以基于CPU
也可以基于GPU
Numba加速Python:在GPU下
Numba
还可以使用GPU
进行加速,目前支持英伟达的CUDA
和AMD的ROC
。
CUDA
是英伟达提供给开发者的一个GPU编程框架
。
CPU+主存
,称为主机(Host
);GPU+显存
,称为设备(Device
)。
CPU
无法直接读取显存数据,GPU
无法直接读取主存数据;主机与设备必须通过总线(Bus
)相互通讯。
安装CUDA及Numba
在进行GPU
编程之前,先确认是否安装了CUDA
工具箱。
# linux环境
echo $CUDA_HOME # 检查CUDA环境变量
# 查看显卡情况
nvidia-smi
可以用Anaconda
里的conda
命令安装CUDA
:
$ conda install cudatoolkit
安装Numba
库:
$ conda install numba
检查CUDA
和Numba
是否安装成功:
from numba import cuda
print(cuda.gpus)
一切顺利,得到<Managed Device 0>...
CUDA程序执行时会独占一张卡,如果机器上有多张GPU卡,CUDA默认会选用0号卡。
CUDA_VISIBLE_DEVICES='5'
此环境变量用来选择某张卡。
目前手上没有GPU,没法测试上述代码。
Numba模拟器
如果手头暂时没有GPU
设备,Numba
提供了一个模拟器,供用户学习和调试,只需要在命令行里添加一个环境变量。
Windows
:
SET NUMBA_ENABLE_CUDASIM = 1
GPU程序与CPU程序的区别
CPU顺序执行
CPU
程序,顺序执行流程:
初始化;
CPU计算;
得到计算结果;
GPU执行
在CUDA
编程中,CPU和主存
被称为主机(Host
),GPU
被称为设备(Device
)。
GPU
计算流程:
初始化,并将必要的数据拷贝到GPU设备的显存上;
CPU
调用GPU
函数,启动GPU
多个核心同时进行计算;
CPU
与GPU
异步计算;
将GPU
计算结果拷贝回主机端,得到计算结果;
GPU程序gpu_print.py解析
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.
解析上述代码,与传统的CPU
运行Python
不同:
- 使用
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层次结构
CUDA
将核函数
(GPU函数
)所定义的运算
称为线程(Thread
),多个线程组成一个块(Block
),多个块组成网格(Grid
)。
这样的方式,一个grid
可以定义成千上万个线程,也就解决了并行执行上万次操作的问题。
例如,把前面的程序改为并行执行8次:可以用2个block
,每个block
中有4个thread
。
原来的代码可以改为gpu_print[2,4]()
,方括号内第一个数字表示整个grid
有多少个blcok
,第二个数字表示,每个block
有多少个thread
。
线程(Thread
)是一个编程上的软件概念。从硬件角度,thread
运行在一个CUDA核心
上,多个thread
组成的blcok
运行在Streaming Multiprocessor
上,多个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,即每个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:
0
1
2
3
4
5
6
7
cpu print:
0
1
2
3
4
5
6
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, blcokDim]
应参考下面的方法:
-
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
都是一维的,实际可以是二维三维。
内存分配
GPU
计算时直接从显存中读取数据,因此每当计算时要将数据从主存拷贝到显存上,用CUDA
的术语来说就是要把数据从主机端拷贝到设备端。
CUDA
强大之处在于它能自动将数据从主机和设备间相互拷贝,不需要程序员在代码中写明。
这种方法对编程者来说非常方便,不必对原有的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多倍。主要原因有:
- 向量加法的这个计算比较简单,
CPU
的numpy
已经优化到了极致,无法突出GPU
的优势,我们要解决实际问题往往比这个复杂得多,当解决复杂问题时,优化后的GPU
代码将远快于CPU
代码。 - 这份代码使用
CUDA
默认的统一内存管理机制,没有对数据的拷贝做优化。CUDA
的统一内存系统是当GPU
运行到某块数据发现不在设备端时,再去主机端中将数据拷贝过来,当执行完核函数后,又将所有的内存拷贝回主存。在上面的代码中,输入的两个向量是只读的,没必要再拷贝回主存。 - 这份代码没有做流水线优化。
CUDA
并非同时计算2千万个数据,一般分批流水线工作:一边对2000
万中的某批数据进行计算,一边将下一批数据从主存拷贝过来。计算占用的是CUDA
核心,数据拷贝占用的是总线,所需资源不同,互相不存在竞争关系。这种机制被称为流水线。
原因2中本该程序员动脑思考的问题交给了CUDA
解决,增加了时间开销,所以CUDA非常方便的统一内存模型缺点是计算速度慢。针对原因2,我们可以继续优化这个程序,告知GPU哪些数据需要拷贝到设备,哪些需要拷贝回主机。
from numba import cuda # 从numba调用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!")
if __name__ == "__main__":
main()
此时,GPU
速度就比CPU
快了。
常用内存分配函数:
cuda.device_array()
:在设备上分配一个空向量,类似于numpy.empty()
;cuda.to_device()
:将主机的数据拷贝到设备;cuda.copy_to_host()
:将设备的数据拷贝回主机;
总结
Python
的Numba
库可以调用CUDA
进行GPU
编程;
CPU
端被称为主机,GPU
端被称为设备;
运行在GPU
上的函数被称为核函数;调用核函数需要有执行配置,用来告诉CUDA以多大的并行力度来计算;
使用GPU
编程需要合理的将数据在主机和设备间互相拷贝。
上图展示了CUDA
编程的基本流程:
- 初始化,并将必要的数据从
CPU主存
拷贝到GPU显存
上; - 使用某个执行配置,以一定的并行粒度调用
CUDA核函数
; -
CPU
和GPU
异步计算; - 将
GPU
计算结果拷贝回主机。