Skip to content

test(cuda): run the deferred CUDA arm for the merged ROCm routing work on a GB10 node #1873

Description

@inureyes

Part of #1801

Problem / Background

The changes listed below are already merged to main. This issue is a verification pass, not a merge gate: it can find a regression that already shipped, and a failure here means opening a fix issue against the specific change, not reverting the set blindly.

The CUDA arm was deferred at merge time because no CUDA host was reachable from the ROCm development machine or from the Metal host, and the self-hosted CUDA CI runner has been down (its jobs sat queued and expired at the 24 hour timeout). Rather than leaving an implied "someone should check CUDA" note in each PR, the whole deferred set is collected here to be run in one pass when a GB10 node is available.

Scope widened on 2026-09-29. The routing work this issue covers grew after it was filed. PRs #2010, #2018, #2026 and #2029 each changed code on the CUDA path, and none was run on CUDA, because the only development host is the gfx1151 machine. PR #2029 says so explicitly in its "Not verified" section ("Metal and CUDA were not run; neither is available on this host"); #2010, #2018 and #2026 have no such section, but each records its validation as gfx1151-only, and #2018 argues CUDA is unaffected by reasoning ("CUDA and Metal are unreachable from every new path, since custom_kernels_available() is true there") rather than by a run. All four are now in scope, and the arm is main at or after d8d34e2b (PR #2029).

What needs the CUDA arm

Change Why CUDA is at risk What would go wrong
refactor(core): route custom kernels by GPU backend kind (#1803, PR #1877, d492c81b; an earlier revision of this issue cited PR #1869, which was closed unmerged) Ten call sites that chose a fused kernel with !mlx::core::metal::is_available() now ask a GpuKernelBackend enum, and four availability predicates now call custom_kernels_available(). On a CUDA build the enum resolves to Cuda and the predicates to true, so the selection is equivalent by construction. (Partly superseded 2026-09-29: since PR #2026, three of those predicates read has_kernel_port on their own table instead, fused_norm.cpp:347, fused_rope_append.cpp:394 and sampling.cpp:377; sampling_rejection.cpp:769 still calls custom_kernels_available().) If it is not, CUDA silently takes a Metal kernel or a graph fallback. A fallback keeps the output identical and only changes speed, so an output comparison alone cannot see it.
The same PR's Result change paged_attention_decode, paged_attention_decode_v2_partial and paged_attention_merge_states now return Result through cxx instead of a bare value, and their launchers refuse when the backend has no port. The refusal branch is unreachable on CUDA, but the Result wrapping is new code on a hot path.
paged_decode_backend() in src/lib/mlxcel-core/src/layers.rs:5534 It now asks custom_kernels_available() (:5543) before Metal and CUDA. If that predicate ever answered false on a CUDA host, every fused paged decode would fall back to gather with identical output and lower throughput.
feat(rocm): port the BitNet bitlinear_matmul kernel (#1862, PR #1870) The kernel selection gained a third arm and a new bitlinear_kernel_available() predicate. Since PR #2029 the three arms are a KernelPorts table, bitlinear_ports() (src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp:240, dispatch :277). CUDA must still pick the CUDA kernel, and the BitNet tests must run rather than skip.
feat(rocm): implement FFT through hipFFT (#1825, PR #1861) Only src/lib/mlx-cpp/patches-rocm/ plus a hipfft entry in the link list and docs. A CUDA build copies no file from that directory. Low risk; included so the run covers the whole merged set.
(Added 2026-09-29) fix(rocm): refuse instead of aborting at the sampler entry points (#1885, PR #2010, 99f48d7a) gumbel_max_sample, fused_sample_rejection and fused_sample_rejection_deferred became Result through cxx, with a port guard ahead of selection in sampling.cpp and sampling_rejection.cpp; both kill-switch tests now accept whichever decline reason applies, and both sampler microbenchmarks handle the Result. CUDA must still select its sampler ports, and tests/sampling_rejection_kill_switch.rs must still attribute the decline to MLXCEL_SAMPLING_REJECTION on CUDA, where a port exists.
(Added 2026-09-29) fix(rocm): guard the four remaining custom-kernel launchers (PR #2018, e2a2ebee) Guards ahead of fused_add_rms_norm, fused_rope_qk_append, ssm_update_kernel and run_fused_moe_two_kernel, and five bridge declarations became Result. The MoE callers (src/models/switch_layers.rs, gemma4.rs, qwen3_next.rs) fold an Err into .ok(); the SSM callers (falcon_h1.rs, granitemoehybrid.rs, nemotron_h.rs, plamo2.rs), layers.rs and fused_moe_parity_tests.rs expect it. A spurious refusal on CUDA is a panic at an expect site, but at a MoE .ok() site it is a silent, output-identical fallback to the SwitchGLU graph path, visible only as lower throughput.
(Added 2026-09-29) refactor(core): choose custom-kernel ports through one helper (PR #2026, e3647a31) The nine two-way launchers now resolve through mlxcel::select_kernel_port (src/lib/mlx-cpp/turbo/kernel_port.cpp:57) on per-kernel KernelPorts tables; a new translation unit kernel_port.cpp is compiled through build.rs; run_fused_moe_two_kernel now takes its caller's entry-point name. The argument that CUDA is unchanged: for a backend with a port, the resolved kernel is the same holder the old use_cuda ? cuda_port : metal_port expression resolved, and the throw path requires a null entry, while all nine tables have a non-null .cuda. The CUDA run is what confirms that argument. A table wired to the wrong getter, or a null .cuda, turns a CUDA fast path into a refusal: a panic at an expect caller, a silent fallback at a .ok() caller.
(Added 2026-09-29) refactor: finish the kernel port standardization (PR #2029, d8d34e2b) Five Metal-only launchers (fused_xielu, mamba1_selective_scan, MoE fc1_relu2, fused_add3_layer_norm, and the turbo4 and sparse-V SDPA launchers) now resolve through tables with .cuda = nullptr, and their bridge declarations became Result; the Rust gates in fused_norm_parity_tests.rs, fused_rope_parity_tests.rs, rms_norm_small_axis_tests.rs and examples/fused_norm_rope_microbench.rs changed from metal_is_available() || cuda_is_available() to the kernel's own *_available() or gpu_backend_available(); src/models/nemotron_h.rs:1359 gained a custom_kernels_available() term and an expect at :1424. On CUDA every changed gate should answer as it did before, so the parity and small-axis tests must still run rather than skip. One traced, not run, consequence: with the opt-in MLXCEL_FUSED_MOE_RELU2 set, Nemotron-H decode reaches select_kernel_port on moe_fc1_relu2_ports() (mlx_cxx_kernels.cpp:2293, .cuda = nullptr) and the expect at nemotron_h.rs:1424 panics. Before PR #2029 that path reached the Metal holder directly, so it was likely already broken on CUDA; record which.

Procedure

This is written for someone with no context from these PRs, on a GB10 node with a working CUDA toolchain and a clone of this repository.

1. Ask the binary which backend it resolved

The merged work ships its own probe, so the central risk of this whole issue does not have to be inferred from counters. gpu_kernel_backend() in src/lib/mlx-cpp/turbo/gpu_backend.cpp resolves the backend once behind a static, and when MLXCEL_DEBUG_KERNEL_BACKEND is set in the environment (gpu_backend.cpp:73) it prints exactly one line to stderr:

[mlxcel] custom kernel backend: cuda

When the resolved backend has no custom kernel ports it appends (no custom kernel ports; using MLX graph fallbacks). Run the arm binary once with the variable set and read that line:

MLXCEL_DEBUG_KERNEL_BACKEND=1 ./target/release/mlxcel generate --model models/mlx/Meta-Llama-3.1-8B-Instruct-4bit --prompt hi -n 8 2>&1 | grep 'custom kernel backend'

Any run that touches a kernel-availability predicate triggers it, and a short generate does, so no special model or flag is needed. (Corrected 2026-09-29: this used to say fused_norm.cpp and sampling.cpp both call custom_kernels_available(); since PR #2026 they call has_kernel_port, which resolves the same gpu_kernel_backend() at kernel_port.cpp:53-54, so the probe still fires.) Pass condition: the line reads cuda with no suffix. That single line settles the risk the rest of this issue is built around, which is that a new predicate answers wrongly and every fused path quietly falls back to an MLX graph with identical output and only lower throughput. It costs one run, so do it first: if it fails, nothing below is worth running, and the failure itself is the bug report.

2. Build both arms

The baseline is 56d82fbc (fix: refuse byte-identical arms in the A/B output gate (#1867)), the last commit on main before this ROCm series began landing; 4641dc7a (feat(rocm): implement FFT through hipFFT, PR #1861) is the first of the series on main. The arm is main at or after d8d34e2b (PR #2029), so it includes the whole widened set. Name the two binaries target/release/mlxcel.before (baseline) and target/release/mlxcel (arm).

git checkout 56d82fbc
cargo build --release --features cuda --bin mlxcel
/bin/cp target/release/mlxcel target/release/mlxcel.before
git checkout main
cargo build --release --features cuda --bin mlxcel 2>&1 | tee /tmp/build-arm.log

The check that matters is in the second build's log: src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp, mlx_cxx_bridge.cpp, and the files under src/lib/mlx-cpp/turbo/ (kernel_port.cpp, paged_attention.cpp, paged_attention_v2.cpp, paged_attention_v2_merge.cpp, sampling.cpp, sampling_rejection.cpp, fused_norm.cpp, fused_rope_append.cpp) must actually recompile. A cargo build that no-ops after a bridge .cpp edit has been observed on this project, and when it happens the arm binary is secretly identical to the baseline, at which point every check below passes for the wrong reason. Grep the log for those filenames and keep the matching lines as evidence.

3. The test gate

This is the strongest of the automated checks, because the tests that exercise these launchers call them directly rather than through a model, so no token floor or default-off gate can quietly skip them.

git checkout 56d82fbc && make verify-test-cuda 2>&1 | tee /tmp/test-before.log
git checkout main && make verify-test-cuda 2>&1 | tee /tmp/test-after.log
grep '^test result:' /tmp/test-before.log > /tmp/result-before.txt
grep '^test result:' /tmp/test-after.log > /tmp/result-after.txt
diff /tmp/result-before.txt /tmp/result-after.txt

Pass condition: the same tests pass in both runs, and nothing that passed before fails after. The groups whose results matter most are paged_v2::*, mla::split_kv, cache::paged_batch_decode and models::bitnet::*, and (added 2026-09-29 for the widened set) sampling_gumbel_tests, sampling_rejection_tests, the sampling_gumbel_kill_switch and sampling_rejection_kill_switch test targets, fused_norm_parity_tests, fused_rope_parity_tests, fused_moe_parity_tests and rms_norm_small_axis_tests. A group that goes from passing to skipped is a failure, not a pass. make verify-test-cuda is the CUDA gate (Makefile:775); there is no separate verify-cuda aggregate target.

(Added 2026-09-29) On the arm, also run the two toolchain-free checkers the widened set added to or relied on, which must pass unchanged on a CUDA host:

make verify-kernel-port-dispatch   # scripts/ci/check_kernel_port_dispatch.py (Makefile:818)
make verify-kernel-dtype-keys      # scripts/ci/check_kernel_dtype_keys.py (Makefile:813)

Pass for the widened set: make verify-test-cuda at its prior baseline (no new failure, no group newly skipped), and the port checker reporting 10 launcher file(s) route through select_kernel_port, 0 exempt, as it does on d8d34e2b.

4. Greedy output equality

./scripts/ab_output_equality.sh \
  --baseline target/release/mlxcel.before \
  --arm target/release/mlxcel \
  --model models/mlx/Qwen3-30B-A3B-4bit \
  --model models/mlx/Meta-Llama-3.1-8B-Instruct-4bit \
  --prompt "Explain how paged attention differs from a contiguous KV cache." \
  -n 64

The script pins --temp 0 and --show-reasoning on both arms, so the comparison covers every generated token and is not a comparison between two samples. It also runs the baseline twice as a self-consistency control and reports INCONCLUSIVE rather than FAIL when the baseline does not reproduce itself. It refuses to start when the two binaries are byte-identical; that refusal is the no-op build from step 2 surfacing, not a fault in the script. Exit status: 0 all pairs equal, 1 at least one pair differs, 2 inconclusive with no differences.

5. Reach the three changed paged bridges

mlxcel generate never reaches any of them: it is single-sequence and dense, and the fused paged path needs a pool-backed cache. Each bridge has its own production route, and they are not the same route, so "run a server" is not a single answer for all three.

Bridge function How production reaches it
paged_attention_decode (v1, src/lib/mlx-cpp/turbo/paged_attention.cpp) Only through mlxcel tune's paged-decode-splits op (src/lib/mlxcel-core/src/autotune/ops/paged_decode_splits.rs:252, driven from src/commands/tune.rs:369). Its other caller is PagedBlockPool::paged_decode_fused (src/lib/mlxcel-core/src/cache/paged.rs:2195, was :2304), whose only non-test caller is paged_decode_attention_pooled (src/lib/mlxcel-core/src/layers.rs:5704, was :5406), and that function has no caller anywhere in the tree outside its own doc comments and tests: #720 and ADR 0001 retired it to a library-only API, and production models call paged_decode_attention_dense_compat (a different, unchanged bridge function) instead.
paged_attention_decode_v2_partial (paged_attention_v2.cpp) The server's pool-backed batched decode, which is the #899 production entry point and default on: src/models/llama3.rs:735 and src/models/qwen3.rs:355 call paged_batch_decode_attention, which reaches PagedBlockPool::paged_decode_batched (src/lib/mlxcel-core/src/cache/paged.rs:1982, was :1957), then launch_v2, run_decode_v2, and V2Context::launch_with_lse (src/lib/mlxcel-core/src/paged_v2/launch.rs:123). It is gated by the paged_v2::dispatch token floors, which MLXCEL_PAGED_ATTENTION_NATIVE=1 bypasses, and it needs paged decode storage, which effective_decode_storage_backend (src/server/batch/scheduler/mod.rs:216) resolves: for a requested Auto or Paged it returns Paged when max_batch_size > 1 && supports_batching && supports_paged_decode_backend, otherwise Dense. All three hold by default here, since --parallel defaults to -1 which resolves to 4 slots (resolve_n_parallel, src/server/cli_input.rs:1217, was :1210) and --max-batch-size inherits it, and Llama3 and Qwen3 both return true from supports_paged_decode_backend (src/models/llama3.rs:1608, src/models/qwen3.rs:1179, were :1556 and :1127). So a default server already selects the paged path; --decode-storage-backend paged only pins it. --kv-unified is a KV budget flag and selects nothing here.
paged_attention_merge_states (paged_attention_v2_merge.cpp) Two routes. The same v2 path as above, but only when the plan emits more than one chunk per request (plan.needs_merge, src/lib/mlxcel-core/src/paged_v2/launch.rs:143, was :152), so a long-context request is needed, not just any request. And DeepSeek MLA: src/lib/mlxcel-core/src/mla/split_kv.rs:305 through absorbed_decode_split_kv, reached from src/models/deepseek_v2.rs:552 behind MLXCEL_MLA_SPLIT_KV, which makes an ordinary generate run reach it.

So run three things rather than one:

# v1 paged_attention_decode: the only production route is the tuner.
MLXCEL_DEBUG_KERNEL_BACKEND=1 ./target/release/mlxcel tune --op paged-decode-splits

# v2 partial (+ merge on long contexts): the batched server path, not --kv-unified.
MLXCEL_PAGED_ATTENTION_NATIVE=1 MLXCEL_DEBUG_KERNEL_BACKEND=1 ./target/release/mlxcel-server --model models/mlx/Qwen3-30B-A3B-4bit --max-batch-size 4 --decode-storage-backend paged --metrics

# merge_states on its own route: DeepSeek MLA split-KV, reachable from plain generate.
MLXCEL_MLA_SPLIT_KV=1 ./target/release/mlxcel generate --model models/mlx/DeepSeek-V2-Lite-4bit --prompt "..." -n 64 --temp 0

For the server arm, send several concurrent chat completions with long prompts so more than one request decodes at once and each clears the dispatch floor, then repeat the identical sequence against the baseline binary. The workspace test gate in step 3 covers all three bridges directly and is the check to trust if a route above cannot be driven on this node. scripts/paged_decode_counter_ab.sh (PR #1881) encodes the server arm and the four traps from the 2026-09-14 comment below, and is backend-agnostic.

The paged arm's pass condition is the decode counters, not output equality

Output equality cannot see this failure. The gather fallback is output-identical to the fused path by construction, so an arm that silently stopped launching the fused kernel scores a clean PASS in step 4. There is an in-tree instrument for exactly this question, so use it rather than inferring.

paged_batch_decode_stats() (src/lib/mlxcel-core/src/cache/paged_batch_decode.rs:213) returns the process-wide counters v2_launches, gather_fallbacks, declines and cascade_launches (plus the cascade token / member sums and cascade_failures). The scheduler publishes them every tick: run_loop.rs:415 calls update_paged_decode_gauges(mlxcel_core::cache::paged_batch_decode_stats()), and /metrics renders three of them (src/server/routes/metrics.rs:225-229), which is why the server command above carries --metrics (the route is mounted unconditionally but the handler answers 501 without the flag, src/server/app.rs:494-499, was :475):

curl -s localhost:8080/metrics | grep mlxcel_paged_decode_launches_total
# mlxcel_paged_decode_launches_total{path="fused_v2"} N
# mlxcel_paged_decode_launches_total{path="gather"}   N
# mlxcel_paged_decode_launches_total{path="cascade"}  N

These are values rather than log lines on purpose, and the doc comment on PagedBlockPool::paged_decode_batched (src/lib/mlxcel-core/src/cache/paged.rs:1973-1976, was :1948-1951) records why: the outcome "is a value rather than a log line because the first cut of this used tracing::debug! and a whole production benchmark sweep then compared gather against gather without anything saying so." That is this issue's failure mode with a different cause, which is what makes the counters the pass condition here rather than a nice-to-have.

Pass condition for the paged server arm, read after the identical request sequence on each arm:

  • v2_launches (path="fused_v2") is non-zero on both arms and moves together across the two runs. Zero on either arm means the fused path never ran and everything else measured gather against gather.
  • gather_fallbacks (path="gather") does not rise on the arm alone. This is the counter the regression this PR could introduce actually moves: PR refactor(core): route custom kernels by GPU backend kind #1877 (earlier cited here as PR refactor(core): route custom kernels by GPU backend kind #1869, which was closed unmerged) added an early return at the top of PagedBlockPool::paged_decode_batched (src/lib/mlxcel-core/src/cache/paged.rs:2006-2013, was :1981-1988) that returns Ok((None, PagedDecodeOutcome::NotServable("the GPU backend has no fused paged-attention kernel port"))) when custom_kernels_available() is false, and paged_batch_decode_attention folds an Ok((None, _)) into gather_fallbacks (paged_batch_decode.rs:375), not into declines.
  • declines (not published to /metrics; read it from a test or a debug build) is the pre-pool reject count, bumped only when batch_is_servable or the single-token shape check rejects the batch before the pool is touched (paged_batch_decode.rs:328). It should not differ between arms either, but the new early return does not touch it, so it is corroboration and not the primary signal.

The server log corroborates the number with a reason. report_once (paged_batch_decode.rs:152) announces the first occurrence of each distinct outcome kind through tracing::info!, deliberately at info rather than debug, so the line paged decode v2: gather: batch not servable (the GPU backend has no fused paged-attention kernel port) appears without anyone enabling a debug filter. Read it as corroboration only: the flags are keyed by outcome kind, and NotServable is one kind (index 8, src/lib/mlxcel-core/src/paged_v2/outcome.rs:130), so an ordinary batched prefill declining with NotServable("not a single-token decode step") earlier in the run can claim the slot and suppress the kernel-port line. The counter cannot be pre-empted that way.

The two instruments answer different questions and both are needed. MLXCEL_DEBUG_KERNEL_BACKEND from step 1 says what the backend predicate resolved to; the counters say whether the fused path actually ran. A build can print cuda and still serve every step from gather (dispatch floor, multi-slab layer, MLXCEL_PAGED_ATTENTION_NATIVE kill switch), and the counters alone cannot say which of those it was. Post both.

Step 3 already exercises the same instrument at test level: tests/cascade_decode_dispatch.rs:185 reads paged_batch_decode_stats() and asserts it agrees with the returned outcome, on the grounds that otherwise the /metrics view of which kernel ran is decorative.

Two corrections to an earlier revision of this issue, recorded so they are not reintroduced. It said the server arm with --kv-unified and MLXCEL_PAGED_ATTENTION_NATIVE=1 exercised paged_attention_decode; it does not, because that bridge sits behind the library-only pooled entry point and the server never calls it. And --kv-unified is not what enables the paged pool: it only collapses the parallel context slots to one for sizing (effective_parallel_context_slots, src/server/startup.rs:838, was :788) and the batch scheduler does not read it at all.

Acceptance Criteria

  • The [mlxcel] custom kernel backend: line from step 1 is posted and reads cuda with no fallback suffix.
  • The two test result: summaries from step 3 are posted and show no new failures and no group that went from passing to skipped.
  • The A/B script's exit status and its per-pair verdict lines from step 4 are posted.
  • The build log evidence from step 2 is posted, showing mlx_cxx_kernels.cpp and the src/lib/mlx-cpp/turbo/ sources actually recompiling in the arm build.
  • Each of the three routes in step 5 is either driven and reported, or explicitly recorded as not runnable on this node with the reason.
  • The mlxcel_paged_decode_launches_total fused / gather / cascade split from step 5 is posted for both arms, showing fused_v2 non-zero on both and gather not rising on the arm alone.
  • Any difference found is filed as its own issue, naming the arm it appeared on and the exact command that produced it.
  • (Added 2026-09-29) The arm is built from main at or after d8d34e2b, so PRs fix(rocm): refuse instead of aborting at the sampler entry points #2010, fix(rocm): guard the four remaining custom-kernel launchers #2018, refactor(core): choose custom-kernel ports through one helper #2026 and refactor: finish the kernel port standardization #2029 are covered, and the posted make verify-test-cuda result shows the groups step 3 names for them passing, not skipped.
  • (Added 2026-09-29) make verify-kernel-port-dispatch and make verify-kernel-dtype-keys output from the GB10 node is posted, both passing.
  • (Added 2026-09-29) The MLXCEL_FUSED_MOE_RELU2 Nemotron-H path is run once on CUDA if a checkpoint is available, and the outcome (panic at nemotron_h.rs:1424, or not) is recorded, with a follow-up issue if it panics.

Note

A plain "all equal" is not sufficient on its own. It needs the step 2 build evidence and the step 1 backend line beside it, because the two most likely failure modes in this set both produce exactly an "all equal" result: a no-op rebuild that leaves both arms pointing at the same code, and a predicate that answers wrongly so every fused path silently takes an MLX graph fallback with identical output and lower throughput. For the paged server arm it also needs the step 5 counter split, which is the only one of the three that observes the fused path actually running.

Refresh log

2026-09-29

Activity

  1. added
    type:testTest related changes
    area:coremlxcel-core: MLX FFI, primitives, KV cache, layers
    platform:linuxLinux (CUDA / packaging) specific
    on Sep 13, 2026
  2. inureyes commented on Sep 14, 2026

    @inureyes
    MemberAuthor

    The Metal half of the same criterion is measured and recorded in #1881: docs/benchmark_results/kernel-backend-kind-metal-m1ultra-2026-09-13.md. Metal selects the same kernels. Four arms, with what was not measured named.

    Four traps cost a re-run each there, and all four apply unchanged here.

    • An EQUAL row whose backend probe stayed silent proves nothing: the two binaries agreed on a path neither changed. Run with MLXCEL_DEBUG_KERNEL_BACKEND=1 and treat a silent arm as failed. Three of five checkpoints were silent on the first pass, because FUSED_ADD_RMSNORM_DEFAULT and FUSED_ROPE_APPEND_DEFAULT are opt-in, falcon-mamba does not use ssm_update_kernel (falcon_h1, granitemoehybrid, nemotron_h, plamo2 do), and split_kv_enabled() short-circuits on absorbed_enabled().
    • Without MLXCEL_PAGED_ATTENTION_NATIVE=1 the paged arm declines at the 4096-token dispatch floor and compares gather against gather, which is the failure the paged_decode_batched doc comment records from the first Route production batched paged decode through the fused v2 kernel, retiring the gather-then-SDPA hot path #899 sweep.
    • /metrics is off unless --metrics is passed, and curl -f fails silently, so a missing counter file reads as a zero.
    • The discriminator is gather_fallbacks, not declines. The gate at the head of paged_decode_batched returns Ok((None, ...)), which the caller folds into gather_fallbacks; declines counts only pre-pool rejections and is not exported to the gauges.

    scripts/paged_decode_counter_ab.sh in #1881 encodes all four and is backend-agnostic, so the CUDA arm should be a single invocation against a GB10-built pair. Also worth knowing before reading its output: mlxcel_paged_decode_launches_total counts layer decode steps, not tokens, so it does not scale with context.

  3. added a commit that references this issue on Sep 14, 2026
  4. added and removed on Oct 6, 2026
  5. inureyes commented on Oct 6, 2026

    @inureyes
    MemberAuthor

    CUDA arm on GB10: steps 1, 2, 4, 5 and the checkers

    Host: spark-101, NVIDIA GB10 (sm_121), Linux, 2026-10-06. Arm: main at a4974fab (after d8d34e2b, so #2010, #2018, #2026 and #2029 are in it). Baseline: 56d82fbc, built in its own detached worktree. Both cargo build --release --features cuda. Every GPU run went through gpu-lock; NV_ERR_NO_MEMORY stayed at 156 for the whole session.

    Step 1: backend probe (pass)

    $ MLXCEL_DEBUG_KERNEL_BACKEND=1 ./target/release/mlxcel generate --model models/mlx/meta-llama-3.1-8b-instruct-4bit --prompt hi -n 8
    [mlxcel] custom kernel backend: cuda
    

    No suffix. The same line appears on every arm run below; the baseline binary has no probe (pre-#1803).

    Step 2: build evidence (pass)

    The arm build ran with CC_ENABLE_DEBUG_OUTPUT=1, which makes cc rerun the build script and log every compile command into the build-script output. All ten files compiled in that build (36 running: lines in total). Object mtimes, build started 19:49:37:

    19:59:48 mlx_cxx_bridge.o        20:00:00 paged_attention_v2.o
    19:59:51 mlx_cxx_kernels.o       20:00:00 paged_attention_v2_merge.o
    19:59:56 kernel_port.o           20:00:01 sampling.o
    19:59:59 paged_attention.o       20:00:02 sampling_rejection.o
    20:00:03 fused_norm.o            20:00:04 fused_rope_append.o
    

    The binaries differ, and arm-only strings are present only in the arm:

    a896a57f...  target/release/mlxcel         probe string 1, "has no fused paged-attention kernel port" 1
    a1052522...  target/release/mlxcel.before  probe string 0, refusal string 0
    

    Step 4: greedy output A/B (pass, exit 0)

    qwen3-30b-a3b-4bit                 Explain how paged attention differs from     EQUAL
    meta-llama-3.1-8b-instruct-4bit    Explain how paged attention differs from     EQUAL
    result: every pair is byte-identical between baseline and arm.
    exit=0
    

    Both arm .err files carry the cuda probe line. The first attempt failed before running anything: on GNU mktemp, mktemp -d -t ab_output_equality needs X's in the template, and the script then wrote to /. Filed as #2141; this run used --out-dir.

    Step 5: the three routes

    v2 partial, server (pass on the counters). scripts/paged_decode_counter_ab.sh with Llama-3.1-8B-Instruct-4bit (Qwen3 MoE does not opt into supports_paged_decode_backend), --out-dir for the same mktemp reason:

    short prompt (11 tokens)          baseline: fused_v2=8352 gather=0 cascade=0   arm: fused_v2=8352 gather=0 cascade=0
    long prompt (3583 tokens)         baseline: fused_v2=8352 gather=0 cascade=0   arm: fused_v2=8352 gather=0 cascade=0   text: EQUAL, exit 0
    arm backend: [mlxcel] custom kernel backend: cuda
    

    The short-prompt run exited 1 on text alone (## Step 2: Describe how vs ... the, token ~60). It is not the routing set. Each binary gives the same text on the fused and gather paths, and the arm's paged text matches both binaries' dense text:

                          fused_v2 (NATIVE=1)   gather (no NATIVE)   dense
    baseline 56d82fbc     71ce23e7              71ce23e7             a5acb8f3
    arm a4974fab          a5acb8f3              a5acb8f3             a5acb8f3
    

    So the baseline's paged path diverged from its own dense path and the arm's does not. Some fix among the other 169 commits since the baseline brought the paged path into line with dense. No issue filed.

    v1 paged_attention_decode (pass). mlxcel tune --op paged-decode-splits -m <llama-8b> --batch-sizes 1,4 --context-lengths 1024,4096 with MLXCEL_CACHE_DIR in a scratch directory: exit 0 on both binaries, all four shapes timed, and the arm printed the cuda probe. The timings ran while a test build was using the CPU, so they are not a performance comparison.

    merge_states via DeepSeek MLA (driven, EQUAL). MLXCEL_MLA_ABSORBED=1 MLXCEL_MLA_SPLIT_KV=1 (split-KV short-circuits without the absorbed flag) on deepseek-v2-lite-4bit with a 3.5k-token prompt: ab_output_equality.sh reports EQUAL with exit 0, and the arm prints the probe. No counter exposes whether the plan cut more than one chunk, and absorbed_decode_split_kv(...).ok() would fall back silently, so the direct coverage for this route is mla::split_kv in step 3. Within the server route, merge is likewise only covered at test level: the outcome log prints once per kind and showed only the first, merge skipped, launch.

    MLXCEL_FUSED_MOE_RELU2 on Nemotron-H (no panic on the arm)

    MLXCEL_FUSED_MOE_RELU2=1 mlxcel generate --model models/mlx/nemotron-h-30b-4bit --prompt "What is the capital of France?" -n 32 --temp 0:

    • arm: exit 0, cuda probe, the same text as the arm without the variable. On main the C++ branch checks fused_moe_relu2_kernels_available() (mlx_cxx_kernels.cpp:1635, reading moe_fc1_relu2_ports(), which has no .cuda) and declines to gather_qmm, so the expect (now nemotron_h.rs:1426) is never reached.
    • baseline: exit 134, terminate called after throwing an instance of 'std::runtime_error' what(): [metal_kernel] No Metal back-end.

    That answers "record which": the path was already broken on CUDA before #2029 (a C++ abort, not a Rust panic), and it works now. No follow-up needed.

    Checkers (pass)

    Kernel port dispatch check passed: 12 launcher file(s) route through select_kernel_port, 0 exempt; 1891 Rust file(s) free of backend-enumerating gates, 0 awaiting a predicate.
    kernel-dtype-keys: OK, 9 in scope (launching cuda_kernel or hip_kernel) of 195 source files scanned.
    all kernel dtype-key checker cases passed
    

    The count is 12 rather than the 10 recorded at d8d34e2b because later launchers on main now route through the helper (for example mlx_cxx_one_bit.cpp, #2035). Zero exempt, as required.

    Step 3 follows in the next comment.

  6. inureyes commented on Oct 6, 2026

    @inureyes
    MemberAuthor

    CUDA arm on GB10: step 3, the test gate

    make verify-test-cuda ran on both trees under gpu-lock, with the test binaries built unlocked first (cargo test --workspace --profile test-fast --features cuda --no-fail-fast --no-run). Both runs exited 2 because of the failures below.

    Totals, summed over every test result: line:

    baseline 56d82fbc   123 targets   11225 passed   3 failed   362 ignored
    arm      a4974fab   146 targets   11931 passed   3 failed   377 ignored
    

    Failures:

    baseline  glm4v_parity           text_only_forward_produces_finite_logits      (passes on the arm)
    baseline  glm_ocr_parity         text_only_forward_produces_finite_logits      (passes on the arm)
    both      iquestloopcoder_parity iquestloopcoder_matches_the_float32_oracle    chunk 97: logit for token 5123 is 29.5, single-pass says 29.125 (same message on both, so it predates this set; also fails with MLX_ENABLE_TF32=0; filed as #2143)
    arm       mlxcel-core lib        ssm_update_parity_tests::ssm_update_kernel_matches_graph_step_bf16
    arm       mlxcel-core lib        ssm_update_parity_tests::ssm_update_kernel_keys_on_a_log_dtype
    

    The two ssm_update_parity_tests failures are not pass-to-fail: the file was added by #2099 (ROCm port, validated on gfx1151 only) and does not exist at the baseline. Both report a bf16-sized error in the new state (normalized RMS 2.5e-3 and 1.8e-3 against an f32 budget of 1e-5) on the granite shape with bf16 A_log. They fail the same way in isolation and with MLX_ENABLE_TF32=0. The ssm_update_kernel launcher is in scope here through #2018 and #2026, so this is filed as #2142.

    Eight tests that passed on the baseline do not appear in the arm log because they were renamed or deleted on main (for example ..._110_fields became ..._111_fields); the renamed ones pass on the arm.

    The groups this issue names (ok / failed / ignored)

    group                          baseline     arm
    paged_v2::                     93/0/0       93/0/0
    mla::split_kv                  9/0/0        9/0/0
    cache::paged_batch_decode      22/0/0       23/0/0
    models::bitnet                 1/0/0        2/0/0
    sampling_gumbel_tests          14/0/0       14/0/0
    sampling_rejection_tests       28/0/0       28/0/0
    fused_norm_parity_tests        8/0/0        10/0/0
    fused_rope_parity_tests        9/0/0        11/0/0
    fused_moe_parity_tests         4/0/0        5/0/0
    rms_norm_small_axis_tests      2/0/0        2/0/0
    tests/sampling_gumbel_kill_switch.rs       1 passed on the arm
    tests/sampling_rejection_kill_switch.rs    1 passed on the arm
    tests/cascade_decode_dispatch.rs           1 passed on the arm
    

    No group lost a test, and nothing moved to ignored. Some of these tests skip with an early return that libtest records as ok, so I reran every group above on the arm with --nocapture: the same counts, and no skipping line anywhere. The sampler and parity groups took seconds (gumbel 3.9 s, rejection 4.4 s, fused_norm 9.5 s, fused_rope 8.2 s, fused_moe 7.4 s), which an early return would not.

    Criteria status

    Steps 1, 2, 4, 5, the checkers and the Nemotron-H run are in the previous comment. Three differences were filed: #2141 (the A/B scripts on GNU mktemp), #2142 (the CUDA SSM parity failure) and #2143 (the iquestloopcoder chunked-prefill mismatch, present on both arms). The routing questions this issue was opened for are settled: the backend resolves to cuda, the fused paged path ran on both arms with no gather fallback, and every named group runs on CUDA.

    I ticked nine of the ten criteria. The second (verify-test-cuda with no new failures) stays open: the two SSM tests are new rather than regressions, but the gate is not at its prior baseline while #2142 is open. Remaining work: once #2142 lands, rerun ssm_update_parity_tests under make verify-test-cuda on GB10, then tick that criterion and close this issue. Status moves to status:blocked on #2142.

  7. added
    status:blockedBlocked by dependencies or other issues
    and removed on Oct 6, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    area:coremlxcel-core: MLX FFI, primitives, KV cache, layersplatform:linuxLinux (CUDA / packaging) specificpriority:highHigh prioritystatus:blockedBlocked by dependencies or other issuestype:testTest related changes

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions