学习CUDA的人几乎都会在某个深夜盯着一堆缩写问出同一个问题SM、SP、GRID、BLOCK、THREAD这几个东西到底是怎么对应的网上文章不少但要么是概念罗列要么直接丢一张芯片架构图让人自己悟。我当年从CPU多线程转到GPU编程时也被这层抽象折磨了好一阵代码里明明是线程到了硬件怎么就变成core了block到底是不是对应SM为什么大家都说要让block数量远超SM数量后来在真实项目里把一个kernel从能跑调到跑得快才算是把这条链路彻底理顺。这篇就把我自己的理解路径完整写出来按软件层和硬件层先分层再沿一次kernel启动走一遍最后落到具体参数计算和调优排错的顺序展开给正在被这些概念折磨的读者一条直达本质的捷径。1. 先给这五个名词分层软件世界的三个名字与硬件世界的两个名字1.1 软件层Thread、Block、Grid是任务怎么切的问题CUDA的编程模型里程序员直接面对的三个层次本质上是任务的切分方式。Thread是执行kernel函数的最小单位。写代码的时候每个线程执行的是同一份代码但通过线程索引区分自己处理哪一份数据。比如threadIdx.x是0的线程处理第0个元素threadIdx.x是1的线程处理第1个元素。一个kernel启动后所有线程跑的是同一段指令数据却各不相同这就是所谓的SIMT单指令多线程风格。Block是若干个线程的集合。block内的线程不只是恰好排在一起它们共享同一块共享内存shared memory可以用__syncthreads()做屏障同步还可以通过warp shuffle指令直接交换寄存器里的数据。也就是说block是CUDA里能够协作的最小单位block和block之间不允许直接同步只能通过全局内存加原子操作这种间接方式协作。Grid则是一次kernel启动产生的全部block的集合。你写kernelgridSize, blockSize的时候实际上就是在声明我要开一个grid这个grid里有gridSize个block每个block里有blockSize个线程。这三层关系可以用一句话记一个Grid包含若干Block一个Block包含若干Thread。这里有个容易忽略的细节这三个维度都支持一维、二维、三维。gridDim.x/y/z、blockIdx.x/y/z、blockDim.x/y/z、threadIdx.x/y/z这些内建变量全是三维的只是大部分例子只用x维。真正写代码时二维索引在图像处理里很常见比如int x blockIdx.x * blockDim.x threadIdx.x;处理列坐标再用y维处理行坐标。但不管维数怎么变底层逻辑都一样用block的编号和block内的线程编号拼出一个全局唯一的线程ID。1.2 硬件层SP是执行单元SM是装着执行单元的车间硬件这边也有两个核心名词而且它们和软件名词特别容易混淆。SPStreaming Processor流处理器是GPU里真正干活的算术单元负责执行浮点、整数、逻辑等标量运算指令。NVIDIA这边通常叫它CUDA CoreAMD那边叫流处理器Stream Processor虽然叫法不同指的都是一个时钟周期内能执行一条标量指令的处理单元。注意标量这个词很关键一个SP一次只处理一个线程的一个数据不是一个向量这和CPU里的SIMD向量指令完全是两码事。SMStreaming Multiprocessor流式多处理器是把一堆SP组织起来的车间。一个SM里除了有几十到上百个SP还有寄存器文件、共享内存、L1缓存、warp调度器、指令分发单元等一整套配套硬件。GPU由多个SM组成比如GA102芯片有82个SMA100有108个SM。SM是整个GPU硬件调度的核心单位无论是block还是warp最终都要落到某个SM上执行。不同架构下每个SM里的SP数量差别很大理解这一点对后面算账很有帮助架构代表芯片每SM的SP/CUDA Core数量FermiGTX 48032KeplerGTX 680192MaxwellGTX 980128PascalGTX 1080128VoltaV10064FP32单元TuringRTX 2080 Ti64FP32单元AmpereA100 / RTX 309064 / 128FP32单元HopperH100128FP32单元看这个表就能明白一件事硬件型号不同同一个Block映射到SM后的物理并行能力完全不一样。这也解释了为什么CUDA程序在不同显卡上的性能表现天差地别。1.3 为什么这五个词总被搞混软件抽象把硬件细节藏起来了搞混的原因其实不在你而在CUDA的设计目标本身——它刻意让程序员只面向Thread/Block/Grid编程把硬件细节藏起来。早期很多教程又喜欢画SP就是Core这种简化图很多人就顺理成章地把Thread直接对应SP、把Block直接对应SM。这个对应方向是对的但漏掉了最关键的中间层Warp。Warp线程束是32个线程组成的一个执行组。block里的线程进入SM后硬件会自动把它们切成若干个warp然后SM以warp为单位取指令、发射指令、等待结果。所以严格来说Thread对应的是SM内某个执行通道lane而一次真正的执行是由一个warp的32个线程在32个并行执行单元上同时完成的。另外还有一个天然造成混乱的点软件层的Thread、Block、Grid是理想化的你可以声明几百万个线程但硬件层的SP、SM是物理的数量有限。中间靠的是硬件调度器把海量软件线程分批映射到有限硬件单元上。把软件抽象和硬件执行这两个层面分开理解后面的所有问题都会顺很多。2. 从一次kernel启动看五个名词如何逐个挂钩2.1 一个最小示例saxpy空谈概念没有用直接看代码。saxpy单精度AXPY是CUDA里最经典的入门例子比hello world更能说明映射关系__global__ void saxpy(float *y, const float *x, float a, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { y[i] a * x[i] y[i]; } } int main() { int n 1 20; // 1048576 个元素 int blockSize 256; // 每个block 256个线程 int gridSize (n blockSize - 1) / blockSize; // 4096 个block saxpygridSize, blockSize(y, x, 2.0f, n); cudaDeviceSynchronize(); return 0; }这里gridSize 4096blockSize 256总线程数正好是4096 × 256 1,048,576每个线程负责一个数组元素。注意int i blockIdx.x * blockDim.x threadIdx.x这一行它就是把三维的block编号和线程编号压平成一维全局索引的经典公式后面所有寻址都从它展开。2.2 kernel之后GPU内部发生了什么从CPU调用kernel到SP上真正执行中间大概经过这四步CPU把kernel的启动配置和参数写入命令缓冲区GPU端的前端调度引擎NVIDIA叫GigaThread引擎收到命令后为这次启动创建一个Grid。GigaThread引擎把Grid里的Block逐一分配给各个SM。注意分配是动态的哪个block去哪个SM由硬件决定程序员控制不了。4096个block会分散到所有SM上某个SM处理完一个block马上领取下一个直到所有block跑完。Block进入SM后SM的线程块调度器把它切分成Warp。blockSize256意味着每个block被切成8个warp256÷328。SM里的warp调度器每个周期挑选一个就绪的warp把它的指令发射给SPCUDA Core阵列执行。SP的个数决定了每个周期能同时让多少个线程的标量指令真正落地。这四步走完五个名词的对应关系就非常清晰了软件/任务层对应硬件单元一句话说明GRID一次kernel整张GPU所有SM一个grid的block会被分发到所有SM上跑完为止BLOCK线程块某个SM一个block整体驻留在同一个SM上绝不能被拆到多个SMTHREAD线程SM内的SP/CUDA Core线程指令最终由SP执行但以warp为单位成组发射WARP32个线程SM内一组并行执行单元硬件真正调度和执行的最小单位2.3 Block这个交接面为什么如此重要整个映射链条里Block是软件和硬件之间的交接面这个设计是刻意的而且非常精妙。首先它保证了可扩展性。你写代码时只需要定义block运行时不关心GPU有几个SM。同一份代码8个SM的显卡能跑108个SM的A100也能跑只是block被分发到更多或更少的SM上并行处理。这种透明的可扩展性是CUDA程序能跨代兼容的根本原因。其次Block是协作的天然边界。block内可以共享内存、同步block之间不允许直接协作。这个限制看起来是约束实际上给了硬件极大的自由既然block之间互相独立硬件就可以把block当作完全自治的工作包随意调度到任意SM上不需要考虑跨SM通信、锁同步这些麻烦事。这也是为什么GPU能靠硬件调度器实现极高吞吐的原因。第三Block的大小直接决定了资源占用的粒度。SM能同时驻留多个block但能驻留多少个受两个硬性约束硬件上限每SM最多2048个线程、最多32个block和资源上限寄存器、共享内存这些会被block瓜分的软资源。block设大了SM上同时驻留的block数就少调度弹性变小block设小了每个block的warp太少又填不满调度器的发射能力。这个博弈下一节用具体数字算给你看。3. 把账算明白用一块RTX 3090推演GRID、BLOCK、THREAD的量化关系3.1 先从全局线程ID说起在深入硬件账本之前先把这个公式刻进脑子里。一维情况下全局线程索引是int tid blockIdx.x * blockDim.x threadIdx.x;二维情况下假设图像宽width、高height通常写成int ix blockIdx.x * blockDim.x threadIdx.x; int iy blockIdx.y * blockDim.y threadIdx.y; int idx iy * width ix; // 压平成一维数组下标理解这个公式的关键在于blockIdx.x是这个block在grid里的编号blockDim.x是每个block沿x方向的线程数两者相乘就是这个block之前已经排了多少个线程再加上threadIdx.x就是block内部第几个线程。所有维度的寻址都是这个思路的推广。很多人写kernel时算错索引原因就是混淆了blockIdx和threadIdx的语义前者是block的编号后者是block内线程的编号两者相差一个blockDim的倍数关系。这个错误在数据量不是blockSize整数倍的时候尤其致命后面排错部分还会提到。3.2 一块真实GPU的硬件账本以RTX 3090GA102芯片为例它的关键硬件参数如下项目数值SM数量82每SM FP32核心数SP128每SM最大驻留线程数2048每SM最大驻留block数32每block最大线程数1024每SM寄存器文件65536个32位寄存器256KB每线程默认最大寄存器数255现在把前面saxpy的配置拿过来算。blockSize256时每个SM最多能驻留多少个block先看线程上限2048÷2568也就是最多8个block再看block数量上限32个832所以线程上限先到顶答案是8个block。这8个block占满后SM上同时驻留的线程数是2048正好100%占满这个SM的线程容量。整个GPU能同时驻留的线程总数是82个SM × 2048 167,936个线程。而saxpy一共有1,048,576个线程所以需要大约6.25波才能全部跑完。这就是一个grid里的线程并非同时执行的最直观证据硬件只能同时容纳这么多剩下的都在排队等着SM空出来。3.3 寄存器与共享内存会改写映射结果上面的计算只考虑了硬件上限实际工程里经常被软资源卡脖子最常见的是寄存器和共享内存。假设saxpy的kernel每线程用了40个寄存器。RTX 3090每个SM有65536个寄存器那么65536÷401638也就是每SM最多只能同时容纳1638个线程的寄存器上下文。如果blockSize2561638÷2566.4向下取整就是6个block总共1536个线程。于是occupancy占用率就从100%掉到1536÷204875%。也就是说block到SM的映射数量不是写死的而是硬件上限和资源约束共同决定的动态结果。寄存器用得越多SM能驻留的block就越少共享内存同理。假设每个block申请40KB共享内存而SM的共享内存总共128KB那最多只能驻留3个block3×40120KB4×40160KB超了。如果blockSize2563个block只有768个线程驻留占用率直接掉到37.5%。这也是为什么共享内存用多了性能反而断崖式下跌。3.4 为什么开发社区常见的BlockSize是128、256、512理解了上面的账就不难明白为什么CUDA社区里大家默认选128、256、512这些值而很少看到100、200、1000这种顺手的数字。第一个原因是warp大小为32。blockSize128正好是4个warp256是8个warp都是整数个warp不会产生残缺warp。blockSize100的话会被切成3个完整的warp96线程加1个只有4线程的残warp后面详说残warp的代价。第二个原因是调度选择空间。SM的warp调度器需要在众多warp中挑就绪的下一个执行驻留的warp越多调度器越不容易因为访存等待而空转。blockSize256时每个SM驻留8个block提供64个warp调度器有非常充足的选择余地。blockSize32时每个SM最多驻留32个block也就是32个warp明显少了一半。第三个原因是同步和资源粒度。block太大比如1024每SM只能驻留2个blockblock内的__syncthreads()会阻塞大量线程等待寄存器分配也不灵活block太小比如32又会让block数量迅速撞上每SM 32个block的上限同样浪费。128到512之间是一个经过广泛验证的甜区256尤其常见8个warp、资源占用适中、边界条件好处理。4. 真正被执行的是Warp给对应关系补上最关键的一环4.1 SIMT32个线程同一条指令前面反复提到warp现在把这一层彻底讲透。Warp是GPU硬件调度和执行的最小单位由32个线程组成。一个warp里的线程执行同一条指令但每个线程操作自己的数据这套执行模型叫SIMTSingle Instruction, Multiple Thread。你可以把warp理解成一列并排前进的士兵指挥官喊一声齐步走一条指令32个士兵同时迈出左脚各自处理各自的数据。如果某个士兵因为条件分歧if-else需要走不同的路径那这个warp就得分别走两条分支每条分支上的线程各执行一次另一部分线程被掩蔽掉。这就是所谓的线程发散thread divergence它会让warp的执行效率打折。在硬件层面一个warp的32个线程对应SM里一组32个并行执行单元lane。RTX 3090的SM有128个SP可以理解为4组32-lane的执行阵列每组由一个warp调度器掌管。每个周期4个调度器各选一个就绪的warp发射到对应阵列上执行。所以一个SM每个周期最多能同时执行4个warp的指令也就是4×32128个线程的标量指令——正好对应128个SP的数量。4.2 隐藏延迟的真相为什么驻留的warp数这么重要GPU性能优化的核心说穿了就两个字延迟。一条内存访问指令从发出到数据返回可能需要几百个时钟周期。如果SM里只驻留了4个warp这4个warp都在等内存数据回来那SP阵列就全部空转性能惨不忍睹。正确的逻辑是SM里驻留的warp数量越多调度器在某个warp等待时切换到其他就绪warp的机会就越多。GPU正是靠这种线程级并行来把访存延迟、指令延迟掩盖掉让SP始终有活干。这也是为什么提高occupancy即驻留warp数占最大容量的比例是CUDA优化最常见的手段。从这个角度回头看blockSize的选择blockSize64每SM驻留32个block、64个warpblockSize256每SM驻留8个block、64个warp。两者在RTX 3090上的occupancy一样但block数量差4倍。block多的好处是block级调度更灵活一个block结束另一个block马上顶上坏处是block切换也有开销而且64线程的block里只有2个warp如果这个block里出现访存停顿调度器的选择范围局限在同一个SM的其他block里。工程实践普遍认为256比64更稳就是这个道理。4.3 Block内线程数为什么最好是32的倍数现在可以回答那个经典问题了为什么block内线程数最好取32的倍数因为GPU切warp是按连续32个线程一组的如果blockSize100切出来的结果是3个满warp0到95号线程加1个只有4个线程的残warp96到99号线程。残warp的28个lane执行通道虽然不参与计算但执行指令时照样占一个warp的发射槽位。也就是说这100个线程实际占了4个warp的执行资源有效利用率只有96÷12875%。更要命的是如果kernel里有__syncthreads()残warp也必须参与同步因为同步的粒度是block内所有线程空转的lane同样要到位。所以两个经验直接给出第一blockSize选32的整数倍第二如果数据量不是32的倍数宁可多开几个线程用if (i n)做边界保护也别让warp残缺。这是每个CUDA项目都该无脑执行的规矩。4.4 那么THREAD与SP的对应到底怎么描述才准确绕了一大圈回到标题的核心问题。准确的链条是这样一句话一个Thread最终由SM内的某个SPCUDA Core执行它的指令但Thread和SP之间不是一对一的固定绑定而是以Warp为单位、由调度器动态分配的临时占用关系。Thread被组织成WarpWarp被调度器发射到SP阵列上一个warp的32个线程在32个执行单元上同时执行同一条指令执行完这一个周期下个周期同一个SP可能执行另一个warp的指令。所以THREAD和SP的关系更像顾客和服务窗口顾客线程数量可以成千上万服务窗口SP只有一百多个顾客必须排队成组warp被叫号叫到谁谁就临时占用窗口。SP不是某个线程的私有财产而是被所有线程时分复用的。这也是为什么一个grid里可以声明几百万个线程而GPU物理上只有几百上千个SP——它们根本不是一一对应的关系靠的是线程池排队调度。5. 实践中的对应关系启动配置、调优与排错的经验5.1 用Occupancy的思路反推启动配置理解了映射关系之后配置kernel启动参数就是一个根据硬件上限倒推的过程。以RTX 3090为例忽略寄存器等软资源不同blockSize对应的理论占用率如下blockSize每SM可驻留block数每SM驻留线程数理论占用率6432受block上限限制2048100%128162048100%25682048100%51242048100%102422048100%看起来全是100%但一旦叠加寄存器或共享内存的约束差距就出来了。还是每线程40个寄存器的例子blockSize128时1638÷12812个block驻留1536线程75%blockSize512时1638÷5123个block驻留1536线程也是75%而blockSize256时同样是6个block、1536线程。三种配置的occupancy一样但block粒度不同实际性能会因调度开销、同步粒度而有差异这就是需要实测的原因。NVIDIA官方提供的Occupancy Calculator在Nsight Compute里能直接算出给定blockSize和资源用量下的期望占用率强烈建议每次调kernel都跑一遍。它最大的价值不是帮你自动选参数而是让你看到是什么在卡占用率——是线程上限、block上限、寄存器还是共享内存对症下药远比瞎试参数高效。5.2 常见的启动配置错误与定位这几年我踩过的坑归纳起来就三类基本覆盖了绝大多数配置相关的故障。第一类启动直接报invalid configuration argument。这种错误九成是blockSize超过了设备上限比如1024或者gridSize、blockSize某个维度设成了0少数情况是gridDim超过2^31-1的上限。排查方法很简单启动前打印一下blockIdx、blockDim这些值或者调用cudaGetLastError()拿到返回码顺着报错信息十秒就能定位。注意kernel这个语法本身是异步的启动错误不会立刻暴露必须在启动后加cudaDeviceSynchronize()才能把错误真正捕获出来。第二类程序能跑但结果不对而且时好时坏。这种情况八成是索引越界或数据漏算。比如把gridSize n / blockSize写成了整除而不是向上取整当n不能被blockSize整除时末尾几个元素就没人处理反过来如果没用if (i n)做边界检查多余的线程会把数组后面不属于它的内存也写一遍轻则数据错乱重则直接导致设备崩溃比如非法内存访问报unspecified launch failure。我在实际项目里的经验是所有涉及数组的kernel都保留边界判断哪怕你确信数据量是blockSize的整数倍也留着成本极低救命次数极多。第三类程序和结果都对但性能远低于预期。这时候别急着调代码先看occupancy。用Nsight Compute看Achieved Occupancy如果长期低于70%要么是寄存器用太多要么是共享内存用太多要么是blockSize设得不符合硬件胃口。有一次我把一个kernel的每线程寄存器压到32个以内占用率从50%提到100%性能直接翻倍而代码逻辑一行没改只改了一个编译限制选项。这类问题靠肉眼是看不出来的必须用工具量化。5.3 几条可以直接用的经验法则结合上面的分析和踩坑经历总结几条可以直接抄走的经验法则blockSize默认取256除非有明确理由。256是32的8倍、8个warp在调度灵活性和资源开销之间平衡最好。数据量小可以用128kernel里shared memory很大时可以降到128或64但别低于64。让grid中的block数量至少是SM数量的几倍。一个简单标准block数至少是SM数的4到8倍。如果82个SM的卡上只开了82个block一旦某个SM分配不均空闲SM就出现了block数远超SM数才能靠动态调度抹平负载差异。在kernel启动后统一加一次cudaGetLastError()检查。CUDA错误是异步的不加这一句错误会在下一次同步点才炸出来定位成本高得多。工程上可以封装一个宏或工具函数统一处理错误返回。优先调资源再调blockSize。寄存器超标就先压寄存器__launch_bounds__或-maxrregcount共享内存超标就减每线程用量这两样是occupancy的隐形杀手比纠结blockSize是256还是512重要得多。每次改完配置用Nsight Compute看一眼Achieved Occupancy和Warp Stall Metrics让数据说话。背下来一张卡的特性参数不如养成算一遍、测一遍、调一遍的习惯。最后再分享一个小技巧。我刚开始写CUDA那阵子特别迷信block越大吞吐越高把blockSize设成1000结果性能比256差一大截Nsight一看占用率被寄存器卡在了40%。后来每次调kernel我都先跑一遍设备信息查询工具deviceQuery看自己卡片的SM数量、每SM线程上限这些硬参数再手算一遍映射账最后用Nsight验证。这套流程走熟了之后基本没有出现不知道block数量应该设多少的迷茫。建议你也把这个方法存成自己的检查清单写kernel之前先算账就能少走我当年那段弯路。
学习CUDA的人几乎都会在某个深夜盯着一堆缩写问出同一个问题SM、SP、GRID、BLOCK、THREAD这几个东西到底是怎么对应的网上文章不少但要么是概念罗列要么直接丢一张芯片架构图让人自己悟。我当年从CPU多线程转到GPU编程时也被这层抽象折磨了好一阵代码里明明是线程到了硬件怎么就变成core了block到底是不是对应SM为什么大家都说要让block数量远超SM数量后来在真实项目里把一个kernel从能跑调到跑得快才算是把这条链路彻底理顺。这篇就把我自己的理解路径完整写出来按软件层和硬件层先分层再沿一次kernel启动走一遍最后落到具体参数计算和调优排错的顺序展开给正在被这些概念折磨的读者一条直达本质的捷径。1. 先给这五个名词分层软件世界的三个名字与硬件世界的两个名字1.1 软件层Thread、Block、Grid是任务怎么切的问题CUDA的编程模型里程序员直接面对的三个层次本质上是任务的切分方式。Thread是执行kernel函数的最小单位。写代码的时候每个线程执行的是同一份代码但通过线程索引区分自己处理哪一份数据。比如threadIdx.x是0的线程处理第0个元素threadIdx.x是1的线程处理第1个元素。一个kernel启动后所有线程跑的是同一段指令数据却各不相同这就是所谓的SIMT单指令多线程风格。Block是若干个线程的集合。block内的线程不只是恰好排在一起它们共享同一块共享内存shared memory可以用__syncthreads()做屏障同步还可以通过warp shuffle指令直接交换寄存器里的数据。也就是说block是CUDA里能够协作的最小单位block和block之间不允许直接同步只能通过全局内存加原子操作这种间接方式协作。Grid则是一次kernel启动产生的全部block的集合。你写kernelgridSize, blockSize的时候实际上就是在声明我要开一个grid这个grid里有gridSize个block每个block里有blockSize个线程。这三层关系可以用一句话记一个Grid包含若干Block一个Block包含若干Thread。这里有个容易忽略的细节这三个维度都支持一维、二维、三维。gridDim.x/y/z、blockIdx.x/y/z、blockDim.x/y/z、threadIdx.x/y/z这些内建变量全是三维的只是大部分例子只用x维。真正写代码时二维索引在图像处理里很常见比如int x blockIdx.x * blockDim.x threadIdx.x;处理列坐标再用y维处理行坐标。但不管维数怎么变底层逻辑都一样用block的编号和block内的线程编号拼出一个全局唯一的线程ID。1.2 硬件层SP是执行单元SM是装着执行单元的车间硬件这边也有两个核心名词而且它们和软件名词特别容易混淆。SPStreaming Processor流处理器是GPU里真正干活的算术单元负责执行浮点、整数、逻辑等标量运算指令。NVIDIA这边通常叫它CUDA CoreAMD那边叫流处理器Stream Processor虽然叫法不同指的都是一个时钟周期内能执行一条标量指令的处理单元。注意标量这个词很关键一个SP一次只处理一个线程的一个数据不是一个向量这和CPU里的SIMD向量指令完全是两码事。SMStreaming Multiprocessor流式多处理器是把一堆SP组织起来的车间。一个SM里除了有几十到上百个SP还有寄存器文件、共享内存、L1缓存、warp调度器、指令分发单元等一整套配套硬件。GPU由多个SM组成比如GA102芯片有82个SMA100有108个SM。SM是整个GPU硬件调度的核心单位无论是block还是warp最终都要落到某个SM上执行。不同架构下每个SM里的SP数量差别很大理解这一点对后面算账很有帮助架构代表芯片每SM的SP/CUDA Core数量FermiGTX 48032KeplerGTX 680192MaxwellGTX 980128PascalGTX 1080128VoltaV10064FP32单元TuringRTX 2080 Ti64FP32单元AmpereA100 / RTX 309064 / 128FP32单元HopperH100128FP32单元看这个表就能明白一件事硬件型号不同同一个Block映射到SM后的物理并行能力完全不一样。这也解释了为什么CUDA程序在不同显卡上的性能表现天差地别。1.3 为什么这五个词总被搞混软件抽象把硬件细节藏起来了搞混的原因其实不在你而在CUDA的设计目标本身——它刻意让程序员只面向Thread/Block/Grid编程把硬件细节藏起来。早期很多教程又喜欢画SP就是Core这种简化图很多人就顺理成章地把Thread直接对应SP、把Block直接对应SM。这个对应方向是对的但漏掉了最关键的中间层Warp。Warp线程束是32个线程组成的一个执行组。block里的线程进入SM后硬件会自动把它们切成若干个warp然后SM以warp为单位取指令、发射指令、等待结果。所以严格来说Thread对应的是SM内某个执行通道lane而一次真正的执行是由一个warp的32个线程在32个并行执行单元上同时完成的。另外还有一个天然造成混乱的点软件层的Thread、Block、Grid是理想化的你可以声明几百万个线程但硬件层的SP、SM是物理的数量有限。中间靠的是硬件调度器把海量软件线程分批映射到有限硬件单元上。把软件抽象和硬件执行这两个层面分开理解后面的所有问题都会顺很多。2. 从一次kernel启动看五个名词如何逐个挂钩2.1 一个最小示例saxpy空谈概念没有用直接看代码。saxpy单精度AXPY是CUDA里最经典的入门例子比hello world更能说明映射关系__global__ void saxpy(float *y, const float *x, float a, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) { y[i] a * x[i] y[i]; } } int main() { int n 1 20; // 1048576 个元素 int blockSize 256; // 每个block 256个线程 int gridSize (n blockSize - 1) / blockSize; // 4096 个block saxpygridSize, blockSize(y, x, 2.0f, n); cudaDeviceSynchronize(); return 0; }这里gridSize 4096blockSize 256总线程数正好是4096 × 256 1,048,576每个线程负责一个数组元素。注意int i blockIdx.x * blockDim.x threadIdx.x这一行它就是把三维的block编号和线程编号压平成一维全局索引的经典公式后面所有寻址都从它展开。2.2 kernel之后GPU内部发生了什么从CPU调用kernel到SP上真正执行中间大概经过这四步CPU把kernel的启动配置和参数写入命令缓冲区GPU端的前端调度引擎NVIDIA叫GigaThread引擎收到命令后为这次启动创建一个Grid。GigaThread引擎把Grid里的Block逐一分配给各个SM。注意分配是动态的哪个block去哪个SM由硬件决定程序员控制不了。4096个block会分散到所有SM上某个SM处理完一个block马上领取下一个直到所有block跑完。Block进入SM后SM的线程块调度器把它切分成Warp。blockSize256意味着每个block被切成8个warp256÷328。SM里的warp调度器每个周期挑选一个就绪的warp把它的指令发射给SPCUDA Core阵列执行。SP的个数决定了每个周期能同时让多少个线程的标量指令真正落地。这四步走完五个名词的对应关系就非常清晰了软件/任务层对应硬件单元一句话说明GRID一次kernel整张GPU所有SM一个grid的block会被分发到所有SM上跑完为止BLOCK线程块某个SM一个block整体驻留在同一个SM上绝不能被拆到多个SMTHREAD线程SM内的SP/CUDA Core线程指令最终由SP执行但以warp为单位成组发射WARP32个线程SM内一组并行执行单元硬件真正调度和执行的最小单位2.3 Block这个交接面为什么如此重要整个映射链条里Block是软件和硬件之间的交接面这个设计是刻意的而且非常精妙。首先它保证了可扩展性。你写代码时只需要定义block运行时不关心GPU有几个SM。同一份代码8个SM的显卡能跑108个SM的A100也能跑只是block被分发到更多或更少的SM上并行处理。这种透明的可扩展性是CUDA程序能跨代兼容的根本原因。其次Block是协作的天然边界。block内可以共享内存、同步block之间不允许直接协作。这个限制看起来是约束实际上给了硬件极大的自由既然block之间互相独立硬件就可以把block当作完全自治的工作包随意调度到任意SM上不需要考虑跨SM通信、锁同步这些麻烦事。这也是为什么GPU能靠硬件调度器实现极高吞吐的原因。第三Block的大小直接决定了资源占用的粒度。SM能同时驻留多个block但能驻留多少个受两个硬性约束硬件上限每SM最多2048个线程、最多32个block和资源上限寄存器、共享内存这些会被block瓜分的软资源。block设大了SM上同时驻留的block数就少调度弹性变小block设小了每个block的warp太少又填不满调度器的发射能力。这个博弈下一节用具体数字算给你看。3. 把账算明白用一块RTX 3090推演GRID、BLOCK、THREAD的量化关系3.1 先从全局线程ID说起在深入硬件账本之前先把这个公式刻进脑子里。一维情况下全局线程索引是int tid blockIdx.x * blockDim.x threadIdx.x;二维情况下假设图像宽width、高height通常写成int ix blockIdx.x * blockDim.x threadIdx.x; int iy blockIdx.y * blockDim.y threadIdx.y; int idx iy * width ix; // 压平成一维数组下标理解这个公式的关键在于blockIdx.x是这个block在grid里的编号blockDim.x是每个block沿x方向的线程数两者相乘就是这个block之前已经排了多少个线程再加上threadIdx.x就是block内部第几个线程。所有维度的寻址都是这个思路的推广。很多人写kernel时算错索引原因就是混淆了blockIdx和threadIdx的语义前者是block的编号后者是block内线程的编号两者相差一个blockDim的倍数关系。这个错误在数据量不是blockSize整数倍的时候尤其致命后面排错部分还会提到。3.2 一块真实GPU的硬件账本以RTX 3090GA102芯片为例它的关键硬件参数如下项目数值SM数量82每SM FP32核心数SP128每SM最大驻留线程数2048每SM最大驻留block数32每block最大线程数1024每SM寄存器文件65536个32位寄存器256KB每线程默认最大寄存器数255现在把前面saxpy的配置拿过来算。blockSize256时每个SM最多能驻留多少个block先看线程上限2048÷2568也就是最多8个block再看block数量上限32个832所以线程上限先到顶答案是8个block。这8个block占满后SM上同时驻留的线程数是2048正好100%占满这个SM的线程容量。整个GPU能同时驻留的线程总数是82个SM × 2048 167,936个线程。而saxpy一共有1,048,576个线程所以需要大约6.25波才能全部跑完。这就是一个grid里的线程并非同时执行的最直观证据硬件只能同时容纳这么多剩下的都在排队等着SM空出来。3.3 寄存器与共享内存会改写映射结果上面的计算只考虑了硬件上限实际工程里经常被软资源卡脖子最常见的是寄存器和共享内存。假设saxpy的kernel每线程用了40个寄存器。RTX 3090每个SM有65536个寄存器那么65536÷401638也就是每SM最多只能同时容纳1638个线程的寄存器上下文。如果blockSize2561638÷2566.4向下取整就是6个block总共1536个线程。于是occupancy占用率就从100%掉到1536÷204875%。也就是说block到SM的映射数量不是写死的而是硬件上限和资源约束共同决定的动态结果。寄存器用得越多SM能驻留的block就越少共享内存同理。假设每个block申请40KB共享内存而SM的共享内存总共128KB那最多只能驻留3个block3×40120KB4×40160KB超了。如果blockSize2563个block只有768个线程驻留占用率直接掉到37.5%。这也是为什么共享内存用多了性能反而断崖式下跌。3.4 为什么开发社区常见的BlockSize是128、256、512理解了上面的账就不难明白为什么CUDA社区里大家默认选128、256、512这些值而很少看到100、200、1000这种顺手的数字。第一个原因是warp大小为32。blockSize128正好是4个warp256是8个warp都是整数个warp不会产生残缺warp。blockSize100的话会被切成3个完整的warp96线程加1个只有4线程的残warp后面详说残warp的代价。第二个原因是调度选择空间。SM的warp调度器需要在众多warp中挑就绪的下一个执行驻留的warp越多调度器越不容易因为访存等待而空转。blockSize256时每个SM驻留8个block提供64个warp调度器有非常充足的选择余地。blockSize32时每个SM最多驻留32个block也就是32个warp明显少了一半。第三个原因是同步和资源粒度。block太大比如1024每SM只能驻留2个blockblock内的__syncthreads()会阻塞大量线程等待寄存器分配也不灵活block太小比如32又会让block数量迅速撞上每SM 32个block的上限同样浪费。128到512之间是一个经过广泛验证的甜区256尤其常见8个warp、资源占用适中、边界条件好处理。4. 真正被执行的是Warp给对应关系补上最关键的一环4.1 SIMT32个线程同一条指令前面反复提到warp现在把这一层彻底讲透。Warp是GPU硬件调度和执行的最小单位由32个线程组成。一个warp里的线程执行同一条指令但每个线程操作自己的数据这套执行模型叫SIMTSingle Instruction, Multiple Thread。你可以把warp理解成一列并排前进的士兵指挥官喊一声齐步走一条指令32个士兵同时迈出左脚各自处理各自的数据。如果某个士兵因为条件分歧if-else需要走不同的路径那这个warp就得分别走两条分支每条分支上的线程各执行一次另一部分线程被掩蔽掉。这就是所谓的线程发散thread divergence它会让warp的执行效率打折。在硬件层面一个warp的32个线程对应SM里一组32个并行执行单元lane。RTX 3090的SM有128个SP可以理解为4组32-lane的执行阵列每组由一个warp调度器掌管。每个周期4个调度器各选一个就绪的warp发射到对应阵列上执行。所以一个SM每个周期最多能同时执行4个warp的指令也就是4×32128个线程的标量指令——正好对应128个SP的数量。4.2 隐藏延迟的真相为什么驻留的warp数这么重要GPU性能优化的核心说穿了就两个字延迟。一条内存访问指令从发出到数据返回可能需要几百个时钟周期。如果SM里只驻留了4个warp这4个warp都在等内存数据回来那SP阵列就全部空转性能惨不忍睹。正确的逻辑是SM里驻留的warp数量越多调度器在某个warp等待时切换到其他就绪warp的机会就越多。GPU正是靠这种线程级并行来把访存延迟、指令延迟掩盖掉让SP始终有活干。这也是为什么提高occupancy即驻留warp数占最大容量的比例是CUDA优化最常见的手段。从这个角度回头看blockSize的选择blockSize64每SM驻留32个block、64个warpblockSize256每SM驻留8个block、64个warp。两者在RTX 3090上的occupancy一样但block数量差4倍。block多的好处是block级调度更灵活一个block结束另一个block马上顶上坏处是block切换也有开销而且64线程的block里只有2个warp如果这个block里出现访存停顿调度器的选择范围局限在同一个SM的其他block里。工程实践普遍认为256比64更稳就是这个道理。4.3 Block内线程数为什么最好是32的倍数现在可以回答那个经典问题了为什么block内线程数最好取32的倍数因为GPU切warp是按连续32个线程一组的如果blockSize100切出来的结果是3个满warp0到95号线程加1个只有4个线程的残warp96到99号线程。残warp的28个lane执行通道虽然不参与计算但执行指令时照样占一个warp的发射槽位。也就是说这100个线程实际占了4个warp的执行资源有效利用率只有96÷12875%。更要命的是如果kernel里有__syncthreads()残warp也必须参与同步因为同步的粒度是block内所有线程空转的lane同样要到位。所以两个经验直接给出第一blockSize选32的整数倍第二如果数据量不是32的倍数宁可多开几个线程用if (i n)做边界保护也别让warp残缺。这是每个CUDA项目都该无脑执行的规矩。4.4 那么THREAD与SP的对应到底怎么描述才准确绕了一大圈回到标题的核心问题。准确的链条是这样一句话一个Thread最终由SM内的某个SPCUDA Core执行它的指令但Thread和SP之间不是一对一的固定绑定而是以Warp为单位、由调度器动态分配的临时占用关系。Thread被组织成WarpWarp被调度器发射到SP阵列上一个warp的32个线程在32个执行单元上同时执行同一条指令执行完这一个周期下个周期同一个SP可能执行另一个warp的指令。所以THREAD和SP的关系更像顾客和服务窗口顾客线程数量可以成千上万服务窗口SP只有一百多个顾客必须排队成组warp被叫号叫到谁谁就临时占用窗口。SP不是某个线程的私有财产而是被所有线程时分复用的。这也是为什么一个grid里可以声明几百万个线程而GPU物理上只有几百上千个SP——它们根本不是一一对应的关系靠的是线程池排队调度。5. 实践中的对应关系启动配置、调优与排错的经验5.1 用Occupancy的思路反推启动配置理解了映射关系之后配置kernel启动参数就是一个根据硬件上限倒推的过程。以RTX 3090为例忽略寄存器等软资源不同blockSize对应的理论占用率如下blockSize每SM可驻留block数每SM驻留线程数理论占用率6432受block上限限制2048100%128162048100%25682048100%51242048100%102422048100%看起来全是100%但一旦叠加寄存器或共享内存的约束差距就出来了。还是每线程40个寄存器的例子blockSize128时1638÷12812个block驻留1536线程75%blockSize512时1638÷5123个block驻留1536线程也是75%而blockSize256时同样是6个block、1536线程。三种配置的occupancy一样但block粒度不同实际性能会因调度开销、同步粒度而有差异这就是需要实测的原因。NVIDIA官方提供的Occupancy Calculator在Nsight Compute里能直接算出给定blockSize和资源用量下的期望占用率强烈建议每次调kernel都跑一遍。它最大的价值不是帮你自动选参数而是让你看到是什么在卡占用率——是线程上限、block上限、寄存器还是共享内存对症下药远比瞎试参数高效。5.2 常见的启动配置错误与定位这几年我踩过的坑归纳起来就三类基本覆盖了绝大多数配置相关的故障。第一类启动直接报invalid configuration argument。这种错误九成是blockSize超过了设备上限比如1024或者gridSize、blockSize某个维度设成了0少数情况是gridDim超过2^31-1的上限。排查方法很简单启动前打印一下blockIdx、blockDim这些值或者调用cudaGetLastError()拿到返回码顺着报错信息十秒就能定位。注意kernel这个语法本身是异步的启动错误不会立刻暴露必须在启动后加cudaDeviceSynchronize()才能把错误真正捕获出来。第二类程序能跑但结果不对而且时好时坏。这种情况八成是索引越界或数据漏算。比如把gridSize n / blockSize写成了整除而不是向上取整当n不能被blockSize整除时末尾几个元素就没人处理反过来如果没用if (i n)做边界检查多余的线程会把数组后面不属于它的内存也写一遍轻则数据错乱重则直接导致设备崩溃比如非法内存访问报unspecified launch failure。我在实际项目里的经验是所有涉及数组的kernel都保留边界判断哪怕你确信数据量是blockSize的整数倍也留着成本极低救命次数极多。第三类程序和结果都对但性能远低于预期。这时候别急着调代码先看occupancy。用Nsight Compute看Achieved Occupancy如果长期低于70%要么是寄存器用太多要么是共享内存用太多要么是blockSize设得不符合硬件胃口。有一次我把一个kernel的每线程寄存器压到32个以内占用率从50%提到100%性能直接翻倍而代码逻辑一行没改只改了一个编译限制选项。这类问题靠肉眼是看不出来的必须用工具量化。5.3 几条可以直接用的经验法则结合上面的分析和踩坑经历总结几条可以直接抄走的经验法则blockSize默认取256除非有明确理由。256是32的8倍、8个warp在调度灵活性和资源开销之间平衡最好。数据量小可以用128kernel里shared memory很大时可以降到128或64但别低于64。让grid中的block数量至少是SM数量的几倍。一个简单标准block数至少是SM数的4到8倍。如果82个SM的卡上只开了82个block一旦某个SM分配不均空闲SM就出现了block数远超SM数才能靠动态调度抹平负载差异。在kernel启动后统一加一次cudaGetLastError()检查。CUDA错误是异步的不加这一句错误会在下一次同步点才炸出来定位成本高得多。工程上可以封装一个宏或工具函数统一处理错误返回。优先调资源再调blockSize。寄存器超标就先压寄存器__launch_bounds__或-maxrregcount共享内存超标就减每线程用量这两样是occupancy的隐形杀手比纠结blockSize是256还是512重要得多。每次改完配置用Nsight Compute看一眼Achieved Occupancy和Warp Stall Metrics让数据说话。背下来一张卡的特性参数不如养成算一遍、测一遍、调一遍的习惯。最后再分享一个小技巧。我刚开始写CUDA那阵子特别迷信block越大吞吐越高把blockSize设成1000结果性能比256差一大截Nsight一看占用率被寄存器卡在了40%。后来每次调kernel我都先跑一遍设备信息查询工具deviceQuery看自己卡片的SM数量、每SM线程上限这些硬参数再手算一遍映射账最后用Nsight验证。这套流程走熟了之后基本没有出现不知道block数量应该设多少的迷茫。建议你也把这个方法存成自己的检查清单写kernel之前先算账就能少走我当年那段弯路。
