Repository navigation
fix(test): make the CUDA merge gate runnable and serialized - #1052
Merged
Merged
Conversation
`cargo test --release --features cuda -p mlxcel-core --lib` has never reached a verdict on GB10. It dies with SIGABRT partway through, at a different test and with a different CUDA error each run, so the abort reads as if whichever test happened to be running is broken. Three full runs on `main` at 6276f55, release, idle machine, MLX pin 2c46b953, locate the cause. | Threads | MLX_USE_CUDA_GRAPHS | Outcome | |---|---|---| | default (20) | on | SIGABRT, `cudaStreamEndCapture ... previous error during capture` | | `--test-threads=1` | on | ran to a verdict: 1406 passed, 4 failed, 88.02s | | default (20) | `0` | SIGABRT, `cuLaunchKernelEx ... invalid argument` | The third run decides the fix and is the one #1048 did not anticipate. Disabling graph capture does not stop the abort, it only changes which CUDA call reports it, so the cause is the libtest harness driving MLX concurrently from many host threads and capture merely selects the symptom. `MLX_USE_CUDA_GRAPHS=0` for the suite, which the issue floated as a candidate remedy, is measurably not one. Capture therefore stays fully on, which also disposes of the criterion about keeping one test that exercises capture: nothing is disabled. Serializing costs 88 seconds for 1410 tests. There was a second, independent hole. No target ran mlxcel-core's tests on this platform at all: `verify-test` pins `--features metal,accelerate`, and `test-fast-cuda` serializes but omits `--workspace`, so a bare `cargo test` under it resolves to `-p mlxcel` and never builds mlxcel-core. That is the blindness #1007 removed on macOS, still present on the other backend. Three changes. `make verify-test-cuda` is the Linux/NVIDIA counterpart of `make verify-test`: `cargo test --workspace --profile test-fast --features cuda --no-fail-fast -- --test-threads=1`. Same scope, same profile, same reporting, with the flag embedded so the documented command cannot be run wrong. Pulling mlxcel-surgery and mlxcel-xla into the CUDA path is intended and close to free: both depend on mlxcel-core, and one workspace invocation unifies that into a single cuda-enabled build of it, so neither triggers a second MLX compile, and mlxcel-xla's `iree` feature stays off. `the_cuda_test_suite_must_run_single_threaded` in the new `src/lib/mlxcel-core/src/cuda_test_serialization_tests.rs` is the part that does not depend on remembering a flag. It fails by name, with the right command, whenever the CUDA suite is running multi-threaded, so a hand-run `cargo test --workspace --features cuda` gets a diagnosis instead of an anonymous abort. Being an ordinary test it is filtered out of narrowed runs, which stay parallel and are unaffected; it is whole-suite runs that abort. It writes to the stderr handle rather than through `eprintln!` so the message survives an abort that beats libtest's end-of-run summary, and `MLXCEL_ALLOW_PARALLEL_CUDA_TESTS=1` downgrades it to a warning. Gated on `feature = "cuda"`, so macOS is untouched; the Metal hazard is the cross-process one already covered by `gpu_exclusivity_tests`. CONTRIBUTING.md gains the CUDA gate command and the reason it must be serialized, and its Linux/CUDA sample command gains the `--test-threads=1` it was missing. docs/installation.md gains the measured table, the `MLX_USE_CUDA_GRAPHS=0` negative result, and the feature-resolution note for the CUDA workspace path. A `.cargo/config.toml` `[env] RUST_TEST_THREADS = "1"` was considered and rejected: it cannot be conditioned on platform or feature, so it would re-scope the macOS gate, where serializing buys nothing because #1008's hazard is between processes, and it would move the gate's behavior into a file the gate's own command line does not mention, which is the objection the root Cargo.toml already records against `default-members`. A committed nextest profile was rejected because per-test processes recreate #1008's cross-process condition and add a required tool. Validated on GB10: the guard passes serialized (5/5) and fails with the intended message at 20 threads, the escape hatch downgrades it, `cargo fmt --all -- --check` is clean, and `cargo tree --workspace --features cuda` resolves all five members. The gate itself was not executed end to end here; `[profile.test-fast]` has never been built on this host and the full suite was out of budget. Refs #1048
CONTRIBUTING.md and docs/installation.md stated flatly that narrowed runs filter `the_cuda_test_suite_must_run_single_threaded` out and stay parallel. That holds because libtest's default filter is a substring match on the test name, so `--lib paged_v2` never selects the guard, but it is a property of the filter rather than an absolute: a filter that does match the guard's name selects it, and on a multicore box the guard then fails a scoped run that would have completed safely. `cuda` is the realistic case, since `--features cuda --lib cuda` is a plausible narrowing and mlxcel-core already holds several genuinely CUDA-named tests (`cuda_l2_is_an_explicitly_unimplemented_stub`, `cuda_graph_cache_default_matches_build_feature`, `selector_cuda_backend_native_on_single_slab_only` among them). Both documents now name the exception and point at `MLXCEL_ALLOW_PARALLEL_CUDA_TESTS=1` for it, so a developer who trips the guard on a scoped run reads that it is expected rather than that their run is invalid. Also narrows "CUDA is gated nowhere in CI" to CUDA tests. release.yml does run `cargo build --release --features cuda --locked` on two self-hosted NVIDIA runners and runs no tests there, which is what step 4 of the same section already refers to as CUDA verification staying exclusive to release.yml; without the qualifier the two sentences read as contradicting each other. Docs only, no behavior change. Refs #1048
Member
Author
Implementation Review SummaryIntentMake the Linux/CUDA merge gate exist and reach a verdict: add Source-level review only. A Verified against the diff
Fixed on this branch (d61f895, docs only)
Left for you, not fixed
Verification
Could not verify without a build
|
`MLXCEL_ALLOW_PARALLEL_CUDA_TESTS`, the escape hatch added with the CUDA serialization guard, was described only in the guard's own source and in docs/installation.md, not in docs/environment-variables.md, which is the file people grep when they want to know what an `MLXCEL_*` variable does. `MLXCEL_ALLOW_CONCURRENT_GPU_TESTS`, the equivalent for the #1008 GPU-exclusivity guard, was missing for the same reason. Both are now in the Test and CI table with what they downgrade and when setting one is legitimate. An existing omission is a weak argument for adding a second. docs/installation.md stated the `cuLaunchKernelEx ... invalid argument` result three times: once in the prose introduction, once in the measured table, and once again in the paragraph interpreting the table. The table is the authoritative record, so the prose now points at it and the interpretation paragraph carries the mechanism (module load racing another thread's stream capture, versus a kernel-configure race) without restating the error string. Also records that the serialization result carries across profiles. The three diagnostic runs were taken under `[profile.release]`, while `make verify-test-cuda` builds under `[profile.test-fast]`; that gate has now run end to end on GB10 at 1411 passed, 4 failed, 89.35s with no abort, so the gate is not resting on an untested profile and the four known failures reproduce under the command a contributor actually runs. Refs #1048
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
cargo test --release --features cuda -p mlxcel-core --libhas never reached a verdict on GB10. It dies with SIGABRT partway through, at a different test and with a different CUDA error each run, so the abort reads as if whichever test happened to be running is broken. This PR makes the CUDA suite reach a verdict, makes serialization the property of the command rather than of the developer's memory, and adds the CUDA test gate that did not exist at all.What the measurements say
Three full runs of that command on
mainat6276f556, release profile,--features cuda, idle GB10 (sm_121), MLX pin2c46b953:MLX_USE_CUDA_GRAPHScudaStreamEndCapture ... failed: operation failed due to a previous error during capture--test-threads=10cuLaunchKernelEx ... failed: invalid argumentRun 3 is the result the issue did not anticipate, and it is the one that decides the fix. Disabling graph capture does not stop the abort. It only changes which CUDA call reports it: under capture the poisoned state surfaces at
cudaStreamEndCapture, without capture the same race surfaces as an invalidcuLaunchKernelEx. So the cause is the libtest harness driving MLX concurrently from many host threads, and graph capture merely selects the symptom.The issue floated "disable capture for the test binary, for example by setting
MLX_USE_CUDA_GRAPHS=0for the suite" as a candidate remedy. It is measurably not one. This PR says so in the guard's panic message, the Makefile comment, CONTRIBUTING.md and docs/installation.md, because a future reader hitting this abort will otherwise reach for exactly that.Two consequences worth stating plainly:
Serialization is cheap: 88 seconds for 1410 tests. There is no throughput argument against making it the gate.
Mechanism, consistent with the note already in docs/installation.md from #642: MLX begins capture with
cudaStreamCaptureModeThreadLocal(mlx/backend/cuda/device.cppat the pinned commit), which restrains only the capturing thread and leaves the others free to issue capture-unsafe calls, while the allocator issuescudaMalloc/cudaFreeon demand from whichever thread needs memory. That accounts for the moving abort site, the varying message, and scoped runs passing on the same binary.The second hole: there was no CUDA gate to fix
make test-fast-cudaalready passed--test-threads=1, but it omits--workspace. Per the workspace comment in the rootCargo.toml, the workspace root is itself themlxcelpackage, so a barecargo testresolves to-p mlxceland never buildsmlxcel-core. That target cannot run a single one of mlxcel-core's 1410 tests.make verify-testis--workspacebut pins--features metal,accelerate, so it is the macOS gate.Net: on Linux/CUDA no target ran mlxcel-core's tests, in the crate that holds the MLX
cxxbridge,layers.rs, the KV cache and the quantization loaders. That is the blindness #1007 removed on macOS, still present on the other backend, and CUDA is gated nowhere in CI, not even nightly, so nothing else was covering it.What changed
make verify-test-cudais the Linux/NVIDIA counterpart ofmake verify-test, following the shape #1007 established so the two gates stay symmetrical:cargo test --workspace --profile test-fast --features cuda --no-fail-fast -- --test-threads=1. Same scope, same profile, same--no-fail-fastreporting,cudain place ofmetal,accelerate, and the serialization flag embedded so the documented command cannot be run wrong.On the
--workspacequestion the issue raised: pulling inmlxcel-surgeryandmlxcel-xlaon the CUDA path is what we want, and it is close to free. Both depend onmlxcel-core, and onecargo test --workspace --features cudaunifies that into a singlecuda-enabled build of it, so neither triggers a second MLX compile.mlxcel-xla'sireefeature stays off, so it remains pure Rust and needs no IREE distribution.mlxcel-mlx-pindoes not depend onmlxcel-coreat all. Their test targets are backend-agnostic Rust, so excluding them would only mean those two crates are gated on macOS and nowhere else. This is written down in docs/installation.md next to the existing macOS explanation.There is deliberately no
verify-clippy-cudain this PR. The CUDA lint hole is a separate hole from the CUDA test hole, #1048 is about the test gate, andcargo clippy --workspace --all-targetsdoes not complete in a usable time on this host, so I would have been adding a target I could not run.A harness-level guard,
the_cuda_test_suite_must_run_single_threadedin the newsrc/lib/mlxcel-core/src/cuda_test_serialization_tests.rs, is the part that does not depend on remembering a flag. It resolves the thread count the way libtest itself does (command line--test-threadswins overRUST_TEST_THREADS, otherwiseavailable_parallelism()) and fails by name, with the right command, whenever the CUDA suite is about to run multi-threaded. A hand-runcargo test --workspace --features cudanow gets a diagnosis instead of an anonymous abort.Three properties worth noting:
eprintln!. libtest captures the print macros and replays them in its end-of-run summary, and in a parallel run there may be no end of run, because the abort takes the process down first. Writing to the handle bypasses the capture. Verified: the message appears on a 20-thread run.MLXCEL_ALLOW_PARALLEL_CUDA_TESTS=1downgrades it to a printed warning, mirroringMLXCEL_ALLOW_CONCURRENT_GPU_TESTSingpu_exclusivity_tests.rs, for someone deliberately reproducing the abort.It is gated on
feature = "cuda", so macOS is untouched. The Metal hazard is the cross-process one in #1008, whichgpu_exclusivity_testsalready covers, and--test-threadscapping is documented there as not helping it.Documentation. CONTRIBUTING.md gains the CUDA gate command, the requirement that it be serialized, and the run-3 reason not to reach for
MLX_USE_CUDA_GRAPHS=0; its Linux/CUDA sample command also gains the--test-threads=1it was missing, so following the contributor docs literally no longer produces a SIGABRT. docs/installation.md gains the measured table, the negative result, the guard's location and the CUDA feature-resolution note. Both state thattest-fast/test-fast-cudaare edit-test-loop targets and not gates. docs/environment-variables.md gainsMLXCEL_ALLOW_PARALLEL_CUDA_TESTSin the Test and CI table, along withMLXCEL_ALLOW_CONCURRENT_GPU_TESTS, which the #1008 guard has always had and which was never documented there either; that file is where people grep for what anMLXCEL_*variable does, and an existing omission is a weak argument for adding a second.Why not the other options
.cargo/config.toml[env] RUST_TEST_THREADS = "1". Cargo's[env]cannot be conditioned on platform or feature, so this would also serialize the macOS gate, where serializing buys nothing because fix(test): mlxcel-core release test binary intermittently hangs forever, blocking later runs via the cargo lock #1008's hazard is between processes rather than between threads, and I cannot measure that cost from this host. It would also move the gate's behavior into a file the gate's own command line does not mention, which is precisely the objection the rootCargo.tomlalready records against usingdefault-membersfor scope.cargo test --workspace --features cudaaborting anonymously, which is the case the guard exists for.Known-failing tests: the gate is not green once it runs
Run 2 reports 4 failures. None are fixed here, and none are marked
#[ignore], because hiding them would recreate exactly the blindness this issue exists to remove.fused_rope_parity_tests::fused_rope_append_matches_graph_rope_across_offsetssrc/lib/mlxcel-core/src/fused_rope_parity_tests.rs:87q offset=131071: normalized rms 2.684e-3 (tol 2.0e-3), normalized max 3.663e-2 (tol 1.2e-2)fused_rope_parity_tests::fused_rope_append_multi_token_positions_are_absolutesrc/lib/mlxcel-core/src/fused_rope_parity_tests.rs:87q seq=3 offset=131069: normalized rms 2.497e-3 (tol 2.0e-3), normalized max 3.698e-2 (tol 1.2e-2)paged_v2::launch::launch_tests::chunk_size_does_not_change_the_answersrc/lib/mlxcel-core/src/paged_v2/launch_tests.rs:324pages_per_chunk 1 relative error 55896.2paged_v2::sparse::sparse_tests::an_f16_cache_matches_the_reference_within_its_own_precisionsrc/lib/mlxcel-core/src/paged_v2/sparse_tests.rs:424max relative error 1.0004048 on an f16 cacheEvery failure a red CUDA gate can currently report maps to a tracking issue, so the next person reading one does not have to work out which of them are accounted for. The two
paged_v2failures were unfiled when this branch was opened and are now #1053 and #1054, split rather than filed together because the triage below says they are different bugs. From a scoped serialized run on this branch (-p mlxcel-core --lib paged_v2 -- --test-threads=1, 91 passed, 1 failed):an_f16_cache_matches_the_reference_within_its_own_precision(fix(paged): paged_v2 sparse f16 cache decode disagrees completely with its reference #1053) reproduces in isolation. A relative error of essentially 1.0 means the output carries no relation to the reference, so this is a correctness failure rather than tolerance drift, against a tolerance already sized for f16 rounding. Theassert!(outcome.is_fused(), ...)atsparse_tests.rs:421runs before the numeric assert at 424, so the fused sparse kernel path is genuinely taken and the failure is purely numeric.chunk_size_does_not_change_the_answer(fix(paged): paged_v2 chunked launch fails only in the full suite, pointing at cross-test state #1054) passes both inside that 92-test scope and run entirely alone (1.73s). It fails only in the full suite. That points at leaked shared mutable state rather than a defect in the chunking path itself, which is a materially different investigation from the first one. The assert atlaunch_tests.rs:324sits insidefor ppc in [1, 2, 3, 4, 7, 16, 64], so the reportedpages_per_chunk 1is only the first iteration and the larger chunk sizes are never reached; the failure is not specific to chunk size 1.Determining whether either also fails on a CPU-only build would need a separate non-CUDA MLX build, which I did not run.
Test plan
The gate itself, on GB10, run end to end.
make verify-test-cuda: the whole workspace under[profile.test-fast]with--features cuda --no-fail-fast -- --test-threads=1, cold-building that profile first.mlxcel-corelib: 1411 passed, 4 failed, 89.35s, no SIGABRT. Every other workspace target passed, roughly 6674 further tests across the other binaries.makeexited 2 on the four failures, which is what a red gate should do.Everything below on the same host, release profile,
--features cuda, narrow targets only.cargo test --release --features cuda -p mlxcel-core --lib cuda_test_serialization -- --test-threads=1gives 5 passed, 0 failed.MLXCEL_ALLOW_PARALLEL_CUDA_TESTS=1on the 20-thread run gives 5 passed with a printed warning.-p mlxcel-core --lib paged_v2 -- --test-threads=1runs 92 tests with the guard filtered out.cargo fmt --all -- --checkclean.cargo clippy -p mlxcel-core --release --tests --features cuda -- -D warningsclean, so the new module's test target carries no lint debt (this is the target--all-targetswithout--workspacewould not have compiled).cargo tree --workspace --features cudaresolves all five members, so the gate's feature selection is valid.make -n verify-test-cudaemits the intended command;make helplists the new target.python3 scripts/ci/check_cross_repo_refs.pyclean.What the end-to-end gate run settles
Two gaps this PR opened with, both now closed by that run.
The serialization result carries across profiles. The three diagnostic runs were taken under
[profile.release], while the gate builds under[profile.test-fast]. Those differ in LTO and codegen units, and this repository has a standing note that a defect can reproduce under one and not the other, so the gate could have been resting on a profile nothing had tested. It is not: 1411 passed, 4 failed, 89.35s, no abort, against 1406/4/88.02s under release.The known-failing list is accurate for the command a contributor actually runs, not only for the release-profile command the diagnosis used. The same four failures reproduce under
test-fast, at the same assertion sites:fused_rope_parity_tests.rs:87twice (#1049),paged_v2/launch_tests.rs:324(#1054),paged_v2/sparse_tests.rs:424(#1053). Nothing else in the workspace is red.One incidental result worth recording: no other test binary aborted either, so the root
mlxcelsuite having no guard of its own did not bite on this run. The guard coversmlxcel-corebecause that is where the aborts were measured; whether the root suite needs one is unresolved rather than answered.Not verified
The fresh-boot half of the acceptance criterion is still not met. The criterion asks for a verdict "both on a fresh boot and after a benchmarking session". Every run here, the three diagnostic ones and the gate run, is on a machine that has been up for days with builds and several full suite runs the same day, which covers the "after a session" half and arguably covers it harder than the criterion intended. A fresh-boot run needs a reboot, which is the maintainer's call, and prior sessions on this host recorded driver-level instability accumulating after heavy runs.
Closes #1048