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 decouples GPU execution from private registers. Mixed-precision LLM serving is 2.25x/1.8x/2.09x faster end-to-end vs Ampere/Hopper/Blackwell; kernels gain up to 2.49x.
Tensor Core peak throughput keeps climbing. Operand supply moved from Ampere's private registers to Hopper and Blackwell's shared memory and tensor memory, which stop warps from reloading the same matrix fragments. For pure GEMM, that path is already close to peak.
LLM kernels are not pure GEMM. Attention, FFN, and W4A16 dequantization interleave Softmax, RoPE, SiLU, and table lookups with matrix multiplies. Classic SIMT hits two walls.
Parallelism is frozen at launch. Thread count and per-thread register budget cannot change inside a kernel. GEMM phases want operand bandwidth, not more threads. Vector phases want more threads, while most of the register file sits idle. In FlashAttention-3's live register trace, instructions that occupy more than 200 registers are only 16.5% of executed instructions. Hopper's setmaxnreg can rebalance registers across warp groups; it cannot change the thread count. CUDA Dynamic Parallelism only nested-launches kernels, too heavy for intra-kernel scaling.
Scheduling is coarse. Since Hopper, tensor dataflow is coordinated at warp-group granularity with phase-level bulk sync: mbarrier, named barriers, DEPBAR. In FlashAttention-4, synchronization stalls already make up 26% of CPI. Fine-grained dependency tracking exists, but only inside a warp's private registers. Cross-warp reuse still goes through shared memory.
FIBER unbinds the execution instance from private register ownership. A fiber carries a PC, an ID, and a few flags, and nothing else. Fibers on one SM share the whole register file. Thirty-two fibers that issue in lockstep form a weft, the analogue of a warp.
The scheduler can then grow or shrink the active fiber count at runtime. Fibers can forward intermediates through registers instead of bouncing them off shared memory. The same shared registers also feed Tensor Cores without fragment duplication.
Three layers make this real.
ISA. SASS has only 13 unused bits, so a full 16-bit operand address does not fit. FIBER extends source operands only, addresses 2048 vector registers at 32-lane alignment, and forbids remote writes. The compiler rewrites those writes as a local write plus a remote read. SETFIBERNUM is a new SM-wide sync that changes the fiber count. An extended DEPBAR gives each vector register empty and full states, close to the Cray MTA empty/full protocol.
Microarchitecture. The physical register file stays put. A selector and a 4x4 crossbar handle cross-subpartition reads; a Read Port Arbiter breaks the new read-port conflicts. A Register Busy Bitmap (empty and full groups of 2048 bits) tracks cross-weft dependencies. Each subpartition can hold 64 weft contexts, versus 48 warps per SM on the baseline. New logic synthesized at TSMC 16nm and scaled to 7nm costs 0.069% of an A100's area and 0.39% of its power, and closes timing at 2.5 GHz.
Software. CUDA-native mode leaves existing kernels almost unchanged and inserts SETFIBERNUM. FIBER-native declares shared tensors with pool and orchestrates dataflow with setfibernum and setbusybit. The compiler does liveness on the 2048 vector registers, places MMA operands to avoid bank conflicts, and rewrites remote writes.
Evaluation uses GTSim, a cycle-level simulator whose MAPE on LLM kernels across GPU generations is 1.2% to 11.3%. Workloads follow Llama: context 4096, dhead 128, hidden 4096, FFN 11008. Attention baselines are FA2, FA3, and FA4 on Ampere, Hopper, and Blackwell.
| Scenario | Ampere | Hopper | Blackwell |
| Mixed-precision LLM serving (e2e) | 2.25x (1.15x on FP16) | 1.8x (1.13x on FP16) | 2.09x (1.11x on FP16) |
| Forward attention | 1.42x | 1.57x | 1.4x |
| Mixed-precision FFN | 2.49x | 1.7x | 1.97x |
| FA backward | 1.54x | 1.50x | 1.59x |
AWQ (compute-heavy W4A16 dequant) is 2.33x versus Ampere; AQLM (table-lookup dequant) is 1.88x. GQA and MQA sit around 1.3-1.5x. Forward attention and FA backward on Hopper and Blackwell both land near or above 95% Tensor Core utilization. Pure projection GEMM barely moves: operand supply already saturates the TC. FlashDecoding is only 1.12x. Its critical path is memory-bound GEMV and Softmax, and it does not use Tensor Cores.
Ablation against Hopper: dynamic parallelism (DP) is the largest term on most non-GEMM-heavy kernels, up to about 80%. Fine-grained scheduling (FS) helps projection the most, because that kernel is mostly synchronization; FFN sees about 69%. Register sharing (RS) contributes about 80% on MQA, where MMA is light and the surrounding ops are not. An H100 oracle splits the three at 1.21x, 1.18x, and 1.1x.
None of this raises Tensor Core peak. It closes the utilization gap.
The claim for kernel and architecture people is specific. The next bottleneck is not a slower Tensor Core. It is SIMT tying an execution instance to a private register file, which blocks per-phase parallelism and register-grain cross-warp dataflow. Hopper already piled on warp specialization, setmaxnreg, TMA, and wgmma; FA4 still spends a quarter of CPI on sync. FIBER trades a tiny fiber context for intra-kernel dynamic parallelism, and replaces stacked barriers with one shared-register view.
CUDA-native is a compatible, incremental change. The 2x-class end-to-end numbers need FIBER-native kernels that unroll compute phases onto specialized wefts at compile time. Hardware extras look negligible if the synthesis and 16nm-to-7nm scaling hold.
Mixed-precision serving is the sweet spot: dequantization packs vector work around GEMM, which is exactly what DP and register dataflow target. Do not expect the same on decode. 1.12x says the memory wall is still there.
Every number is from GTSim, not silicon. The simulator's own error band goes up to 11.3%, which matters for claims like 95% TC utilization.
There is no tape-out and no real nvcc backend. FIBER-native is a source-to-source SASS rewrite. The compiler only maps dataflow that is compile-time deterministic and independent of runtime information. Whether dynamic shapes or irregular control flow punch through the busy bitmap and SETFIBERNUM's SM-wide sync is untested.
Decode is 1.12x. Pure GEMM projection is essentially flat on Hopper and Blackwell. The 2.25x end-to-end figure is tied to a W4A16 serving mix aligned with vLLM, not to FP16 training or decode-heavy online serving.
Under a tight wiring budget the design falls back to intra-subpartition sharing and sends cross-SP traffic through shared memory, which weakens the SM-wide register-dataflow story. Remote writes are illegal in the ISA and exist only as compiler rewrites. There is an ablation chart, not an error-case study of messy kernels.
Workloads are almost entirely Llama-style prefill, decode, and backward. No diffusion, no MoE expert routing, no non-LLM HPC kernels.