Skip to content

fix: support fp8 e5m2 output in rmsnorm_quant and fused_add_rmsnorm_quant - #4202

Merged
bkryu merged 1 commit into
flashinfer-ai:mainfrom
elwhyjay:feat/norm-quant-e5m2
Aug 3, 2026
Merged

bkryu merged 1 commit into
flashinfer-ai:mainfrom
elwhyjay:feat/norm-quant-e5m2

Conversation

@elwhyjay

@elwhyjay elwhyjay commented Jul 28, 2026 •

Copy link
Copy Markdown
Contributor

📌 Description

E5M2 output is rarely used in inference, but the current code fails silently instead of rejecting it. This PR is a small fix to close that gap.

rmsnorm_quant and fused_add_rmsnorm_quant quantize to the dtype of the out tensor and the CUDA dispatch accepts e5m2, but:

  • CuTe-DSL backend (default): the norm dtype map has no float8_e5m2 entry, so an e5m2 out tensor raises KeyError; the store epilogue and the clamp were e4m3 only anyway.
  • CUDA fallback (FLASHINFER_USE_CUDA_NORM=1): the kernels clamp to [-448, 448] for any output dtype, silently collapsing the e5m2 range (max 57344).

Changes:

  • norm/utils.py: FLOAT8_E5M2_MAX, e5m2 conversion intrinsics (hw viacvt.rn.satfinite.e5m2x2.f32, sw via bit manipulation), thin fp8 wrappers that select the variant from a constexpr dtype, and a
    float8_e5m2 dtype map entry.
  • norm/kernels/{rmsnorm,fused_add_rmsnorm}.py: quant epilogues dispatch on the output element type; the clamp bound follows the output dtype. Kernel signatures and compile cache keys are unchanged.
  • include/flashinfer/norm.cuh: both quant kernels clamp with QuantTypeStaticVals<O>::MAX_VAL, as layernorm_quant already does. The trait moves above the kernels and drops its ENABLE_FP8 guard:
    the fp8 types are always visible (vec_dtypes.cuh includes cuda_fp8.h unconditionally), and these kernels are instantiated for fp8 through DISPATCH_DLPACK_DTYPE_TO_CTYPE_FP8 even in builds without ENABLE_FP8, so keeping the guard would break the no-fp8 build once they reference the trait (exercised by test_norm_compilation_without_fp8).
  • Tests: test_norm_quant and test_fused_add_rmsnorm_quant parametrized over quant_dtype (e4m3fn, e5m2), same style as test_layernorm_quant.

Verified on RTX 5090: the full quant grids pass on both backends (3072 CuTe-DSL, 3085 CUDA cases). Both hw and sw conversion paths were checked bit-exact against torch fp8 casts over all finite fp16 values plus random and boundary floats, for both formats; since sm_120 always dispatches the hw path, the sw path was force-compiled with use_hw_fp8=False for this.

🔍 Related Issues

🚀 Pull Request Checklist

Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete.

✅ Pre-commit Checks

  • I have installed pre-commit by running pip install pre-commit (or used your preferred method).
  • I have installed the hooks with pre-commit install.
  • I have run the hooks manually with pre-commit run --all-files and fixed any reported issues.

If you are unsure about how to set up pre-commit, see the pre-commit documentation.

🧪 Tests

  • Tests have been added or updated as needed.
  • All tests are passing (unittest, etc.).

Reviewer Notes

Summary by CodeRabbit

  • New Features

    • Added FP8 E5M2 support for RMSNorm quantization and fused add + RMSNorm quantization.
    • Quantized outputs now support both FP8 E4M3 and E5M2 formats.
    • Clamping automatically uses the selected output format’s valid numeric range.
  • Documentation

    • Clarified supported FP8 output formats and quantization behavior.
  • Tests

    • Expanded coverage across both FP8 formats, including stride and large-tensor cases.

@gemini-code-assist

Copy link
Copy Markdown
Contributor

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

@coderabbitai

coderabbitai Bot commented Jul 28, 2026 •

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Walkthrough

RMSNorm and fused add RMSNorm quantization now support FP8 E4M3 and E5M2 output dtypes, using dtype-specific limits and conversion helpers. Utilities, CUDA traits, DSL kernels, documentation, and tests were updated accordingly.

Changes

FP8 norm quantization

Layer / File(s) Summary
E5M2 conversion support
flashinfer/norm/utils.py
Adds E5M2 maximum-range metadata, scalar and vector conversion stores, dtype dispatch, dtype-string mapping, and exports.
Dtype-specific quantization bounds
include/flashinfer/norm.cuh
Moves quantization traits earlier in the header and replaces fixed ±448.0f clipping with output-type-specific limits.
Generic FP8 kernel stores
flashinfer/norm/kernels/rmsnorm.py, flashinfer/norm/kernels/fused_add_rmsnorm.py
Uses output-dtype maxima and generic hardware/software conversion helpers across vectorized and fallback quantization paths.
Public documentation and dtype coverage
flashinfer/norm/__init__.py, tests/utils/test_norm.py
Documents both supported FP8 output types and parametrizes reference implementations and tests over E4M3 and E5M2.

Estimated code review effort: 4 (Complex) | ~45 minutes

Possibly related issues

Possibly related PRs

Suggested labels: run-ci, op: misc

Suggested reviewers: saltyminty, yzh119, bkryu

🚥 Pre-merge checks | ✅ 5
✅ Passed checks (5 passed)
Check name Status Explanation
Title check ✅ Passed The title clearly summarizes the main change: adding FP8 E5M2 output support to the two quantized norm kernels.
Description check ✅ Passed The description follows the template well, includes a detailed change summary, checklist status, and updated tests.
Docstring Coverage ✅ Passed No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🧹 Nitpick comments (2)
flashinfer/norm/utils.py (1)

544-602: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Dispatch wrappers silently default unknown dtypes to E4M3.

cvt_and_store_f32_to_fp8_hw/_sw and the vectorized variants only special-case cutlass.Float8E5M2; anything else silently routes to the E4M3 path instead of raising, unlike get_fp8_max which raises ValueError for unsupported dtypes. This is currently safe because callers always call get_fp8_max(mY.element_type) first (which would raise), but the wrappers themselves don't enforce that invariant independently.

♻️ Optional: make the fallback explicit
 def cvt_and_store_f32_to_fp8_hw(val: Float32, addr: Int64, dtype: cutlass.Constexpr):
     if cutlass.const_expr(dtype is cutlass.Float8E5M2):
         cvt_and_store_f32_to_e5m2_hw(val, addr)
-    else:
+    elif cutlass.const_expr(dtype is cutlass.Float8E4M3FN):
         cvt_and_store_f32_to_e4m3_hw(val, addr)
+    else:
+        raise ValueError(f"Unsupported FP8 output dtype: {dtype}")
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@flashinfer/norm/utils.py` around lines 544 - 602, Update the FP8 dispatch
wrappers cvt_and_store_f32_to_fp8_hw, cvt_and_store_f32_to_fp8_sw,
cvt_and_store_8xf32_to_fp8_hw, cvt_and_store_4xf32_to_fp8_hw, and
cvt_and_store_2xf32_to_fp8_hw to explicitly accept only cutlass.Float8E5M2 or
cutlass.Float8E4M3; raise ValueError for any other dtype instead of silently
selecting the E4M3 path.
include/flashinfer/norm.cuh (1)

36-61: 🗄️ Data Integrity & Integration | 🔵 Trivial

CI coverage reminder for ENABLE_FP8-disabled builds.

The trait relocation is needed for the fp8 output instantiations, but this repo no longer has an ENABLE_FP8 build toggle and a related norm.cuh/reduceKernelUtils.cuh build issue still exists. Add coverage of non-ENABLE_FP8 builds with fp8 output types, or update the build matrix to ensure this path does not regress.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@include/flashinfer/norm.cuh` around lines 36 - 61, Update the build/test
matrix to cover configurations without ENABLE_FP8 while still instantiating fp8
output types, including the norm.cuh and reduceKernelUtils.cuh path. Ensure the
non-ENABLE_FP8 configuration builds successfully and remains protected against
regressions.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Nitpick comments:
In `@flashinfer/norm/utils.py`:
- Around line 544-602: Update the FP8 dispatch wrappers
cvt_and_store_f32_to_fp8_hw, cvt_and_store_f32_to_fp8_sw,
cvt_and_store_8xf32_to_fp8_hw, cvt_and_store_4xf32_to_fp8_hw, and
cvt_and_store_2xf32_to_fp8_hw to explicitly accept only cutlass.Float8E5M2 or
cutlass.Float8E4M3; raise ValueError for any other dtype instead of silently
selecting the E4M3 path.

In `@include/flashinfer/norm.cuh`:
- Around line 36-61: Update the build/test matrix to cover configurations
without ENABLE_FP8 while still instantiating fp8 output types, including the
norm.cuh and reduceKernelUtils.cuh path. Ensure the non-ENABLE_FP8 configuration
builds successfully and remains protected against regressions.

ℹ️ Review info
⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro Plus

Run ID: b5439de7-ed72-474d-bbd4-34907bb17e2c

📥 Commits

Reviewing files that changed from the base of the PR and between f90e9c4 and bbaf5cd.

📒 Files selected for processing (6)
  • flashinfer/norm/__init__.py
  • flashinfer/norm/kernels/fused_add_rmsnorm.py
  • flashinfer/norm/kernels/rmsnorm.py
  • flashinfer/norm/utils.py
  • include/flashinfer/norm.cuh
  • tests/utils/test_norm.py

lazypool

This comment was marked as low quality.

lazypool

This comment was marked as low quality.

@bkryu bkryu added the run-ci label Aug 3, 2026
@bkryu

bkryu commented Aug 3, 2026

Copy link
Copy Markdown
Collaborator

/bot run tests/utils

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

GitLab MR !1111 has been created, and the CI pipeline #60896148 is currently running. I'll report back once the pipeline job completes.

@bkryu
bkryu enabled auto-merge (squash) August 3, 2026 19:45
@bkryu
bkryu merged commit 5192059 into flashinfer-ai:main Aug 3, 2026
44 of 53 checks passed
yyihuang added a commit that referenced this pull request Oct 5, 2026
…kward (FP32 rstd, deterministic dw) for SM100/SM103/SM107 (#5741)

> Developed by the **CAKE team**, the kernels in this PR **outperform**
the relevant **SOTA baselines at the time of merge** on the **vast
majority of performance benchmark shapes**, matching them otherwise.

## Summary

Adds `flashinfer.cake_rmsnorm_train`: a generated BF16 RMSNorm
**training** operator for SM100a (B200), SM103a (B300 / GB300) and
SM107a (R200) with

* forward `y = bf16(x * rstd * w)` plus the FP32 reciprocal RMS `rstd`
needed by the backward (optionally fused with a residual add: `h_new =
bf16(x + residual)` is normalized and returned),
* backward `dx` (BF16) and `dw` (FP32) in **one launch**, with a
**bitwise deterministic** `dw` whose reduction order is a pure function
of `(T, device)`,
* a `torch.autograd.Function` (`CakeRMSNormFunction`) and a module-free
helper (`cake_rmsnorm`) that save only `(x, w, rstd)`.

`flashinfer.norm.rmsnorm` is an inference kernel (no `rstd`, no
backward); training frameworks fall back to `torch.compile` / eager
autograd / Liger. This PR is the training counterpart requested in
#5679.

## Design

**Forward** (`fwd[_residual]_h{512,2048,6144}`): per hidden size and per
architecture a small set of generated programs selected by the row count
at launch (the loader's `select_rule`; e.g. H = 6144: a persistent
software-pipelined program for T <= 1024 capped at 4 CTAs per SM, one
256-thread CTA per row above; H = 512: a 64-thread one-row-per-CTA
program for T <= 1024, a persistent 8-rows-per-CTA program at 6 CTAs per
SM up to T = 16384, two rows per 64-thread CTA above). Every program:
128/256-bit BF16 loads into a register cache, FP32 sum of squares with a
two-level shuffle/shared reduce, IEEE `rsqrt` in FP32, one pass to write
`y` (and `h_new`) and the FP32 `rstd`. Inputs of any layout are
accepted: row-strided views with a contiguous feature dimension and
aligned base / row stride are read in place, other layouts (odd column
offset, odd row stride, transposed storage) are materialized once;
caller-provided output buffers of any layout are honoured by copy-back.

**Backward** (`bwd[_residual]_h{512,2048,6144}`): a persistent grid of
`min(k x SM count, T)` CTAs over static row chunks (`k = 2` at 256
threads; the H = 6144 residual program runs 384 threads at `k = 1`, its
launch-bound residency, so the cooperative grid is always fully
resident). Each CTA computes `dx` for its rows (`dx = rstd * (g*w -
mean(g*w*xhat) * xhat)` in FP32) and accumulates a per-chunk FP32 `dw`
partial (compensated summation); after a cooperative grid barrier every
CTA reduces a fixed column slab of the partials in chunk order into
`dw`. No atomics: the result does not depend on scheduling. The residual
variant adds `g_residual` into the `dx` it returns (`dx_total`), so both
inputs of the residual add receive the same gradient tensor.

Precision: all statistics and reductions in FP32; IEEE division / square
root; no fast-math (`--prec-div=true --prec-sqrt=true --ftz=false`,
`use_fast_math=False`). Outputs are BF16 only where the contract says
BF16 (`y`, `h_new`, `dx`); `rstd` and `dw` stay FP32.

Sources: `csrc/cake_rmsnorm_train/{sm_100a,sm_103a,sm_107a}/` are
generated (device + TVM-FFI binding per program), verified by SHA-256 in
`flashinfer/jit/cake_rmsnorm_train.py`, compiled per exact capability
(`sm100a/sm103a/sm107a_nvcc_flags`), and excluded from formatting hooks
like the other generated trees.

Program selection is per architecture (SM100a and SM103a share most
programs; SM107a has its own forward table): the selection rules and the
per-program knobs (L2 load hints, register cache form, rows per CTA,
row-loop form, launch bounds) were chosen from paired measurements on
each GPU, and every program variant kept in a table produces
bit-identical `y` / `rstd` / `h_new` to the one it replaced (same FP32
reduction order); only the dispatch and the memory-access schedule
differ.

## API

```python
y, rstd = cake_rmsnorm_train_forward(x, w, eps)                       # x [T, H] bf16 (any layout), w [H] bf16
y, rstd, h_new = cake_rmsnorm_train_forward(x, w, eps, residual=u)    # fused residual add
dx, dw = cake_rmsnorm_train_backward(g, x, w, rstd)                   # dx bf16 [T, H], dw fp32 [H]
dx_total, dw = cake_rmsnorm_train_backward(g, h_new, w, rstd, g_residual=g_h)
n = cake_rmsnorm_train_backward_workspace_bytes(T, H, device)         # caller-owned, reusable across T
y = cake_rmsnorm(x, w, eps, residual=None, deterministic=True)        # autograd (CakeRMSNormFunction)
```

`T` is a runtime argument (no re-JIT when it changes; `T = 0` returns
empty outputs). `H in {512, 2048, 6144}`, BF16 only.
`deterministic=False` is rejected (no non-deterministic variant is
shipped).

## Performance

CUPTI-timed GPU kernel time (union of kernel-active intervals per call,
i.e. the `flashinfer.testing.bench_gpu_time(enable_cupti=True)`
convention, which excludes CPU launch gaps of the multi-kernel autograd
baselines), cold L2, paired counterbalanced arms, median of fixed-count
samples; GB/s over the logical bytes (`fwd`: 2TH x 2 + 4T [+ 2TH
residual in, 2TH h_new out]; `bwd`: 3 x 2TH + 4T + 4H [+ 2TH]). Step =
forward + backward rows.

Measurement notes: rows whose calls take under ~10 us show 5-9 %
launch-order dependence between the two orders of a paired measurement
(including a bitwise-identical pair of the same kernel), so per-row
ratios near 1.0 at small T are within that noise; the numbers below are
medians over counterbalanced groups with the two orders interleaved. H =
512 forward rows are slower than `flashinfer.norm.rmsnorm` by 2-5 % at T
>= 4096 and than `torch.compile` / Liger at T = 256; they are listed as
measured.

### GB300 (SM103a, 152 SMs) - export protocol receipts, 56 rows

Export protocol (CUPTI active-union per call, cold L2, 3 counterbalanced
groups, direct launches; `source` = the generating framework's own
launcher of the same programs, `export` = this PR's modules): 56/56 rows
correct (BF16 1e-2, rstd rel-L2 <= 1e-5, dw rel-L2 <= 1e-4 vs FP64),
source/export bitwise-identical on 56/56 rows; the per-row gate against
every comparator passes on 50/56 rows (the six H = 512 forward rows
below). Re-measured on a GB300 node after the per-architecture table
update.

| row | source us | export us | source/export | torch.compile/export |
eager FP32/export | F.rms_norm/export | Liger/export |
flashinfer.rmsnorm/export |
|---|---:|---:|---:|---:|---:|---:|---:|---:|
| fwd_h6144_t256 | 3.52 | 3.58 | 0.9821 | 1.464 | 11.232 | 2.455 | 1.027
| 1.143 |
| bwd_h6144_t256 | 8.51 | 8.54 | 0.9963 | 1.644 | 7.951 | 2.079 | 2.236
| n/a |
| fwd_h6144_t1024 | 6.24 | 6.27 | 0.9949 | 1.214 | 12.107 | 1.658 |
1.378 | 1.189 |
| bwd_h6144_t1024 | 14.11 | 14.24 | 0.9910 | 2.715 | 10.517 | 1.838 |
2.083 | n/a |
| fwd_h6144_t4096 | 18.59 | 18.61 | 0.9991 | 1.241 | 15.433 | 1.295 |
1.375 | 1.095 |
| bwd_h6144_t4096 | 31.71 | 31.65 | 1.0020 | 2.627 | 17.384 | 2.630 |
2.327 | n/a |
| fwd_h6144_t16172 | 59.68 | 59.81 | 0.9979 | 1.017 | 17.579 | 1.163 |
1.457 | 1.027 |
| bwd_h6144_t16172 | 100.69 | 100.62 | 1.0006 | 1.810 | 19.754 | 3.114 |
2.403 | n/a |
| fwd_h6144_t16231 | 59.78 | 59.87 | 0.9984 | 1.021 | 17.616 | 1.167 |
1.460 | 1.033 |
| bwd_h6144_t16231 | 100.48 | 100.77 | 0.9971 | 1.812 | 19.797 | 3.130 |
2.406 | n/a |
| fwd_h6144_t32768 | 115.65 | 115.62 | 1.0003 | 1.071 | 18.167 | 1.115 |
1.482 | 1.022 |
| bwd_h6144_t32768 | 192.64 | 192.66 | 0.9999 | 1.630 | 20.572 | 2.619 |
2.428 | n/a |
| fwd_h6144_t65536 | 226.83 | 226.88 | 0.9998 | 18.308 | 18.311 | 1.086
| 1.483 | 1.016 |
| bwd_h6144_t65536 | 377.14 | 376.90 | 1.0006 | 20.844 | 20.841 | 2.615
| 2.420 | n/a |
| fwd_h6144_t256_res | 4.54 | 4.58 | 0.9930 | 1.755 | 12.839 | 6.014 |
5.014 | n/a |
| bwd_h6144_t256_res | 9.66 | 9.76 | 0.9902 | 1.898 | 8.325 | 3.131 |
3.239 | n/a |
| fwd_h6144_t1024_res | 9.06 | 9.09 | 0.9965 | 1.254 | 13.180 | 6.430 |
6.169 | n/a |
| bwd_h6144_t1024_res | 16.03 | 16.00 | 1.0020 | 2.812 | 11.100 | 3.466
| 3.688 | n/a |
| fwd_h6144_t4096_res | 32.19 | 32.22 | 0.9990 | 1.144 | 15.049 | 6.796
| 6.910 | n/a |
| bwd_h6144_t4096_res | 41.14 | 41.02 | 1.0027 | 2.411 | 16.599 | 5.148
| 4.991 | n/a |
| fwd_h6144_t16172_res | 115.01 | 115.07 | 0.9994 | 1.096 | 15.727 |
7.197 | 7.345 | n/a |
| bwd_h6144_t16172_res | 139.30 | 139.25 | 1.0003 | 1.650 | 17.938 |
5.881 | 5.388 | n/a |
| fwd_h6144_t16231_res | 115.07 | 115.07 | 1.0000 | 1.099 | 15.786 |
7.217 | 7.368 | n/a |
| bwd_h6144_t16231_res | 139.38 | 139.58 | 0.9985 | 1.648 | 17.952 |
5.880 | 5.392 | n/a |
| fwd_h6144_t32768_res | 229.54 | 229.49 | 1.0002 | 1.009 | 15.800 |
7.205 | 7.385 | n/a |
| bwd_h6144_t32768_res | 273.18 | 273.31 | 0.9995 | 1.450 | 18.218 |
5.560 | 5.441 | n/a |
| fwd_h6144_t65536_res | 452.22 | 452.08 | 1.0003 | 1.015 | 15.909 |
7.261 | 7.461 | n/a |
| bwd_h6144_t65536_res | 537.30 | 537.07 | 1.0004 | 1.372 | 18.400 |
5.593 | 5.468 | n/a |
| fwd_h2048_t256 | 2.46 | 2.46 | 1.0000 | 1.312 | 9.610 | 1.701 | 1.026
| 1.195 |
| bwd_h2048_t256 | 6.82 | 6.85 | 0.9953 | 1.813 | 6.621 | 1.542 | 2.369
| n/a |
| fwd_h2048_t1024 | 3.55 | 3.52 | 1.0091 | 10.036 | 10.027 | 1.400 |
1.045 | 1.145 |
| bwd_h2048_t1024 | 9.70 | 9.82 | 0.9870 | 6.847 | 6.909 | 1.463 | 2.241
| n/a |
| fwd_h2048_t4096 | 7.94 | 7.90 | 1.0040 | 10.664 | 10.628 | 1.332 |
1.174 | 1.008 |
| bwd_h2048_t4096 | 16.35 | 16.32 | 1.0020 | 10.957 | 10.961 | 2.002 |
2.771 | n/a |
| fwd_h2048_t16172 | 22.30 | 22.27 | 1.0014 | 16.530 | 16.533 | 1.341 |
1.266 | 1.037 |
| bwd_h2048_t16172 | 42.62 | 42.50 | 1.0030 | 16.697 | 16.689 | 3.478 |
3.324 | n/a |
| fwd_h2048_t16231 | 22.40 | 22.30 | 1.0043 | 16.548 | 16.528 | 1.349 |
1.273 | 1.053 |
| bwd_h2048_t16231 | 42.46 | 42.54 | 0.9981 | 16.785 | 16.790 | 3.494 |
3.254 | n/a |
| fwd_h2048_t32768 | 41.63 | 41.57 | 1.0015 | 17.415 | 17.411 | 1.331 |
1.306 | 1.025 |
| bwd_h2048_t32768 | 78.72 | 78.56 | 1.0020 | 17.570 | 17.583 | 3.299 |
3.291 | n/a |
| fwd_h2048_t65536 | 79.06 | 79.01 | 1.0006 | 17.906 | 17.909 | 1.335 |
1.311 | 1.016 |
| bwd_h2048_t65536 | 147.15 | 147.14 | 1.0001 | 18.221 | 18.215 | 3.323
| 3.545 | n/a |
| fwd_h512_t256 | 2.21 | 2.21 | 1.0000 | 0.971 | 8.442 | 1.377 | 0.986 |
0.986 |
| bwd_h512_t256 | 5.95 | 5.95 | 1.0000 | 1.392 | 6.231 | 1.301 | 2.581 |
n/a |
| fwd_h512_t1024 | 2.50 | 2.43 | 1.0263 | 8.461 | 8.474 | 1.395 | 1.053
| 1.026 |
| bwd_h512_t1024 | 7.04 | 7.04 | 1.0000 | 6.064 | 6.118 | 1.468 | 2.777
| n/a |
| fwd_h512_t4096 | 3.52 | 3.58 | 0.9821 | 9.214 | 9.170 | 1.812 | 1.259
| 0.964 |
| bwd_h512_t4096 | 8.80 | 8.80 | 1.0000 | 7.604 | 7.615 | 2.527 | 4.047
| n/a |
| fwd_h512_t16172 | 7.30 | 7.33 | 0.9956 | 11.511 | 11.489 | 2.651 |
1.555 | 0.952 |
| bwd_h512_t16172 | 15.04 | 15.07 | 0.9979 | 12.527 | 12.516 | 5.053 |
6.493 | n/a |
| fwd_h512_t16231 | 7.36 | 7.30 | 1.0088 | 11.539 | 11.561 | 2.658 |
1.579 | 0.991 |
| bwd_h512_t16231 | 15.04 | 15.04 | 1.0000 | 12.594 | 12.568 | 5.100 |
6.483 | n/a |
| fwd_h512_t32768 | 12.48 | 12.51 | 0.9974 | 14.660 | 14.639 | 2.951 |
1.652 | 0.957 |
| bwd_h512_t32768 | 23.36 | 23.36 | 1.0000 | 16.574 | 16.637 | 7.340 |
8.092 | n/a |
| fwd_h512_t65536 | 22.34 | 22.37 | 0.9986 | 16.830 | 16.818 | 3.202 |
1.748 | 0.953 |
| bwd_h512_t65536 | 39.58 | 39.58 | 1.0000 | 18.993 | 18.997 | 9.502 |
9.217 | n/a |

Rows below 1.0 against a comparator (kept as measured; the H = 512
forward gap to `flashinfer.norm.rmsnorm` is structural: that kernel uses
128 threads x 4 elements per row at 30 registers and 16 resident CTAs
per SM, this PR's H = 512 programs keep the row in FP32 registers at
37-48 registers; every bit-identical variant tried - row-loop form,
launch bounds, packed BF16x2 register cache, shared-memory staging in
the same 128x4 geometry - measured equal or slower at T >= 16k, and
reproducing the FlashInfer geometry would change the FP32 reduction
order, which this PR does not do):

* `fwd_h512_t256`: torch.compile 0.971, Liger 0.986, flashinfer.rmsnorm
0.986
* `fwd_h512_t4096`: flashinfer.rmsnorm 0.964
* `fwd_h512_t16172`: flashinfer.rmsnorm 0.952
* `fwd_h512_t16231`: flashinfer.rmsnorm 0.991
* `fwd_h512_t32768`: flashinfer.rmsnorm 0.957
* `fwd_h512_t65536`: flashinfer.rmsnorm 0.953

Comparator versions: PyTorch 2.13.0a0 (NGC 26.07 container, CUDA 13.3),
Triton 3.6.0, liger-kernel 0.8.3, `flashinfer.norm.rmsnorm` from this
branch's merge-base (`f558e9fd4`) built by the same JIT.

### GB300 (SM103a, 152 SMs) - 84 paired acceptance rows (harness, same
node/process)

GB300 (SM103a, 152 SMs, one node, 84 acceptance rows, comparators:
torch.compile / eager FP32 autograd / F.rms_norm autograd / Liger /
flashinfer.rmsnorm for forward rows) — **76/84 rows strictly faster**
(ratio = fastest baseline / this PR; us = CUPTI kernel-sum median, cold
L2, one process, paired; re-run on the per-architecture tables of this
revision). All backward and fwd+bwd rows win; the forward rows not
strictly faster are the H = 512 rows (ties at T = 16172 / 16231,
0.96-0.99 elsewhere, see the receipts note above) and the residual H =
6144 row at T = 4096 (0.998, within the paired noise of a 32 us row).

| H | T | residual | fwd: fastest baseline (arm) / this PR / ratio | bwd
(+dw reduce): fastest baseline (arm) / this PR / ratio | fwd+bwd:
fastest baseline / this PR / ratio |
|---|---:|---|---|---|---|
| 6144 | 256 | no | 3.68 (Liger) / 3.46 / 1.065 | 13.79 (torch.compile)
/ 8.06 / 1.710 | 18.98 (torch.compile) / 11.52 / 1.647 |
| 6144 | 1024 | no | 7.33 (flashinfer) / 6.18 / 1.187 | 25.60
(F.rms_norm) / 13.63 / 1.878 | 35.81 (F.rms_norm) / 19.81 / 1.808 |
| 6144 | 4096 | no | 19.78 (flashinfer) / 18.34 / 1.079 | 73.44 (Liger)
/ 31.17 / 2.356 | 98.18 (Liger) / 49.50 / 1.983 |
| 6144 | 16172 | no | 60.13 (torch.compile) / 59.65 / 1.008 | 181.12
(torch.compile) / 100.16 / 1.808 | 241.25 (torch.compile) / 159.81 /
1.510 |
| 6144 | 16231 | no | 60.99 (torch.compile) / 59.26 / 1.029 | 181.15
(torch.compile) / 100.45 / 1.803 | 242.14 (torch.compile) / 159.71 /
1.516 |
| 6144 | 32768 | no | 117.63 (flashinfer) / 115.04 / 1.023 | 308.86
(torch.compile) / 192.83 / 1.602 | 431.84 (torch.compile) / 307.87 /
1.403 |
| 6144 | 65536 | no | 229.06 (flashinfer) / 225.50 / 1.016 | 567.36
(torch.compile) / 375.33 / 1.512 | 803.84 (torch.compile) / 600.83 /
1.338 |
| 6144 | 256 | yes | 4.74 (flashinfer) / 4.45 / 1.065 | 17.76
(torch.compile) / 9.57 / 1.856 | 25.57 (torch.compile) / 14.02 / 1.824 |
| 6144 | 1024 | yes | 10.91 (flashinfer) / 9.18 / 1.188 | 32.48
(F.rms_norm) / 15.81 / 2.055 | 49.12 (F.rms_norm) / 24.99 / 1.965 |
| 6144 | 4096 | yes | 32.29 (flashinfer) / 32.16 / 1.004 | 93.31
(torch.compile) / 41.47 / 2.250 | 130.05 (torch.compile) / 73.63 / 1.766
|
| 6144 | 16172 | yes | 115.81 (flashinfer) / 114.43 / 1.012 | 204.51
(torch.compile) / 138.98 / 1.472 | 329.89 (torch.compile) / 253.41 /
1.302 |
| 6144 | 16231 | yes | 115.68 (flashinfer) / 115.23 / 1.004 | 205.09
(torch.compile) / 138.66 / 1.479 | 330.53 (torch.compile) / 253.89 /
1.302 |
| 6144 | 32768 | yes | 228.48 (flashinfer) / 227.26 / 1.005 | 336.86
(torch.compile) / 271.23 / 1.242 | 568.90 (torch.compile) / 498.50 /
1.141 |
| 6144 | 65536 | yes | 453.76 (flashinfer) / 452.38 / 1.003 | 621.89
(torch.compile) / 531.52 / 1.170 | 1077.18 (torch.compile) / 983.90 /
1.095 |
| 2048 | 256 | no | 2.50 (Liger) / 2.50 / 1.000 **(not faster)** | 10.82
(F.rms_norm) / 6.34 / 1.707 | 14.88 (F.rms_norm) / 8.83 / 1.685 |
| 2048 | 1024 | no | 3.58 (Liger) / 3.52 / 1.018 | 14.46 (F.rms_norm) /
9.09 / 1.592 | 19.26 (F.rms_norm) / 12.61 / 1.528 |
| 2048 | 4096 | no | 7.90 (flashinfer) / 7.58 / 1.042 | 25.98
(torch.compile) / 15.87 / 1.637 | 35.07 (torch.compile) / 23.46 / 1.495
|
| 2048 | 16172 | no | 22.78 (flashinfer) / 22.08 / 1.032 | 80.74
(torch.compile) / 41.86 / 1.929 | 104.80 (torch.compile) / 63.94 / 1.639
|
| 2048 | 16231 | no | 22.85 (flashinfer) / 22.11 / 1.033 | 81.60
(torch.compile) / 41.79 / 1.953 | 105.66 (torch.compile) / 63.90 / 1.653
|
| 2048 | 32768 | no | 42.46 (flashinfer) / 41.79 / 1.016 | 112.77
(torch.compile) / 77.22 / 1.460 | 156.70 (torch.compile) / 119.01 /
1.317 |
| 2048 | 65536 | no | 80.19 (flashinfer) / 78.72 / 1.019 | 211.94
(torch.compile) / 146.27 / 1.449 | 295.17 (torch.compile) / 224.99 /
1.312 |
| 512 | 256 | no | 2.11 (flashinfer) / 2.18 / 0.971 **(not faster)** |
8.19 (F.rms_norm) / 5.60 / 1.463 | 10.43 (torch.compile) / 7.78 / 1.342
|
| 512 | 1024 | no | 2.40 (flashinfer) / 2.43 / 0.987 **(not faster)** |
10.59 (F.rms_norm) / 6.40 / 1.655 | 13.86 (F.rms_norm) / 8.83 / 1.569 |
| 512 | 4096 | no | 3.36 (flashinfer) / 3.42 / 0.981 **(not faster)** |
14.88 (torch.compile) / 8.19 / 1.816 | 18.75 (torch.compile) / 11.62 /
1.614 |
| 512 | 16172 | no | 7.07 (flashinfer) / 7.07 / 1.000 **(not faster)** |
33.22 (torch.compile) / 14.37 / 2.312 | 40.80 (torch.compile) / 21.44 /
1.903 |
| 512 | 16231 | no | 7.10 (flashinfer) / 7.10 / 1.000 **(not faster)** |
33.38 (torch.compile) / 14.40 / 2.318 | 40.80 (torch.compile) / 21.50 /
1.897 |
| 512 | 32768 | no | 11.78 (flashinfer) / 12.29 / 0.958 **(not faster)**
| 28.00 (torch.compile) / 22.69 / 1.234 | 40.51 (torch.compile) / 34.98
/ 1.158 |
| 512 | 65536 | no | 21.31 (flashinfer) / 21.79 / 0.978 **(not faster)**
| 48.83 (torch.compile) / 38.69 / 1.262 | 71.52 (torch.compile) / 60.48
/ 1.183 |

### B200 (SM100a, 148 SMs) - export protocol receipts (56 rows) and 84
paired acceptance rows on a B200 node

Modules JIT-compiled on the node (nvcc 13.3, `compute_100a`); module
test suite 133 passed / 28 skipped on the B200. The same frozen export
protocol as on GB300, measured on the B200 for the SM100a rows: 56/56
rows correct, source/export bitwise-identical on 56/56 rows.

| row | source us | export us | source/export | torch.compile/export |
eager FP32/export | F.rms_norm/export | Liger/export |
flashinfer.rmsnorm/export |
|---|---:|---:|---:|---:|---:|---:|---:|---:|
| fwd_h6144_t256__sm_100a | 3.65 | 3.58 | 1.0179 | 1.080 | 11.758 |
2.429 | 1.062 | 1.179 |
| bwd_h6144_t256__sm_100a | 8.67 | 8.77 | 0.9891 | 1.624 | 8.106 | 2.157
| 2.241 | n/a |
| fwd_h6144_t1024__sm_100a | 6.40 | 6.37 | 1.0050 | 1.236 | 12.467 |
1.638 | 1.412 | 1.206 |
| bwd_h6144_t1024__sm_100a | 14.40 | 14.42 | 0.9989 | 2.730 | 10.744 |
1.869 | 2.111 | n/a |
| fwd_h6144_t4096__sm_100a | 18.98 | 18.98 | 1.0001 | 1.253 | 15.688 |
1.336 | 1.401 | 1.091 |
| bwd_h6144_t4096__sm_100a | 31.94 | 31.94 | 1.0000 | 2.646 | 17.811 |
2.701 | 2.391 | n/a |
| fwd_h6144_t16172__sm_100a | 61.76 | 61.70 | 1.0010 | 1.001 | 17.638 |
1.183 | 1.493 | 1.029 |
| bwd_h6144_t16172__sm_100a | 102.77 | 102.85 | 0.9992 | 1.871 | 19.951
| 3.135 | 2.424 | n/a |
| fwd_h6144_t16231__sm_100a | 61.89 | 61.86 | 1.0005 | 0.998 | 17.650 |
1.190 | 1.492 | 1.019 |
| bwd_h6144_t16231__sm_100a | 103.26 | 103.30 | 0.9997 | 1.863 | 19.941
| 3.141 | 2.423 | n/a |
| fwd_h6144_t32768__sm_100a | 119.18 | 119.20 | 0.9999 | 1.089 | 18.220
| 1.136 | 1.527 | 1.013 |
| bwd_h6144_t32768__sm_100a | 199.57 | 199.57 | 1.0000 | 1.674 | 20.501
| 2.655 | 2.414 | n/a |
| fwd_h6144_t65536__sm_100a | 231.87 | 231.94 | 0.9997 | 18.524 | 18.523
| 1.112 | 1.541 | 1.016 |
| bwd_h6144_t65536__sm_100a | 390.32 | 390.19 | 1.0003 | 20.776 | 20.774
| 2.632 | 2.400 | n/a |
| fwd_h6144_t256_res__sm_100a | 4.64 | 4.64 | 1.0000 | 1.028 | 13.027 |
6.041 | 5.103 | n/a |
| bwd_h6144_t256_res__sm_100a | 9.79 | 9.76 | 1.0032 | 1.905 | 8.675 |
3.193 | 3.357 | n/a |
| fwd_h6144_t1024_res__sm_100a | 9.76 | 9.76 | 1.0000 | 1.151 | 12.908 |
6.266 | 6.052 | n/a |
| bwd_h6144_t1024_res__sm_100a | 16.22 | 16.26 | 0.9980 | 2.805 | 11.474
| 3.577 | 3.805 | n/a |
| fwd_h6144_t4096_res__sm_100a | 33.44 | 33.44 | 1.0000 | 1.001 | 14.980
| 6.776 | 6.896 | n/a |
| bwd_h6144_t4096_res__sm_100a | 41.82 | 41.76 | 1.0015 | 2.394 | 16.838
| 5.229 | 5.062 | n/a |
| fwd_h6144_t16172_res__sm_100a | 117.82 | 117.81 | 1.0001 | 1.011 |
15.883 | 7.280 | 7.432 | n/a |
| bwd_h6144_t16172_res__sm_100a | 142.67 | 142.82 | 0.9990 | 1.682 |
18.034 | 5.879 | 5.411 | n/a |
| fwd_h6144_t16231_res__sm_100a | 118.21 | 118.14 | 1.0005 | 1.013 |
15.897 | 7.284 | 7.442 | n/a |
| bwd_h6144_t16231_res__sm_100a | 142.88 | 142.85 | 1.0002 | 1.675 |
18.079 | 5.904 | 5.415 | n/a |
| fwd_h6144_t32768_res__sm_100a | 234.24 | 234.02 | 1.0010 | 1.008 |
16.009 | 7.310 | 7.503 | n/a |
| bwd_h6144_t32768_res__sm_100a | 278.91 | 279.12 | 0.9993 | 1.491 |
18.393 | 5.624 | 5.473 | n/a |
| fwd_h6144_t65536_res__sm_100a | 464.99 | 465.45 | 0.9990 | 1.011 |
15.985 | 7.300 | 7.512 | n/a |
| bwd_h6144_t65536_res__sm_100a | 548.24 | 548.13 | 1.0002 | 1.405 |
18.590 | 5.670 | 5.508 | n/a |
| fwd_h2048_t256__sm_100a | 2.50 | 2.53 | 0.9877 | 1.278 | 9.608 | 1.671
| 1.013 | 1.203 |
| bwd_h2048_t256__sm_100a | 6.94 | 6.91 | 1.0048 | 1.820 | 6.806 | 1.533
| 2.445 | n/a |
| fwd_h2048_t1024__sm_100a | 3.46 | 3.49 | 0.9908 | 10.633 | 10.642 |
1.422 | 1.064 | 1.174 |
| bwd_h2048_t1024__sm_100a | 9.76 | 9.76 | 1.0000 | 7.184 | 7.184 |
1.508 | 2.325 | n/a |
| fwd_h2048_t4096__sm_100a | 8.13 | 8.06 | 1.0079 | 10.880 | 10.896 |
1.365 | 1.218 | 1.020 |
| bwd_h2048_t4096__sm_100a | 16.61 | 16.55 | 1.0038 | 11.347 | 11.318 |
2.104 | 2.843 | n/a |
| fwd_h2048_t16172__sm_100a | 22.94 | 22.91 | 1.0014 | 16.647 | 16.643 |
1.404 | 1.324 | 1.046 |
| bwd_h2048_t16172__sm_100a | 43.78 | 43.81 | 0.9993 | 16.730 | 16.712 |
3.439 | 3.261 | n/a |
| fwd_h2048_t16231__sm_100a | 22.66 | 22.66 | 1.0000 | 16.836 | 16.842 |
1.424 | 1.322 | 1.041 |
| bwd_h2048_t16231__sm_100a | 43.81 | 43.90 | 0.9978 | 16.735 | 16.744 |
3.429 | 3.266 | n/a |
| fwd_h2048_t32768__sm_100a | 42.56 | 42.53 | 1.0008 | 17.620 | 17.603 |
1.421 | 1.360 | 1.023 |
| bwd_h2048_t32768__sm_100a | 80.45 | 80.51 | 0.9992 | 17.679 | 17.674 |
3.295 | 3.442 | n/a |
| fwd_h2048_t65536__sm_100a | 80.11 | 79.98 | 1.0016 | 18.317 | 18.318 |
1.451 | 1.390 | 1.024 |
| bwd_h2048_t65536__sm_100a | 152.48 | 152.42 | 1.0004 | 18.172 | 18.167
| 3.349 | 3.484 | n/a |
| fwd_h512_t256__sm_100a | 2.27 | 2.27 | 1.0000 | 0.986 | 8.465 | 1.380
| 0.972 | 0.958 |
| bwd_h512_t256__sm_100a | 6.24 | 6.08 | 1.0263 | 1.384 | 6.347 | 1.337
| 2.634! | n/a |
| fwd_h512_t1024__sm_100a | 2.53 | 2.53 | 1.0000 | 8.468 | 8.430 | 1.392
| 1.013 | 1.000 |
| bwd_h512_t1024__sm_100a | 7.36 | 7.26 | 1.0132 | 6.224 | 6.247 | 1.485
| 2.797 | n/a |
| fwd_h512_t4096__sm_100a | 3.65 | 3.58 | 1.0176 | 9.604 | 9.622 | 1.919
| 1.330 | 0.991 |
| bwd_h512_t4096__sm_100a | 9.22 | 9.15 | 1.0070 | 7.721 | 7.693 | 2.472
| 4.074 | n/a |
| fwd_h512_t16172__sm_100a | 7.36 | 7.33 | 1.0045 | 11.954 | 11.954 |
2.917 | 1.655 | 1.000 |
| bwd_h512_t16172__sm_100a | 15.30 | 15.30 | 1.0000 | 12.966 | 12.954 |
5.282 | 6.633 | n/a |
| fwd_h512_t16231__sm_100a | 7.46 | 7.42 | 1.0043 | 11.866 | 11.815 |
2.888 | 1.642 | 0.974 |
| bwd_h512_t16231__sm_100a | 15.30 | 15.42 | 0.9917 | 12.913 | 12.894 |
5.307 | 6.627 | n/a |
| fwd_h512_t32768__sm_100a | 12.67 | 12.70 | 0.9975 | 15.139 | 15.106 |
3.196 | 1.748 | 0.965 |
| bwd_h512_t32768__sm_100a | 23.68 | 23.71 | 0.9987 | 16.965 | 16.972 |
7.557 | 8.283 | n/a |
| fwd_h512_t65536__sm_100a | 23.39 | 23.46 | 0.9973 | 16.608 | 16.607 |
3.352 | 1.784 | 0.958 |
| bwd_h512_t65536__sm_100a | 40.61 | 40.53 | 1.0020 | 19.085 | 19.084 |
9.080 | 9.319 | n/a |

Rows below 1.0 against a comparator (same H = 512 structural gap as on
GB300):

* `fwd_h6144_t16231__sm_100a`: torch.compile 0.998
* `fwd_h512_t256__sm_100a`: torch.compile 0.986, Liger 0.972,
flashinfer.rmsnorm 0.958
* `fwd_h512_t4096__sm_100a`: flashinfer.rmsnorm 0.991
* `fwd_h512_t16231__sm_100a`: flashinfer.rmsnorm 0.974
* `fwd_h512_t32768__sm_100a`: flashinfer.rmsnorm 0.965
* `fwd_h512_t65536__sm_100a`: flashinfer.rmsnorm 0.958

B200 (SM100a, 148 SMs, SM clock 1965 MHz, one node, 84 acceptance rows,
same comparators and thresholds re-measured on the B200) — **76/84 rows
strictly faster** (ratio = fastest baseline / this PR; us = CUPTI
kernel-sum median, cold L2, one process, paired; re-run on the
per-architecture tables of this revision).

| H | T | residual | fwd: fastest baseline (arm) / this PR / ratio | bwd
(+dw reduce): fastest baseline (arm) / this PR / ratio | fwd+bwd:
fastest baseline / this PR / ratio |
|---|---:|---|---|---|---|
| 6144 | 256 | no | 3.78 (Liger) / 3.52 / 1.073 | 14.14 (torch.compile)
/ 8.29 / 1.707 | 17.95 (torch.compile) / 11.81 / 1.520 |
| 6144 | 1024 | no | 7.58 (flashinfer) / 6.27 / 1.209 | 26.58
(F.rms_norm) / 14.02 / 1.896 | 36.94 (F.rms_norm) / 20.29 / 1.821 |
| 6144 | 4096 | no | 21.02 (flashinfer) / 18.97 / 1.108 | 75.97 (Liger)
/ 31.49 / 2.413 | 102.56 (Liger) / 50.46 / 2.032 |
| 6144 | 16172 | no | 61.70 (torch.compile) / 61.44 / 1.004 | 191.49
(torch.compile) / 102.56 / 1.867 | 253.18 (torch.compile) / 164.00 /
1.544 |
| 6144 | 16231 | no | 62.02 (torch.compile) / 61.63 / 1.006 | 191.84
(torch.compile) / 102.62 / 1.869 | 253.86 (torch.compile) / 164.26 /
1.545 |
| 6144 | 32768 | no | 121.34 (flashinfer) / 118.53 / 1.024 | 330.40
(torch.compile) / 198.69 / 1.663 | 459.42 (torch.compile) / 317.22 /
1.448 |
| 6144 | 65536 | no | 235.07 (flashinfer) / 231.78 / 1.014 | 604.87
(torch.compile) / 388.61 / 1.556 | 852.55 (torch.compile) / 620.38 /
1.374 |
| 6144 | 256 | yes | 4.58 (torch.compile) / 4.61 / 0.993 **(not
faster)** | 18.21 (torch.compile) / 9.57 / 1.903 | 22.78 (torch.compile)
/ 14.18 / 1.607 |
| 6144 | 1024 | yes | 11.20 (torch.compile) / 9.60 / 1.167 | 33.22
(F.rms_norm) / 16.22 / 2.047 | 50.53 (F.rms_norm) / 25.82 / 1.957 |
| 6144 | 4096 | yes | 33.44 (torch.compile) / 32.74 / 1.022 | 95.30
(torch.compile) / 42.08 / 2.265 | 128.74 (torch.compile) / 74.82 / 1.721
|
| 6144 | 16172 | yes | 118.59 (flashinfer) / 117.79 / 1.007 | 215.17
(torch.compile) / 142.24 / 1.513 | 335.01 (torch.compile) / 260.03 /
1.288 |
| 6144 | 16231 | yes | 118.69 (flashinfer) / 118.50 / 1.002 | 215.58
(torch.compile) / 142.11 / 1.517 | 334.58 (torch.compile) / 260.61 /
1.284 |
| 6144 | 32768 | yes | 234.56 (flashinfer) / 234.05 / 1.002 | 354.34
(torch.compile) / 279.01 / 1.270 | 590.34 (torch.compile) / 513.06 /
1.151 |
| 6144 | 65536 | yes | 464.90 (flashinfer) / 463.52 / 1.003 | 651.55
(torch.compile) / 546.43 / 1.192 | 1121.66 (torch.compile) / 1009.95 /
1.111 |
| 2048 | 256 | no | 2.53 (Liger) / 2.46 / 1.026 | 10.98 (F.rms_norm) /
6.50 / 1.690 | 15.04 (F.rms_norm) / 8.96 / 1.679 |
| 2048 | 1024 | no | 3.71 (Liger) / 3.42 / 1.084 | 15.04 (F.rms_norm) /
9.38 / 1.604 | 20.00 (F.rms_norm) / 12.80 / 1.562 |
| 2048 | 4096 | no | 8.29 (flashinfer) / 8.03 / 1.032 | 24.74
(torch.compile) / 16.38 / 1.510 | 34.40 (torch.compile) / 24.42 / 1.409
|
| 2048 | 16172 | no | 23.55 (flashinfer) / 22.75 / 1.035 | 71.81
(torch.compile) / 43.71 / 1.643 | 97.41 (torch.compile) / 66.46 / 1.466
|
| 2048 | 16231 | no | 23.62 (flashinfer) / 22.82 / 1.035 | 75.65
(torch.compile) / 43.68 / 1.732 | 101.47 (torch.compile) / 66.50 / 1.526
|
| 2048 | 32768 | no | 43.49 (flashinfer) / 42.05 / 1.034 | 96.22
(torch.compile) / 80.74 / 1.192 | 143.30 (torch.compile) / 122.78 /
1.167 |
| 2048 | 65536 | no | 81.76 (flashinfer) / 80.06 / 1.021 | 171.90
(torch.compile) / 152.32 / 1.129 | 261.25 (torch.compile) / 232.39 /
1.124 |
| 512 | 256 | no | 2.18 (flashinfer) / 2.24 / 0.971 **(not faster)** |
8.38 (torch.compile) / 5.70 / 1.472 | 10.56 (torch.compile) / 7.94 /
1.331 |
| 512 | 1024 | no | 2.43 (flashinfer) / 2.50 / 0.974 **(not faster)** |
11.30 (F.rms_norm) / 6.85 / 1.650 | 14.62 (F.rms_norm) / 9.34 / 1.565 |
| 512 | 4096 | no | 3.49 (flashinfer) / 3.58 / 0.973 **(not faster)** |
15.26 (torch.compile) / 8.74 / 1.747 | 19.26 (torch.compile) / 12.32 /
1.564 |
| 512 | 16172 | no | 7.20 (flashinfer) / 7.33 / 0.983 **(not faster)** |
35.30 (torch.compile) / 15.04 / 2.347 | 43.39 (torch.compile) / 22.37 /
1.940 |
| 512 | 16231 | no | 7.23 (flashinfer) / 7.33 / 0.987 **(not faster)** |
35.49 (torch.compile) / 15.07 / 2.355 | 43.65 (torch.compile) / 22.40 /
1.948 |
| 512 | 32768 | no | 12.54 (flashinfer) / 12.90 / 0.973 **(not faster)**
| 28.77 (torch.compile) / 23.58 / 1.220 | 41.63 (torch.compile) / 36.48
/ 1.141 |
| 512 | 65536 | no | 22.14 (flashinfer) / 23.36 / 0.948 **(not faster)**
| 50.05 (torch.compile) / 39.97 / 1.252 | 73.38 (torch.compile) / 63.33
/ 1.159 |

### R200 (SM107a, 212 SMs) - 84 paired acceptance rows on a Rubin node

Rubin R200 node (212 SMs, CC 10.7, 4x R200; CUDA 13.5 container with
`compute_107` nvcc/NVRTC: torch 2.13.0a0, Triton 3.7.0, liger-kernel,
flashinfer 0.7.1 source checkout = this PR's base; TransformerEngine
2.18 importable but not timeable under the CUPTI harness), harness v2,
same methodology as above. Two passes:

- **Shipped programs (this PR's `sm_107a` sources and host table,
kernels unchanged): 65/84 rows strictly faster** - every backward and
step row wins; 19 forward rows lose to flashinfer `rmsnorm` /
`fused_add_rmsnorm` (H=6144 res=0 T>=4096 0.871-0.944, H=6144 res=1
T=256 and T>=4096 0.931-0.983, H=2048 T=4096/16172/16231 0.847-0.985,
H=512 T=256/1024/4096/32768/65536 0.961-0.998). On R200 the torch copy
ceiling (6.4-6.6 TB/s at H=6144) is below both kernels, so these are not
roofline ties: the sm_107a forward host table had been copied from the
GB300 tuning.
- **Re-tuned forward table (per-arch program selection only: L2 load
hints, f32 vs packed register cache, rows per CTA, T dispatch; y / r
bit-identical to the shipped programs, verified per variant): 76/84 rows
strictly faster** - bwd 1.26-2.98x, fwd+bwd 1.19-2.41x on every row;
H=6144 res=1 forward wins at every T (1.02-1.10x), H=6144 res=0 at T <=
4096, H=2048 at T <= 4096 and T >= 32768, H=512 at
T=1024/4096/16172/16231/65536. The `sm_107a` sources and loader table in
this PR are regenerated from the re-tuned selection (this push): y / r
are bit-identical to the first pass' programs for every row, SM100 /
SM103 sources are byte-identical to the previous push, and the
regenerated SM107 modules were JIT-built (nvcc 13.5, compute_107) and
tested on an R200 node (`tests/norm/test_cake_rmsnorm_train.py`: 133
passed, 28 skipped). A kernel that the dispatch launches under several
launch policies is recorded once per policy in the loader (records
sharing the kernel's sources and JIT cache name). The table below is the
re-tuned pass.

| H | T | residual | fwd: fastest baseline (arm) / this PR / ratio | bwd
(+dw reduce): fastest baseline (arm) / this PR / ratio | fwd+bwd:
fastest baseline / this PR / ratio |
|---|---:|---|---|---|---|
| 6144 | 256 | no | 3.30 (Liger) / 3.26 / 1.010 | 12.54 (torch.compile)
/ 7.01 / 1.790 | 16.00 (torch.compile) / 10.27 / 1.558 |
| 6144 | 1024 | no | 5.47 (flashinfer) / 4.86 / 1.125 | 20.42
(F.rms_norm) / 11.30 / 1.807 | 29.83 (F.rms_norm) / 16.16 / 1.846 |
| 6144 | 4096 | no | 14.14 (flashinfer) / 12.90 / 1.097 | 59.30 (Liger)
/ 24.64 / 2.406 | 80.35 (torch.compile) / 37.54 / 2.141 |
| 6144 | 16172 | no | 41.66 (flashinfer) / 44.83 / 0.929 **(not
faster)** | 183.01 (Liger) / 76.48 / 2.393 | 258.31 (Liger) / 121.31 /
2.129 |
| 6144 | 16231 | no | 41.73 (flashinfer) / 45.22 / 0.923 **(not
faster)** | 183.62 (Liger) / 77.20 / 2.378 | 258.75 (torch.compile) /
122.42 / 2.114 |
| 6144 | 32768 | no | 80.03 (flashinfer) / 88.96 / 0.900 **(not
faster)** | 349.28 (Liger) / 145.12 / 2.407 | 445.09 (torch.compile) /
234.08 / 1.901 |
| 6144 | 65536 | no | 156.26 (flashinfer) / 173.09 / 0.903 **(not
faster)** | 679.43 (Liger) / 282.15 / 2.408 | 861.70 (torch.compile) /
455.24 / 1.893 |
| 6144 | 256 | yes | 3.87 (flashinfer) / 3.78 / 1.025 | 15.07
(torch.compile) / 7.84 / 1.922 | 19.10 (torch.compile) / 11.62 / 1.645 |
| 6144 | 1024 | yes | 7.62 (flashinfer) / 6.91 / 1.102 | 26.21
(F.rms_norm) / 12.13 / 2.161 | 41.22 (F.rms_norm) / 19.04 / 2.165 |
| 6144 | 4096 | yes | 23.07 (flashinfer) / 22.72 / 1.016 | 74.98
(torch.compile) / 30.30 / 2.474 | 102.11 (torch.compile) / 53.02 / 1.926
|
| 6144 | 16172 | yes | 81.15 (flashinfer) / 78.53 / 1.033 | 205.83
(torch.compile) / 100.48 / 2.048 | 301.99 (torch.compile) / 179.01 /
1.687 |
| 6144 | 16231 | yes | 81.41 (flashinfer) / 78.75 / 1.034 | 207.33
(torch.compile) / 100.03 / 2.073 | 303.36 (torch.compile) / 178.79 /
1.697 |
| 6144 | 32768 | yes | 160.16 (flashinfer) / 154.88 / 1.034 | 357.79
(torch.compile) / 196.03 / 1.825 | 549.51 (torch.compile) / 350.92 /
1.566 |
| 6144 | 65536 | yes | 317.22 (flashinfer) / 307.23 / 1.032 | 696.20
(torch.compile) / 378.15 / 1.841 | 1074.89 (torch.compile) / 685.38 /
1.568 |
| 2048 | 256 | no | 2.50 (Liger) / 2.43 / 1.026 | 9.66 (F.rms_norm) /
5.73 / 1.687 | 13.73 (F.rms_norm) / 8.16 / 1.682 |
| 2048 | 1024 | no | 3.33 (flashinfer) / 2.95 / 1.130 | 12.74
(F.rms_norm) / 7.91 / 1.611 | 17.25 (F.rms_norm) / 10.85 / 1.590 |
| 2048 | 4096 | no | 5.70 (flashinfer) / 5.60 / 1.017 | 29.79
(F.rms_norm) / 12.67 / 2.351 | 39.01 (torch.compile) / 18.27 / 2.135 |
| 2048 | 16172 | no | 16.26 (flashinfer) / 16.51 / 0.984 **(not
faster)** | 92.87 (torch.compile) / 32.06 / 2.896 | 115.20
(torch.compile) / 48.58 / 2.372 |
| 2048 | 16231 | no | 16.26 (flashinfer) / 16.54 / 0.983 **(not
faster)** | 92.74 (torch.compile) / 31.17 / 2.975 | 114.94
(torch.compile) / 47.71 / 2.409 |
| 2048 | 32768 | no | 29.31 (flashinfer) / 29.15 / 1.005 | 141.03
(torch.compile) / 56.64 / 2.490 | 181.70 (torch.compile) / 85.79 / 2.118
|
| 2048 | 65536 | no | 54.56 (flashinfer) / 53.79 / 1.014 | 270.53
(torch.compile) / 105.76 / 2.558 | 346.91 (torch.compile) / 159.55 /
2.174 |
| 512 | 256 | no | 2.08 (flashinfer) / 2.18 / 0.956 **(not faster)** |
7.65 (F.rms_norm) / 4.90 / 1.562 | 10.08 (torch.compile) / 7.07 / 1.425
|
| 512 | 1024 | no | 2.34 (flashinfer) / 2.30 / 1.014 | 10.05
(F.rms_norm) / 6.18 / 1.627 | 13.09 (F.rms_norm) / 8.48 / 1.543 |
| 512 | 4096 | no | 2.88 (flashinfer) / 2.82 / 1.023 | 12.32
(torch.compile) / 7.30 / 1.689 | 15.49 (torch.compile) / 10.11 / 1.532 |
| 512 | 16172 | no | 6.15 (flashinfer) / 6.14 / 1.000 | 30.11
(torch.compile) / 11.71 / 2.571 | 37.47 (torch.compile) / 17.86 / 2.099
|
| 512 | 16231 | no | 6.24 (flashinfer) / 6.21 / 1.005 | 30.40
(torch.compile) / 11.65 / 2.610 | 37.66 (torch.compile) / 17.86 / 2.109
|
| 512 | 32768 | no | 10.53 (flashinfer) / 10.69 / 0.985 **(not faster)**
| 23.42 (torch.compile) / 17.95 / 1.305 | 35.14 (torch.compile) / 28.64
/ 1.227 |
| 512 | 65536 | no | 18.69 (flashinfer) / 18.37 / 1.017 | 37.70
(torch.compile) / 29.82 / 1.264 | 57.41 (torch.compile) / 48.19 / 1.191
|

Rows not strictly faster after the re-tune (all forward, vs flashinfer
`rmsnorm`): H=6144 res=0 T=16172 / 16231 / 32768 / 65536 at 0.929 /
0.923 / 0.900 / 0.903 (this PR 8.8-9.3 TB/s vs 9.6-10.3 TB/s). The
comparator here is FlashInfer's CuTe-DSL `RMSNormKernel`
(`flashinfer/norm/kernels/rmsnorm.py`, the default `rmsnorm` dispatch on
SM100/SM103/SM107; at H=6144 it runs 64 threads per row, two rows per
128-thread CTA, twelve 16-byte loads per thread). A plain-CUDA port of
that structure reproduces its bandwidth on R200 (41.8-42.0 us at T=16231
vs 41.1 us; 152-158 us at T=65536 vs 155.9 us) with or without
shared-memory staging, while the same port with this PR's
256-threads-per-row group runs at 47.2-48.7 us: the remaining gap is the
per-row thread group (fewer, longer per-thread load streams), not the
exact rsqrt chain (an `rsqrt.approx` probe recovers only 0.5-1.6%). A
64- or 128-thread row group changes the FP32 reduction order, so it is
not adopted in this PR, H=2048 T=16172 / 16231 at 0.984 / 0.983, H=512
T=256 at 0.956 and T=32768 at 0.985.

Device gates on R200: `tests/norm/test_cake_rmsnorm_train.py` 133 passed
/ 28 skipped with the sm_107a modules JIT-built by nvcc 13.5
(compute_107); compute-sanitizer memcheck + synccheck 0 errors on every
forward/backward benchmark row and on 26 direct runs of the re-tuned
program selections; correctness matrix 785/785 on both passes; dw
bitwise-identical across runs.

### R200 (sm_107a) round 3: H=512 forward small program without the row
loop

The SM107 H=512 non-residual forward program for T <= 4096 (256-thread
CTAs, eight rows per CTA) now runs its rows without the grid-stride loop
(every CTA owns exactly eight rows in that range); `y` and `rstd` stay
bit-identical (same 32-thread reduction tree). One generated
kernel/binding pair is renamed for its new digest and the module table
follows; the other 90 generated files are unchanged.

Paired CUPTI timings on one R200 node (cold L2, FlashInfer 0.7.1
`rmsnorm` as the threshold): T=256 2.21 -> 2.14/2.18 us (threshold
2.14), T=1024 2.40 -> 2.34 us (2.40), T=4096 2.91 -> 2.88/2.91 us
(2.91-2.94). Re-measuring the full 84-row table on that node gives 72/84
rows strictly faster for both the previous and the new sources (the
table above was measured on a different R200 node at 76/84; the marginal
rows H=512 T >= 16172 and H=2048 T=4096 move within +/-1.5 % between
nodes and runs). Rows not strictly faster on this node: H=6144 fwd T >=
16172 (0.88-0.92), H=2048 fwd T=4096 (tie), T=16172/16231 (0.977-0.983),
H=512 fwd T=256 (0.985-1.000), T=16172/16231 (0.961-0.985), T=32768
(0.957), T=65536 (0.975-0.978). All 56 backward and fwd+bwd rows are
faster (1.17-2.9x).

Export protocol run on the R200 device for all 56 sm_107a shapes (the
first device-gated run of these sources): 56/56 receipts, correctness
56/56, source/export bit-identical 56/56, 43/56 all gates (rows below
1.0 against FlashInfer `rmsnorm` in the public-call regime: H=6144 fwd T
>= 16172 0.88-0.93, H=2048 fwd T=16172/16231 0.96-0.98, H=512 fwd
T=256/4096/16172/32768/65536 0.96-0.985; bwd H=512 T <= 1024 fail only
the launch-order-dependence gate).
`tests/norm/test_cake_rmsnorm_train.py`: 133 passed, 28 skipped on R200
with the regenerated sources.

### Round 4 (all three platforms): ordered per-arch forward program
tables, H=512 `quad` program

The forward dispatch table is now an ordered per-arch list of (program,
T bound) entries instead of three fixed slots. One program was
added on every architecture — `quad`: 128-thread CTAs owning 4 rows (one
warp per row, one trip, no grid-stride loop), the same
reduction tree as the programs it displaces, so `y` / `rstd` stay
bit-identical — for 1024 < T <= 4096 on SM100a / SM103a and for
T <= 256 on SM107a. All 42 previously shipped programs render
byte-identically (whole-table render check on the Rubin node: 45 vs 42
programs, 3 added lines, 0 changed); the generated sources gain exactly
one unit per architecture (JIT role `fwd_quad`), nothing is
retired. A second candidate (an evict-first looped SM103a residual
program for 1024 < T <= 4096) was measured in the same round and
not adopted (no gain over the existing main program: 32.29 us in both
builds vs 32.16 us).

Validation per platform (same node, the previous module run as a control
on another GPU at the same time, two independent builds
of the new module; 84 paired acceptance rows, CUPTI cold-L2 kernel-sum,
threshold = fastest comparator per row):

| platform | gates | A/B control (previous module) | A/B new, build 1 |
A/B new, build 2 | compute-sanitizer | export / module tests |
|---|---|---|---|---|---|---|
| GB300 (SM103a, 152 SMs) | facade matrix 785/785, GPU e2e 16/16,
bench-regression 25 names 0 crash, forward static analysis clean | 77/84
| 79/84 | 77/84 (incl. the dropped looped row) | memcheck + synccheck 0
errors on 6 shapes in the changed windows + the registered `h512_t4096`
launcher | 56 export rows, correctness 56/56, exporter perf gates 50/56
(the six H=512 fwd rows below); `tests/norm/test_cake_rmsnorm_train.py`
133 passed / 28 skipped |
| B200 (SM100a, 148 SMs) | same gate set, all green | 76/84 | 76/84 |
76/84 | memcheck + synccheck 0 errors on the same 6 shapes (the
registered-launcher tool run deadlocks in the x86 image's JIT step, as
in earlier rounds; the GB300 and R200 runs of that launcher are clean) |
56 export rows, correctness 56/56, exporter perf gates 49/56 (the six
H=512 fwd rows + H=6144 residual fwd T=4096 vs torch.compile under the
export protocol); generated sources byte-identical to the GB300 and R200
exports (99 files); module tests 133 passed / 28 skipped |
| R200 (SM107a, 212 SMs) | same gate set, all green | 73/84 | 74/84 |
73/84 | memcheck + synccheck 0 errors on T=1 / 200 / 256 / 257 / 4096 +
the registered `h512_t256` / `t4096` / `t16231` launchers | 56 export
rows, correctness 56/56, exporter perf gates 48/56 (the H=6144 / H=2048
forward rows above + two sub-10 us bwd rows failing the protocol's
order-sensitivity gate); generated sources byte-identical to the GB300
export (99 files); module tests 133 passed / 28 skipped |

Rows whose program changed (us, cake vs the threshold arm = FlashInfer
`rmsnorm`):

| row | platform | previous program (control) | `quad`, build 1 / build
2 |
|---|---|---|---|
| H=512 fwd T=4096 | GB300 | 3.42 vs 3.36 (0.981) | 3.26 vs 3.36
(**1.029**) / 3.30 vs 3.36 (**1.018**) |
| H=512 fwd T=4096 | B200 | 3.58 vs 3.49 (0.973) | 3.49 vs 3.49 (1.000)
/ 3.52 vs 3.49 (0.991) |
| H=512 fwd T=256 | R200 | 2.18 vs 2.14 (0.985) | 2.14 vs 2.14 (1.000) /
2.14 vs 2.14 (1.000) |

Final standing with this table: GB300 78-79/84, B200 76/84, R200
73-74/84. Every backward and training-step row wins on all three
platforms; the rows that are not strictly faster are forward-only
comparisons against inference kernels: H=512 fwd (launch-latency
regime at T <= 1024 — FlashInfer's `rsqrt.approx` vs this module's exact
IEEE `sqrt`/`rcp`, and its 128-thread smem-staged row group
at T >= 16k), H=6144 fwd T >= 16k on R200 (FlashInfer's 64-thread row
group), H=2048 fwd T~16k on R200 (1-2 %), and two sub-5 us
ties (H=6144 residual fwd T=256 on B200, H=6144 fwd T=256 on R200).
Closing them would change the FP32 reduction order or the
rounding of the reciprocal square root, which this module deliberately
does not do.

### Round 5 (SM107a only): pipelined persistent programs for the
throughput windows

Round 4 left the largest numerics-neutral gaps on the Rubin platform in
the throughput windows of the three non-residual widths, where
the shipped programs were the plain grid-stride loops shared with
Blackwell. Own-process paired sweeps on the same node (CUPTI cold L2,
FlashInfer `rmsnorm` measured in the same process, every candidate
checked bit-identical to the program it displaces) showed that the
software-pipelined persistent program — until now only a latency program
(4 CTAs/SM for T <= 4096) and a mid program at 6 / 16 CTAs/SM —
wins those windows once its CTA budget per SM is sized for them. The
SM107a table changes to:

| width | previous SM107a table | round-5 SM107a table |
|---|---|---|
| H=6144 | `small` <= 4096 pipeline 4/SM, `mid` <= 16384 pipeline 16/SM,
`main` f32 256x1 | `small` <= 1024 pipeline 4/SM; `mid` <= 4096 pipeline
4/SM plain x; `large` <= 12288 pipeline 16/SM plain x; `wide` <= 24576
pipeline 32/SM plain x; `main` pipeline 48/SM plain x |
| H=2048 | `small` <= 1024 128-thr; `mid` <= 4096 two rows/thread;
`main` plain 64-thr bf16x16 | `small` <= 1024; `mid` <= 8192; `large` <=
12288 pipelined 64-thr 32/SM; `wide` <= 24576 pipelined 64-thr 64/SM;
`main` plain 64-thr |
| H=512 | `quad` <= 256; `small` <= 4096 256x8 single trip; `mid` <=
16384 pipeline 6/SM; `main` 128x4 | `quad` <= 256; `small` <= 12288;
`mid` <= 16384 pipeline 6/SM; `main` pipeline 8/SM |

Every program of a width keeps that width's per-row thread group, so `y`
/ `rstd` are bit-identical across the table (the sweeps
verified `torch.equal` for every candidate). Sweep gains over the
previous program at the acceptance T (us, same process): H=6144
42.82 vs 46.46 at 16231, 82.47 vs 89.25 at 32768, 166.10 vs 174.72 at
65536, 12.86 vs 13.10 at 4096; H=2048 16.13 vs 16.66 at 16231;
H=512 10.27 vs 10.85 at 32768, 18.05 vs 18.53 at 65536. The two new slot
names are shared by the two widths that need five windows,
so the module gains exactly two JIT roles (`fwd_large`, `fwd_wide`); the
SM100a / SM103a tables and their 45 generated programs are
unchanged (whole-table render check on both a Blackwell and a Rubin
node: 45 vs 49 programs, the diff is confined to the SM107a entries;
45 vs 49 programs, 6 lines removed / 10 added, none outside SM107a,
identical result on the Blackwell and the Rubin node).

Validation on R200 (SM107a, 212 SMs; same node, the previous module run
as a control on another GPU at the same time, two independent
builds of the new module; 84 paired acceptance rows, CUPTI cold-L2
kernel-sum, threshold = fastest comparator per row):

| gates | A/B control (previous module) | A/B new, build 1 | A/B new,
build 2 | compute-sanitizer | export / module tests |
|---|---|---|---|---|---|
| facade matrix 785/785, GPU e2e 16/16, bench-regression 25 names 0
crash, forward static analysis clean (also on B200 for the unchanged
Blackwell programs) | 73/84 | 77/84 | 76/84 | memcheck + synccheck 0
errors on 15 shapes at and inside the new window edges (H=6144 T=4097 /
8192 / 8193 / 12288 / 16231 / 16385 / 32768, H=2048 T=4097 / 8192 /
16231 / 16385, H=512 T=16385 / 20480 / 32768 / 65536) and on the 7
registered launchers x 2 tools (44 runs, 44 x 0 errors) | device export
protocol 182aea87f680…, 56 rows: correctness 56/56, build checks passed,
exporter gates 49/56 (the open forward rows + one informational
launch-order flag), FI module pytest 133 passed / 28 skipped |

Rows whose program changed (us, cake vs the threshold arm):

| row (fwd) | previous program (control) | round-5 program, build 1 /
build 2 |
|---|---|---|
| H=6144 T=4096 | 13.15 vs 14.21 (1.080) | 12.99 / 12.93 vs 14.21 (1.091
/ 1.092) |
| H=6144 T=16172 | 45.44 vs 41.95 (0.923) | 42.72 / 42.75 vs 41.95
(0.983 / 0.981) |
| H=6144 T=16231 | 45.54 vs 42.18 (0.926) | 42.75 / 42.72 vs 42.18
(0.986 / 0.986) |
| H=6144 T=32768 | 89.86 vs 80.77 (0.899) | 82.94 / 82.91 vs 80.77
(0.975 / 0.974) |
| H=6144 T=65536 | 173.12 vs 157.82 (0.912) | 165.51 / 165.03 vs 157.82
(0.953 / 0.954) |
| H=2048 T=16172 | 16.61 vs 16.42 (0.988) | 16.29 / 16.26 vs 16.42
(**1.010 / 1.006**) |
| H=2048 T=16231 | 16.64 vs 16.42 (0.987) | 16.22 / 16.19 vs 16.42
(**1.014 / 1.010**) |
| H=512 T=32768 | 10.72 vs 10.66 (0.994) | 10.21 / 10.24 vs 10.66
(**1.034 / 1.031**) |
| H=512 T=65536 | 18.56 vs 18.91 (1.019) | 17.86 / 18.02 vs 18.91 (1.048
/ 1.037) |

The training-step rows of these shapes gain 0.4-3.2 % on top; the 58
rows whose program did not change stay within the 3.2 % same-node band
(none of them slower than the control beyond it).

Final standing on R200 with this table: 76-77/84 (control 73). Rows not
strictly faster, all forward-only comparisons against inference kernels:
H=6144 fwd T=16172 / 16231 / 32768 / 65536 (0.98 / 0.99 / 0.97 / 0.95,
up from 0.92 / 0.93 / 0.90 / 0.91 — FlashInfer's 64-thread row group
remains the open numerics-changing lever), H=512 fwd T=256 (2.14 vs
2.14, launch floor), H=512 fwd T=16172 / 16231 (0.99 / 1.00, within 1
%), and the H=6144 fwd T=256 tie (1.000 / 1.010). GB300 78-79/84 and
B200 76/84 stand unchanged from round 4 (programs byte-identical).

#### Blackwell check of the round-5 lever (no change)

The same pipelined-grid hypothesis was paired-tested on SM100a and
SM103a with the shipped programs as the base (own-process sweeps, every
arm bit-identical): for H=512 above 16k rows the shipped programs
(pipeline at 6 CTAs/SM to 16384, 64x2 above) remain the best arms or
within the 3 % same-node noise band (SM103a 7.10 / 7.23 / 12.29 / 22.56
us at T=16172 / 16231 / 32768 / 65536; SM100a 7.42 / 7.39 / 13.09 /
23.36 us), and for H=6144 the pipelined plain-x grid is 0.96-1.00 of the
shipped packed main on both. The SM100a / SM103a tables and generated
sources are unchanged. A follow-up check of 256-bit single loads
(vec=16) for every H=512 program shape on SM100a / SM103a found them not
bit-identical to the 128-bit form (the per-thread accumulation order
follows the load grouping) and not faster, so the 128-bit programs stay.

### Round 7 — SM107a table refinement (H=6144 largest T, H=2048 mid
window)

Paired own-process sweeps on a Rubin R200 node (every arm bit-identical
to the shipped program): the H=6144 software-pipelined persistent
program at 64 CTAs/SM beats the 48 CTAs/SM program from 32768 rows
upward (81.95 vs 82.85 us @32768, 100.63 vs 103.07 @40960, 161.09 vs
165.91 @65536) while 48/SM keeps 26624-28672 rows, so the SM107a table
gains an `xwide` window (48 CTAs/SM for 24576 < T <= 30720) and `main`
moves to 64 CTAs/SM — the same generated kernel, only the persistent
grid differs. The H=2048 mid program (64 threads, two rows per thread,
1024 < T <= 8192) takes evict-first `x` loads, which are never slower
(4.32 vs 4.45 us @2048, tie @4096, 8.83 vs 8.90 @6144, 9.82 vs 10.10
@8192). Wider CTAs (512x2, 1024x4), 80 / 96 CTAs/SM, launch bounds,
no-allocate `y` stores, plain `w` loads, single-trip / early-`w` /
four-rows-per-thread mid forms were measured and lose or stay within
noise. SM100a / SM103a tables and generated sources are unchanged.

Same-node paired validation on the R200 node (three simultaneous 84-row
runs, CUPTI cold L2, per-row threshold = fastest baseline arm): previous
module 74 / 84 rows strictly faster, new module 76 / 84 and 76 / 84 in
two independent builds. The six rows in the changed windows are faster
or tied in both builds and no row is slower than the previous module by
more than 0.5 %; the 78 unchanged rows move at most 1.6 % (same-node
run-to-run band).

| H | T | mode | previous us | new us (build 1 / build 2) | vs previous
| threshold us (arm) | x threshold (prev / new) |
|---|---|---|---|---|---|---|---|
| 6144 | 32768 | fwd | 82.72 | 82.24 / 82.18 | 1.006 / 1.007 | 80.58
(flashinfer rmsnorm) | 0.974 / 0.982, 0.979 |
| 6144 | 32768 | fwd+bwd | 225.89 | 225.60 / 225.38 | 1.001 / 1.002 |
441.26 (torch.compile) | 1.953 / 1.973, 1.988 |
| 6144 | 65536 | fwd | 166.40 | 161.80 / 161.57 | 1.028 / 1.030 | 156.90
(flashinfer rmsnorm) | 0.943 / 0.968, 0.972 |
| 6144 | 65536 | fwd+bwd | 445.26 | 442.15 / 441.10 | 1.007 / 1.009 |
853.00 (torch.compile) | 1.916 / 1.940, 1.964 |
| 2048 | 4096 | fwd | 5.66 | 5.66 / 5.60 | 1.000 / 1.011 | 5.73
(flashinfer rmsnorm) | 1.011 / 1.006, 1.017 |
| 2048 | 4096 | fwd+bwd | 18.59 | 18.43 / 18.46 | 1.009 / 1.007 | 39.04
(torch.compile) | 2.100 / 2.068, 2.087 |

Gates on the Rubin node for the new table: facade correctness 785 / 785,
GPU e2e 16 / 16, forward static analysis clean, compute-sanitizer
memcheck + synccheck 0 errors on the eight shapes of the changed windows
(H=6144 T=24577 / 32768 / 40960 / 65536, H=2048 T=2048 / 4096 / 4097 /
8192), export protocol correctness 56 / 56 with source-vs-export bitwise
56 / 56, and this module's pytest on the regenerated sources 133 passed
/ 28 skipped.

## Accuracy

Reference: FP64 recomputation of the same formulas. Tolerances: `y` /
`dx` BF16 `atol = rtol = 1e-2` after unit-scale normalization; `rstd`
rel-L2 <= 1e-5; `dw` rel-L2 <= 1e-4 per case, and the kernel's `dw`
error is not worse than the eager FP32 path `(g*x*rstd).sum(0)` as a
median ratio over the whole matrix (both sit at the FP32 rounding floor,
so a per-case comparison is rounding-order noise). Test matrix: `T in
{0, 1, 2, 7, 63, 64, 65, 127, 128, 129, 255, 257, 16172, 16231}`, row
strides `{H, H + 64, 2H}`, odd column offset and odd row stride (`H +
1`) on every input/output, scales `{1, 1e-3, 1e3}`, all-zero rows, a
single extreme row, residual on/off, gradient views (expanded / sliced),
end-to-end autograd vs FP32 eager autograd.

| output | max abs err (worst row) | rel-L2 (worst row) | bound |
|---|---|---|---|
| y | 1.56e-2 (fwd H=6144 T=65536) | 1.67e-3 (fwd H=512 T=1024) |
elementwise atol 1e-2 + rtol 1e-2 x abs(ref) |
| rstd | - | 5.1e-8 (fwd H=6144 T=256) | 1e-5 |
| dx | 2.78e-2 (bwd residual H=6144 T=16172) | 1.66e-3 (bwd H=512
T=4096) | elementwise atol 1e-2 + rtol 1e-2 x abs(ref) |
| dw | 1.20e-4 (bwd H=6144 T=65536; values O(1e2)) | 1.26e-7 (bwd H=6144
T=65536) | 1e-4; kernel/eager-FP32-path rel-L2 ratio median 0.60, max
0.82 over the 28 backward rows (<= 1) |

## Memory

| tensor | bytes | note |
|---|---|---|
| forward saved for backward | 2TH (x or h_new) + 2H (w) + 4T (rstd) |
no normalized copy is kept |
| backward workspace | 4 x n_chunks x H + counters (<= 4 x 2 x SM x H +
256 B) | caller-owned, reusable for every T; GB300 H = 6144: 7,471,872 B
(4 x 304 x 6144 partials + 192 x 4 B counters) |
| peak extra allocation in backward | 0 (with a supplied workspace) |
verified by the allocation probe in the export validation |
| peak intermediate allocation vs baselines (paired harness, requested
bytes minus declared outputs) | backward/fwd+bwd rows: below every
baseline on all 56 rows per platform; forward rows: exactly 4T bytes
(the saved `rstd`) | equal to what every training forward (torch.compile
/ eager / F.rms_norm / Liger) saves; above only the inference-only
`flashinfer.rmsnorm` / `fused_add_rmsnorm`, which save nothing; no [T,H]
FP32 temporaries in any path |

## Determinism

`dw` is bitwise identical across repeated calls (test: three runs with a
NaN-poisoned workspace between them), and the reduction order depends
only on `(T, device SM count)`; `dx` is elementwise and trivially
deterministic. Evidence: 56/56 export-protocol shapes bitwise-identical
between the generating framework's launcher and these modules; `dw`
bitwise-identical across 3 repeats on all 84 acceptance rows on GB300
and on B200.

## Baselines and their source PRs

Merge-base of this branch: `f558e9fd4e56807b6d12f147591efd44901d1d23`
(upstream `main`, 2026-09-30). All comparisons were measured against
this revision of the baseline files; `git log f558e9fd4..origin/main --
<files>` is empty at publication time (re-verified at publication).

* `flashinfer.norm.rmsnorm` (forward-only inference baseline, plain
forward rows): `flashinfer/norm/__init__.py` dispatching to the CuTe-DSL
kernel in `flashinfer/norm/kernels/rmsnorm.py` on every measured GPU
(the CUDA fallback `csrc/norm.cu` / `include/flashinfer/norm.cuh` is
only used with `FLASHINFER_USE_CUDA_NORM=1`); last upstream changes
#5305 (`012542c85`, 2026-09-18), #4480 (`6d5ef77f8`, 2026-08-26), #4202
(`519205914`, 2026-08-04).
* `torch.compile` (inductor, `dynamic=False`, per-shape) and eager FP32
autograd of the manual RMSNorm formula with an FP32 master weight:
PyTorch 2.13.0a0 (NGC 26.07 container, CUDA 13.3), Triton 3.6.0.
* `torch.nn.functional.rms_norm` autograd (BF16 weight gradient, upcast
for comparison): PyTorch 2.13.0a0.
* Liger `LigerRMSNormFunction` (Triton, `casting_mode="llama"`,
`in_place=False`): liger-kernel 0.8.3.
* Test oracle: FP64 references in
`tests/norm/test_cake_rmsnorm_train.py` (this PR).

## Export refresh: architecture-neutral delivery from the current Cake
exporter (2026-10-03)

The generated module was re-exported from the current Cake main (merged
into the producing branch) with the architecture-neutral layout and the
schedule-constant audit enabled; the kernels are unchanged.

- **Layout.** One generated device/binding pair per distinct program
text, shared by every architecture that builds it
(`csrc/cake_rmsnorm_train/` holds 21 programs instead of 46
per-architecture copies; 45 files / 12,602 lines instead of 92 files /
27,525 lines). The loader compiles the shared source with the exact flag
set of the device it runs on and keys the JIT cache by architecture;
routes are keyed per architecture to the shared program.
- **Binding.** Thin launcher on `tvm_ffi_utils.h`
(`ffi::CUDADeviceGuard`, `get_stream`, the target's check helpers); the
previous private device guard / mutex / context cache / check copies are
gone (binding smells 46 -> 0 in the export audit).
- **Constants.** The exporter's audit rejects delivered kernels that
differ only in an undeclared constant. The forward programs now name the
two knob literals as `#define ROWS_PER_CTA` / `GROUP_THREADS` and spell
`__launch_bounds__` with `LAUNCH_BOUNDS_*` macros; the declared physical
schedule axes of this family are `THREADS`, `LAUNCH_BOUNDS_*`,
`ROWS_PER_CTA`, `GROUP_THREADS`, `NUM_MAIN_STAGES` and the derived
`SMEM_*` layout defines. Two delivered kernel pairs differ only in
`THREADS` / `ROWS_PER_CTA` (the one-row and two-row variants of a
shape); the hidden size is a compile-time axis of this family
(per-thread vector counts, register caches and the partial-sum layout
derive from it), so programs of different widths differ structurally and
are not clones.
- **Kernel identity proof.** Every forward program of the
per-architecture tables was rendered from the previous and the refreshed
producer tree, compiled with `nvcc -cubin -O3 --prec-div=true
--prec-sqrt=true --ftz=false`, and compared with `cuobjdump -sass`
(addresses and symbol names normalised): sm_100a 14/14, sm_103a 15/15,
sm_107a 20/20 SASS-identical (text differs; no compile failures). The
producer's GPU correctness slice passes 16/16 on B200, GB300 and R200.
- **Host path.** Tensor arguments carry their strides (the only copy is
the documented aligned-copy fallback for a view that violates the
16-element row-stride / alignment contract); the per-call
device-capability query is unchanged.
- **Stage receipts (export protocol, 168 rows = 56 per architecture,
measured on each architecture's own hardware).** sm_107a (R200) 56/56
measured, 49 pass; sm_103a (GB300) 56/56, 47 pass; sm_100a (B200) 56/56,
47 pass — 143 of 168 rows pass every gate (correctness, source/export
parity, order dependence, every baseline ratio). The 25 non-passing
rows: the forward rows already reported above against
`flashinfer.norm.rmsnorm` (R200 H=6144 / H=2048 at T >= 16k, 0.97-0.99;
B200 and GB300 H=512 at 0.89-0.99), the small-row launch-floor ties
against `torch.compile` / Liger (H=512 T=256 / 1024, H=6144 residual
T=256 / 4096 at 0.98-1.00), and three 2-5 us rows whose source-first /
export-first timings disagree by more than the protocol's 10 % (R200 bwd
H=512 T=256, GB300 bwd H=2048 T=256 and fwd H=6144 T=256) — reproducible
across two passes. Each architecture is a separate stage on its own
hardware; the two largest B200 residual rows were measured in a fresh
process on a second B200 (the single-process run exhausts device memory
after 54 rows).
- **Downstream tests.** `tests/norm/test_cake_rmsnorm_train.py`: 133
passed / 28 skipped on R200 and on GB300 against the applied delivery;
the B200 configurations run in CI.

The rows that do not beat every baseline are the same open rows reported
above (the R200 H=6144 / H=2048 forward rows against
`flashinfer.norm.rmsnorm` at 0.97-0.99 and the small-row launch-floor
ties); nothing in this refresh changes a kernel or a measurement.

## Tests

```
pytest tests/norm/test_cake_rmsnorm_train.py -q
```

Skips unless the device capability is exactly (10, 0), (10, 3) or (10,
7).

Closes #5679

### Round 8 — SM107a H=512 program for T > 6144

Paired own-process sweeps on a Rubin R200 node (every arm bit-identical
to the shipped program): a 512-thread CTA that normalizes 16 rows (one
warp per row, so the reduction tree is unchanged) in the
software-pipelined persistent loop beats the shipped 256x8 programs at
every T from 8192 up, and the best persistent budget is 2 CTA(s) per SM:
4.03 vs 4.58 us @8192, 5.92 vs 6.34 @16231, 9.38 vs 10.50 @32768, 16.71
vs 18.43 @65536 (`flashinfer.norm.rmsnorm` 4.22 / 6.21 / 10.59 / 18.56).
The SM107a H=512 table becomes `quad` (T <= 256), `small` (single-trip
256x8, T <= 6144) and the new `main`; the H=6144 and H=2048 weight-hint
/ launch-bound / budget variants measured in the same round stay within
noise and are not adopted. SM100a / SM103a tables and generated sources
are unchanged.

Same-node paired A/B against the previous module running simultaneously
on another GPU (CUPTI cold-L2, 84 rows, two independent builds of the
new module): H=512 forward above 6144 rows now runs a 512-thread, 16-row
pipelined persistent CTA (one warp per row, so `y` / `rstd` are
bit-identical) at 2 CTAs per SM. fwd 16172 / 16231 / 32768 / 65536: +9.9
/ +9.8 / +11.3 / +8.6 % vs the previous module in build 1 and +9.3 /
+9.8 / +11.3 / +8.8 % in build 2; the fused step rows follow at +2.5 …
+3.7 %; no other row moved by more than the 0.5 % band. Against
`flashinfer.norm.rmsnorm` the four rows stand at 1.066-1.133 (previously
0.975-1.012).

| H=512 res=0 row | control (round 7) | build 1 | build 2 | threshold
arm | control x | build 1 x | build 2 x |
|---|---|---|---|---|---|---|---|
| fwd 16172 | 6.40 us | 5.82 us (+9.9 %) | 5.86 us (+9.3 %) | 6.24 us
FlashInfer rmsnorm | 0.975 | 1.066 | 1.071 |
| step 16172 | 18.27 | 17.70 (+3.2 %) | 17.82 (+2.5 %) | 38.27
torch.compile | 2.095 | 2.136 | 2.153 |
| fwd 16231 | 6.43 | 5.86 (+9.8 %) | 5.86 (+9.8 %) | 6.34 FlashInfer
rmsnorm | 0.985 | 1.071 | 1.082 |
| step 16231 | 18.27 | 17.70 (+3.3 %) | 17.79 (+2.7 %) | 38.30
torch.compile | 2.096 | 2.134 | 2.164 |
| fwd 32768 | 10.40 | 9.34 (+11.3 %) | 9.34 (+11.3 %) | 10.53 FlashInfer
rmsnorm | 1.012 | 1.123 | 1.133 |
| step 32768 | 28.58 | 27.55 (+3.7 %) | 27.59 (+3.6 %) | 35.10
torch.compile | 1.228 | 1.259 | 1.277 |
| fwd 65536 | 18.27 | 16.83 (+8.6 %) | 16.80 (+8.8 %) | 18.50 FlashInfer
rmsnorm | 1.012 | 1.101 | 1.105 |
| step 65536 | 48.16 | 46.75 (+3.0 %) | 46.72 (+3.1 %) | 57.38
torch.compile | 1.191 | 1.220 | 1.236 |

Validation on the R200 node: 785 / 785 facade correctness checks, GPU
e2e 16 / 16, compute-sanitizer memcheck + synccheck 0 errors on the
eight changed-window shapes (T = 6144 / 6145 / 8192 / 16231 / 16385 /
32768 / 65536 / 65537) and on the registered 16231 launcher; export
protocol correctness 56 / 56 and source-vs-export bitwise 56 / 56;
`tests/norm/test_cake_rmsnorm_train.py` 133 passed / 28 skipped on the
regenerated module. Generated sm_107a unit 2ded65dc7d99a604cec2 →
5e5665877a507da46c2c (the H=512 fwd mid role is gone).

### Round 11 — SM107a H=2048: the remaining forward gap is the per-row
FP32 `r` store; strided-input test coverage extended (no kernel change)

A plain-CUDA port of the shipped 64-thread H=2048 row group runs at
1.057-1.068 of `flashinfer.norm.rmsnorm` when it writes no per-row scale
and at 0.983-1.018 once it writes the training forward's FP32 `r[row]`
(one 4-byte store per row, +4-8 % at 16172 / 16231 / 32768 / 65536
rows); the port with every feature of the shipped kernel is bitwise
`y`-identical and times the shipped kernel. Store cache policies,
gathering the sector inside multi-row CTAs, persistent or
sector-interleaved row mappings and a cluster-gathered sector (eight
one-row CTAs per cluster) were all measured and are slower (the cluster
launch alone costs 31-36 % on this part). Writing `r` as one complete
32-byte sector per row (a 256-bit store of `{r, 0, 0, 0, 0, 0, 0, 0}`,
i.e. a padded `r` with stride 8) removes the whole penalty: on the
shipped SM107a programs it measures 1.058-1.062 of `rmsnorm` at 16172 /
16231 rows and 1.040-1.048 at 32768 / 65536 (shipped 0.987-1.010), and
1.011-1.019 on the H=6144 pipelined programs at 16k rows (shipped
0.975); it changes the storage layout of `rstd` (values unchanged) and
is therefore not adopted in this PR without an API decision; on SM100 /
SM103 the same padded layout is neutral to slightly negative
(x0.953-1.010 / x0.950-1.007 of the shipped programs, bitwise), so the
per-row `r` store costs nothing measurable there. Tables and generated
sources are unchanged.

Tests: the row-strided cases now include an 8-element pad (a
16-byte-only row stride), which exercises the wrapper's contiguous-copy
path for the forward's 32-byte alignment floor
(`FORWARD_ALIGNMENT_BYTES`) — the kernels' 256-bit row loads must never
see a row start that is only 16-byte aligned. The previous pads (64
elements and 2H) were all 32-byte multiples and never reached that path.

🤖 Generated with [Claude Code](https://claude.com/claude-code)



<!-- This is an auto-generated comment: release notes by coderabbit.ai
-->
## Summary by CodeRabbit

* **New Features**
* Added fused BF16 RMSNorm training forward and backward operations, an
autograd wrapper, and optional residual fusion for supported NVIDIA GPU
architectures.
* Added workspace sizing and allocation helpers, and made the RMSNorm
training APIs available through the package’s public interface.
* **Documentation**
* Documented the training RMSNorm API and its forward and backward
capabilities.
* **Chores**
* Added a benchmark comparing training RMSNorm performance across
implementations.
* Generated kernel sources carry a per-directory
`csrc/cake_rmsnorm_train/.clang-format` (`DisableFormat: true`) so
pre-commit leaves their receipt-attested bytes unchanged; the global
pre-commit configuration is not modified.
<!-- end of auto-generated comment: release notes by coderabbit.ai -->

Co-authored-by: Yingyi Huang <averyh@nvidia.com>
Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>
yyihuang added a commit that referenced this pull request Oct 12, 2026
…RMSNorm + RoPE overlapped with the Ulysses head exchange (SM90 / SM100 / SM103) (#6332)

## 📌 Description

Part of #4254. Implements the overlap requested in #6304 (assumptions
stated below where the issue thread has not confirmed a choice).

Adds `flashinfer.comm.QkvNormRopeExchangeOverlap`: one prepared session
that fuses the attention front end of a
sequence-parallel (Ulysses) layer — the BF16 QKV projection, per-head
RMSNorm of Q and K, interleaved RoPE from a
per-token cos/sin table, and the head exchange — into a single call per
layer. The Cake backend stores every
destination's head band straight from the GEMM epilogue into that rank's
symmetric receive slot over NVLink and
publishes per-chunk readiness words, so the exchange overlaps the
projection instead of following it; a one-CTA wait
kernel on a consumer stream acquires the words and the caller's stream
joins it.

* Architectures: SM90 (WGMMA producer), SM100 and SM103 (tcgen05
producer; 2-CTA pair tile for `block_n = 256` and
rows of at least 4096 tokens). Single-node NVLink groups of 2, 4 or 8
ranks, `head_dim = 128`, BF16 operands,
  fp32 cos/sin table.
* Layouts: `"bshd"` returns `(q, k, v)` views of `[B, S_global, H_local,
D]`; `"fused"` returns one
`[B, S_global, H_local, 3 * D]` payload. Outputs are views of the
current receive slot and stay valid until the
slot is reused `num_slots` calls later; `out=` copies into caller-owned
tensors.
* Backend selection: `backend="auto"` goes through
`resolve_ulysses_backend` (NVLink mesh probe) and the device
capability; otherwise the NCCL path (cuBLAS projection,
`flashinfer.norm.rmsnorm`, RoPE,
`UlyssesCommunicator.scatter_qkv_head_chunk`). `backend="nvlink"` raises
`UlyssesBackendError` when the mesh or the
architecture is not supported; the NCCL path reports `graph_capturable
== False`.
* CUDA graphs: a prepared call allocates nothing and runs no host
collective; captured calls publish a sentinel
epoch and clear their readiness words with a stream memset after the
wait kernel, so replays interleave with
eager calls. A call with more rows than `max_tokens`, or a row count
first seen inside a capture, raises instead of
allocating. At most `num_slots` calls of one session can be captured
into one graph without an eager call between.
* `heads % world_size != 0` raises `ValueError` deterministically on
every rank before any collective (same
  contract as `UlyssesCommunicator.scatter_heads`).

### Assumptions (open points in #6304)

The issue thread has not yet confirmed the following choices; this PR
implements the defaults below and labels them
here so reviewers and the requester can object:

* **Exchange algorithm**: Ulysses-style QKV all-to-all (each rank holds
a token shard and the full QKV weight; every
rank receives its head band of all tokens). Ring-style KV exchange is
out of scope.
* **Operator math**: the public FLUX definitions — fp32 RMSNorm with
`eps = 1e-6` and a shared `[head_dim]` weight per
Q / K, interleaved (pairwise) RoPE applied in fp32 from a per-token
cos/sin table, in the order projection → QK-norm →
  RoPE; BF16 operands and outputs (no FP8 producer).
* **Receive layout**: both the `(q, k, v)` `[B, S_global, H_local, D]`
views and the fused `[B*S_global, 3*H_local*D]`
payload are supported; the numbers below are measured with the `bshd`
layout on every architecture.
* **Tolerance**: `atol = rtol = 1e-2` against an fp32 step-by-step
reference. Isolated single-ulp BF16 differences
(a value landing on the other side of a BF16 rounding boundary after a
different fp32 accumulation order, max_abs
0.03125) occur both in this kernel and in the existing per-operator
path; the tests classify those elements explicitly
  instead of loosening the tolerance.

### Files

| Path | Content |
|---|---|
| `flashinfer/comm/qkv_norm_rope_exchange_overlap.py` | public class,
request validation, NCCL path |
| `flashinfer/comm/cake_qkv_norm_rope_exchange_overlap.py` | Cake
backend session: symmetric slots, barrier -> producer -> wait per call,
capture protocol |
| `flashinfer/comm/cake_qkv_exchange_plan.py` | tile plan and plan
policy (host only) |
| `flashinfer/jit/cake_qkv_norm_rope_exchange_overlap.py` | JIT loader:
program and route tables, `gen_jit_spec` per program and architecture |
| `csrc/cake_qkv_norm_rope_exchange_overlap/` | generated device and
binding translation units (one `run` launcher per program; excluded from
pre-commit like the other generated trees) |
| `tests/comm/test_qkv_norm_rope_exchange_overlap.py` | host-only plan /
validation / loader tests; 2-, 4-, 8-GPU correctness, capture/replay,
refusals, NCCL path |
| `benchmarks/comm/bench_qkv_norm_rope_exchange_overlap.py` | torchrun
benchmark against the per-operator paths (CUPTI, cold L2) |
| `docs/api/comm.rst`, `flashinfer/comm/__init__.py` | docs section and
re-export |

## 🔍 Related Issues

- #6304 — the feature request this PR implements; the choices the thread
has not confirmed yet are labelled under "Assumptions" above.
- Part of #4254.

## 🚀 Pull Request Checklist

Thank you for contributing to FlashInfer! Before we review your pull
request, please make sure the following items are complete.

### ✅ Pre-commit Checks

- [x] I have installed `pre-commit` by running `pip install pre-commit`
(or used your preferred method).
- [ ] I have installed the hooks with `pre-commit install`.
- [x] I have run the hooks manually with `pre-commit run --all-files`
and fixed any reported issues.

> If you are unsure about how to set up `pre-commit`, see [the
pre-commit documentation](https://pre-commit.com/).

The hooks were run as `pre-commit run --files <the 31 files changed by
this PR>` with the hook versions pinned in `.pre-commit-config.yaml`
(unchanged by this PR); all 12 hooks pass, including clang-format, mypy,
ruff check and ruff format.

## 🧪 Tests

- [x] Tests have been added or updated as needed.
- [ ] All tests are passing (`unittest`, etc.).

Every path is checked against an fp32 step-by-step reference
(projection, RMSNorm with the bf16 cast before the
bf16 scale, interleaved RoPE, bf16 output) redistributed with
`all_gather`, `atol = rtol = 1e-2`, on both layouts
including a tail row count (`S_local = 200`), after CUDA-graph replays
with fresh inputs, with two sessions in flight on
one group, and on the NCCL path. The multi-GPU tests spawn 2, 4 and 8
ranks on one node and skip on fewer devices or
an unsupported architecture.

The host-only tests of
`tests/comm/test_qkv_norm_rope_exchange_overlap.py` (`-k "not
multi_gpu"`, 21 passed) ran on the pre-check node; the 2-, 4- and 8-GPU
tests need a multi-GPU node and are left to CI.

### Test plan

- [ ] `pytest tests/comm/test_qkv_norm_rope_exchange_overlap.py -q` on
an 8-GPU SM100 node
- [ ] same on an 8-GPU SM103 node
- [ ] same on an 8-GPU SM90 node
- [ ] `torchrun --standalone --nproc-per-node=8
benchmarks/comm/bench_qkv_norm_rope_exchange_overlap.py --rows
representative --json bench_w8.json` per architecture
- [x] `pre-commit run --files <new files>` (31 files; 12 / 12 hooks
pass)

## 🔬 Experimental Track

<!-- Only for PRs submitted under the experimental policy
(CONTRIBUTING.md → "Experimental APIs and Backends").
     Leave this section untouched for normal PRs. -->

- [ ] This PR is **experimental**: it adds or changes code under
`flashinfer/experimental/` and/or an `@flashinfer_experimental_api`.
Tracking issue: #
- [ ] The tracking issue names an owner, the reason for the experimental
path, and a graduation plan with a target release.
- [ ] Core changes are limited to a thin entry point (signature, shared
validation, feature-gate check, backend selection, handoff).
- [ ] Tests live in `tests/experimental/` and were validated on the
intended hardware; a runnable example is included.
- [ ] Nothing is registered in `flashinfer/aot.py`, and no experimental
backend is reachable from `backend="auto"` without
`FLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS=1`. (Calling an
`@flashinfer_experimental_api` or naming a backend explicitly is itself
the opt-in and needs no environment variable.)
- [ ] **Test scope declared below.** The experimental CI lane runs
exactly these targets, so keep them as narrow as the change allows.

<!-- Required for experimental PRs. Replace the commented lines below
with your targets.
Do not delete the fence or change its `experimental-tests` tag — the
experimental-track
watcher reads it verbatim to decide which targets to ask CI for. -->

```experimental-tests
# One target per line: a directory or a file. (A pytest ::selector is not
# supported -- the sharding runner cannot consume one.) Must be under
# tests/experimental/ and must exist. Delete these comment lines and add yours, e.g.
#
#   tests/experimental/test_my_backend.py
#   tests/experimental/my_backend/
#
# Declaring the whole tree (tests/experimental/) is allowed but means every
# experimental PR pays for every other feature's tests, in every matrix cell.
```

## Reviewer Notes

- The device and binding translation units under
`csrc/cake_qkv_norm_rope_exchange_overlap/` are generated and attested
by content hash; the host path
(`flashinfer/comm/qkv_norm_rope_exchange_overlap.py`,
`cake_qkv_norm_rope_exchange_overlap.py`, `cake_qkv_exchange_plan.py`),
the JIT loader and the tests are hand-written and are the parts to read
first.
- The Speedup section separates the benchmark table from the
export-protocol receipts; rows that failed the protocol's validity gates
(power-capped clocks) are listed, not reported as measured.

## Speedup

Measured with `benchmarks/comm/bench_qkv_norm_rope_exchange_overlap.py`
(CUPTI activity timing, cold L2, per-rank
maximum then median; `--rows representative`) against the fastest
per-operator path on the same node. The
per-operator baselines are: B1 three `scatter_heads` collectives, B2 one
`scatter_qkv_head_chunk` over the whole head
band, B3 the default head-chunk schedule, B4
`ulysses_a2a_in(fuse_qkv=True)` over NCCL. On SM90 the Ulysses
communicator of B1-B3 runs on its NCCL backend (the NVLink Ulysses
kernels target SM100 / SM103).

### Measured speedup (72 rows per architecture)

Grid: `W ∈ {2, 4, 8}` ranks on one node × `hidden ∈ {3072, 4096, 6144,
8192}` (`head_dim = 128`) × `S_global ∈ {4096 … 131072}`,
`B = 1`, BF16. Each row was timed on one 8-GPU node with CUPTI activity
tracing (cold L2, per-rank maximum then median,
five A/B/B/A rounds of ten samples, an A/A control arm within ±1 % or
the row was re-measured) against the fastest of
the four per-operator paths above measured in the same process. speedup
= fastest baseline / fused. Every measured row is
faster than its fastest baseline.

| GPU | rows faster than the fastest baseline | geometric mean
(A/A-valid rows) | W = 8 / 4 / 2 | minimum |
|---|---|---|---|---|
| B200 (SM100) | 72 / 72 | 1.831 | 1.824 / 1.849 / 1.820 | 1.261 |
| B300 (SM103) | 72 / 72 | 1.906 | 1.911 / 1.891 / 1.916 | 1.241 |
| H100 (SM90) | 71 / 71 | 1.332 | 1.339 / 1.348 / 1.310 | 1.135 |

Notes: on B300 the row `W = 4, hidden 3072, S_global 8192` (2.26–2.31×
in three measurements) exceeded the ±1 % A/A
control in all three and is excluded from the geometric mean; on H100
(80 GB) the row `W = 2, hidden 8192,
S_global 131072` does not fit in device memory with the reference and
baseline buffers and is not measured.

<details>
<summary>B200 (SM100) — per-row table</summary>

| W | hidden | S_global | fused ms | fastest baseline ms | speedup | A/A
% |
|---|---|---|---|---|---|---|
| 8 | 3072 | 4096 | 0.0722 | 0.2389 | 3.307 | 0.11 |
| 8 | 3072 | 8192 | 0.1014 | 0.2493 | 2.458 | 0.88 |
| 8 | 3072 | 16384 | 0.1562 | 0.2821 | 1.807 | 0.09 |
| 8 | 3072 | 32768 | 0.2713 | 0.4512 | 1.663 | 0.20 |
| 8 | 3072 | 65536 | 0.5131 | 0.8815 | 1.718 | 0.10 |
| 8 | 3072 | 131072 | 0.9575 | 1.7075 | 1.783 | 0.14 |
| 8 | 4096 | 4096 | 0.0968 | 0.2440 | 2.522 | 0.71 |
| 8 | 4096 | 8192 | 0.1193 | 0.2576 | 2.160 | 0.34 |
| 8 | 4096 | 16384 | 0.1834 | 0.3626 | 1.977 | 0.67 |
| 8 | 4096 | 32768 | 0.3086 | 0.6454 | 2.091 | 0.25 |
| 8 | 4096 | 65536 | 0.5906 | 1.2448 | 2.108 | 0.05 |
| 8 | 4096 | 131072 | 1.1808 | 2.4446 | 2.070 | 0.11 |
| 8 | 6144 | 4096 | 0.1224 | 0.2503 | 2.045 | 0.62 |
| 8 | 6144 | 8192 | 0.1902 | 0.3208 | 1.686 | 0.69 |
| 8 | 6144 | 16384 | 0.3327 | 0.5789 | 1.740 | 0.32 |
| 8 | 6144 | 32768 | 0.6438 | 1.1486 | 1.784 | 0.06 |
| 8 | 6144 | 65536 | 1.3623 | 2.3209 | 1.704 | 0.17 |
| 8 | 6144 | 131072 | 2.8806 | 4.4242 | 1.536 | 0.01 |
| 8 | 8192 | 4096 | 0.1897 | 0.2766 | 1.458 | 0.26 |
| 8 | 8192 | 8192 | 0.3278 | 0.4798 | 1.464 | 0.54 |
| 8 | 8192 | 16384 | 0.5804 | 0.9063 | 1.562 | 0.22 |
| 8 | 8192 | 32768 | 1.2319 | 1.7849 | 1.449 | 0.08 |
| 8 | 8192 | 65536 | 2.5323 | 3.5145 | 1.388 | 0.06 |
| 8 | 8192 | 131072 | 5.1799 | 6.9348 | 1.339 | 0.00 |
| 4 | 3072 | 4096 | 0.0943 | 0.2456 | 2.605 | 0.25 |
| 4 | 3072 | 8192 | 0.1316 | 0.2755 | 2.093 | 0.67 |
| 4 | 3072 | 16384 | 0.2077 | 0.4497 | 2.165 | 0.31 |
| 4 | 3072 | 32768 | 0.3794 | 0.8772 | 2.312 | 0.19 |
| 4 | 3072 | 65536 | 0.6863 | 1.6978 | 2.474 | 0.06 |
| 4 | 3072 | 131072 | 1.4040 | 3.3245 | 2.368 | 0.10 |
| 4 | 4096 | 4096 | 0.1200 | 0.2506 | 2.088 | 0.23 |
| 4 | 4096 | 8192 | 0.1768 | 0.3601 | 2.037 | 0.50 |
| 4 | 4096 | 16384 | 0.3100 | 0.6342 | 2.046 | 0.34 |
| 4 | 4096 | 32768 | 0.5834 | 1.2666 | 2.171 | 0.13 |
| 4 | 4096 | 65536 | 1.1508 | 2.4901 | 2.164 | 0.01 |
| 4 | 4096 | 131072 | 2.3583 | 4.9367 | 2.093 | 0.01 |
| 4 | 6144 | 4096 | 0.1901 | 0.3145 | 1.654 | 0.87 |
| 4 | 6144 | 8192 | 0.3349 | 0.5856 | 1.748 | 0.37 |
| 4 | 6144 | 16384 | 0.6606 | 1.1643 | 1.763 | 0.13 |
| 4 | 6144 | 32768 | 1.3538 | 2.3620 | 1.745 | 0.04 |
| 4 | 6144 | 65536 | 2.8794 | 4.4967 | 1.562 | 0.14 |
| 4 | 6144 | 131072 | 5.8146 | 8.9359 | 1.537 | 0.01 |
| 4 | 8192 | 4096 | 0.3220 | 0.4808 | 1.493 | 0.05 |
| 4 | 8192 | 8192 | 0.5756 | 0.9014 | 1.566 | 0.03 |
| 4 | 8192 | 16384 | 1.2447 | 1.8077 | 1.452 | 0.03 |
| 4 | 8192 | 32768 | 2.5719 | 3.5700 | 1.388 | 0.10 |
| 4 | 8192 | 65536 | 5.2743 | 7.0276 | 1.332 | 0.01 |
| 4 | 8192 | 131072 | 10.2687 | 14.3224 | 1.395 | 0.23 |
| 2 | 3072 | 4096 | 0.1234 | 0.2666 | 2.160 | 0.92 |
| 2 | 3072 | 8192 | 0.1972 | 0.4282 | 2.171 | 0.52 |
| 2 | 3072 | 16384 | 0.3494 | 0.8643 | 2.474 | 0.04 |
| 2 | 3072 | 32768 | 0.6568 | 1.6784 | 2.555 | 0.16 |
| 2 | 3072 | 65536 | 1.3126 | 3.2883 | 2.505 | 0.20 |
| 2 | 3072 | 131072 | 2.7374 | 6.5161 | 2.380 | 0.13 |
| 2 | 4096 | 4096 | 0.1788 | 0.3531 | 1.975 | 0.69 |
| 2 | 4096 | 8192 | 0.3104 | 0.6411 | 2.065 | 0.40 |
| 2 | 4096 | 16384 | 0.5834 | 1.2492 | 2.141 | 0.11 |
| 2 | 4096 | 32768 | 1.1547 | 2.4569 | 2.128 | 0.06 |
| 2 | 4096 | 65536 | 2.3984 | 4.8868 | 2.038 | 0.10 |
| 2 | 4096 | 131072 | 4.8963 | 9.6901 | 1.979 | 0.01 |
| 2 | 6144 | 4096 | 0.3260 | 0.5658 | 1.735 | 0.40 |
| 2 | 6144 | 8192 | 0.6637 | 1.1501 | 1.733 | 0.14 |
| 2 | 6144 | 16384 | 1.3643 | 2.3287 | 1.707 | 0.01 |
| 2 | 6144 | 32768 | 2.8530 | 4.4548 | 1.561 | 0.14 |
| 2 | 6144 | 65536 | 5.7498 | 8.8343 | 1.536 | 0.26 |
| 2 | 6144 | 131072 | 11.7871 | 17.7952 | 1.510 | 0.15 |
| 2 | 8192 | 4096 | 0.5740 | 0.8991 | 1.566 | 0.21 |
| 2 | 8192 | 8192 | 1.2396 | 1.7892 | 1.443 | 0.01 |
| 2 | 8192 | 16384 | 2.5594 | 3.5306 | 1.379 | 0.08 |
| 2 | 8192 | 32768 | 5.2429 | 6.9682 | 1.329 | 0.23 |
| 2 | 8192 | 65536 | 10.5083 | 14.1865 | 1.350 | 0.21 |
| 2 | 8192 | 131072 | 23.0629 | 29.0725 | 1.261 | 0.63 |

</details>

<details>
<summary>B300 (SM103) — per-row table</summary>

| W | hidden | S_global | fused ms | fastest baseline ms | speedup | A/A
% |
|---|---|---|---|---|---|---|
| 8 | 3072 | 4096 | 0.0774 | 0.2697 | 3.487 | 0.42 |
| 8 | 3072 | 8192 | 0.1027 | 0.2863 | 2.788 | 0.66 |
| 8 | 3072 | 16384 | 0.1579 | 0.3131 | 1.982 | 0.88 |
| 8 | 3072 | 32768 | 0.2759 | 0.4525 | 1.640 | 0.39 |
| 8 | 3072 | 65536 | 0.5072 | 0.8430 | 1.662 | 0.14 |
| 8 | 3072 | 131072 | 0.9269 | 1.6214 | 1.749 | 0.02 |
| 8 | 4096 | 4096 | 0.0982 | 0.2790 | 2.842 | 0.36 |
| 8 | 4096 | 8192 | 0.1223 | 0.2896 | 2.367 | 0.54 |
| 8 | 4096 | 16384 | 0.1788 | 0.3498 | 1.956 | 0.49 |
| 8 | 4096 | 32768 | 0.2904 | 0.6136 | 2.113 | 0.04 |
| 8 | 4096 | 65536 | 0.5417 | 1.1756 | 2.170 | 0.12 |
| 8 | 4096 | 131072 | 1.0635 | 2.3010 | 2.164 | 0.05 |
| 8 | 6144 | 4096 | 0.1258 | 0.2846 | 2.263 | 0.61 |
| 8 | 6144 | 8192 | 0.1809 | 0.3167 | 1.751 | 0.51 |
| 8 | 6144 | 16384 | 0.3056 | 0.5489 | 1.797 | 0.12 |
| 8 | 6144 | 32768 | 0.5970 | 1.0776 | 1.805 | 0.01 |
| 8 | 6144 | 65536 | 1.2046 | 2.0970 | 1.741 | 0.06 |
| 8 | 6144 | 131072 | 2.4691 | 4.1316 | 1.673 | 0.01 |
| 8 | 8192 | 4096 | 0.1852 | 0.2919 | 1.576 | 0.01 |
| 8 | 8192 | 8192 | 0.3020 | 0.4518 | 1.496 | 0.74 |
| 8 | 8192 | 16384 | 0.5334 | 0.8476 | 1.589 | 0.19 |
| 8 | 8192 | 32768 | 1.0841 | 1.6653 | 1.536 | 0.10 |
| 8 | 8192 | 65536 | 2.1971 | 3.2764 | 1.491 | 0.07 |
| 8 | 8192 | 131072 | 4.4165 | 6.4274 | 1.455 | 0.03 |
| 4 | 3072 | 4096 | 0.0945 | 0.2622 | 2.774 | 0.69 |
| 4 | 3072 | 8192 | 0.1346 | 0.3073 | 2.283 (A/A > 1 %) | 4.33 |
| 4 | 3072 | 16384 | 0.2104 | 0.4479 | 2.128 | 0.17 |
| 4 | 3072 | 32768 | 0.3899 | 0.8362 | 2.145 | 0.33 |
| 4 | 3072 | 65536 | 0.6896 | 1.6060 | 2.329 | 0.11 |
| 4 | 3072 | 131072 | 1.3541 | 3.1379 | 2.317 | 0.05 |
| 4 | 4096 | 4096 | 0.1207 | 0.2752 | 2.280 | 0.66 |
| 4 | 4096 | 8192 | 0.1735 | 0.3476 | 2.003 | 0.09 |
| 4 | 4096 | 16384 | 0.2905 | 0.6134 | 2.111 | 0.01 |
| 4 | 4096 | 32768 | 0.5332 | 1.1986 | 2.248 | 0.06 |
| 4 | 4096 | 65536 | 1.0389 | 2.3463 | 2.258 | 0.02 |
| 4 | 4096 | 131072 | 2.0955 | 4.6322 | 2.211 | 0.01 |
| 4 | 6144 | 4096 | 0.1834 | 0.3102 | 1.691 | 0.39 |
| 4 | 6144 | 8192 | 0.3131 | 0.5571 | 1.779 | 0.09 |
| 4 | 6144 | 16384 | 0.6136 | 1.0952 | 1.785 | 0.33 |
| 4 | 6144 | 32768 | 1.2197 | 2.1381 | 1.753 | 0.01 |
| 4 | 6144 | 65536 | 2.4356 | 4.1972 | 1.723 | 0.06 |
| 4 | 6144 | 131072 | 4.9588 | 8.3107 | 1.676 | 0.03 |
| 4 | 8192 | 4096 | 0.3039 | 0.4521 | 1.488 | 0.15 |
| 4 | 8192 | 8192 | 0.5342 | 0.8521 | 1.595 | 0.06 |
| 4 | 8192 | 16384 | 1.1044 | 1.6901 | 1.530 | 0.09 |
| 4 | 8192 | 32768 | 2.1950 | 3.3179 | 1.512 | 0.08 |
| 4 | 8192 | 65536 | 4.4062 | 6.5251 | 1.481 | 0.10 |
| 4 | 8192 | 131072 | 9.4225 | 13.3648 | 1.418 | 0.28 |
| 2 | 3072 | 4096 | 0.1199 | 0.2954 | 2.463 | 0.11 |
| 2 | 3072 | 8192 | 0.1977 | 0.4372 | 2.211 | 0.15 |
| 2 | 3072 | 16384 | 0.3375 | 0.8262 | 2.448 | 0.06 |
| 2 | 3072 | 32768 | 0.5991 | 1.5968 | 2.665 | 0.04 |
| 2 | 3072 | 65536 | 1.1957 | 3.1269 | 2.615 | 0.01 |
| 2 | 3072 | 131072 | 2.3627 | 6.1780 | 2.615 | 0.10 |
| 2 | 4096 | 4096 | 0.1690 | 0.3380 | 1.999 | 0.06 |
| 2 | 4096 | 8192 | 0.2898 | 0.6124 | 2.113 | 0.01 |
| 2 | 4096 | 16384 | 0.5334 | 1.1898 | 2.230 | 0.10 |
| 2 | 4096 | 32768 | 1.0432 | 2.3309 | 2.234 | 0.09 |
| 2 | 4096 | 65536 | 2.0717 | 4.6073 | 2.224 | 0.01 |
| 2 | 4096 | 131072 | 4.2197 | 9.1462 | 2.167 | 0.05 |
| 2 | 6144 | 4096 | 0.2977 | 0.5423 | 1.822 | 0.04 |
| 2 | 6144 | 8192 | 0.6118 | 1.0884 | 1.779 | 0.04 |
| 2 | 6144 | 16384 | 1.1907 | 2.1221 | 1.782 | 0.02 |
| 2 | 6144 | 32768 | 2.3919 | 4.1900 | 1.752 | 0.04 |
| 2 | 6144 | 65536 | 4.9604 | 8.2729 | 1.668 | 0.34 |
| 2 | 6144 | 131072 | 11.4356 | 16.7756 | 1.467 | 0.23 |
| 2 | 8192 | 4096 | 0.5286 | 0.8535 | 1.615 | 0.11 |
| 2 | 8192 | 8192 | 1.0660 | 1.6769 | 1.573 | 0.01 |
| 2 | 8192 | 16384 | 2.1406 | 3.3157 | 1.549 | 0.22 |
| 2 | 8192 | 32768 | 4.5065 | 6.5023 | 1.443 | 0.24 |
| 2 | 8192 | 65536 | 9.5439 | 13.3589 | 1.400 | 0.16 |
| 2 | 8192 | 131072 | 22.1867 | 27.5295 | 1.241 | 0.03 |

</details>

<details>
<summary>H100 (SM90) — per-row table</summary>

| W | hidden | S_global | fused ms | fastest baseline ms | speedup | A/A
% |
|---|---|---|---|---|---|---|
| 8 | 3072 | 4096 | 0.1070 | 0.1708 | 1.595 | 0.54 |
| 8 | 3072 | 8192 | 0.1587 | 0.2288 | 1.442 | 0.13 |
| 8 | 3072 | 16384 | 0.2722 | 0.3818 | 1.403 | 0.23 |
| 8 | 3072 | 32768 | 0.5272 | 0.7191 | 1.364 | 0.16 |
| 8 | 3072 | 65536 | 1.0614 | 1.4336 | 1.351 | 0.57 |
| 8 | 3072 | 131072 | 2.1269 | 2.8805 | 1.354 | 0.34 |
| 8 | 4096 | 4096 | 0.1444 | 0.1875 | 1.298 | 0.33 |
| 8 | 4096 | 8192 | 0.2046 | 0.3137 | 1.533 | 0.45 |
| 8 | 4096 | 16384 | 0.3804 | 0.5538 | 1.456 | 0.29 |
| 8 | 4096 | 32768 | 0.7378 | 1.1132 | 1.509 | 0.87 |
| 8 | 4096 | 65536 | 1.4582 | 2.1800 | 1.495 | 0.03 |
| 8 | 4096 | 131072 | 2.9126 | 4.3244 | 1.485 | 0.13 |
| 8 | 6144 | 4096 | 0.2603 | 0.3047 | 1.171 | 0.50 |
| 8 | 6144 | 8192 | 0.4252 | 0.5419 | 1.274 | 0.27 |
| 8 | 6144 | 16384 | 0.7434 | 1.0295 | 1.385 | 0.33 |
| 8 | 6144 | 32768 | 1.5361 | 2.0466 | 1.332 | 0.46 |
| 8 | 6144 | 65536 | 3.2696 | 4.0454 | 1.237 | 0.11 |
| 8 | 6144 | 131072 | 6.4700 | 8.0205 | 1.240 | 0.01 |
| 8 | 8192 | 4096 | 0.3395 | 0.4407 | 1.298 | 0.41 |
| 8 | 8192 | 8192 | 0.6529 | 0.8201 | 1.256 | 0.36 |
| 8 | 8192 | 16384 | 1.3057 | 1.6042 | 1.229 | 0.16 |
| 8 | 8192 | 32768 | 2.6329 | 3.2563 | 1.237 | 0.08 |
| 8 | 8192 | 65536 | 5.4181 | 6.4155 | 1.184 | 0.25 |
| 8 | 8192 | 131072 | 11.4591 | 13.0004 | 1.135 | 0.23 |
| 4 | 3072 | 4096 | 0.1570 | 0.2114 | 1.347 | 0.18 |
| 4 | 3072 | 8192 | 0.2545 | 0.3688 | 1.449 | 0.02 |
| 4 | 3072 | 16384 | 0.4623 | 0.7117 | 1.540 | 0.05 |
| 4 | 3072 | 32768 | 0.8927 | 1.3835 | 1.550 | 0.02 |
| 4 | 3072 | 65536 | 1.7488 | 2.7224 | 1.557 | 0.08 |
| 4 | 3072 | 131072 | 3.4939 | 5.3476 | 1.531 | 0.11 |
| 4 | 4096 | 4096 | 0.1987 | 0.2856 | 1.437 | 0.02 |
| 4 | 4096 | 8192 | 0.3767 | 0.5468 | 1.452 | 0.09 |
| 4 | 4096 | 16384 | 0.7412 | 1.0631 | 1.434 | 0.06 |
| 4 | 4096 | 32768 | 1.4287 | 2.0803 | 1.456 | 0.01 |
| 4 | 4096 | 65536 | 2.8435 | 4.1326 | 1.453 | 0.02 |
| 4 | 4096 | 131072 | 5.7662 | 8.1828 | 1.419 | 0.27 |
| 4 | 6144 | 4096 | 0.4104 | 0.5240 | 1.277 | 0.08 |
| 4 | 6144 | 8192 | 0.7647 | 1.0093 | 1.320 | 0.07 |
| 4 | 6144 | 16384 | 1.4989 | 1.9741 | 1.317 | 0.11 |
| 4 | 6144 | 32768 | 3.0037 | 3.9102 | 1.302 | 0.05 |
| 4 | 6144 | 65536 | 6.1564 | 7.8220 | 1.271 | 0.11 |
| 4 | 6144 | 131072 | 13.0916 | 15.5928 | 1.191 | 0.02 |
| 4 | 8192 | 4096 | 0.6562 | 0.8133 | 1.239 | 0.12 |
| 4 | 8192 | 8192 | 1.2991 | 1.5901 | 1.224 | 0.00 |
| 4 | 8192 | 16384 | 2.5754 | 3.1704 | 1.231 | 0.00 |
| 4 | 8192 | 32768 | 5.2416 | 6.2735 | 1.197 | 0.35 |
| 4 | 8192 | 65536 | 11.1591 | 12.7948 | 1.147 | 0.10 |
| 4 | 8192 | 131072 | 22.6522 | 26.1679 | 1.155 | 0.07 |
| 2 | 3072 | 4096 | 0.2538 | 0.3503 | 1.380 | 0.19 |
| 2 | 3072 | 8192 | 0.4533 | 0.6655 | 1.468 | 0.04 |
| 2 | 3072 | 16384 | 0.8781 | 1.2965 | 1.476 | 0.06 |
| 2 | 3072 | 32768 | 1.7103 | 2.5706 | 1.503 | 0.00 |
| 2 | 3072 | 65536 | 3.4275 | 4.9401 | 1.441 | 0.03 |
| 2 | 3072 | 131072 | 7.0964 | 9.8401 | 1.387 | 0.17 |
| 2 | 4096 | 4096 | 0.3736 | 0.5182 | 1.387 | 0.19 |
| 2 | 4096 | 8192 | 0.7116 | 1.0057 | 1.413 | 0.06 |
| 2 | 4096 | 16384 | 1.3972 | 1.9768 | 1.415 | 0.00 |
| 2 | 4096 | 32768 | 2.7507 | 3.9208 | 1.425 | 0.02 |
| 2 | 4096 | 65536 | 5.5682 | 7.6357 | 1.371 | 0.06 |
| 2 | 4096 | 131072 | 11.8074 | 15.2643 | 1.293 | 0.08 |
| 2 | 6144 | 4096 | 0.7606 | 0.9606 | 1.263 | 0.06 |
| 2 | 6144 | 8192 | 1.4823 | 1.8973 | 1.280 | 0.02 |
| 2 | 6144 | 16384 | 2.9507 | 3.7534 | 1.272 | 0.09 |
| 2 | 6144 | 32768 | 5.9493 | 7.3609 | 1.237 | 0.13 |
| 2 | 6144 | 65536 | 12.6436 | 14.8480 | 1.174 | 0.12 |
| 2 | 6144 | 131072 | 26.1080 | 30.5084 | 1.169 | 0.46 |
| 2 | 8192 | 4096 | 1.2689 | 1.5342 | 1.209 | 0.09 |
| 2 | 8192 | 8192 | 2.5373 | 3.0496 | 1.202 | 0.12 |
| 2 | 8192 | 16384 | 5.1135 | 6.0219 | 1.178 | 0.22 |
| 2 | 8192 | 32768 | 10.6399 | 12.2423 | 1.151 | 0.22 |
| 2 | 8192 | 65536 | 21.6914 | 25.0241 | 1.154 | 0.07 |
| 2 | 8192 | 131072 | — | — | not measured (device memory) | — |

</details>

### Export-protocol validation (per-row receipts, three architectures)

Independently of the table above, every row was re-measured by the
export protocol that produced the generated
sources in this PR: same CUPTI method (cold L2, per-rank maximum then
median, counterbalanced A/B/B/A groups), the
generated programs loaded through this PR's JIT path, the four
per-operator baselines measured in the same process,
plus the validity gates per row (source/export agreement >= 0.95,
directional disagreement <= 0.12, endpoint
drift <= 0.15, loaded SM clock >= 50 % of the maximum boost clock). A
row counts only when all gates pass.

**Protocol change (2026-10-11): loaded-SM-clock floor 0.60 -> 0.50.**
The floor used to be 0.60. At the largest
`W = 2 / 4, S_global = 131072` rows the B200 (1000 W) and B300 (1100 W)
board power limits hold the loaded SM clock
at 0.48 - 0.60 of the maximum on every node tried (five nodes per SKU,
three rounds), so those rows were rejected
unconditionally. The floor is now 0.50; the directional, drift and
source gates are unchanged. Every previously
accepted row was measured under the stricter 0.60 floor and remains
valid as measured; under the relaxed floor one
further row was accepted (`corr_w2_max` on B300, a correctness row). The
rows that still fail are listed below.

| GPU | perf rows with a valid receipt | all valid rows faster than the
fastest baseline | geometric mean (W = 8 / 4 / 2) | minimum | rows
without a valid receipt |
|---|---|---|---|---|---|
| B200 (SM100) | 69 / 72 | yes | 1.686 (1.535 / 1.716 / 1.832) | 1.170 |
3 |
| B300 (SM103) | 69 / 72 | yes | 1.700 (1.511 / 1.737 / 1.889) | 1.109 |
3 |
| H100 (SM90) | 70 / 72 | yes | 1.306 (1.296 / 1.319 / 1.301) | 1.093 |
2 |

Correctness / negative rows of the export protocol (13 per GPU,
including the largest-shape correctness row
`corr_w2_max`): H100 13 / 13, B300 13 / 13, B200 12 / 13 (`corr_w2_max`
reads 0.496 and 0.492 of the maximum clock
on two nodes, below the 0.50 floor; its correctness is covered by the
harness correctness table above).

Rows without a valid receipt (disclosed, not counted; the raw ratio of
an invalid attempt is not a measurement):

| GPU | W | hidden | S_global | why the receipt is invalid | attempts
(nodes) | raw ratio of the last attempt |
|---|---|---|---|---|---|---|
| B200 | 2 | 8192 | 131072 | rounds 1 - 2: loaded SM clock 847 / 1170 /
1147 MHz of 1965 under the 1000 W limit; round 3 at the 0.50 floor: B1
endpoint drift 0.30 / 0.22 and B2 directional disagreement > 0.12 | 4 (3
nodes) | 1.37 - 1.69 |
| B200 | 4 | 8192 | 131072 | rounds 1 - 2: loaded SM clock below the 60
% floor (960 / 1177 MHz); round 3: B1 directional disagreement > 0.12 |
3 (3 nodes) | 1.55 - 1.89 |
| B200 | 2 | 6144 | 131072 | directional disagreement of the B1 / B3
baselines > 0.12 (candidate gates pass) | 2 (2 nodes) | 1.54 - 1.90 |
| B300 | 2 | 8192 | 131072 | rounds 1 - 2: loaded SM clock 1117 / 1140 /
1170 MHz of 2032 under the 1100 W limit; round 3: B2 / B3 directional
disagreement > 0.12 | 4 (3 nodes) | 1.44 - 1.84 |
| B300 | 4 | 8192 | 131072 | rounds 1 - 2: loaded SM clock 1147 / 1215
MHz of 2032; round 3: B1 directional disagreement, B2 directional
disagreement + endpoint drift | 4 (3 nodes) | 1.47 - 1.90 |
| B300 | 2 | 6144 | 131072 | directional disagreement of the B2 / B3
baselines > 0.12 | 3 (2 nodes) | 1.70 - 2.16 |
| H100 | 8 | 8192 | 131072 | directional disagreement of the B4 (NCCL)
baseline 0.15 - 0.18 > 0.12 in all five attempts (candidate and B1 - B3
gates pass); the raw B2 ratio of the attempts that reached the arms is
0.97 - 0.99 | 5 (5 nodes) | 0.97 - 1.15 |
| H100 | 2 | 8192 | 131072 | does not fit in 80 GB with the reference
and the four baseline arms (not measurable) | - | - |

The largest shapes run power-capped on the B200 / B300 nodes available
for this work (clocks at 48 - 60 % of the
maximum boost); once the clock floor no longer rejects them, the order
effect between the baseline arms (each tens of
milliseconds, with power-state transitions between arms) exceeds the
0.12 directional gate on every node tried. The
per-operator table above measured those rows on other nodes of the same
type, where the row
`W = 2, hidden 8192, S_global 131072` reached 23 ms vs 28.9 ms. The
export protocol is not relaxed further for them;
they are listed here rather than reported as measured.

### Baselines and their source PRs

| Baseline | Path | Source PRs |
|---|---|---|
| B1 ordinary | cuBLAS projection → `flashinfer.norm.rmsnorm` →
`apply_rope_with_cos_sin_cache` → `UlyssesCommunicator.scatter_heads` ×
3 | #3820, #4240 (communicator), #4480, #4202 (rmsnorm), #5952, #5956,
#6188 (RoPE) |
| B2 whole-QKV head chunk | one `scatter_qkv_head_chunk` collective over
the whole head band | #5027 |
| B3 head-chunk schedule | `scatter_qkv_head_chunk` per head band on
three streams, the band schedule of
`benchmarks/comm/bench_ulysses_head_chunk.py` (default schedule per row)
| #5027, #5263, #5953 |
| B4 NCCL all-to-all | `parallel_wrapper.ulysses_a2a_in(fuse_qkv=True)`
| #2630 |

Baseline implementation files (all at the merge-base):
`flashinfer/comm/ulysses.py` (B1 `scatter_heads`, B2 and B3
`scatter_qkv_head_chunk`, and `UlyssesCommunicator`),
`benchmarks/comm/bench_ulysses_head_chunk.py` (B3 band schedule),
`flashinfer/norm/` (B1 `rmsnorm`), `flashinfer/rope.py` (B1
`apply_rope_with_cos_sin_cache`),
`flashinfer/parallel_attention/parallel_wrapper.py` (B4
`ulysses_a2a_in`). This branch is based on `34bb66e4` (its merge-base
with `main`); `git log 34bb66e..main -- <those files>` is empty as of
`main` = `65c38d2a`, so the baselines measured here are the
implementations currently in `main`. Source PR merge commits: #3820
`8a94642d8`, #4240 `668a1ba1c`, #4480 `6d5ef77f8`, #4202 `519205914`,
#5952 `a973cfc92`, #5956 `335c4eb4e`, #6188 `d71cc4b38`, #5027
`15e83b7bb`, #5263 `8c54f3446`, #5953 `7e73a7470`, #2630 `ce02358d2`.

🤖 Generated with [Claude Code](https://claude.com/claude-code)


<!-- This is an auto-generated comment: release notes by coderabbit.ai
-->
## Summary by CodeRabbit

* **New Features**
* Added a prepared sequence-parallel API combining QKV projection, Q/K
normalization, rotary embeddings, and head exchange.
* Automatically selects the optimized NVLink backend when supported,
with an NCCL fallback for other configurations.
* Supports CUDA graph capture on the optimized backend and reusable
output slots.
* **Documentation**
* Added API guidance on supported configurations, backend selection,
outputs, and capture behavior.
* **Benchmarks**
* Added a distributed benchmark comparing the fused operation with
separate operator paths, including output checks and optional JSON
results.
<!-- end of auto-generated comment: release notes by coderabbit.ai -->

---------

Co-authored-by: Yingyi Huang <averyh@nvidia.com>
Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants