企业官网建设流程全解析

1. 为什么我劝你先啃透 warp 调度与延迟隐藏

先说一个我这两年里反复遇到的场景：一个看着很朴素的带宽型 kernel，比如向量加、归约、或者转置，用 ncu 一看 occupancy 已经 100% 了，Achieved Occupancy 跑到 64 warps/SM 满值，但 DRAM Throughput 只有峰值的 40% 出头。这时候很多人第一反应是"是不是 block 开少了""是不是 grid 不够大"，于是把 grid 从 108 拉到 10800，结果数字纹丝不动。真正的原因藏在smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active这个指标里——它在告诉你，绝大多数 warp 其实都卡在等全局内存返回，调度器每一轮都在面对一堆"没资格发射"的 warp。

CUDA里所有性能话题最终都会落到两个地方：SM里的warp 调度器在每个周期挑谁发射，以及当被挑中的 warp 需要等 500 个周期才能拿到数据时，怎么让别的 warp 把这段空档填满，也就是访存延迟隐藏。这两件事是同一枚硬币的两面：调度机制决定"谁能上"，延迟隐藏决定"上去之后能不能把流水线喂饱"。搞不清楚这两层，你做算子优化基本就是瞎调参数。

这篇文章我打算按自己的理解路径来写：先把 SM 的物理结构拆开看看有哪些部件、各自的预算上限是多少；再钻进调度器的状态机，看它每个周期到底在做什么决策；然后用 Little's Law 把"我到底需要多少 warp 才能藏住延迟"算成一道小学算术题；最后拿一个真实的带宽型 kernel 做实录，一边改代码一边看指标，把 42% 拉到 88%。中间会穿插 Nsight Compute 的指标速查表和我这些年踩过的坑。适合已经能写能跑 CUDA kernel、但性能始终卡在天花板下的朋友，也适合刚接触 GPU 性能分析、想建立一套系统判断框架的人。不需要你懂电路，但需要你会写for循环。

1.1 一个反直觉的现象：occupancy 拉满，带宽只有四成

我拿归约 kernel 举过太多次例子了。假设你写了一个最朴素的版本：每个线程用网格跨步循环累加，最后atomicAdd到全局。这个 kernel 的 occupancy 通常能轻松打满，因为每线程只用了十几个寄存器，共享内存为零，块大小设成 256 的话，一个 SM 能塞下 8 个块、64 个 warp，正好顶到硬件上限。

但它的性能往往很难看。原因不在"线程不够多"，而在"每个线程只携带了极少的内存并行度"。具体说：线程执行s += x[i]这条语句时，会先发起一次 32-bit 的全局加载，然后立刻需要这个值来做加法。加载指令在 LSU 里排队、穿过 L1、L2、到 HBM，再原路返回来，这一趟在 A100 上大约要 400 到 600 个周期。返回之前，这条线程所在的 warp 处于long scoreboard阻塞状态，调度器不会选它。如果所有 64 个 warp 都处于这种"发一条、等回来"的节奏，那么整个 SM 在每个瞬间真正在飞的访存请求数其实非常有限。

粗算一下：一个 warp 发一条LDG.E.32指令，请求 128 字节（32 线程 × 4 字节）。64 个 warp 全部各有一条在飞，也就 8KB 的数据在路上。而 A100 满带宽时，每 SM 每周期需要的字节吞吐是十几个字节，500 周期的延迟意味着你需要大约 6 到 7KB 的数据持续在飞才能把管道填满——看起来 8KB 够了？问题是现实中没有这么理想，地址计算、循环变量更新、请求返回后的消费阶段都会占掉时间片，实际同时悬挂的请求数往往只有理论值的一半甚至三分之一。这就是为什么 occupancy 100% 却只能跑出四成带宽。

1.2 这篇文章会带你走完的完整链路

整条链路的逻辑顺序是这样的：SM 被切成四个子分区，每个子分区有自己的 warp 调度器和寄存器堆；调度器每周期从"就绪"的 warp 里挑一个发射，被选中的前提是该 warp 的下一条指令的操作数已经准备好；如果没准备好，warp 就带着一个 stall 原因挂起，等资源就绪后重新进入候选池；而所谓"隐藏延迟"，本质就是让调度器手上永远有足够多的备用 warp 可选，这样即使一批 warp 在等内存，另一批仍然能把发射端口占满。

把这套逻辑想通之后你会发现，几乎所有 CUDA 优化手段都能被归到三类里：增加并发 warp 数（提高 TLP，靠 occupancy）、增加每个 warp 的在飞请求数（提高 MLP，靠展开和向量化）、降低单次访问的延迟或提高带宽利用率（靠合并访问、共享内存、缓存策略）。剩下的都是这三类的变体。后面每一章我都会落到具体的公式、代码或者指标上，不停留在概念。

2. 拆开 SM：调度器到底在什么样的物理环境里工作

很多人学 CUDA 时把 SM 当成一个黑盒，只知道"线程块扔进去、结果吐出来"。但要理解调度，必须先知道调度器手上有多少资源可分配、每个周期能花多少钱。这一章我们把 SM 的图纸画出来，顺便把几个关键的数字钉死，后面算延迟隐藏的时候要用。

2.1 SM 分四个 SMSP，每个 SMSP 是独立小工厂

从 Volta 开始，NVIDIA 的 SM 结构就固定成了"四分之一分区"的形态。一个 SM 内部被切成 4 个处理块，文档里叫 processing block 或者 sub-partition，我习惯直接叫 SMSP。每个 SMSP 内部包含：一个 warp 调度器、一个指令分发单元（dispatch unit）、一份独立的寄存器堆、一个 L0 指令缓存，以及一组执行单元。

执行单元这一块，从 Volta 到 Hopper 大致是这样演进的：Volta 每个 SMSP 有 16 个 FP32 单元加 16 个 INT32 单元；Turing 把 FP32 和 INT32 合并成 16 个可双用的单元再加 16 个 FP32；Ampere 的 GA10x 每 SM 有 128 个 FP32，分摊到 SMSP 就是 32 个；Hopper 也维持 128 FP32/SM 的规模。除此之外，每个 SMSP 还带一组 LSU（加载存储单元）、SFU（特殊函数单元）、以及 Tensor Core。

关键的一点是：这四个 SMSP 之间是高度独立的。它们各自调度各自的 warp，各自发射各自的指令，互不干涉。一个 SMSP 因为所有 warp 都在等内存而空转，另外三个 SMSP 完全不受影响。这个特性在分析性能时有实际意义：如果你发现 4 个 SMSP 的利用率严重不均，那通常是 warp 分配或者分支分歧导致的问题，而不是带宽问题。

寄存器堆也是按 SMSP 划分的。A100 每个 SM 有 65536 个 32-bit 寄存器，摊到 4 个 SMSP 就是每个 16384 个。这个数字直接决定了你能塞进多少 warp——后面算 occupancy 的时候会用到。

2.2 warp 为什么必须是 32 个线程

这个问题我在很多次分享里被问到过。答案不复杂：GPU 的执行单元是按"通道"组织的，一个 SMSP 里的 32 个 FP32 单元排成一条流水线，一条指令进来，32 个操作数同时分发到 32 个通道上并行执行。这就是 SIMT 的本意。所以一个 warp 自然就是 32 个线程，它对应一次指令发射所能覆盖的最大并行宽度。

换句话说，warp 不是一个抽象概念，它是硬件发射粒度的直接映射。调度器每次做决策，选的是一个 warp，而不是单个线程；指令缓存里存的也是 warp 级的指令流；寄存器分配的物理单位同样是按 warp 组织的。这也解释了为什么当 warp 内部出现分支分歧时，代价那么大——32 个通道里有 16 个走 if 分支、16 个走 else，硬件没有"部分发射"这种能力，只能把两条路径串行执行两遍，各屏蔽掉一半通道。

从 Volta 开始，NVIDIA 引入了独立线程调度（Independent Thread Scheduling），每个线程有了自己的 PC 和调用栈，理论上可以让分歧的两半真正并行推进。但这不等于没有代价：重汇聚点变得不确定，__syncwarp()从"基本没用"变成了"必须显式加"。如果你在 Volta 之后的架构上写依赖 warp 同步的代码（比如 warp 级归约）却没有加__syncwarp()，是有概率踩到正确性问题的。

2.3 每个周期的发射预算与资源上限

调度器的"预算"这个词很贴切：每个 SMSP 的调度器每周期最多只能发射1 条指令。注意是 1 条，不是 2 条。早期 Fermi 时代某些 SM 有双发射能力，但那个设计早被砍掉了，从 Kepler 之后主流架构都是每调度器每周期 1 条。所以一个 SM 每周期最多发射 4 条指令（4 个 SMSP 各一条）。

这个上限意味着什么？意味着指令发射吞吐本身就是一种稀缺资源。如果你的 kernel 每条指令能干的活太少（比如全是标量 4 字节加载），那即使内存带宽没跑满，发射端口也可能先被打满。这是smsp__issue_active和smsp__inst_executed这类指标的意义所在。

我把几个关键上限整理成一张表，方便后面查阅。不同架构的具体数值会有差异，下表以 Ampere GA100（A100）为主，其他架构我会在括号里标注：

注意：寄存器分配有粒度限制。即使你的 kernel 平均每线程只用 20 个寄存器，编译器也会按 warp 粒度（通常是 256 个寄存器）对齐分配，所以算出来的 occupancy 往往比手算值低一档。这一点在做 occupancy 精算时特别容易忽略。

举个具体的数：如果你用__launch_bounds__(256)不限制寄存器，编译器给了 40 个寄存器/线程。一个块 256 线程 = 8 个 warp，每 warp 分配 40×32 = 1280 个寄存器，向上取整到 1280（已经是 256 的倍数，5 个 256 单位）。8 个 warp 就是 10240 个寄存器。65536 / 10240 = 6.4，所以只能放 6 个块 = 48 个 warp，occupancy 是 48/64 = 75%。如果你把寄存器压到 32 个，每块 8192 个，能放 8 个块 = 64 warp，正好 100%。这中间差的就是那 8 个寄存器。

3. warp 调度器：每周期只做一件小事的决策者

现在进入核心。调度器每个周期都在重复同一套动作：扫描候选 warp、判断是否就绪、挑一个、发射一条指令。听起来简单，但里面的策略细节决定了整个 GPU 的行为特征。这一章我把状态机、调度策略和分支处理讲清楚。

3.1 warp 的四种状态：resident / eligible / stalled / selected

任何时刻，一个 SMSP 里的 16 个 warp 槽位中，每个 warp 都处于以下某种状态。理解这四个状态是所有性能分析的基础。

Resident（常驻）指的是这个 warp 已经被分配到这个 SM 上，占着寄存器和槽位。常驻不等于能跑，它可能正在等任何东西。

Eligible（就绪）指的是这个 warp 的下一条指令所需的所有操作数都已经准备好了，理论上可以立刻发射。调度器只会在 eligible 集合里做选择。

Stalled（阻塞）指的是 warp 在等某个资源，典型的有：等全局内存数据回来（long scoreboard）、等共享内存或寄存器依赖（short scoreboard）、等固定延迟的运算结果（wait）、等 barrier（__syncthreads）、等指令缓存填充（no instruction）。每个 stall 都会被硬件计数，这就是 ncu 里那些 stall 指标的来源。

Selected（被选中）指的是这一周期调度器挑了它，指令被送进分发单元。选中的下一周期，这个 warp 大概率会进入 stalled 状态（如果它下一条指令依赖本次结果），也可能继续 eligible（如果下一条是独立的）。

这里有个容易混淆的点：stall_not_selected这个指标代表"warp 已经就绪，但这一周期没被选中"。它高不是坏事，恰恰说明你的并行度很充裕，调度器挑不过来。真正要警惕的是 long scoreboard 和 short scoreboard 占比高。

状态之间的迁移是硬件自动完成的，程序员无法直接干预，但可以通过改变代码结构去间接影响：比如把一条访问拆成多路独立的访问，就能让 warp 在等第一路数据的同时，第二路请求已经发出去了，从而缩短 stalled 的时间占比。

3.2 贪心-最老优先（GTO）：为什么不用简单轮询

调度策略这件事，NVIDIA 从来没有官方公开过细节，但学术界用微基准测试做过大量逆向分析，结论比较一致：实际行为非常接近贪心-最老优先策略，也就是 Greedy-Then-Oldest。

具体怎么理解？假设当前周期有 3 个 warp 处于 eligible 状态。如果上一周期被选中的那个 warp 这一周期仍然 eligible，调度器会继续选它——这就是"贪心"的部分。这个设计的动机很好理解：连续发射同一个 warp 的指令，可以复用上一周期已经译码的结果，省掉一次指令译码的开销，也避免在多个 warp 之间来回切换上下文。

那什么时候换人？当当前 warp 因为依赖阻塞而变成 stalled 时，调度器就把它从候选里踢出去，转向下一个。这时候的选择依据是优先级，而优先级的核心因素是"这个 warp 有多久没发射过指令了"。等待时间最长的那个 warp 优先获得发射机会，这就是"最老"的部分。

为什么要加"最老"这一条？如果只贪心不轮转，理论上可能出现某个活跃 warp 长期霸占发射端口，其他 warp 迟迟拿不到机会，极端情况下会导致部分线程饿死。加入年龄因素后，硬件保证了公平性下限。

这个策略对写代码有两个直接的启示。第一，长的依赖链对调度器是友好的：一个 warp 连续执行十几条独立的计算指令，会被连续发射，效率很高。第二，不要让一个 warp 的指令流频繁地在不同资源之间跳跃，比如算一条、访一次存、再算一条、再访一次，这样每次访存都会强制上下文切换，白白损失发射效率。这实际上是"把访存批量攒起来再一起发"这种优化手法的理论依据。

3.3 双发射的兴衰与 Volta 之后的独立线程调度

我在网上看到过不少人还在讲"SM 每周期可以发射 2 条指令"，这个说法在 Fermi 的某些型号上有过对应的硬件设计，但早就不是主流了。从 Kepler 开始，主流 SM 设计就是每个调度器每周期发射 1 条指令。所以任何基于"双发射"的性能模型推导出来的结论，在现代架构上都不成立。

Volta 带来的真正重要的变化是独立线程调度。在这之前，一个 warp 里的 32 个线程共享一个 PC，遇到分歧只能串行化；Volta 之后每个线程有自己的 PC 和调用栈，warp 在硬件层面被进一步细分成了 4 组、每组 8 个线程来管理，调度器维护的是所谓的"线程束内组"级别的状态。

这带来的实际影响有两个。其一是分歧的代价模型变了：以前是"两条路径串行执行、时间相加"，现在是"两组线程可以交替推进，硬件在每个周期选择一组发射"，总时间仍然会长于无分歧情况，但不完全是简单相加了。其二是同步语义变了：__syncwarp()从"可选"变成了"必需"。我在一个 warp 级 shuffle 归约的代码里就踩过这个坑：在 Kepler 上跑得好好的，换到 Ampere 上偶发结果错误，加了__syncwarp()之后问题消失。

3.4 warp 分歧、重汇聚与 ITL 的实际代价

关于分歧，我想补充一个容易被忽略的细节：分歧的代价不仅取决于分支的条数，还取决于分歧的持续时间。如果 if/else 两个分支各自只有两三条指令，硬件处理起来很快；如果其中一个分支里有一整个循环，那这个循环的每一轮迭代都要在屏蔽状态下执行，代价就很大了。

常见的处理手法有这么几类。第一类是把分歧改成无分支计算，用算术或位运算替代条件跳转，比如用x = flag ? a : b这种三元表达式，编译器通常会生成SEL指令而不是分支，避免了分歧。第二类是按数据重排，把一个块内的线程按照某种属性排序，让同一个 warp 里的 32 个线程尽量走同一条路径，这在处理不规则数据（比如稀疏矩阵、粒子模拟）时很有效。第三类是warp 级协作，把原本需要分支的负载均衡问题，改成所有线程合作为所有数据分流，典型代表是 GPU 上的基数排序。

从 Volta 开始的硬件在分支处理上还有一个新特性：硬件会维护一个主动线程掩码（active mask）栈，自动处理嵌套分歧的重汇聚。但收敛点不再保证一定在分支结束处，所以如果你的代码里有依赖"分支结束后所有线程自动同步"这个假设的逻辑，必须显式加__syncwarp()。

4. 访存延迟隐藏的数学：用 Little's Law 算出你到底要多少 warp

这一章是全文的核心。前面所有关于调度器的讨论，最终都要回答一个问题：要多少并行度，才能让内存带宽跑满。这个问题有精确的数学答案，不需要凭感觉调参。

4.1 各级存储的延迟量级

先把数据准备好。下表是我在 A100 和 RTX 系列上实测得到的量级参考，单位是 SM 时钟周期。不同架构、不同频率下会有浮动，但数量级是稳的：

看这张表要抓住一个关键对比：HBM 延迟是 L1 命中的 20 倍以上。这就解释了为什么很多 kernel 的性能瓶颈看起来是"内存带宽"，实际上是"延迟没藏住、带宽被浪费"——如果每个 warp 发完请求就傻等，LSU 和内存控制器之间会出现大量空档，带宽利用率自然上不去。

另外一个值得关注的点是延迟和频率的关系。HBM 延迟的绝对时间（纳秒）基本由内存控制器和 DRAM 时序决定，但换算成周期数时，SM 频率越高，周期数就越大。所以同一张卡在标称频率和降频状态下，能藏住延迟所需的 warp 数是不一样的。这也是为什么笔记本 GPU 和桌面版同型号的性能差异有时会比纸面参数更大。

实测小技巧：想量准你自己机器上的内存延迟，可以写一个指针追逐（pointer chasing）微基准，让一个 warp 沿随机地址链跳转，用clock64()计时，跳若干次后除以跳数。这个方法比查文档可靠得多，因为实际延迟受显存配置、ECC、温度影响。

4.2 把带宽需求折算成"在飞字节数"

现在上公式。Little's Law 在排队论里的标准形式是：系统中的平均并发数 = 到达率 × 平均停留时间。搬到 GPU 的场景里翻译一下就是：

需要的在飞字节数 = 目标带宽 × 访存延迟

这个式子非常好用，因为它把"我需要多少并发"变成了两个已知量的乘积。让我用 A100 的实际数字算一遍。

A100 SXM 版本的 HBM2e 带宽峰值是 2039 GB/s，SM 数量 108 个，加速频率约 1.41 GHz。先算每个 SM 每个周期需要搬运多少字节：

每 SM 每周期字节数 = 2039e9 / (108 × 1.41e9) ≈ 13.4 bytes/cycle

也就是说，为了让 HBM 带宽满载，每个 SM 平均每个周期要向内存系统提交 13.4 字节的请求。现在乘以延迟，取 500 周期这个中间值：

在飞字节数 = 13.4 × 500 ≈ 6700 bytes

结论：每个 SM 需要大约 6.7KB 的数据持续悬停在访存管道里，才能把 HBM 带宽吃满。

现在把这个数字折算成 warp 级请求。一条 128-bit 的 warp 访存指令（LDG.E.128）请求的字节数是：

32 线程 × 16 bytes = 512 bytes

所以需要的并发请求数是：

6700 / 512 ≈ 13 条 warp 级访存指令同时在路上

分摊到 4 个 SMSP，每个 SMSP 大约需要3.3 条访存指令同时在飞。

这个数字看起来很小！但注意它的前提是"每条在飞指令必须携带 512 字节"。如果你用的是 32-bit 标量加载，一条指令只带 128 字节，那需要的并发指令数就要乘 4，变成 13 条×4 = 52 条每 SM，也就是 13 条每 SMSP。而每个 SMSP 最多只有 16 个 warp 槽位，这意味着你必须让几乎每一个 warp 都时刻挂着一条访存请求，才能勉强达到带宽峰值。没有任何容错余量。

这个推导直接给出了两个优化方向：要么增加每个 warp 携带的字节数（向量化，让一条指令带 512 字节而不是 128 字节），要么增加每个 warp 同时悬挂的请求数（展开循环造 MLP）。两个方向本质是同一件事的两种做法。

4.3 TLP 与 MLP：两条腿走路

上一节算出来的 6.7KB 在飞数据，可以通过两种方式凑出来。这两条路在文献里分别叫 TLP（Thread Level Parallelism，线程级并行）和 MLP（Memory Level Parallelism，内存级并行），我更喜欢把它们叫做"多找几个人排队"和"让每个人多拿几件东西"。

TLP 路线就是提高 occupancy，也就是增加常驻 warp 数。每个 warp 发一条访存指令然后等，靠 warp 数量堆出并发请求。这条路的天花板很清楚：每个 SMSP 最多 16 个 warp，每 warp 一条在飞指令，总共 16 条，如果每条带 128 字节，就是 2KB——远低于 6.7KB 的需求。这解释了为什么纯靠 occupancy 打不满带宽。而且现实更糟，因为常驻 warp 里总有一部分在处理非访存指令（地址计算、循环判断、结果消费），真正挂着访存请求的比例更低。

MLP 路线是让每个 warp 一次发出多条独立的访存请求，然后等所有数据都回来再一起消费。做法有两种：循环展开（unroll）和向量化（vectorized load）。

循环展开的效果最直观。原本的循环是"发请求 → 等 → 用 → 下一轮"，展开 4 路之后变成"发请求 1、2、3、4 → 等 → 四个都用"。四个请求是背靠背发出的，不需要等前一个返回，于是同一个 warp 的在飞字节数从 128 涨到了 512。如果每个 SMSP 有 8 个 warp 处于这种状态，就是 4KB 在飞，配合另外 8 个 warp 的其他活动，基本能顶到目标值。

向量化是另一条更"便宜"的路径。把 4 次 32-bit 加载合并成 1 次 128-bit 加载，指令数少了 4 倍，但在飞字节数不变——因为 512 字节还是一次请求。等等，这里要澄清一个常见误解：向量化的主要收益不是增加在飞字节数，而是减少指令数和提高访存效率。同样的 512 字节，用一条LDG.128发出，只占一个 LSU 流水线槽位、一次地址计算；用四条LDG.32发出，占四个槽位、四次地址计算，还会消耗四倍的发射带宽。所以在发射端口成为瓶颈的场景下，向量化收益巨大；在纯延迟受限的场景下，向量化配合展开才是最优解。

我用一个具体的代码片段说明展开的写法。这是一个累加 kernel 的两种实现：

// 版本 A：无展开，每个 warp 每次只有一条在飞请求 __global__ void sum_naive(const float* __restrict__ x, float* __restrict__ out, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; int stride = gridDim.x * blockDim.x; float s = 0.0f; for (; i < n; i += stride) { s += x[i]; // 加载后立即依赖，warp 立刻 stalled } atomicAdd(out, s); } // 版本 B：4 路展开，一个 warp 同时悬挂 4 条独立请求 __global__ void sum_unroll4(const float* __restrict__ x, float* __restrict__ out, int n) { int base = blockIdx.x * blockDim.x * 4 + threadIdx.x; int stride = gridDim.x * blockDim.x * 4; float s0 = 0.f, s1 = 0.f, s2 = 0.f, s3 = 0.f; for (int i = base; i + 3 * blockDim.x < n; i += stride) { s0 += x[i]; s1 += x[i + blockDim.x]; s2 += x[i + 2 * blockDim.x]; s3 += x[i + 3 * blockDim.x]; // 四条加载互相独立，硬件可以背靠背发射，全部在飞 } atomicAdd(out, (s0 + s1) + (s2 + s3)); }

版本 B 的关键点在于四个累加器s0到s3是独立的。编译器看到s0 += x[i]和s1 += x[i + blockDim.x]之间没有依赖关系，就会把四条加载指令连续排布，全部发出去之后再等结果。硬件层面，四条请求同时在 LSU 和内存系统里穿行，有效延迟被"摊薄"了。如果你写成s += x[i] + x[i+1] + x[i+2] + x[i+3]，那就退化成串行依赖了，因为有单条依赖链。

这里有个细节要提醒：四个累加器不能让编译器优化掉或者合并成一个，所以必须显式写成四个独立变量，且在最后汇总。用#pragma unroll让编译器自动展开循环也能达到类似效果，但编译器展开后的调度顺序不一定符合你的预期，关键路径上我还是倾向于手写。

4.4 occupancy 是入场券，不是成绩单

讲到这里我想强调一个反直觉但很重要的观点：occupancy 只在它限制你的并发请求数时才有意义。

当你的 kernel 处于"带宽受限、MLP 不足"的状态时，提高 occupancy 有帮助，因为更多 warp 意味着更多并发请求。但当你的 kernel 已经通过展开和向量化获得了足够的 MLP，继续提高 occupancy 收益就很小了，甚至可能有害——因为寄存器压力会让编译器减少寄存器用量，导致指令调度空间变小、出现溢出访问（local memory spill），反而拖慢速度。

我在做矩阵乘法的时候深有体会。一个经典的 128×128 tile 分块实现，用__launch_bounds__(256, 2)把每个 SM 限制成 2 个块，寄存器给到 128 个，occupancy 只有 25%（16 个 warp / 64）。但它的性能远好于一个 occupancy 100% 的朴素版本，因为它每个线程扛着 8×8 的输出累加器，计算访存比极高，共享内存访问模式也经过了精心设计。这时候 occupancy 低反而说明寄存器用得好。

所以正确的思路是：先用 ncu 判断瓶颈类型（带宽受限 / 延迟受限 / 计算受限 / 发射受限），再决定往哪个方向使劲。如果smsp__issue_active已经接近 100%，说明发射端口满了，提高 occupancy 没用，得减少指令数。如果dram__throughput只有一半而 long scoreboard 很高，那才是 MLP/TLP 的问题。

5. 实操：把一个带宽跑不满的 kernel 从 42% 拉到 88%

前面都是理论，这一章上真家伙。我拿一个实际手头做过的归约 kernel 做样本，把每一步改动、每一次指标变化都记录下来。测试环境是一张 A100 40GB PCIe 卡，CUDA 12.1，驱动 530 系列，数据量 256M 个 float（1GB 数据）。

5.1 先搭测量基线，别急着改代码

第一步永远是量化。我用 ncu 跑一遍完整指标集，重点关注三类数据：吞吐、occupancy、stall 分布。

# 编译时务必带上行号信息，否则 ncu 无法把指标映射回源码 nvcc -O3 -lineinfo -arch=sm_80 reduce_bench.cu -o reduce_bench # 抓两组指标：全局吞吐 + per-SMSP 的调度与 stall 分布 ncu --kernel-name regex:sum_naive \ --metrics dram__throughput.avg.pct_of_peak_sustained_elapsed,\ sm__warps_active.avg.pct_of_peak_sustained_active,\ smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio,\ smsp__average_warps_issue_stalled_not_selected_per_issue_active.ratio,\ smsp__issue_active.avg.pct_of_peak_sustained_active \ ./reduce_bench

关于-lineinfo：这个选项会略微增大二进制体积、对性能有极小影响，但它让你在 ncu 里能把每一行源码和指标对应起来，性价比远高于那一点点开销。生产构建里可以去掉。

基线数据是这样的（为了可读性我把指标名简化了）：

这组数据把问题定性得非常清楚：occupancy 满、发射半空、stall 集中在 long scoreboard、并且没有多余的就绪 warp 可以调度。典型的 MLP 不足 + TLP 无余量的组合。这时候继续加 block 数量是无效的，因为 occupancy 已经到顶了，得从每个 warp 的在飞请求数下手。

5.2 第一刀：展开循环，把 MLP 从 1 提到 4

按 4.3 节的思路，我把 kernel 改成 4 路展开，四个独立累加器。同时为了防止编译器把展开后的加载顺序打乱，我用__restrict__明确告诉编译器x和out不重叠，给它更多重排自由。

改完之后立刻复测。指标变化是这样的：

long scoreboard 降低是因为同一个 warp 的等待时间被多条请求平摊了——原本等一次 500 周期的数据才做 1 次加法，现在等一次能填满 4 次加法，单位有效工作对应的等待比例下降。not selected 上升是好消息，说明调度器现在有得挑。

但 71.5% 还没到位。继续往下挖。

5.3 第二刀：向量化 + 参数约束

下一刀是两个动作叠加。第一，把 4 次 32-bit 加载换成 1 次 128-bit 加载，减少指令数、降低发射压力。第二，加__launch_bounds__(256)约束块大小，让编译器把寄存器控制在合理范围。

向量化这部分有个必须注意的细节：地址必须 16 字节对齐。我用了__builtin_assume_aligned和显式类型转换来保证这一点，同时因为数据量是 256M 个 float，整除 4，所以尾部不用特殊处理。

__global__ void __launch_bounds__(256) sum_vec4(const float4* __restrict__ x4, float* __restrict__ out, int n4) { int base = blockIdx.x * blockDim.x * 4 + threadIdx.x; int stride = gridDim.x * blockDim.x * 4; float s0 = 0.f, s1 = 0.f, s2 = 0.f, s3 = 0.f; // 每个线程一次拿 4 个 float4，即 16 个 float，凑出 64 字节在飞 for (int i = base; i + 3 * blockDim.x < n4; i += stride) { float4 v0 = x4[i]; float4 v1 = x4[i + blockDim.x]; float4 v2 = x4[i + 2 * blockDim.x]; float4 v3 = x4[i + 3 * blockDim.x]; s0 += (v0.x + v0.y) + (v0.z + v0.w); s1 += (v1.x + v1.y) + (v1.z + v1.w); s2 += (v2.x + v2.y) + (v2.z + v2.w); s3 += (v3.x + v3.y) + (v3.z + v3.w); } atomicAdd(out, (s0 + s1) + (s2 + s3)); }

这个版本每个 warp 每次循环悬挂的字节数是：32 线程 × 4 条 float4 × 16 字节 = 2048 字节。对比基线的 128 字节，提升了 16 倍。按 4.2 节算的需求（每 SM 6.7KB），只要有三四个 warp 处于这个状态就能满足。

实测结果：

注意发射活跃度从 74.8% 掉到了 62.1%，但带宽反而涨了。这完全符合预期：向量化把指令数砍掉了 62%，同样的工作量需要的发射次数大幅减少，所以发射端口可以更闲，而每个周期搬运的字节数更多了。这是一个从"发射受限"转向"带宽受限"的典型信号。

5.4 第三刀：网格配置与尾部效应

88.4% 已经很接近了，剩下的 11.6% 去哪儿了？我抓了一次 timeline 视图，发现两个问题。

第一个是网格跨步循环的尾部效应。当 1GB 数据被 grid-stride 循环切分时，最后一轮迭代中只有一部分线程还在工作，其余都退出了。这时候 SM 上的活跃 warp 数骤降，带宽自然掉下来。解决办法是让网格大小正好整除数据量，或者让每个线程处理的元素数固定，通过多次 kernel launch 覆盖剩余部分。我用的是第二种：算出每个 SM 需要多少块、每块处理多少元素，凑成一个不产生尾部的配置。

第二个是块大小对 warp 调度的影响。我把块从 256 调到 512 再调到 128 各测了一遍，得到的数据挺有意思：

128 略低是因为块数多、调度开销略大；512 差距不大；1024 明显掉下来，因为每 SM 只能塞 2 个块，块数太少导致调度器在不同阶段难以找到足够的就绪 warp，而且寄存器压力让 occupancy 掉到了 75%。结论是 256 到 512 是这类内存密集型 kernel 的甜点区，和 CUDA 官方 best practices 里的建议一致。

最终定格在 88.4%，剩下那 11.6% 主要是 DRAM 刷新开销、ECC 校验、以及 PCIe 版本的带宽上限（PCIe 卡的实际可达带宽通常比 SXM 版本低一些）。

一个容易被忽略的点：dram__throughput这个指标的分母是"理论峰值"，而实际可用的持续带宽通常只有理论值的 90% 到 93%。所以当你的指标跑到 88% 时，很可能已经接近物理极限了，没必要再折腾。判断方法是对比一次cudaMemcpy的带宽——纯拷贝能达到的上限，就是你的 kernel 能达到的上限。

5.5 把优化前后的数据整理成对照

为了让你有个整体印象，我把四个版本的完整对比列在这里。所有数值都是同一张卡、同一份数据、连跑五次的稳定值：

从 1.00 到 0.47，一倍多的提速，全程没有换算法、没有换精度，纯粹是把并发请求数从 128 字节提到了 2048 字节。这就是延迟隐藏在实践里的分量。

6. Nsight Compute 指标怎么读：stall 原因速查表

理论讲完了，实操也做了，但我发现真正卡住大多数人的不是不懂原理，而是打开 ncu 看到一屏指标不知道怎么下手。这一章把最常用的 stall 指标整理成一张速查表，再补充一些使用上的注意事项。

6.1 stall 原因对照：每个数字背后是什么

ncu 会给出每个 SMSP 的平均 stall 分布，这些指标名都很长，我在下面统一简写。所有 stall 指标都是"每条发射指令平均挂着多少个这种状态的 warp"，数值越大说明这种等待越普遍：

用这张表的时候有一个重要技巧：先看占比最高的那一项，而不是看绝对值。如果 long scoreboard 占全部 stall 的 60%，那它就是主因；如果它只占 15%，而 math throttle 占 50%，那你该去优化算术而不是访存。新手最容易犯的错误是看到 long scoreboard 数字大就去改访存，结果发现另一项才是真正的瓶颈。

另外，ncu 里还有一组smsp__average_warps_issue_stalled_*_per_issue_active的指标，和上面这些smsp__average_warps_issue_stalled_*_per_issue_active.ratio是同一族，命名规则在不同版本里有变化，但含义一致：都是"平均每条发射指令背后，有多少个 warp 因为某个原因在等待"。

6.2 采样、replay 与测量误差

ncu 的工作原理是内核重放：它会拦住你的 kernel 执行，采集一部分指标，然后把 kernel 重新跑一遍再采集另一组。一个 kernel 可能被重放几十次。这带来几个实际影响，必须知道。

第一，计时结果不代表真实耗时。ncu 报告的耗时可能比实际运行高几倍，因为它要停下来做采集。想测真实时间必须用cudaEvent或者多次 launch 取平均，别信 ncu 的 duration。

第二，重放会改变缓存状态。如果指标采集需要多次重放，那第一次重放时 L2 是冷的，后续几次可能命中了上一次留下的数据。ncu 通常会尝试在每次重放前清空缓存（用--cache-control all），但这本身也有开销。所以在看 L2 命中率这类指标时要留意这个因素。

第三，--set full非常慢。一个 kernel 全量采集可能被重放上百次，跑几分钟很正常。我的习惯是先跑--set speedoflight或者只抓几个关键指标，定位方向之后再针对性地做深度采集。这样迭代速度快很多。

# 快速定位瓶颈类型：只要 speedoflight 视图里的三行 ncu --set speedoflight -k regex:reduce ./reduce_bench # 定向采集：只看访存延迟相关 ncu -k regex:reduce \ --metrics dram__bytes.sum,lts__t_sector_hit_rate.pct,\ smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio \ ./reduce_bench # 采集并导出报告，方便反复看 ncu -o report_01 --set detailed ./reduce_bench ncu-ui report_01.ncu-rep

关于 ncu 版本：不同 CUDA 版本附带的 ncu 支持的指标名会变。如果你从教程里抄来的指标名报错，先跑一次ncu --query-metrics | grep scoreboard看看实际支持的名称。这个坑我踩过不止一次，尤其是从 CUDA 11 的教程抄到 12 上跑。

6.3 我常用的一套排查顺序

经过这几年的使用，我形成了固定的排查流程，写下来供你参考。

第一步，看 Speed of Light。这个视图给了三个百分比：Compute、Memory、DRAM。谁的百分比最高，谁就是瓶颈候选。如果三个都低于 60%，那多半是延迟问题或者同步问题。

第二步，看 Achieved Occupancy。如果它远低于 Theoretical Occupancy，说明你在启动配置上限制了并行度（块大小、共享内存、寄存器）。如果它和理论值一样高但性能还是差，那不是 occupancy 的问题。

第三步，看 warp state 统计。也就是前面那张 stall 表。找到占比最高的两项，它们大概率能解释 80% 的性能损失。

第四步，看访存效率。l1tex__t_sectors_per_request这个比值能告诉你访存是否合并。理想值是 4（一次 warp 的 128 字节请求正好覆盖 4 个 32 字节 sector）。如果显著大于 4，说明有跨行访问或者随机访问。

第五步，看源码级指标。有了-lineinfo之后，ncu 会把指标映射到具体代码行。切换到 Source 视图，按 stall 数排序，直接能看到是哪几行的访存拖了后腿。这一步往往是最高效的，因为它跳过了所有中间推断，直接告诉你问题在哪一行。

7. 踩坑记录与高频疑问

这一章写点不那么"教科书"的东西。前面几章讲的是方法论，但真正让我掉头发的是那些文档里不会写的细节。

7.1 几个我实际踩过的坑

坑一：--use_fast_math让访存优化失效。有一次我为了提升浮点性能打开了这个开关，结果发现带宽反而降了。排查半天才明白，-use_fast_math会启用一系列激进的优化，包括把float累加重排成允许更早收缩的形式，编译器借此把原本独立的四个累加器重新串成了一条依赖链，MLP 直接归零。这个教训是：数学优化开关和性能优化目标不一定同向，打开之前先测一次。

坑二：PTX 版本不匹配导致的性能漂移。在一次跨机器部署时，我发现同一个二进制在两台配置相同的机器上性能差了 30%。原因是我编译时只指定了-arch=sm_80而没有指定 PTX 版本，导致在一台驱动较早的机器上走了 JIT 编译路径。JIT 出来的 SASS 和 AOT 编译的 SASS 在指令调度上可能有差异，进而影响延迟隐藏效果。正确做法是同时指定 SASS 和 PTX：

nvcc -gencode arch=compute_80,code=sm_80 \ -gencode arch=compute_80,code=compute_80 \ -O3 kernel.cu -o kernel

这样既有针对 sm_80 的原生二进制，又保留了 compute_80 的 PTX 供更高版本架构 JIT，兼顾了兼容性和性能。如果你的卡是 4060Ti（Ada，sm_89）或者 4090，编译目标至少要包含compute_89，否则会走 PTX JIT 路径。

坑三：块大小改一下就从 88% 掉到 79%。前面表格里那个 1024 块大小的数据就是真实的教训。当时我天真地以为"块越大越好，调度开销越小"，结果 occupancy 从 100% 掉到 75%，带宽掉了一大截。在内存密集型 kernel 上，块大小的甜点区通常比计算密集型更窄，因为你需要足够多的块来给调度器留出调度余量。

坑四：环境层面的坑不全是环境问题。有时候你会遇到cuda .run gzip: stdin: invalid compressed>