From c1309a0e1876122eb08b80f4b986b7636b9efe28 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Wed, 7 Oct 2026 02:39:57 +0900 Subject: [PATCH 1/4] update(rocm): enable fused add-RMSNorm and RoPE-append by default on ROCm FUSED_ADD_RMSNORM_DEFAULT and FUSED_ROPE_APPEND_DEFAULT become cfg!(feature = "rocm"): on in a ROCm build, off on Metal and CUDA. The HIP ports from #2063 are byte-identical to the graph and were never slower on gfx1151. The env vars still override in both directions. The Llama 3.1 rope_scaling bypass notice now fires only for an explicit truthy MLXCEL_FUSED_ROPE_APPEND, so default ROCm runs stay silent. New layers::fused_flag_explicit_value separates a setting from the default. Tests pin the per-build defaults and the notice decision; docs updated. Closes #2145 --- .../rocm-correctness-gfx1151-2026-09-30.md | 2 +- docs/benchmarks.md | 3 +- docs/environment-variables.md | 4 +- docs/installation.md | 2 +- src/lib/mlxcel-core/src/layers.rs | 128 +++++++++++++----- src/models/llama3.rs | 26 +++- src/models/llama3_tests.rs | 22 ++- 7 files changed, 141 insertions(+), 46 deletions(-) diff --git a/docs/benchmark_results/rocm-correctness-gfx1151-2026-09-30.md b/docs/benchmark_results/rocm-correctness-gfx1151-2026-09-30.md index ea57332ff..e21ddabfb 100644 --- a/docs/benchmark_results/rocm-correctness-gfx1151-2026-09-30.md +++ b/docs/benchmark_results/rocm-correctness-gfx1151-2026-09-30.md @@ -277,6 +277,6 @@ The Metal traces were produced with the loop in `benchmarks/logit_traces/rocm_gf - The second acceptance criterion of #1809 (decided-position mismatch reported per model and width against Metal) is now met for all nine models: fifteen of these sixteen pairs pass and `nemotron-3-nano-30b-a3b` `w1` is inconclusive, with the same checkpoint's `w8` and `w256` pairs passing. The two `w1ctx512` rows, which cover the fused SSM update kernel on Metal, also pass. - Kernel against kernel for the SSM update: the `w1ctx512` rows here compare Metal's kernel with ROCm's graph. The HIP port (#2067) added kernel-against-kernel rows, with zero decided-position mismatches for both hybrids, in [rocm-ssm-update-kernel-gfx1151-2026-10-04.md](rocm-ssm-update-kernel-gfx1151-2026-10-04.md#model-logits-kernel-against-kernel). - Two MoE implementations in the Nemotron-H pairs: since the HIP fused MoE ports (#2065) ROCm's gate reads `moe_down_kernel_available()` and takes `fused_moe_forward`, as Metal does. New Nemotron-H traces on that path pass against these Metal traces at `w8` and `w1ctx512` with zero decided-position mismatches, in [rocm-fused-moe-gfx1151-2026-10-05.md](rocm-fused-moe-gfx1151-2026-10-05.md#model-logits). -- The opt-in fused add-RMSNorm and RoPE + append kernels: these traces ran with both off, the default. Their HIP ports (#2063) are byte-identical to the graph they replace: Llama 3.1 and Qwen2.5 traces with them on match the traces with them off byte for byte at `w1`, `w8` and `w1ctx512`, so for the models that call them, turning them on does not change ROCm's side of these comparisons ([rocm-fused-norm-rope-gfx1151-2026-10-05.md](rocm-fused-norm-rope-gfx1151-2026-10-05.md)). +- The fused add-RMSNorm and RoPE + append kernels: these traces ran with both off, the default at the time. They are on by default on ROCm since lablup/mlxcel#2145. Their HIP ports (#2063) are byte-identical to the graph they replace: Llama 3.1 and Qwen2.5 traces with them on match the traces with them off byte for byte at `w1`, `w8` and `w1ctx512`, so for the models that call them, turning them on does not change ROCm's side of these comparisons ([rocm-fused-norm-rope-gfx1151-2026-10-05.md](rocm-fused-norm-rope-gfx1151-2026-10-05.md)). - `src/models/nemotron_h.rs` prints its loading messages to stdout, which puts five non-trace lines into every Nemotron-H trace and stops `compare_logit_traces.py`. Either the loader should log to stderr or the comparison script should skip them. - No noise floor was measured; the first run's reasoning for not needing one still applies. diff --git a/docs/benchmarks.md b/docs/benchmarks.md index 6249afc71..30c56e4bb 100644 --- a/docs/benchmarks.md +++ b/docs/benchmarks.md @@ -341,7 +341,8 @@ A speedup below 1.00 for either op is the signal to flip `FUSED_ADD_RMSNORM_DEFAULT` or `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs`, which leaves the kernel available and opt-in through `MLXCEL_FUSED_ADD_RMSNORM=1` / `MLXCEL_FUSED_ROPE_APPEND=1` -instead of removing it. +instead of removing it. Both constants are `cfg!(feature = "rocm")`: off on +Metal and CUDA builds, on in a ROCm build (#2145). The op-level number is a lower bound on the end-to-end effect. Both fusions also remove a full-width intermediate from the MLX graph per call, which shows up as diff --git a/docs/environment-variables.md b/docs/environment-variables.md index 94ef7ef07..cadf80fc1 100644 --- a/docs/environment-variables.md +++ b/docs/environment-variables.md @@ -571,8 +571,8 @@ recommended as normal deployment settings. | `MLXCEL_SSM_KERNEL` | `0` disables; any other value or unset enables | on | Gates the fused single-token SSM decode kernel (`ssm_update_kernel`) on every backend that has a port: Metal, CUDA (issue #631) and ROCm (issue #2067). It replaces the ~55-op SSD scan graph (`ssm_step`) with one launch per SSM layer per decode step for the hybrid SSM models (granite-4.0-h, falcon-h1, plamo-2, nemotron-h). `ssm_kernel_available()` answers from the kernel's port table (`ssm_ports()`), so a backend with no port takes the graph path on its own. It is also false when the default device is the CPU (`MLXCEL_DEVICE=cpu`). Set to `0` to force the graph path (A/B, rollback). Measured on GB10 (CUDA): 2.6-4.5x decode on the granite/falcon family, see `docs/benchmark_results/hybrid-ssm-decode-cuda-kernel-gb10-2026-07-10.md`; on gfx1151 (ROCm): see `docs/benchmark_results/rocm-ssm-update-kernel-gfx1151-2026-10-04.md`. | | `MLXCEL_SSM_CUDA_KERNEL` | `0` disables; any other value or unset enables | on | Older name of `MLXCEL_SSM_KERNEL`, from when the CUDA port was the only one off Apple, kept as an alias so existing scripts still work. Either variable set to `0` forces the graph path, now on every backend (Metal and ROCm included). | | `MLXCEL_FUSED_QK_NORM` | `1`/`true`/`on`/`yes` enable; any other value or unset disables | off (opt-in) | Fused single-token QKV projection + Q/K RMSNorm + RoPE kernel (#326) for Qwen3 and Qwen3-MoE decode. Opt-in: set to `1` to enable. Matches the graph path within RMS < 5e-3 (the reduction is over the transpose-invariant head_dim axis), but greedy temp-0 is not byte-identical over long generation; on CUDA the graph path is itself non-deterministic run-to-run from GPU FP-reduction order, while the fused path is deterministic, so its output stays inside the graph baseline's own envelope. The kernel cuts Rust/C++ FFI crossings rather than MLX op count, so it does not speed up the GPU/bandwidth-bound decode loop: on M1 Ultra it measured 1 to 3.4% slower (qwen3-0.6b 275 vs 284, qwen3-8b 82.3 vs 83.2 tok/s); on GB10/CUDA (SM 12.1) it is also slower (qwen3-0.6b 0.96x, qwen3-8b ~1.0x, qwen3-30b-a3b 0.92x fused/graph; see `docs/benchmark_results/fused-qk-norm-decode-gb10.md`), so there is no per-backend win. Ships as a reusable shared primitive for the deferred QK-norm families and stays opt-in (default off) on every measured backend (M1 Ultra, M5 Max, GB10/CUDA), mirroring the opt-in `MLXCEL_FUSED_MOE_RELU2`. Active only when `l == 1` (decode) and weights are quantized. | -| `MLXCEL_FUSED_ADD_RMSNORM` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (opt-in) | Fused residual-add + RMSNorm decode kernel (#905). Computes `new_residual = x + residual` and `normed = rms_norm(new_residual) * (weight_bias + weight)` in one dispatch instead of an elementwise `Add` followed by `fast::rms_norm`, removing one kernel launch and one full-width intermediate at every pre-norm residual join. Adopted by the Llama3-family block (which also serves Qwen2/Qwen2.5 and every VLM whose text backbone is `Llama3Model`/`Qwen2Model`) and by Gemma, where the `(1 + w)` convention rides in as a scalar `weight_bias = 1.0` folded in the weight's own dtype rather than as a second weight tensor. Numerics are pinned to MLX's own `rms_norm` kernel (fp32 accumulation over the dtype-rounded sum, `precise::rsqrt`, the `x * inv_mean` product rounded to the activation dtype before the weight multiply), so the fused and unfused paths agree to within a small multiple of one ulp; greedy temp-0 is not guaranteed byte-identical over long generation, because a near-tie argmax can land on either side of that difference. Set to `0` to force the `add` + `fast_rms_norm` pair (A/B, rollback). The compiled-in default lives in `FUSED_ADD_RMSNORM_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off)**: the op-level microbench on M1 Ultra (Metal, f16, hidden {2048, 4096, 8192} x batch {1, 4, 8}, three repetitions) measured the fused path at parity, speedups scattering 0.77x to 1.14x with no structure and the extremes not reproducing. Per-iteration time stayed near 280-410us across a 4x change in hidden size, so the harness is dominated by dispatch cost rather than the op. Issue #905's measure-then-keep policy lands a fusion unwired absent a demonstrated win, matching how `MLXCEL_FUSED_QK_NORM` (#326) shipped. The production decode path agrees: on M1 Ultra at pp512/tg128, Llama 3.1 8B decodes at 113.18 tok/s off and 113.01 on, Qwen2.5 7B at 115.38 and 115.34, with prefill unchanged. Set to `1` to opt in. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-add-rmsnorm-decode-m1ultra-2026-09-27.md`. On ROCm the kernel runs as a HIP port since lablup/mlxcel#2063 whose output is byte-identical to the `add` + `fast_rms_norm` graph (Llama 3.1 traces with it on and off match), and it stays off there too: Llama 3.1 8B decode on gfx1151 measured 37.85 off and 37.95 tok/s on, medians of five ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). -| `MLXCEL_FUSED_ROPE_APPEND` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (opt-in) | Fused q/k RoPE + KV-append-layout decode kernel (#905). Reads the row-contiguous fused-QKV projection output whole, applies rotary embedding to the q and k blocks, and emits q, k and v already in their consumers' layouts, replacing three trailing-axis slices, three reshape/transpose pairs and two `fast::rope` calls with one dispatch. The rotation is transcribed from MLX's own `rope.metal`, so it agrees with `fast::rope` including the `traditional` convention and partial `rope_dims`. The trig call is the one thing chosen per backend rather than transcribed: each body calls whatever MLX's own RoPE kernel for that backend calls, so the Metal body uses `metal::fast::cos`/`sin` after `rope.metal` and the CUDA body uses libdevice `cosf`/`sinf` after `rope.cu`. Using the CUDA fast intrinsics `__cosf`/`__sinf` there instead cost about 1.2e-2 rad of applied phase at absolute position 131071 and failed the parity suite at the two large-offset cases (issue #1049). Adopted by the Llama3-family dense decode path, except for a checkpoint whose `rope_scaling` block selects a frequency table: the kernel derives its frequencies from `rope_base` and has no parameter for a table, so such a model takes the graph path (#1355), and a one-time notice on stderr says so when this variable asked for the kernel. `traditional` and the position scale are real parameters and are honored, so neither triggers that notice for this kernel. The paged block-pool destination layout the kernel also supports is implemented and tested but not wired, because the batched paged decode path belongs to issue #899. The kernel emits the append payload rather than writing the cache slab, because MLX custom-kernel outputs are always freshly allocated while the existing `slice_update` append donates the slab buffer, so an in-kernel write would trade an O(1) donated update for an O(capacity) copy. Set to `0` to force the reshape/transpose/`fast_rope` graph (A/B, rollback). The compiled-in default lives in `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off)** for the same measured reason as `MLXCEL_FUSED_ADD_RMSNORM`, with one cell consistently below parity rather than merely noisy: hidden 8192 at batch 1 measured 0.94x, 0.90x and 0.89x across three repetitions. Set to `1` to opt in. On GB10/CUDA the same harness measures 1.02x to 2.06x with every cell at or above parity, a different picture from Apple Silicon that has not been acted on; flipping the default is its own decision with its own evidence bar. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-rope-cuda-trig-gb10-2026-08-06.md`. On ROCm the kernel runs as a HIP port since lablup/mlxcel#2063 that calls `sincosf` as `rope.hip` does and reproduces the graph's FMA forms and f16 rounding, so its output is byte-identical to the graph; it stays off there too, because Qwen2.5 7B decode on gfx1151 gained 0.6% with it alone and 1.7% with both fusions (medians of seven), a gain whose size drifted between runs as much as the off arm's own spread ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). +| `MLXCEL_FUSED_ADD_RMSNORM` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (opt-in) | Fused residual-add + RMSNorm decode kernel (#905). Computes `new_residual = x + residual` and `normed = rms_norm(new_residual) * (weight_bias + weight)` in one dispatch instead of an elementwise `Add` followed by `fast::rms_norm`, removing one kernel launch and one full-width intermediate at every pre-norm residual join. Adopted by the Llama3-family block (which also serves Qwen2/Qwen2.5 and every VLM whose text backbone is `Llama3Model`/`Qwen2Model`) and by Gemma, where the `(1 + w)` convention rides in as a scalar `weight_bias = 1.0` folded in the weight's own dtype rather than as a second weight tensor. Numerics are pinned to MLX's own `rms_norm` kernel (fp32 accumulation over the dtype-rounded sum, `precise::rsqrt`, the `x * inv_mean` product rounded to the activation dtype before the weight multiply), so the fused and unfused paths agree to within a small multiple of one ulp; greedy temp-0 is not guaranteed byte-identical over long generation, because a near-tie argmax can land on either side of that difference. Set to `0` to force the `add` + `fast_rms_norm` pair (A/B, rollback). The compiled-in default lives in `FUSED_ADD_RMSNORM_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` and depends on the build: `cfg!(feature = "rocm")`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off) on Metal and CUDA**: the op-level microbench on M1 Ultra (Metal, f16, hidden {2048, 4096, 8192} x batch {1, 4, 8}, three repetitions) measured the fused path at parity, speedups scattering 0.77x to 1.14x with no structure and the extremes not reproducing. Per-iteration time stayed near 280-410us across a 4x change in hidden size, so the harness is dominated by dispatch cost rather than the op. Issue #905's measure-then-keep policy lands a fusion unwired absent a demonstrated win, matching how `MLXCEL_FUSED_QK_NORM` (#326) shipped. The production decode path agrees: on M1 Ultra at pp512/tg128, Llama 3.1 8B decodes at 113.18 tok/s off and 113.01 on, Qwen2.5 7B at 115.38 and 115.34, with prefill unchanged. Set to `1` to opt in. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-add-rmsnorm-decode-m1ultra-2026-09-27.md`. On ROCm it is **on by default** (lablup/mlxcel#2145): the kernel runs as a HIP port (lablup/mlxcel#2063) whose output is byte-identical to the `add` + `fast_rms_norm` graph (Llama 3.1 traces with it on and off match), and on gfx1151 it was never slower in a median (Llama 3.1 8B decode 37.85 tok/s off and 37.95 on, medians of five; Qwen2.5 7B +1.5%), so the change costs no accuracy and the maintainer turned it on there. Set to `0` on ROCm to opt out. Gemma and IQuest Loop Coder take the same kernel by default on ROCm but have no checkpoint measured on that backend; their `(1 + w)` convention is covered by the parity tests only ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). +| `MLXCEL_FUSED_ROPE_APPEND` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (opt-in) | Fused q/k RoPE + KV-append-layout decode kernel (#905). Reads the row-contiguous fused-QKV projection output whole, applies rotary embedding to the q and k blocks, and emits q, k and v already in their consumers' layouts, replacing three trailing-axis slices, three reshape/transpose pairs and two `fast::rope` calls with one dispatch. The rotation is transcribed from MLX's own `rope.metal`, so it agrees with `fast::rope` including the `traditional` convention and partial `rope_dims`. The trig call is the one thing chosen per backend rather than transcribed: each body calls whatever MLX's own RoPE kernel for that backend calls, so the Metal body uses `metal::fast::cos`/`sin` after `rope.metal` and the CUDA body uses libdevice `cosf`/`sinf` after `rope.cu`. Using the CUDA fast intrinsics `__cosf`/`__sinf` there instead cost about 1.2e-2 rad of applied phase at absolute position 131071 and failed the parity suite at the two large-offset cases (issue #1049). Adopted by the Llama3-family dense decode path, except for a checkpoint whose `rope_scaling` block selects a frequency table: the kernel derives its frequencies from `rope_base` and has no parameter for a table, so such a model takes the graph path (#1355), and a one-time notice on stderr says so when this variable is explicitly set to a truthy value (the ROCm default alone does not print it). `traditional` and the position scale are real parameters and are honored, so neither triggers that notice for this kernel. The paged block-pool destination layout the kernel also supports is implemented and tested but not wired, because the batched paged decode path belongs to issue #899. The kernel emits the append payload rather than writing the cache slab, because MLX custom-kernel outputs are always freshly allocated while the existing `slice_update` append donates the slab buffer, so an in-kernel write would trade an O(1) donated update for an O(capacity) copy. Set to `0` to force the reshape/transpose/`fast_rope` graph (A/B, rollback). The compiled-in default lives in `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` and depends on the build: `cfg!(feature = "rocm")`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off) on Metal and CUDA** for the same measured reason as `MLXCEL_FUSED_ADD_RMSNORM`, with one cell consistently below parity rather than merely noisy: hidden 8192 at batch 1 measured 0.94x, 0.90x and 0.89x across three repetitions. Set to `1` to opt in. On GB10/CUDA the same harness measures 1.02x to 2.06x with every cell at or above parity, a different picture from Apple Silicon that has not been acted on; flipping the default is its own decision with its own evidence bar. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-rope-cuda-trig-gb10-2026-08-06.md`. On ROCm it is **on by default** (lablup/mlxcel#2145): the kernel runs as a HIP port (lablup/mlxcel#2063) that calls `sincosf` as `rope.hip` does and reproduces the graph's FMA forms and f16 rounding, so its output is byte-identical to the graph, and Qwen2.5 7B decode on gfx1151 gained 0.6% with it alone and 1.7% with both fusions (medians of seven), inside run-to-run drift but never slower. Set to `0` on ROCm to opt out. Llama 3.1 still takes the graph path there because of its `rope_scaling` table; the stderr notice about that fires only when this variable is set to a truthy value, not for the ROCm default ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). | `MLXCEL_COMPILED_QGELU_MLP` | `0`/`false`/`off`/`no` disable; any other value or unset enables | on | Compiles the affine-quantized GeGLU MLP (gate/up/gelu/down) into one `mx::compile` graph so the tanh-approx GELU's ~14 element-wise ops per layer collapse into a single fused `Compiled` primitive. Covers the Gemma family (gemma/gemma2/gemma3/gemma4). The group_size=64/bits=4 case was already compiled on every shape and is unchanged; other affine quantizations (notably the group_size=64/bits=8 MLP weights in Gemma 4 mixed-precision checkpoints, issue #680) are compiled only on the single-token decode call, because compiling the 8-bit prefill GEMM measured 8-9% slower and +0.2-0.7 GB peak on GB10 (the shapeless fused graph forces a decode-oriented qmm kernel onto the large prefill matmul). GB10 gemma-4-12b decode is weight-bandwidth-bound (~94% GPU-busy), so this primitive-count cut is throughput-neutral there (measured <1% from a ~530-primitive/step reduction); it still removes real CPU dispatch and helps op-count-bound backends. Set to `0` to force the op-at-a-time fallback (A/B + rollback). It also gates the NVFP4 scaled fused MLP variant (`MLXCEL_DISABLE_FUSED_GLOBAL_SCALE`), forcing its eager fold when set to `0`. | | `MLXCEL_DISABLE_FUSED_GLOBAL_SCALE` | `1`/`true`/`on`/`yes` (case-insensitive) disable; unset or any other value keeps the fold | off (fold on) | Rolls back the NVFP4 global-scale fold for Gemma 4 (issues #698/#705). By default the fused MLP and per-layer-input-gate C++ paths fold each per-projection `weight_scale_2` sidecar (from the direct ModelOpt transcode, issue #693/#697) into the fused kernel at the mathematically correct points: the gate scale before the GeGLU activation, the up scale on the up product, and the down scale on the fused output, each reproducing `apply_global_scale` byte-for-byte. Native NVFP4 prefill now uses a shape-specific scaled MLP graph, and standalone `UnifiedLinear` sidecar projections can apply qmm + global scale + dense bias through one C++ helper. When set, sidecar-carrying paths fall back to the op-at-a-time `UnifiedLinear::forward` scalar application (the pre-#698 bypass). Greedy temp-0 decode is token-identical across the two paths on `gemma-4-31b-it-nvfp4`; the fold removes element-wise dispatches on the gemma-4 path where CUDA graphs are disabled (#688). See `docs/benchmark_results/nvfp4-direct-transcode-gb10-2026-07-08.md` and `docs/benchmark_results/nvfp4-native-prefill-m1ultra-2026-07-09.md`. | | `MLXCEL_ONE_BIT_KERNEL` | `0`/`false`/`off`/`no` (case-insensitive) disable; unset or any other value keeps the kernels | on (Metal) | Affine 1-bit checkpoints (`quantization.bits == 1`, issue #1370). MLX has no 1-bit kernel, so the `quantized_matmul`, `quantized_linear_forward`, `dequantize` and `quantized_embedding` bridge primitives route a 1-bit affine triple to mlxcel's own fused Metal qmv (fewer than 16 rows) or tiled qmm kernel. Set to `0` to force the dequantize graph (`x @ dequant(W)^T`) instead, which is the parity oracle for the kernels and the path CUDA and ROCm always take. Read once in C++ at first use and cached for the process lifetime. | diff --git a/docs/installation.md b/docs/installation.md index 770994a64..daf1b1e50 100644 --- a/docs/installation.md +++ b/docs/installation.md @@ -555,7 +555,7 @@ so `mlxcel inspect` and `--estimate-memory` give the same answers as before. | NVFP4 checkpoints | No native kernel. ModelOpt NVFP4 checkpoints (for example the Gemma 4 NVFP4 exports) are converted to affine 4-bit at load with no environment variables, and the load log names the route and the reason. A layer that cannot be converted fails the load with its name and the reason. MLX-native NVFP4 exports (`"mode": "nvfp4"` in `config.json`, such as `mlx-community/*-nvfp4`) have no load-time conversion and are refused at load; use an affine export instead (lablup/mlxcel#1806). | | Affine MoE models (for example Qwen3-30B-A3B) | Run natively. Single-token decode takes the fused two-kernel MoE path (gate-up, then down) through its HIP port (lablup/mlxcel#2065), as on Metal and CUDA: Qwen3-30B-A3B-4bit decode on `gfx1151` went from 61.3 to 62.5 tok/s, with logits matching Metal's fused kernel at every decided position ([results](benchmark_results/rocm-fused-moe-gfx1151-2026-10-05.md)); `MLXCEL_FUSED_MOE=0` selects the `gather_qmm` graph path instead. Experts wider than `MLXCEL_FUSED_MOE_MAX_DFF` (8192 by default, as on CUDA; Mixtral-8x7B's 14336) stay on `gather_qmm`. | | Model-specific fused kernels (Apertus xIELU, Cohere2 add3 + LayerNorm, Mamba1 selective scan, Nemotron-H's opt-in `MLXCEL_FUSED_MOE_RELU2`) | Run as HIP ports (lablup/mlxcel#2069). xIELU and the add3 LayerNorm are byte-identical to ROCm's own unfused graph; the Mamba1 scan keeps its state in f32 as on Metal; the relu2 path is performance-neutral on Nemotron-H and stays opt-in. Apertus, Cohere2 and the Mamba1 families have no checkpoint on the tested host, so they are covered by kernel tests only ([results](benchmark_results/rocm-metal-only-ports-gfx1151-2026-10-05.md)). | -| Fused residual-add RMSNorm and RoPE + KV append (`MLXCEL_FUSED_ADD_RMSNORM`, `MLXCEL_FUSED_ROPE_APPEND`) | Off by default, as on Metal and CUDA. Opted in, both run as HIP ports (lablup/mlxcel#2063) whose output is byte-identical to the graph they replace, so turning them on changes no logits; decode on `gfx1151` moved by +0.3% on Llama 3.1 8B and +1.7% on Qwen2.5 7B, not clearly above the noise, so the defaults stay off ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). | +| Fused residual-add RMSNorm and RoPE + KV append (`MLXCEL_FUSED_ADD_RMSNORM`, `MLXCEL_FUSED_ROPE_APPEND`) | On by default on ROCm (lablup/mlxcel#2145), off on Metal and CUDA. Both run as HIP ports (lablup/mlxcel#2063) whose output is byte-identical to the graph they replace, so they change no logits; decode on `gfx1151` was never slower with them (+0.3% on Llama 3.1 8B, +1.7% on Qwen2.5 7B, inside run-to-run drift). Set either variable to `0` to fall back to the graph. Llama 3.1 keeps the graph RoPE because of its `rope_scaling` table; Gemma and IQuest Loop Coder also take the norm port but have no checkpoint measured on ROCm ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). | | Paged attention (the server's batched paged decode, MLA split-KV, sparse paged decode) | The v1 decode, v2 partial and merge kernels run as HIP ports (lablup/mlxcel#2068), so these paths take the fused kernels instead of gather-then-SDPA, as on Metal and CUDA; the kernels' tests match the gather and host references on `gfx1151`. Server decode on `gfx1151` with Meta-Llama-3.1-8B-Instruct-4bit, against the gather path on the same binary: 1.28x per-request decode at batch 4 and ~16K tokens, 1.40x and 1.77x for one sequence at ~16K and ~32K, no measurable change at batch 4 and ~4K ([results](benchmark_results/rocm-paged-attention-gfx1151-2026-10-05.md)). `MLXCEL_PAGED_ATTENTION_NATIVE=0` forces the gather path. | | Sampled decode (`temperature > 0`) | The fused Gumbel-max sampler (no filter) and the dual-pivot rejection sampler (top-p, and the other filter combinations its routing admits) run as HIP kernels, as on Metal and CUDA (lablup/mlxcel#2064). On `gfx1151` with Meta-Llama-3.1-8B-Instruct-4bit, decode at `--temperature 0.8 --top-p 0.95` went from 33.98 to 37.56 tok/s (medians of three; the before runs drifted from 36.26 down to 32.81, the after runs held 37.06 to 37.82); at `--temperature 0.8` alone the change is inside the noise ([results](benchmark_results/rocm-samplers-gfx1151-2026-10-05.md)). `MLXCEL_SAMPLING_GUMBEL=0` and `MLXCEL_SAMPLING_REJECTION=0` force the graph. | | Hybrid SSM decode (granite-4.0-h, falcon-h1, plamo-2, Nemotron-H) | The single-token Mamba2 SSM update runs the fused HIP kernel instead of the ~55-op SSD graph (lablup/mlxcel#2067): decode on `gfx1151` went from 60.4 to 88.4 tok/s on granite-4.0-h-tiny-4bit and from 51.4 to 74.4 tok/s on Nemotron-3-Nano-30B-A3B-4bit, with model logits matching the graph and Metal's kernel at every decided position ([results](benchmark_results/rocm-ssm-update-kernel-gfx1151-2026-10-04.md)). Measured on granite and Nemotron-H; falcon-h1 and plamo-2 take the same gate but have not been run on ROCm. `MLXCEL_SSM_KERNEL=0` forces the graph. | diff --git a/src/lib/mlxcel-core/src/layers.rs b/src/lib/mlxcel-core/src/layers.rs index 60db577ff..de07c0040 100644 --- a/src/lib/mlxcel-core/src/layers.rs +++ b/src/lib/mlxcel-core/src/layers.rs @@ -816,11 +816,17 @@ fn dense_gemm( /// Default for the fused residual-add + RMSNorm decode path. /// /// **This is the one place to flip if the measurement does not justify the -/// fusion.** `MLXCEL_FUSED_ADD_RMSNORM=0` disables it at runtime without a -/// rebuild; setting this constant to `false` makes off the default and -/// `MLXCEL_FUSED_ADD_RMSNORM=1` the opt-in. +/// fusion.** The default is per build: off on Metal and CUDA builds, on in a +/// `rocm` build. `MLXCEL_FUSED_ADD_RMSNORM=0` disables it at runtime without a +/// rebuild and `MLXCEL_FUSED_ADD_RMSNORM=1` enables it, on every backend. /// -/// Default-OFF: measured, and the measurement did not justify wiring it on. +/// The feature flag is the backend here, the same way PR #2098 set the ROCm +/// `MLXCEL_FUSED_MOE_SGY` default under `MLXCEL_BRIDGE_ROCM_BACKEND`: a `rocm` +/// build compiles no Metal or CUDA backend, so no runtime backend comparison is +/// needed (and `scripts/ci/check_kernel_port_dispatch.py` rejects one). +/// +/// Metal and CUDA: default-OFF. Measured, and the measurement did not justify +/// wiring it on. /// /// Op-level microbench on Apple M1 Ultra (Metal, f16, hidden {2048, 4096, /// 8192} x batch {1, 4, 8}, three repetitions) put the fused path at roughly @@ -846,29 +852,35 @@ fn dense_gemm( /// See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and /// `docs/benchmark_results/fused-add-rmsnorm-decode-m1ultra-2026-09-27.md`. /// -/// ROCm (#2063) measured the same: with the HIP port, whose output is -/// byte-identical to the graph, Llama 3.1 8B decode on gfx1151 was 37.85 tok/s -/// off and 37.95 on (medians of five), and Qwen2.5 7B gained 1.5% with a size -/// that drifted between runs as much as the off arm's spread. Not a clear win, -/// so ROCm keeps the shared default. See -/// `docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md`. -pub(crate) const FUSED_ADD_RMSNORM_DEFAULT: bool = false; +/// ROCm: default-ON (#2145). The HIP port (#2063, PR #2107) is byte-identical +/// to the ROCm graph it replaces (every on/off logit-trace pair identical), and +/// on gfx1151 it was never slower in a median: Llama 3.1 8B decode 37.85 tok/s +/// off and 37.95 on (medians of five), Qwen2.5 7B +1.5% alone and +1.7% with +/// the RoPE fusion. The gains sit inside run-to-run drift, but the change costs +/// no accuracy, so the maintainer turned both fusions on for ROCm. See +/// `docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md`, which +/// also has the re-measurement of the default-on build. +pub(crate) const FUSED_ADD_RMSNORM_DEFAULT: bool = cfg!(feature = "rocm"); /// Default for the fused q/k RoPE + KV-append-layout decode path. Same -/// flip-here contract and the same measured outcome as -/// [`FUSED_ADD_RMSNORM_DEFAULT`], so it also ships opt-in via -/// `MLXCEL_FUSED_ROPE_APPEND=1`. -/// -/// One cell was consistently below parity rather than merely noisy: hidden -/// 8192 at batch 1 measured 0.94x, 0.90x and 0.89x across three repetitions. -/// That is the single reproducible signal in the sweep and it points the wrong -/// way, which is the stronger reason to leave this unwired until a backend or -/// shape is found where it wins. -/// -/// ROCm (#2063) did not find one: on gfx1151 Qwen2.5 7B decode gained 0.6% with -/// this kernel alone and 1.7% with both fusions (medians of seven), inside the -/// drift between runs. Same results page as [`FUSED_ADD_RMSNORM_DEFAULT`]. -pub(crate) const FUSED_ROPE_APPEND_DEFAULT: bool = false; +/// flip-here contract and the same per-build default as +/// [`FUSED_ADD_RMSNORM_DEFAULT`]: off on Metal and CUDA (opt in with +/// `MLXCEL_FUSED_ROPE_APPEND=1`), on in a `rocm` build (opt out with `=0`). +/// +/// Metal: one cell was consistently below parity rather than merely noisy: +/// hidden 8192 at batch 1 measured 0.94x, 0.90x and 0.89x across three +/// repetitions. That is the single reproducible signal in the sweep and it +/// points the wrong way, which is the stronger reason to leave this unwired +/// there until a shape is found where it wins. +/// +/// ROCm (#2145): the HIP port (#2063) is byte-identical to the graph, and on +/// gfx1151 Qwen2.5 7B decode gained 0.6% with this kernel alone and 1.7% with +/// both fusions (medians of seven), inside the drift between runs but never +/// slower, so it is on by default with the norm fusion. Llama 3.1's +/// `rope_scaling` frequency table still routes around it (see +/// `report_fused_rope_bypass_once` in `src/models/llama3.rs`). Same results +/// page as [`FUSED_ADD_RMSNORM_DEFAULT`]. +pub(crate) const FUSED_ROPE_APPEND_DEFAULT: bool = cfg!(feature = "rocm"); /// Whether the fused residual-add + RMSNorm path (#905) is enabled. /// @@ -880,7 +892,9 @@ pub(crate) const FUSED_ROPE_APPEND_DEFAULT: bool = false; /// over long generation. The two differ only by the rounding of the residual /// sum feeding the sum of squares (see `fused_norm.cpp`), which is far below /// the argmax-flip scale, but a near-tie argmax can still land on the other -/// side. That is a tie-break difference, not a regression. +/// side. That is a tie-break difference, not a regression. The ROCm HIP port +/// (#2063) measured byte-identical to the ROCm graph instead: every on/off +/// logit-trace pair in its results page matched. pub fn fused_add_rmsnorm_enabled() -> bool { static ENABLED: std::sync::OnceLock = std::sync::OnceLock::new(); *ENABLED.get_or_init(|| { @@ -906,13 +920,22 @@ pub fn fused_rope_append_enabled() -> bool { /// value takes `default` rather than silently disabling: a typo in a deployment /// script should not quietly change the decode graph. fn fused_flag_enabled_from(value: Option<&str>, default: bool) -> bool { - match value { - Some(v) => match v.trim().to_ascii_lowercase().as_str() { - "0" | "false" | "off" | "no" => false, - "1" | "true" | "on" | "yes" => true, - _ => default, - }, - None => default, + fused_flag_explicit_value(value).unwrap_or(default) +} + +/// The explicit setting a #905 kill-switch value spells, if any: `Some(false)` +/// for `0`/`false`/`off`/`no`, `Some(true)` for `1`/`true`/`on`/`yes` +/// (case-insensitive, trimmed), and `None` when the variable is unset or the +/// value is unrecognised, which is when the compiled-in default applies. +/// +/// Public so a caller can tell "the user asked for this" apart from "the build +/// defaults to this": the Llama 3.1 RoPE bypass notice (#2145) must fire only +/// for the former, now that the ROCm default is on. +pub fn fused_flag_explicit_value(value: Option<&str>) -> Option { + match value?.trim().to_ascii_lowercase().as_str() { + "0" | "false" | "off" | "no" => Some(false), + "1" | "true" | "on" | "yes" => Some(true), + _ => None, } } @@ -9385,6 +9408,45 @@ mod tests { /// recognised disable string turns the fusion off, an unset variable leaves /// the compiled-in default, and an unrecognised value must not silently /// change the decode graph. + #[test] + fn fused_905_defaults_are_on_for_rocm_builds_only() { + // #2145: a `rocm` build turns both fusions on by default; Metal and CUDA + // builds keep them off. Pinned against the feature rather than a + // literal so each backend's build checks its own default. + assert_eq!(FUSED_ADD_RMSNORM_DEFAULT, cfg!(feature = "rocm")); + assert_eq!(FUSED_ROPE_APPEND_DEFAULT, cfg!(feature = "rocm")); + #[cfg(feature = "rocm")] + { + assert!(fused_add_rmsnorm_enabled_from(None)); + assert!(fused_rope_append_enabled_from(None)); + // The kill switch still turns them off on ROCm. + assert!(!fused_add_rmsnorm_enabled_from(Some("0"))); + assert!(!fused_rope_append_enabled_from(Some("0"))); + } + #[cfg(not(feature = "rocm"))] + { + assert!(!fused_add_rmsnorm_enabled_from(None)); + assert!(!fused_rope_append_enabled_from(None)); + // And `=1` still opts in elsewhere. + assert!(fused_add_rmsnorm_enabled_from(Some("1"))); + assert!(fused_rope_append_enabled_from(Some("1"))); + } + } + + #[test] + fn fused_flag_explicit_value_separates_a_setting_from_the_default() { + assert_eq!(fused_flag_explicit_value(None), None); + for v in ["", "maybe", "2"] { + assert_eq!(fused_flag_explicit_value(Some(v)), None, "{v:?}"); + } + for v in ["0", "false", "off", "no", " OFF "] { + assert_eq!(fused_flag_explicit_value(Some(v)), Some(false), "{v:?}"); + } + for v in ["1", "true", "on", "yes", " Yes "] { + assert_eq!(fused_flag_explicit_value(Some(v)), Some(true), "{v:?}"); + } + } + #[test] fn fused_905_flags_follow_the_default_and_respect_explicit_values() { // Unset means the compiled-in default, whichever way it is set. Asserting diff --git a/src/models/llama3.rs b/src/models/llama3.rs index 2695dce5c..84db19fd6 100644 --- a/src/models/llama3.rs +++ b/src/models/llama3.rs @@ -257,17 +257,27 @@ const FUSED_ROPE_ENV_VARS_WITHOUT_TABLE_SUPPORT: usize = 2; /// /// The first two are presence-enabled, so `is_ok()` is the whole test. /// `MLXCEL_FUSED_ROPE_APPEND` is tri-state (`0` disables, `1` enables, unset -/// takes the compiled-in default), so presence would report a lost kernel to -/// someone who set it to `0` precisely to not use one. Asking the same gate the -/// launcher asks keeps the notice tied to what actually happens. +/// takes the compiled-in default), so it counts only when it holds an explicit +/// truthy value: presence would report a lost kernel to someone who set it to +/// `0` precisely to not use one, and asking the gate would report one to every +/// Llama 3.1 run on a `rocm` build, where the default is on (#2145) and the +/// user never asked for anything. fn fused_rope_launcher_requested(name: &str) -> bool { if name == FUSED_ROPE_APPEND_ENV { - mlxcel_core::layers::fused_rope_append_enabled() + fused_rope_append_explicitly_requested(std::env::var(name).ok().as_deref()) } else { std::env::var(name).is_ok() } } +/// Pure decision behind the `MLXCEL_FUSED_ROPE_APPEND` arm of +/// [`fused_rope_launcher_requested`]: only an explicit truthy value counts. +/// Unset and unrecognised values take the compiled-in default, which is not a +/// request. +pub(crate) fn fused_rope_append_explicitly_requested(value: Option<&str>) -> bool { + mlxcel_core::layers::fused_flag_explicit_value(value) == Some(true) +} + /// Print, at most once per process, that a checkpoint has been routed around a /// fused RoPE launcher that an environment variable asked for. /// @@ -289,8 +299,9 @@ fn fused_rope_launcher_requested(name: &str) -> bool { /// Correctness does not depend on this: [`Attention::forward`] bypasses the /// launchers regardless. What it buys is that the fallback is reported rather /// than silent, which is the one real cost the bypass carries. The notice fires -/// only when a bypass reason holds *and* a variable asked for a fused path, so a -/// user who never sets the variables never sees it. +/// only when a bypass reason holds *and* a variable explicitly asked for a fused +/// path, so a user who never sets the variables never sees it, including on a +/// `rocm` build where `MLXCEL_FUSED_ROPE_APPEND` defaults to on. /// /// `eprintln!` rather than `tracing::warn!` on purpose: only the server installs /// a `tracing` subscriber, so a `warn!` here is a no-op in the `mlxcel` CLI, @@ -536,7 +547,8 @@ impl Attention { /// 2. [`FusedQKVLinear::forward_split_rope`] (quantized projection, opt-in /// through `MLXCEL_ENABLE_FUSED_QKV_SPLIT_ROPE`), /// 3. [`FusedQKVLinear::forward_fused_rope_append`] (dense decode, opt-in - /// through `MLXCEL_FUSED_ROPE_APPEND`), + /// through `MLXCEL_FUSED_ROPE_APPEND` on Metal and CUDA, on by default + /// in a `rocm` build), /// 4. the graph fallback, which goes through [`Self::apply_rope`] and so /// calls `fast_rope_with_freqs` when a table is present and `fast_rope` /// otherwise. diff --git a/src/models/llama3_tests.rs b/src/models/llama3_tests.rs index f4bcc047d..505f4a60a 100644 --- a/src/models/llama3_tests.rs +++ b/src/models/llama3_tests.rs @@ -39,7 +39,7 @@ use super::{ Attention, FUSED_CAUSAL_PREFILL_ENV, FUSED_QKV_SPLIT_ROPE_ENV, FUSED_ROPE_APPEND_ENV, - FUSED_ROPE_ENV_VARS, Llama3Model, ModelArgs, + FUSED_ROPE_ENV_VARS, Llama3Model, ModelArgs, fused_rope_append_explicitly_requested, }; use crate::models::rope_utils::{RopeScalingKind, RopeScalingSpec}; use crate::test_support::env_lock::env_lock; @@ -386,6 +386,26 @@ fn the_fused_rope_env_vars_are_the_ones_the_code_actually_reads() { assert_eq!(FUSED_ROPE_APPEND_ENV, FUSED_ROPE_ENV_VARS[2]); } +#[test] +fn the_rope_append_bypass_notice_needs_an_explicit_truthy_value() { + // #2145: `MLXCEL_FUSED_ROPE_APPEND` defaults to on in a `rocm` build, so the + // bypass notice must not treat the default as a request. Unset is the case + // every default ROCm Llama 3.1 run hits; it must stay silent on every build. + assert!(!fused_rope_append_explicitly_requested(None)); + for v in ["", "maybe", "2", "0", "false", "off", "no"] { + assert!( + !fused_rope_append_explicitly_requested(Some(v)), + "{v:?} is not a request for the fused RoPE + append launcher" + ); + } + for v in ["1", "true", "on", "yes", " ON "] { + assert!( + fused_rope_append_explicitly_requested(Some(v)), + "{v:?} explicitly requests the fused RoPE + append launcher" + ); + } +} + /// A genuinely quantized attention block, which is the only shape that can /// reach either fused launcher, built with the given rotation convention. fn quantized_attention(rope_traditional: bool) -> Attention { From 3f5783c65654f3be4f25f96367241eb7ef6f8251 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Wed, 7 Oct 2026 08:16:56 +0900 Subject: [PATCH 2/4] update(rocm): record the default-on re-measurement for the fused norm and RoPE ports --- ...o-gfx1151_2026-10-07_fused-default-off.csv | 7 +++++ ...lo-gfx1151_2026-10-07_fused-default-on.csv | 7 +++++ ...rocm-fused-norm-rope-gfx1151-2026-10-05.md | 31 +++++++++++++++---- 3 files changed, 39 insertions(+), 6 deletions(-) create mode 100644 benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-off.csv create mode 100644 benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-on.csv diff --git a/benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-off.csv b/benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-off.csv new file mode 100644 index 000000000..4408d14ac --- /dev/null +++ b/benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-off.csv @@ -0,0 +1,7 @@ +model,model_path,prompt_tokens,generated_tokens,prefill_ms,prefill_tok_s,decode_ms,decode_tok_s,date,hardware,mlxcel_version,build_type,max_tokens,prompt,prompt_target_len,mlxcel_commit,mlx_commit,mlx_rocm_overlay_commit,hip_version +Meta-Llama-3.1-8B-Instruct-4bit,models/mlx/Meta-Llama-3.1-8B-Instruct-4bit,512,128,506.66,1010.54,3382.72,37.84,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Meta-Llama-3.1-8B-Instruct-4bit,models/mlx/Meta-Llama-3.1-8B-Instruct-4bit,512,128,499.94,1024.12,3317.53,38.58,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Meta-Llama-3.1-8B-Instruct-4bit,models/mlx/Meta-Llama-3.1-8B-Instruct-4bit,512,128,504.75,1014.37,3355.18,38.15,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Qwen2.5-7B-Instruct-4bit,models/mlx/Qwen2.5-7B-Instruct-4bit,512,128,331.59,1544.08,2694.05,47.51,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Qwen2.5-7B-Instruct-4bit,models/mlx/Qwen2.5-7B-Instruct-4bit,512,128,328.36,1559.28,2666.65,48.00,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Qwen2.5-7B-Instruct-4bit,models/mlx/Qwen2.5-7B-Instruct-4bit,512,128,333.50,1535.22,2705.67,47.31,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 diff --git a/benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-on.csv b/benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-on.csv new file mode 100644 index 000000000..a4ad1c4c1 --- /dev/null +++ b/benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-on.csv @@ -0,0 +1,7 @@ +model,model_path,prompt_tokens,generated_tokens,prefill_ms,prefill_tok_s,decode_ms,decode_tok_s,date,hardware,mlxcel_version,build_type,max_tokens,prompt,prompt_target_len,mlxcel_commit,mlx_commit,mlx_rocm_overlay_commit,hip_version +Meta-Llama-3.1-8B-Instruct-4bit,models/mlx/Meta-Llama-3.1-8B-Instruct-4bit,512,128,495.60,1033.08,3380.83,37.86,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Meta-Llama-3.1-8B-Instruct-4bit,models/mlx/Meta-Llama-3.1-8B-Instruct-4bit,512,128,499.68,1024.65,3360.17,38.09,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Meta-Llama-3.1-8B-Instruct-4bit,models/mlx/Meta-Llama-3.1-8B-Instruct-4bit,512,128,504.77,1014.32,3359.47,38.10,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Qwen2.5-7B-Instruct-4bit,models/mlx/Qwen2.5-7B-Instruct-4bit,512,128,318.02,1609.96,2672.40,47.90,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Qwen2.5-7B-Instruct-4bit,models/mlx/Qwen2.5-7B-Instruct-4bit,512,128,319.59,1602.06,2650.37,48.30,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 +Qwen2.5-7B-Instruct-4bit,models/mlx/Qwen2.5-7B-Instruct-4bit,512,128,312.08,1640.59,2708.01,47.27,2026-10-07,AMD_Radeon_8060S_Graphics_gfx1151_ROCm10.0.0_96GB,0.7.0,release,128,"Hello, how are you today?",512,57b59629,81ba1c6a,75915908,7.15.26333 diff --git a/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md b/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md index 75d111d6e..f66e5abba 100644 --- a/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md +++ b/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md @@ -1,6 +1,6 @@ # HIP ports of the fused add-RMSNorm and RoPE + append kernels on gfx1151 (2026-10-05) -lablup/mlxcel#2063, part of #1814. `fused_add_rms_norm` (#905) and `fused_rope_qk_append` (#905) had Metal and CUDA ports only, so on ROCm `MLXCEL_FUSED_ADD_RMSNORM=1` and `MLXCEL_FUSED_ROPE_APPEND=1` quietly kept the graph path. Both fusions ship off on every backend (`FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs`), and the decode profile ([rocm-decode-profile-gfx1151-2026-09-30.md](rocm-decode-profile-gfx1151-2026-09-30.md)) put the most they could take over at 0.83 to 0.89% of Llama 3.1 decode GPU time. This page records what the ports match, what turning them on does to decode, and why the ROCm default stays off. +lablup/mlxcel#2063, part of #1814. `fused_add_rms_norm` (#905) and `fused_rope_qk_append` (#905) had Metal and CUDA ports only, so on ROCm `MLXCEL_FUSED_ADD_RMSNORM=1` and `MLXCEL_FUSED_ROPE_APPEND=1` quietly kept the graph path. Both fusions ship off on every backend (`FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs`), and the decode profile ([rocm-decode-profile-gfx1151-2026-09-30.md](rocm-decode-profile-gfx1151-2026-09-30.md)) put the most they could take over at 0.83 to 0.89% of Llama 3.1 decode GPU time. This page records what the ports match, what turning them on does to decode, and the ROCm default: the measurement below recommended keeping it off, and the maintainer then turned both fusions on for ROCm (#2145, see [Decision](#decision-2145)). ## Environment @@ -61,13 +61,30 @@ How to read it: - **Qwen3-30B-A3B** calls neither kernel (`qwen3_moe.rs` has its own block), so its row is a control: +0.3%, noise. - **Qwen2.5 7B** is the one model that reaches both ports. Every on arm is above off in 6 of 7 rounds, but the size does not hold still: rounds 1 to 3 gave +0.95 to +3.5 tok/s, rounds 4 to 7 (run about two hours later on the same binary) mostly -0.4 to +1 (with -2.7 for `both` in round 4 and +2.7 for `add` in round 7), and `both` is not above `add`, which it would be if both ports saved time independently. An earlier run on the branch's first commit gave the same picture (7 rounds: off 46.54, add 46.84, rope 46.95, both 47.58 tok/s). The off arm's own spread (45.75 to 46.95) is about as wide as the median gain. A gain of 1 to 2% on Qwen2.5 is plausible from the removed dispatches (the RoPE port alone stands in for three slices, three reshape and transpose pairs and two `fast_rope` calls per layer, and #2099 showed that GPU-time shares understate paths made of many small ops), but this data does not establish it. -## Recommendation +## Recommendation (#2063) -Keep `FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` at `false` on ROCm, as on Metal and CUDA. The measurement does not clearly support a flip: no change on Llama 3.1, and on Qwen2.5 a small positive median whose size changes with the time of day. Because the ports are byte-identical to the graph, opting in on ROCm costs nothing in output, so `MLXCEL_FUSED_ADD_RMSNORM=1 MLXCEL_FUSED_ROPE_APPEND=1` is safe for a deployment that measures a win on its own model. What would justify a ROCm default of on: a longer interleaved run on a quiet GPU (20 or more pairs) that shows Qwen2.5-class models (no `rope_scaling`, so both ports run) gaining at least 1% with the paired differences consistently positive, and Llama-class models not losing. +The original recommendation, superseded by the decision below: keep `FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` at `false` on ROCm, as on Metal and CUDA. The measurement does not clearly support a flip: no change on Llama 3.1, and on Qwen2.5 a small positive median whose size changes with the time of day. Because the ports are byte-identical to the graph, opting in on ROCm costs nothing in output, so `MLXCEL_FUSED_ADD_RMSNORM=1 MLXCEL_FUSED_ROPE_APPEND=1` is safe for a deployment that measures a win on its own model. What would justify a ROCm default of on: a longer interleaved run on a quiet GPU (20 or more pairs) that shows Qwen2.5-class models (no `rope_scaling`, so both ports run) gaining at least 1% with the paired differences consistently positive, and Llama-class models not losing. + +## Decision (#2145) + +The maintainer turned both fusions on by default on ROCm: the ports cost no accuracy (every on/off trace pair is byte-identical) and no median above was slower. `FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` are now `cfg!(feature = "rocm")`, so a ROCm build runs both kernels with no environment variables set, Metal and CUDA builds keep them off, and `MLXCEL_FUSED_ADD_RMSNORM=0` / `MLXCEL_FUSED_ROPE_APPEND=0` still restore the graph. The Llama 3.1 `rope_scaling` bypass notice now prints only when `MLXCEL_FUSED_ROPE_APPEND` is explicitly set to a truthy value, so default ROCm runs stay quiet. + +Re-measurement of the default-on build (2026-10-07, commit `57b59629` on `main` `08e22698`, same host and toolchain as above): `scripts/bench_decode.sh` at pp512/tg128, default (no variables) against `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0`, arm order alternated per round, every run through `scripts/rocm_gpu_guard.sh` with its host-wide lock (#2146), 12 of 12 clean on the first attempt. Raw rows: `benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-{on,off}.csv`. + +| Model | Arm | Decode tok/s, by round | Median | vs off | Paired (on minus off), by round | Prefill median | +|---|---|---|---|---|---|---| +| Llama-3.1-8B | default (on) | 37.86 38.09 38.10 | 38.09 | -0.2% | +0.02 -0.49 -0.05 | 1024.65 | +| | off | 37.84 38.58 38.15 | 38.15 | | | 1014.37 | +| Qwen2.5-7B | default (on) | 47.90 48.30 47.27 | 47.90 | +0.8% | +0.39 +0.30 -0.04 | 1609.96 | +| | off | 47.51 48.00 47.31 | 47.51 | | | 1544.08 | + +Same reading as before: Llama 3.1, where only the norm port runs, is flat (its -0.2% comes from one off run at 38.58, above every other run of either arm); Qwen2.5, where both run, is slightly above off in two of three rounds, inside the run-to-run spread. Qwen2.5 prefill was 2.7 to 6.9% higher with the default in every round, a pattern the earlier run did not show and three rounds cannot separate from drift. No result here argues against the decision. + +Kernel trace (`rocprofv3 --kernel-trace`, 8-token `mlxcel generate`, no variables set): Qwen2.5 dispatched `fused_add_rms_norm` and `fused_rope_qk_append` 252 times each (28 layers x 9 forward passes) and no graph RoPE kernel; with both variables at `0` it dispatched `rope_single_1d` instead. Llama 3.1 dispatched the norm port 288 times (32 x 9) and kept `rope_single_freqs_1d` / `rope_freqs`. `logit_trace` `w8` traces of both models with the default and with both variables at `0` are byte-identical, file for file. ## Not covered -Metal and CUDA (not available on this host); their kernel sources and table entries are untouched. Their predicates gained the GPU-device term: GPU behavior is unchanged, and `MLXCEL_DEVICE=cpu` now takes the graph path instead of throwing at the first launch. Gemma and IQuest Loop Coder also call `fused_add_rms_norm` but have no checkpoint here; the Gemma `(1 + w)` convention is covered by the tolerance tests only. Wave64 (CDNA) is untested. +Metal and CUDA (not available on this host); their kernel sources and table entries are untouched. Their predicates gained the GPU-device term: GPU behavior is unchanged, and `MLXCEL_DEVICE=cpu` now takes the graph path instead of throwing at the first launch. Gemma and IQuest Loop Coder also call `fused_add_rms_norm`, and on ROCm now do so by default, but have no checkpoint here; the Gemma `(1 + w)` convention is covered by the tolerance tests only. Wave64 (CDNA) is untested. ## Reproducing @@ -75,8 +92,10 @@ Metal and CUDA (not available on this host); their kernel sources and table entr cargo build --release --features rocm --bin mlxcel --bin mlxcel-bench-decode --example logit_trace cargo test --release --features rocm -p mlxcel-core --lib -- --test-threads=1 fused_norm_parity_tests fused_rope_parity_tests cargo test --release --features rocm --test cpu_device_custom_kernel_gates -MLXCEL_FUSED_ADD_RMSNORM=1 MLXCEL_FUSED_ROPE_APPEND=1 \ - scripts/rocm_gpu_guard.sh -- scripts/bench_decode.sh models/mlx/Qwen2.5-7B-Instruct-4bit --output /tmp/on.csv +# On ROCm the fusions are on by default since #2145; the off arm sets both to 0. +scripts/rocm_gpu_guard.sh -- scripts/bench_decode.sh models/mlx/Qwen2.5-7B-Instruct-4bit --output /tmp/on.csv +MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0 \ + scripts/rocm_gpu_guard.sh -- scripts/bench_decode.sh models/mlx/Qwen2.5-7B-Instruct-4bit --output /tmp/off.csv ``` The trace commands are in `benchmarks/logit_traces/rocm_gfx1151_3f0e51af/README.md`. From 9874c3c4786d74bf13735739064b85a910aff2b7 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Wed, 7 Oct 2026 08:36:04 +0900 Subject: [PATCH 3/4] update(rocm): align README, docs and the decode profile with the ROCm fused defaults --- README.md | 2 +- .../rocm-fused-norm-rope-gfx1151-2026-10-05.md | 2 +- docs/environment-variables.md | 4 ++-- scripts/rocm_decode_profile.py | 9 +++++---- src/lib/mlxcel-core/src/layers.rs | 10 +++++----- tests/test_rocm_decode_profile.py | 9 ++++++--- 6 files changed, 20 insertions(+), 16 deletions(-) diff --git a/README.md b/README.md index 1f28a61ef..bceddfd0e 100644 --- a/README.md +++ b/README.md @@ -16,7 +16,7 @@ Apple Silicon is the primary target. Linux/CUDA is supported as a secondary targ The current `main` branch is v0.7.0 plus unreleased work. Install from source to use features that have not reached the latest tagged release; see the [changelog](CHANGELOG.md) for the release-by-release record. -- **Experimental AMD GPU (ROCm) support.** `--features rocm` builds mlxcel on Linux against an MLX ROCm backend vendored into the source tree as the ROCm part of mlxcelverse, applied on the same pinned MLX commit as the Metal and CUDA builds. Affine, mxfp8, and mxfp4 checkpoints, including gpt-oss MoE experts, run on RDNA 3.5 (`gfx1151`); the fused kernels with a HIP port (BitLinear, the samplers, the SSM update step, the Mamba1 scan, the decode-MoE kernels, xIELU, the add3 LayerNorm, and the opt-in fused add-RMSNorm and RoPE + append) run natively and the rest fall back to MLX graphs. See [Linux with AMD ROCm](docs/installation.md#linux-with-amd-rocm-experimental) for the build and the open gaps. +- **Experimental AMD GPU (ROCm) support.** `--features rocm` builds mlxcel on Linux against an MLX ROCm backend vendored into the source tree as the ROCm part of mlxcelverse, applied on the same pinned MLX commit as the Metal and CUDA builds. Affine, mxfp8, and mxfp4 checkpoints, including gpt-oss MoE experts, run on RDNA 3.5 (`gfx1151`); the fused kernels with a HIP port (BitLinear, the samplers, the SSM update step, the Mamba1 scan, the decode-MoE kernels, xIELU, the add3 LayerNorm, and the fused add-RMSNorm and RoPE + append, which are on by default on ROCm) run natively and the rest fall back to MLX graphs. See [Linux with AMD ROCm](docs/installation.md#linux-with-amd-rocm-experimental) for the build and the open gaps. - **Verified `llama-server` compatibility.** A frozen b10621 manifest classifies all 376 pinned options, routes, and native request fields, with no deferred entries. Native completion, embedding, tokenization, template, infill, props, slots, metrics, resumable-stream, router, and LoRA surfaces are implemented or explicitly classified instead of being silently ignored. - **OpenAI, Anthropic, and Vertex-compatible serving.** Chat Completions, Completions, Responses, Embeddings, Reranking, Audio, and Anthropic Messages are available alongside the native `llama-server` routes. Optional Vertex AI custom-container routing is supported through the standard `AIP_*` variables. - **Multi-model routing and live adapters.** Router mode discovers checkpoints from the model store, a model directory, or INI presets; loads them on demand; and bounds the resident set with LRU eviction. Multiple LoRA adapters can remain unfused for per-request or live scale changes, or be fused for zero decode overhead. diff --git a/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md b/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md index f66e5abba..29c4bd6cb 100644 --- a/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md +++ b/docs/benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md @@ -69,7 +69,7 @@ The original recommendation, superseded by the decision below: keep `FUSED_ADD_R The maintainer turned both fusions on by default on ROCm: the ports cost no accuracy (every on/off trace pair is byte-identical) and no median above was slower. `FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` are now `cfg!(feature = "rocm")`, so a ROCm build runs both kernels with no environment variables set, Metal and CUDA builds keep them off, and `MLXCEL_FUSED_ADD_RMSNORM=0` / `MLXCEL_FUSED_ROPE_APPEND=0` still restore the graph. The Llama 3.1 `rope_scaling` bypass notice now prints only when `MLXCEL_FUSED_ROPE_APPEND` is explicitly set to a truthy value, so default ROCm runs stay quiet. -Re-measurement of the default-on build (2026-10-07, commit `57b59629` on `main` `08e22698`, same host and toolchain as above): `scripts/bench_decode.sh` at pp512/tg128, default (no variables) against `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0`, arm order alternated per round, every run through `scripts/rocm_gpu_guard.sh` with its host-wide lock (#2146), 12 of 12 clean on the first attempt. Raw rows: `benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-{on,off}.csv`. +Re-measurement of the default-on build (2026-10-07, commit `57b59629`, the pre-rebase SHA of PR #2189's first commit, on `main` `08e22698`, same host and toolchain as above): `scripts/bench_decode.sh` at pp512/tg128, default (no variables) against `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0`, arm order alternated per round, every run through `scripts/rocm_gpu_guard.sh` with its host-wide lock (#2146), 12 of 12 clean on the first attempt. Raw rows: `benchmarks/rocm_strixhalo-gfx1151_2026-10-07_fused-default-{on,off}.csv`. | Model | Arm | Decode tok/s, by round | Median | vs off | Paired (on minus off), by round | Prefill median | |---|---|---|---|---|---|---| diff --git a/docs/environment-variables.md b/docs/environment-variables.md index cadf80fc1..ebe304c61 100644 --- a/docs/environment-variables.md +++ b/docs/environment-variables.md @@ -571,8 +571,8 @@ recommended as normal deployment settings. | `MLXCEL_SSM_KERNEL` | `0` disables; any other value or unset enables | on | Gates the fused single-token SSM decode kernel (`ssm_update_kernel`) on every backend that has a port: Metal, CUDA (issue #631) and ROCm (issue #2067). It replaces the ~55-op SSD scan graph (`ssm_step`) with one launch per SSM layer per decode step for the hybrid SSM models (granite-4.0-h, falcon-h1, plamo-2, nemotron-h). `ssm_kernel_available()` answers from the kernel's port table (`ssm_ports()`), so a backend with no port takes the graph path on its own. It is also false when the default device is the CPU (`MLXCEL_DEVICE=cpu`). Set to `0` to force the graph path (A/B, rollback). Measured on GB10 (CUDA): 2.6-4.5x decode on the granite/falcon family, see `docs/benchmark_results/hybrid-ssm-decode-cuda-kernel-gb10-2026-07-10.md`; on gfx1151 (ROCm): see `docs/benchmark_results/rocm-ssm-update-kernel-gfx1151-2026-10-04.md`. | | `MLXCEL_SSM_CUDA_KERNEL` | `0` disables; any other value or unset enables | on | Older name of `MLXCEL_SSM_KERNEL`, from when the CUDA port was the only one off Apple, kept as an alias so existing scripts still work. Either variable set to `0` forces the graph path, now on every backend (Metal and ROCm included). | | `MLXCEL_FUSED_QK_NORM` | `1`/`true`/`on`/`yes` enable; any other value or unset disables | off (opt-in) | Fused single-token QKV projection + Q/K RMSNorm + RoPE kernel (#326) for Qwen3 and Qwen3-MoE decode. Opt-in: set to `1` to enable. Matches the graph path within RMS < 5e-3 (the reduction is over the transpose-invariant head_dim axis), but greedy temp-0 is not byte-identical over long generation; on CUDA the graph path is itself non-deterministic run-to-run from GPU FP-reduction order, while the fused path is deterministic, so its output stays inside the graph baseline's own envelope. The kernel cuts Rust/C++ FFI crossings rather than MLX op count, so it does not speed up the GPU/bandwidth-bound decode loop: on M1 Ultra it measured 1 to 3.4% slower (qwen3-0.6b 275 vs 284, qwen3-8b 82.3 vs 83.2 tok/s); on GB10/CUDA (SM 12.1) it is also slower (qwen3-0.6b 0.96x, qwen3-8b ~1.0x, qwen3-30b-a3b 0.92x fused/graph; see `docs/benchmark_results/fused-qk-norm-decode-gb10.md`), so there is no per-backend win. Ships as a reusable shared primitive for the deferred QK-norm families and stays opt-in (default off) on every measured backend (M1 Ultra, M5 Max, GB10/CUDA), mirroring the opt-in `MLXCEL_FUSED_MOE_RELU2`. Active only when `l == 1` (decode) and weights are quantized. | -| `MLXCEL_FUSED_ADD_RMSNORM` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (opt-in) | Fused residual-add + RMSNorm decode kernel (#905). Computes `new_residual = x + residual` and `normed = rms_norm(new_residual) * (weight_bias + weight)` in one dispatch instead of an elementwise `Add` followed by `fast::rms_norm`, removing one kernel launch and one full-width intermediate at every pre-norm residual join. Adopted by the Llama3-family block (which also serves Qwen2/Qwen2.5 and every VLM whose text backbone is `Llama3Model`/`Qwen2Model`) and by Gemma, where the `(1 + w)` convention rides in as a scalar `weight_bias = 1.0` folded in the weight's own dtype rather than as a second weight tensor. Numerics are pinned to MLX's own `rms_norm` kernel (fp32 accumulation over the dtype-rounded sum, `precise::rsqrt`, the `x * inv_mean` product rounded to the activation dtype before the weight multiply), so the fused and unfused paths agree to within a small multiple of one ulp; greedy temp-0 is not guaranteed byte-identical over long generation, because a near-tie argmax can land on either side of that difference. Set to `0` to force the `add` + `fast_rms_norm` pair (A/B, rollback). The compiled-in default lives in `FUSED_ADD_RMSNORM_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` and depends on the build: `cfg!(feature = "rocm")`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off) on Metal and CUDA**: the op-level microbench on M1 Ultra (Metal, f16, hidden {2048, 4096, 8192} x batch {1, 4, 8}, three repetitions) measured the fused path at parity, speedups scattering 0.77x to 1.14x with no structure and the extremes not reproducing. Per-iteration time stayed near 280-410us across a 4x change in hidden size, so the harness is dominated by dispatch cost rather than the op. Issue #905's measure-then-keep policy lands a fusion unwired absent a demonstrated win, matching how `MLXCEL_FUSED_QK_NORM` (#326) shipped. The production decode path agrees: on M1 Ultra at pp512/tg128, Llama 3.1 8B decodes at 113.18 tok/s off and 113.01 on, Qwen2.5 7B at 115.38 and 115.34, with prefill unchanged. Set to `1` to opt in. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-add-rmsnorm-decode-m1ultra-2026-09-27.md`. On ROCm it is **on by default** (lablup/mlxcel#2145): the kernel runs as a HIP port (lablup/mlxcel#2063) whose output is byte-identical to the `add` + `fast_rms_norm` graph (Llama 3.1 traces with it on and off match), and on gfx1151 it was never slower in a median (Llama 3.1 8B decode 37.85 tok/s off and 37.95 on, medians of five; Qwen2.5 7B +1.5%), so the change costs no accuracy and the maintainer turned it on there. Set to `0` on ROCm to opt out. Gemma and IQuest Loop Coder take the same kernel by default on ROCm but have no checkpoint measured on that backend; their `(1 + w)` convention is covered by the parity tests only ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). -| `MLXCEL_FUSED_ROPE_APPEND` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (opt-in) | Fused q/k RoPE + KV-append-layout decode kernel (#905). Reads the row-contiguous fused-QKV projection output whole, applies rotary embedding to the q and k blocks, and emits q, k and v already in their consumers' layouts, replacing three trailing-axis slices, three reshape/transpose pairs and two `fast::rope` calls with one dispatch. The rotation is transcribed from MLX's own `rope.metal`, so it agrees with `fast::rope` including the `traditional` convention and partial `rope_dims`. The trig call is the one thing chosen per backend rather than transcribed: each body calls whatever MLX's own RoPE kernel for that backend calls, so the Metal body uses `metal::fast::cos`/`sin` after `rope.metal` and the CUDA body uses libdevice `cosf`/`sinf` after `rope.cu`. Using the CUDA fast intrinsics `__cosf`/`__sinf` there instead cost about 1.2e-2 rad of applied phase at absolute position 131071 and failed the parity suite at the two large-offset cases (issue #1049). Adopted by the Llama3-family dense decode path, except for a checkpoint whose `rope_scaling` block selects a frequency table: the kernel derives its frequencies from `rope_base` and has no parameter for a table, so such a model takes the graph path (#1355), and a one-time notice on stderr says so when this variable is explicitly set to a truthy value (the ROCm default alone does not print it). `traditional` and the position scale are real parameters and are honored, so neither triggers that notice for this kernel. The paged block-pool destination layout the kernel also supports is implemented and tested but not wired, because the batched paged decode path belongs to issue #899. The kernel emits the append payload rather than writing the cache slab, because MLX custom-kernel outputs are always freshly allocated while the existing `slice_update` append donates the slab buffer, so an in-kernel write would trade an O(1) donated update for an O(capacity) copy. Set to `0` to force the reshape/transpose/`fast_rope` graph (A/B, rollback). The compiled-in default lives in `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` and depends on the build: `cfg!(feature = "rocm")`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off) on Metal and CUDA** for the same measured reason as `MLXCEL_FUSED_ADD_RMSNORM`, with one cell consistently below parity rather than merely noisy: hidden 8192 at batch 1 measured 0.94x, 0.90x and 0.89x across three repetitions. Set to `1` to opt in. On GB10/CUDA the same harness measures 1.02x to 2.06x with every cell at or above parity, a different picture from Apple Silicon that has not been acted on; flipping the default is its own decision with its own evidence bar. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-rope-cuda-trig-gb10-2026-08-06.md`. On ROCm it is **on by default** (lablup/mlxcel#2145): the kernel runs as a HIP port (lablup/mlxcel#2063) that calls `sincosf` as `rope.hip` does and reproduces the graph's FMA forms and f16 rounding, so its output is byte-identical to the graph, and Qwen2.5 7B decode on gfx1151 gained 0.6% with it alone and 1.7% with both fusions (medians of seven), inside run-to-run drift but never slower. Set to `0` on ROCm to opt out. Llama 3.1 still takes the graph path there because of its `rope_scaling` table; the stderr notice about that fires only when this variable is set to a truthy value, not for the ROCm default ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). +| `MLXCEL_FUSED_ADD_RMSNORM` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (Metal, CUDA) / on (ROCm) | Fused residual-add + RMSNorm decode kernel (#905). Computes `new_residual = x + residual` and `normed = rms_norm(new_residual) * (weight_bias + weight)` in one dispatch instead of an elementwise `Add` followed by `fast::rms_norm`, removing one kernel launch and one full-width intermediate at every pre-norm residual join. Adopted by the Llama3-family block (which also serves Qwen2/Qwen2.5 and every VLM whose text backbone is `Llama3Model`/`Qwen2Model`) and by Gemma, where the `(1 + w)` convention rides in as a scalar `weight_bias = 1.0` folded in the weight's own dtype rather than as a second weight tensor. Numerics are pinned to MLX's own `rms_norm` kernel (fp32 accumulation over the dtype-rounded sum, `precise::rsqrt`, the `x * inv_mean` product rounded to the activation dtype before the weight multiply), so the fused and unfused paths agree to within a small multiple of one ulp; greedy temp-0 is not guaranteed byte-identical over long generation, because a near-tie argmax can land on either side of that difference. Set to `0` to force the `add` + `fast_rms_norm` pair (A/B, rollback). The compiled-in default lives in `FUSED_ADD_RMSNORM_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` and depends on the build: `cfg!(feature = "rocm")`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off) on Metal and CUDA**: the op-level microbench on M1 Ultra (Metal, f16, hidden {2048, 4096, 8192} x batch {1, 4, 8}, three repetitions) measured the fused path at parity, speedups scattering 0.77x to 1.14x with no structure and the extremes not reproducing. Per-iteration time stayed near 280-410us across a 4x change in hidden size, so the harness is dominated by dispatch cost rather than the op. Issue #905's measure-then-keep policy lands a fusion unwired absent a demonstrated win, matching how `MLXCEL_FUSED_QK_NORM` (#326) shipped. The production decode path agrees: on M1 Ultra at pp512/tg128, Llama 3.1 8B decodes at 113.18 tok/s off and 113.01 on, Qwen2.5 7B at 115.38 and 115.34, with prefill unchanged. Set to `1` to opt in. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-add-rmsnorm-decode-m1ultra-2026-09-27.md`. On ROCm it is **on by default** (lablup/mlxcel#2145): the kernel runs as a HIP port (lablup/mlxcel#2063) whose output is byte-identical to the `add` + `fast_rms_norm` graph (Llama 3.1 traces with it on and off match), and on gfx1151 it was never slower in a median (Llama 3.1 8B decode 37.85 tok/s off and 37.95 on, medians of five; Qwen2.5 7B +1.5%), so the change costs no accuracy and the maintainer turned it on there. Set to `0` on ROCm to opt out. Gemma and IQuest Loop Coder take the same kernel by default on ROCm but have no checkpoint measured on that backend; their `(1 + w)` convention is covered by the parity tests only ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). +| `MLXCEL_FUSED_ROPE_APPEND` | `0`/`false`/`off`/`no` disable; `1`/`true`/`on`/`yes` enable; unset or unrecognised keeps the default | off (Metal, CUDA) / on (ROCm) | Fused q/k RoPE + KV-append-layout decode kernel (#905). Reads the row-contiguous fused-QKV projection output whole, applies rotary embedding to the q and k blocks, and emits q, k and v already in their consumers' layouts, replacing three trailing-axis slices, three reshape/transpose pairs and two `fast::rope` calls with one dispatch. The rotation is transcribed from MLX's own `rope.metal`, so it agrees with `fast::rope` including the `traditional` convention and partial `rope_dims`. The trig call is the one thing chosen per backend rather than transcribed: each body calls whatever MLX's own RoPE kernel for that backend calls, so the Metal body uses `metal::fast::cos`/`sin` after `rope.metal` and the CUDA body uses libdevice `cosf`/`sinf` after `rope.cu`. Using the CUDA fast intrinsics `__cosf`/`__sinf` there instead cost about 1.2e-2 rad of applied phase at absolute position 131071 and failed the parity suite at the two large-offset cases (issue #1049). Adopted by the Llama3-family dense decode path, except for a checkpoint whose `rope_scaling` block selects a frequency table: the kernel derives its frequencies from `rope_base` and has no parameter for a table, so such a model takes the graph path (#1355), and a one-time notice on stderr says so when this variable is explicitly set to a truthy value (the ROCm default alone does not print it). `traditional` and the position scale are real parameters and are honored, so neither triggers that notice for this kernel. The paged block-pool destination layout the kernel also supports is implemented and tested but not wired, because the batched paged decode path belongs to issue #899. The kernel emits the append payload rather than writing the cache slab, because MLX custom-kernel outputs are always freshly allocated while the existing `slice_update` append donates the slab buffer, so an in-kernel write would trade an O(1) donated update for an O(capacity) copy. Set to `0` to force the reshape/transpose/`fast_rope` graph (A/B, rollback). The compiled-in default lives in `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` and depends on the build: `cfg!(feature = "rocm")`. Read once at first use and cached for the process lifetime. | Ships **opt-in (default off) on Metal and CUDA** for the same measured reason as `MLXCEL_FUSED_ADD_RMSNORM`, with one cell consistently below parity rather than merely noisy: hidden 8192 at batch 1 measured 0.94x, 0.90x and 0.89x across three repetitions. Set to `1` to opt in. On GB10/CUDA the same harness measures 1.02x to 2.06x with every cell at or above parity, a different picture from Apple Silicon that has not been acted on; flipping the default is its own decision with its own evidence bar. See `docs/benchmark_results/fused-norm-rope-m1ultra-2026-07-31.md` and `docs/benchmark_results/fused-rope-cuda-trig-gb10-2026-08-06.md`. On ROCm it is **on by default** (lablup/mlxcel#2145): the kernel runs as a HIP port (lablup/mlxcel#2063) that calls `sincosf` as `rope.hip` does and reproduces the graph's FMA forms and f16 rounding, so its output is byte-identical to the graph, and Qwen2.5 7B decode on gfx1151 gained 0.6% with it alone and 1.7% with both fusions (medians of seven), inside run-to-run drift but never slower. Set to `0` on ROCm to opt out. Llama 3.1 still takes the graph path there because of its `rope_scaling` table; the stderr notice about that fires only when this variable is set to a truthy value, not for the ROCm default ([results](benchmark_results/rocm-fused-norm-rope-gfx1151-2026-10-05.md)). | `MLXCEL_COMPILED_QGELU_MLP` | `0`/`false`/`off`/`no` disable; any other value or unset enables | on | Compiles the affine-quantized GeGLU MLP (gate/up/gelu/down) into one `mx::compile` graph so the tanh-approx GELU's ~14 element-wise ops per layer collapse into a single fused `Compiled` primitive. Covers the Gemma family (gemma/gemma2/gemma3/gemma4). The group_size=64/bits=4 case was already compiled on every shape and is unchanged; other affine quantizations (notably the group_size=64/bits=8 MLP weights in Gemma 4 mixed-precision checkpoints, issue #680) are compiled only on the single-token decode call, because compiling the 8-bit prefill GEMM measured 8-9% slower and +0.2-0.7 GB peak on GB10 (the shapeless fused graph forces a decode-oriented qmm kernel onto the large prefill matmul). GB10 gemma-4-12b decode is weight-bandwidth-bound (~94% GPU-busy), so this primitive-count cut is throughput-neutral there (measured <1% from a ~530-primitive/step reduction); it still removes real CPU dispatch and helps op-count-bound backends. Set to `0` to force the op-at-a-time fallback (A/B + rollback). It also gates the NVFP4 scaled fused MLP variant (`MLXCEL_DISABLE_FUSED_GLOBAL_SCALE`), forcing its eager fold when set to `0`. | | `MLXCEL_DISABLE_FUSED_GLOBAL_SCALE` | `1`/`true`/`on`/`yes` (case-insensitive) disable; unset or any other value keeps the fold | off (fold on) | Rolls back the NVFP4 global-scale fold for Gemma 4 (issues #698/#705). By default the fused MLP and per-layer-input-gate C++ paths fold each per-projection `weight_scale_2` sidecar (from the direct ModelOpt transcode, issue #693/#697) into the fused kernel at the mathematically correct points: the gate scale before the GeGLU activation, the up scale on the up product, and the down scale on the fused output, each reproducing `apply_global_scale` byte-for-byte. Native NVFP4 prefill now uses a shape-specific scaled MLP graph, and standalone `UnifiedLinear` sidecar projections can apply qmm + global scale + dense bias through one C++ helper. When set, sidecar-carrying paths fall back to the op-at-a-time `UnifiedLinear::forward` scalar application (the pre-#698 bypass). Greedy temp-0 decode is token-identical across the two paths on `gemma-4-31b-it-nvfp4`; the fold removes element-wise dispatches on the gemma-4 path where CUDA graphs are disabled (#688). See `docs/benchmark_results/nvfp4-direct-transcode-gb10-2026-07-08.md` and `docs/benchmark_results/nvfp4-native-prefill-m1ultra-2026-07-09.md`. | | `MLXCEL_ONE_BIT_KERNEL` | `0`/`false`/`off`/`no` (case-insensitive) disable; unset or any other value keeps the kernels | on (Metal) | Affine 1-bit checkpoints (`quantization.bits == 1`, issue #1370). MLX has no 1-bit kernel, so the `quantized_matmul`, `quantized_linear_forward`, `dequantize` and `quantized_embedding` bridge primitives route a 1-bit affine triple to mlxcel's own fused Metal qmv (fewer than 16 rows) or tiled qmm kernel. Set to `0` to force the dequantize graph (`x @ dequant(W)^T`) instead, which is the parity oracle for the kernels and the path CUDA and ROCm always take. Read once in C++ at first use and cached for the process lifetime. | diff --git a/scripts/rocm_decode_profile.py b/scripts/rocm_decode_profile.py index 8ea9256bd..2ce063751 100644 --- a/scripts/rocm_decode_profile.py +++ b/scripts/rocm_decode_profile.py @@ -367,12 +367,13 @@ def reach(unit: str, config: dict, sampled: bool) -> tuple[tuple[str, ...], tupl if unit == "2063": if mt != "llama": return (), (), f"{mt} never calls fused_add_rms_norm or forward_fused_rope_append" - optin = ("add_rms_join_post_attn",) + (() if rope_table else ("rope_append",)) - note = ("both paths ship off (FUSED_ADD_RMSNORM_DEFAULT / FUSED_ROPE_APPEND_DEFAULT " - "false); only the post-attention join calls the fused norm") + reached = ("add_rms_join_post_attn",) + (() if rope_table else ("rope_append",)) + note = ("both paths are on by default in a rocm build (FUSED_ADD_RMSNORM_DEFAULT / " + "FUSED_ROPE_APPEND_DEFAULT = cfg!(feature = \"rocm\"), #2145); only the " + "post-attention join calls the fused norm") if rope_table: note += "; rope_scaling builds a frequency table, which the RoPE kernel cannot take" - return (), optin, note + return reached, reached, note if unit == "2064": if sampled: return ("sampler_tail",), ("sampler_tail",), "sampled run: the draw is the fallback chain" diff --git a/src/lib/mlxcel-core/src/layers.rs b/src/lib/mlxcel-core/src/layers.rs index de07c0040..a491c7cc0 100644 --- a/src/lib/mlxcel-core/src/layers.rs +++ b/src/lib/mlxcel-core/src/layers.rs @@ -9403,11 +9403,6 @@ mod tests { } } - /// The #905 kill switches are default-ON, which is the inverse of the - /// opt-in `MLXCEL_FUSED_QK_NORM` flag above, so they get their own pin: a - /// recognised disable string turns the fusion off, an unset variable leaves - /// the compiled-in default, and an unrecognised value must not silently - /// change the decode graph. #[test] fn fused_905_defaults_are_on_for_rocm_builds_only() { // #2145: a `rocm` build turns both fusions on by default; Metal and CUDA @@ -9447,6 +9442,11 @@ mod tests { } } + /// The #905 kill switches default per build (on for `rocm`, off elsewhere), + /// unlike the always-opt-in `MLXCEL_FUSED_QK_NORM` flag above, so they get + /// their own pin: a recognised disable string turns the fusion off, an + /// unset variable leaves the compiled-in default, and an unrecognised value + /// must not silently change the decode graph. #[test] fn fused_905_flags_follow_the_default_and_respect_explicit_values() { // Unset means the compiled-in default, whichever way it is set. Asserting diff --git a/tests/test_rocm_decode_profile.py b/tests/test_rocm_decode_profile.py index fe15891a6..076a219e5 100644 --- a/tests/test_rocm_decode_profile.py +++ b/tests/test_rocm_decode_profile.py @@ -386,12 +386,15 @@ def test_mamba_mixer(self): class ReachTests(unittest.TestCase): - def test_fused_norm_and_rope_ship_off_and_rope_tables_never_reach(self): + def test_fused_norm_and_rope_ship_on_for_rocm_and_rope_tables_never_reach(self): + # #2145: both fusions are on by default in a rocm build, so the roles + # they take over are reached with the shipped defaults. default, optin, _ = rdp.reach("2063", {"model_type": "llama", "rope_scaling": {"rope_type": "llama3"}}, False) - self.assertEqual(default, ()) + self.assertEqual(default, ("add_rms_join_post_attn",)) self.assertEqual(optin, ("add_rms_join_post_attn",)) - _, optin, _ = rdp.reach("2063", {"model_type": "llama"}, False) + default, optin, _ = rdp.reach("2063", {"model_type": "llama"}, False) + self.assertEqual(default, ("add_rms_join_post_attn", "rope_append")) self.assertEqual(optin, ("add_rms_join_post_attn", "rope_append")) def test_moe_reach_follows_the_caller(self): From 29f993d0d60f01353587f7c34bd4e07015f83382 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Wed, 7 Oct 2026 09:12:50 +0900 Subject: [PATCH 4/4] docs: add technical report for PR #2189 --- ...cm-fused-norm-rope-defaults-20261007.en.md | 145 ++++++++++++++++++ ...cm-fused-norm-rope-defaults-20261007.ko.md | 145 ++++++++++++++++++ 2 files changed, 290 insertions(+) create mode 100644 TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.en.md create mode 100644 TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.ko.md diff --git a/TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.en.md b/TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.en.md new file mode 100644 index 000000000..700c356a8 --- /dev/null +++ b/TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.en.md @@ -0,0 +1,145 @@ +# Technical Report: PR #2189 - Enable Fused Add-RMSNorm and RoPE-Append by Default on ROCm + +**Date**: 2026-10-07 + +**Status**: Implemented and verified on the gfx1151 host (ROCm/HIP); head `9874c3c4` (rebased onto origin/main `fc21784f`), PR open, pending merge. Closes #2145 (part of #1801). + +**Languages**: Rust (mlxcel-core `layers.rs`; mlxcel `models/llama3.rs`, `models/llama3_tests.rs`), Python (`scripts/rocm_decode_profile.py`, `tests/test_rocm_decode_profile.py`), Markdown (README, `docs/environment-variables.md`, `docs/installation.md`, `docs/benchmarks.md`, two results pages), CSV (two benchmark files) + +**Risk Level**: Low. The two HIP ports this PR turns on are byte-identical to the ROCm graph they replace, so the default change cannot move a logit on the models that were traced. Metal and CUDA builds compile the same `false` constants as before. The residual risk is in model families that now take the norm port by default on ROCm without a local checkpoint (Gemma, IQuest Loop Coder). + +## Executive Summary + +PR #2107 (closing #2063) ported `fused_add_rms_norm` and `fused_rope_qk_append` to HIP and left both opt-in, because on gfx1151 the decode gains (Llama 3.1 8B +0.3%, Qwen2.5 7B up to +1.7%) sat inside run-to-run drift. The maintainer then decided in #2145 to turn both on by default on ROCm only. The reasoning is asymmetric: the ports cost no accuracy (every on/off logit-trace pair was byte-identical) and no median was slower, so there is nothing to lose by defaulting them on, even though the gain cannot be distinguished from noise. This PR is the implementation of that decision, not a performance claim. + +The change is two lines of logic: `FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` in `src/lib/mlxcel-core/src/layers.rs` become `cfg!(feature = "rocm")`. The environment variables keep overriding in both directions. One side effect needed a real fix: the Llama 3.1 `rope_scaling` bypass notice asked the gate whether the RoPE kernel was requested, so with the default on it would have printed on every Llama 3.1 run on ROCm. A new `layers::fused_flag_explicit_value` lets the notice fire only for an explicit truthy `MLXCEL_FUSED_ROPE_APPEND`. + +The default-on build was re-measured on gfx1151: Llama 3.1 -0.2% and Qwen2.5 +0.8% against `=0`, both noise. `w8` logit traces are byte-identical between the default and both variables at `0`, and `rocprofv3` shows the default path dispatching the ported kernels. The unit's `make verify-rocm` at `9874c3c4` printed OK with 11,996 passed, 0 failed and 383 ignored; the orchestrator re-ran the full gate on the same head before merge. + +## 1. Problem Statement + +### 1.1 The state after #2107 + +`FUSED_ADD_RMSNORM_DEFAULT` and `FUSED_ROPE_APPEND_DEFAULT` were plain `false` for every backend. Their doc comments recorded why: on Metal (M1 Ultra) the op-level microbench and the decode measurement did not justify the fusions, and the RoPE fusion had one reproducible below-parity cell (hidden 8192, batch 1, 0.89x to 0.94x). The #2063 ROCm measurement was appended to those comments with the conclusion "not a clear win, so ROCm keeps the shared default". + +The ROCm numbers from #2107: + +| Model | Arm | Off | On | Change | +|---|---|---|---|---| +| Llama-3.1-8B | norm only (RoPE bypassed by `rope_scaling`) | 37.85 | 37.95 | +0.3% | +| Qwen2.5-7B | norm / RoPE / both | 46.53 | 47.23 / 46.80 / 47.32 | +1.5% / +0.6% / +1.7% | +| Qwen3-30B-A3B | control | | | +0.3% | + +On beat off in 6 of 7 Qwen2.5 rounds, but the size of the gain drifted between runs as much as the off arm's spread. + +### 1.2 The decision + +The maintainer's call in #2145 rests on correctness, not speed. Both ports produce output byte-identical to the graph on ROCm (every pair at `w1`, `w8` and `w1ctx512`, and 0/230 and 0/274 decided mismatches against Metal). A default that cannot change a token and was never slower in a median has no downside on that backend, and running the ported kernels by default also means they get exercised by every ROCm user instead of only by people who know the variables exist. On Metal the picture is different (the RoPE fusion has a measured regression cell, and the Metal norm kernel differs from the graph by a rounding step), so Metal and CUDA stay off. + +## 2. Change Summary + +| Area | Change | +|---|---| +| `mlxcel-core/src/layers.rs` | Both `FUSED_*_DEFAULT` constants become `cfg!(feature = "rocm")`; doc comments rewritten per backend with the #2107 numbers and the decision. `fused_flag_enabled_from` now delegates to the new public `fused_flag_explicit_value(value) -> Option`. New tests `fused_905_defaults_are_on_for_rocm_builds_only` and `fused_flag_explicit_value_separates_a_setting_from_the_default` | +| `src/models/llama3.rs` | `fused_rope_launcher_requested` uses the new `fused_rope_append_explicitly_requested(value)` for `MLXCEL_FUSED_ROPE_APPEND` instead of the gate; notice and launcher-order doc comments updated | +| `src/models/llama3_tests.rs` | `the_rope_append_bypass_notice_needs_an_explicit_truthy_value` | +| `scripts/rocm_decode_profile.py`, `tests/test_rocm_decode_profile.py` | The #2063 roles count as reached with the shipped defaults; the test asserts the default tuple instead of `()` | +| Docs | README ROCm bullet; `environment-variables.md` Default column "off (Metal, CUDA) / on (ROCm)"; ROCm row in `installation.md`; `benchmarks.md`; the note in `rocm-correctness-gfx1151-2026-09-30.md`; a Decision section with the re-measurement in `rocm-fused-norm-rope-gfx1151-2026-10-05.md` | +| `benchmarks/` | `rocm_strixhalo-gfx1151_2026-10-07_fused-default-{on,off}.csv` | + +Three commits: the default flip with the notice fix and tests (`c1309a0e`), the re-measurement (`3f5783c6`), and README, docs and decode-profile alignment (`9874c3c4`). 13 files, 197 insertions and 65 deletions. + +## 3. Design + +### 3.1 `cfg!(feature = "rocm")` on the constants + +```rust +pub(crate) const FUSED_ADD_RMSNORM_DEFAULT: bool = cfg!(feature = "rocm"); +pub(crate) const FUSED_ROPE_APPEND_DEFAULT: bool = cfg!(feature = "rocm"); +``` + +The `rocm` feature here is mlxcel-core's own (`src/lib/mlxcel-core/Cargo.toml`), which the root `rocm` feature enables. A ROCm build compiles no Metal or CUDA backend, so the build flag already identifies the backend. This follows PR #2098, which set the ROCm `MLXCEL_FUSED_MOE_SGY` default to 2 under `#ifdef MLXCEL_BRIDGE_ROCM_BACKEND` on the C++ side for the same reason. + +The alternatives the issue ruled out: + +- **A runtime backend comparison** (`gpu_kernel_backend() == Rocm`) or a `metal_is_available() || cuda_is_available()` gate. `scripts/ci/check_kernel_port_dispatch.py` (the `verify-kernel-port-dispatch` gate) rejects those patterns in kernel-port dispatch code, and they would add a runtime branch to answer a question the compiler already knows the answer to. +- **A separate ROCm constant or a ROCm-only gate function.** That would duplicate the `OnceLock` gates and the parser and give the two backends a chance to drift. With `cfg!` there is one constant per fusion and one place to flip it, which keeps the "flip here" contract the doc comments describe. + +`cfg!` rather than `#[cfg]` items keeps both values type-checked on every build and lets the tests compare against `cfg!(feature = "rocm")` directly, so each backend's test run checks its own default. + +### 3.2 Env override in both directions + +`fused_flag_enabled_from` keeps its contract. `0`/`false`/`off`/`no` disable and `1`/`true`/`on`/`yes` enable on every backend (case-insensitive, trimmed); unset or unrecognised takes the compiled-in default, so a typo does not silently change the decode graph. On ROCm, `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0` restores the graph; on Metal and CUDA, `=1` still opts in. The gates still read each variable once into a `OnceLock`, so the setting must be in place before the first decode. + +The refactor splits the parser: `fused_flag_explicit_value(value)` returns `Some(false)`, `Some(true)` or `None`, and `fused_flag_enabled_from` is now `fused_flag_explicit_value(value).unwrap_or(default)`. Behavior is identical; the split exists so a caller can ask "did the user set this?" separately from "is it on?". + +### 3.3 The Llama 3.1 bypass notice + +Llama 3.1's `rope_scaling` builds a frequency table, which the fused RoPE kernel cannot take, so `Attention::forward` routes around it regardless of the flag. `report_fused_rope_bypass_once` prints a one-time stderr notice when a bypass reason holds and a variable asked for a fused RoPE launcher. Its doc comment promised that a user who never sets the variables never sees it. + +Before this PR, `fused_rope_launcher_requested` answered the `MLXCEL_FUSED_ROPE_APPEND` arm by asking the gate (`fused_rope_append_enabled()`), which was correct while the default was off: the gate was true only if the user had opted in, and asking the gate instead of checking presence kept `=0` from being reported as a lost kernel. With the default on in a ROCm build, the gate is true for every unset run, so every Llama 3.1 run on ROCm would have printed the notice. + +The fix asks a different question. `fused_rope_append_explicitly_requested(value)` is `fused_flag_explicit_value(value) == Some(true)`, so the notice fires only for an explicit truthy value. Unset, unrecognised and `0` stay silent on every backend. The other two variables in the list (`MLXCEL_FUSED_CAUSAL_PREFILL`, `MLXCEL_ENABLE_FUSED_QKV_SPLIT_ROPE`) are presence-enabled and unchanged. + +One consequence to note: an unrecognised value such as `MLXCEL_FUSED_ROPE_APPEND=maybe` on ROCm takes the default (on) but does not trigger the notice. That matches the intent, since the user did not clearly ask for the kernel. + +### 3.4 Tests + +- `fused_905_defaults_are_on_for_rocm_builds_only` asserts both constants equal `cfg!(feature = "rocm")`, then under `#[cfg(feature = "rocm")]` checks unset is on and `0` is off, and under `#[cfg(not(...))]` checks unset is off and `1` is on. Each backend's test run therefore pins its own default and the override in the direction that matters there. +- `fused_flag_explicit_value_separates_a_setting_from_the_default` covers `None`, unrecognised values (`""`, `maybe`, `2`), both spellings of each recognised set, and whitespace and case. +- `the_rope_append_bypass_notice_needs_an_explicit_truthy_value` covers the unset case every default ROCm Llama 3.1 run hits, plus the falsy and truthy sets. +- The existing `fused_905_flags_follow_the_default_and_respect_explicit_values` and the parity tests already read the constant rather than a literal, so they needed no change. + +### 3.5 The decode profile script + +`scripts/rocm_decode_profile.py` classifies which decode roles each ROCm kernel-port unit reaches. For #2063 it returned `()` for the default and listed the roles only as opt-in. It now returns the same tuple for both, `add_rms_join_post_attn` plus `rope_append` when the checkpoint has no `rope_scaling` table, with a note pointing at the `cfg!` constants and #2145. + +## 4. Re-measurement + +gfx1151, `scripts/bench_decode.sh` at pp512/tg128, default (no variables) against `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0`, arm order alternated per round, medians of 3. Every run went through `scripts/rocm_gpu_guard.sh` with the host-wide lock from #2146; 12 of 12 runs were clean on the first attempt. The measurement ran on the pre-rebase SHA of the first commit (`57b59629`). + +| Model | Default (on) | `=0` (off) | Change | Paired (on minus off), by round | +|---|---|---|---|---| +| Llama-3.1-8B-Instruct-4bit | 38.09 | 38.15 | -0.2% | +0.02 / -0.49 / -0.05 | +| Qwen2.5-7B-Instruct-4bit | 47.90 | 47.51 | +0.8% | +0.39 / +0.30 / -0.04 | + +Both changes are inside run-to-run noise. Llama 3.1 runs only the norm port, and its -0.2% comes from one off run at 38.58, above every other run of either arm. Qwen2.5 runs both ports and is slightly above off in two of three rounds. Neither number is a speedup claim, and neither argues against the decision, which never depended on a speedup. + +Qwen2.5 prefill was 2.7% to 6.9% higher with the default in every round (median 1609.96 against 1544.08). The #2107 run did not show this, and three rounds cannot separate it from drift, so the results page records it without drawing a conclusion. + +## 5. Verification + +All on gfx1151: + +- **Byte identity.** `logit_trace` `w8` for both models, default against both variables at `0`: the files are byte-identical (`cmp`). +- **The default path reaches the kernels.** `rocprofv3 --kernel-trace` on an 8-token `mlxcel generate` with no variables set: Qwen2.5 dispatched `fused_add_rms_norm` and `fused_rope_qk_append` 252 times each (28 layers x 9 forward passes) and no graph RoPE kernel; with both variables at `0` it dispatched `rope_single_1d` instead. Llama 3.1 dispatched the norm port 288 times (32 x 9) and kept `rope_single_freqs_1d` / `rope_freqs`. This is what proves the flip took effect in the binary, not only in the constant. +- **Bypass notice.** Absent with the default on Llama 3.1, present with `MLXCEL_FUSED_ROPE_APPEND=1`, absent with `=0`. +- **Targeted tests.** `layers::tests::fused`, `fused_norm_parity_tests` and `fused_rope_parity_tests` under `--features rocm`: 32 passed. `models::llama3::`: 27 passed, including `greedy_decode_is_token_identical_to_the_unfused_baseline`, which also passes with both variables at `0`. `tests.test_rocm_decode_profile`: 26 passed. +- **Script gates.** `make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt`: pass. `verify-kernel-port-dispatch` passing is the check that no runtime backend comparison was introduced. +- **Unit's `make verify-rocm`** at `9874c3c4` (rebased onto `fc21784f`): OK. Clippy and the ROCm smoke test passed; 11,996 passed, 0 failed, 383 ignored across 152 test binaries. +- **Orchestrator's `make verify-rocm`**: re-run on the same head before merge. + +Not verified: Metal and CUDA, which are not available on this host. On those builds both constants evaluate to `false` exactly as before, and the notice change only narrows when it prints. + +## 6. Technical Decisions + +- **Default on for ROCm on correctness grounds, not speed.** A byte-identical port that is never slower in a median has no downside; the decision is explicitly not a claim that the fusion is faster. +- **Keep Metal and CUDA off.** The Metal measurements included a reproducible regression cell for the RoPE fusion, and the Metal norm kernel can differ from its graph by the rounding of the residual sum. Nothing in #2107 or this PR changes that evidence. +- **The build feature is the backend.** `cfg!(feature = "rocm")` follows PR #2098's precedent, adds no runtime branch, and keeps `verify-kernel-port-dispatch` passing. +- **One constant per fusion, one parser.** No ROCm-only duplicate of the gate, so the two backends cannot drift apart in parsing or caching. +- **Separate "explicitly set" from "enabled".** `fused_flag_explicit_value` exists because the bypass notice needs the first question and the gate answers the second. Once the default differs by backend, they are no longer the same. +- **Re-measure the shipped default, not the opt-in.** The new run compares the default against `=0`, which is the comparison a ROCm user now faces, and records it next to the original #2107 numbers. + +## 7. Residual Risks and Follow-ups + +- **Unverified model families.** Gemma and IQuest Loop Coder also call `fused_add_rms_norm` (`src/models/gemma.rs`, `src/models/iquestloopcoder.rs`) and now take the norm port by default on ROCm, but no checkpoint for either is available on this host. Gemma's `(1 + w)` weight convention is covered only by the tolerance tests in `fused_norm_parity_tests`, not by a logit trace or a generation run. VLMs whose text backbone is `Llama3Model` or the Qwen2 path also inherit the defaults; they share the traced Llama/Qwen2.5 code path but were not run individually. A trace on a Gemma checkpoint is the obvious follow-up. +- **Wave64 (CDNA) is untested.** All measurements and traces are on RDNA 3.5 gfx1151 (wave32). The default now applies to any ROCm target the build supports. +- **Qwen2.5 prefill shift.** The 2.7% to 6.9% prefill increase is unexplained and may be drift; a longer run would settle it. +- **The gain is noise-level.** If a later ROCm change makes the graph path faster than the ports, the default should be revisited. `MLXCEL_FUSED_ADD_RMSNORM=0` / `MLXCEL_FUSED_ROPE_APPEND=0` remain the kill switches, and the constants remain the one place to flip. + +## 8. Learning Points + +- **A byte-identical port changes the default question.** When a replacement kernel cannot change a single logit, "is it clearly faster?" is the wrong bar; "is it ever slower?" is enough, and the default can follow. +- **A notice tied to a gate changes meaning when the default changes.** Asking the gate was right while the gate meant "the user opted in". Flipping the default silently changed what the gate meant, and anything that inferred user intent from it had to switch to reading the explicit value. +- **Prove the default path in the binary.** Unit tests pin the constant, but the `rocprofv3` dispatch counts (252 and 288) are what show a default build actually launches the ported kernels and that Llama 3.1 still keeps the graph RoPE. +- **Let the compiler answer build-time questions.** When the build flag already determines the backend, `cfg!` is simpler and safer than a runtime backend check, and the repository's dispatch gate enforces that pattern. diff --git a/TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.ko.md b/TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.ko.md new file mode 100644 index 000000000..94cf7b724 --- /dev/null +++ b/TECHNICAL_REPORTS/2189-rocm-fused-norm-rope-defaults-20261007.ko.md @@ -0,0 +1,145 @@ +# 기술 보고서: PR #2189 - ROCm에서 Fused Add-RMSNorm과 RoPE-Append를 기본으로 켜기 + +**날짜**: 2026-10-07 + +**상태**: gfx1151 호스트(ROCm/HIP)에서 구현 및 검증 완료. 헤드 `9874c3c4`(origin/main `fc21784f` 위로 리베이스), PR 열림, 머지 대기. #2145를 닫음(#1801의 일부). + +**언어**: Rust(mlxcel-core `layers.rs`; mlxcel `models/llama3.rs`, `models/llama3_tests.rs`), Python(`scripts/rocm_decode_profile.py`, `tests/test_rocm_decode_profile.py`), Markdown(README, `docs/environment-variables.md`, `docs/installation.md`, `docs/benchmarks.md`, 결과 페이지 두 개), CSV(벤치마크 파일 두 개) + +**위험도**: 낮음. 이 PR이 켜는 두 HIP 포트는 대체하는 ROCm 그래프와 바이트 단위로 동일하므로, 트레이스를 뜬 모델에서는 기본값 변경이 logit을 하나도 바꿀 수 없습니다. Metal과 CUDA 빌드는 이전과 같은 `false` 상수를 컴파일합니다. 남은 위험은 로컬 체크포인트 없이 ROCm에서 기본으로 norm 포트를 타게 된 모델 계열(Gemma, IQuest Loop Coder)에 있습니다. + +## 요약 + +PR #2107(#2063 해결)은 `fused_add_rms_norm`과 `fused_rope_qk_append`를 HIP로 포팅했지만 둘 다 opt-in으로 남겼습니다. gfx1151에서 디코드 이득(Llama 3.1 8B +0.3%, Qwen2.5 7B 최대 +1.7%)이 실행 간 편차 안에 있었기 때문입니다. 이후 메인테이너는 #2145에서 둘 다 ROCm에서만 기본으로 켜기로 결정했습니다. 근거는 비대칭입니다. 포트는 정확도 비용이 없고(on/off logit 트레이스 쌍이 모두 바이트 동일), 느려진 중앙값이 하나도 없었으므로, 이득을 노이즈와 구분할 수 없더라도 기본으로 켜서 잃을 것이 없습니다. 이 PR은 그 결정을 구현한 것이지 성능 향상 주장이 아닙니다. + +로직 변경은 두 줄입니다. `src/lib/mlxcel-core/src/layers.rs`의 `FUSED_ADD_RMSNORM_DEFAULT`와 `FUSED_ROPE_APPEND_DEFAULT`가 `cfg!(feature = "rocm")`이 됩니다. 환경 변수는 여전히 양방향으로 덮어씁니다. 부작용 하나는 실제로 고쳐야 했습니다. Llama 3.1 `rope_scaling` 우회 안내는 RoPE 커널이 요청됐는지를 게이트에 물었기 때문에, 기본값이 켜지면 ROCm의 모든 Llama 3.1 실행에서 출력됐을 것입니다. 새 `layers::fused_flag_explicit_value`로 이 안내는 `MLXCEL_FUSED_ROPE_APPEND`가 명시적으로 참 값일 때만 출력됩니다. + +기본값이 켜진 빌드를 gfx1151에서 다시 측정했습니다. `=0` 대비 Llama 3.1 -0.2%, Qwen2.5 +0.8%로 둘 다 노이즈입니다. `w8` logit 트레이스는 기본값과 두 변수 모두 `0`인 경우가 바이트 동일하고, `rocprofv3`는 기본 경로가 포팅된 커널을 디스패치함을 보여 줍니다. `9874c3c4`에서 유닛의 `make verify-rocm`은 OK를 출력했고 11,996 통과, 0 실패, 383 무시였습니다. 오케스트레이터는 머지 전에 같은 헤드에서 전체 게이트를 다시 실행했습니다. + +## 1. 문제 정의 + +### 1.1 #2107 이후의 상태 + +`FUSED_ADD_RMSNORM_DEFAULT`와 `FUSED_ROPE_APPEND_DEFAULT`는 모든 백엔드에서 그냥 `false`였습니다. 문서 주석에 이유가 기록되어 있었습니다. Metal(M1 Ultra)에서는 op 수준 마이크로벤치와 디코드 측정 모두 fusion을 정당화하지 못했고, RoPE fusion에는 재현되는 패리티 미만 셀이 하나 있었습니다(hidden 8192, batch 1, 0.89x에서 0.94x). #2063의 ROCm 측정은 "확실한 이득이 아니므로 ROCm도 공통 기본값을 유지한다"는 결론과 함께 그 주석에 덧붙여졌습니다. + +#2107의 ROCm 수치: + +| 모델 | 구성 | Off | On | 변화 | +|---|---|---|---|---| +| Llama-3.1-8B | norm만(RoPE는 `rope_scaling` 때문에 우회) | 37.85 | 37.95 | +0.3% | +| Qwen2.5-7B | norm / RoPE / 둘 다 | 46.53 | 47.23 / 46.80 / 47.32 | +1.5% / +0.6% / +1.7% | +| Qwen3-30B-A3B | 대조군 | | | +0.3% | + +Qwen2.5의 7개 라운드 중 6개에서 on이 off보다 높았지만, 이득의 크기는 off 쪽 편차만큼 실행마다 흔들렸습니다. + +### 1.2 결정 + +#2145에서 메인테이너의 판단은 속도가 아니라 정확성에 근거합니다. 두 포트는 ROCm에서 그래프와 바이트 동일한 출력을 냅니다(`w1`, `w8`, `w1ctx512`의 모든 쌍, Metal 대비 판정 불일치 0/230과 0/274). 토큰을 바꿀 수 없고 중앙값 기준으로 한 번도 느리지 않았던 기본값은 그 백엔드에서 손해가 없으며, 기본으로 켜면 변수의 존재를 아는 사람만이 아니라 모든 ROCm 사용자가 포팅된 커널을 실행하게 됩니다. Metal은 상황이 다릅니다(RoPE fusion에 측정된 회귀 셀이 있고, Metal norm 커널은 반올림 한 단계만큼 그래프와 다를 수 있습니다). 그래서 Metal과 CUDA는 꺼진 상태로 둡니다. + +## 2. 변경 요약 + +| 영역 | 변경 | +|---|---| +| `mlxcel-core/src/layers.rs` | 두 `FUSED_*_DEFAULT` 상수가 `cfg!(feature = "rocm")`이 됨. 문서 주석을 백엔드별로 다시 쓰고 #2107 수치와 결정을 기록. `fused_flag_enabled_from`은 새 공개 함수 `fused_flag_explicit_value(value) -> Option`에 위임. 새 테스트 `fused_905_defaults_are_on_for_rocm_builds_only`, `fused_flag_explicit_value_separates_a_setting_from_the_default` | +| `src/models/llama3.rs` | `fused_rope_launcher_requested`가 `MLXCEL_FUSED_ROPE_APPEND`에 대해 게이트 대신 새 `fused_rope_append_explicitly_requested(value)`를 사용. 안내와 런처 순서 문서 주석 갱신 | +| `src/models/llama3_tests.rs` | `the_rope_append_bypass_notice_needs_an_explicit_truthy_value` | +| `scripts/rocm_decode_profile.py`, `tests/test_rocm_decode_profile.py` | #2063 역할이 기본값으로 도달하는 것으로 집계됨. 테스트는 `()` 대신 기본 튜플을 검사 | +| 문서 | README ROCm 항목, `environment-variables.md` 기본값 열 "off (Metal, CUDA) / on (ROCm)", `installation.md`의 ROCm 행, `benchmarks.md`, `rocm-correctness-gfx1151-2026-09-30.md`의 메모, `rocm-fused-norm-rope-gfx1151-2026-10-05.md`에 재측정을 담은 Decision 섹션 | +| `benchmarks/` | `rocm_strixhalo-gfx1151_2026-10-07_fused-default-{on,off}.csv` | + +커밋은 세 개입니다. 안내 수정과 테스트를 포함한 기본값 전환(`c1309a0e`), 재측정(`3f5783c6`), README, 문서, 디코드 프로파일 정렬(`9874c3c4`). 13개 파일, 197줄 추가, 65줄 삭제입니다. + +## 3. 설계 + +### 3.1 상수에 `cfg!(feature = "rocm")` + +```rust +pub(crate) const FUSED_ADD_RMSNORM_DEFAULT: bool = cfg!(feature = "rocm"); +pub(crate) const FUSED_ROPE_APPEND_DEFAULT: bool = cfg!(feature = "rocm"); +``` + +여기의 `rocm` feature는 mlxcel-core 자체의 것(`src/lib/mlxcel-core/Cargo.toml`)이고, 루트 `rocm` feature가 이를 켭니다. ROCm 빌드에는 Metal이나 CUDA 백엔드가 컴파일되지 않으므로 빌드 플래그가 곧 백엔드를 식별합니다. 이는 같은 이유로 C++ 쪽에서 `#ifdef MLXCEL_BRIDGE_ROCM_BACKEND` 아래 ROCm `MLXCEL_FUSED_MOE_SGY` 기본값을 2로 정한 PR #2098을 따른 것입니다. + +이슈가 배제한 대안: + +- **런타임 백엔드 비교**(`gpu_kernel_backend() == Rocm`) 또는 `metal_is_available() || cuda_is_available()` 게이트. `scripts/ci/check_kernel_port_dispatch.py`(`verify-kernel-port-dispatch` 게이트)가 커널 포트 디스패치 코드에서 이런 패턴을 거부하며, 컴파일러가 이미 아는 답을 위해 런타임 분기를 추가하게 됩니다. +- **별도의 ROCm 상수나 ROCm 전용 게이트 함수.** `OnceLock` 게이트와 파서를 중복시키고 두 백엔드가 어긋날 여지를 만듭니다. `cfg!`를 쓰면 fusion마다 상수 하나, 뒤집을 곳 하나가 유지되어 문서 주석이 말하는 "여기서 뒤집는다" 계약이 지켜집니다. + +`#[cfg]` 항목 대신 `cfg!`를 쓰면 모든 빌드에서 두 값이 타입 검사되고, 테스트가 `cfg!(feature = "rocm")`과 직접 비교할 수 있어 각 백엔드의 테스트 실행이 자기 기본값을 검사합니다. + +### 3.2 양방향 환경 변수 오버라이드 + +`fused_flag_enabled_from`의 계약은 그대로입니다. 모든 백엔드에서 `0`/`false`/`off`/`no`는 끄고 `1`/`true`/`on`/`yes`는 켭니다(대소문자 무시, 공백 제거). 설정되지 않았거나 인식할 수 없는 값은 컴파일된 기본값을 따르므로 오타가 디코드 그래프를 조용히 바꾸지 않습니다. ROCm에서는 `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0`으로 그래프로 돌아가고, Metal과 CUDA에서는 여전히 `=1`로 opt-in합니다. 게이트는 여전히 각 변수를 `OnceLock`에 한 번 읽으므로 설정은 첫 디코드 전에 되어 있어야 합니다. + +리팩터링은 파서를 나눕니다. `fused_flag_explicit_value(value)`는 `Some(false)`, `Some(true)`, `None` 중 하나를 반환하고, `fused_flag_enabled_from`은 이제 `fused_flag_explicit_value(value).unwrap_or(default)`입니다. 동작은 동일하며, 호출자가 "켜져 있는가"와 별개로 "사용자가 설정했는가"를 물을 수 있도록 나눈 것입니다. + +### 3.3 Llama 3.1 우회 안내 + +Llama 3.1의 `rope_scaling`은 주파수 테이블을 만들고, fused RoPE 커널은 이를 받을 수 없으므로 `Attention::forward`는 플래그와 상관없이 커널을 우회합니다. `report_fused_rope_bypass_once`는 우회 사유가 성립하고 어떤 변수가 fused RoPE 런처를 요청했을 때 stderr에 한 번 안내를 출력합니다. 그 문서 주석은 변수를 설정하지 않은 사용자는 안내를 보지 않는다고 약속했습니다. + +이 PR 이전에 `fused_rope_launcher_requested`는 `MLXCEL_FUSED_ROPE_APPEND`에 대해 게이트(`fused_rope_append_enabled()`)에 물었습니다. 기본값이 꺼져 있을 때는 이것이 맞았습니다. 게이트는 사용자가 opt-in했을 때만 참이었고, 존재 여부 대신 게이트에 물었기 때문에 `=0`을 잃어버린 커널로 보고하지 않았습니다. ROCm 빌드에서 기본값이 켜지면 설정하지 않은 모든 실행에서 게이트가 참이므로, ROCm의 모든 Llama 3.1 실행이 안내를 출력했을 것입니다. + +수정은 다른 질문을 던집니다. `fused_rope_append_explicitly_requested(value)`는 `fused_flag_explicit_value(value) == Some(true)`이므로 안내는 명시적인 참 값에서만 출력됩니다. 미설정, 인식 불가 값, `0`은 모든 백엔드에서 조용합니다. 목록의 나머지 두 변수(`MLXCEL_FUSED_CAUSAL_PREFILL`, `MLXCEL_ENABLE_FUSED_QKV_SPLIT_ROPE`)는 존재만으로 켜지는 변수이며 바뀌지 않았습니다. + +한 가지 결과: ROCm에서 `MLXCEL_FUSED_ROPE_APPEND=maybe` 같은 인식 불가 값은 기본값(on)을 따르지만 안내를 출력하지 않습니다. 사용자가 커널을 분명히 요청한 것이 아니므로 의도와 맞습니다. + +### 3.4 테스트 + +- `fused_905_defaults_are_on_for_rocm_builds_only`는 두 상수가 `cfg!(feature = "rocm")`과 같은지 확인한 뒤, `#[cfg(feature = "rocm")]`에서는 미설정이 on이고 `0`이 off인지, `#[cfg(not(...))]`에서는 미설정이 off이고 `1`이 on인지 확인합니다. 각 백엔드의 테스트 실행이 자기 기본값과 그 백엔드에서 중요한 방향의 오버라이드를 고정합니다. +- `fused_flag_explicit_value_separates_a_setting_from_the_default`는 `None`, 인식 불가 값(`""`, `maybe`, `2`), 인식되는 각 집합의 모든 철자, 공백과 대소문자를 다룹니다. +- `the_rope_append_bypass_notice_needs_an_explicit_truthy_value`는 기본 ROCm Llama 3.1 실행이 항상 겪는 미설정 경우와 거짓, 참 집합을 다룹니다. +- 기존 `fused_905_flags_follow_the_default_and_respect_explicit_values`와 패리티 테스트는 이미 리터럴이 아니라 상수를 읽으므로 수정이 필요 없었습니다. + +### 3.5 디코드 프로파일 스크립트 + +`scripts/rocm_decode_profile.py`는 각 ROCm 커널 포트 유닛이 어떤 디코드 역할에 도달하는지 분류합니다. #2063에 대해 기본값은 `()`를 반환하고 역할은 opt-in으로만 나열했습니다. 이제 둘 다 같은 튜플을 반환합니다. `add_rms_join_post_attn`에, 체크포인트에 `rope_scaling` 테이블이 없으면 `rope_append`가 더해지며, 메모는 `cfg!` 상수와 #2145를 가리킵니다. + +## 4. 재측정 + +gfx1151, `scripts/bench_decode.sh` pp512/tg128, 기본값(변수 없음) 대 `MLXCEL_FUSED_ADD_RMSNORM=0 MLXCEL_FUSED_ROPE_APPEND=0`, 라운드마다 순서 교대, 3회 중앙값. 모든 실행은 #2146의 호스트 전역 잠금과 함께 `scripts/rocm_gpu_guard.sh`를 거쳤고, 12회 모두 첫 시도에 깨끗했습니다. 측정은 첫 커밋의 리베이스 전 SHA(`57b59629`)에서 했습니다. + +| 모델 | 기본값(on) | `=0`(off) | 변화 | 쌍별 차이(on - off), 라운드별 | +|---|---|---|---|---| +| Llama-3.1-8B-Instruct-4bit | 38.09 | 38.15 | -0.2% | +0.02 / -0.49 / -0.05 | +| Qwen2.5-7B-Instruct-4bit | 47.90 | 47.51 | +0.8% | +0.39 / +0.30 / -0.04 | + +두 변화 모두 실행 간 노이즈 안에 있습니다. Llama 3.1은 norm 포트만 실행하며, -0.2%는 어느 쪽의 다른 모든 실행보다 높은 off 실행 하나(38.58)에서 나왔습니다. Qwen2.5는 두 포트를 모두 실행하며 세 라운드 중 두 번 off보다 약간 높았습니다. 어느 수치도 속도 향상 주장이 아니고, 애초에 속도 향상에 기대지 않은 결정에 반하는 것도 아닙니다. + +Qwen2.5 prefill은 모든 라운드에서 기본값 쪽이 2.7%에서 6.9% 높았습니다(중앙값 1609.96 대 1544.08). #2107 측정에서는 보이지 않았고 세 라운드로는 편차와 구분할 수 없으므로, 결과 페이지는 결론 없이 기록만 합니다. + +## 5. 검증 + +모두 gfx1151에서: + +- **바이트 동일성.** 두 모델의 `logit_trace` `w8`, 기본값 대 두 변수 모두 `0`: 파일이 바이트 동일(`cmp`). +- **기본 경로가 커널에 도달함.** 변수 없이 8토큰 `mlxcel generate`에 `rocprofv3 --kernel-trace`: Qwen2.5는 `fused_add_rms_norm`과 `fused_rope_qk_append`를 각각 252회(28 레이어 x 9 forward) 디스패치했고 그래프 RoPE 커널은 없었습니다. 두 변수가 `0`이면 대신 `rope_single_1d`를 디스패치했습니다. Llama 3.1은 norm 포트를 288회(32 x 9) 디스패치했고 `rope_single_freqs_1d` / `rope_freqs`를 유지했습니다. 이것이 상수만이 아니라 바이너리에서 전환이 효과를 냈다는 증거입니다. +- **우회 안내.** Llama 3.1에서 기본값일 때 없음, `MLXCEL_FUSED_ROPE_APPEND=1`일 때 있음, `=0`일 때 없음. +- **대상 테스트.** `--features rocm`에서 `layers::tests::fused`, `fused_norm_parity_tests`, `fused_rope_parity_tests`: 32 통과. `models::llama3::`: 27 통과, 두 변수가 `0`일 때도 통과하는 `greedy_decode_is_token_identical_to_the_unfused_baseline` 포함. `tests.test_rocm_decode_profile`: 26 통과. +- **스크립트 게이트.** `make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt`: 통과. `verify-kernel-port-dispatch` 통과는 런타임 백엔드 비교가 들어오지 않았다는 확인입니다. +- **유닛의 `make verify-rocm`** (`9874c3c4`, `fc21784f` 위로 리베이스): OK. Clippy와 ROCm 스모크 테스트 통과, 152개 테스트 바이너리에서 11,996 통과, 0 실패, 383 무시. +- **오케스트레이터의 `make verify-rocm`**: 머지 전에 같은 헤드에서 다시 실행. + +검증하지 않음: Metal과 CUDA(이 호스트에 없음). 그 빌드에서는 두 상수가 이전과 똑같이 `false`로 평가되고, 안내 변경은 출력 조건을 좁히기만 합니다. + +## 6. 기술적 선택과 그 이유 + +- **속도가 아니라 정확성 근거로 ROCm 기본값 켜기.** 바이트 동일하고 중앙값으로 한 번도 느리지 않은 포트는 손해가 없습니다. 이 결정은 fusion이 더 빠르다는 주장이 명시적으로 아닙니다. +- **Metal과 CUDA는 끈 채로 유지.** Metal 측정에는 RoPE fusion의 재현되는 회귀 셀이 있었고, Metal norm 커널은 잔차 합의 반올림만큼 그래프와 다를 수 있습니다. #2107도 이 PR도 그 근거를 바꾸지 않습니다. +- **빌드 feature가 곧 백엔드.** `cfg!(feature = "rocm")`은 PR #2098의 선례를 따르고, 런타임 분기를 추가하지 않으며, `verify-kernel-port-dispatch`를 통과시킵니다. +- **fusion마다 상수 하나, 파서 하나.** 게이트의 ROCm 전용 복제가 없으므로 두 백엔드가 파싱이나 캐싱에서 어긋날 수 없습니다. +- **"명시적으로 설정됨"과 "켜짐"을 분리.** 우회 안내는 앞의 질문이 필요하고 게이트는 뒤의 질문에 답하기 때문에 `fused_flag_explicit_value`가 생겼습니다. 기본값이 백엔드마다 달라지면 두 질문은 더 이상 같지 않습니다. +- **opt-in이 아니라 출시되는 기본값을 재측정.** 새 측정은 기본값을 `=0`과 비교합니다. 이제 ROCm 사용자가 마주하는 비교가 이것이며, 원래 #2107 수치 옆에 기록됩니다. + +## 7. 남은 위험과 후속 작업 + +- **검증되지 않은 모델 계열.** Gemma와 IQuest Loop Coder도 `fused_add_rms_norm`을 호출하고(`src/models/gemma.rs`, `src/models/iquestloopcoder.rs`) 이제 ROCm에서 기본으로 norm 포트를 타지만, 이 호스트에는 둘 다 체크포인트가 없습니다. Gemma의 `(1 + w)` 가중치 규약은 logit 트레이스나 생성 실행이 아니라 `fused_norm_parity_tests`의 허용 오차 테스트로만 다뤄집니다. 텍스트 백본이 `Llama3Model`이나 Qwen2 경로인 VLM도 기본값을 물려받습니다. 트레이스를 뜬 Llama/Qwen2.5 코드 경로를 공유하지만 개별 실행은 하지 않았습니다. Gemma 체크포인트로 트레이스를 뜨는 것이 분명한 후속 작업입니다. +- **Wave64(CDNA) 미검증.** 모든 측정과 트레이스는 RDNA 3.5 gfx1151(wave32)에서 했습니다. 기본값은 이제 빌드가 지원하는 모든 ROCm 타깃에 적용됩니다. +- **Qwen2.5 prefill 변화.** 2.7%에서 6.9%의 prefill 증가는 설명되지 않았고 편차일 수 있습니다. 더 긴 측정으로 판가름할 수 있습니다. +- **이득은 노이즈 수준.** 이후 ROCm 변경으로 그래프 경로가 포트보다 빨라지면 기본값을 다시 봐야 합니다. `MLXCEL_FUSED_ADD_RMSNORM=0` / `MLXCEL_FUSED_ROPE_APPEND=0`은 여전히 차단 스위치이고, 상수는 여전히 뒤집을 유일한 곳입니다. + +## 8. 학습 포인트 + +- **바이트 동일한 포트는 기본값의 질문을 바꾼다.** 대체 커널이 logit 하나도 바꿀 수 없다면 "확실히 더 빠른가"는 잘못된 기준이고 "한 번이라도 느린가"로 충분하며, 기본값은 그에 따를 수 있습니다. +- **게이트에 묶인 안내는 기본값이 바뀌면 의미가 바뀐다.** 게이트가 "사용자가 opt-in했다"를 뜻하는 동안에는 게이트에 묻는 것이 맞았습니다. 기본값을 뒤집자 게이트의 의미가 조용히 바뀌었고, 게이트에서 사용자 의도를 추론하던 코드는 명시적 값을 읽도록 바꿔야 했습니다. +- **기본 경로를 바이너리에서 증명하라.** 단위 테스트는 상수를 고정하지만, 기본 빌드가 실제로 포팅된 커널을 실행하고 Llama 3.1이 그래프 RoPE를 유지한다는 것을 보여 주는 것은 `rocprofv3` 디스패치 횟수(252와 288)입니다. +- **빌드 시점의 질문은 컴파일러에게 맡겨라.** 빌드 플래그가 이미 백엔드를 결정한다면 `cfg!`가 런타임 백엔드 검사보다 단순하고 안전하며, 저장소의 디스패치 게이트가 그 패턴을 강제합니다.