通过Numba调用CUDA用GPU为Python加速：进阶理解网格跨步、多流、共享内存

前导知识

理解本文需要先了解：
1. 计算机底层基础知识，CPU、机器码、编译等《编译型语言与解释型语言如何在计算机底层运行》《计算机底层运转机制：多核、缓存、CPU、CU、ALU、Cache》
2. Python代码与GPU加速的关系《Python程序如何用GPU加速：Tesla、CUDA、Numba》
3. 在CPU入门numba《Python代码在CPU下加速：Numba入门》
4. 在GPU入门numba《Python通过Numba实现GPU加速》
CUDA优化方向

CPU + GPU 是一种异构计算的组合，各有独立的内存，GPU的优势是更多的计算核心。该架构在并行计算上有很大优势，但是数据需要从主机和设备间相互拷贝，会造成一定的延迟。因此，要从下面两个方面来优化GPU程序：
- 充分利用GPU的多核心，最大化并行执行度
- 优化内存使用，最大化数据吞吐量，减少不必要的数据拷贝
哪个方向有更大收益，最终还是要看具体的计算场景。

英伟达提供了非常强大的性能分析器nvprof和可视化版nvvp，使用性能分析器能监控到当前程序的瓶颈。分析器只支持C/C++编译后的可执行文件，Python Numba目前应该不支持。
并行计算优化

CUDA的执行配置：[gridDim, blockDim]中的blockDim最大只能是1024，gridDim最大为一个32位整数的最大值（2,147,483,648）,即20多亿。
网格跨度

如果对并行计算的维度有更高需求，就需要网格跨度。

以[2,4]的执行配置为例，该执行配置中整个grid只能并行执行8个线程，如上图，如果我们需要并行计算的数据是32，会发现8-32号共计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()
```
跨步大小：网格(grid)中线程(thread)总数=gridDim.x * blockDim.x，即for循环中的step，也是网格跨步名称的由来。

举例，如果网格总线程数(跨步大小)为1024，那么0号线程将计算第0、1024、2048…号的数据。

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

前面讨论的并行，都是线程级别的：即CUDA开启多个线程，并行执行核函数内的代码。

GPU最多上千个核心，同一时间只能并行执行上千个任务，当处理千万级别的大数据时，整个任务无法被GPU一次执行，所有的计算任务需要放在一个队列中，排队顺序执行。

CUDA将放入队列顺序执行的一系列操作称为流（Stream）。

由于异构计算的硬件特性，CUDA中以下操作是相互独立的，通过编程，是可以操作他们并发地执行的：
- 主机端上的计算
- 设备端的计算（核函数）
- 数据从主机和设备间相互拷贝
- 数据从设备内拷贝或转移
- 数据从多个GPU设备间拷贝或转移
针对这种互相独立的硬件架构（上图：数据拷贝和计算的重叠），CUDA使用多流作为一种高并发的方案：把一个大任务中的上述几部分拆分开，放到多个流中，每次只对一部分数据进行拷贝、计算和回写，并把这个流程做成流水线。

因为数据拷贝不占用计算资源，计算不占用数据拷贝的总线（Bus）资源，因此计算和数据拷贝完全可以并发执行。

如下图(默认流与多流)所示，将数据拷贝和函数计算重叠起来的，形成流水线，能获得非常大的性能提升。实际上，流水线作业的思想被广泛应用于CPU和GPU等计算机芯片设计上，以加速程序。

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

默认情况下，CUDA使用0号流，又称默认流。不使用多流时，所有任务都在默认流中顺序执行，效率较低。在使用多流之前，必须先了解多流的一些规则：
- 给定流内的所有操作会按序执行。
- 非默认流之间的不同操作，无法保证其执行顺序。
- 所有非默认流执行完后，才能执行默认流；默认流执行完后，才能执行其他非默认流。
参照上图，可将这三个规则解释为：
- 非默认流1中，根据进流的先后顺序，核函数1和2是顺序执行的。
- 无法保证核函数2与核函数4的执行先后顺序，因为他们在不同的流中。他们执行的开始时间依赖于该流中前一个操作结束时间，例如核函数2的开始依赖于核函数1的结束，与核函数3、4完全不相关。
- 默认流有阻塞的作用。如图中红线所示，如果调用默认流，那么默认流会等非默认流都执行完才能执行；同样，默认流执行完，才能再次执行其他非默认流。
某个流内的操作是顺序的，非默认流之间是异步的，默认流有阻塞作用。

多流的使用

# 如果想使用多流时，必须先定义流
stream = numba.cuda.stream()
# 核函数调用的地方除了要写清执行配置，还要加一项stream参数
kernel[blocks_per_grid, threads_per_block, stream=0]
# CUDA的数据拷贝以及核函数都有专门的stream参数来接收流，以告知该操作放入哪个流中执行
numba.cuda.to_device(obj, stream=0, copy=True, to=None)
numba.cuda.copy_to_host(self, ary=None, 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.xif idx < n :result[idx] = a[idx] + b[idx]def main():n = 20000000x = np.arange(n).astype(np.int32)y = 2 * xstart = time()x_device = cuda.to_device(x)y_device = cuda.to_device(y)out_device = cuda.device_array(n)threads_per_block = 1024blocks_per_grid = math.ceil(n / threads_per_block)# 使用默认流gpu_add[blocks_per_grid, threads_per_block](x_device, y_device, out_device, n)gpu_result = out_device.copy_to_host()cuda.synchronize()print("gpu vector add time " + str(time() - start))start = time()# 使用5个streamnumber_of_streams = 5# 每个stream处理的数据量为原来的 1/5# 符号//得到一个整数结果segment_size = n // number_of_streams# 创建5个cuda streamstream_list = list()for i in range (0, number_of_streams):stream = cuda.stream()stream_list.append(stream)threads_per_block = 1024# 每个stream的处理的数据变为原来的1/5blocks_per_grid = math.ceil(segment_size / threads_per_block)streams_out_device = cuda.device_array(segment_size)streams_gpu_result = np.empty(n)# 启动多个streamfor i in range(0, number_of_streams):# 传入不同的参数，让函数在不同的流执行x_i_device = cuda.to_device(x[i * segment_size : (i + 1) * segment_size], stream=stream_list[i])y_i_device = cuda.to_device(y[i * segment_size : (i + 1) * segment_size], stream=stream_list[i])gpu_add[blocks_per_grid, threads_per_block, stream_list[i]](x_i_device, y_i_device, streams_out_device,segment_size)streams_gpu_result[i * segment_size : (i + 1) * segment_size] = streams_out_device.copy_to_host(stream=stream_list[i])cuda.synchronize()print("gpu streams vector add time " + str(time() - start))if __name__ == "__main__":main()

是否使用多流的计算时间差距非常大：

gpu vector add time 9.33862018585205
gpu streams vector add time 1.4097239971160889

在上面的程序中，将向量分拆成了5份，同时也创建了5个流，每个流执行1/5的“拷贝、计算、回写”操作，多个流之间异步执行，最终得到非常大的性能提升。

内存优化

CPU和GPU组成异构计算架构，如果想从内存上优化程序：
1. 必须尽量减少主机与设备间的数据拷贝；
2. 将更多计算从主机端转移到设备端；
3. 尽量在设备端初始化数据，并计算中间数据；
4. 尽量不做无意义的数据回写。
GPU的内存结构如上图所示：
- GPU的计算核心都在Streaming Multiprocessor（SM）上；
- Multiprocessor里有计算核心可直接访问的寄存器（Register）和共享内存（Shared Memory）；
- 多个SM可以读取显卡上的显存，包括全局内存（Global Memory）；
- 每个Multiprocessor上的Shared Memory相当于该Multiprocessor上的一个缓存，一般都很小，当前最强的GPU Telsa V100的Shared Memory也只有96KB。
Shared Memory和Global Memory的字面上都有共享的意思，但是不要将两者的概念混淆：
- Shared Memory离计算核心更近，延迟很低；
- Global Memory是整个显卡上的全局内存，延迟高。
更详细的划分如上图（英伟达GPU存储结构），CUDA的线程可以访问不同级别的存储：
- 每个Thread有独立的私有内存；
- 每个Block中多个Thread都可以在该Block的Shared Memory中读写数据；内存优化一般在这个层面；
- 整个Grid中所有Thread都可以读写Global Memory；
二维和三维执行配置

之前使用的threadIdx 和blockIdx变量都是一维的，实际上，CUDA允许这两个变量最多为三维，一维、二维和三维的大小配置可以适应向量、矩阵和张量等不同的场景。
Shared Memory内存优化

多维配置及内存优化的加速过程，是将数据从Global Memory拷贝到Shared Memory上，后续多次重复利用Shared Memory上的数据。

具体实现，后续有具体需求时再来深入研究。
总结

从两个方向：
1. 增大并行度：网格跨步、多流
2. 充分利用内存：Shared Memory内存优化
来对CUDA进行优化。
References

CUDA 官方推荐教程：C Programming Guide
GPU加速03:多流和共享内存—让你的CUDA程序如虎添翼的优化技术！

通过Numba调用CUDA用GPU为Python加速：进阶理解网格跨步、多流、共享内存相关推荐

CUDA 网格级并发-流(1)
CUDA中的程序并发可以分为两种内核级并发网格级并发内核级并发是开发程序中经常使用到的,即通过划分block和thread实现同一个内核在GPU上同时并发,将同一个核分别部署到不同的SP上进行同 ...
Android系统匿名共享内存（Anonymous Shared Memory）C++调用接口分析
出自:http://blog.csdn.net/luoshengyang/article/details/6939890 在Android系统中,针对移动设备内存空间有限的特点,提供了一种在进程间共享 ...
Android调用C++实现共享内存（Native层）
Android下匿名共享内存java层接口利用MemoryFile实现进程间内存共享:利用MemoryFile可以参考这篇文章:https://blog.csdn.net/qq_24451593/ar ...
CUDA 多GPU调用实现
当设备存在多块GPU时,为了高效利用GPU,我们常常需要使用多卡计算.本例中我们使用OpenMP来进行多线程调用多GPU运行,初学者无须详细了解OpenMP,只需知道一两句命令就行. 详细步骤如下: ...
基于CUDA的GPU并行计算技术实现网课课表编排
这篇文章是用来填这个坑的:https://blog.csdn.net/xinew4712/article/details/108276264 上篇文末设想的是用天灾和定向改造机制来提高排课运算的效率, ...
win10 1050显卡torch调用不到GPU问题处理
目录先看我装的环境: pytorch torchvision python的版本对应: 使用conda命令指向清华源下载并总结经验先看我装的环境: 测试: import torch print( ...
3维线程格 gpu_基于CUDA的GPU并行优化重力三维反演
重力勘探由于其成本较低.施工方法方便等, 被广泛应用于大尺度的地质异常体勘查.大范围找矿普查.以及小比例尺密度三维地质建模等工作中.目前常用的反演方法有两种, 2.5维联合3维界面反演[和三维物性反演 ...
如何在MFC中调用CUDA
如何在MFC中调用CUDA 有时候,我们需要在比较大的项目中调用CUDA,这就涉及到MFC+CUDA的环境配置问题,以矩阵相乘为例,在MFC中调用CUDA程序.我们参考罗振东iylzd@163.com ...
2021.08.24学习内容torch.utils.data.DataLoader以及CUDA与GPU的关系
pytorch数据加载: ①totchvision 的包,含有支持加载类似Imagenet,CIFAR10,MNIST 等公共数据集的数据加载模块 torchvision.datasets impor ...
Qt 调用CUDA静态库和动态库生成与配置
前言通过将CUDA相关计算操作放在库中,方便在项目中调用,省去了每次编译cu文件的麻烦,也便于集成到其他平台上. 关于部署CUDA加速的程序时,往往对CUDA加速的程序编译为动态链接库或者静态链接库 ...

通过Numba调用CUDA用GPU为Python加速：进阶理解网格跨步、多流、共享内存

前导知识

CUDA优化方向

并行计算优化

网格跨度

多流

主机端上的计算

设备端的计算（核函数）

数据从主机和设备间相互拷贝

数据从设备内拷贝或转移

数据从多个`GPU`设备间拷贝或转移

多流的规则

给定流内的所有操作会按序执行。

非默认流之间的不同操作，无法保证其执行顺序。

所有非默认流执行完后，才能执行默认流；默认流执行完后，才能执行其他非默认流。

多流的使用

内存优化

二维和三维执行配置

Shared Memory内存优化

总结

增大并行度：网格跨步、多流

充分利用内存：`Shared Memory`内存优化

References

通过Numba调用CUDA用GPU为Python加速：进阶理解网格跨步、多流、共享内存相关推荐

最新文章

热门文章

通过Numba调用CUDA用GPU为Python加速：进阶理解网格跨步、多流、共享内存

前导知识

CUDA优化方向

并行计算优化

网格跨度

多流

主机端上的计算

设备端的计算（核函数）

数据从主机和设备间相互拷贝

数据从设备内拷贝或转移

数据从多个GPU设备间拷贝或转移

多流的规则

给定流内的所有操作会按序执行。

非默认流之间的不同操作，无法保证其执行顺序。

所有非默认流执行完后，才能执行默认流；默认流执行完后，才能执行其他非默认流。

多流的使用

内存优化

二维和三维执行配置

Shared Memory内存优化

总结

增大并行度：网格跨步、多流

充分利用内存：Shared Memory内存优化

References

通过Numba调用CUDA用GPU为Python加速：进阶理解网格跨步、多流、共享内存相关推荐

最新文章

热门文章

数据从多个`GPU`设备间拷贝或转移

充分利用内存：`Shared Memory`内存优化