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 c52e05d30..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 @@ -766,9 +772,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/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() diff --git a/docs/installation.md b/docs/installation.md index 55a8297b9..b33ffaf9b 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,28 @@ 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` 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. 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 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} 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..8dc757a8a 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 @@ -595,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. @@ -623,7 +631,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| { @@ -647,17 +656,53 @@ 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`. +/// +/// 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: `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. /// -/// 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. +/// 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..067953c2c 100644 --- a/src/lib/mlxcel-core/src/cuda_arch.rs +++ b/src/lib/mlxcel-core/src/cuda_arch.rs @@ -43,11 +43,29 @@ 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 -/// 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", @@ -325,16 +343,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..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"; @@ -474,10 +477,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 +495,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"); }