A Thread-Register Decoupled GPU Execution Model for Efficient Tensor Computation
Zihan Liu, Jingwen Leng, Yangjie Zhou, Yitong Ding, Guanlin Zhu, Yilu Huang, Chiheng Jin, Chen Zhang, Shixuan Sun, Yu Feng, Anbang Wu, Minyi Guo, Jian Weng, Jiajin Tu, Junsong Wang
cs.AR
2026-08-20
上海交大等提出 FIBER：把 GPU 执行实例从私有寄存器解耦成共享视图的 fiber，在混精度 LLM 服务上相对 Ampere、Hopper、Blackwell 端到端加速 2.25、1.8、2.09 倍，kernel 最高 2.49 倍。
Tensor Core 的峰值吞吐一代比一代高。操作数供给也从 Ampere 的私有寄存器，改成 Hopper、Blackwell 的共享内存和 Tensor Memory，专门消掉 warp 之间重复加载矩阵碎片的浪费。纯 GEMM 这条路已经接近峰值。
现代 LLM kernel 不是纯 GEMM。Attention、FFN、W4A16 反量化会把 Softmax、RoPE、SiLU、查表插在矩阵乘中间。现有 SIMT 卡在两处。
并行度在 launch 时锁死。线程数和每线程寄存器配额不能在 kernel 内改。GEMM 阶段不缺线程，缺操作数带宽；Softmax 这类向量阶段缺线程级并行，寄存器却大量闲着。FlashAttention-3 的现场寄存器轨迹里，占用超过 200 个寄存器的指令只占执行指令的 16.5%。Hopper 的 setmaxnreg 能在 warp group 之间重分寄存器，改不了线程数。CUDA Dynamic Parallelism 只能嵌套 launch kernel，开销太大，撑不起 kernel 内调并行度。
调度太粗。Hopper 之后，TC 相关数据流按 warp group 做阶段式 bulk 同步，要同时对付 mbarrier、named barrier、DEPBAR。FlashAttention-4 里同步停顿已经占到 CPI 的 26%。细粒度依赖追踪其实存在，但只覆盖单 warp 的私有寄存器，跨 warp 只能走共享内存。
FIBER 把执行实例从私有寄存器所有权里拆出来。新的并行单位叫 fiber：只带 PC、编号和少量控制位，不占私有寄存器。同一 SM 上的 fiber 共享整份寄存器堆。按 32 路锁步跑的一组 fiber 叫 weft，对应现在的 warp。
调度器因此可以在运行时增减活跃 fiber 数，不再被 launch 时的寄存器配额绑死。fiber 之间能直接在寄存器里转发中间结果，不必每步绕共享内存。共享寄存器同时充当无冗余的 TC 操作数通路。
落地分三层。
ISA。SASS 指令只剩 13 个未用比特，撑不起给每个操作数加到 16 位。FIBER 只扩展源操作数的高位，按 32-lane 对齐寻址 2048 个向量寄存器，禁止远程写，编译器改写成本地写加远程读。新增 SETFIBERNUM，全 SM 同步后改 fiber 数量。扩展原有 DEPBAR，给每个向量寄存器维护 empty/full 两态，做法接近 Cray MTA 的空满语义。
微结构。保留现有分 bank 寄存器堆，加选择器、4×4 crossbar 做跨 subpartition 读，再加 Read Port Arbiter 处理读端口冲突。一份 Register Busy Bitmap（empty 和 full 各 2048 bit）做跨 weft 依赖。每个 subpartition 最多挂 64 个 weft 上下文，对照基线每 SM 48 个 warp。新增单元在 TSMC 16nm 综合后缩放到 7nm，相对 A100 面积 0.069%、功耗 0.39%，时序走得过 2.5 GHz。
软件。CUDA-native 模式几乎不用改现有 kernel，插入 SETFIBERNUM 即可。FIBER-native 用 pool 声明共享张量，用 setfibernum / setbusybit 编排数据流。编译器做活跃区间分配、按 GEMM 访问模式摆操作数避免 bank 冲突，并把远程写改写掉。
评测用 GTSim 周期级模拟器，对多代 GPU 上 LLM kernel 的 MAPE 在 1.2%–11.3%。工作负载按 Llama 家族：context 4096、dhead 128、hidden 4096、FFN 11008。Attention 在三代基线上分别对照 FA2、FA3、FA4。
| 场景 | Ampere | Hopper | Blackwell |
| 混精度 LLM 服务端到端 | 2.25×（FP16 为 1.15×） | 1.8×（FP16 为 1.13×） | 2.09×（FP16 为 1.11×） |
| 前向 Attention | 1.42× | 1.57× | 1.4× |
| 混精度 FFN | 2.49× | 1.7× | 1.97× |
| FA backward | 1.54× | 1.50× | 1.59× |
W4A16 的 AWQ（计算密集反量化）相对 Ampere 是 2.33×，AQLM（查表反量化）是 1.88×。GQA、MQA 大约 1.3–1.5×。Hopper、Blackwell 上前向 Attention 和 backward 都报到接近或超过 95% 的 TC 利用率。纯投影 GEMM 几乎没增益，操作数供给已经把 TC 喂饱。FlashDecoding 只有 1.12×，主路径是吃内存的 GEMV 和 Softmax，不用 Tensor Core。
相对 Hopper 的消融：动态并行度（DP）在多数非纯 GEMM 负载上贡献最大，最高约 80%；细粒度调度（FS）对投影这种几乎全是同步的 kernel 最明显，FFN 约 69%；寄存器共享（RS）对 MMA 很轻、其他操作很重的 MQA 贡献约 80%。H100 上的 oracle 把三项拆开，分别是 DP 1.21×、FS 1.18×、RS 1.1×。
加速来自把利用率拉向峰值，峰值 TC 吞吐没有加。
给做 kernel 和架构的人看，这篇的判断很具体：下一代瓶颈不在 Tensor Core 算得不够快，而在 SIMT 把执行实例和私有寄存器绑死之后，没法按阶段改并行度，也没法在寄存器粒度做跨 warp 数据流。Hopper 用 warp specialization、setmaxnreg、TMA、wgmma 已经把软件编排推到很复杂，FA4 同步开销还是 CPI 的四分之一。FIBER 用更轻的 fiber 上下文换来 kernel 内动态并行，编程面上用一份共享寄存器视图替换多层 barrier。
如果只写 CUDA-native，这是一次兼容性很好的渐进改动。要吃满约 2 倍的端到端数字，得把 kernel 重写成 FIBER-native，把计算阶段在编译期摊开到特化 weft 上。硬件增量按综合数字可以忽略，前提是模拟器和 16nm 到 7nm 的缩放成立。
混精度服务是主战场：反量化把大量向量操作塞进 GEMM 周围，正好打中 DP 和寄存器数据流。Decode 阶段别指望同样的倍数，1.12× 说明内存墙还在。
全部数字来自 GTSim，不是硅上测量。模拟器自己报的误差上限是 11.3%，对「95% TC 利用率」这类贴顶数字要留余量。
没有流片，也没有完整 nvcc 后端，FIBER-native 靠源到源改 SASS。编译器明确只映射「编译期可确定、不依赖运行时信息」的数据流。动态 shape、不规则控制流会不会把 busy bitmap 和 SETFIBERNUM 的全 SM 同步打穿，论文没测。
Decode 只有 1.12×，纯 GEMM 投影在 Hopper、Blackwell 上几乎为零。端到端 2.25× 高度依赖 W4A16 那种「GEMM 周围塞满逐元素操作」的 serving 设定，和 vLLM 对齐，但不代表 FP16 训练或 decode 为主的在线服务能看到同样的倍数。
连线预算紧的时候，方案退化成 subpartition 内共享，跨 SP 走共享内存，SM 级寄存器数据流这条卖点就打折。远程写被 ISA 禁止，全靠编译器改写。复杂 kernel 的正确性和性能是否仍然干净，只有消融图，没有错误案例分析。
工作负载几乎全是 Llama 结构的 prefill、decode、backward，没有扩散模型、MoE expert 路由，也没有非 LLM 的 HPC kernel。