GPU计算的基本概念

好未来技术 2023-03-15

前言

GPU计算已广泛应用于深度学习领域，极大的促进了深度学习在各行业的应用，并反过来促进了计算型GPU的快速发展。本文希望通过串联一些GPU计算的基本概念，对GPU计算做一个直观和简单的介绍，达到了解GPU运行原理的目的。本文中所有图片均来自网络。

什么是GPU呢？首先说一下什么是GPU，GPU(Graphics Processing Unit)即图形处理器，又称显示核心，也就是显卡的核心，本来专门用于图形图像处理。从2003年开始，人们开始研究将GPU用于通用计算，并提出了GPGPU概念，前一个GP则表示通用目的(General Purpose)，所以GPGPU一般也被称为通用图形处理器或通用GPU。我们这里讨论通用计算型GPU。

什么时候用GPU呢？GPU为什么比CPU计算能力强大？下面这张图比较直观的显示了GPU的强大计算能力。CPU计算核心少，顶多数十个，GPU计算核心多，几千个计算核心很常见。但两者的计算核心有显著不同，CPU的计算核心有复杂而强大的控制单元Control，但是有较少的计算单元ALU；GPU相反，控制单元简单，但是计算单元ALU数量众多。可以简单的理解为，GPU可以完成很复杂的任务，计算仅是其“副业”；而GPU一般仅做简单的计算任务，但是由于计算单元众多，可把计算任务瓜分，人多力量大，很快把任务做完。GPU特别适合大量并行的数据运算（高运算密度）。

GPU强于计算，但并不能独自完成复杂的任务，所以GPU通用计算通常采用所谓CPU+GPU异构模式，由CPU负责执行复杂逻辑处理和事务处理，由GPU负责计算密集型的大规模数据并行计算。CPU（加主机内存）充当主机Host的角色，GPU（加显存）充当设备Device的角色。Host与Device之间可以PCI-E总线进行通信，这样它们之间可以进行数据拷贝，协同工作。

GPU硬件结构

计算部分

我们先来了解一下GPU的硬件结构。这里以NVIDIA 显卡为例。NVIDIA不同系列的GPU结构存在着一定的差别，随着GPU的发展，GPU的结构也在发生变化。这里，我们介绍一下GPU的一般结构。

从计算角度，GPU的最小的计算单元是SP(Streaming Processor)，也称为CUDA Core。众多的SP组成一个SM(Streaming Multiprocessor)，若干SM组成GPU（GPU架构示意图1）。有一些架构的GPU，将若干SM组成TPC（Texture Processing Clusters），若干TPC组成GPU(GPU架构示意图2)。无论怎么组织，一般从计算角度可以认为，GPU是SM的集合，GPU中的每一个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，所以在一个GPU上并发执行数千个线程是有可能的。

GPU架构示意图1

GPU架构示意图2

SP介绍（即CUDA Core）

具体的指令和任务都是在SP上处理的。SP包括控制单元，以及浮点计算单元、整数计算单元，另外还包括计算结果队列等，相当于微型CPU。可以类比CPU的计算核心，CPU计算核心上同一时间运行一个CPU线程，GPU计算核心SP运行一个GPU线程(thread)。GPU线程相比于CPU线程，轻量级的多，所谓轻量，是指维持线程运行的资源、线程状态等规模较小，线程切换的代价很低。好点的GPU，SP的数量有几千个，所以能在物理上并行执行的GPU线程数儿理论上也能到到几千个。

SM介绍

多个SP加上一些其他资源组成一个SM(见SM逻辑结构图)。

图中绿色的就是CUDA Core，也就是SP。每一个SM的其他资源包括：指令缓存(Instruction Cache)，L1缓存，共享内存、寄存器组(Register File)等组成的内存系统。其中，共享内存和寄存器由CUDA core共用，访问速度极快，但大小是K级别的，属稀缺资源，限制了实际能并发的线程数儿(比如，一个SM有8K个寄存器，一个SM最多有768个线程，如果要达到这个线程最大数，每个线程只能用8K/768=10个寄存器，如果每个线程要用11个寄存器，那么线程数就会减少)，这两种内存均可以在编程中使用。LD/ST 是load store unit，用来内存操作的。SFU是SPecial function unit，用来做特殊函数计算，硬件实现类似__cos()这样的函数。Warp调度器(Warp Scheduler)比较特殊，用于Wrap调度，下面介绍Warp的概念。

Warp介绍

CPU调度的单位是线程，与CPU不同，GPU执行程序时的调度单位是Wrap，Warp实际上是一个和硬件相关的概念，通常一个SM中的SP(thread执行者)会分成几个Warp(也就是SP在SM中是进行分组的，物理上进行的分组)。目前各种架构的GPU，Warp的大小均为32。Warp内的SP执行相同的指令（就是执行一份代码），处理不同的数据。如果没有这么多thread分配给Warp，那么这个Warp中的一些SP是不工作的。因为所有同一个Warp中的thread必须执行相同的指令，那么如果这些线程在遇到控制流语句时，如果进入不同的分支，那么同一时刻除了正在执行的分支之外，其余分支都被阻塞了(对应的SP未处于活跃状态，等待轮到自己执行)，十分影响性能，这类问题就是Warp divergence。

内存部分

上面的介绍，是从计算角度分析的GPU的逻辑结构。其中介绍SM的时候，提到了其内存系统，在SM之外，还有全局内存，我们可以从内存层面看一下GPU。GPU的内存(也就是显存)，与编程有关的主要有：全局内存(global memory)、共享内存(shared memory)、寄存器内存(Registers)。其中global memory是我们通常理解的显存，通常比较大，比如有的显卡显存是4G或16G，主机Host在执行运算前，会将要处理的数据从主机内存拷贝到global memory，global memory对全部的SP计算单元可见，SP可以从global memory取需要计算的数据，计算完成，也可以把计算结果写回global memory，然后主机再将global memory中的结果拷贝回主机内存。共享内存shared memory属于SM，在SM内可见，SM中的SP可以借助shared memory完成高速、低延迟的交互。寄存器内存(Registers)类似CPU中的通用寄存器，被SP使用。L1和L2缓存，对开发者不可见（不能编程使用），是GPU自己内存调度时使用。

不同的内存，大小差异很大，越接近SP，访问速度越快（是数量级的快）。但越快的内存越小，越稀缺。控制好对稀缺内存的占用，能提高运算的并发度（如果一个计算线程需要占用很多这种稀缺内存，那么实际能并行执行的线程数儿就会大大减少）。下图是一个GPU带宽和延迟的一份直观数据。寄存器内存的带宽可以达到TB/s级，延迟仅需一个执行周期，而全局内存要低得多。

GPU计算的软件概念

粗略介绍完GPU的硬件结构，下面介绍一下计算任务如何分配到GPU硬件上。在程序中，一个计算任务可能需要十万、甚至百万级的GPU线程，GPU的SP虽然多，但感觉不够用啊。实际应用时，也应该尽量减少硬件差异对编程的影响。这样，在软件层面，计算任务的众多线程会被合理划分，也采用两级结构：Grid、Block。划分的依据是任务特性和GPU本身的硬件特性。

一个计算任务，就是一个Grid；一个Grid分为多个Block；一个Block包含多个Thread。下图中，Kernel就是Host分配的一个计算任务，对应软件层面的一个Grid，Grid划分成众多的Block，每个Block对应一组线程。

在任务运行时，可以简单的理解为：Thread运行于SP，Block运行于SM，Grid运行于GPU，如下图所示：

线程运行于SP上是显而易见的。Block也会被整体分配到一个SM上，GPU保证同一个Block的线程运行在同一个SM上，因为SM资源有限(寄存器、共享内存)，所以Block中的线程数量受到限制(比如我台式机GeForce GTX 1050 Ti显卡，一个Block最大线程数儿1024)，如果Block中的线程需要申请较多的寄存器或共享内存，那么一个Block实际可承载的线程数儿会更低，Block中线程数儿的典型值是128~512个。

这里可能有一个疑问：一个Grid运行在一个GPU上，如果Block的数量太多(最大数目一般为65535，是GPU的一个硬件参数)，大大超过了SM的数量（虽然一个SM有能力同时运行多个Block），或者说线程数儿远超SP数量，会怎样？GPU会将Block分配到SM上，SM将Block划分成Warp执行运算过程。GPU对Block的处理，不一定是同时的，Block太多时，会对Block排队处理，直到把这个Grid中的所有Block执行完毕，这次计算任务才算完成。显然，GPU包含的SM越多，处理一个计算任务的总耗时就越短，速度就越快。如下图所示。

Grid和Block的索引：Grid-id和Block-id都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此Grid-id和Block-id可以灵活地定义为1-dim (x)，2-dim (x,y)以及3-dim (x,y,z)结构。通过dim3变量，在Grid中定个位某个Block，在Block中定位某个线程，比如三维(Dx, Dy, Dz)块, 索引为(x, y, z) 对应ID为 (x + y Dx+ z Dx Dy)。实际应用中，可以灵活的利用dim3结构，将线程进行一个合理的划分，每个线程在执行时，能得到自己的索引，根据自己的索引，去定位并处理对应的任务数据，完成自己的计算任务。如何合理的将计算任务分配，是编程者的责任，也是能将GPU最大化利用的关键。

CUDA架构

CUDA简介

具体使用上，得通过CUDA。首先介绍CUDA(Compute Unified Device Architecture)，CUDA是一种将GPU作为数据并行计算设备的软硬件体系，硬件上NVIDIA GeForce 8系列以后的GPU(包括GeForce、ION、Quadro、Tesla系列)已经采用支持CUDA的架构，软件开发包上CUDA也已经发展到CUDA Toolkit 10，并且支持Windows、Linux、MacOS三种主流操作系统。CUDA采用比较容易掌握的类C语言进行开发。CUDA的架构如下图所示。

如上图所示，CUDA软件架构由：CUDA硬件驱动API，CUDA Runtime API，CUDA软件库和GPU应用四层组成。GPU显卡，作为一个设备，必然有设备驱动，CUDA提供了硬件驱动的API，同时也提供了运行时API。相比之下，驱动程序API更难编程，但提供了对CUDA使用方式的更多控制。程序员必须直接处理初始化，模块加载等。显然，可以通过驱动程序API查询比通过运行时API更详细的设备信息。例如，只能通过驱动程序API查询设备上可用的空闲内存。一种更高级的API，称为CUDA运行时API，它是基于CUDA驱动程序实现的API。CUDA硬件驱动API和CUDA Runtime API是互斥的：应用程序只能使用其中的一种。基于CUDA运行时API，开发出一系列通用软件库，比如：如：cuSPARSE线性代数库，cuBLAS是CUDA标准的线代库，cuFFT傅里叶变换库，cuRAND随机数库等。再上层就是我们的CUDA应用程序了。看上去我们离GPU很远，但也可以很近，我们可以直接在应用程序中调用CUDA硬件驱动API操作GPU硬件。

CUDA编程模式

在 CUDA 标准编程模式中，同一份源文件中，即包含主机端运行的代码，又包含device端(GPU端)执行的代码，为了标明哪些代码需要在device上运行，定义了一些函数声明关键字：

用__global__关键字声明函数，表明此函数在 CPU 上调用，在 GPU 中执行。这也是大名鼎鼎的 “核函数”(前文提到的kernel)，其函数代码实现一个计算任务。

用__device__关键字声明函数，表明此函数只能在 GPU 中被调用，在GPU中执行。显然，这类函数只能被__global__ 类型函数或 device 类型函数调用。

这样的源文件，一般后缀名为.cu，不能用通常的C/C++编译器编译，而是使用nvcc编译器编译，自动编译出在主机端运行的指令和在GPU端执行的指令。

程序执行时，代码在主机(CPU执行主机端代码)和Device(GPU执行GPU端代码)端交替执行，如下图所示。

CUDA程序的典型执行过程如下：

1.定义一个需要在 device 端执行的核函数。( 函数声明前加 golbal 关键字 )

2.主程序调用cudaMalloc 函数，在显存中为待运算的数据以及需要存放结果的变量开辟显存空间(global memory)。

3.在主程序中，调用cudaMemcpy等函数，将待运算的数据传输进显存(global memory)。

4.主程序中调用核函数，同时指定需要创建的线程数，示例如下：
dim3 Grid(3, 2);
dim3 Block(5, 3);
kernel_fun<<< Grid, Block >>>(prams…);
此时，启动了核函数在GPU中的运行，主程序不用等待，直接返回。

一个核函数就是一个计算任务，也就是前文说的Grid。

Grid执行时，GPU线程根据自己的位置索引，从global memory取分配给自己的任务数据，按照kernel_fun完成处理任务，任务执行中，可以使用SM的shared memory与同一Block中的其它线程交互数据(kernel_fun中的代码可以分配和使用共享内存)，最后将处理结果放回global memory。

当然，这里可以调用多个核函数，下发多个计算任务，但任务是顺序执行的。

5.主程序调用cudaMemcpy等函数从显存中获取计算结果。cudaMemcpy是隐含的同步调用，会等4中的计算任务完成后返回。当然还有异步版函数cudaMemcpyAsync，调用cudaMemcpyAsync后，主程序不会同步等待，拷贝内存的操作被当成任务队列里的一项任务。这种情况下，就可能需要CUDA提供的一些类似event的同步方法来进行程序流程控制了，这里不详述。

6.主程序cudaFree释放申请的显存空间，显存空间像堆内存一样，必须显示释放。

上面的操作类似一个任务队列：数据拷贝、核函数执行、核函数执行、数据拷贝、核函数执行…

在GPU上有一个Stream流的概念：流（Stream）是一系列顺序执行的命令，流之间相对无序或并发的执行他们的命令。上面主程序调用的函数都有一个参数，指定属于哪个Stream，同一个Stream中的任务顺序执行。程序可以创建多个Stream，将不相关的任务分配到不同的Stream中运行，最大化利用GPU的计算能力。现在常见的GPU一般都支持多Stream并行运行。

后记

以上是对GPU计算很简单的介绍，主要是想把GPU计算的原理说清楚，并不纠结于细节。但实际使用中，有太多的细节和技巧需要掌握，才能真正了解和运用GPU计算。共同努力。

end

作者简介