[midagedev]

cuda-oxide now comes straight from NVIDIA/cuda-rust at 9df9e7ce, on stable Rust 1.99.1. The fork (midagedev/cuda-oxide) is no longer a dependency; its last pin is tagged pin/87bc9518.

What changes

• Cargo.toml: cuda-device, cuda-host and cuda-core by rev 9df9e7ce; the [patch] to the fork is gone. rust-toolchain.toml: 1.99.1. .cargo/config.toml sets RUSTC_BOOTSTRAP=1, as cargo oxide new does.
• One backend patch is carried until upstream merges it: tools/oxide-backend/cuda-rust-1648.patch, the diff of perf(cute): speed up the CuTe verifier's provenance analysis NVIDIA/cuda-rust#1648 (juliankang4, fixes #1647) unchanged. Without it the CuTe verifier added in #1598 is cubic in a kernel's block count, and bloomery-gpu does not build in 1700 s.
  • tools/oxide-backend.sh builds the backend from the pin's checkout plus that patch, and writes the patch's sha256 into PROVENANCE.
  • tools/oxide-host.sh refuses (rc 70, by name) a backend without it, before every cargo oxide build.
• The fork's other jobs move into the repo:
  • A requires takes integer literals, each held to its constant by a const _: () = assert!(..).
  • Per-triple host flags: tools/oxide-host.sh generates .cargo/cuda-oxide.toml from .cargo/config.toml's table, with the llc major check.
• Six ptx-shapes.tsv rows are re-pinned at the measured spill and JIT local (PIN(2026-10-26)).
• Clippy on 1.99.1 is fixed in host and test code; just lint is at 0 warnings.

Evidence

• Cold cargo oxide build … --bin generate_ds41, box:

  Backend	Time
  0143af5c (before #1598)	388 s
  9df9e7ce unpatched	over 1700 s (rc 124)
  9df9e7ce + #1648	480 s

• ptx-scan --no-jit md5 of generate_ds41 (437 entries) and gate_kquant (372): equal between the #1648 backend and a backend that skips the verifier in scopes with no CuTe operation.

• just canon-check against the old pin: 1507 of 1585 entries differ, 1483 only by ld.global.nc loads (upstream marks read-only reference arguments noalias). The safe launchers cannot alias an input with an output. The two in-place add callers are element-wise.

• The full fixture batch is running; its DONE line will follow as a comment.

🤖 Generated with Claude Code

[github-actions]

Verdict: MUST_FIX: Concurrent gate preparation can abort a valid build.

Must fix

1. Give each preparation process its own temporary file

tools/oxide-host.sh:97 Both preparations use .cargo/cuda-oxide.toml.tmp. With the generated config absent, both can pass the existence check and finish writing that name before either renames it. The first os.replace removes the name; the second raises FileNotFoundError, and the shell exits 70 before Cargo runs.

This path is reachable: tools/gate-batch.sh:7155 starts lanes A and B concurrently. tools/box.sh:1258 takes a shared directory lock, and its host preparation runs before the requested Cargo command. Cargo's build lock cannot serialize these writes.

Fixing the temporary name now prevents intermittent gate failures. Replace line 97 with:

tmp = f"{path}.{os.getpid()}.tmp"

Add a two-writer case to the self-test; its current sequential calls do not expose the race.

Non-blocking

⚠ Attach the codegen validation: Cargo.toml:12. This compiler pin changes PTX; the new spill-pin comments also report occupancy changes, including gdn_conv_prep going from 9 to 7 blocks per SM. The supplied PR body says the fixture batch is still running and contains no timed A/B. Attach the selected gates' completed output and the occupancy class's same-lease A/B, with the card/driver line, both commits, model, commands and output.

Checked

Diagnostics: dryer reports zero qualifying pairs in the crates it scanned; GPU crates were outside that scan. Crapper reports no changed function crossing its thresholds.
Concurrency: the self-test's repeated writes are sequential; the directory lock permits concurrent preparations.
Gates: all six changed spill/JIT rows carry dated PIN reasons.
Escape catalog: the reported race does not match C1–C4.
Validation: an in-memory interleaving of the actual write/rename statements fails with the shared name and passes with distinct names. The inline suggestion targets added line 97 and leaves the embedded Python syntactically valid. GPU gates were not run.

_{Card check: needs-gpu. This pull request changes crates/gpu-deepseek41/src/swap.rs, crates/gpu-gates/src/bin/gate_deepseek41_attn.rs, crates/gpu-gates/src/bin/gate_deepseek41_chain_attn.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Concurrent gate preparations share this temporary name; the second rename fails and aborts its gate.

Suggested change

	tmp = path + ".tmp"
	tmp = f"{path}.{os.getpid()}.tmp"

[github-actions]

Since the last review:

• Must-fix 1 remains open: tools/oxide-host.sh:97.

Verdict: MUST_FIX: Concurrent preparation can still abort a valid build.

Must fix

1. Give each preparation process its own temporary file

tools/oxide-host.sh:97 Both writers use the same temporary path. With the generated config absent, both can finish writing before either renames it. The first rename removes the temporary name; the second raises FileNotFoundError, and preparation exits 70 before Cargo runs.

tools/gate-batch.sh:7155 starts concurrent lanes. tools/box.sh:1258 permits concurrent preparations with a shared directory lock. Cargo's build lock starts after preparation.

Fixing this now prevents intermittent gate failures. Replace line 97 with:

tmp = f"{path}.{os.getpid()}.tmp"

Add a two-writer self-test; the current sequential calls pass with this race present.

Checked

Duplicates: 0 new definitions in the incremental diff; dryer reports 0 qualifying pairs.
Complexity: crapper flags no changed function; this increment changes no function.
Silent failure: 0 new sites in the incremental diff.
Tests: none changed. The current write/rename statements fail a deterministic in-memory two-writer interleaving; the proposed replacement passes it.
Gates: the one incremental spill/JIT row has a dated PIN, measurement conditions and reason. No GPU gates were run.
Escape catalog: no C1–C4 pattern matches the incremental diff or the retained finding.

Notes

Dryer excluded crates/gpu and crates/gpu-gates from its scan.

_{Card check: needs-gpu. This pull request changes crates/gpu-deepseek41/src/swap.rs, crates/gpu-gates/src/bin/gate_deepseek41_attn.rs, crates/gpu-gates/src/bin/gate_deepseek41_chain_attn.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[midagedev]

Fixture tier full batch on 651fc4c, box, lead ledger:
DONE total=261 red=0 tier=fixture skipped=2 wall=4324s laneA=3592s laneB=3600s laneX=723s deferred=266 real_only=50 no_fixture=5 real_file=37 needs_env=7
The first try stopped at the canary: gate-ptx-spill red on hc_gated_up_mma_mix_4_bf16 alone (spill 628, jit_local 200 against 676/192), re-pinned with its PIN line.
The review's temp-file race is fixed in the next commit; its self-test is red on the old shared name (rc 70) and green with the fix.

[github-actions]

Since the last review:

• Must-fix 1 is fixed: tools/oxide-host.sh:99 gives each concurrent writer its own temporary file.

Verdict: CLEAN: The preparation race is fixed, and the incremental changes introduce no blocking defect.

Checked

Duplicates: 0 new definitions; dryer reports 0 qualifying pairs.
Complexity: crapper flags no changed function. The filename change and added self-test case add no decision branches.
Silent failure: 0 new sites; write and rename errors still terminate preparation with exit 70.
Tests: restoring the shared temporary name at tools/oxide-host.sh:99 makes the new case at line 202 fail because line 201 plants a directory there. The equal-file early return remains. Shell syntax and both Python heredocs compile; the self-test was not rerun during validation.
Gates: no pins, thresholds or existing cases change. This increment touches build preparation, outside the inference hot path.
Escape catalog: no C1–C4 patterns match the incremental changes.

Notes

Dryer excluded crates/gpu and crates/gpu-gates; this increment changes neither.
The PR carries no complete GPU gate evidence.

_{Card check: needs-gpu. This pull request changes crates/gpu-deepseek41/src/swap.rs, crates/gpu-gates/src/bin/gate_deepseek41_attn.rs, crates/gpu-gates/src/bin/gate_deepseek41_chain_attn.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{Not approved: this pull request changes .cargo/config.toml, .cargo/cuda-oxide.toml, AGENTS.md …, which changes the bot, CI or the toolchain and needs a person's approval.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Verdict: MUST_FIX: a relative backend directory installs the backend at a different path from the one subsequent builds select.

Must fix

1. Resolve the backend root before changing directories

tools/oxide-backend.sh:49 accepts BLOOMERY_OXIDE_BACKENDS=target/oxide. After build changes directory at line 122, its relative target and out point inside the copied source. The build can succeed there, while a later --path names an absent backend.
Fixing the directory owner now prevents misplaced installations and repeated setup failures. Replace the definition below. Add a stub case that builds with a relative root and checks the path returned by a separate invocation.

backends_root() {
  local root=${BLOOMERY_OXIDE_BACKENDS:-${CARGO_HOME:-$HOME/.cargo}/cuda-oxide-bloomery}
  case $root in
    /*) printf '%s\n' "$root" ;;
    *) printf '%s/%s\n' "$ROOT" "$root" ;;
  esac
}

2. Remove the duplicate backend identity check

tools/box.sh:1196 repeats the file/revision checks now owned by tools/oxide-backend.sh:63. The new OXCHECK already calls that owner through oxide-host.sh before every matching build. This is the same policy twice (R14).
Removing the old branch now leaves one place to maintain backend validation and removes elements. Replace lines 1194–1199, including the obsolete local-mode comment, with:

OXCLI=${OXCLI//OXREV/$OXREV}
[ "$LOCAL" = 0 ] || OXIDE="unset CUDA_OXIDE_BACKEND && "

3. Correct the local setup instructions

docs/contrib/local-check.md:235 says local mode builds its own backend. The new tools/oxide-host.sh:241 instead refuses a missing patched backend before Cargo starts.
Correcting this now gives new contributors the required setup step. Replace the sentence beginning with PATH on line 234 through line 235 with:

   `PATH` and `CUDA_TOOLKIT_PATH` as well. Local mode unsets `CUDA_OXIDE_BACKEND`; first run
   `tools/oxide-backend.sh` to install the patched backend. Before each oxide build, `tools/oxide-host.sh`
   selects it in `.cargo/cuda-oxide.toml` and refuses a missing or unpatched backend.

Non-blocking

⚠ Complete the compiler-change evidence: Cargo.toml:12 changes code generation; the supplied scan reports changed occupancy and spills. Attach the completed selected gates and the occupancy class's timed A/B on the contributor's card, with the card/driver line, both hashes, model, commands and output. The body currently says the fixture batch is running.

Checked

Duplicates: 28 new functions/types and 9 script constants checked; backend validation duplicates identified above. Dryer reported 0 pairs; its excluded GPU crates were searched manually. The carried upstream patch was read for correctness.
Complexity: crapper flagged no changed function; no complexity finding.
Silent failure: 8 added fallback/filter sites followed. Missing backend records, backend output, triple and LLVM banner reach named refusals; the alias fallback computes the original check; cleanup does not decide a result.
Tests: static mutation review: dropping flag emission, LLVM-major rejection, patch-hash rejection or host preparation makes the new script cases red. Capping opaque origins at one or accepting escaped cells makes the carried assertions red. Iterator rewrites retain their assertions and caller-proven lengths. The relative-root case is absent. Shell syntax and an in-memory path reproduction passed; build and gate suites were not run.
Gates: launch literals match compile-time assertions; float reduction order is preserved. Spill/JIT pins carry dated reasons.
Escape catalog: no C1–C4 pattern matched; the unquoted stub heredoc deliberately expands its fixture paths and escapes runtime variables.

_{Card check: needs-gpu. This pull request changes crates/gpu-deepseek41/src/swap.rs, crates/gpu-gates/src/bin/gate_deepseek41_attn.rs, crates/gpu-gates/src/bin/gate_deepseek41_chain_attn.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[github-actions]

Since the last review:

• Relative backend root: fixed at tools/oxide-backend.sh:50, with regression cases.
• Duplicate backend validation: fixed at tools/box.sh:1195. Host preparation retains the shared check.
• Local setup instructions: fixed at docs/contrib/local-check.md:234.

Verdict: CLEAN: all earlier must-fix items are fixed, and the incremental changes introduce no must-fix defect.

Checked

Duplicates: 0 new functions, types or constants in this increment; the existing path resolver changed. Duplicate validation was removed. Dryer reported 0 pairs.
Complexity: crapper flagged none. The path resolver adds a branch for relative versus absolute paths.
Silent failure: 0 new fallback sites. Removed identity guards remain enforced through oxide-host.sh and oxide-backend.sh --check; invalid records produce exit 70.
Tests: shell syntax checks passed for oxide-backend.sh, oxide-host.sh and box.sh. Static mutation check: removing the ROOT prefix at line 54 fails the new relative-build assertion; replacing ROOT with PWD fails the cross-directory check.
Gates: none moved in this increment.
Escape catalog: no C1–C4 pattern matched the new code.

Notes

Validation used source inspection and shell syntax checks. The validator did not rerun the stub suites or run a CUDA build or GPU gate.

_{Card check: needs-gpu. This pull request changes crates/gpu-deepseek41/src/swap.rs, crates/gpu-gates/src/bin/gate_deepseek41_attn.rs, crates/gpu-gates/src/bin/gate_deepseek41_chain_attn.rs … and carries no gate run on a card (CONTRIBUTING's evidence list); a maintainer adds gpu-verified after one.}

_{Not approved: this pull request changes .cargo/config.toml, .cargo/cuda-oxide.toml, AGENTS.md …, which changes the bot, CI or the toolchain and needs a person's approval.}

_{🔨 Shingler, bloomery's review bot: approves a clean pull request and requests changes on a must-fix (how it works)}

[midagedev]

Review fixes in 0c54c84 and 2b5d3b1:

1. backends_root resolves a relative root against the tree. The new self-test cases (relative build, then --check from /) are red on the old function and green now.
2. box.sh's duplicate backend check is gone; oxide-host.sh owns it. The box.sh header line is corrected too.
3. local-check.md names the setup step.
   On the non-blocking note: the fixture-tier full batch (261 gates, red 0) is in the comment above. No timed A/B was run: this PR moves the compiler, and the occupancy rows are re-pinned with their PIN lines.