1. 从“串行”到“并行”为什么我们需要CUDA如果你写过代码尤其是处理过一些计算量大的任务比如图像处理、科学模拟或者机器学习训练那你一定对“程序跑得太慢”这件事深有体会。在单核CPU上你的程序就像一条单车道所有车辆数据都得一辆接一辆地排队通过。当数据量爆炸式增长时这条单车道就成了最大的瓶颈。这时候你可能会想到多线程。没错多线程确实能让CPU的多个核心同时工作把单车道拓宽成几条车道。但CPU的设计初衷是“通用计算”它要处理复杂的逻辑判断、分支预测、中断响应它的核心数量有限通常几个到几十个每个核心都非常“聪明”但“昂贵”。对于海量数据中那些简单、重复、但数量极其庞大的计算任务比如对一张1000万像素的图片每个像素都做同样的滤镜计算用CPU的多线程来处理就像是让一群博士去流水线上拧螺丝效率不高成本还大。于是GPU图形处理器登场了。GPU最初是为图形渲染设计的它的任务极其规律对屏幕上成千上万个像素点顶点、片元执行几乎相同的着色器程序。这种“单指令多数据”SIMD的计算模式催生了GPU高度并行的架构。它拥有成百上千个更简单、更专注的计算核心虽然每个核心的“智商”不如CPU但“人多力量大”在并行处理海量同质化数据时能爆发出惊人的吞吐量。CUDACompute Unified Device Architecture就是NVIDIA公司为它的GPU打造的一套通用并行计算平台和编程模型。它让开发者能够使用熟悉的C/C等语言直接编写在GPU上运行的程序从而将GPU的强大并行计算能力从图形领域解放出来应用到科学计算、深度学习、金融分析等更广阔的领域。简单来说CUDA让你能指挥GPU这支“千军万马”的部队去完成那些适合大规模并行处理的计算任务。而要指挥好这支军队你首先得了解它的“军制”和“术语”。这篇文章我们就来彻底厘清CUDA编程中最核心的那些名词概念这是你踏入GPU并行计算世界的第一步也是避免后续“踩坑”的关键。2. 硬件架构视角GPU是如何组织它的计算资源的理解CUDA编程必须从理解GPU的硬件架构开始。你不能把GPU当成一个黑盒子只知道它“快”而要明白它为什么快以及它的能力边界在哪里。2.1 流式多处理器SMGPU的“计算兵团”你可以把GPU想象成一个庞大的计算军团。这个军团的基本作战单位不是单个士兵而是一个个“连队”在NVIDIA的术语里这个“连队”就叫流式多处理器。SM是GPU上真正执行指令和计算的核心部件。一个GPU芯片由多个SM组成。例如NVIDIA Tesla V100有80个SM而消费级的RTX 4090则有128个SM。每个SM内部又包含了CUDA核心这是执行整数和单精度浮点运算的基本单元。你可以粗略地理解为“一个计算士兵”。一个SM里集成了几十到上百个CUDA核心。张量核心从Volta架构开始引入专门用于执行矩阵乘加运算在深度学习训练和推理中速度极快。寄存器文件SM内部的高速存储供正在SM上执行的线程快速存取私有变量。速度极快但容量有限每个线程能分到的寄存器数量是重要限制。共享内存一块被同一个SM内所有线程共享的、可编程的高速缓存。它比全局内存快得多是优化性能的关键。调度器/分发单元负责将线程块后面会讲调度到SM上执行并管理线程束。关键理解你的CUDA程序内核是被分配到各个SM上执行的。一个SM可以同时处理多个线程块但一个线程块只能在一个SM上执行不能被拆分到多个SM。SM的数量直接决定了GPU的并行处理能力上限。2.2 内存层次结构数据的“高速公路与乡间小道”在GPU上数据存放在不同速度和容量的存储器中形成了一个层次结构。理解这个层次是进行有效内存访问优化的基础。全局内存GPU的“主内存”也就是我们常说的显存。容量最大几GB到几十GB但延迟最高带宽也相对较慢虽然比CPU内存带宽高很多。所有SM都可以访问全局内存。在主机CPU代码中通过cudaMalloc分配的就是这块内存。常量内存位于显存中但有特殊的缓存机制。适合存储所有线程都需要读取、且在核函数执行期间不会改变的数据如常数、查找表。访问常量内存如果缓存命中速度极快。纹理内存同样是具有缓存的只读内存最初为图形纹理设计优化了具有空间局部性的访问模式比如图像中相邻像素的读取。在某些访问模式下比全局内存高效。共享内存位于SM内部速度堪比寄存器。由同一个线程块内的所有线程共享。这是手动性能优化的主战场。你可以把需要频繁读写、在线程间需要通信的中间数据放在共享内存中从而避免昂贵的全局内存访问。寄存器位于SM内部速度最快。每个线程都有自己私有的寄存器。编译器会尽可能将自动变量如循环索引、临时变量分配到寄存器。寄存器资源是稀缺的如果一个线程使用了太多寄存器会导致SM上能同时驻留的线程数量减少可能影响并行度。本地内存实际上位于全局内存中。当线程的私有数据如大的局部数组、寄存器溢出的变量无法完全放入寄存器时编译器会将其放入本地内存。访问速度很慢应尽量避免。一个生动的类比把SM比作一个工厂车间计算单元寄存器就是每个工人手边的工作台极快但空间小。共享内存是这个车间里的公共工具墙很快车间内共享。全局内存则是工厂外的大型中央仓库容量大但来回取货慢。高效的程序要尽量让工人在工作台寄存器完成操作频繁使用的工具放在工具墙共享内存只有原材料和最终成品才去仓库全局内存存取。2.3 线程束SM执行的基本单元这是CUDA模型中最精妙也最容易让人困惑的概念之一。GPU的SM并不是以单个线程为单位进行调度和执行的。线程束是SM执行指令的基本单位。目前一个线程束包含32个连续的线程。这32个线程被“捆绑”在一起以锁步的方式执行同一条指令。也就是说在任何一个时钟周期线程束中的所有32个线程都在执行相同的指令只是操作的数据可能不同。这源于GPU的SIMD单指令多数据架构。这种设计极大地简化了控制逻辑和调度开销。但这也带来了一个核心约束分支发散。分支发散如果线程束中的线程在执行时遇到了条件判断如if-else并且这32个线程的走向不一致一部分走if一部分走else那么线程束就必须串行化执行所有分支路径。先执行走if的线程走else的线程等待再执行走else的线程。这会导致性能严重下降。注意编写CUDA内核时要尽量避免线程束内的分支发散。例如尽量让相邻的线程它们的threadIdx连续执行相同的控制流。可以通过重构算法或使用类似__shfl_sync的束内原语来减少发散。3. 编程模型视角如何用代码组织你的并行任务硬件架构决定了GPU的能力而CUDA编程模型则提供了我们组织计算任务的方法。这是你编写.cu文件时直接打交道的抽象层。3.1 网格、线程块与线程三级并行层次这是CUDA编程模型的骨架。它以一种层次化的方式组织并行线程完美映射到GPU的硬件层次。线程最小的执行单元。每个线程都独立运行内核函数的一份副本并通过内置的threadIdx、blockIdx等变量来区分自己从而处理不同的数据。线程块一组线程的集合。一个线程块内的线程会被调度到同一个SM上执行。可以通过共享内存进行高效通信与协作。可以通过__syncthreads()函数进行同步确保块内所有线程都执行到某个点后再继续。线程块的大小每个块包含多少线程在启动内核时由开发者指定如numBlocks, threadsPerBlock中的第二个参数。通常线程块的大小是线程束大小32的整数倍例如128、256、512。网格所有线程块的集合。一个内核启动就对应一个网格。网格中的线程块可以被调度到任意可用的SM上执行且执行顺序是不确定的、并行的。它们的关系与硬件映射一个网格被启动后其包含的所有线程块被分发到GPU的各个SM上等待执行。一个SM可以同时执行多个线程块具体数量受限于SM的资源如寄存器、共享内存总量。一个线程块一旦被分配到一个SM上就会一直驻留直到执行完毕。在SM上线程块被进一步划分为线程束来调度执行。如何确定网格和线程块的尺寸这是一个经验与性能分析相结合的过程。一个常见的启发式方法是线程块大小通常设为256或512。太小如64可能无法充分利用SM太大如1024可能因为寄存器限制而减少SM上同时驻留的块数。网格大小根据总数据量N和线程块大小blockSize计算gridSize (N blockSize - 1) / blockSize向上取整。确保有足够多的线程块至少是SM数量的几倍来隐藏内存访问延迟并让所有SM都保持忙碌。3.2 内核函数在GPU上执行的代码内核函数是CUDA编程的核心它定义了每个线程要执行的操作。它用__global__关键字声明由主机CPU调用在设备GPU上执行。// 一个简单的向量加法内核 __global__ void vectorAdd(const float* A, const float* B, float* C, int numElements) { // 计算当前线程的全局索引 int i blockDim.x * blockIdx.x threadIdx.x; // 确保索引不越界 if (i numElements) { C[i] A[i] B[i]; // 每个线程负责一个加法 } }内核启动语法kernelNamegridDim, blockDim, sharedMemSize, stream(arguments...);gridDim网格的维度可以是dim3类型指定了线程块在x, y, z三个方向上的数量。blockDim线程块的维度同样是dim3类型指定了每个线程块中线程在x, y, z三个方向上的数量。sharedMemSize可选动态分配的共享内存大小字节。stream可选关联的CUDA流。3.3 主机与设备CPU与GPU的协同CUDA编程是异构编程涉及两个处理器主机指CPU及其内存主机内存。设备指GPU及其显存设备内存。它们有各自独立的内存空间。因此数据必须在主机和设备之间进行传输这是CUDA程序中的一个主要开销来源。基本流程如下在主机上分配并初始化数据。使用cudaMalloc在设备上分配内存。使用cudaMemcpy将数据从主机复制到设备。启动内核函数在设备上处理数据。使用cudaMemcpy将结果从设备复制回主机。使用cudaFree释放设备内存。关键API与概念cudaMalloc/cudaFree设备内存的分配与释放。cudaMemcpy内存复制。方向由cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost等参数指定。固定内存使用cudaMallocHost分配的主机内存该内存页被锁定不可被操作系统交换出去。GPU可以通过DMA直接访问固定内存从而在主机到设备的数据传输中获得更高的带宽。对于频繁传输的数据应使用固定内存。统一内存从CUDA 6.0开始引入通过cudaMallocManaged分配。系统自动管理数据在主机和设备间的迁移简化了编程模型但开发者需要对访问模式有一定理解以获得最佳性能。4. 执行与调度模型GPU如何幕后管理成千上万的线程理解了静态的编程模型我们还需要了解动态的执行过程。GPU如何调度成千上万个线程以实现极高的吞吐量4.1 隐藏延迟GPU高性能的秘诀GPU计算核心的执行速度非常快但访问全局内存的延迟非常高需要几百个时钟周期。如果线程在发出内存加载请求后只是空等那么大部分时间计算核心都会处于闲置状态性能会极其低下。GPU解决这个问题的方法是大量并行线程的快速切换以隐藏内存访问延迟。当一个线程束因为等待内存数据而停滞时SM的调度器会立刻切换到另一个就绪的线程束去执行。由于SM上驻留着成百上千个线程来自多个线程块调度器可以确保几乎在任何时刻都有线程束可以执行计算指令从而让计算核心始终保持忙碌将内存延迟“隐藏”在计算之下。这就要求你的内核启动必须有足够的并行度。即同时启动的线程总数网格大小 × 线程块大小要远远大于GPU的物理核心数通常是几万甚至几十万上百万才能充分隐藏延迟。4.2 占用率一个重要的性能指标占用率是指每个SM上活跃的线程束数量与SM支持的最大线程束数量之比。高占用率意味着SM上有更多的线程束可以参与调度有助于更好地隐藏延迟。然而高占用率并不总是等于高性能。影响占用率的主要因素有线程块大小线程块越大每个块提供的线程束越多但SM上能同时驻留的块可能越少受资源限制。寄存器使用量每个线程使用的寄存器数量。SM上的寄存器总量是固定的。如果每个线程使用很多寄存器那么SM上能同时驻留的线程总数就会减少从而降低占用率。共享内存使用量每个线程块使用的共享内存大小。同样SM的共享内存总量固定使用过多会限制同时驻留的线程块数量。有时为了使用更多的寄存器或共享内存来优化单个线程的性能或减少全局内存访问可以接受较低的占用率。这是一个需要权衡和通过性能分析工具来评估的过程。NVIDIA提供的CUDA Occupancy Calculator可以帮助你分析这些约束。4.3 同步在正确的时间点协调线程并行计算中线程间的协调至关重要。__syncthreads()这是一个线程块级别的屏障同步。调用该函数后线程块内的所有线程都必须执行到此位置然后才会继续执行后面的指令。非常重要必须确保线程块内所有线程都能到达这个同步点否则会导致死锁。例如不能在只有部分线程满足的if条件内调用__syncthreads()。原子操作当多个线程需要读写同一个全局内存或共享内存地址时为了避免竞争条件需要使用原子操作如atomicAdd、atomicExch等。原子操作保证该操作是“不可分割”的。但原子操作是串行的会严重影响性能应尽量避免或减少使用。网格级别的同步在内核函数内部没有直接提供网格级别的同步原语。因为线程块执行顺序不确定且可能在任何时候结束。如果需要全局同步通常需要将计算拆分为多个内核启动利用CUDA流和事件进行更复杂的控制。5. 内存访问模式为什么对齐与合并访问如此重要即使你启动了足够多的线程如果内存访问模式很糟糕性能也会一落千丈。GPU的全局内存带宽虽然高但要高效利用它必须满足其特定的访问模式。5.1 内存事务与合并访问GPU的全局内存控制器是以内存事务为单位来服务内存请求的。一次事务可以读取32字节、64字节或128字节对齐的连续内存数据。理想情况下一个线程束32个线程的所有内存请求应该合并成少数几个甚至一个内存事务。合并访问当一个线程束中的所有线程访问全局内存中一片连续的、对齐的数据块时它们的访问请求可以被“合并”成一个或几个内存事务从而最大化内存带宽利用率。未合并访问如果线程束中的线程访问的内存地址是分散的、不连续的那么每个线程的请求都可能需要单独的内存事务导致有效带宽急剧下降。5.2 实践中的访问模式优化确保对齐访问分配设备内存时尽量使用cudaMalloc它保证至少256字节对齐。对于自定义数据结构可以使用__align__关键字或alignas来确保对齐。设计线程索引映射数据索引这是最关键的一点。要让相邻的线程threadIdx.x连续访问相邻的内存地址。优化前跨步访问差int tid blockDim.x * blockIdx.x threadIdx.x; int index tid * stride; // 如果stride很大相邻线程访问的地址相隔很远优化后连续访问好int tid blockDim.x * blockIdx.x threadIdx.x; int index tid; // 相邻线程访问连续的index对于多维数组要特别注意行主序/列主序确保内层循环对应的线程索引是连续的。利用共享内存作为中转当数据访问模式无法做到全局内存完美合并时一个经典的优化模式是让线程块中的线程以合并访问的方式将数据从全局内存加载到共享内存。在共享内存中进行需要随机访问或线程间共享数据的计算。最后再将结果以合并访问的方式写回全局内存。 共享内存的访问延迟远低于全局内存且对访问模式不敏感这完美地解决了问题。5.3 一个简单的性能对比思考假设你要对一个矩阵的每一行元素求和。矩阵在内存中是按行存储的。方案A每个线程负责一行。那么线程束中的32个线程每个线程访问的是不同行的首元素地址间隔为“一行的大小”这是最糟糕的未合并访问。方案B每个线程负责一列。那么线程束中的32个线程访问的是同一行的前32个连续元素。这实现了完美的合并访问。当然这需要在线程块内进行归约求和但内存访问效率的提升是巨大的。理解并应用这些概念是写出高性能CUDA程序的关键。它不仅仅是“让程序跑起来”而是“让程序飞起来”的必经之路。在后续的实际编码中你会反复用到这些概念来分析和优化你的内核。