跳转至

02 Warp 硬件调度器与 Scoreboard 依赖控制

1. Warp 调度器微架构与延迟隐藏 (Latency Hiding)

现代 GPU 的 SM 内部划分为 4 个物理 Sub-Core(Processing Block)。每个 Sub-Core 配备独立的 Dual-Issue Warp 调度器,管理多达 8~16 个活跃 Warp:

graph TD
    subgraph SubCore["SM Sub-Core (Processing Block)"]
        WarpPool["活跃 Warp 候选池 (Active Warps Pool: Warp 0 ~ 15)"]
        Scoreboard["Scoreboard 硬件记分板 (RAW 依赖检查)"]
        Scheduler["Dual-Issue Warp 调度器 (挑选 2 条无冲突就绪指令)"]
        Dispatch["Dispatch Unit (分派至 ALU / Tensor / Memory Pipe)"]
    end

    WarpPool --> Scoreboard
    Scoreboard --> Scheduler
    Scheduler --> Dispatch

1. 零开销上下文轮转 (Zero-Overhead Context Switching)

  • CPU 乱序执行的代价:乱序重排序缓冲(ROB)、保留站(Reservation Station)与复杂重命名寄存器消耗了巨大的芯片面积与功耗。
  • GPU 吞吐设计:硬件将每个活跃 Warp 的寄存器状态直接常驻在片上物理寄存器堆中。当 Warp 0 遇到全局显存 Load 产生长延迟阻塞(~400 cycles)时,调度器在单个时钟周期内无缝切换至 Warp 1、Warp 2 发射计算指令,实现硬件级的零开销延迟隐藏。

2. Scoreboard 依赖追踪与流水线暂停 (Warp Stall)

GPU 硬件通过记分板(Scoreboard)精确追踪操作数依赖:

sequenceDiagram
    autonumber
    participant WS as Warp 调度器
    participant SB as Scoreboard 记分板
    participant Pipe as 计算流水线
    participant Mem as L2/HBM 显存单元

    WS->>Pipe: 发射指令: LDG.E R4, [R2] (全局内存加载)
    Pipe->>SB: 将目标寄存器 R4 置为 Dirty (Pending)
    WS->>SB: 尝试发射: FADD R5, R4, R1 (依赖 R4)
    SB-->>WS: 寄存器 R4 处于 Pending 状态 -> 触发 Stall Long Scoreboard
    Note over WS: 调度器切换至其他就绪 Warp 发射指令...
    Mem-->>Pipe: 内存数据返回,写入物理寄存器 R4
    Pipe->>SB: 清除 R4 Pending 标志
    WS->>Pipe: 成功发射 FADD R5, R4, R1

核心 Warp Stall 指标归因表

Stall 原因分类 对应微架构事件 根因剖析 推荐工程优化手段
Stall Long Scoreboard 等待全局显存/L2 数据加载 显存读取延迟高、未命中 L2 Cache 增加数据预取、提升 Tiling 块大小、增大线程并发度
Stall Short Scoreboard 等待 Shared Memory / MIO 读返回 共享内存 Bank 冲突或局部 RAW 依赖 消除 Bank 冲突、调整指令发射顺序展开循环(ILP)
Stall Barrier Warp 阻塞在 __syncthreads() Thread Block 内部线程负载不均衡 减少不必要的同步屏障、拆分轻量级 Sub-Warp 同步
Stall Not Selected Warp 已就绪但未能抢占发射槽 发射槽位受限或双发射冲突 优化混合指令发射配比(如 FP32 + INT32 混合发射)