CUDA介绍

CUDA介绍GPU介绍GPU(Graphics Processing Unit)最初是作为3D图像设计的专用处理器其作为固定功能的硬件用于加速实时3D渲染中的并行运算目前来说GPU逐渐可以通过编程来并行的针对3D场景或图像各个组件。而在之后NVIDIA退出了一套统一的计算设备架构即CUDA其使得任何计算任务都独立于图形API充分利用GPU的吞吐能力。简单来说GPU就是显卡中的计算芯片它本身便是硬件里面含有大量适合并行计算的核心而CUDA则是NVIDIA推出的专门针对GPU并行计算的平台。GPU优势与同等价格和功耗范围内的CPU相比GPU在指令吞吐量和内存带宽方面具有显著优势。其他计算设备如FPGA虽然也非常节能但编程灵活性远低于GPU。GPU和CPU的设计目标不同。CPU旨在尽可能快速地执行一系列连续的操作并能并行处理几十个这样的线程而GPU则专注于并行执行数千个线程通过牺牲单线程性能来实现更高的整体吞吐量。GPU专为高度并行的计算而设计将更多晶体管用于数据处理单元而CPU则将更多晶体管用于数据缓存和流量控制。编程模型异构系统异构系统就是指一个系统里面同时又不同类型的计算单元并可以让这些不同的计算单元处理各自擅长的事情如电脑中的CPU、GPU以及之前所使用得到RK3588的CPU、GPU、NPU。拿平时的电脑来说CPU及其直接连接的内存分别称为主机和主机内存。GPU及其直接连接的内存则分别称为设备和设备内存。在某些系统上这些组件可能集成在同一芯片上因此这样的电脑一般叫做核显而大多数配置更好更高的电脑都是将这些组件分开的甚至可能含有多个CPU、GPU。异构编译内核通常来说应用程序都是从CPU开始运行的但是它的部分代码会在GPU上执行运行在CPU上的主机代码可以通过CUDA API在主机内存与设备内存之间复制数据、启动GPU上的代码执行。应用程序在GPU上执行的代码称为设备代码而被调用在GPU上执行的函数则被称为内核。启动内核的过程称为启动内核具体细节后续说明。GPU硬件模型在CUDA编程中GPU通常可视为多组流式多处理器(Streaming Multiprocessors)--SMs这些处理器都含有一个本地寄存器(local register file)、统一的数据缓存(data cache)以及若干执行计算的功能单元。统一数据缓存为共享内存和一级缓存提供了物理资源。统一数据缓存对一级缓存和共享内存的分配可在运行时进行配置。不同类型的内存大小以及每个处理器内部功能单元的数量可能因GPU架构而异。线程块和网格当应用程序启动内核时会同时创建大量线程通常可达数百万个。这些线程被组织成多个线程块。不出所料一个线程被称为“线程块”。所有线程块按一定方式组成网格且每个网格中的线程块大小和维度相同。当启动一个内核时会使用特定的执行配置来指定网格和线程块的维度。该执行配置还可包含可选参数例如线程块簇(thread cluster)、流(stream)以及SM配置设置这些内容将在后续章节中介绍。通过内置变量执行内核的每个线程都可以确定其在所属线程块中的位置以及该线程块在包含网格中的位置。线程还可利用这些内置变量来确定线程块的维度以及内核被启动所用的网格。这使得每个线程在运行内核的所有线程中拥有唯一的标识符。该标识符常用于确定某个线程负责哪些数据或操作。线程块中的所有线程都在同一个SM中执行这使得线程块内的线程能够高效地相互通信和同步。线程块内的所有线程均可访问片上共享内存可用于线程块内各线程之间的信息交换。一个网格可能包含数百万个线程块而执行该网格的GPU通常仅有几十到几百个SM。线程块中的所有线程均由单个流处理器执行并且在大多数情况下这些线程会在该流处理器上运行至完成。由于线程块之间无法保证调度顺序因此一个线程块不能依赖其他线程块的结果因为其他线程块可能要等到当前线程块完成才能被调度。CUDA编程模型支持在任意大小的GPU上运行任意规模的网格。因此CUDA编程模型要求不同线程块中的线程之间不存在数据依赖关系即同一个线程块内的所有线程同时在同一个SM上运行。CUDA编程模型要求能够以任意顺序、并行或串行地执行线程块。线程块簇线程块簇是比线程块更高一层的组织方式在compute capability 9.0后CUDA增加了一种可选的层级也就是线程块簇。块簇是由多个线程块组成的集合与线程块和网格一样可以沿一维、二维或三维进行布局。通过指定簇组将相邻的线程块组合成簇并在簇层级上提供额外的同步和通信机会。具体而言簇中的所有线程块都在同一个GPU核心内执行。由于线程块同时且在同一GPC内被调度因此同一簇内不同线程块中的线程可以通过协作组提供的软件接口相互通信和同步。簇内的线程可以访问该簇中所有线程块共享的内存这种内存被称为分布式共享内存(distributed shared memory)。也就是可以跨SM共享数据。簇的最大大小取决于硬件不同设备之间存在差异。传统CUDAThread-Warp-Thread Block-GridSCompute Capability 9.0开始Thread-Warp-Thread Block-Thread Block Cluster-GridIMT和SIMTWarp是GPU中线程实际被调度执行的基本单位而SIMT则是其执行机构在一条线程块中32个线程被组织为Warp。每个Warp以单指令多线程SIMT方式执行内核代码。在SIMT架构中同一个Warp中的所有线程都执行相同的内核代码但每条线程可能在代码中选择不同的分支路径。当一个Warp执行线程时这些线程会被分配到一个Warp lane上。Warp lane编号为0至31线程块中的线程会按照硬件多线程机制中详细说明的可预测方式被分配到各个战列中。Warp中的所有线程同时执行同一条指令。如果Warp中某些线程在执行过程中遵循控制流分支而其他线程则不遵循该分支那么不遵循分支的线程将被屏蔽而遵循分支的线程则会被正常执行。这就是Warp Divergence。if (threadIdx.x 16) { A(); } else { B(); }比如这个代码一个Warp有0~31编号的线程而这个而非if-else可能导致该Warp内部线程执行不同指令这就会出现Warp Divergence假如0~15执行A()代码16~31执行B()代码那么GPU就要分别处理两条执行路径而0~15在执行A()代码时16~31就无法进行工作反之16~31在执行B()代码时0~15也无法进行工作这就导致了效率的下降。在SIMT模型中 Warp中的所有线程以锁步方式依次执行内核代码。硬件执行可能有所不同。因此Warp执行时线程块最好指定其总线数为32的倍数。原因如上。同时不建议利用Warp执行在实际硬件上的映射知识进行优化。因为CUDA编程模型和SIMT指出Warp中的所有线程会同时执行代码。只要遵循编程模型硬件可以以对程序透明的方式优化掩码的lane。如果程序违反了该模型则可能导致未定义行为且不同GPU硬件上表现可能不同。Tile Programming除了上面SIMT编程模式CUDA还支持一种Tile Programming。在这种编程模式中程序员以整个线程块来编写代码。Tile Programming的内核在由块构成的网格上启动每块执行自己对应的内核并可查询其在网格中的位置以确定它负责的数据区域而程序员只需要指定网格的维度即可。简单来说SIMT就是以线程为基础单位去编程而Tile Programming则是以线程块为基础单位去编程对Tile数据进行操作再由编译器把这些操作映射到块内线程。支持标准的控制流结构如条件判断和循环但由于块遵循单一的控制流因此不存在Warp Divergence上述所提到的SIMT中的错误概念。标量操作由块中的单个线程执行例如计算索引或循环边界而Tile操作例如逐元素相加两个Tile由块中所有线程并行执行。需要注意的是不要将执行单元“块”与数据单元“Tile”混淆。一个块可以创建并操作多个不同形状和数据类型的Tile。Arrays和TilesTile内核处理两中类型的数据Array和Tile。Array是存储在设备内存中的多维元素容器Array是可变的其内容可通过内核内的存储操作进行修改因此Array也具有形状和数据类型。Tile则是一种多维的数值集合其仅存在于Tile编程模型中并且其属于单个块的局部范围类似于SIMT中线程块于线程的关系。但是Tile是不可变的因此对Tile进行任何的操作都会导致生成一个新的Tile同时Tile在内存中由具体的表示其通常由编译器决定Tile数据如何存储如使用寄存器、共享内存或者其他SM(流式多处理器)资源等。同时Tile的每个维度必须是2的幂次方并且必须在编译时已知因为Tile不可变。Tile空间和数据移动数据通过加载和存储操作在Array和Tile之间进行传输这些操作都需要使用Tile空间(Tile Space)该空间是将Array在概念上划分为大小相等、不重叠的Tile。比如我现在有一个形状大小为[M,N]的二维Array如果我使用一个加载操作指定了一个形状大小为[m,n]的二维Tile这个Array就会被分为若干[M/m]行[N/n]列的二维Tile在这个空间进行(i,j)的索引时就是指定在这个Array中(i,j)位置的数值进行加载操作。而其将会返回一个[m,n]的Tile这个Tile中包含数组中对应位置的元素。而当某个Tile超出Array边界时加载操作会指定如何处理越界元素例如用零填充。存储操作则执行相反的操作给定一个Tile和该Tile空间中的索引它将Tile的元素写入Array的相应区域任何超出Array边界的数据写入操作都会被静默丢弃Tile程序还支持聚集和散列操作用于从Array中任意位置加载或存储数据。对Tile进行操作Tile程序提供了一组内置操作用于处理Tile包括逐元素算术运算、矩阵乘法、沿一个或多个维度进行的归约如求和或最大值、形状操作如重塑或转置以及类型转换。当两个不同形状的Tile在操作中合并时较小的Tile会自动扩展以匹配较大的Tile瓷砖类似于张量的广播然后再执行操作。与SIMT编程的关系Tile编程和SIMT编程在CUDA中并存一个应用程序可以同时包含SIMT和Tile内核且这两种内核均可对设备内存中的相同数据进行操作。编程模型的选择是针对每个内核单独做出的决定。Tile编程并不取代SIMT编程SIMT提供了对单个线程的细粒度控制。而Tile编程则提供了一种更高层次的抽象有助于简化内核开发由于线程级别的决策由编译器负责同一个tile内核可以在不同的GPU架构上运行而无需修改源代码。两种模型均基于相同的底层硬件——即前文所述的SM、线程块和网格。此外两种模型都使用相同的设备内存空间该部分将在下文介绍。GPU内存在现代计算系统中高效利用内存与充分发挥执行计算功能单元的作用同等重要。异构系统包含多个内存空间而GPU除了缓存外还具备多种类型的可编程片上内存。异构系统中的RAM内存GPU和CPU都直接连接有DRAM芯片。在配备多个GPU的系统中每个GPU都有独立的内存从设备代码的角度来看连接到GPU的DRAM被称为全局内存因为它可被GPU中的所有SM(流处理器)访问。DRAM芯片就是动态随机存取存储器芯片(Dynamic Random Access Memory)而GPU显卡本质上就是由一颗颗DRAM芯片组成的。我们常说的DDR4和DDR5就是DRAM芯片包括显卡常用的GDDR6、GDDR6X、GDDR7高端AI GPU还常用HBM、HBM2等系列。但是DRAM通常很慢因此在编程时通常需要尽量少的访问DRAM而多复用已经搬到更快存储中的数据。Register - Shared Memory - L1 Cache - L2 Cache - DRAM这个关系链容量从小到大但是延迟速度从快到慢。与CPU类似GPU也使用虚拟内存寻址。在当前支持的所有系统中CPU和GPU共享一个统一的虚拟内存空间。这意味着系统中每个GPU的虚拟内存地址范围都是唯一的并且与CPU以及系统中其他任何GPU的地址范围不同。因此对于给定的虚拟内存地址可以判断该地址位于GPU内存还是系统内存而在多GPU系统中还可以确定是哪个GPU的内存存储了该地址。CUDA API提供了分配GPU内存、CPU内存的功能以及在CPU和GPU之间、同一GPU内部或多GPU系统中进行内存复制的操作。当需要时可以显式控制数据的局部性。GPU中的片上存储器(On-Chip Memory)除了全局内存外每个GPU还拥有部分片上存储器。每个SM都具有独立的寄存器文件和共享内存。这些内存属于SM的一部分可由在SM内运行的线程极快地访问。片上存储器包含Register、Shared Memory和Cache也就是相较于上文中更加适合编程使用的存储器。寄存器文件用于存储线程局部变量通常由编译器分配。共享内存则可被线程块或线程簇内的所有线程之间访问可用于线程间的数据交换。SM中的寄存器文件和统一数据缓存具有有限的容量。各SM的寄存器文件大小、统一数据缓存大小以及如何配置统一数据缓存以实现L1缓存与共享内存之间的平衡。寄存器文件、共享内存空间和L1缓存是线程块中所有线程共享的资源。要将一个线程块调度到某个SM上必须满足每个线程所需寄存器总数乘以该线程块中的线程数小于或等于该SM可用的寄存器数量。如果线程块所需的寄存器数量超过了寄存器文件的大小内核将无法启动此时必须减少线程块中的线程数量以使线程块可启动。共享内存的分配是在线程块级别进行的。也就是说与按线程分配寄存器不同共享内存的分配是整个线程块共用的。Caches除了可编程内存外GPU还具有L1和L2缓存。每个SM都拥有一个L1缓存该缓存属于统一数据缓存的一部分。更大的L2缓存由GPU内的所有SM共享每个SM还拥有独立的常量缓存用于缓存在全局内存中被声明为在整个内核执行期间保持不变的值编译器也可以将内核参数存储到常量内存中。这可以通过将内核参数与L1数据缓存分开缓存在SM中以提升内核性能。统一内存当应用程序在GPU或CPU上显式分配内存时该内存仅对运行在该设备上的代码可访问。也就是说CPU内存只能由CPU代码访问GPU内存只能由运行在GPU内核的代码访问。但是通过使用CUDA API在CPU与GPU之间复制内存可以在正确的时间将数据显式地复制到目标内存位置。CUDA的一项名为Unified Memory(统一内存)的功能该功能使应用程序能够分配可被CPU和GPU同时访问的内存。CUDA运行时环境或底层硬件会在需要时自动提供访问权限或将数据迁移至正确的位置。即使使用了统一内存仍可通过将内存迁移次数降至最低并尽可能直接从与内存所在位置连接的处理器访问数据以实现最佳性能。简单来说统一内存就是由CUDA和硬件负责映射、迁移和管理的可以给CPU和GPU都提供统一的数据访问的模型但是统一内存不是一定就会更快因为数据可能发生频繁迁移因此CUDA还提供了cudaMemPrefetchAsync()函数来提前告诉CUDA和硬件这块数据马上要被谁使用这也是上面第二段中提到的自动提供的访问权限。CUDA平台NVIDIA CUDA 平台由众多软件和硬件组件以及多项关键技术构成旨在支持异构系统上的计算。计算机能力(CC)和SM版本每个CUDA平台都有一个Compute Capability(计算机能力简称CC)该能力表示GPU支持的功能并指定了GPU的硬件能力。CC以X.Y格式表示其中X为主版本号Y为次版本号。例如CC 12.0的主版本为12次版本为0。CC直接对应SM的版本号。例如CC 12.0的GPU中的SM版本为sm_120。CUDA工具包和NVIDIA驱动程序NVIDIA 驱动程序可被视为 GPU 的操作系统。它是一个必须安装在主机系统操作系统的软件组件并对于所有 GPU 使用场景包括显示和图形功能都是必需的。NVIDIA 驱动程序是 CUDA 平台的基础。除了 CUDA 之外NVIDIA 驱动程序还支持其他 GPU 使用方式例如 Vulkan 和 Direct3D。NVIDIA 驱动程序有如 r580 等版本号。CUDA 工具包是一组用于编写、构建和分析利用 GPU 计算的软件的库、头文件和工具的集合。CUDA 工具包是一个独立于 NVIDIA 驱动程序的软件产品。CUDA runtime是 CUDA 工具包提供的某个库的特例。CUDA runtime提供了 API 及部分语言扩展用于处理诸如内存分配、在不同 GPU 或 GPU 与 CPU 之间复制数据以及启动内核等常见任务。CUDA runtime运行时的 API 组件被称为 CUDA runtime API。总结来说NVIDIA 驱动程序就是一个驱动平台类似于串口而CUDA Toolkit则是开发和编译CUDA程序的工具包。CUDA Toolkit中有需要东西其中较为重要的便是CUDA Runtime它提供了需要供开发者使用的API称为如CUDA Runtime API如cudaMalloc()、cudaMemcpy()和cudaFree()等。而NVIDIA 驱动程序也有自己底层的API称为Driver API如cuInit()、cuMemAlloc()等。可以看出Runtime API和Driver API在前缀上都不相同Runtime API以cuda为前缀而Driver API以cu为前缀。并行线程指令集(PTX)在CUDA平台中有一个基础但是不常使用的虚拟指令集Parallel Thread Execution(并行线程指令集简称PTX)。其本质还是汇编语言为“面向GPU的汇编语言”。针对高级语言的领域特定语言和编译器可以将PTX代码生成为中间表示IR然后使用NVIDIA的离线或即时编译工具生成可执行的GPU二进制代码。这使得CUDA平台能够通过除NVIDIA提供的工具之外的语言进行编程。流程如下CUDA C-PTX-SASS-Cubin-Fatbin打包-NVIDIA Driver-GPU执行CUDA c和c之间的区别// CPU代码c int main() { ... } // GPU代码CUDA c __global__ void kernel() { ... }Cubins和FatbinsCUDA程序和库通常使用C等高级语言编写。然后PTX再被ptxas编译器编译为针对GPU的真实机器码即CUDA二进制文件简称Cubins。使用GPU计算的可执行文件和库二进制文件同时包含CPU和GPU代码GPU代码存储在一个名为Fatbin的容器中。Fatbins可以包含多个不同目标平台的cubins(CUDA二进制文件)和PTX(并行线程指令集)。这样一个应用程序就可以在不同的GPU框架下运行(GPU的不同主要是SM版本的不同)。通常来说NVIDIA都可以保证二进制文件的兼容性。只要CC的主版本号相同那么更新的此版本可以向下兼容。同时PTX文件也有一定的兼容性。当GPU中的代码被编译成以某个cc为目标的PTX后只要后续GPU中cc的版本等于或高于编译时的版本这个代码就可以被实时编译。实时编译(JIT)应用程序在运行时加载PTX代码并由设备驱动程序编译为二进制代码这个过程称为Just-in-time Compilation(实时编译简称JIT)。实时编译会增加应用程序首次加载的时间但是其可以利用新版驱动中的编译器进行优化同时由于PTX的兼容性因此它可以更好的适应许多更新版本甚至还未发布的GPU型号。其只会增加首次加载的时间这是因为实时编译在编译后会自动缓存生成的二进制代码副本这个缓存被称为计算缓存但是其在设备驱动程序升级时会自动失效依次来获得新设备驱动程序中的编译器带来的优化。二进制最终化Binary Finalizsation(二进制最终化)指CUDA驱动在程序运行/加载时把一个还需要针对当前GPU做最后处理的Cubin最终整理为当前GPU可以真正执行的最终机器代码。二进制最终化和前面的PTX的实时编译(JIT)十分像其目的都是为了更好的兼容但是前者的输入已经是二进制代码只需要针对实际GPU做最后的最终化处理后者则仍然还是PTX虚拟指令还需要再进一步编译为当前GPU的Cubin因此运行上比实时编译快很多。