feat(rdna2): default the MoE quantized GEMM epilogue to fp32

csrc/rocm/moe_accum_rdna2.cuh

@@ -6,16 +6,21 @@
 // ``moe_q_gemm_rdna2.cu`` (W4A16) and ``moe_w4a8_rdna2.cu`` (W4A8) emit the
 // same output contract: one token-block row of 4 consecutive N columns with
 // the router weight already folded in fp32, atomically reduced across experts
-// when ``output_topk > 0``. Two accumulation modes:
+// when ``output_topk > 0``. Accumulation modes:
 //
-//   * fp32 (C_T == float, default): native v_global_atomic_add_f32 into a
-//     cached fp32 scratch -- no CAS, no intermediate fp16 rounding. A single
-//     fp32 -> fp16 cast runs once at the end, so the result does not depend
-//     on the atomic ordering at fp16 precision (run-to-run stable in
+//   * fp32 (C_T == float, default): atomicAdd of each partial into a cached
+//     fp32 scratch, then one fp32 -> fp16 cast. No intermediate fp16
+//     rounding. The source calls HIP atomicAdd(float*), not
+//     v_global_atomic_add_f32. gfx1030 has no native global f32 atomic add
+//     (that instruction is gfx11), so the call compiles to a
+//     compare-and-swap loop. A gfx1100 compile of this helper may emit
+//     v_global_atomic_add_f32. One final cast means the result does not
+//     depend on atomic ordering at fp16 precision (run-to-run stable in
 //     practice).
-//   * fp16 CAS (C_T == 16-bit half, opt-in): packed fp16 CAS-64 atomic add
-//     into the caller's pre-zeroed fp16 output. Order-dependent at fp16
-//     precision.
+//   * fp16 CAS (C_T == 16-bit half, opt out with VLLM_RDNA2_MOE_FP32_ACCUM=0):
+//     packed fp16 CAS-64 atomic add into the caller's pre-zeroed fp16 output.
+//     Order-dependent at fp16 precision. gfx1030 has no native
+//     v_global_atomic_pk_add_f16.
 //   * cache fixed point (C_T == long long): native integer atomic add with
 //     32 fractional bits into caller-zeroed [rows,N+1] storage. The last
 //     column is a sticky row status. Nonfinite/out-of-bound partials set it
@@ -51,7 +56,9 @@ namespace vllm {
 namespace gptq_rdna2 {
 
 // Emit one token's 4 N-columns into accumulator row ``out``.
-//   C_T == float          -> native fp32 atomic add (no CAS) into the scratch.
+//   C_T == float          -> atomicAdd(float*). gfx1030: CAS loop (no native
+//                            v_global_atomic_add_f32). gfx1100: may be the
+//                            native instruction.
 //   C_T == long long      -> caller-owned fixed point with 32 fractional bits.
 //   C_T == 16-bit half    -> packed fp16 CAS-64 atomic add into the output.
 // The 16-bit case accepts any layout-compatible half (HIP ``half`` /

csrc/rocm/moe_q_gemm_rdna2.cu

@@ -18,9 +18,11 @@
 //      the upstream MoE path. gfx1030 has no WMMA ISA, so a prefill
 //      MoE kernel on gfx1030 would also be a scalar V_DOT2 design.
 //   2. Output accumulation supports two modes (see moe_accum_rdna2.cuh):
-//      the default fp32 path uses native v_global_atomic_add_f32 into a
-//      cached fp32 scratch and casts to fp16 once at the end; the opt-in
-//      fp16 path uses a packed atomic add emulated with
+//      the default fp32 path atomicAdds into a cached fp32 scratch and
+//      casts to fp16 once at the end. On gfx1030 that atomicAdd is a CAS
+//      loop (no v_global_atomic_add_f32; the instruction is gfx11, so a
+//      gfx1100 compile may use it). The opt-out fp16 path
+//      (VLLM_RDNA2_MOE_FP32_ACCUM=0) uses a packed atomic add emulated with
 //      global_atomic_cmpswap_b64 plus retry (gfx1030 has no native
 //      v_global_atomic_pk_add_f16).
 //
@@ -228,8 +230,9 @@ __global__ void moe_gemm_q4_kernel_rdna2(
   }
 
   // --- Epilogue: apply topk_weight, then accumulate. Shared with the W4A8
-  //     MoE kernel: C_T == float uses fp32 atomics into a scratch, a 16-bit
-  //     C_T uses the packed fp16 CAS into the pre-zeroed output. ---
+  //     MoE kernel: C_T == float atomicAdds into the fp32 scratch (CAS loop
+  //     on gfx1030; gfx1100 may use v_global_atomic_add_f32). A 16-bit C_T
+  //     uses the packed fp16 CAS into the pre-zeroed output. ---
   #pragma unroll
   for (int m = 0; m < BLOCK_SIZE_M; ++m) {
     int32_t token_id = sorted_token_ids[offset_m_base + m];
@@ -377,8 +380,10 @@ void dispatch_moe_gemm_q4(
 //   block_size_m           int (1, 2, 4, or 8)
 //   mul_topk_weight        bool
 //   fp32_accum             bool: accumulate partials in a cached fp32 scratch
-//                          (native fp32 atomics, no CAS) and round to fp16
-//                          once at the end. Default in _custom_ops.py.
+//                          via atomicAdd(float*) and round to fp16 once.
+//                          gfx1030 lowers that atomic to a CAS loop; gfx1100
+//                          may use v_global_atomic_add_f32. Default is on
+//                          (VLLM_RDNA2_MOE_FP32_ACCUM, _custom_ops.py).
 
 void moe_gptq_gemm_rdna2(torch::Tensor a, torch::Tensor c,
                          torch::Tensor b_q_weight, torch::Tensor b_scales,
@@ -445,8 +450,9 @@ void moe_gptq_gemm_rdna2(torch::Tensor a, torch::Tensor c,
           mul_topk_weight, (int)output_topk, use_v2_format ? 0 : 1, stream);
       C10_CUDA_KERNEL_LAUNCH_CHECK();
     } else if (fp32_accum) {
-      // fp32 accumulation: partials land in a cached fp32 scratch via the
-      // native v_global_atomic_add_f32 (no CAS), then one elementwise cast
+      // fp32 accumulation: partials land in a cached fp32 scratch via
+      // atomicAdd(float*) (a CAS loop on gfx1030, possibly native
+      // v_global_atomic_add_f32 on gfx1100), then one elementwise cast
       // rounds to fp16. The scratch pointer is stable across calls and the
       // per-call zero is a hipMemsetAsync, so this stays cudagraph-safe.
       TORCH_CHECK(c.is_contiguous(),

csrc/rocm/moe_w4a8_rdna2.cu

@@ -25,9 +25,10 @@
 //   clamp to tile 0/row 0 and are skipped by the epilogue.
 // * Epilogue is byte-identical to ``moe_gptq_gemm_rdna2`` (shared helper in
 //   ``moe_accum_rdna2.cuh``): router-weight multiply in fp32, then either the
-//   default fp32 accumulation (native fp32 atomics into a cached scratch +
-//   one fp32 -> fp16 cast) or the opt-in packed fp16 CAS-64 atomic add,
-//   with ``output_topk`` row reduce fused.
+//   default fp32 scratch (atomicAdd(float*): CAS loop on gfx1030, native
+//   v_global_atomic_add_f32 possible on gfx1100, then one fp32 -> fp16 cast)
+//   or the opt-out packed fp16 CAS-64 atomic add
+//   (VLLM_RDNA2_MOE_FP32_ACCUM=0), with ``output_topk`` row reduce fused.
 //
 // Weight format matches the W4A16 MoE kernel: [E, K/8, N] uint32 shuffled,
 // [E, groups, N] scales, [E, groups, N/8] packed zeros. ``use_v2_format``
@@ -525,8 +526,10 @@ void moe_w4a8_gemm_rdna2(torch::Tensor a, torch::Tensor c,
       static_cast<int>(output_topk)};
 
   if (fp32_accum) {
-    // fp32 accumulation: partials land in the cached fp32 scratch, then one
-    // elementwise cast rounds to fp16. Shared with the W4A16 MoE kernel.
+    // fp32 accumulation: partials land in the cached fp32 scratch via
+    // atomicAdd(float*) (CAS loop on gfx1030; gfx1100 may use native
+    // v_global_atomic_add_f32), then one elementwise cast rounds to fp16.
+    // Shared with the W4A16 MoE kernel.
     if (!c.is_contiguous()) {
       fallback();
       return;

docs/rdna2/moe-epilogue-accum.md

@@ -1,67 +1,75 @@
 # MoE epilogue accumulation mode for gfx1030 (`VLLM_RDNA2_MOE_FP32_ACCUM`)
 
 Branch: `rdna_extras`. Applies to the fused RDNA2 MoE kernels
-`moe_gptq_gemm_rdna2` (W4A16) and `moe_w4a8_gemm_rdna2` (W4A8) on gfx1030.
+`moe_gptq_gemm_rdna2` (W4A16) and `moe_w4a8_gemm_rdna2` (W4A8).
 The knob selects how the per-expert partial sums are accumulated into the
-pre-zeroed output `c`.
+output `c`.
 
 ## The two modes
 
-### Default OFF — legacy packed-fp16 CAS-64 (`fp32_accum=False`)
+### Default ON — fp32 scratch accumulator (`fp32_accum=True`)
 
-Each partial sum is accumulated with the native packed-fp16 CAS atomic
-(`atomic_add_pk4_f16`) directly into the fp16 output. This is the historical
-behaviour and the **production default**.
+Partial sums accumulate in an **fp32 scratch** via HIP `atomicAdd(float*)`,
+then a single fp32→fp16 cast. The source does not call
+`v_global_atomic_add_f32`.
 
-* **Faster** — measured +2.5 % (MTP=0) and **+12–13 %** (MTP=2) at c=1 vs the
-  fp32 path (see the numbers below).
-* **Order-dependent** — the packed-fp16 atomic add is not associative, so the
-  accumulation order across concurrent expert blocks can differ between runs,
-  which can flip the last 1–2 fp16 ULPs and, on close-valued cells, an argmax.
-  This is the run-to-run non-determinism the fp32 mode exists to remove.
-
-### Opt-in ON — fp32 scratch accumulator (`fp32_accum=True`)
-
-Partial sums accumulate in an **fp32 scratch** via native `global_atomic_add_f32`
-(no CAS), with a single fp32→fp16 cast at the end.
-
-* **Run-to-run deterministic in practice** — one rounding step at the end; the
-  residual fp32 atomic re-ordering is bounded to a few fp16 ULPs and does not
-  flip an argmax on any plausible threshold.
+* **gfx1030 lowers that atomic to a compare-and-swap loop.** RDNA2 has no
+  native global f32 atomic add. `v_global_atomic_add_f32` exists on gfx11;
+  a gfx1100 compile of the same helper may emit it. The device body of
+  these kernels is instantiated for gfx1030 (the gfx1100 image is the stub).
+* **One rounding step at the end.** There is no intermediate fp16 rounding.
+  Residual fp32 atomic reordering stays inside a few fp16 ULPs and does not
+  flip an argmax on any plausible threshold (run-to-run stable in practice).
 * **cudagraph-safe** — the fp32 scratch is a persistent per-(rows, n, device)
   allocation allocated eagerly before graph capture; the per-call zero is a
   capture-legal `hipMemsetAsync`.
-* **Costs a little at c=1** — see the numbers below; at c=8 the two modes are
-  at parity.
+* **May be slower than packed-fp16 CAS-64 on gfx1030**, because the fp32
+  atomic is itself a CAS loop. See the 2026-09-30 V620 numbers below.
+
+### Opt-out OFF — packed-fp16 CAS-64 (`fp32_accum=False`)
+
+Set `VLLM_RDNA2_MOE_FP32_ACCUM=0`. Each partial is accumulated with
+`atomic_add_pk4_f16` (a `global_atomic_cmpswap_b64` retry loop; gfx1030 has
+no native `v_global_atomic_pk_add_f16`) directly into the pre-zeroed fp16
+output.
+
+* **Faster on the 2026-09-30 V620 matrix** — +2.5 % (MTP=0) and **+12–13 %**
+  (MTP=2) at c=1 versus the fp32 path (see the numbers below).
+* **Order-dependent** — the packed-fp16 atomic add is not associative, so the
+  accumulation order across concurrent expert blocks can differ between runs,
+  which can flip the last 1–2 fp16 ULPs and, on close-valued cells, an argmax.
 
 ## Knobs
 
 | Knob | Value | Effect |
 |---|---|---|
-| env `VLLM_RDNA2_MOE_FP32_ACCUM` | unset / `0` | **Default:** CAS packed-fp16 epilogue. |
-| env `VLLM_RDNA2_MOE_FP32_ACCUM` | `1` | fp32 scratch accumulator. |
+| env `VLLM_RDNA2_MOE_FP32_ACCUM` | unset / `1` | **Default:** fp32 scratch accumulator. |
+| env `VLLM_RDNA2_MOE_FP32_ACCUM` | `0` | Packed-fp16 CAS-64 epilogue. |
 | per-call `fp32_accum=` (in `_custom_ops.py` wrapper) | `True` / `False` | Overrides the env-resolved default for a single call. |
 
-The wrapper reads the env var **once at import** into
-`_custom_ops._RDNA2_MOE_FP32_ACCUM` (no env access in traced/captured code) and
-passes the resolved `fp32_accum` per call. The per-call override wins over the
-env default. Both `moe_gptq_gemm_rdna2` and `moe_w4a8_gemm_rdna2` share the
-same flag.
+The pybind schemas for `moe_gptq_gemm_rdna2` and `moe_w4a8_gemm_rdna2` also
+default `fp32_accum=True`. Serving goes through the Python wrapper, which
+reads the env var **once at import** into `_custom_ops._RDNA2_MOE_FP32_ACCUM`
+(no env access in traced/captured code) and passes the resolved `fp32_accum`
+per call. The per-call override wins over the env default. Both kernels share
+the same flag. Recipes and `tools/rdna/serve_rdna.sh` do not set it, so a
+launch follows this default. Revert a launch with `VLLM_RDNA2_MOE_FP32_ACCUM=0`.
 
 ```bash
-# Production: CAS (default)
-export VLLM_RDNA2_MOE_FP32_ACCUM=0        # or leave unset
-
-# Reproducible debugging / A-B: fp32
+# Default: fp32 scratch (leave unset, or set explicitly)
 export VLLM_RDNA2_MOE_FP32_ACCUM=1
+
+# Opt out: packed-fp16 CAS-64
+export VLLM_RDNA2_MOE_FP32_ACCUM=0
 ```
 
 ## Measured trade-off
 
 Source: `bench_results/2026-09-30_fp32-accum-matrix/SUMMARY.md` (4× Radeon PRO
 V620, TP=4, FULL_AND_PIECEWISE cudagraphs, prefix caching, FA-RDNA2, frozen
 TunableOp rows, Flash-Next AWQ-W4A16). `Δ` = CAS relative to fp32, **positive =
-CAS faster**.
+CAS faster**. These cells predate the default flip; they are the cost of
+leaving fp32 on.
 
 | Arm / cell | out tok/s fp32 | out tok/s CAS | Δ out | prefill fp32→CAS | Δ prefill |
 |---|---:|---:|---:|---|---:|
@@ -70,30 +78,33 @@ CAS faster**.
 | MTP=2 16k/1k c=1 | 39.91 | 44.60 | **+11.8 %** | 1608.0→1684.7 | +4.8 % |
 | MTP=2 1k/512 c=1 | 66.71 | 75.65 | **+13.4 %** | 1571.7→1698.5 | +8.1 % |
 
-* The fp32 path is **not free**: CAS is 2.5 % faster at MTP=0 and 12–13 %
-  faster at MTP=2, all at c=1.
-* At c=8 the two modes are at parity (16k/1k c=8: 66.64 vs 66.08 out tok/s).
+* CAS was 2.5 % faster at MTP=0 and 12–13 % faster at MTP=2, all at c=1.
+* At c=8 the two modes were at parity (16k/1k c=8: 66.64 vs 66.08 out tok/s).
 * No correctness regression on either side — greedy outputs coherent on both
   (France→Paris, 2+2→4), no garbage flags, no PCI SERR.
+* Re-measure decode tok/s and worst-case ITL on V620 before treating those
+  deltas as the current tree. Revert with `VLLM_RDNA2_MOE_FP32_ACCUM=0`.
 
 ## When to use which
 
-* **CAS (default)** — production serving. Higher throughput, especially for
-  MTP-2 at low concurrency. Accept the tiny run-to-run output variation.
-* **fp32 (`=1`)** — reproducible debugging and A/B measurement where you need
-  bitwise-stable output across runs, or a long cudagraph session where
-  accumulated ULP noise must not drift into an argmax flip.
+* **fp32 (default)** — the production epilogue. One fp16 rounding at the end,
+  in exchange for the CAS-loop cost on gfx1030.
+* **CAS (`=0`)** — opt out when that throughput matters more than the last
+  fp16 ULPs, especially MTP-2 at low concurrency.
 
 ## Tests
 
+* `test_rdna2_moe_fp32_accum_defaults_on` (`tests/test_envs.py`) — unset env
+  resolves to fp32; `=0` resolves to CAS. No GPU.
 * `test_fp32_accum_run_to_run_stable` — fp32 is run-to-run stable within the
   fp16-noise budget.
 * `test_fp32_accum_rel_l2_vs_cas` — fp32 and CAS agree within 5 % rel-L2.
 * `test_fp32_accum_cudagraph_capture_stable` — fp32 scratch survives a graph
   capture-replay cycle.
-* `test_fp32_accum_default_off_byte_identical_to_cas` — the **default** (unset
-  env) resolves to CAS and is byte-identical to an explicit `fp32_accum=False`.
-* `test_fp32_accum_env_on_enables_fp32_scratch` — `=1` resolves the default to
-  the fp32 scratch path (run-to-run stable on a contention shape).
+* `test_fp32_accum_flag_off_byte_identical_to_cas` — with the module flag
+  forced off, `fp32_accum=None` is byte-identical to explicit CAS.
+* `test_fp32_accum_env_on_enables_fp32_scratch` — the module flag on (the
+  unset default) resolves `fp32_accum=None` to the fp32 scratch path
+  (run-to-run stable on a contention shape).
 
-All in `tests/kernels/quantization/test_rdna2_moe_w4a16.py`.
+GPU tests are in `tests/kernels/quantization/test_rdna2_moe_w4a16.py`.

docs/rdna2/tuning.md

@@ -377,7 +377,7 @@ QSA ring alignment requirements remain independent of graph sizes.
 
 | Env | Default | Notes |
 | --- | --- | --- |
-| `VLLM_RDNA2_MOE_FP32_ACCUM` | 0 | 0 = packed-fp16 CAS epilogue (faster, run-to-run ULP flips). 1 = fp32 scratch (deterministic, for debugging/A-B). |
+| `VLLM_RDNA2_MOE_FP32_ACCUM` | 1 | Default on: fp32 scratch, then one fp32→fp16 cast. On gfx1030 the fp32 atomic is a CAS loop (`v_global_atomic_add_f32` is gfx11; gfx1100 may use it). Better accuracy than packed-fp16 CAS-64; may be slower (2026-09-30 V620: CAS was +2.5% at MTP=0 and +12–13% at MTP=2, c=1). `=0` selects the packed-fp16 CAS epilogue. |
 | `VLLM_RDNA2_MOE_BSM_LARGE`, `_BSM_MIN_TOKENS` | 8, 512 | Prefill tile width (rows) and the token count where it kicks in. 0 keeps the decode tile. |
 | `VLLM_ROCM_MOE_SKINNY`, `_MAX_M` | 1, 8 | Wave-per-row skinny GEMV for **Triton WNA16** MoE weights only (max 16). Not used by the shuffled RDNA2 MoE path. |
 | `VLLM_RDNA_MOE_RESIDENT`, `_SKINNY` | 0 | Keep RDNA2 W4A16 MoE weights in the native shuffled layout. Disables W4A8 MoE. |

tests/kernels/quantization/test_rdna2_moe_w4a16.py

@@ -830,13 +830,14 @@ def reference():
 
 
 # ---------------------------------------------------------------------------
-# MoE epilogue fp32 accumulator (opt-in, VLLM_RDNA2_MOE_FP32_ACCUM=1; default
-# off = legacy packed-fp16 CAS). Two distinct contracts:
-#   1. run-to-run stable: same inputs in two back-to-back runs produce a
-#      bitwise identical output (no fp16 CAS reordering).
-#   2. fp32 vs CAS: fp32 accumulation and the legacy CAS path both produce
-#      a valid output; the rel-L2 divergence is bounded by the fp16
-#      accumulation error of the CAS path.
+# MoE epilogue accumulator. Unset VLLM_RDNA2_MOE_FP32_ACCUM is the fp32
+# scratch (tests/test_envs.py::test_rdna2_moe_fp32_accum_defaults_on). =0
+# selects packed-fp16 CAS-64. These GPU tests monkeypatch the import-time
+# snapshot `_RDNA2_MOE_FP32_ACCUM`. Two contracts:
+#   1. run-to-run stable: same inputs in two back-to-back fp32 runs produce
+#      a bitwise identical output (no fp16 CAS reordering).
+#   2. fp32 vs CAS: both paths produce a valid output; the rel-L2
+#      divergence is bounded by the fp16 accumulation error of the CAS path.
 # ---------------------------------------------------------------------------
 
 
@@ -1160,9 +1161,9 @@ def test_fp32_accum_cudagraph_capture_stable():
 
 
 @gfx1030_only
-def test_fp32_accum_default_off_byte_identical_to_cas(monkeypatch):
-    """Module default (VLLM_RDNA2_MOE_FP32_ACCUM unset) is the CAS epilogue:
-    the None-default must be byte-identical to an explicit fp32_accum=False.
+def test_fp32_accum_flag_off_byte_identical_to_cas(monkeypatch):
+    """With `_RDNA2_MOE_FP32_ACCUM` forced off (VLLM_RDNA2_MOE_FP32_ACCUM=0),
+    fp32_accum=None takes the CAS epilogue and matches explicit False.
 
     A single-K-block shape (K=256, grid.z=1) keeps the CAS order deterministic
     (one block writes each output element), so the two invocations are bitwise
@@ -1197,18 +1198,18 @@ def test_fp32_accum_default_off_byte_identical_to_cas(monkeypatch):
         output_topk=0,
     )
     assert torch.equal(out_default, out_cas), (
-        "fp32_accum=None (module default off) must be byte-identical to "
+        "fp32_accum=None with the module flag off must be byte-identical to "
         "explicit fp32_accum=False (CAS)"
     )
 
 
 @gfx1030_only
 def test_fp32_accum_env_on_enables_fp32_scratch(monkeypatch):
-    """VLLM_RDNA2_MOE_FP32_ACCUM=1 resolves the None-default to the fp32
-    scratch path: on a contention shape (K=2048, grid.z=8) the None-default is
-    run-to-run stable, which the order-dependent CAS epilogue is not. If the
-    None-default resolved to CAS instead, the second run would reorder and the
-    max diff would blow past the fp16-noise budget.
+    """Module flag True (the unset default, and =1) resolves fp32_accum=None
+    to the fp32 scratch: on a contention shape (K=2048, grid.z=8) the
+    None-default is run-to-run stable, which the order-dependent CAS
+    epilogue is not. If None resolved to CAS, the second run would reorder
+    and the max diff would blow past the fp16-noise budget.
     """
     monkeypatch.setattr(ops, "_RDNA2_MOE_FP32_ACCUM", True)
     E, K, N, top_k, group_size, block_size_m, M = 16, 2048, 512, 8, 32, 4, 4
@@ -1241,6 +1242,6 @@ def test_fp32_accum_env_on_enables_fp32_scratch(monkeypatch):
     assert torch.isfinite(out_a).all() and torch.isfinite(out_b).all()
     diff = (out_a.float() - out_b.float()).abs().max().item()
     assert diff <= 4.0 / 1024.0, (
-        f"None-default with VLLM_RDNA2_MOE_FP32_ACCUM=1 not run-to-run stable "
+        f"None-default with the fp32 module flag on not run-to-run stable "
         f"(max diff {diff} > 4 fp16 ULPs)"
     )