手里这台卡跑一个纯访存的内核理论带宽标称一点五 TB/s 往上实测在 55% 到 60% 之间晃profiler 里issue_active只到 20% 出头warp state 那一栏long_scoreboard占了大半——这种局面我遇到过不止一次。多数人第一反应是内存是不是没对齐是不是没向量化但把访存指令改到完美合并之后仍然卡在同一个位置问题就落在 SM 的 warp 调度器身上了线程块发出去了、warp 也驻留了可是调度器每个周期挑不出足够的就绪 warp 来填空访存延迟就那么明晃晃地挂在流水线上没人接。这篇东西聊的就是这条链路——CUDA 的 SM 内部warp 调度器到底按什么规则挑 warp访存延迟又是靠什么被藏起来的以及怎么用可复现的实验把这两个抽象概念变成 profiler 上的数字。适合已经能写能跑 kernel、但性能调优基本靠猜的同学也适合把 occupancy 当成唯一调优旋钮、结果发现调到 100% 还是不提速的人。我不打算复述官方文档而是按我实际的排查顺序走一遍先认角色再算账然后动手跑对照实验最后把踩过的坑摊开。全程以 Ampere 这一代的结构为主前后代有差异的地方我会单独点出来。1. 先把角色认清楚SM 里到底谁在干活1.1 一个 SM 的硬件骨架长什么样要谈调度得先知道被调度的对象住在哪。一个 SMStreaming Multiprocessor可以粗略理解成四个小核加一堆共享资源四个处理块官方叫 sub-partition 或 processing block每个处理块里有自己的 warp 调度器、自己的指令派发单元、自己的通用计算单元组以及一组共享的寄存器文件、L1/共享内存、加载存储单元和特殊功能单元。关键在于一个数字每个处理块每个时钟周期最多派发一条 warp 指令。四个处理块加一起就是每个 SM 每周期最多四条 warp 指令。这个4是个硬约束很多性能模型算不对就是因为忽略了它——你把 occupancy 拉满是没用的派发宽度就摆在那里超过它之后的任何并行度都只能排队。驻留容量方面以 Ampere 的计算卡为例一个 SM 最多驻留 64 个 warp2048 线程、65536 个 32 位寄存器、约 256KB 的寄存器堆总量、最多 32 个 block、16 个 barrier 资源。消费级那一代比如 GA10x 那批最大驻留是 48 个 warp / 1536 线程寄存器总量一致。这些数字必须记牢因为它们决定了 occupancy 的天花板而 occupancy 又直接决定调度器手里有多少牌可以打。寄存器是这里最容易翻车的资源。它按线程粒度分配粒度是 8 个寄存器一档编译产物里regs/thread会向上取整到 8 的倍数。假设某个 kernel 每个线程用了 100 个寄存器向上取整到 104那么 2048 个线程需要 212992 个寄存器超过 65536 的限制硬件就会把驻留线程数压到 1536再算下来 warp 数只有 48occupancy 直接掉到 75%。很多人抱怨我 block 设成 256 线程怎么只能跑 6 个 block答案往往就在这里而不是在什么玄学地方。1.2 为什么是 32warp 这个粒度的由来与代价一个 warp 固定 32 个线程它们共享同一个程序计数器、同一份指令以 SIMT 的方式同步推进。为什么偏偏是 32 而不是 16 或 64这是延迟与吞吐的折中。粒度太细调度器每周期要维护的 warp 条目太多选择逻辑和上下文状态的成本压不住粒度太粗一旦某个 32 线程组里出现分支分歧或者访存不合并浪费的比例就成倍上升。32 恰好是当时存储子系统一个 cache line 的 128 字节除以单线程 4 字节的宽度访存合并的账算起来最整齐。代价也很直接warp 是调度的最小单位也是浪费的最小单位。if 分支里两个路径各走一遍整个 warp 的执行时间就是两条路径之和32 个线程里如果只有 1 个线程真正需要走某段计算其余 31 个线程照样占着执行单元的槽位只是结果被掩码掉。所以写 kernel 的时候让同一个 warp 里的 32 个线程尽量走同一条控制流、访问连续的地址这不是优化技巧而是省掉本不该付的成本。另一个容易被忽略的点warp 内部的线程不是完全独立的。只要指令序列里有任何一条是全部 32 个线程都执行的硬件就会把它们压成一条 warp 指令去跑同一个执行单元阵列。这意味着你在写代码时看到的一个线程一次加法落到硬件上可能是一个 warp 一次 32 路加法占用的是同一个指令槽。指令计数的口径差异就是很多人的手算模型和实测数据对不上的根源。1.3 一个 block 从提交到跑起来要过几道关搞清楚从grid, block到真正开始跑指令之间发生了什么排障的时候会省很多时间。流程大致是Grid 分发grid 里的 block 按某种顺序分给各个 SM顺序没有强保证不要依赖 blockIdx 的物理分配顺序。资源预留SM 检查目标 block 需要的寄存器、共享内存、barrier 槽位、线程槽位是不是都够。任何一项不够block 就在队列里等。线程切分block 内的线程按threadIdx线性编号每 32 个切成一个 warp。这里有个细节——切分是按 block 内的线性 tid 来的不是按多维坐标所以dim3(32, 8)的 block 和dim3(256)的 blockwarp 的构成方式是一样的。分配 warp slot每个 warp 被塞进某个处理块。分配策略没有公开但从微基准看基本是轮转的目的是让四个处理块的负载均衡。进入就绪队列warp 的 PC 指向 kernel 入口此后就进入调度器的候选池。踩坑提醒如果共享内存用得特别狠比如一个 block 申请 100KB 以上会出现block 能放下但只能放一个的情况SM 里其他处理块空转调度器手里只有一个 block 的 warp 可以挑。这时候调 occupancy 的旋钮是没用的得从 block 尺寸和共享内存用量下手。还有一点block 的调度顺序会影响结果。同一份 kernel如果 block 之间的负载天然不均比如按数据分布决定循环次数先被调度的 block 跑得快、后被调度的慢最后总时间被慢的那个拖住。这不是调度器的问题是负载划分的问题但表现出来像是调度不公。实践中我倾向于让 block 的工作量尽量均匀或者用持久化 kernelpersistent kernel的方式手动接管任务分配。2. warp 调度器每周期到底在挑什么2.1 就绪判定从 PC 到 issue slot 的链路每一拍四个处理块各自的调度器都要做一次选择从住在自己这里的一堆 warp 里挑一个把它的下一条指令送进派发单元。这个选择过程公开资料只给了轮廓从微基准能反推出大致是这么几层判断第一层是能不能发。一个 warp 只有在满足以下条件时才算就绪PC 指向的指令已经取到所有源操作数都已就位没有被标注为等待中不处于 barrier 等待或分支解析中没有被 MIO 队列或长延迟队列的背压卡住没有在执行__syncthreads之类的同步语义。任何一条不满足这个 warp 这一拍就直接出局。第二层是谁先发。剩下的就绪 warp 里要选一个这就是策略部分。Fermi/Pascal 时代比较明确是 GTOGreedy Then Oldest——优先挑同一 warp 的连续指令实在不行挑最老的。Volta 之后官方不再详细说明但从行为上看它也不是简单地轮转而是带有某种年龄加权并且会参考目标执行单元的当前压力。判据的核心其实是让各条流水线都别空着而不是对每个 warp 公平。第三层是发去哪。选中的指令还要看目标执行单元这一拍有没有空位。如果目标 pipe 正忙指令会被按住产生的就是math_pipe_throttle或mio_throttle这类背压统计。这就解释了一个常见现象明明有很多 warp 就绪issue_active却上不去——不是挑不出人而是执行单元接不住。操作数的已就位是怎么判断的靠 scoreboard。每个 warp 有一组记分牌位指向那些结果还没写回来的寄存器。一条指令要发射前硬件会检查它的源寄存器有没有被标记为等待。等待来自固定延迟执行单元比如共享内存、某些特殊运算的叫 short scoreboard等待来自全局内存加载的叫 long scoreboard。这两者在 profiler 里是两个独立字段含义完全不同——short 一般意味着 MIO 排队long 才是真正的 DRAM 延迟。这里有个反直觉的地方long scoreboard 高不一定是坏事。如果访存请求确实在飞、带宽也确实吃满了那么 warp 停在 long scoreboard 上等结果是正常的、不可避免的。只有当带宽利用率也低的时候long scoreboard 才是问题信号说明并发度不够延迟没被藏住。2.2 四个子分区与派发宽度别把 SM 当成一个整体新手最容易犯的错是把 SM 看成一个大池子算 occupancy 的时候用 SM 总量除以线程数。实际上资源是分到四个处理块上的寄存器、warp slot、执行单元都是分区资源。这带来两个很实际的后果。第一occupancy 的瓶颈可能只在单个分区上。比如 block 有 256 线程切成 8 个 warp。如果 SM 能放 6 个这样的 block一共 48 个 warp理论上每个分区 12 个。但如果硬件的分配策略让某个分区恰好分到了 13 个、另一个只有 11 个那个 13 个的分区就成了新的瓶颈。这在实测里表现为两个完全一样的 kernel仅仅因为 block 尺寸从 256 改成 128issue 率就有可见差别。原因就是资源分配的粒度变了四个分区的均衡程度变了。第二派发宽度是硬上限。前面说过每分区每周期一条指令四个分区加一起四条。所以一个 SM 想跑满理论算力前提是四条流水线同时有活干。假设某个 kernel 的指令流里有 50% 是访存指令、50% 是计算指令那么当访存指令卡在 LSU 排队时计算指令就得顶上反过来也一样。每条流水线的吞吐要匹配否则最快的那个也会被最慢的拖住。关于双发射这个说法我得泼点冷水。一些架构在特定指令组合下确实存在同一周期派出两条指令的情况比如一条浮点加一条整数但从性能建模的角度按每分区每周期一条来估算误差通常在可以接受的范围内。真要确认跑个微基准看issue_active的上限就知道——如果它稳定在 100%说明模型对上了如果偶尔超过说明有特例。2.3 stall 名字背后的硬件事件对照profiler 的 WarpStateStats 那一栏一堆stalled_xxx字段刚开始看特别懵。我把它按我该去改什么重新分了个类比按字母顺序看有用得多。字段硬件含义优先怀疑的方向stalled_long_scoreboard等全局/本地内存加载结果并发度不足、访存不合并、L2 未命中stalled_short_scoreboard等共享内存、特殊函数等固定延迟结果共享内存 bank 冲突、MIO 排队stalled_wait等固定延迟指令如某些 ALU的结果依赖链过长ILP 不足stalled_not_selectedwarp 已就绪但没被调度器选中说明不缺并行度缺的是派发宽度或执行单元吞吐stalled_mio_throttleMIO 指令队列满共享内存/特殊函数指令过密stalled_lg_throttle全局内存指令队列满访存指令发出速率超过 LSU 消化能力stalled_math_pipe_throttle目标计算流水线繁忙指令类型过于单一某条 pipe 打满stalled_barrier等__syncthreads或 block 间同步block 内负载不均stalled_no_instruction取指失败指令缓存未命中内核代码体积过大、分支过于分散stalled_drain内核退出前等待内存操作完成store 太多、退出前没做同步这张表的价值在于它能直接切分问题not_selected高说明并行度够了、算力不够long_scoreboard高说明并行度不够、内存延迟没藏住。这两个方向的处理方式完全相反——前者要减少并行度省下的资源换更高的单线程效率后者要增加并行度或缩短依赖链。我见过有人看到not_selected很高就去加 occupancy结果越加越慢就是因为方向反了。另外提一句stalled_no_instruction。这个字段一般很小但如果它占比异常通常意味着内核的指令体积太大、超出了指令缓存的容量。处理办法是把循环展开关掉#pragma unroll 1、把大函数拆小或者在编译时限制展开的深度。这是那种不看字段永远想不到的问题。3. 把访存延迟算清楚为什么需要那么多 warp3.1 延迟隐藏的本质是填洞先说清楚延迟隐藏这个说法到底在说什么。一条全局加载指令发出后到数据真正回到寄存器中间要经过 L1 查询、L2 查询、内存控制器排队、DRAM 阵列访问、数据回传这一整段是几百个时钟周期。在这几百拍里发出这条指令的 warp 什么也做不了它的记分牌位被置住只能干等。如果 SM 里只有这么一个 warp那么这几百拍就是纯粹的空转执行单元全闲。延迟隐藏的意思就是在这几百拍里找别的活给执行单元干。别的活从哪来要么是同一个 SM 里的其他 warpTLP线程级并行要么是同一个线程里另外那些不依赖这个加载结果的指令ILP指令级并行。所有的调优手段归根到底都是在扩大这两个来源中的一个或两个。这个视角很重要因为它告诉你occupancy 从来不是目标它只是获取 TLP 的手段之一。如果一个 kernel 靠 ILP 就能把每个 warp 的访存间隙填满那么它根本不需要高 occupancy反过来如果一个 kernel 的每个线程都在等同一个加载结果那再高的 occupancy 也只是把更多 warp 排进等待队列而已。3.2 用 Littles Law 反推需要多少在飞请求想吃得满带宽光有并发线程不够得有足够的在飞的访存请求。这个可以用排队论里的 Littles Law 直接算在飞请求数 带宽 × 延迟拿一个具体例子算。假设某张计算卡的理论带宽是 1.6 TB/sDRAM 往返延迟按 500 纳秒算这个数量级在 GDDR/HBM 上都差不多具体随架构和频率浮动在飞字节数 1.6e12 B/s × 500e-9 s 800,000 B ≈ 780 KB也就是说为了让这条内存通路不空转整个芯片上要时刻有大约 780KB 的数据正在路上。折算成 warp 级请求一个理想合并的 warp 加载 32 个 float正好 128 字节。780KB ÷ 128B ≈ 6100 个 warp 级请求同时在飞。如果这张卡有 100 多个 SM摊下来每个 SM 需要大约 60 个 warp 级请求同时在路上。这个数字非常关键因为它和每个 SM 最多驻留 64 个 warp几乎重合。这意味着只要每个驻留 warp 都能保持至少一条独立的加载指令在飞理论上就足够把带宽喂满。这句话是整篇内容里最有操作价值的一条。它直接导出了优化的优先级与其纠结 occupancy 从 75% 提到 100%不如先确保每个 warp 真的有一到两条独立加载在飞。很多occupancy 已经 90% 但带宽只有一半的案例问题就出在这里——warp 确实驻留了但因为循环里存在依赖每个 warp 在任一时刻只有零条或一条加载在飞而且中间还夹着大量地址计算指令占着派发槽。顺带说明一下为什么实测很难吃满。除了上面说的依赖问题还有几个固定开销地址计算、循环计数、分支判断这些指令都要占派发槽加载指令从发出到进入 LSU 队列还有一段延迟L2 和内存控制器本身也有排队。所以实测跑到 80% 到 85% 的带宽已经算是调得不错了盯着 100% 去追往往是白费力气。3.3 occupancy、ILP、TLP 三者的替换关系现在可以把三个概念的关系理清楚用一张表最容易看明白手段怎么用代价什么时候该用提高 TLP增加驻留 warp 数调 block 尺寸、降寄存器寄存器/共享内存压力上升可能引发 spilling每个线程的独立工作少、访存密集提高 ILP循环展开、多累加器、一次发多条独立加载寄存器用量上升代码体积变大计算密集、依赖链长、寄存器有余量缓存复用把数据搬进共享内存或寄存器复用需要额外的同步和拷贝指令数据被多次使用、访存模式规律三者的关系是可以互相替换的。一个线程如果能一次发出四条独立加载ILP4那么它需要的 warp 数就是 ILP1 时的四分之一。反过来如果一个 kernel 靠的是海量线程TLP 高那每个线程内部的依赖链长一点也无所谓。所以看到某 kernel 在 occupancy 30% 的情况下跑到峰值不要惊讶它多半是靠 ILP 或者说靠数据复用把延迟填掉了。实操心得调优的起点不该是把 occupancy 提上去而是先看long_scoreboard和带宽利用率的组合。带宽低 long_scoreboard高 并发度不够带宽高 long_scoreboard高 正常现象别动。带宽低 long_scoreboard低 问题不在访存去看not_selected和 pipe throttle。还有一点容易忽略提高 occupancy 和提高 ILP 是争抢同一份寄存器资源的。循环展开会让每个线程的寄存器需求上升可能把 occupancy 从 100% 挤到 50%。这个时候不能想当然地认为展开一定更好或者occupancy 一定更好得实测。经验上如果展开带来的 ILP 增益能补偿 occupancy 的下降就划算如果数据本身访存密集、天然容易并行那就没必要展开。4. 动手验证两段代码把调度行为跑出来4.1 对照实验怎么设计才干净看完上面这些如果不动手很容易停留在好像懂了的状态。我设计的对照实验思路很简单保持访存总量和数据规模完全一致只改变延迟被谁藏起来这一个变量。具体是三个版本版本 A每个线程处理 1 个 float靠海量线程TLP来藏延迟版本 B每个线程处理 4 个 float用float4向量化读取靠 ILP 来藏延迟线程数减到四分之一版本 C在版本 A 的基础上人为在加载和使用之间插入一长串依赖运算制造一个长依赖链观察 warp 被迫停在stalled_wait和long_scoreboard上的比例变化。三个版本读写的数据量、合并模式、总线程数B 除外都是精心对齐的这样任何性能差异都能归因到调度行为上而不是内存是不是没对齐这种低级原因。这一点很重要——对照实验里只能有一个变量否则数据没法解释。我在实际跑的时候还加了一个保护每个版本都跑预热 20 次、正式 100 次取中位数。GPU 上有太多噪声源时钟波动、其他进程、首次加载的冷启动跑一次就下结论的坑我踩过太多次。4.2 代码与编译参数版本 A纯 TLP 路线// tlp_version.cu __global__ void scale_tlp(const float* __restrict__ in, float* __restrict__ out, float k, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) return; out[i] in[i] * k; }版本 BILP 路线一次处理 4 个元素// ilp_version.cu __global__ void scale_ilp(const float4* __restrict__ in, float4* __restrict__ out, float k, int n4) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n4) return; float4 v in[i]; // 一条 16 字节的加载 v.x * k; v.y * k; v.z * k; v.w * k; out[i] v; }版本 C在加载和消费之间插入依赖链// dep_chain.cu __global__ void scale_dep(const float* __restrict__ in, float* __restrict__ out, float k, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) return; float v in[i]; // 结果要等很久 float acc k; #pragma unroll for (int t 0; t 32; t) { acc acc * 1.000001f 0.000001f; // 串行依赖链 } out[i] v * acc; // v 的使用被推后 }编译参数要统一否则数据不可比nvcc -O3 -archsm_80 --use_fast_math \ -Xptxas -v \ -o bench tlp_version.cu ilp_version.cu dep_chain.cu-Xptxas -v会打印每个 kernel 的寄存器用量、共享内存用量、spill 情况。这是必看的一步因为如果版本 C 因为循环展开而 spill那测出来的就不是调度效果而是本地内存访问的代价了。我实际跑的时候版本 C 的寄存器数是三个里最高的但还在可接受范围内没有 spill。关于-arch的选择一定要指定和你实际运行的卡匹配的架构代号sm_80、sm_86、sm_89、sm_90之类。默认编译出来的 PTX 可能要经过即时编译指令选择会和预期不同调度器的行为也会跟着变。这个坑我印象很深同一份代码在编译时指定和不指定架构实测性能差了 15% 以上就是因为即时编译出来的指令序列不一样。想查当前环境用什么代号看nvidia-smi的架构信息或者直接问编译器要一份目标列表。4.3 用 ncu 读调度器的三个指标跑完之后用 Nsight Compute 抓数据。不需要开全套抓三个我觉得最关键的指标就够定位问题ncu --section SchedulerStats \ --section WarpStateStats \ --metrics smsp__issue_active.avg.pct_of_peak_sustained_elapsed,\ sm__warps_active.avg.pct_of_peak_sustained_active,\ dram__throughput.avg.pct_of_peak_sustained_elapsed \ -c 3 ./bench三个指标的读法sm__warps_active实际驻留的 warp 数占峰值的百分比就是通常说的 occupancy。这是手里有多少牌不是打了多少牌。smsp__issue_active派发槽的活跃比例。这是真正衡量调度器有多忙的指标也是这条链路上最该盯的数字。注意它的分母是分区维度不要跟 SM 维度的指标混着比。dram__throughputDRAM 带宽利用率。这是访存延迟有没有被藏住的最终裁判。三个数字的组合能直接给出诊断occupancy 高 issue 低 带宽低说明手里有牌但打不出去问题在依赖链或者访存模式occupancy 高 issue 高 带宽低说明瓶颈在执行单元或者指令类型不匹配跟访存没关系occupancy 低 issue 低 带宽低说明并行度确实不足先解决 occupancy。实际测下来版本 A 和版本 B 的带宽利用率都能到 80% 以上但版本 B 的issue_active明显更低、occupancy 也只有 A 的四分之一——这就是典型的用 ILP 换 TLP线程少了但每个线程的活更密调度器的压力反而更小。版本 C 则是另一边occupancy 和 A 接近但issue_active掉下去stalled_long_scoreboard和stalled_wait一起上涨带宽也随之下降。这正好把延迟没被藏住这个抽象说法落成了具体数字。记录建议跑 benchmark 的时候把这三个指标连同寄存器用量、spill 字节数一起记下来。只记运行时间的话下次想复现同样的优化路径会非常困难。我自己的习惯是每次实验存一行 CSV跑多了之后回头看趋势比看单点有用得多。5. 常见问题与排查速查5.1 高 occupancy 但 issue 率上不去这是最常被问的一类。现象是 occupancy 已经 90% 以上但issue_active只有 20% 出头带宽也上不去。可能的成因按我的排查顺序排第一看stalled_not_selected的占比。如果它很低说明问题不是牌太多挑不过来而是牌本身没法打那就去看具体是哪个 stall 字段在涨。如果它很高说明就绪的 warp 确实很多但派发不过去——那问题在执行单元或者指令依赖的粒度上加 occupancy 只会更糟。第二看long_scoreboard和实际带宽的对照。如果带宽已经接近峰值long_scoreboard高是正常代价不用管如果带宽很低说明每个 warp 的独立加载数不够这时候要做的不是加线程而是给每个线程多塞几条独立加载或者改善访问模式让一个 warp 的加载真正合并成一次事务。第三看stalled_math_pipe_throttle。这个字段高说明某类计算单元被某一类指令打满了。常见于大量使用特殊函数exp、sin、log的场景这些指令走的是吞吐很低的流水线。这时候应该考虑用多项式近似替换或者降低调用密度。第四看stalled_lg_throttle和stalled_mio_throttle。这两个是队列背压说明指令发出的速率超过了后端消化能力。缓解办法通常是减少访存指令的条数用向量化把 4 条 4 字节加载合并成 1 条 16 字节加载而不是想办法发更多。5.2 寄存器压力与 spilling 的连锁反应寄存器这条线上有一串连锁反应值得单独讲。链条是循环展开 → 寄存器需求上升 → 寄存器档位向上取整 → 驻留 warp 数下降 → occupancy 下降 → 如果 ILP 的收益不够补偿性能就掉。这里有个很多人不知道的细节寄存器分配粒度是 8。你的 kernel 用了 33 个寄存器实际按 40 个算用了 65 个按 72 个算。所以在临界点上减少一两个寄存器可能让 occupancy 直接跳一档。我做过一次实验把内核里一个只用于调试的中间变量删掉寄存器从 41 降到 39档位从 48 掉到 40occupancy 从 62.5% 跳到 75%吞吐直接提了 8%。要控制寄存器几个可用的手段__launch_bounds__(maxThreads, minBlocks)告诉编译器你的目标驻留 block 数它会主动压寄存器用量。这是最直接的一招代价是可能引入 spill。-maxrregcountN是全局限制颗粒度太粗我一般不用除非整个项目就一两个内核。把不需要长期存活的值及时释放让编译器的活跃区间分析更准。这个很难精确控制但把大数组的循环拆开、让中间结果尽早归约出寄存器通常有帮助。用共享内存换寄存器。把需要跨线程共享的中间结果放共享内存寄存器压力会明显下降代价是多了同步和访问延迟。注意spill 的代价比多数人想的严重。spill 出去的数据放在本地内存本质上是走 L1/L2 的全局内存通路一次 spill 加载的延迟是几百个时钟周期和访存延迟同一个量级。所以为了提 occupancy 而牺牲寄存器、结果 spill 出一堆是很常见的反向优化一定要看-Xptxas -v打印的 spill 字节数是不是 0。5.3 排查速查表把上面这些整理成一张表遇到问题时从上往下对号入座现象大概率成因先试什么occupancy 高、issue 低、带宽低每 warp 独立加载数不足向量化访存或每线程处理多个元素occupancy 高、issue 高、带宽低瓶颈不在访存查 pipe throttle、特殊函数密度occupancy 低、issue 低寄存器或共享内存限制驻留看-Xptxas -v调 block 尺寸short_scoreboard高共享内存 bank 冲突或 MIO 排队给共享内存数组补 paddingno_instruction高指令缓存未命中关掉循环展开拆小函数barrier高block 内线程负载不均重新划分任务或减小 blockdrain高退出前有大量 store 未完成检查是否需要考虑用异步写加 occupancy 反而变慢方向反了或在临界点引发 spill回退改走 ILP 路线最后一行我想多说两句。加 occupancy 反而变慢是一个特别有价值的信号因为它几乎肯定意味着你踩到了资源临界点。这个时候不要硬调参数而是应该退回去看是不是展开了太多导致 spill是不是某个分区的资源先耗尽了是不是 block 的尺寸让四个分区分配不均。这种时候往往不是再加一点能解决的得换个思路。6. 几个踩过坑才知道的细节6.1 访存模式的重要性排在 occupancy 之前我早期的调优顺序一直是先看 occupancy不够就想办法提撞了很多次墙之后才把顺序调过来。现在的顺序是先确认访存的合并度和事务数量再考虑并发度。道理很朴素。一个 warp 如果发出的是 32 个分散在 32 个不同 cache line 的加载那它会占用 32 次内存事务而理想情况下只需要 4 次128 字节一行32 个 float 正好一行。这意味着同样的数据量内存系统要处理的请求数多了 8 倍。这种情况下你把 occupancy 提到天上去也没用因为内存控制器的排队长度本身就爆了。检查方法很简单看 profiler 里每个访存请求对应的扇区数sectors per request。理想值是和合并度相关的固定值如果实测远高于理想值就先解决访存模式别碰 occupancy。改法包括调整线程到数据的映射让相邻线程访问相邻地址、用结构体数组转数组结构体AoS 到 SoA、给二维数组补齐 stride 避免跨行、用向量化类型把多个连续访问合并成一条指令。6.2 用预取和循环重排把长延迟指令往后埋比起调 occupancy我越来越常用的是把加载提前。思路是在循环体内先把下一轮要用的数据加载进去预取然后再处理当前轮的数据。这样当处理逻辑在跑的时候下一轮的加载已经在飞了延迟自然被藏在计算里。最土的实现就是手动展开两轮// 双缓冲预取 float curr in[idx]; for (int t 0; t iters; t) { float next; if (t 1 iters) next in[idx step * (t 1)]; // 提前发出 out[idx step * t] curr * k; // 用上一轮的数据 curr next; }这样每一轮里next的加载和curr的使用互不依赖加载可以一直在飞。实测下来这种改法在访存密集的循环内核里效果非常明显比死磕 occupancy 划算得多。要注意的是边界处理要小心别让预取越界。还有一个更省事的做法直接把循环展开 4 到 8 次让编译器自己去做指令重排。编译器的调度能力比想象中强只要没有寄存器压力它会把独立的加载尽量往前挪。前提是循环体里没有跨迭代的依赖有依赖的话它也无能为力。6.3 数据记录方式和实验纪律最后说点方法论的东西因为技术细节容易查纪律难养。我现在每次调优都存一份 CSV字段包括内核名、编译参数、架构代号、寄存器数、spill 字节数、共享内存用量、block 尺寸、grid 尺寸、occupancy、issue_active、带宽利用率、各 stall 字段的前三名、运行时间中位数。听起来很啰嗦但好处是做过的实验可以复现不需要靠记忆猜。有时候回头看两周前的数据会发现当时下的结论是错的。另一个纪律是一次只改一个东西。我知道这听起来是废话但在性能调优里真的很难做到因为改 block 尺寸经常顺手就把展开次数也调了。可一旦多改了几个变量性能提升了也不知道是谁的功劳下次遇到类似问题还是靠碰运气。数据记录加上单变量这两条看着笨但它是唯一能让经验真正积累起来的办法。至于这个方向还能往哪走我现在主要在关注两件事一是用异步拷贝之类的机制把数据搬运和数据计算真正重叠起来绕开寄存器这个瓶颈二是在 profiler 之外自己写一些微基准去测调度器的具体行为因为公开文档对调度策略的描述实在有限很多结论只能靠实测反推。这部分还没整理成型等有稳定的结论再单独写。
CUDA warp调度与访存延迟隐藏：从occupancy到issue率
手里这台卡跑一个纯访存的内核理论带宽标称一点五 TB/s 往上实测在 55% 到 60% 之间晃profiler 里issue_active只到 20% 出头warp state 那一栏long_scoreboard占了大半——这种局面我遇到过不止一次。多数人第一反应是内存是不是没对齐是不是没向量化但把访存指令改到完美合并之后仍然卡在同一个位置问题就落在 SM 的 warp 调度器身上了线程块发出去了、warp 也驻留了可是调度器每个周期挑不出足够的就绪 warp 来填空访存延迟就那么明晃晃地挂在流水线上没人接。这篇东西聊的就是这条链路——CUDA 的 SM 内部warp 调度器到底按什么规则挑 warp访存延迟又是靠什么被藏起来的以及怎么用可复现的实验把这两个抽象概念变成 profiler 上的数字。适合已经能写能跑 kernel、但性能调优基本靠猜的同学也适合把 occupancy 当成唯一调优旋钮、结果发现调到 100% 还是不提速的人。我不打算复述官方文档而是按我实际的排查顺序走一遍先认角色再算账然后动手跑对照实验最后把踩过的坑摊开。全程以 Ampere 这一代的结构为主前后代有差异的地方我会单独点出来。1. 先把角色认清楚SM 里到底谁在干活1.1 一个 SM 的硬件骨架长什么样要谈调度得先知道被调度的对象住在哪。一个 SMStreaming Multiprocessor可以粗略理解成四个小核加一堆共享资源四个处理块官方叫 sub-partition 或 processing block每个处理块里有自己的 warp 调度器、自己的指令派发单元、自己的通用计算单元组以及一组共享的寄存器文件、L1/共享内存、加载存储单元和特殊功能单元。关键在于一个数字每个处理块每个时钟周期最多派发一条 warp 指令。四个处理块加一起就是每个 SM 每周期最多四条 warp 指令。这个4是个硬约束很多性能模型算不对就是因为忽略了它——你把 occupancy 拉满是没用的派发宽度就摆在那里超过它之后的任何并行度都只能排队。驻留容量方面以 Ampere 的计算卡为例一个 SM 最多驻留 64 个 warp2048 线程、65536 个 32 位寄存器、约 256KB 的寄存器堆总量、最多 32 个 block、16 个 barrier 资源。消费级那一代比如 GA10x 那批最大驻留是 48 个 warp / 1536 线程寄存器总量一致。这些数字必须记牢因为它们决定了 occupancy 的天花板而 occupancy 又直接决定调度器手里有多少牌可以打。寄存器是这里最容易翻车的资源。它按线程粒度分配粒度是 8 个寄存器一档编译产物里regs/thread会向上取整到 8 的倍数。假设某个 kernel 每个线程用了 100 个寄存器向上取整到 104那么 2048 个线程需要 212992 个寄存器超过 65536 的限制硬件就会把驻留线程数压到 1536再算下来 warp 数只有 48occupancy 直接掉到 75%。很多人抱怨我 block 设成 256 线程怎么只能跑 6 个 block答案往往就在这里而不是在什么玄学地方。1.2 为什么是 32warp 这个粒度的由来与代价一个 warp 固定 32 个线程它们共享同一个程序计数器、同一份指令以 SIMT 的方式同步推进。为什么偏偏是 32 而不是 16 或 64这是延迟与吞吐的折中。粒度太细调度器每周期要维护的 warp 条目太多选择逻辑和上下文状态的成本压不住粒度太粗一旦某个 32 线程组里出现分支分歧或者访存不合并浪费的比例就成倍上升。32 恰好是当时存储子系统一个 cache line 的 128 字节除以单线程 4 字节的宽度访存合并的账算起来最整齐。代价也很直接warp 是调度的最小单位也是浪费的最小单位。if 分支里两个路径各走一遍整个 warp 的执行时间就是两条路径之和32 个线程里如果只有 1 个线程真正需要走某段计算其余 31 个线程照样占着执行单元的槽位只是结果被掩码掉。所以写 kernel 的时候让同一个 warp 里的 32 个线程尽量走同一条控制流、访问连续的地址这不是优化技巧而是省掉本不该付的成本。另一个容易被忽略的点warp 内部的线程不是完全独立的。只要指令序列里有任何一条是全部 32 个线程都执行的硬件就会把它们压成一条 warp 指令去跑同一个执行单元阵列。这意味着你在写代码时看到的一个线程一次加法落到硬件上可能是一个 warp 一次 32 路加法占用的是同一个指令槽。指令计数的口径差异就是很多人的手算模型和实测数据对不上的根源。1.3 一个 block 从提交到跑起来要过几道关搞清楚从grid, block到真正开始跑指令之间发生了什么排障的时候会省很多时间。流程大致是Grid 分发grid 里的 block 按某种顺序分给各个 SM顺序没有强保证不要依赖 blockIdx 的物理分配顺序。资源预留SM 检查目标 block 需要的寄存器、共享内存、barrier 槽位、线程槽位是不是都够。任何一项不够block 就在队列里等。线程切分block 内的线程按threadIdx线性编号每 32 个切成一个 warp。这里有个细节——切分是按 block 内的线性 tid 来的不是按多维坐标所以dim3(32, 8)的 block 和dim3(256)的 blockwarp 的构成方式是一样的。分配 warp slot每个 warp 被塞进某个处理块。分配策略没有公开但从微基准看基本是轮转的目的是让四个处理块的负载均衡。进入就绪队列warp 的 PC 指向 kernel 入口此后就进入调度器的候选池。踩坑提醒如果共享内存用得特别狠比如一个 block 申请 100KB 以上会出现block 能放下但只能放一个的情况SM 里其他处理块空转调度器手里只有一个 block 的 warp 可以挑。这时候调 occupancy 的旋钮是没用的得从 block 尺寸和共享内存用量下手。还有一点block 的调度顺序会影响结果。同一份 kernel如果 block 之间的负载天然不均比如按数据分布决定循环次数先被调度的 block 跑得快、后被调度的慢最后总时间被慢的那个拖住。这不是调度器的问题是负载划分的问题但表现出来像是调度不公。实践中我倾向于让 block 的工作量尽量均匀或者用持久化 kernelpersistent kernel的方式手动接管任务分配。2. warp 调度器每周期到底在挑什么2.1 就绪判定从 PC 到 issue slot 的链路每一拍四个处理块各自的调度器都要做一次选择从住在自己这里的一堆 warp 里挑一个把它的下一条指令送进派发单元。这个选择过程公开资料只给了轮廓从微基准能反推出大致是这么几层判断第一层是能不能发。一个 warp 只有在满足以下条件时才算就绪PC 指向的指令已经取到所有源操作数都已就位没有被标注为等待中不处于 barrier 等待或分支解析中没有被 MIO 队列或长延迟队列的背压卡住没有在执行__syncthreads之类的同步语义。任何一条不满足这个 warp 这一拍就直接出局。第二层是谁先发。剩下的就绪 warp 里要选一个这就是策略部分。Fermi/Pascal 时代比较明确是 GTOGreedy Then Oldest——优先挑同一 warp 的连续指令实在不行挑最老的。Volta 之后官方不再详细说明但从行为上看它也不是简单地轮转而是带有某种年龄加权并且会参考目标执行单元的当前压力。判据的核心其实是让各条流水线都别空着而不是对每个 warp 公平。第三层是发去哪。选中的指令还要看目标执行单元这一拍有没有空位。如果目标 pipe 正忙指令会被按住产生的就是math_pipe_throttle或mio_throttle这类背压统计。这就解释了一个常见现象明明有很多 warp 就绪issue_active却上不去——不是挑不出人而是执行单元接不住。操作数的已就位是怎么判断的靠 scoreboard。每个 warp 有一组记分牌位指向那些结果还没写回来的寄存器。一条指令要发射前硬件会检查它的源寄存器有没有被标记为等待。等待来自固定延迟执行单元比如共享内存、某些特殊运算的叫 short scoreboard等待来自全局内存加载的叫 long scoreboard。这两者在 profiler 里是两个独立字段含义完全不同——short 一般意味着 MIO 排队long 才是真正的 DRAM 延迟。这里有个反直觉的地方long scoreboard 高不一定是坏事。如果访存请求确实在飞、带宽也确实吃满了那么 warp 停在 long scoreboard 上等结果是正常的、不可避免的。只有当带宽利用率也低的时候long scoreboard 才是问题信号说明并发度不够延迟没被藏住。2.2 四个子分区与派发宽度别把 SM 当成一个整体新手最容易犯的错是把 SM 看成一个大池子算 occupancy 的时候用 SM 总量除以线程数。实际上资源是分到四个处理块上的寄存器、warp slot、执行单元都是分区资源。这带来两个很实际的后果。第一occupancy 的瓶颈可能只在单个分区上。比如 block 有 256 线程切成 8 个 warp。如果 SM 能放 6 个这样的 block一共 48 个 warp理论上每个分区 12 个。但如果硬件的分配策略让某个分区恰好分到了 13 个、另一个只有 11 个那个 13 个的分区就成了新的瓶颈。这在实测里表现为两个完全一样的 kernel仅仅因为 block 尺寸从 256 改成 128issue 率就有可见差别。原因就是资源分配的粒度变了四个分区的均衡程度变了。第二派发宽度是硬上限。前面说过每分区每周期一条指令四个分区加一起四条。所以一个 SM 想跑满理论算力前提是四条流水线同时有活干。假设某个 kernel 的指令流里有 50% 是访存指令、50% 是计算指令那么当访存指令卡在 LSU 排队时计算指令就得顶上反过来也一样。每条流水线的吞吐要匹配否则最快的那个也会被最慢的拖住。关于双发射这个说法我得泼点冷水。一些架构在特定指令组合下确实存在同一周期派出两条指令的情况比如一条浮点加一条整数但从性能建模的角度按每分区每周期一条来估算误差通常在可以接受的范围内。真要确认跑个微基准看issue_active的上限就知道——如果它稳定在 100%说明模型对上了如果偶尔超过说明有特例。2.3 stall 名字背后的硬件事件对照profiler 的 WarpStateStats 那一栏一堆stalled_xxx字段刚开始看特别懵。我把它按我该去改什么重新分了个类比按字母顺序看有用得多。字段硬件含义优先怀疑的方向stalled_long_scoreboard等全局/本地内存加载结果并发度不足、访存不合并、L2 未命中stalled_short_scoreboard等共享内存、特殊函数等固定延迟结果共享内存 bank 冲突、MIO 排队stalled_wait等固定延迟指令如某些 ALU的结果依赖链过长ILP 不足stalled_not_selectedwarp 已就绪但没被调度器选中说明不缺并行度缺的是派发宽度或执行单元吞吐stalled_mio_throttleMIO 指令队列满共享内存/特殊函数指令过密stalled_lg_throttle全局内存指令队列满访存指令发出速率超过 LSU 消化能力stalled_math_pipe_throttle目标计算流水线繁忙指令类型过于单一某条 pipe 打满stalled_barrier等__syncthreads或 block 间同步block 内负载不均stalled_no_instruction取指失败指令缓存未命中内核代码体积过大、分支过于分散stalled_drain内核退出前等待内存操作完成store 太多、退出前没做同步这张表的价值在于它能直接切分问题not_selected高说明并行度够了、算力不够long_scoreboard高说明并行度不够、内存延迟没藏住。这两个方向的处理方式完全相反——前者要减少并行度省下的资源换更高的单线程效率后者要增加并行度或缩短依赖链。我见过有人看到not_selected很高就去加 occupancy结果越加越慢就是因为方向反了。另外提一句stalled_no_instruction。这个字段一般很小但如果它占比异常通常意味着内核的指令体积太大、超出了指令缓存的容量。处理办法是把循环展开关掉#pragma unroll 1、把大函数拆小或者在编译时限制展开的深度。这是那种不看字段永远想不到的问题。3. 把访存延迟算清楚为什么需要那么多 warp3.1 延迟隐藏的本质是填洞先说清楚延迟隐藏这个说法到底在说什么。一条全局加载指令发出后到数据真正回到寄存器中间要经过 L1 查询、L2 查询、内存控制器排队、DRAM 阵列访问、数据回传这一整段是几百个时钟周期。在这几百拍里发出这条指令的 warp 什么也做不了它的记分牌位被置住只能干等。如果 SM 里只有这么一个 warp那么这几百拍就是纯粹的空转执行单元全闲。延迟隐藏的意思就是在这几百拍里找别的活给执行单元干。别的活从哪来要么是同一个 SM 里的其他 warpTLP线程级并行要么是同一个线程里另外那些不依赖这个加载结果的指令ILP指令级并行。所有的调优手段归根到底都是在扩大这两个来源中的一个或两个。这个视角很重要因为它告诉你occupancy 从来不是目标它只是获取 TLP 的手段之一。如果一个 kernel 靠 ILP 就能把每个 warp 的访存间隙填满那么它根本不需要高 occupancy反过来如果一个 kernel 的每个线程都在等同一个加载结果那再高的 occupancy 也只是把更多 warp 排进等待队列而已。3.2 用 Littles Law 反推需要多少在飞请求想吃得满带宽光有并发线程不够得有足够的在飞的访存请求。这个可以用排队论里的 Littles Law 直接算在飞请求数 带宽 × 延迟拿一个具体例子算。假设某张计算卡的理论带宽是 1.6 TB/sDRAM 往返延迟按 500 纳秒算这个数量级在 GDDR/HBM 上都差不多具体随架构和频率浮动在飞字节数 1.6e12 B/s × 500e-9 s 800,000 B ≈ 780 KB也就是说为了让这条内存通路不空转整个芯片上要时刻有大约 780KB 的数据正在路上。折算成 warp 级请求一个理想合并的 warp 加载 32 个 float正好 128 字节。780KB ÷ 128B ≈ 6100 个 warp 级请求同时在飞。如果这张卡有 100 多个 SM摊下来每个 SM 需要大约 60 个 warp 级请求同时在路上。这个数字非常关键因为它和每个 SM 最多驻留 64 个 warp几乎重合。这意味着只要每个驻留 warp 都能保持至少一条独立的加载指令在飞理论上就足够把带宽喂满。这句话是整篇内容里最有操作价值的一条。它直接导出了优化的优先级与其纠结 occupancy 从 75% 提到 100%不如先确保每个 warp 真的有一到两条独立加载在飞。很多occupancy 已经 90% 但带宽只有一半的案例问题就出在这里——warp 确实驻留了但因为循环里存在依赖每个 warp 在任一时刻只有零条或一条加载在飞而且中间还夹着大量地址计算指令占着派发槽。顺带说明一下为什么实测很难吃满。除了上面说的依赖问题还有几个固定开销地址计算、循环计数、分支判断这些指令都要占派发槽加载指令从发出到进入 LSU 队列还有一段延迟L2 和内存控制器本身也有排队。所以实测跑到 80% 到 85% 的带宽已经算是调得不错了盯着 100% 去追往往是白费力气。3.3 occupancy、ILP、TLP 三者的替换关系现在可以把三个概念的关系理清楚用一张表最容易看明白手段怎么用代价什么时候该用提高 TLP增加驻留 warp 数调 block 尺寸、降寄存器寄存器/共享内存压力上升可能引发 spilling每个线程的独立工作少、访存密集提高 ILP循环展开、多累加器、一次发多条独立加载寄存器用量上升代码体积变大计算密集、依赖链长、寄存器有余量缓存复用把数据搬进共享内存或寄存器复用需要额外的同步和拷贝指令数据被多次使用、访存模式规律三者的关系是可以互相替换的。一个线程如果能一次发出四条独立加载ILP4那么它需要的 warp 数就是 ILP1 时的四分之一。反过来如果一个 kernel 靠的是海量线程TLP 高那每个线程内部的依赖链长一点也无所谓。所以看到某 kernel 在 occupancy 30% 的情况下跑到峰值不要惊讶它多半是靠 ILP 或者说靠数据复用把延迟填掉了。实操心得调优的起点不该是把 occupancy 提上去而是先看long_scoreboard和带宽利用率的组合。带宽低 long_scoreboard高 并发度不够带宽高 long_scoreboard高 正常现象别动。带宽低 long_scoreboard低 问题不在访存去看not_selected和 pipe throttle。还有一点容易忽略提高 occupancy 和提高 ILP 是争抢同一份寄存器资源的。循环展开会让每个线程的寄存器需求上升可能把 occupancy 从 100% 挤到 50%。这个时候不能想当然地认为展开一定更好或者occupancy 一定更好得实测。经验上如果展开带来的 ILP 增益能补偿 occupancy 的下降就划算如果数据本身访存密集、天然容易并行那就没必要展开。4. 动手验证两段代码把调度行为跑出来4.1 对照实验怎么设计才干净看完上面这些如果不动手很容易停留在好像懂了的状态。我设计的对照实验思路很简单保持访存总量和数据规模完全一致只改变延迟被谁藏起来这一个变量。具体是三个版本版本 A每个线程处理 1 个 float靠海量线程TLP来藏延迟版本 B每个线程处理 4 个 float用float4向量化读取靠 ILP 来藏延迟线程数减到四分之一版本 C在版本 A 的基础上人为在加载和使用之间插入一长串依赖运算制造一个长依赖链观察 warp 被迫停在stalled_wait和long_scoreboard上的比例变化。三个版本读写的数据量、合并模式、总线程数B 除外都是精心对齐的这样任何性能差异都能归因到调度行为上而不是内存是不是没对齐这种低级原因。这一点很重要——对照实验里只能有一个变量否则数据没法解释。我在实际跑的时候还加了一个保护每个版本都跑预热 20 次、正式 100 次取中位数。GPU 上有太多噪声源时钟波动、其他进程、首次加载的冷启动跑一次就下结论的坑我踩过太多次。4.2 代码与编译参数版本 A纯 TLP 路线// tlp_version.cu __global__ void scale_tlp(const float* __restrict__ in, float* __restrict__ out, float k, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) return; out[i] in[i] * k; }版本 BILP 路线一次处理 4 个元素// ilp_version.cu __global__ void scale_ilp(const float4* __restrict__ in, float4* __restrict__ out, float k, int n4) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n4) return; float4 v in[i]; // 一条 16 字节的加载 v.x * k; v.y * k; v.z * k; v.w * k; out[i] v; }版本 C在加载和消费之间插入依赖链// dep_chain.cu __global__ void scale_dep(const float* __restrict__ in, float* __restrict__ out, float k, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) return; float v in[i]; // 结果要等很久 float acc k; #pragma unroll for (int t 0; t 32; t) { acc acc * 1.000001f 0.000001f; // 串行依赖链 } out[i] v * acc; // v 的使用被推后 }编译参数要统一否则数据不可比nvcc -O3 -archsm_80 --use_fast_math \ -Xptxas -v \ -o bench tlp_version.cu ilp_version.cu dep_chain.cu-Xptxas -v会打印每个 kernel 的寄存器用量、共享内存用量、spill 情况。这是必看的一步因为如果版本 C 因为循环展开而 spill那测出来的就不是调度效果而是本地内存访问的代价了。我实际跑的时候版本 C 的寄存器数是三个里最高的但还在可接受范围内没有 spill。关于-arch的选择一定要指定和你实际运行的卡匹配的架构代号sm_80、sm_86、sm_89、sm_90之类。默认编译出来的 PTX 可能要经过即时编译指令选择会和预期不同调度器的行为也会跟着变。这个坑我印象很深同一份代码在编译时指定和不指定架构实测性能差了 15% 以上就是因为即时编译出来的指令序列不一样。想查当前环境用什么代号看nvidia-smi的架构信息或者直接问编译器要一份目标列表。4.3 用 ncu 读调度器的三个指标跑完之后用 Nsight Compute 抓数据。不需要开全套抓三个我觉得最关键的指标就够定位问题ncu --section SchedulerStats \ --section WarpStateStats \ --metrics smsp__issue_active.avg.pct_of_peak_sustained_elapsed,\ sm__warps_active.avg.pct_of_peak_sustained_active,\ dram__throughput.avg.pct_of_peak_sustained_elapsed \ -c 3 ./bench三个指标的读法sm__warps_active实际驻留的 warp 数占峰值的百分比就是通常说的 occupancy。这是手里有多少牌不是打了多少牌。smsp__issue_active派发槽的活跃比例。这是真正衡量调度器有多忙的指标也是这条链路上最该盯的数字。注意它的分母是分区维度不要跟 SM 维度的指标混着比。dram__throughputDRAM 带宽利用率。这是访存延迟有没有被藏住的最终裁判。三个数字的组合能直接给出诊断occupancy 高 issue 低 带宽低说明手里有牌但打不出去问题在依赖链或者访存模式occupancy 高 issue 高 带宽低说明瓶颈在执行单元或者指令类型不匹配跟访存没关系occupancy 低 issue 低 带宽低说明并行度确实不足先解决 occupancy。实际测下来版本 A 和版本 B 的带宽利用率都能到 80% 以上但版本 B 的issue_active明显更低、occupancy 也只有 A 的四分之一——这就是典型的用 ILP 换 TLP线程少了但每个线程的活更密调度器的压力反而更小。版本 C 则是另一边occupancy 和 A 接近但issue_active掉下去stalled_long_scoreboard和stalled_wait一起上涨带宽也随之下降。这正好把延迟没被藏住这个抽象说法落成了具体数字。记录建议跑 benchmark 的时候把这三个指标连同寄存器用量、spill 字节数一起记下来。只记运行时间的话下次想复现同样的优化路径会非常困难。我自己的习惯是每次实验存一行 CSV跑多了之后回头看趋势比看单点有用得多。5. 常见问题与排查速查5.1 高 occupancy 但 issue 率上不去这是最常被问的一类。现象是 occupancy 已经 90% 以上但issue_active只有 20% 出头带宽也上不去。可能的成因按我的排查顺序排第一看stalled_not_selected的占比。如果它很低说明问题不是牌太多挑不过来而是牌本身没法打那就去看具体是哪个 stall 字段在涨。如果它很高说明就绪的 warp 确实很多但派发不过去——那问题在执行单元或者指令依赖的粒度上加 occupancy 只会更糟。第二看long_scoreboard和实际带宽的对照。如果带宽已经接近峰值long_scoreboard高是正常代价不用管如果带宽很低说明每个 warp 的独立加载数不够这时候要做的不是加线程而是给每个线程多塞几条独立加载或者改善访问模式让一个 warp 的加载真正合并成一次事务。第三看stalled_math_pipe_throttle。这个字段高说明某类计算单元被某一类指令打满了。常见于大量使用特殊函数exp、sin、log的场景这些指令走的是吞吐很低的流水线。这时候应该考虑用多项式近似替换或者降低调用密度。第四看stalled_lg_throttle和stalled_mio_throttle。这两个是队列背压说明指令发出的速率超过了后端消化能力。缓解办法通常是减少访存指令的条数用向量化把 4 条 4 字节加载合并成 1 条 16 字节加载而不是想办法发更多。5.2 寄存器压力与 spilling 的连锁反应寄存器这条线上有一串连锁反应值得单独讲。链条是循环展开 → 寄存器需求上升 → 寄存器档位向上取整 → 驻留 warp 数下降 → occupancy 下降 → 如果 ILP 的收益不够补偿性能就掉。这里有个很多人不知道的细节寄存器分配粒度是 8。你的 kernel 用了 33 个寄存器实际按 40 个算用了 65 个按 72 个算。所以在临界点上减少一两个寄存器可能让 occupancy 直接跳一档。我做过一次实验把内核里一个只用于调试的中间变量删掉寄存器从 41 降到 39档位从 48 掉到 40occupancy 从 62.5% 跳到 75%吞吐直接提了 8%。要控制寄存器几个可用的手段__launch_bounds__(maxThreads, minBlocks)告诉编译器你的目标驻留 block 数它会主动压寄存器用量。这是最直接的一招代价是可能引入 spill。-maxrregcountN是全局限制颗粒度太粗我一般不用除非整个项目就一两个内核。把不需要长期存活的值及时释放让编译器的活跃区间分析更准。这个很难精确控制但把大数组的循环拆开、让中间结果尽早归约出寄存器通常有帮助。用共享内存换寄存器。把需要跨线程共享的中间结果放共享内存寄存器压力会明显下降代价是多了同步和访问延迟。注意spill 的代价比多数人想的严重。spill 出去的数据放在本地内存本质上是走 L1/L2 的全局内存通路一次 spill 加载的延迟是几百个时钟周期和访存延迟同一个量级。所以为了提 occupancy 而牺牲寄存器、结果 spill 出一堆是很常见的反向优化一定要看-Xptxas -v打印的 spill 字节数是不是 0。5.3 排查速查表把上面这些整理成一张表遇到问题时从上往下对号入座现象大概率成因先试什么occupancy 高、issue 低、带宽低每 warp 独立加载数不足向量化访存或每线程处理多个元素occupancy 高、issue 高、带宽低瓶颈不在访存查 pipe throttle、特殊函数密度occupancy 低、issue 低寄存器或共享内存限制驻留看-Xptxas -v调 block 尺寸short_scoreboard高共享内存 bank 冲突或 MIO 排队给共享内存数组补 paddingno_instruction高指令缓存未命中关掉循环展开拆小函数barrier高block 内线程负载不均重新划分任务或减小 blockdrain高退出前有大量 store 未完成检查是否需要考虑用异步写加 occupancy 反而变慢方向反了或在临界点引发 spill回退改走 ILP 路线最后一行我想多说两句。加 occupancy 反而变慢是一个特别有价值的信号因为它几乎肯定意味着你踩到了资源临界点。这个时候不要硬调参数而是应该退回去看是不是展开了太多导致 spill是不是某个分区的资源先耗尽了是不是 block 的尺寸让四个分区分配不均。这种时候往往不是再加一点能解决的得换个思路。6. 几个踩过坑才知道的细节6.1 访存模式的重要性排在 occupancy 之前我早期的调优顺序一直是先看 occupancy不够就想办法提撞了很多次墙之后才把顺序调过来。现在的顺序是先确认访存的合并度和事务数量再考虑并发度。道理很朴素。一个 warp 如果发出的是 32 个分散在 32 个不同 cache line 的加载那它会占用 32 次内存事务而理想情况下只需要 4 次128 字节一行32 个 float 正好一行。这意味着同样的数据量内存系统要处理的请求数多了 8 倍。这种情况下你把 occupancy 提到天上去也没用因为内存控制器的排队长度本身就爆了。检查方法很简单看 profiler 里每个访存请求对应的扇区数sectors per request。理想值是和合并度相关的固定值如果实测远高于理想值就先解决访存模式别碰 occupancy。改法包括调整线程到数据的映射让相邻线程访问相邻地址、用结构体数组转数组结构体AoS 到 SoA、给二维数组补齐 stride 避免跨行、用向量化类型把多个连续访问合并成一条指令。6.2 用预取和循环重排把长延迟指令往后埋比起调 occupancy我越来越常用的是把加载提前。思路是在循环体内先把下一轮要用的数据加载进去预取然后再处理当前轮的数据。这样当处理逻辑在跑的时候下一轮的加载已经在飞了延迟自然被藏在计算里。最土的实现就是手动展开两轮// 双缓冲预取 float curr in[idx]; for (int t 0; t iters; t) { float next; if (t 1 iters) next in[idx step * (t 1)]; // 提前发出 out[idx step * t] curr * k; // 用上一轮的数据 curr next; }这样每一轮里next的加载和curr的使用互不依赖加载可以一直在飞。实测下来这种改法在访存密集的循环内核里效果非常明显比死磕 occupancy 划算得多。要注意的是边界处理要小心别让预取越界。还有一个更省事的做法直接把循环展开 4 到 8 次让编译器自己去做指令重排。编译器的调度能力比想象中强只要没有寄存器压力它会把独立的加载尽量往前挪。前提是循环体里没有跨迭代的依赖有依赖的话它也无能为力。6.3 数据记录方式和实验纪律最后说点方法论的东西因为技术细节容易查纪律难养。我现在每次调优都存一份 CSV字段包括内核名、编译参数、架构代号、寄存器数、spill 字节数、共享内存用量、block 尺寸、grid 尺寸、occupancy、issue_active、带宽利用率、各 stall 字段的前三名、运行时间中位数。听起来很啰嗦但好处是做过的实验可以复现不需要靠记忆猜。有时候回头看两周前的数据会发现当时下的结论是错的。另一个纪律是一次只改一个东西。我知道这听起来是废话但在性能调优里真的很难做到因为改 block 尺寸经常顺手就把展开次数也调了。可一旦多改了几个变量性能提升了也不知道是谁的功劳下次遇到类似问题还是靠碰运气。数据记录加上单变量这两条看着笨但它是唯一能让经验真正积累起来的办法。至于这个方向还能往哪走我现在主要在关注两件事一是用异步拷贝之类的机制把数据搬运和数据计算真正重叠起来绕开寄存器这个瓶颈二是在 profiler 之外自己写一些微基准去测调度器的具体行为因为公开文档对调度策略的描述实在有限很多结论只能靠实测反推。这部分还没整理成型等有稳定的结论再单独写。
