[pedrozaz]

Summary

The WGMMA PTX fallback could report PASS after a failed Cargo run or when the expected PTX was missing or empty. Require a successful run and a nonempty artifact before accepting the fallback.

Fixes #1414

Changes

• Check Cargo exit status before accepting the WGMMA fallback.
• Require nonempty PTX before the optional ptxas gate.
• Add CPU-only regression tests for verdicts and the runner's failure summary and exit status.
• Preserve tolerated non-Hopper PTX load failures and missing/old ptxas skips.

Testing

bash cuda-oxide/scripts/check-example-smoketest-contract.sh passed:

OK: 13 WGMMA verdict cases and failed-run summary/exit status.

Bash syntax, SPDX headers, and git diff --check also passed. In temporary copies, the tests rejected the original bug and removal of either new check.

• just -f cuda-oxide/Justfile check passes
• cargo oxide run wgmma or GPU smoketest passes

---

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[Copilot]

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

[pedrozaz]

Local validation: just -f cuda-oxide/Justfile check passed (exit code 0).

[pedrozaz]

WGMMA smoketest passed on NVIDIA GeForce RTX 5060 Ti (sm_120).

• Command, from cuda-oxide/: bash scripts/smoketest.sh -o '^wgmma$' --no-color --keep-logs
• Result: PASS (wgmma, PTX compiled; ptxas sm_90a ok)
• Summary: 1 passed, 0 failed.
• Exit code: 0.

This validates the non-Hopper PTX fallback and ptxas assembly.

[pedrozaz]

I forgot the contribution PRs were paused, I'm sorry.

[0z5a]

Tested this on a GH200 (Hopper, aarch64), which adds a real-Hopper run next to the RTX 5060 Ti result.

Environment: GH200 480GB (cc 9.0), driver 595.58, CUDA 13.2, Linux aarch64, Rust 1.99.1 (repo pin). The smoketest runs below used main fc747196; the second run used this PR's smoketest.sh change applied on top.

CPU contract test (this PR's head):

$ bash cuda-oxide/scripts/check-example-smoketest-contract.sh
OK: 13 WGMMA verdict cases and failed-run summary/exit status.
OK: 16 *_EXAMPLES arrays name only real examples, no example is claimed by two categories, and all 200 standard examples print a success marker.

20 repeated runs on aarch64: 0 failures.

GPU smoketest, Hopper subset, main:

GPU: NVIDIA GH200 480GB, 9.0
[ 1/ 9] cluster                          ... PASS [18s]
[ 2/ 9] elect_leader                     ... PASS [7s]
[ 3/ 9] iket_trace                       ... PASS (iket, executed on sm_90) [3s]
[ 4/ 9] select_unpredictable             ... PASS [3s]
[ 5/ 9] shared_stage_select              ... PASS [3s]
[ 6/ 9] tma_copy                         ... PASS [3s]
[ 7/ 9] tma_multicast                    ... PASS [3s]
[ 8/ 9] wgmma                            ... PASS (wgmma, executed) [3s]
[ 9/ 9] wgmma_mma_bf16                   ... PASS (compiled, no launch) [3s]
Pass: 9 / 9

WGMMA pair with this PR's verdict:

[ 1/ 2] wgmma                            ... PASS (wgmma, executed) [3s]
[ 2/ 2] wgmma_mma_bf16                   ... PASS (compiled, no launch) [3s]

On Hopper, wgmma takes the executed path, so the stricter fallback checks don't change its verdict. That is the expected no-regression result. wgmma_mma_bf16 is compile-only by design (NO_LAUNCH_EXAMPLES).

• contract test passes on aarch64
• wgmma executes and passes on GH200, before and after