GPU并行计算相关概念有哪些？

前言
并行计算是指什么？
GPU并行计算实际上是基于CPU+GPU的异构计算架构。在这种架构里，GPU和CPU借PCIe总线相连协同工作，CPU所在之处被称为主机端（host），GPU所在之处被叫做设备端（device），所示。

GPU含有较多运算核心（绿色），非常适合像大型矩阵运算这类数据并行的计算密集型任务。CPU运算核心（绿色）较少，不过能进行复杂逻辑运算（黄色），所以适合控制密集型任务。

类比来说，CPU如同头脑发达（可进行复杂逻辑运算，类似黄色部分占比大）但四肢简单（能实现的并行处理少，类似绿色部分占比小）；GPU则是头脑简单（逻辑运算简单，黄色占比小）而多肢发达（并行计算量大，绿色占比大）。
GPU有哪些性能指标？
GPU的核心数、内存容量、计算峰值（每秒单精度或双精度运算能力）以及内存带宽。
CUDA是何物？
英伟达公司推出了CUDA（Compute Unified Device Architecture），这是一个基于其GPU的通用高性能计算平台与编程模型。借助CUDA，能够充分发挥英伟达GPU强大的计算能力。
CUDAC为ANSI C的扩展，CUDA平台有驱动层接口（Driver API）以及运行时接口（Runtime API）。
CUDAC编程时只需编写顺序执行程序，代码里无需进行多线程处理。

每个内核（kernel）运行时产生的所有线程合起来被称作grid，grid包含多个block，而每个block又由多个线程thread组成，同一个block里的线程能够借助同步机制与块内共享内存来进行数据交互。

线程的标识。
在CUDA平台中，每个线程都能被blockIdx和threadIdx唯一确定。线程标识由平台内置并分配，在内核程序里可被访问。线程标识（blockIdx、threadIdx）为三维向量，能通过下标x、y、z来访问线程模型维度，而线程模型维度由内置变量blockDim和gridDim来标识。
利用GPU做并行计算时，threadIdx.x、blockIdx.x和blockDim.x这三个变量很重要。threadIdx.x能表明线程在自己所属线程块中的位置，blockIdx.x可标识线程块处于网格的何处，blockDim.x则规定了线程块的大小，它们的区别如下：

它表示的是线程在自身所属线程块（block）里的索引编号。在一个线程块当中，每个线程都有独一无二的threadIdx.x值，这个值能将不同线程区分开来。threadIdx.x的取值是从0到blockDim.x - 1，比如blockDim.x为32时，threadIdx.x就能够取0至31这些数值。

用于表示线程块在所属网格中的索引编号。网格包含多个线程块，每个线程块在网格里都有唯一的位置，这一位置通过blockIdx.x来标记。其取值范围由网格中的线程块总数决定，没有固定的上限数值，而是在程序执行时依据网格大小确定。

它规定了线程块在x维度的大小，也就是线程块所含的线程数量。比如，当blockDim.x为32时，该线程块在x方向就有32个线程。

grid里x方向的block数量。
类比：就像yolo里划分网格那样，又或是ViT为使用transformer而划分的patch。

在CV网络里，每个grid就如同输入图像。yolo中划分的网格或者ViT里的patch相当于每个block，而每个thread则类似像素点。
threadIdx.x就相当于patch里的第x个像素点，blockIdx.x则等同于第x个patch，而blockDim.x就相当于patch的宽度。这样一来，某个线程（也就是像素点）在全局的坐标便是：blockIdx.x×blockDim.x + threadIdx.x。同理，要是thread和block是多维的，那就是：blockIdx.y×blockDim.y + threadIdx.y。
GPU有多种用途各异的内存，分别为寄存器、全局内存、共享内存、线程私有本地内存、常量内存以及纹理内存。
，全局、常量、纹理内存与主机可互相访问，这是各内存的可访问关系。

寄存器
GPU里访问速度最快的内存中，内核代码里声明且无其他修饰符的自动变量一般存于寄存器。寄存器内存为每个线程私有，用来存放频繁访问的变量。
本地的内存。
在内核函数里，那些不能存进寄存器的变量会存于本地内存。像编译代码时无法确定下标的数组、大型数据结构、不符合内核寄存器限制条件的变量都是如此。
共享内存是一种什么情况？
或者
共享内存是怎么回事？
内核函数里被_shared_修饰的变量都存于共享内存。共享内存是片上存储空间，有低延迟、高带宽的特性。
常量的内存。
被_constant_修饰的变量存于常量内存，该内存可供所有内核代码访问。
全局的内存。
全局内存数量最多、用量最大、延迟最高。
静态分配用_device_关键字，动态分配在主机中使用内存管理函数。
CUDA的内存管理涵盖GPU内存的分配、释放，以及数据在主机和设备间的传输。标准C内存管理函数与CUDA的有如下对应关系。
nvcc是指什么？
nvcc从本质来讲是编译器驱动程序（compiler driver），其会依据传入的命令参数执行诸多命令工具，以完成程序编译各阶段的工作。
CUDA内核代码的编写可采用ISA（即CUDA指令集架构，也叫PTX），或者使用扩展的C语言来编写。
nvcc可把PTX或C语言编写的代码编译成可执行程序。
nvcc将CUDA程序复杂的编译、链接过程封装起来，让程序开发者能轻松完成CUDA程序编译。
nvcc的工作流程是怎样的？
CUDA程序在默认情况下采用全程序编译模式（whole program compilation mode）进行编译。
从源文件里分离出GPU相关的内核代码，编译成cubin或者PTX中间文件，然后保存到fatbinary中。
从源文件里分离出主机相关代码，用系统中能使用的编译器编译，再把fatbinary嵌入进去。
NVIDIA CUDA编译器驱动：官方文档。
GPU上的线程会并发执行内核函数（kernel function）。
定义内核函数时，使用__global__。
CUDA程序里的函数修饰符（function qualifier）。

内核函数仅能访问GPU内存，需返回void，不得使用变长参数、静态变量与函数指针，且具备异步性。
调用内核函数时，用<<<…>>>来指定线程配置以便执行。执行内核的每个线程都有唯一的ID，能通过内置变量在内核中访问。
例如，下面的示例代码运用内置变量threadIdx，把大小为N的向量A和B相加，结果存到向量C里：
如上述情况，存在1个block和N个threads，也就是定义一个block，其执行的线程维度为N。
比如，如下代码把两个NxN大小的矩阵A、B相加，结果存于矩阵C：

使用nvcc进行编译：

若提示如下错误：
按自己的版本，将相关路径添加到环境变量（需确保cl.exe在该路径内）。
错误：
错误是由Windows SDK缺失而导致的。
要是能找到windows kits的路径，就打开注册表regedit并搜索KitsRoot10，会得出两个结果。因为vs2019是32位的，所以会优先查找WOW6432Node下的Installed Roots路径，而2022则查找Microsoft的路径。于是，下面这两个路径都得修改：HKEY_LOCAL_MACHINE\SOFTWARE\WOW6432Node\Microsoft\Windows Kits\Installed Roots和HKEY_LOCAL_MACHINE\SOFTWARE\Microsoft\Windows Kits\Installed Roots。把这两个路径都改成windows kits的路径：....Windows Kits10。
若没有windows kits，可使用visual studio insteller重新安装，重新下载windows kits。
最后重新启动软件，便可使用nvcc进行编译了。
SP是最基本的处理单元，即streaming processor，也叫CUDA core。最终具体的指令与任务都由SP处理。GPU进行并行计算时，众多SP会同时处理，一个SP能够执行一个thread。
多个SP与其他一些资源组合起来就形成一个streaming multiprocessor，它也被称为GPU大核，其他资源包括warp scheduler、register、shared memory等。SM可被视作GPU的核心（类似CPU核心），register和shared memory属于SM的稀缺资源。CUDA会把这些资源分配给SM中的所有线程。所以，这些有限的资源对每个SM中的活跃warp有着非常严格的限制，进而限制了并行能力。
GPU的并行性是借助流处理器SM（streaming multiprocessor）达成的。
SM（流多处理器）包含以下结构：CUDA核、共享内存与L1缓存、注册文件、加载和存储单元、特殊函数单元以及Warps调度器。这些结构在SM中各自发挥着重要作用，共同协作以实现高效的计算任务处理等功能。
在整个GPU架构里，有众多的流式多处理器（Streaming Multiprocessor，SM），其在GPU中占据了大部分空间。每个SM具备自己的寄存器文件、共享内存与缓存等资源，也拥有许多Core资源，能够同时执行多个线程，SM可被视为GPU里可独立运算的单元。
线程束（warp）是最基本的执行单元。
在NVIDIA的GPU架构里，一个流多处理器（SM）正常能同时执行多个线程束（warp）（通常可同时处理32个warp），每个warp有32个线程，且这些线程只能执行同样的指令。
Warp Scheduler是什么？
杠精：有时候，借助warp的调度能够在某种程度上达成看似超过1024个线程的执行效果，不过这并非严格意义上同时执行超过1024个线程。
若一个SM最多只能存储1024个线程的信息，而一个SM可拥有超1024个线程，此时在SM内存中，warp线程的调度单元便是Warp Scheduler。
GPU是基于SIMT（单指令多线程）架构的，其SM（流多处理器）能迅速切换执行不同的warp。若某个warp因等待数据之类的情况而被阻塞，SM就可调度别的warp来执行，如此一来，在一段时间内可处理的线程数量，似乎就超出了SM理论上能同时执行的线程数量。
然而，在任一时刻，SM实际上依旧是同时执行多个warp，每个warp有32个线程，所以严格意义上讲，并非同时执行超1024个线程。某些新的GPU架构与技术或许会凭借更复杂的调度和资源管理方法来提升处理线程的能力，但也不会无限制地让超1024个线程同时执行。
由于资源有限，一个SM能同时并发的warp数量是有限的。SM要给每个线程块分配共享内存，还得给每个线程束中的线程分配单独的寄存器，所以SM的配置会对其支持的线程块和warp并发数量产生影响。
CUDA编程模型与GPU的硬件架构联系紧密，弄清楚二者的映射关系对编写高效CUDA程序有益。在CUDA编程模型里，会把问题分解成众多并行任务，这些任务由线程执行，且线程会以特定形式映射到GPU硬件上。
每个SM可同时执行的warp数上限，受硬件限制、kernel参数设置以及线程与线程块资源使用情况导致的实际数量限制影响。
若资源充足，且多于线程和线程块所使用的资源，此时每个流式多处理器（SM）执行的线程束（warp）数量会受内核（kernel）设置参数的限制。例如，每个线程块的线程数过少，由于SM同时执行的线程块数量有限，会使SM同时执行的线程数不足。
线程与线程块资源的使用会限制实际可执行的数量。要是线程块使用过多共享内存，例如一个线程块就占用了一半以上的共享内存，那么一个流式多处理器（SM）最多只能执行一个线程块。
要让一个SM能同时执行多个线程块，很明显，每个线程块只能使用每个SM总共享内存的几分之一。
假设存储资源足够，无需考虑内存大小。某GPU含2个SM，每个SM最多能执行32个warps，而每个warp有32个线程，这样总共就有1024个线程。
合理的设定：
第一种设置：grid的值为1，block的值为64。
第二种设置：grid是(2,2)，block为(16,16)。
第三种设置：grid设为8，block设为256。（输入数据充足，可实现100%的利用率）
4个block未必会被分配至同一个SM。从理论来讲，1个SM是能够处理这些block的，不过GPU调度器也许会把它们分配到多个SM，从而达成更优的负载均衡与资源利用。若GPU资源许可，有可能在1个SM上执行所有block；要是存在多个空闲的SM，GPU或许会把block分配到多个SM来加快执行速度。
卧龙：
设置：grid的值为40，block的值为20。
Grid的大小为40，这表明总共有40个block。每个block的大小是20个线程，那么总线程数就是800个。所需的warp数为40个（一个warp只能处理一个block中的线程）。而完成所有工作所需的SM数量为2个（因为每个SM能够支持32个warp）。从这个情况来看，一个SM最多可实现1024个线程，这里800个线程使用了两个SM，并且第二个SM还占用了8个warp，这8个warp仅控制了160个线程。
凤雏：
设置：grid的值为20，block的值设为36。
Grid规模：grid的值为20，这表示总共有20个block。Block规模：每个block包含36个线程。总线程数量：达到720个线程。所需warp数量：由于共有20个block，每个block需要2个warp，所以总共需要40个warp（多个warp能够同时执行同一个block中的线程）。所需SM数量：40个warp需要2个SM，因为每个SM只能同时支持32个warp。评价：一个SM最多可实现1024个线程，而这里720个线程就用了2个SM，并且第二个SM还占用了8个warp，其中20个warp仅控制80个线程，即每个warp控制4个线程。
GPU架构一般决定着每个block的线程数上限。在很多现代NVIDIA GPU里，block中的线程数量最多通常为1024。
比如，每个SM（流多处理器）最多能处理32个线程块、64个warp（也就是2048个线程），每个块可处理1024个线程。
当申请资源时若把线程块大小设为32，总共要申请2048个线程的话，那就需要64个这样的线程块。每个流多处理器（SM）一次只能处理32个块，即便一个SM完全有能力运行2048个线程，可受线程块的限制，每个SM只能运行32×32个线程，所以得用两个SM来运行，且每个只能占用50%的GPU资源。
每个SM也具备65536个寄存器。若要同时执行2048个线程，每个线程最多可有32个寄存器（计算方式为65536除以2048等于32）。要是在我们的程序里，一个内核中的每个线程需要64个寄存器，那么每个SM就只能运行1024个线程，这同样会造成50%的占用率。
一个kenel对应一个grid，得有足够的线程块才能充分发挥整个GPU全部SM的效能。一方面，一个SM需驻留多个线程块，那要将整个GPU几十上百个SM用满，所需的线程块数量得乘以一个较大的倍数才行。
这里举一个深度学习中reduce/layer_norm计算的实际例子：
若计算一个tensor最内层维度每行的reduce mean，若按朴素想法让每个线程计算一行，那总共的线程数仅有200。
这样只能生成一两个线程块，仅能使用一两个SM，性能显然非常差。要是用一个warp来计算一行的话，就会有200个warp，若一个线程块包含4个warp，那就会有50个线程块，这样就能用到大部分SM了。
当然也能用一个线程块计算一行，这样就会有200个线程块，SM的利用率也会更高。
PTX（并行线程执行）是CUDA平台针对基于GPU的通用计算所定义的虚拟机与指令集。
CUDA程序能以PTX相关指令来编写。与用C编写程序相比，PTX更接近底层，可把PTX当作CUDA平台的汇编语言。
CUDA程序若由C编写，会先转成PTX指令集，经优化后再变为特定GPU架构的指令集。通常，nvcc编译程序时，得指定目标虚拟机架构与真实GPU架构。
PTX不是物理GPU，而是对GPU计算能力与特性的抽象性定义。
，c开发的程序编译分两步，先转换为ptx，再得到cubin二进制文件，以供cuda执行。

PTX有何作用？
为GPU的演化提供稳固的指令集架构（ISA）支持。
让应用程序获得和GPU本地指令相当的性能。
为上层C程序开发提供不依赖GPU架构的ISA支持。
为PTX代码生成器与转换器进行源码级的优化。
程序结构步骤（基本适用于所有cuda程序）：
设置GPU设备，初始化矩阵，定义CUDA内核，分配GPU内存，传入数据计算后在主机获取结果。
用CUDA编程来达成两个向量的加法。在进行编程实现之前，先简要介绍下CUDA编程里的内存管理API。第一个是在设备（device）上分配内存的cudaMalloc函数：
此函数与C语言里的malloc相近，不过它是在device（设备）上申请指定字节大小的显存，devPtr为指向所分配内存的指针。若要释放已分配的内存，需使用cudaFree函数，该函数和C语言中的free函数相对应。还有一个重要函数是负责host（主机）和device之间数据通信的cudaMemcpy函数。
这里src代表数据源，dst为目标区域，count表示要复制的字节数量，kind则用于控制复制的方向，它有cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost和cudaMemcpyDeviceToDevice这几种情况，就像cudaMemcpyHostToDevice是把host（主机）上的数据拷贝到device（设备）上。
此处grid和block均设为1维，首先对kernel作如下定义：
stride代表整个grid的线程数。当向量元素数量很多时，可在每个线程实现多个元素（即元素总数除以线程总数）的加法，这就如同运用多个grid来处理，这是grid - stride loop方式。但下面例子中一个线程仅处理一个元素，所以kernel里的循环不执行。下面来具体实现向量加法。

在这里，向量大小为1<<20，block大小是256，如此一来grid大小为4096，kernel的线程层级结构所示：

通过nvprof工具能够分析kernel的运行状况，如下所示的结果表明，kernel函数耗时大概1.5ms。
对block的大小进行调整，比较不同配置下kernel的运行状况。我测试发现，block为128时，kernel耗时约1.6毫秒；block为512时，kernel耗时约1.7毫秒；block为64时，kernel耗时约2.3毫秒。可见，block并非越大越好，要进行适当选取。
统一内存的优化。
在CUDA里，可通过cudaMallocManaged函数来分配托管内存。
借助统一内存，上述程序可简化如下：
输入矩阵是A与B，需得到C = A×B。其实现思路为：每个线程计算C的一个元素值$C_{i,j}$，在进行矩阵运算时，应选用二维的grid和block。
首先对矩阵的结构体进行定义：
核函数可实现矩阵乘法，在此定义了两个辅助的__device__函数，分别用于获取矩阵元素值和给矩阵元素赋值，具体代码如下：
最后用统一内存编写矩阵相乘的测试实例。
此处矩阵的大小是，所设计线程的block大小为(32, 32)，这样的话grid大小即为(32, 32)。
参考