很多刚接触CUDA的人都会背这么一句GRID是网格BLOCK是线程块THREAD是线程SM是流式多处理器SP是流处理器。背得滚瓜烂熟面试官随口加问一句“那它们到底怎么对应起来的”就支支吾吾了。我当年也是这样把名词都记得清清楚楚却被一句“为什么一个kernel能开几百万个线程而GPU上显存的SP数不过几千个”问得当场卡壳。后来自己翻架构文档、调kernel踩坑才把这块彻底理顺。这篇文章就把这套对应关系从头到尾拆开讲一遍哪些是软件层的抽象哪些是硬件层的物理单元block怎么被塞进SMthread又是怎么跑到SP上去执行的以及在决定block尺寸和grid大小时这些对应关系怎么直接影响你的程序能跑多快。1. 先看全景软件层的线程组织和硬件层的执行单元1.1 软件层GRID、BLOCK、THREAD是什么关系CUDA的线程组织是个三层嵌套结构。最外层是GRID一个GRID对应一次kernel启动里面装着若干个BLOCK每个BLOCK里面又装着若干个THREAD。写代码的时候用这样的形式启动kernelkernelgridSize, blockSize(args);gridSize是dim3类型表示这一个GRID里有多少个BLOCKblockSize也是dim3类型表示一个BLOCK里有多少个THREAD。可以是1D、2D、3D比如kernel1024, 256表示GRID里有1024个BLOCK每个BLOCK里有256个THREAD总共就是1024乘以256等于262144个线程。BLOCK里的线程是允许互相协作的。协作手段是__shared__共享内存和__syncthreads()同步屏障。什么意思呢同一个BLOCK里的线程可以把中间结果写进一块共享内存里然后互相看对方算出来的数据但这些协作只在BLOCK内部有效跨BLOCK没法直接共享内存也没有轻量级同步原语。这段设计背后其实有两个非常实际的考虑。第一是资源局部性共享内存在硬件上离计算单元非常近访问速度远快于全局内存所以把可以局部化的数据放在BLOCK内部协作效率最高。第二是调度问题BLOCK是硬件能调度的最小单位一个BLOCK固定被分到一个SM上整个BLOCK就在那个SM上跑完不会中途被拆走。这样共享内存和同步才有物理基础——都在同一个SM上。这个逻辑可以类比成工厂生产。GRID是一张大订单BLOCK是一个车间THREAD是车间里的工人。一个订单可以同时让很多车间开工一个车间里有很多工人互相打配合但订单都统一由一个调度室分派车间内部的事不会跨车间商量。1.2 硬件层SM和SP的物理身份软件层的GRID、BLOCK、THREAD只是程序员眼里的抽象模型GPU硬件本身根本不知道什么是GRID也不认识BLOCK。硬件认识的是SM和SP。一块GPU上分布着几十上百个SMStreaming Multiprocessor流式多处理器。每个SM才是一个真正完整的“计算单元”它有自己的指令调度器、寄存器文件、共享内存、缓存和一堆计算单元。这堆计算单元就是SPStreaming Processor流处理器也就是我们常说的CUDA core。NVIDIA的每一代架构里SM的内部结构都有变化但SP总是那一堆最简单、数量最多的算术执行单元。比如Pascal架构的GTX 1080 Ti一个SM里塞了128个SPTuring架构的RTX 2080 Super一个SM里是64个SPA100的GA100核心一个SM里也有64个SP。有一个常见误区是拿CPU去类比SM像CPU的核SP像CPU核里的ALU。这个类比方向对但别太较真。CPU一个核跑一个线程靠复杂的分支预测和乱序执行压榨性能GPU的SM靠大量SP并行执行批量的简单线程单个线程的智能程度很低但是线程数可以堆到成千上万靠吞吐量取胜。而SP本身的职责非常纯粹执行一条算术指令比如算个浮点加法、整数乘法。它不关心线程是怎么调度的只负责计算。谁把指令发给它它就算谁的数据。2. 对应关系拆解从线程到SM再到SP完整走一遍2.1 一个kernel从启动开始硬件里发生了什么把五层概念串起来的关键是搞清楚一条完整的调度链路。我直接用实际流程说。第一步host端调用kernelgridSize, blockSize驱动把这个启动请求和kernel的代码一起提交给GPU。第二步GPU内部的硬件线程块分配器在不少文档里叫GigaThread Engine登场。它把GRID里的BLOCK一个个挑出来分发给各个SM。硬件会尽量把BLOCK均匀分布在各个SM上但程序员无法控制某一个BLOCK到底去哪颗SM。第三步SM收下BLOCK之后把BLOCK里的THREAD按每32个一组切分成warp。warp是GPU执行指令的真正基本单位注意不是thread是warp。BLOCK尺寸如果不是32的整数倍最后一组会凑不满32空出的lane闲置。第四步SM里的warp scheduler每个时钟周期从一堆可调度的warp里挑一个把它当前要执行的指令发射出去。发射的时候这条指令会被同时分发到SM里的一组SP上。一个warp有32个线程所以这条指令实际上是在32个SP上并行执行每个SP处理一个线程的数据。所以最终落到硅片上的景象是一个SM同时装着几个BLOCK每个BLOCK被切成了若干个warpwarp由调度器排队裁决胜出的warp把指令派给SP去算。这整个链路里没有哪个环节是程序员能干预的。BLOCK到哪个SM、warp什么时候被发射、指令在哪个SP上执行全是硬件自动完成。2.2 五个概念一张表看清对应关系概念所属层次本质是什么和对方的对应关系GRID软件/线程层次一次kernel启动产生的全部线程集合一个GRID对应一次kernel调用BLOCK软件/线程层次GRID中的线程块块内线程可协作一个BLOCK被整体调度到某一个SM上执行THREAD软件/线程层次最小执行单位执行同一个kernel函数线程被切进warp由SP执行指令SM硬件执行单元GPU里的物理计算单元拥有独立寄存器和共享内存同时执行多个BLOCK内部管理若干warpSP硬件执行单元SM内部的流处理器也叫CUDA core具体执行某条指令的一个线程的数据这个表是全文的总纲后面所有内容其实都在解释这张表。有一个点必须特别注意表格里BLOCK对应SMTHREAD对应SP看似是两对映射但对应级别完全不同。BLOCK到SM是“完整的、不可分割的分配关系”带排他性的THREAD到SP是“执行时的一次性问题”不存在哪个线程绑定哪个SP。线程只是某个时刻被调度执行在某个SP上下一时刻可能换到另一个SP上虽然实际硬件里通常是固定lane但软件视角上不要理解为绑定。2.3 为什么BLOCK必须完整地待在一个SM上这是理解整套对应关系的关键约束值得单独说透。一个BLOCK在运行期间内部线程会通过共享内存交换数据还会用__syncthreads()做同步。如果BLOCK被拆开分给两个SM跑那这两拨线程的共享内存就物理隔离了互相看不见同步也要跨SM做而GPU的硬件原语里根本没有跨SM屏障这种东西还得靠全局内存在软件层面模拟性能直接崩掉。所以硬件把BLOCK设计成调度的最小分配单位——要分配就整个分给一个SM。这也带来一个限制一个SM能同时装多少个BLOCK不仅要看SM自己的线程上限还要看寄存器、共享内存等资源够不够这些BLOCK分配。还有个程序员容易误解的点既然硬件分配的是BLOCK那是不是把BLOCK尺寸设大一点SM就能多接任务不是的SM能同时容纳BLOCK的数量有硬上限常见是32个能同时容纳的线程数也有硬上限常见是2048个。BLOCK太大会让资源占用过快导致SM上并发块的数目反而减少总并发warp数量下降后面性能就会出问题。3. 用真实GPU算一笔账资源和性能如何被对应关系锁死3.1 四个决定生死存亡的资源上限既然BLOCK整体占SM那任何意义上的“规划线程规模”本质上都是在算资源账。每次启动一个kernel之前最好心里有这几条线。第一每SM线程数上限。典型的Pascal到Ampere架构都是2048个线程也就是64个warp。超过这个值的BLOCK压根不会被分配进去。第二每SM并发BLOCK数上限。通常是32个不管你的BLOCK是不是小得只有64个线程。第三寄存器文件大小。比如GTX 1080 Ti的每个SM有65536个32位寄存器。所有并发线程合起来最多只能用这么多。一个线程占的寄存器越多能并发的线程就越少。第四共享内存容量。典型值是每个SM96KB或164KB。BLOCK里申请的总共享内存超过这个量SM就只能装更少的BLOCK。任何一条爆了都会直接影响你程序的占用率。下面拿一块具体GPU算一笔账。3.2 实例计算GTX 1080 Ti上的资源分配拿GTX 1080 Ti开刀。它的参数是28个SM每个SM128个SP每SM最大2048线程最大32个BLOCK65536个寄存器共享内存96KB。假设我写了这样一个kernel__global__ void example_kernel(float* data) { int tid blockIdx.x * blockDim.x threadIdx.x; data[tid] data[tid] * 2.0f; }然后启动参数是example_kernel128, 256。总线程数是128乘以256等于32768。整个GPU能同时容纳的最大线程数是28乘以2048等于57344所以从总量上看一次全部装下没问题。但每个SM分几个BLOCK这是调度器自己决定的。理想情况是28个SM128个BLOCK每个SM分到四五个BLOCK每个SM上的线程数在1024到1280之间占用率50%以上已经不错了。再换一个每线程吃寄存器的例子。假设某个kernel每线程需要48个寄存器。每个SM有65536个寄存器满并发2048线程需要2048乘以48等于98304个寄存器超了。所以每SM最多能容纳的线程数按寄存器算只有65536除以48等于1365再向下取整到warp边界就是1344个线程。如果BLOCK尺寸设为256每个BLOCK需要256乘48等于12288个寄存器。65536除以12288等于5.33所以每SM最多塞下5个BLOCK也就是1280个线程占用率1280除以2048等于62.5%。有人会想那我把BLOCK尺寸降到128是不是能塞更多线程算一下每个BLOCK需要128乘48等于6144个寄存器65536除以6144等于10.66最多10个BLOCK总共还是1280个线程。发现没寄存器总量是硬约束改变BLOCK大小只是改变了BLOCK数量总线程数被寄存器卡死依然是62.5%占用率。如果程序里还申请了共享内存比如每个BLOCK申请20KB那每SM最多96除以20等于4个BLOCK。再叠加线程数限制就要看BLOCK尺寸大小了BLOCK是256线程的话4个BLOCK总共1024线程占用率50%BLOCK是512线程的话4个BLOCK总共2048线程刚好满占用。这个案例里反而大BLOCK更能充分利用SM。从这个计算能体会到一个重要推论不存在万能的固定参数每次都要对着具体kernel的资源消耗算一遍。3.3 占用率不是越高越好warp调度才是幕后黑手算占用率的意义在哪里在于隐藏访存延迟。GPU的线程执行有个特点当某个warp里的线程访问全局内存要等几百个周期的数据返回这时如果SM里只有这一个warp在跑那就只能干等如果SM里同时躺着几十个warp调度器会立刻切换到一个不需要等内存的warp上去执行。寄存器全部存在SM本地切换warp不需要保存和恢复现场基本零开销。这也是GPU能容忍高延迟的原因不是因为它没有延迟而是它有足够多的并发warp把延迟掩盖掉。所以一般来说占用率越高越有可能把访存延迟藏住SM上的活跃warp越多调度器每次选warp时的选择空间也越大。但是占用率真的越高越好吗不一定。为了强行把每线程寄存器压到很低编译器可能会把本来放在寄存器里的数据溢出到local memory而local memory本质上是全局内存访问速度慢得惊人。这种“溢出换占用率”的操作常常得不偿失。实际调优应该用Nsight Compute看真实的stall原因和吞吐指标而不是盯着占用率这个数字自我感动。4. 高频问题与实战排错记录4.1 SP和THREAD到底是不是一一对应这个问题几乎是每次交流必被问到我这里明确回答不是。SP是物理执行单元THREAD是逻辑执行体。一个SP在同一时刻只执行一个线程的一条指令但一个线程不绑定任何特定SP。GPU上几千个SP却要执行几十万甚至几百万个线程靠的是时间片轮转。warp调度器把一个个warp送上SP执行几条指令后又切换到别的warp。这和操作系统在单核CPU上切换进程是一个道理区别在于GPU切换warp的开销几乎为零。顺带说一句SIMT的“单指令多线程”模型让一个warp里的32个线程在执行同一行代码。如果warp里出现了分支分叉比如一半线程走if另一半走else硬件会把这两拨人分别执行这个现象叫warp divergence。所以写代码时尽量保证同一个warp内的线程走向一致别让32个线程背着不同的数据乱跑。4.2 BLOCK尺寸到底设多大合适面试和实际项目里最常被追问的配置问题。没有标准答案但有靠谱的参考区间。BLOCK太小比如设成64每个SM能塞32个BLOCK但每个BLOCK只有2个warpSM上总共也就64个warp好像已经满了。可问题是BLOCK尺寸小会导致任务切分过碎很多资源被block的管理结构白白占用负载均衡也容易被极端情况扰乱。BLOCK太大比如1024每个BLOCK占32个warp如果每线程寄存器吃得狠SM只能装下一个BLOCK那整个SM的warp数量可能远低于64访存延迟掩盖能力大打折扣。实践里128到512之间的BLOCK尺寸最常见尤其是256。它兼顾了warp数量、资源弹性、负载均衡三件事。需要让BLOCK的内部线程协作复杂的时候尺寸倾向512甚至1024访存密集、共享内存用得少的128更灵活。还有个优先级更高的原则优先保证SM上有足够的并发warp比如48到64个然后再去微调BLOCK尺寸。4.3 报错out of resources怎么定位运行时报too many resources requested for launch或者cudaErrorLaunchOutOfResources本质就是启动参数超过了硬件上限。我排查这类问题的顺序一般是这样的第一步查BLOCK尺寸和GRID尺寸。看看是不是单个BLOCK申请了超过1024个线程或者总共发出来的BLOCK数超过了grid上限。第二步查共享内存。cudaFuncSetAttribute或者kernel内部的动态共享内存是不是申请了大于SM共享内存容量。第三步查寄存器。每线程寄存器占用可以通过编译时加--ptxas-options-v看到。如果某个kernel每线程用了超过64个寄存器那它天然就不可能满占用率有时候还会直接超限。第四步如果确认是寄存器爆了用__launch_bounds__(maxThreadsPerBlock, minBlocksPerSM)告诉编译器要节省寄存器或者用maxrregcount整体压一下。这个报错还有个常见变种不是在launch时报而是在cudaMemcpy或cudaDeviceSynchronize时报。别被误导了错误源头就是之前的kernel launchAPI调用是同步暴露出错误而已。4.4 二维三维的时候索引别搞混很多人写二维block时栽在索引上。归纳一个通用公式int blockId blockIdx.x blockIdx.y * gridDim.x blockIdx.z * gridDim.x * gridDim.y; int tid blockId * blockDim.x threadIdx.x;如果是二维grid加上二维block比如处理图像像素通常做法是int x blockIdx.x * blockDim.x threadIdx.x; int y blockIdx.y * blockDim.y threadIdx.y; int index y * width x;把二维索引换算成一维数组下标时一定要想清楚数组行优先顺序。这地方错了不是报错而是数据错得神不知鬼不觉排查起来非常费劲。我的经验是写完后先用小尺寸跑一遍把index打印出来和CPU串行版本对比确认无误再放大数据。另外还有个搜索引擎带来的小插曲很多人搜grid最后翻到的是CSS里的grid布局那个做网页排版的GRID和CUDA的GRID除了英文单词一样没有任何关系。别把display:grid和kernelgrid, block搅在一个脑回路里。我个人在实际调试中的体会是不要光靠感觉去改block和grid参数。拿一张纸把每线程寄存器数、共享内存申请量、SM容量写下来手算一遍每个SM能装几个block、几个warp再上机器验证一趟下来你对这套对应关系的理解会比背十遍概念都扎实。尤其碰到性能问题时这种“先纸上计算到底层映射再归因到具体瓶颈”的思路比瞎猜参数能省出好几个通宵。
