From 388795d4f1c1fc2653d7c5f48f382209469d55aa Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 18:42:25 +0900 Subject: [PATCH 1/7] chore(build): stop auto-detect at the shipped Blackwell spelling With MLX_CUDA_ARCHITECTURES unset, `sm_arch_with_suffix` appended CUDA's architecture-specific `a` suffix to every SM from 90 up, so a default `cargo build --features cuda` on a GB10 resolved the whole MLX target to `121a`. That is the whole-target form PR #1938 decided against: every translation unit compiles under `__CUDA_ARCH_SPECIFIC__`, CUTLASS keys `CUTLASS_ARCH_MMA_SM121A_ENABLED` on the same macro, and the per-source injection in `src/lib/mlx-cpp/CMakeLists.txt` never fires at all, because an entry already carrying `a` is recorded as architecture-specific and skipped rather than given a duplicate `--generate-code`. A local build and a shipped build then differ in the machine code their decode kernels run, not merely in architecture coverage. Both functions that spell a list from a detected capability now stop suffixing at a shared `FIRST_PLAIN_SM` boundary: `sm_arch_with_suffix` in the build script, and `CudaArchMismatch::suggested_architecture`, whose startup message was telling an operator on a Blackwell card to rebuild with `MLX_CUDA_ARCHITECTURES=121a`, the one value this repository rejects everywhere else. Hopper keeps `90a` because `release.yml` ships `90a`. The stale rationale for it is replaced rather than repeated: upstream MLX commit `44540d12` moved `qmm_sm80`, `qmm_sm90` and `gather_gemm` to runtime NVRTC compilation and deleted the `MLX_CUDA_SM90A_ENABLED` definition the old comments cited, and `jit_module.cpp` now derives the NVRTC `--gpu-architecture` from the running device, appending `a` itself from compute capability 9 up. `check_cuda_arch_lists.py` grows rule 6 for the auto-detect path, which rule 3 cannot see because no workflow writes that list down, and `check_cuda_arch_lists_test.sh` grows three cases for it. Refs #1943, #1934, PR #1938. --- scripts/ci/check_cuda_arch_lists.py | 143 ++++++++++++++++++++- scripts/ci/check_cuda_arch_lists_test.sh | 109 +++++++++++++++- src/lib/mlxcel-core/build.rs | 65 ++++++++-- src/lib/mlxcel-core/src/cuda_arch.rs | 29 ++++- src/lib/mlxcel-core/src/cuda_arch_tests.rs | 14 +- 5 files changed, 329 insertions(+), 31 deletions(-) diff --git a/scripts/ci/check_cuda_arch_lists.py b/scripts/ci/check_cuda_arch_lists.py index 84228a934..29312546f 100755 --- a/scripts/ci/check_cuda_arch_lists.py +++ b/scripts/ci/check_cuda_arch_lists.py @@ -39,9 +39,9 @@ 40-minute CUDA job. 3. No Blackwell entry (major >= 10) may carry the ``a`` suffix, because that is the whole-target form the per-source injection exists to avoid. Hopper's - ``90a`` is untouched: it is load-bearing for MLX's own quantized-kernel gate - and cannot reach these converters anyway, which also require compute - capability 10.0. + ``90a`` is untouched: these converters require compute capability 10.0 and so + cannot be reached from 9.0 however it is spelled, and ``90a`` is what the + release lists ship. Then, once for the repository: @@ -52,6 +52,18 @@ ``docs/installation.md``. Documentation drift is how the original defect stayed invisible: the document names the architectures the published archives carry, and nothing made it move when the workflow did. +6. The two places that spell an architecture list from a detected capability, + rather than reading one out of a workflow, must stop suffixing at the same + Blackwell boundary. A list nobody wrote down is not covered by rule 3: with + ``MLX_CUDA_ARCHITECTURES`` unset, ``build.rs`` auto-detects from + ``nvidia-smi``, and on a Blackwell host an unbounded suffix rule resolves the + whole target to ``121a``, which is the shape rule 3 rejects and which stops + the per-source injection from firing at all (an entry already carrying ``a`` + is skipped, because a duplicate ``--generate-code`` is an nvcc error). The + startup mismatch message spells one the same way, and it is the only value + the project hands an operator to paste. Guarded statically here because the + job has no CUDA toolkit, and because ``cargo test`` does not run build + scripts (lablup/mlxcel#1943). """ from __future__ import annotations @@ -78,6 +90,42 @@ #: runtime for the startup mismatch check. ENTRY = re.compile(r"^(?P\d{2,})(?P[af]?)(?P-real|-virtual)?$") +#: The Rust functions that spell an architecture list from a detected compute +#: capability instead of reading one out of a workflow, as (path, function, +#: consequence) triples relative to the repository root. Both must stop +#: suffixing at ``FIRST_PLAIN_SM``; see rule 6 in the module docstring. The +#: consequence is what an unbounded rule actually does at that site, so the +#: failure names the damage rather than the style. +SUFFIX_RULES = ( + ( + Path("src/lib/mlxcel-core/build.rs"), + "sm_arch_with_suffix", + "a Blackwell host with MLX_CUDA_ARCHITECTURES unset then builds the whole MLX target as " + "`121a`: every translation unit compiles under __CUDA_ARCH_SPECIFIC__, CUTLASS keys " + "CUTLASS_ARCH_MMA_SM121A_ENABLED on the same macro, and the per-source injection in " + "src/lib/mlx-cpp/CMakeLists.txt never fires at all, because an entry already carrying `a` " + "is skipped rather than given a duplicate --generate-code. That is the whole-target shape " + "rule 3 rejects in a workflow, reached by a path no workflow names", + ), + ( + Path("src/lib/mlxcel-core/src/cuda_arch.rs"), + "suggested_architecture", + "the startup mismatch message then tells an operator on a Blackwell card to rebuild with " + "`MLX_CUDA_ARCHITECTURES=121a`, which is the one value this repository rejects everywhere " + "else and the only architecture list the project ever hands someone to paste", + ), +) + +#: The shared boundary constant, which must equal ``BLACKWELL_MAJOR * 10``. It +#: is declared once per file because a build script cannot import from the crate +#: it builds, so this check is what keeps the two copies equal. +BOUNDARY_CONST = re.compile(r"const\s+FIRST_PLAIN_SM\s*:\s*u32\s*=\s*(?P\d+)\s*;") + +#: A suffix rule with no upper bound, i.e. the pre-#1943 shape. Matches both +#: spellings the two functions used: a match guard (``Ok(n) if n >= 90 =>``) and +#: an ``if`` expression (``if sm >= 90 { "a" }``). +UNBOUNDED_SUFFIX = re.compile(r"\b[a-z_]+\s*>=\s*(?P\d+)") + @dataclass(frozen=True) class Entry: @@ -207,6 +255,94 @@ def check_injection(cmakelists: Path) -> list[str]: ] +def _function_body(text: str, name: str) -> str | None: + """The braced body of ``fn ``, or ``None`` when it is not there. + + Brace counting rather than indentation, so it reads a free function and an + ``impl`` method the same way. Safe for these two bodies because neither + holds an unbalanced brace inside a string literal; ``format!("{sm}a")`` and + friends are balanced. + """ + start = text.find(f"fn {name}") + if start < 0: + return None + opening = text.find("{", start) + if opening < 0: + return None + depth = 0 + for index in range(opening, len(text)): + if text[index] == "{": + depth += 1 + elif text[index] == "}": + depth -= 1 + if depth == 0: + return text[opening : index + 1] + return None + + +def check_auto_detect(root: Path) -> list[str]: + """The suffix rule the auto-detected and suggested lists are spelled with. + + Rule 3 reads workflow files, so it cannot see the list a developer's own + machine produces when ``MLX_CUDA_ARCHITECTURES`` is unset. That list comes + from the functions named in ``SUFFIX_RULES``, and this is what holds them to + the same Blackwell boundary the shipped lists use. + """ + problems: list[str] = [] + expected_boundary = BLACKWELL_MAJOR * 10 + for relative, function, consequence in SUFFIX_RULES: + path = root / relative + if not path.exists(): + problems.append( + f"{relative} does not exist, so the {function} suffix rule cannot be checked. " + "If it moved, point SUFFIX_RULES at the new location rather than dropping the " + "check: an unguarded rule is how a local build silently stopped matching a " + "shipped one (issue #1943)." + ) + continue + text = path.read_text() + + boundary = BOUNDARY_CONST.search(text) + if boundary is None: + problems.append( + f"{relative} does not define `FIRST_PLAIN_SM`, the boundary {function} stops " + f"appending CUDA's architecture-specific `a` suffix at. Declare it as " + f"`const FIRST_PLAIN_SM: u32 = {expected_boundary};` (issue #1943)." + ) + elif int(boundary.group("value")) != expected_boundary: + problems.append( + f"{relative} sets FIRST_PLAIN_SM to {boundary.group('value')}, but the " + f"converter gate in nvfp4_quantize.cuh is `__CUDA_ARCH__ >= {expected_boundary}0`, " + f"so the boundary is {expected_boundary}. A mismatch makes an auto-detected build " + "disagree with the release lists on which architectures are named plainly " + "(issue #1943)." + ) + + body = _function_body(text, function) + if body is None: + problems.append( + f"{relative} no longer defines `{function}`, which is where a detected compute " + "capability is spelled as an architecture-list entry. This check cannot verify " + "the Blackwell boundary without it (issue #1943)." + ) + continue + + if "FIRST_PLAIN_SM" not in body: + problems.append( + f"{relative}: `{function}` does not bound its suffix rule with FIRST_PLAIN_SM. " + f"Unbounded, it appends `a` to Blackwell too, and {consequence}. See issue #1943." + ) + unbounded = UNBOUNDED_SUFFIX.search(body) + if unbounded is not None: + problems.append( + f"{relative}: `{function}` still compares with `>= {unbounded.group('threshold')}` " + "and no upper bound, which suffixes every architecture from there up. Use a " + "bounded range ending at FIRST_PLAIN_SM so Hopper keeps the `90a` the release " + "lists ship and Blackwell stays plain (issue #1943)." + ) + return problems + + def check_docs(release_lists: list[ArchList], installation: Path) -> list[str]: """Release lists that `docs/installation.md` does not quote verbatim.""" if not installation.exists(): @@ -254,6 +390,7 @@ def main() -> int: failures.append(f"{relative}:{arch_list.line}: {problem}") failures.extend(check_injection(args.root / "src" / "lib" / "mlx-cpp" / "CMakeLists.txt")) + failures.extend(check_auto_detect(args.root)) release_lists = [a for a in arch_lists if a.path.name == "release.yml"] failures.extend(check_docs(release_lists, args.root / "docs" / "installation.md")) diff --git a/scripts/ci/check_cuda_arch_lists_test.sh b/scripts/ci/check_cuda_arch_lists_test.sh index 6f4c3d336..df1be513d 100755 --- a/scripts/ci/check_cuda_arch_lists_test.sh +++ b/scripts/ci/check_cuda_arch_lists_test.sh @@ -22,12 +22,91 @@ trap 'rm -rf "$work"' EXIT failures=0 +# Writes the two Rust sources that spell an architecture list from a detected +# capability, in the shape named by $2 (build.rs) and $3 (cuda_arch.rs): +# `bounded` is the shipped rule, `unbounded` is the pre-#1943 one that suffixes +# Blackwell too, and `wrong-boundary` keeps the constant but moves it off the +# first Blackwell capability. Only the parts the checker reads are reproduced. +write_suffix_rules() { + local root="$1" build_rs="$2" cuda_arch_rs="$3" + mkdir -p "$root/src/lib/mlxcel-core/src" + case "$build_rs" in + bounded) + cat > "$root/src/lib/mlxcel-core/build.rs" <<'RS' +const FIRST_PLAIN_SM: u32 = 100; + +fn sm_arch_with_suffix(sm: &str) -> String { + match sm.parse::() { + Ok(n) if (90..FIRST_PLAIN_SM).contains(&n) => format!("{sm}a"), + _ => sm.to_string(), + } +} +RS + ;; + wrong-boundary) + cat > "$root/src/lib/mlxcel-core/build.rs" <<'RS' +const FIRST_PLAIN_SM: u32 = 120; + +fn sm_arch_with_suffix(sm: &str) -> String { + match sm.parse::() { + Ok(n) if (90..FIRST_PLAIN_SM).contains(&n) => format!("{sm}a"), + _ => sm.to_string(), + } +} +RS + ;; + *) + cat > "$root/src/lib/mlxcel-core/build.rs" <<'RS' +fn sm_arch_with_suffix(sm: &str) -> String { + match sm.parse::() { + Ok(n) if n >= 90 => format!("{sm}a"), + _ => sm.to_string(), + } +} +RS + ;; + esac + if [ "$cuda_arch_rs" = "bounded" ]; then + cat > "$root/src/lib/mlxcel-core/src/cuda_arch.rs" <<'RS' +const FIRST_PLAIN_SM: u32 = 100; + +impl CudaArchMismatch { + pub fn suggested_architecture(&self) -> String { + let (major, minor) = self.device; + let sm = major * 10 + minor; + let suffix = if (90..FIRST_PLAIN_SM).contains(&sm) { + "a" + } else { + "" + }; + format!("{sm}{suffix}") + } +} +RS + else + cat > "$root/src/lib/mlxcel-core/src/cuda_arch.rs" <<'RS' +impl CudaArchMismatch { + pub fn suggested_architecture(&self) -> String { + let (major, minor) = self.device; + let sm = major * 10 + minor; + let suffix = if sm >= 90 { "a" } else { "" }; + format!("{sm}{suffix}") + } +} +RS + fi +} + # Builds a root with one workflow carrying $2 as its architecture list, and an # installation document quoting $3 (default: the same list, so a case that is -# not about documentation drift never trips the documentation rule). +# not about documentation drift never trips the documentation rule). $5 and $6 +# select the shape of the two auto-detect suffix rules, both correct by default +# so a case that is not about them never trips rule 6. setup_root() { local root="$1" list="$2" documented="${3-$2}" injection="${4-present}" + local build_rs="${5-bounded}" cuda_arch_rs="${6-bounded}" mkdir -p "$root/.github/workflows" "$root/docs" "$root/src/lib/mlx-cpp" + write_suffix_rules "$root" "$build_rs" "$cuda_arch_rs" cat > "$root/.github/workflows/release.yml" </dev/null || rm -f "$wor printf 'jobs:\n build:\n runs-on: ubuntu-latest\n' > "$work/empty/.github/workflows/ci.yml" expect "no architecture pin at all fails" "$work/empty" 1 "no MLX_CUDA_ARCHITECTURES assignment" +# 11. The list nobody writes down. With MLX_CUDA_ARCHITECTURES unset, build.rs +# spells one from nvidia-smi, and the pre-#1943 rule suffixed everything from +# SM 90 up, so this host's own default build resolved to `121a`: the shape +# case 4 rejects in a workflow, reached by a path no workflow names. +setup_root "$work/autodetect" "90a;100;121" "90a;100;121" "present" "unbounded" +expect "unbounded auto-detect suffix rule fails" "$work/autodetect" 1 "sm_arch_with_suffix" + +# 12. The same rule on the runtime side. This one is user-facing: the startup +# mismatch message is the only place the project hands an operator an +# MLX_CUDA_ARCHITECTURES value to paste, and unbounded it pastes `121a`. +setup_root "$work/suggestion" "90a;100;121" "90a;100;121" "present" "bounded" "unbounded" +expect "unbounded startup suggestion fails" "$work/suggestion" 1 "MLX_CUDA_ARCHITECTURES=121a" + +# 13. A boundary that exists but sits somewhere other than the first Blackwell +# capability. `120` looks plausible and still suffixes sm_100. +setup_root "$work/boundary" "90a;100;121" "90a;100;121" "present" "wrong-boundary" +expect "misplaced plain-spelling boundary fails" "$work/boundary" 1 "sets FIRST_PLAIN_SM to 120" + +# The other half of the #1943 decision, that Hopper keeps the `90a` the release +# lists ship, is not checkable here: this script reads text and cannot evaluate +# the rule. `the_suggested_rebuild_matches_the_shipped_spelling` in +# src/lib/mlxcel-core/src/cuda_arch_tests.rs asserts the produced spelling for +# H100 and GB10 instead, and runs under `cargo test`. + if [ "$failures" -ne 0 ]; then echo "$failures case(s) failed" >&2 exit 1 diff --git a/src/lib/mlxcel-core/build.rs b/src/lib/mlxcel-core/build.rs index f43ae6242..ad2b2caba 100644 --- a/src/lib/mlxcel-core/build.rs +++ b/src/lib/mlxcel-core/build.rs @@ -413,13 +413,16 @@ fn build_mlx(expected_commit: &str, cuda_architectures: &str, rocm_architectures // honored verbatim (escape hatch); otherwise we auto-detect via nvidia-smi // and fall back to Hopper's sm_90a. // - // The `a` suffix is load-bearing: MLX only defines MLX_CUDA_SM90A_ENABLED - // (which compiles the dedicated Hopper `qmm_sm90` quantized kernel) when - // "90a" is in the arch list. MLX's own CMake appends that suffix for - // cc >= 90, but only inside its `if(NOT DEFINED MLX_CUDA_ARCHITECTURES)` - // branch. Because we always pass MLX_CUDA_ARCHITECTURES explicitly, that - // branch never runs, so we apply the same rule ourselves here and in - // detect_cuda_arch. See docs/installation.md (CUDA architecture selection). + // MLX's own CMake would spell a detected capability itself, appending the + // architecture-specific `a` suffix for cc >= 90, but only inside its + // `if(NOT DEFINED MLX_CUDA_ARCHITECTURES)` branch. We always pass the + // variable explicitly, so that branch never runs and `sm_arch_with_suffix` + // spells it here instead. It deliberately stops suffixing at + // `FIRST_PLAIN_SM`, where MLX's rule would not: on Blackwell the suffix + // belongs on one translation unit rather than on the whole target, which + // is what `src/lib/mlx-cpp/CMakeLists.txt` arranges and what the release + // lists carry. See `sm_arch_with_suffix` and docs/installation.md (CUDA + // architecture selection). // // The value is resolved in `main` by `resolve_cuda_architectures` and // passed in, so the list CMake compiles for and the list recorded in @@ -647,17 +650,51 @@ fn detect_cuda_arch() -> Option { } } -/// Append CUDA's architecture-specific `a` suffix for SM >= 90, mirroring MLX's -/// own CMake logic (`MLX_CUDA_ARCHITECTURES GREATER_EQUAL 90` -> append `a`). +/// First SM number this project names plainly even though CUDA offers an +/// architecture-specific spelling for it: Blackwell's `sm_100`. /// -/// The `a` suffix enables architecture-specific features (e.g. Hopper wgmma/TMA) -/// the dedicated quantized kernels rely on, and it gates MLX_CUDA_SM90A_ENABLED on -/// "90a" (not "90"). SM < 90 (e.g. Ampere sm_80/sm_86) has no `a` variant and is -/// returned unchanged. +/// Both release lists in `.github/workflows/release.yml` stop suffixing here, +/// and `scripts/ci/check_cuda_arch_lists.py` rejects any workflow list that +/// does not. Auto-detection uses the same boundary so that a local build and a +/// shipped build differ in which architectures they cover and never in what +/// machine code those architectures get (lablup/mlxcel#1943). +#[cfg(feature = "cuda")] +const FIRST_PLAIN_SM: u32 = 100; + +/// Spell one detected SM number the way the shipped architecture lists spell +/// it: `a`-suffixed for Hopper, plain from `FIRST_PLAIN_SM` up. +/// +/// Suffixing Blackwell is what that boundary exists to prevent. The whole +/// target would compile under `__CUDA_ARCH_SPECIFIC__`, CUTLASS keys +/// `CUTLASS_ARCH_MMA_SM121A_ENABLED` on the same macro, and every decode kernel +/// would get a second, architecture-specific image that the driver prefers on a +/// matching device. `src/lib/mlx-cpp/CMakeLists.txt` injects that image into +/// `fp_quantize.cu` alone instead, the one translation unit that can reach +/// MLX's hardware block-float converters, and that injection only fires for +/// plain entries: an entry already carrying `a` is recorded as +/// architecture-specific and skipped, because a duplicate `--generate-code` is +/// an nvcc error rather than a no-op. A plain `121` is therefore what makes the +/// shipped mechanism work (lablup/mlxcel#1934, PR #1938). +/// +/// Hopper keeps `90a` because `release.yml` ships `90a` and this function's job +/// is to agree with it. The suffix no longer gates anything at the current MLX +/// pin: upstream commit `44540d12` moved `qmm_sm80`, `qmm_sm90` and +/// `gather_gemm` to runtime NVRTC compilation and deleted the +/// `MLX_CUDA_SM90A_ENABLED` definition an earlier revision of this comment +/// cited, and `jit_module.cpp` now derives the NVRTC `--gpu-architecture` from +/// the running device, appending `a` itself from compute capability 9 up. +/// Cross-compiling the pinned tree at `90` and at `90a` agrees: the quantized +/// translation units, `qmm_sm90.cu` included, produce identical SASS either +/// way. `CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on +/// `__CUDA_ARCH_FEAT_SM90_ALL`, so a later pin can make the suffix matter +/// again, which is a second reason not to drop what the release list carries. +/// +/// SM < 90 (Ampere `sm_80`, `sm_86`) has no `a` variant and is returned +/// unchanged. See `docs/installation.md` (CUDA architecture selection). #[cfg(feature = "cuda")] fn sm_arch_with_suffix(sm: &str) -> String { match sm.parse::() { - Ok(n) if n >= 90 => format!("{sm}a"), + Ok(n) if (90..FIRST_PLAIN_SM).contains(&n) => format!("{sm}a"), _ => sm.to_string(), } } diff --git a/src/lib/mlxcel-core/src/cuda_arch.rs b/src/lib/mlxcel-core/src/cuda_arch.rs index 31d0b4cdd..e15c085e1 100644 --- a/src/lib/mlxcel-core/src/cuda_arch.rs +++ b/src/lib/mlxcel-core/src/cuda_arch.rs @@ -43,6 +43,15 @@ use std::fmt; use std::sync::OnceLock; +/// First SM number this project names plainly even though CUDA offers an +/// architecture-specific spelling for it: Blackwell's `sm_100`. +/// +/// The same boundary as `FIRST_PLAIN_SM` in `build.rs`, which auto-detection +/// spells architecture lists with. Duplicated rather than shared because a +/// build script cannot import from the crate it builds; the guard in +/// `scripts/ci/check_cuda_arch_lists.py` keeps the two from drifting apart. +const FIRST_PLAIN_SM: u32 = 100; + // ── Compiled architecture list ──────────────────────────────────────────────── /// The `MLX_CUDA_ARCHITECTURES` list this binary's MLX device code was @@ -325,16 +334,24 @@ impl CudaArchMismatch { /// The `MLX_CUDA_ARCHITECTURES` value that would build for this device. /// /// Applies the same rule `build.rs`'s `sm_arch_with_suffix` applies when it - /// auto-detects: CUDA's architecture-specific `a` suffix from SM 90 up. The - /// suffix is load-bearing on Hopper and newer, where MLX compiles its - /// dedicated quantized kernel only when the list says `90a` rather than - /// `90`. Suggesting anything else here would send an operator to a rebuild - /// that differs from the one auto-detection would have produced. + /// auto-detects: CUDA's architecture-specific `a` suffix on Hopper, plain + /// from Blackwell up. This message is the one place the project hands an + /// operator an `MLX_CUDA_ARCHITECTURES` value to paste, so a suffix here + /// that auto-detection would not produce sends them to a build the project + /// rejects everywhere else: `121a` compiles every translation unit + /// architecture-specific rather than letting `src/lib/mlx-cpp/CMakeLists.txt` + /// give that image to `fp_quantize.cu` alone, and + /// `scripts/ci/check_cuda_arch_lists.py` fails any workflow list spelled + /// that way (lablup/mlxcel#1934, #1943). #[must_use] pub fn suggested_architecture(&self) -> String { let (major, minor) = self.device; let sm = major * 10 + minor; - let suffix = if sm >= 90 { "a" } else { "" }; + let suffix = if (90..FIRST_PLAIN_SM).contains(&sm) { + "a" + } else { + "" + }; format!("{sm}{suffix}") } } diff --git a/src/lib/mlxcel-core/src/cuda_arch_tests.rs b/src/lib/mlxcel-core/src/cuda_arch_tests.rs index c4272c952..1d77b3931 100644 --- a/src/lib/mlxcel-core/src/cuda_arch_tests.rs +++ b/src/lib/mlxcel-core/src/cuda_arch_tests.rs @@ -474,10 +474,14 @@ fn the_mismatch_message_names_both_sides_and_a_working_rebuild() { } #[test] -fn the_suggested_rebuild_carries_the_hopper_suffix() { - // `90` and `90a` are not interchangeable: MLX only compiles its dedicated - // Hopper quantized kernel when the list says `90a`, so a suggestion that - // dropped the suffix would rebuild into a slower binary. +fn the_suggested_rebuild_matches_the_shipped_spelling() { + // This message is the only place the project hands an operator an + // MLX_CUDA_ARCHITECTURES value to paste, so it has to name the spelling the + // release workflow builds and auto-detection produces. Hopper keeps `90a`, + // which is what `release.yml` ships. Blackwell stays plain: `121a` would + // compile every translation unit architecture-specific instead of letting + // the per-source injection give that image to `fp_quantize.cu` alone, and + // the CI guard rejects a workflow list spelled that way (#1934, #1943). let suggest = |device| { CudaArchMismatch { device, @@ -488,5 +492,5 @@ fn the_suggested_rebuild_carries_the_hopper_suffix() { assert_eq!(suggest(V100), "70"); assert_eq!(suggest(A100), "80"); assert_eq!(suggest(H100), "90a"); - assert_eq!(suggest(GB10), "121a"); + assert_eq!(suggest(GB10), "121"); } From 1fa102b585e20dba9202b76a24172ac3144a163e Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 18:43:38 +0900 Subject: [PATCH 2/7] docs(installation): describe one CUDA architecture rule, not two The section stated the auto-detect rule as suffixing every SM from 90 up, citing a Hopper quantized-kernel gate, and then told operators not to set `a` on Blackwell and that building `121a` is a step backwards. Both passages described the same code path in opposite terms, and `check_docs` in the guard script only verifies that the release lists appear verbatim, so nothing caught it. It now states the rule once, the way the build script and the release workflow both spell it, and explains what `121a` costs and why the per-source injection makes it unnecessary. The Hopper paragraph records what is actually true at the current MLX pin: upstream `44540d12` moved the Hopper quantized kernel to runtime NVRTC compilation and removed the `MLX_CUDA_SM90A_ENABLED` definition this section used to cite, `90a` is kept because both release lists ship it, and `CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on `__CUDA_ARCH_FEAT_SM90_ALL` so a later pin can make it matter again. Refs #1943. --- docs/installation.md | 51 +++++++++++++++++++++++++++++--------------- 1 file changed, 34 insertions(+), 17 deletions(-) diff --git a/docs/installation.md b/docs/installation.md index 55a8297b9..6310a894a 100644 --- a/docs/installation.md +++ b/docs/installation.md @@ -164,18 +164,23 @@ CUDA_HOME=/opt/cuda cargo build --release --features cuda `src/lib/mlxcel-core/build.rs` reads `MLX_CUDA_ARCHITECTURES`. If it is unset, -the build script tries to detect the compute capability with `nvidia-smi` and -falls back to `90a` when detection fails. For SM 90 and above it appends CUDA's -architecture-specific `a` suffix (so `90` becomes `90a`), because the dedicated -Hopper quantized kernel (`qmm_sm90`) is only compiled when `90a` is in the arch -list. An explicitly set `MLX_CUDA_ARCHITECTURES` is used verbatim, so include the -suffix yourself for Hopper (`90a`). - -On Blackwell (`sm_100`, `sm_120`, `sm_121`) the `a` suffix decides something -else, and you should not set it. MLX compiles the hardware block-float -converters in `mlx/backend/cuda/quantized/nvfp4_quantize.cuh`, the ones that -issue `cvt.rn.satfinite.e2m1x2.f32`, only when nvcc is compiling for an -architecture-specific target, which it signals with `__CUDA_ARCH_SPECIFIC__`. +the build script detects the compute capability with `nvidia-smi` and spells it +the way the release workflow spells it: Hopper gets CUDA's architecture-specific +`a` suffix (`90` becomes `90a`), and Blackwell (`sm_100`, `sm_120`, `sm_121`) +stays plain. Detection failure falls back to `90a`. An auto-detected build and a +published one therefore differ in which architectures they cover and never in +what machine code those architectures get, which is the point: before issue +#1943 the rule suffixed everything from SM 90 up, so a default build on a GB10 +produced `121a` while the release archives for the same card carried `121`. + +An explicitly set `MLX_CUDA_ARCHITECTURES` is used verbatim, so spell it the same +way by hand: `90a` for Hopper, plain for Blackwell. The rest of this section is +why Blackwell is plain. + +On Blackwell the `a` suffix decides something else. MLX compiles the hardware +block-float converters in `mlx/backend/cuda/quantized/nvfp4_quantize.cuh`, the +ones that issue `cvt.rn.satfinite.e2m1x2.f32`, only when nvcc is compiling for +an architecture-specific target, which it signals with `__CUDA_ARCH_SPECIFIC__`. A plain `121` build therefore compiles them out, and both NVFP4 and MXFP4 quantization fall back to a scalar CUTLASS conversion sequence with nothing in the build output saying so. @@ -192,13 +197,25 @@ So the list stays plain and `src/lib/mlx-cpp/CMakeLists.txt` adds the architecture-specific image to `fp_quantize.cu` alone, which is the only translation unit that can reach those converters. You get the hardware path without asking for it, and without it reaching anything else. Building with -`121a` is a step backwards, not a step forwards; `121f` is worse still, because -it satisfies the converters' dispatcher gate but not their own gate and fails -to compile outright. +`121a` is a step backwards, not a step forwards, and it is also self-defeating: +the injection skips a capability the list already names architecture-specific, +because asking nvcc for the same `--generate-code` twice is an error rather than +a no-op. `121f` is worse still, because it satisfies the converters' dispatcher +gate but not their own gate and fails to compile outright. + +Hopper's `90a` is a different case: it is what both release lists ship, so the +auto-detected spelling matches it. At the current MLX pin the suffix no longer +gates anything on its own. Upstream commit `44540d12` moved `qmm_sm80`, +`qmm_sm90` and `gather_gemm` to runtime NVRTC compilation and removed the +`MLX_CUDA_SM90A_ENABLED` definition an earlier version of this section cited, +and `jit_module.cpp` now derives the NVRTC `--gpu-architecture` from the running +device, appending `a` itself from compute capability 9 up. Cross-compiling the +pinned tree at `90` and at `90a` produces identical SASS for the quantized +translation units. `CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on +`__CUDA_ARCH_FEAT_SM90_ALL`, so a later pin can make it matter again. ```bash -# Hopper / GH200-style target. The `a` suffix is required for the Hopper -# quantized kernel; plain `90` builds without it. +# Hopper / GH200-style target, spelled the way the release workflow spells it. MLX_CUDA_ARCHITECTURES=90a cargo build --release --features cuda # GB10 / DGX Spark-style target used by the release workflow. Plain: the From 8291de1a198f715001b5516ebf78ed5c335f4a4e Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 18:49:33 +0900 Subject: [PATCH 3/7] chore: state the Hopper cross-compile result precisely The first pass said the quantized translation units produce identical SASS at `90` and `90a`. Measured across all five, that was too strong in one direction and too weak in another. `qmm_sm90.cu`, `qmm.cu` and `qmm_sm80.cu` emit no device function at either spelling, because upstream `44540d12` moved those kernels to runtime NVRTC compilation; their objects carry 21 lines of fatbin header and zero kernel symbols. `qmv.cu` (2898 kernels) and `fp_qmv.cu` (54) do emit device code, and it is instruction-identical, but their cubins are not byte-identical: the `90a` build adds `EF_CUDA_ACCELERATORS` to the `.headerflags` line, which is the ELF flag marking a cubin architecture-specific. So the suffix changes no instruction at this pin, which is what the decision rests on, and the text now says that rather than implying the objects match byte for byte. Also drops the two remaining statements of the pre-#1943 rule that the new guard does not read: the inline comment in `detect_cuda_arch`, and the `90a is load-bearing` rationale on the x86_64 release job, which cited the same deleted macro. The architecture lists themselves are untouched. Refs #1943. --- .github/workflows/release.yml | 11 ++++++++--- docs/installation.md | 9 ++++++--- src/lib/mlxcel-core/build.rs | 11 +++++++---- 3 files changed, 21 insertions(+), 10 deletions(-) diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index c52e05d30..6629d1cd8 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -766,9 +766,14 @@ jobs: permissions: contents: write - # One fat binary covering Ampere and later. 90a is load-bearing: MLX gates - # its Hopper quantized kernel on "90a" being in the list (MLX_CUDA_SM90A_ENABLED), - # and build.rs forwards MLX_CUDA_ARCHITECTURES verbatim (no auto-suffix). + # One fat binary covering Ampere and later. Hopper is spelled 90a, and + # build.rs forwards MLX_CUDA_ARCHITECTURES verbatim, so this list is exactly + # what nvcc is asked for. At the current MLX pin that suffix gates nothing + # by itself: upstream 44540d12 moved the Hopper quantized kernel to runtime + # NVRTC compilation and deleted MLX_CUDA_SM90A_ENABLED, which the rationale + # here used to cite. It stays because CUTLASS_ARCH_MMA_SM90A_ENABLED still + # keys on __CUDA_ARCH_FEAT_SM90_ALL, so a later pin can make it matter, and + # because auto-detection now mirrors this list (issue #1943). # Blackwell stays plain for the reason spelled out in the aarch64 job # above: the hardware NVFP4/MXFP4 converters are added to `fp_quantize.cu` # alone by `src/lib/mlx-cpp/CMakeLists.txt` rather than to the whole target. diff --git a/docs/installation.md b/docs/installation.md index 6310a894a..b33ffaf9b 100644 --- a/docs/installation.md +++ b/docs/installation.md @@ -210,9 +210,12 @@ gates anything on its own. Upstream commit `44540d12` moved `qmm_sm80`, `MLX_CUDA_SM90A_ENABLED` definition an earlier version of this section cited, and `jit_module.cpp` now derives the NVRTC `--gpu-architecture` from the running device, appending `a` itself from compute capability 9 up. Cross-compiling the -pinned tree at `90` and at `90a` produces identical SASS for the quantized -translation units. `CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on -`__CUDA_ARCH_FEAT_SM90_ALL`, so a later pin can make it matter again. +pinned tree at `90` and at `90a` agrees: `qmm_sm90.cu`, `qmm.cu` and +`qmm_sm80.cu` emit no device function at either spelling, and `qmv.cu` and +`fp_qmv.cu` emit identical SASS apart from the `EF_CUDA_ACCELERATORS` header +flag that marks a cubin architecture-specific. +`CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on `__CUDA_ARCH_FEAT_SM90_ALL`, so +a later pin can make it matter again. ```bash # Hopper / GH200-style target, spelled the way the release workflow spells it. diff --git a/src/lib/mlxcel-core/build.rs b/src/lib/mlxcel-core/build.rs index ad2b2caba..c1c0af267 100644 --- a/src/lib/mlxcel-core/build.rs +++ b/src/lib/mlxcel-core/build.rs @@ -626,7 +626,8 @@ fn detect_cuda_arch() -> Option { .ok()?; let caps = String::from_utf8_lossy(&output.stdout); // Parse "X.Y" compute capabilities, convert to SM number (e.g. "9.0" -> "90"), - // and append the architecture-specific "a" suffix for cc >= 90 (e.g. "90a"). + // and let `sm_arch_with_suffix` spell each one the way the release lists + // spell it: "90" becomes "90a", Blackwell stays plain. let archs: Vec = caps .lines() .filter_map(|line| { @@ -683,9 +684,11 @@ const FIRST_PLAIN_SM: u32 = 100; /// `MLX_CUDA_SM90A_ENABLED` definition an earlier revision of this comment /// cited, and `jit_module.cpp` now derives the NVRTC `--gpu-architecture` from /// the running device, appending `a` itself from compute capability 9 up. -/// Cross-compiling the pinned tree at `90` and at `90a` agrees: the quantized -/// translation units, `qmm_sm90.cu` included, produce identical SASS either -/// way. `CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on +/// Cross-compiling the pinned tree at `90` and at `90a` agrees: `qmm_sm90.cu`, +/// `qmm.cu` and `qmm_sm80.cu` emit no device function at either spelling, and +/// `qmv.cu` and `fp_qmv.cu` emit identical SASS apart from the +/// `EF_CUDA_ACCELERATORS` header flag that marks a cubin architecture-specific. +/// `CUTLASS_ARCH_MMA_SM90A_ENABLED` still keys on /// `__CUDA_ARCH_FEAT_SM90_ALL`, so a later pin can make the suffix matter /// again, which is a second reason not to drop what the release list carries. /// From 420761a349202d592949d21aa6b2fd868ad67484 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 19:03:02 +0900 Subject: [PATCH 4/7] docs(build): say why the no-detection fallback stays 90a The issue tied the `90a` fallback literal to the Hopper decision, and the cross-compile shows the suffix buys no instruction at this pin, so leaving the literal in place needed a reason rather than silence. The reason is the same one that keeps Hopper suffixed at all: `release.yml` ships `90a`, and a fallback spelled `90` would hand a host without `nvidia-smi` a build shape no published archive uses, which is the divergence this issue closed on Blackwell. The doc comment on `resolve_cuda_architectures` said only that `90a` is the last resort. Refs #1943. --- src/lib/mlxcel-core/build.rs | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/src/lib/mlxcel-core/build.rs b/src/lib/mlxcel-core/build.rs index c1c0af267..8dc757a8a 100644 --- a/src/lib/mlxcel-core/build.rs +++ b/src/lib/mlxcel-core/build.rs @@ -598,10 +598,15 @@ fn cmake_bool_from_env(name: &str) -> Option<&'static str> { /// /// An explicitly set `MLX_CUDA_ARCHITECTURES` wins verbatim (the documented /// escape hatch); otherwise `nvidia-smi` detection decides, and `90a` is the -/// last resort. That fallback is why the runtime mismatch check exists: on a -/// host without `nvidia-smi` it produces a binary that cannot run on its own -/// build machine, and without the check the only symptom is an opaque CUDA -/// load failure at the first kernel launch. +/// last resort. It stays spelled `90a` for the same reason `sm_arch_with_suffix` +/// still suffixes Hopper: that is what `release.yml` ships, and a fallback that +/// said `90` would hand a host without `nvidia-smi` a build shape no published +/// archive uses, which is the divergence lablup/mlxcel#1943 closed on Blackwell. +/// +/// That fallback is why the runtime mismatch check exists: on a host without +/// `nvidia-smi` it produces a binary that cannot run on its own build machine, +/// and without the check the only symptom is an opaque CUDA load failure at the +/// first kernel launch. /// /// Empty string on a non-CUDA build, which the runtime reads as "no compiled /// architecture list" and skips the check entirely. From 1e2877364516ab6f82a819481430ad22741a7e82 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 19:17:41 +0900 Subject: [PATCH 5/7] docs(core): stop the compiled-list example showing a rejected shape The doc comment on `compiled_cuda_architectures` illustrated the string with `"80;86;89;90a;100a-real;100;120a-real;120"`, which names Blackwell architecture-specific twice. A doc comment on a parser is where a reader copies a list from, and after this branch the project's own guard rejects that shape, so the file's example and the file's rule disagreed. It now shows the real x86_64 release list, which exercises the same parser features that mattered, pre-Blackwell plain entries and a Hopper `a` entry, and cannot drift into recommending something the build refuses. The `-real` and `-virtual` forms are still described on `parse_cuda_arch_entry` and round-tripped in `entry_display_round_trips_through_the_parser`, which is where an illustration of the grammar belongs. The same pass corrects the header on the release-list constants in the test module, which still justified Hopper's `90a` by the quantized-kernel gate that upstream deleted. Refs #1943. --- src/lib/mlxcel-core/src/cuda_arch.rs | 13 +++++++++++-- src/lib/mlxcel-core/src/cuda_arch_tests.rs | 7 +++++-- 2 files changed, 16 insertions(+), 4 deletions(-) diff --git a/src/lib/mlxcel-core/src/cuda_arch.rs b/src/lib/mlxcel-core/src/cuda_arch.rs index e15c085e1..067953c2c 100644 --- a/src/lib/mlxcel-core/src/cuda_arch.rs +++ b/src/lib/mlxcel-core/src/cuda_arch.rs @@ -55,8 +55,17 @@ const FIRST_PLAIN_SM: u32 = 100; // ── Compiled architecture list ──────────────────────────────────────────────── /// The `MLX_CUDA_ARCHITECTURES` list this binary's MLX device code was -/// compiled for, as a CMake-style semicolon-separated string -/// (`"80;86;89;90a;100a-real;100;120a-real;120"`). +/// compiled for, as a CMake-style semicolon-separated string. An x86_64 release +/// build records `"80;86;89;90a;100;120"`, the list +/// `.github/workflows/release.yml` passes it. +/// +/// The example is a real shipped list rather than a tour of the grammar, so it +/// cannot drift into recommending a spelling the project rejects: Blackwell is +/// plain because the architecture-specific image belongs on `fp_quantize.cu` +/// alone (lablup/mlxcel#1934, #1943). The `a`, `f`, `-real` and `-virtual` +/// forms this module still has to parse are described on +/// `parse_cuda_arch_entry` and round-tripped in +/// `entry_display_round_trips_through_the_parser`. /// /// Empty on any build without the `cuda` feature, and empty on a CUDA build /// whose build script predates this record. Callers treat empty as "unknown", diff --git a/src/lib/mlxcel-core/src/cuda_arch_tests.rs b/src/lib/mlxcel-core/src/cuda_arch_tests.rs index 1d77b3931..ee43b2d25 100644 --- a/src/lib/mlxcel-core/src/cuda_arch_tests.rs +++ b/src/lib/mlxcel-core/src/cuda_arch_tests.rs @@ -32,8 +32,11 @@ use crate::cuda_arch::{ /// architecture-specific target, but `src/lib/mlx-cpp/CMakeLists.txt` gives that /// to `fp_quantize.cu` alone rather than to the whole `mlx` target, so these /// lists stay plain and every other kernel keeps the code it had (issue #1934). -/// Hopper's `90a` is unrelated and load-bearing for MLX's own quantized-kernel -/// gate. +/// Hopper's `90a` is unrelated: these converters need compute capability 10.0, +/// so 9.0 cannot reach them however it is spelled. It is kept because it is +/// what the release lists ship and what auto-detection mirrors, not because it +/// gates a kernel at the current MLX pin, where the macro that once did was +/// deleted upstream (issue #1943). const RELEASE_AARCH64: &str = "90a;100;121"; const RELEASE_X86_64: &str = "80;86;89;90a;100;120"; From cc2158a965e15930e9915a13a705b66044574922 Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 19:22:08 +0900 Subject: [PATCH 6/7] chore(ci): correct the reasons that named the pre-#1943 auto-detect shape Four comments justified themselves by what auto-detection used to produce. Three CUDA jobs in `ci.yml` said their pin exists because auto-detection "yields `121a` here"; after this branch it yields `121`. Each now states the reason that survives, that a gate should name the list it compiles rather than infer it from whichever GPU the runner has, so the pin keeps a rationale instead of becoming a bare value someone deletes as redundant. `scripts/bench_block_width.sh` is the one that inverted rather than going stale. It told a reader that a row measured on an auto-detected build is not comparable with records pinned at plain `121`, which after this branch is backwards, since an auto-detected GB10 build now resolves to exactly `121`. A wrong comparability rule in a benchmark harness produces discarded runs or published numbers someone trusts, so it now says which rows are actually suspect: those recorded between #1934 and #1943 on an unpinned build. The aarch64 release job carried the same deleted-macro rationale the x86_64 one did, missed in the earlier pass because only one of the two was grepped. Refs #1943. --- .github/workflows/ci.yml | 29 +++++++++++++++++------------ .github/workflows/release.yml | 12 +++++++++--- scripts/bench_block_width.sh | 11 ++++++++--- 3 files changed, 34 insertions(+), 18 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index f50cb9589..95de0b862 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -557,10 +557,13 @@ jobs: contents: read timeout-minutes: 90 env: - # What releases build for this runner's architecture. Pinned rather than - # left to auto-detection, which yields `121a` here and would compile every - # kernel architecture-specific (issue #1934), so the artifact this job - # verifies carries the same device code the published one does. + # What releases build for this runner's architecture, pinned rather than + # inferred so the artifact this job verifies carries the same device code + # the published one does whatever GPU the runner turns out to have. Since + # issue #1943 auto-detection produces this same `121` on a GB10, so the pin + # no longer corrects a shape mismatch: it states the list instead of + # inferring it, which is what makes the job reviewable and what keeps it + # right if this ever runs on something other than a GB10. MLX_CUDA_ARCHITECTURES: "121" steps: - uses: actions/checkout@v7 @@ -929,10 +932,11 @@ jobs: contents: read timeout-minutes: 120 env: - # The value shipped for this runner's architecture. Auto-detection yields - # `121a`, which is not what releases build and would compile every kernel - # architecture-specific (issue #1934), so the gate compiles what ships. The - # hardware NVFP4/MXFP4 converters come from the per-source injection in + # The value shipped for this runner's architecture, so the gate compiles + # what ships rather than whatever the runner's GPU implies. Since issue + # #1943 auto-detection produces this same `121` on a GB10; the pin remains + # because a gate should name the list it compiles instead of inferring it. + # The hardware NVFP4/MXFP4 converters come from the per-source injection in # `src/lib/mlx-cpp/CMakeLists.txt`, not from this list. MLX_CUDA_ARCHITECTURES: "121" # Every rustc lint, not just the `unused_imports` this job started with. @@ -1113,10 +1117,11 @@ jobs: cancel-in-progress: ${{ github.ref != 'refs/heads/main' }} timeout-minutes: 120 env: - # The value shipped for this runner's architecture. Auto-detection yields - # `121a`, which is not what releases build and would compile every kernel - # architecture-specific (issue #1934), so the gate links what ships. The - # hardware NVFP4/MXFP4 converters come from the per-source injection in + # The value shipped for this runner's architecture, so the gate links what + # ships rather than whatever the runner's GPU implies. Since issue #1943 + # auto-detection produces this same `121` on a GB10; the pin remains + # because a gate should name the list it compiles instead of inferring it. + # The hardware NVFP4/MXFP4 converters come from the per-source injection in # `src/lib/mlx-cpp/CMakeLists.txt`, not from this list. MLX_CUDA_ARCHITECTURES: "121" steps: diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index 6629d1cd8..7dce84fe0 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -534,9 +534,15 @@ jobs: # One aarch64 fat binary covering the NVIDIA aarch64 targets in a single # build: GH200 (Grace Hopper, sm_90a), GB200 (Grace Blackwell, sm_100), and # GB10 / DGX Spark (Blackwell, sm_121), cross-compiled on the GB10 runner. - # 90a is load-bearing for MLX's Hopper quantized kernel gate, and build.rs - # forwards the list verbatim. The CLI and server (~347 MB each) ship as - # separate archives, both with the CCCL headers. + # build.rs forwards the list verbatim, so this is exactly what nvcc is asked + # for. Hopper is spelled 90a because that is the spelling this project has + # shipped and the one auto-detection mirrors, not because it gates a kernel + # at the current MLX pin: upstream 44540d12 moved the Hopper quantized kernel + # to runtime NVRTC compilation and deleted MLX_CUDA_SM90A_ENABLED, which this + # comment used to cite. CUTLASS_ARCH_MMA_SM90A_ENABLED still keys on + # __CUDA_ARCH_FEAT_SM90_ALL, so a later pin can make it matter again + # (issue #1943). The CLI and server (~347 MB each) ship as separate archives, + # both with the CCCL headers. # # Blackwell stays plain here even though MLX's hardware NVFP4/MXFP4 # converters are gated on `__CUDA_ARCH_SPECIFIC__` and a plain target diff --git a/scripts/bench_block_width.sh b/scripts/bench_block_width.sh index 7467bbc1d..009a1499c 100755 --- a/scripts/bench_block_width.sh +++ b/scripts/bench_block_width.sh @@ -111,9 +111,14 @@ MEM=$(( $(sysctl -n hw.memsize 2>/dev/null || echo 0) / 1073741824 )) # 2026-09-20 and the resulting driver change was invisible to every harness # here, so a cross-session comparison had nothing to key on (issue #1797). # These three travel together as the host identity. The architecture belongs -# with them because build.rs auto-detects `121a` on GB10 while every earlier -# GB10 record pins plain `121`, so a row measured on an auto-detected build is -# not directly comparable with those records. +# with them because it decides what machine code the kernels are, so two rows +# built for different lists are not comparable whatever else matches. Since +# issue #1943 an auto-detected GB10 build resolves to plain `121`, the same list +# every earlier GB10 record pins, so an auto-detected row IS comparable with +# those records. That was not true before #1943, when auto-detection produced +# `121a` and gave every kernel a second, architecture-specific image; rows +# recorded between #1934 and #1943 on an unpinned build are the ones to treat +# with suspicion. The field is recorded rather than assumed either way. KERNEL=$(uname -r) DRIVER=$(nvidia-smi --query-gpu=driver_version --format=csv,noheader 2>/dev/null | head -1) ARCHES=${MLX_CUDA_ARCHITECTURES:-auto-detected} From b538e2758c2d2bc61b4038bd9b3b33354e09378d Mon Sep 17 00:00:00 2001 From: Jeongkyu Shin Date: Mon, 21 Sep 2026 19:28:41 +0900 Subject: [PATCH 7/7] chore(bench): correct the reused width harness on auto-detected builds The `draft-block-width-gb10-2026-09-20/harness/` directory is a live tool despite its dated path: the post-#1939 width sweep and the #1943 architecture A/B both re-used it, and `draft-block-width-post-1939-gb10-2026-09-21/` ships no `harness/` of its own for exactly that reason. Its four files told a reader that `build.rs` auto-detects `121a` here, so a row from an unpinned build is not comparable with the pinned records. After this branch that is backwards: auto-detection resolves to plain `121`, the same list those records pin. Each now keeps its pin and its reason, states that the pin no longer corrects a shape but still names what was compiled, and identifies the rows that are genuinely suspect, those recorded between #1934 and #1943 on an unpinned build, rather than implying every auto-detected build is. The record prose under `docs/benchmark_results/` and the `TECHNICAL_REPORTS` entry are left as written; they document what was true when they were measured. Refs #1943. --- .../harness/README.md | 2 +- .../harness/identity.sh | 13 ++++++++----- .../harness/run_session.sh | 14 +++++++++----- .../harness/sweep_server_widths.py | 9 ++++++--- 4 files changed, 24 insertions(+), 14 deletions(-) diff --git a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/README.md b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/README.md index fb05595e8..e1ec575f6 100644 --- a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/README.md +++ b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/README.md @@ -21,7 +21,7 @@ docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/run_sessio **The binary is renamed before it runs.** `run_session.sh` copies it to `mlxcel1797-server`. The imported host gate blocks on a foreign inference process by `/proc//comm`, matching `mlxcel` and `mlxcel-server`, which is what keeps a peer session's model off the GPU during a timed run. A server started under the stock name matches that list and gates itself forever. -**`MLX_CUDA_ARCHITECTURES=121` is pinned for the build and for every server.** `build.rs` auto-detects `121a` on this host, and every earlier GB10 record here pins plain `121`. The suffix does not change the decode kernels: in the pinned MLX tree only the NVFP4 weight-quantization converter keys on `__CUDA_ARCH_SPECIFIC__`. The pin is for comparability with those records, and it happens to match the release configuration as of 2026-09-20. +**`MLX_CUDA_ARCHITECTURES=121` is pinned for the build and for every server.** Every earlier GB10 record here uses plain `121`, and the pin states what was compiled rather than leaving a reader to infer it. Since issue #1943 `build.rs` auto-detects that same plain `121` on this host, so a run of this harness without the pin is comparable with those records; before #1943 auto-detection yielded `121a`, which is why rows recorded between #1934 and #1943 on an unpinned build are the ones that do not line up. The suffix does not change the decode kernels either way: in the pinned MLX tree only the NVFP4 weight-quantization converter keys on `__CUDA_ARCH_SPECIFIC__`. **The host gate is imported from `../../sdpa-plan-bucket-gb10-2026-09-12/harness/hostgate.py`, not copied.** It encodes three predicates that were each paid for with a lost measurement: contention means processes that can consume CPU now (match `/proc//comm` exactly, drop state `T`, exclude the CI runner's permanent `RunnerService.js` and `Runner.Listener` daemons, never a `%cpu` threshold because `ps` reports a lifetime average); a foreign model process is a reason to wait; and the cumulative `NV_ERR_NO_MEMORY` count is the freeze precursor on this host and does not decay. A second copy would drift from the original. diff --git a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/identity.sh b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/identity.sh index 8b22c66dc..605ae51b5 100755 --- a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/identity.sh +++ b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/identity.sh @@ -28,11 +28,14 @@ printf 'nvidia_driver: %s\n' \ "$(nvidia-smi --query-gpu=driver_version --format=csv,noheader 2>/dev/null | head -1)" printf 'gpu: %s\n' \ "$(nvidia-smi --query-gpu=name,compute_cap --format=csv,noheader 2>/dev/null | head -1)" -# The architecture list the build was PINNED to, which is not what a default -# build produces here: build.rs auto-detects and yields `121a`, while every -# prior GB10 record pins plain `121`. The suffix does not change the decode -# kernels (only the NVFP4 weight-quantization converter keys on -# __CUDA_ARCH_SPECIFIC__); it is recorded so two sessions can be told apart. +# The architecture list the build was compiled for, recorded so two sessions can +# be told apart rather than assumed comparable. Since issue #1943 a default +# build here auto-detects plain `121`, the same list every GB10 record pins, so +# an auto-detected row and a pinned row agree. Before #1943 auto-detection +# yielded `121a`, so rows recorded between #1934 and #1943 on an unpinned build +# are the ones to check this field on. The suffix does not change the decode +# kernels either way (only the NVFP4 weight-quantization converter keys on +# __CUDA_ARCH_SPECIFIC__). printf 'MLX_CUDA_ARCHITECTURES: %s\n' "${MLX_CUDA_ARCHITECTURES:-unset (auto-detected)}" # nvcc is not on the default PATH on this host; the CUDA install root is. NVCC=$(command -v nvcc || echo /usr/local/cuda/bin/nvcc) diff --git a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/run_session.sh b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/run_session.sh index 3eb666f95..2fe13a175 100644 --- a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/run_session.sh +++ b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/run_session.sh @@ -15,11 +15,15 @@ # unambiguously not the name a peer's server carries. # # 2. `MLX_CUDA_ARCHITECTURES=121` is exported for the build AND for every -# server process. build.rs auto-detects `121a` on this host, which is not -# what any earlier GB10 record used. The suffix does not change the decode -# kernels this harness measures (only the NVFP4 weight-quantization -# converter keys on __CUDA_ARCH_SPECIFIC__); it is pinned so these numbers -# sit alongside every earlier GB10 record's. +# server process, so these numbers sit alongside every earlier GB10 record's +# and the list is stated rather than inferred. Since issue #1943 a default +# build here auto-detects this same plain `121`, so the pin no longer +# corrects a shape; it is kept because a harness should name what it +# compiled. Before #1943 auto-detection yielded `121a`, so rows recorded +# between #1934 and #1943 on an unpinned build are the ones that are not +# directly comparable. The suffix does not change the decode kernels this +# harness measures either way (only the NVFP4 weight-quantization converter +# keys on __CUDA_ARCH_SPECIFIC__). set -uo pipefail BIN=${1:?usage: run_session.sh [outdir]} D="$(cd "$(dirname "$0")" && pwd)" diff --git a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/sweep_server_widths.py b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/sweep_server_widths.py index f5ed5719e..59db6555b 100755 --- a/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/sweep_server_widths.py +++ b/docs/benchmark_results/data/draft-block-width-gb10-2026-09-20/harness/sweep_server_widths.py @@ -190,9 +190,12 @@ def run_arm(a, width, prompt, tag): # left unset, so the record states it instead of the reader assuming it. env.setdefault("MLX_ENABLE_TF32", "1") env.setdefault("RUST_LOG", "info") - # Pinned, not auto-detected: build.rs yields `121a` here while every prior - # GB10 record uses plain `121`. The suffix does not change the decode - # kernels; the pin is for comparability with those records. + # Stated rather than inferred, for comparability with every prior GB10 + # record, which uses plain `121`. Since issue #1943 auto-detection here + # produces that same `121`, so the pin no longer corrects a shape; before + # #1943 it yielded `121a`, which is why rows recorded between #1934 and + # #1943 on an unpinned build are the ones that do not line up. The suffix + # does not change the decode kernels either way. env.setdefault("MLX_CUDA_ARCHITECTURES", "121") env.update(arm_env(width)) t0 = time.time()