All You Need to Know About GPU Scoreboarding

从低延迟到高吞吐：两种处理器组织“时间”的不同方式。

大白话说：CPU 在一条指令流后面找活，GPU 在许多 Warp 之间换一个能跑的。

一条加载指令迟迟不返回，CPU 和 GPU 都不能凭空造出结果。它们只能在等待期间找别的工作。真正的差别在于去哪里找：CPU 沿着当前线程继续向后找，GPU 则从已经驻留的其他 Warp 中换一个。

我最初理解 GPU Scoreboard 时，总想先弄清楚它有多少 bit、多少 slot。后来发现顺序反了。Scoreboard 为什么存在，要先从 Warp Scheduler 每个周期面对的问题说起。

本文说的 CPU，是典型的现代乱序 CPU 核心；本文说的 GPU，是典型的面向吞吐的 SIMT GPU。CPU 只作为最后的对照，用来解释 GPU 为什么选择这套机制，不是说每颗 CPU、每颗 GPU 都长成同一个样子。

Warp Scheduler 每个周期都在找谁

一个 kernel 启动后，SM 上通常同时驻留多个 Warp。它们的 PC、活动掩码和线程寄存器已经留在片上，所以从 W0 换到 W1，不需要像操作系统切换线程一样先保存、恢复整套现场。CUDA Programming Guide 对 resident Warp、寄存器资源和 Warp 选择的关系有完整说明。[7]

每个发射机会到来时，Scheduler 都要回答一个问题：现在选哪个 Warp？

不是所有驻留的 Warp 都能选。假设当前有三个 Warp：

W0：下一条指令需要 R1，但 R1 还没写好
W1：下一条指令的操作数都已经准备好
W2：已经走到 __syncthreads()，还在等别的 Warp

W0、W1、W2 都是 resident，也就是现场还在 SM 上；但只有 W1 能参与这一轮选择，这才叫 eligible。

这里很容易把几个状态混在一起。可以先这么理解：

• resident：Warp 的现场还在 SM 上；
• eligible：Active Warp 当前没有被依赖、同步或资源条件挡住，可以交给 Scheduler 仲裁；
• selected：多个 eligible Warp 中，这一轮真正被选中的那个。

所以 W1 本轮没被选中，不代表它 stalled。它可能一直 eligible，只是 Scheduler 这次选了别人。

Scheduler 也不是等某个 Warp 报告“我 stall 了”以后才开始工作。概念上，它每个周期都在本地候选 Warp 中考虑谁能发射。W0 从 eligible 变成 not eligible，只是候选集合少了一个，不是给 Scheduler 发了一次中断。

现实中的 SM 往往还会分成多个处理分区，每个分区有自己的 Warp Scheduler 和候选 Warp。H100 的官方框图把一个 SM 画成四个 processing block，每个 block 都有 Warp Scheduler；Hopper Tuning Guide 给出的上限是每个 SM 同时驻留 64 个 Warp。[12][13]

图 1：NVIDIA H100 Tensor Core GPU Architecture Whitepaper，Figure 7。图中可以直接看到一个 GH100 SM 内画出了四个重复的 processing block，每块都有自己的 Warp Scheduler、Dispatch Unit 和 Register File。[12]

这两个数字不能继续推导成“每个 Scheduler 永远固定管理 16 个 Warp”，也不能证明 Block 内的 Warp 必然按某种简单轮询方式分配。CUDA 编程模型把 Thread Block 分配到 SM，而不是向软件承诺它与某个 Scheduler 一一绑定；同一个 SM 可以同时驻留多个 Block，Scheduler 的候选集合因此也不应理解成某个 Block 的私有队列。精确映射和 Warp 是否迁移属于微架构实现，不是 CUDA 接口的一部分。

问题来了：Scheduler 怎么知道 W0 的 R1 还没写好？

这就需要 Scoreboard。

一条 Load 怎样把 Scoreboard 引出来

先看一段很短的指令流。为了突出依赖，暂时不管真实指令集的拼写：

I0: R1 = load [A]
I1: R2 = R1 + R3
I2: R4 = R5 * R6

I0 从地址 A 读数据，结果写进 R1。I1 马上读取 R1，所以它必须等 I0。I2 只读 R5、R6，与前两条没有数据关系。

设 W0 正在执行这段代码，另外还有 W1、W2：

cycle 0：W0 发射 I0，登记 W0:R1 尚未就绪
cycle 1：W0 的下一条是 I1，I1 需要 R1，因此 W0 不能发射
cycle 1：Scheduler 选择 W1，发射它的 mul
cycle 2：Scheduler 选择 W2，发射它的 add
...
cycle N：I0 返回，W0:R1 就绪，W0 再次具备发射资格

Scoreboard 在这里做的事情很朴素：I0 发射时记一笔“W0 的 R1 还在被生产”；I1 准备发射时查一眼，发现源操作数 R1 仍然 pending，于是不能发射；I0 完成时把这笔账清掉。

即使这次 Load 命中 L1，依赖也不会凭空消失。命中只意味着不用继续访问更远的 L2 或 HBM，数据从 L1TEX 返回并写好目的寄存器仍然需要时间。如果 consumer 在结果就绪前走到队首，它照样要等，只是通常等得更短。这里说的是生产—消费关系，不能据此把 profiler 中的 Short/Long Scoreboard 直接按 cache 层级分类。[9]

GPGPU-Sim 3.x 的模型把这三个动作叫作 reserve、check、release：生产者发射时预留目的寄存器，后续指令发射前检查冲突，生产者写回时释放。[4] 这套说法很适合建立直觉，但它是模拟器模型，不是某代 NVIDIA 芯片公开的门级电路。Bakhoda 等人的论文可以支持当时 GPGPU-Sim 的 SIMT 和 Warp 执行背景，不能替真实芯片证明 Scoreboard 的每一个实现细节。[3]

图 2：W0 的 R1 还没写好，因此眼前的消费者不能发射；W1 的操作数齐全，Scheduler 可以选择 W1。

Producer 发射，不等于整个 Warp 立刻停住

把代码稍微换一下：

I0: R1 = load [A]
I2: R4 = R5 * R6       # 不读取 R1
I1: R2 = R1 + R3       # 读取 R1

I0 发射后，Scoreboard 会登记 R1 pending，但 W0 不一定立刻 stalled。I2 不读取 R1，只要执行管线等其他条件允许，它仍然可以继续发射。只有 I1 走到队首，而 R1 仍然没有返回时，W0 才会因为这项依赖失去 eligible 资格。

所以，Scoreboard 不是给整个 Warp 粗暴地贴一个“前面有慢指令，全部停下”的标签。它登记的是尚未解决的 producer，真正检查的是当前下一条指令会不会撞上这些依赖。

这也解释了 GPU 的 in-order issue。它不是：

I0 complete -> I2 issue

而是：后面的指令不能越过一条尚未 issue 的前序指令。已经 issue 的 I0 可以继续在内存系统里跑，I2 随后发射，甚至先完成。

I0 issue ---------------------------------> complete
         I2 issue ----> complete

因此，一个 Warp 可以同时有多条 in-flight 指令，也可以出现 out-of-order completion。它仍然不是 CPU 意义上的乱序调度：如果队首 I1 因 R1 pending 而不能 issue，GPU 通常不会跳过 I1，再到同一个 Warp 后面寻找 I2。编译器可以提前把独立指令移到 I1 前面；运行时撞到 I1 后，Scheduler 主要去换另一个 Warp。

Register File 与依赖状态

另一个常见误解是把 Register File 和 Scoreboard 当成同一块东西。

Register File 保存真正的数据。一个 Warp 有 32 个线程，每个线程都有自己的 R0、R1、R2……。例如每个线程用了 128 个 32-bit 寄存器，那么这一个 Warp 的数据就占 32 × 128 个寄存器槽位。

但一条 Warp 指令写的是统一的逻辑目的编号：

LDG R1, [R2]

它在活动的 lane 上分别写各线程自己的 R1。对于本文讨论的 Warp 级依赖抽象，Scoreboard 可以记成：

(Warp 0, R1) = pending

而不必先把它理解成 32 条可以独立调度的指令，再复制 32 套完整的 pending 表。若指令带 predicate，可能只有部分 lane 真正写 R1；硬件仍可以保守地把这个 Warp 的逻辑 R1 当成未解决依赖。具体芯片是否结合 active mask 做得更细，公开资料不足，不能从抽象模型反推所有产品的比特布局。

于是 W0:R1 和 W1:R1 也就很好理解了。它们只是逻辑编号相同，实际属于不同线程上下文：

W0:R1 pending
W1:R1 ready

W0 正在写自己的 R1，不会因为名字相同就让 W1 读取自己的 R1 时 stall。

这里还要把依赖跟踪和寄存器分配拆开。寄存器分配回答的是“这个 Warp 的寄存器数据占 Register File 的哪一片资源”，Scoreboard 回答的是“这片上下文里的 R1 现在能不能被下一条指令读取”。前者会影响一个 SM 能驻留多少 Warp，后者影响某个已驻留 Warp 眼前是否 eligible。Scoreboard 里的 (Warp, logical GPR) 只是依赖账目的索引方式，不等于给整个 SM 发明了一个大家共享的“物理 R1”。

不同 Warp 当然仍会互相影响。它们可能访问同一块 shared/global memory，等待同一个 barrier，争抢同一条执行管线、cache 或内存带宽。但这些不是“两个 Warp 都叫 R1”产生的寄存器依赖。需要跨 Warp 传数据时，程序要通过共享／全局内存、原子操作以及正确作用域的同步明确表达顺序。[7]

到这里也不能把“寄存器 ready”和“Warp eligible”画等号。一条指令的寄存器全准备好了，也可能在等同步点、执行管线、发射队列、依赖跟踪项或编译器编码的 stall 周期。Scoreboard 只提供依赖是否满足的信息，Eligible 判断还要合并资源、同步和时序等条件。

Scoreboard 的三层实现：从朴素硬件到软硬件协同

知道它要解决的问题以后，才适合问“这本账怎么记”。这里没有一张能套在所有 GPU 上的永恒电路图，但几种机制之间有一条很清楚的演进线：

第一层：每个 GPR 一个 pending bit
        先解决“值没回来，consumer 不能发射”的正确性问题
        问题：状态多，运行时比较也多

第二层：只记录真正 pending 的 GPR range
        先把大量 ready 状态从表里拿掉
        问题：存储省了，运行时仍要比较 GPR ID

第三层：编译工具链提前编码依赖，硬件只跟踪动态完成状态
        固定延迟用静态调度，变化延迟用 dependency barrier

这三层不是“三代 GPU 各选一种”的产品年表，也不是彼此互斥的单选题。它们是三种逐步减少运行时工作的设计思路；真实芯片完全可能组合使用。

第一层：朴素 Scoreboard——每个 GPR 一个 pending bit

假设一个 Warp 使用 R0 到 R127，最直观的表是：

Warp 0

     R0 R1 R2 R3 R4 ... R127
      0  1  1  0  0 ...    0

1 表示这个逻辑寄存器还有 producer 没完成。下一条指令读取 R1、R6，硬件查到 R1 pending，就不让它发射。

这里不只会碰到 RAW。假设一条慢指令还没有把旧结果写进 R1，后面另一条指令也准备写 R1。虽然两条指令按程序顺序发射，完成顺序却可能反过来；年轻指令若先写回，随后老指令再覆盖它，程序最后看到的值就错了。因此这类直接 Scoreboard 往往也要检查目的寄存器冲突，也就是 WAW。GPGPU-Sim 的模型同时检查 RAW 和 WAW，原因就在这里。[4]

第一层先把正确性解决了：producer 未完成时，consumer 绝不能拿旧值或垃圾值继续跑。但它把代价全部留给硬件。真正 pending 的寄存器通常只占编号空间的一小部分；而发射路径不仅要保存整张状态表，还要快速检查一条指令的多个源、目的寄存器。Warp 越多，需要保存和比较的状态也越多。

于是第一个优化问题很自然：既然大部分寄存器都是 ready，为什么还要给它们永久留位置？

第二层：稀疏 Scoreboard——只记录 pending range

一种自然的压缩办法是只保存正在忙的寄存器。如果连续的 R0、R1、R2、R3 都 pending，可以用一项表示：

slot 0 = { base: R0, count: 4 }

如果当前 pending 集合是 R0-R3 和 R5，单个 base + count 显然表达不了整个集合。抽象的数据结构要么使用两项，要么换一种编码：

slot 0 = R0-R3
slot 1 = R5

这只是说明 range encoding 面对不连续集合时的限制，不代表已经证明某代 NVIDIA GPU 会让一条指令占用多个 slot。US8225076B1 的具体实施例更窄：一条指令若写一串连续目的寄存器，可以只保存第一个寄存器编号和 size indicator，例如用 R0 + size 4 表示 R0-R3。[10]

Producer 完成后，对应项释放。问题随之而来：slot 数量有限，如果所有项都还在使用，新指令发射后又需要登记新的 pending destination，那么即便源操作数 ready、执行单元也空着，它仍然不能安全发射。原因不是 RAW，而是“发出去以后没地方记账”，属于依赖跟踪资源的结构性压力。

这里也不需要把 slot 消耗乘以 32 个线程。这种模型利用 SIMT 指令共享逻辑寄存器编号的特点，按 Warp 的 GPR range 记依赖，不是让每个 lane 各消耗一项。

这项专利能说明压缩连续目的寄存器的思路确实存在，也明确讨论了 storage area 太少会造成 issue bottleneck；但专利不等于每一块量产 GPU 都采用相同结构。

图 3：Bitmap 给整个编号空间留位置；稀疏表示只保存 pending 区间，但发射时仍要做 GPR ID 或区间比较。

第二层解决的是存储，不是依赖发现。当前指令读 R2，硬件仍要判断 R2 是否落在某个 pending range 中。于是又有一个问题：编译器本来就知道谁读 R2、谁写 R2，为什么让硬件运行时重新发现一次？

第三层：软硬件协同——工具链提前编码依赖

第三层不再只琢磨“硬件里的表怎样压缩”，而是重新划分工作。编译器后端拥有整条指令流，天然知道谁读谁写；assembler 可以把调度和依赖信息写进控制字段；运行时硬件只保留那些必须等程序真正执行以后才能知道的状态。

这里首先要把延迟分成两类。依赖关系和完成时间都能在编译期确定时，可以尽量静态解决；依赖关系已知、完成时间却随 cache 和拥塞变化时，才需要动态 completion tracking。

固定延迟：静态调度加 stall cycles

假设 producer 的延迟可以预测：

I0: OP   R1, ...
I1: ADD  R2, R1, R3

编译器后端知道 I1 依赖 I0，也知道两者之间要隔多少周期。它可以先把无关指令塞进中间：

I0: produce R1
J0: independent
J1: independent
J2: independent
I1: consume R1

如果独立指令还不够覆盖延迟，assembler 再把剩余发射间隔编码进 control information 的 stall cycles。它不是机械插入若干 NOP，而是先用有用工作填空档，剩下的时间才明确等待。

GK110/GK210 白皮书回顾 Fermi 时提到，Fermi 对数学管线做更复杂的动态依赖检查；在文档所说的 GK110/GK210 上，固定延迟数学指令更多依靠编译器提供调度信息，而纹理、加载等长延迟操作仍保留寄存器 Scoreboarding。[8] 这句话只适用于资料点名的芯片，不能一句话推广到所有 NVIDIA GPU。

可变延迟：Dependency Barrier

Load 不一样。编译器后端知道 I1 读取 I0 产生的 R1，但不知道 I0 什么时候完成：可能命中 cache，也可能走到更远的内存，还会受运行时拥塞影响。

也就是说：

dependency      编译器后端知道
completion time 编译期不知道

只按静态距离排指令不能保证正确性，于是需要运行时的 dependency barrier。可以把它想成几个编号很少的等待牌：B0、B1……

I0: LDG R1, [A]       write barrier B2
J0: independent       不等待 B2
J1: independent       不等待 B2
I1: ADD R2, R1, R3    wait B2

I0 发射后，B2 变成 busy；J0、J1 不等 B2，可以继续发射；I1 走到队首时检查 B2。若 Load 已完成，B2 已释放，I1 直接发射；若 B2 仍 busy，I1 留在原地，Scheduler 换其他 Warp。

一条 consumer 还可能同时依赖多个 producer：

I0: LDG R1, [A]       write B1
I1: LDG R2, [B]       write B4
I2: ADD R3, R1, R2    wait B1 | B4

因此 consumer 携带的是 Wait Barrier Mask，而不只是一个编号。硬件做的检查接近：

WaitMask & BusyBarrierMask == 0 ?

只有需要等待的 barrier 全部释放，I2 才能发射。相比把每个源 GPR 与所有 pending range 比较，这类控制信息让编译工具链提前编码“谁依赖谁”，运行时只维护少量 barrier 的 busy/free 状态。

图 4：B2 busy 并不会立刻冻住整个 Warp；独立指令照常发射，只有真正等待 B2 的 consumer 到达队首才会停。

Write Barrier 主要保护“结果还没有产生”的 RAW 关系。Read Barrier 处理的是另一个方向。Jia 等人给出的典型场景是未缓冲的内存写操作：它把寄存器内容写向内存，在操作结束前，源寄存器必须保持不变。[5]

I0: STG [A], R1       read barrier B3
I1: MOV R1, R5        wait B3

I0 还在使用旧 R1 时，I1 不能先覆盖 R1，这属于 WAR。Assembler 在 I0 上设置 read barrier，并让后面的 writer 等 B3 释放。

Barrier slot 也不可能无限多。Jia 等人的 Maxwell、Pascal、Volta 结果中有 6 个 dependency barrier。若 B0 到 B5 都仍被需要独立跟踪的操作占用，assembler 不能把新的 producer 随便复用到 busy slot，否则 consumer 会分不清自己到底在等谁。因此 barrier allocation 本身就是有限资源分配，也会限制某些可变延迟操作能够独立跟踪的并发数。[5]

注意，它和前面的稀疏 GPR slot 不是一回事：

• GPR range slot 保存“R0-R3 还不能用”，按寄存器描述依赖；
• dependency barrier 保存“某个 producer 还没完成”，按操作描述依赖；
• consumer 用 wait mask 说明自己等哪些 producer。

Jia 等人对 Maxwell、Pascal、Volta 的逆向研究展示了 stall、yield、wait mask、read/write barrier 等控制字段；Maxwell、Pascal 的控制字组织和 Volta 每条 128 位指令携带控制信息的方式也不同。[5] 这说明编译工具链编码的依赖控制确实是理解这些代际的重要入口，但逆向结果不能自动外推到后来的每一代 GPU。中文资料可作为问题线索，具体硬件结论仍以论文、官方文档和专利为准。[11]

Control code 不只是 Dependency Barrier

把这些字段放在一起看，很容易统称为“软件 Scoreboard”，但它们负责的事情不同：

• wait/read/write barrier：保证依赖正确；
• stall cycles：表达固定延迟下的发射间隔；
• reuse flag：让反复使用的源操作数留在小型 operand reuse buffer，减少 Register File 读取和 bank conflict；
• yield flag：给 Warp 选择提供调度提示。

Reuse 优化的是操作数从哪里取，Yield 影响多个 Warp 都 eligible 时更倾向选谁，它们都不是寄存器依赖正确性的核心。

图 5：Control Code 不是一个笼统的“软件 Scoreboard”。Barrier 字段管依赖正确性，Stall Cycles 管固定延迟时序，Reuse 管操作数读取，Yield 给 Scheduler 提示；具体位域布局随 GPU 代际变化。

现在回头看，第三层也不是“软件取代硬件”。编译器后端给固定延迟指令安排距离，assembler 把 stall 和 dependency barrier 写入控制字段；硬件维护运行时 barrier、资源和必要的依赖状态；Scheduler 最后消费这些结果。

图 6：静态已知的依赖和固定延迟尽量在编译工具链中处理；只有运行时才能知道的完成时刻由硬件 barrier 跟踪。

真正的分界不是“软件还是硬件”

这一段是我觉得最值得琢磨的地方。

编译器后端看得见完整指令流，因此很擅长回答静态问题：谁写 R1，谁读 R1，哪些指令互不依赖，consumer 最早能被排到哪里。它可以据此构造数据依赖图。让运行时硬件再从寄存器编号里推导一遍，很多时候是在重复劳动。

但是编译期看不见这一次 Load 究竟会命中哪一级 cache，也不知道它发射时内存系统有多拥塞。这些信息只有指令真正运行以后才出现。于是合理的边界不是“全交给软件”或者“全交给硬件”，而是：

依赖关系是谁造成的      -> 编译器后端分析
固定延迟怎样填满        -> 后端静态调度
控制字段怎样编码        -> assembler
可变延迟什么时候完成    -> 运行时硬件报告
当前有资格发射哪个 Warp -> Scheduler 仲裁

Dependency barrier 的价值也正在这里。它不是把 Scoreboard 从硬件搬进软件，而是把静态依赖关系压成少量运行时 token。Producer 占住一个 token，完成事件释放它，consumer 只声明自己等哪些 token。工具链把已经知道的关系编码进去，硬件只保留必须等到运行时才能知道的那一小部分状态。

从编译器后端和 assembler 的角度，可以把 barrier index allocation 类比成一个很小的寄存器分配问题：producer 建立一段依赖存活期，consumer 等待它，等这段关系不再需要以后，编号才能安全复用。区别是普通寄存器装的是值，barrier 装的是“某件事完成没有”。编号很少时，工具链既要尽量重叠独立的长延迟操作，又不能让两个仍然活跃的依赖错误地共用同一编号。这里是帮助理解的类比，不是说两种分配算法完全相同。

这也解释了为什么指令排序很重要。若编译器把真正的 consumer 过早放到前面，Warp 一撞上 busy barrier 就停住；写在 consumer 后面的独立指令再多，也不能越过它。编译器不是在锦上添花，它是在决定这些独立工作能不能在 Warp 被卡住以前暴露给硬件。

编译器、依赖硬件和 Scheduler 怎样接起来

继续用前面的 Load，把完整过程走一遍：

I0: LDG  R1, [A]        # variable latency
I2: IADD R4, R5, R6     # independent
I3: IMAD R7, R8, R9     # independent
I1: FADD R2, R1, R3     # depends on I0

编译器后端先分析出 I0 -> I1，并尽量把 I2、I3 移到 consumer 前面；assembler 再让 I0 占用一个 write barrier，并在 I1 的 wait mask 中写入同一个 barrier。

运行时，I0 发射，barrier 变成 busy，Load 进入内存系统。I2、I3 不等它，可以依次发射。这时 I0 没有完成并不妨碍同一个 Warp 继续前进。

等 I1 走到队首，依赖硬件检查 wait mask。如果 I0 已经返回，barrier 是 free，I1 可以参与本轮发射；如果没返回，I1 不能发射。由于同一个 Warp 按序 issue，后面的指令也不能越过 I1。

Scheduler 不需要重新推导“I1 读 R1、I0 写 R1，所以两者有依赖”。编译工具链已经把“等谁”编码出来。Scheduler 看到的是更直接的结果：stall counter 是否归零、wait barrier 是否释放、执行资源是否可用，以及其他条件是否满足。

所以整条链是：

编译器后端分析依赖并安排独立指令
        |
        v
assembler 编码 stall / barrier
        |
        v
运行时硬件维护 pending/barrier/completion
        |
        v
形成当前 eligible Warp 集合
        |
        v
Warp Scheduler 选择并发射

现在终于可以准确回答 Scoreboard 的位置：它不决定“发 W1 还是 W2”，而是参与回答“W1、W2 现在有没有资格发”。

CPU 为什么还要 Tomasulo

GPU 有许多 Warp 可以换，CPU 经常没有这个条件。服务器的一条请求正在查哈希表，桌面线程正在响应一次点击，编译器正在沿指针遍历一棵树。当前线程可能很重要，核心上却没有成百上千份同构工作已经驻留，随时可以替换它。

还是开头的三条指令：

I0: R1 = load [A]
I1: R2 = R1 + R3
I2: R4 = R5 * R6

GPU 的 W0 遇到 I1 以后，主要去换 W1。乱序 CPU 则把 I0、I1、I2 放进一个更深的指令窗口。I1 在等 I0，I2 的两个源操作数已经就绪，于是硬件让 I2 越过尚未执行的 I1，先占用乘法单元。

图 7：CPU 在同一条流里找到 I2；GPU 让 W0 原地等，换 W1、W2 填上发射空档。

要做到这一点，CPU 不能只知道队首指令 ready 不 ready。它要同时保存许多尚未完成的指令，知道每个源操作数在等哪个结果，并在结果到来时叫醒所有消费者。

Tomasulo 的核心做法是使用保留站、标签和结果转发。现代乱序 CPU 又在这条思想主线上加入更完整的物理寄存器重命名、统一调度队列、ROB、分支预测与恢复等结构。

先看寄存器依赖。I0 写 R1，I1 读 R1，这是 RAW，属于真实数据依赖。把 R1 改名也不能让 I1 在数据产生前开始计算。

若较老的 I0 还要读旧 R1，较年轻的 I1 却要写 R1，这是 WAR；若两条指令都写 R1，这是 WAW。后两者只是复用了“R1”这个名字。重命名可以让两次写分别落到 P42、P57，消掉假依赖，再由按序退休保证程序最后看到正确结果。

接着是 wakeup/select。加载结果回来时，标签 P42 被送到等待队列，与许多源标签比较；匹配的指令被唤醒。选择逻辑再从所有已就绪候选中，挑出符合执行端口条件的几条。

结果可以直接转发给消费者，不必先写回再读。指令可以乱序执行和完成，但现代 CPU 通常借助 ROB 按程序顺序退休，从而维持精确异常和正确的软件可见状态。必须说明，1967 年原始 Tomasulo 方案本身没有后来现代 CPU 常见的 ROB；今天说“Tomasulo 风格乱序核心”，通常说的是沿着它发展出来的一整套组合。[1]

图 8：I1 等待 P42，独立的 I2 可以先执行；结果回来后唤醒 I1，ROB 负责现代 CPU 的按序退休边界。

这套硬件并不便宜。窗口越大，标签比较、ready 选择、物理寄存器、旁路网络和错误恢复的成本越高，布线和时序也越难。但对 CPU 来说，只要能让眼前这一条重要线程少等一些，这笔面积和能耗就是值得的。

历史上的 CDC 6600 Scoreboard 也会集中追踪功能单元和冒险，但没有 Tomasulo 后来的寄存器重命名。[2] GPU 依赖 Scoreboard 又主要服务 Warp 队首资格判断。三者都在谈“谁准备好了”，记录对象、调度范围和依靠的并行来源却不同。因此，GPU Scoreboard 不是廉价版 Tomasulo，更不能因为名字相似就把 Tomasulo 叫成“CPU Scoreboard”。

CPU 和 GPU 在赢两场不同的比赛

CPU 花硬件在一个线程的未来里寻找工作。GPU 让程序先暴露大量线程，再从许多 Warp 的现在里寻找工作。

CPU 常见的压力是不规则控制流、指针追逐、缓存 miss 和单个请求的尾延迟。运行时会发生什么很难提前知道，所以硬件负责预测分支、重命名、动态调度，猜错后再恢复。

GPU 常见的是图像、矩阵、科学计算等批量任务。程序本身能提供大量 TLP，编译器还能在一个 Warp 内安排 ILP。硬件没有必要给每个 Warp 复制一套很深的乱序窗口，而是保留许多 Warp 的寄存器、PC、活动掩码和依赖状态，谁能跑就选谁。

两边都要付钱，只是钱花在不同地方。CPU 每多看一条未来指令，就可能多占一个窗口项、ROB 项、物理目的寄存器和若干标签。GPU 每多驻留一个 Warp，就要保存更多线程寄存器和执行状态。寄存器使用量过高会减少同一 SM 可容纳的 Warp，压得太低又可能 spill 到 local memory。GPU Register File 的容量、能耗和并发本来就是一组互相牵制的设计。[6]

它们失败的方式也不同。CPU 的窗口若全被一次长 miss 堵住，或者真正独立的指令还在窗口之外，再复杂的 select 也选不出工作。GPU 若只有很少 Warp 驻留，或者所有 Warp 同时等内存、同步点或同一条管线，Scheduler 也无法凭空造出 eligible Warp。

最朴素的总结是：CPU 更关心这一份工作多快做完，GPU 更关心一段时间内能做完多少份工作。前者主要挖单线程 ILP，后者主要消费多 Warp TLP。这不是绝对分工，但足以解释两种中心设计为什么都合理。

两条路线也在互相靠近

现实产品不会站在教科书坐标轴的两个端点。CPU 使用 SMT，让一个核心同时保留多个硬件线程；向量指令一次处理更多数据；多个未决内存请求还能提供 MLP。CPU 并不只靠单线程 ILP。

GPU 也不是“Warp 内毫无 ILP”。前面已经看到，独立的 I2、I3 可以在 Load 完成前发射。编译器会重排固定延迟指令，程序员可以做软件流水、双缓冲和异步拷贝；Warp specialization 还会让不同 Warp 分别负责加载、计算或传递数据。

但靠近不等于中心消失。CPU 的关键投入仍是让少数通用线程在不规则控制流中尽快前进；GPU 的关键投入仍是让大量显式并行线程高吞吐地共享执行资源。CPU 加入 SMT，不会让 ROB 和重命名失去意义；GPU 加入更复杂的流水，也不会让 resident Warp、依赖资格判断和 Warp Scheduler 消失。

怎样理解 Nsight 中的 Scoreboard Stall

理解上面的链条以后，再看 profiler 会少走很多弯路。

第一，某个 Warp 出现 Scoreboard stall，只能说明它眼前的指令在等依赖，不代表整个 Scheduler 当时没有工作。若其他 Warp 仍然 eligible，发射槽可能一直有人使用。要同时看 Eligible Warps、发射利用率和 No Eligible 等统计。

第二，long scoreboard 不能简单翻译成“等 HBM”，short scoreboard 也不是“L1/L2 命中，所以等得短”。Nsight Compute 的定义是：Long Scoreboard 通常表示等待 L1TEX 路径上的 global、local、surface、texture 等操作；Short Scoreboard 通常表示等待 MIO 路径，常见的是 shared memory，也可能包括特殊数学或动态分支。[9]

因此，看到 long scoreboard，要继续看访存是否合并、cache 行利用率和数据局部性；看到 short scoreboard，要重点排查 shared-memory bank conflict 等 MIO 问题。名字中的 long、short 不能直接拿来判断数据最终来自 L1、L2 还是 HBM。

图 9：Long/Short 描述的是等待所归属的硬件路径，不是把 cache 命中层级粗暴分成“短”和“长”。

第三，优化的目标不是把 occupancy 数字拉满，而是让等待期间还有别的工作。可以更早发出 Load，在 consumer 前安排独立计算，使用软件流水或异步拷贝；也可以降低寄存器、共享内存压力，让更多 Warp 驻留。但如果压寄存器导致 spill，反而会增加 local-memory 访问。如果所有 Warp 都堵在同一个内存瓶颈上，再多几个 Warp 也只是一起排队。

回到最开始的问题：Scoreboard 的作用不是让慢操作变快，而是避免消费者拿到尚未产生的数据，并让 Scheduler 知道这个 Warp 眼下能不能发射。CPU 则愿意用更复杂的 Tomasulo 风格硬件，绕过同一线程中被挡住的指令，继续寻找后面的独立工作。

一种往同一线程后面找，一种换到另一个能跑的 Warp。看懂它们从哪里找工作，后面的寄存器、barrier、控制字段和 profiler 指标就都能串起来了。

参考资料

1. [论文] R. M. Tomasulo，IBM Journal of Research and Development，1967，DOI：10.1147/rd.111.0025，An Efficient Algorithm for Exploiting Multiple Arithmetic Units。
2. [专著] James E. Thornton，Scott, Foresman and Company，1970，Design of a Computer: The Control Data 6600。
3. [论文] Ali Bakhoda 等，ISPASS 2009，DOI：10.1109/ISPASS.2009.4919648；本文仅将其用于模拟器基线、SIMT 与 Warp 上下文，Analyzing CUDA Workloads Using a Detailed GPU Simulator。
4. [模拟器资料] GPGPU-Sim 3.x 模型文档，路径标识 gpgpu-sim_distribution/wiki/test1，GPGPU-Sim 3.x test1 wiki。
5. [逆向研究] Zhe Jia 等，arXiv:1804.06826，§2.1 分析 Maxwell、Pascal、Volta 的指令控制信息，Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking。
6. [技术报告] Hyeran Jeon、Murali Annavaram，USC CENG-2014-05，GPGPU Register File Management by Hardware: Co-operated Register Reallocation。
7. [官方文档] NVIDIA，CUDA Programming Guide：Hardware Multithreading、Independent Thread Scheduling、CUDA C++ Memory Model／Data Races。
8. [官方文档] NVIDIA Kepler GK110/GK210 Architecture Whitepaper，NVIDIA Kepler GK110/GK210 Architecture Whitepaper。
9. [官方文档] NVIDIA，Nsight Compute Profiling Guide：Warp Stall Reasons。
10. [专利] NVIDIA US8225076B1，Tracking register usage during multithreaded processing using a scoreboard having separate memory regions and storing sequential register size indicators。
11. [二手资料] 知乎文章，ID 1987631409838454164，仅作中文延伸阅读与线索，GPU Scoreboard 相关文章。
12. [官方文档] NVIDIA，H100 Tensor Core GPU Architecture Whitepaper。
13. [官方文档] NVIDIA，Hopper Tuning Guide：Occupancy。