[0z5a]

Summary

Closes #1628.

The tcgen05 example's compile-only coverage kernels include tcgen05.mma ... kind::i8, which sm_103a does not provide. On a cc 10.3 GPU, the backend therefore builds an sm_100a module even with the detected sm_103a hint. The runtime gate accepted cc 10.3, the driver could not JIT the PTX (CUDA_ERROR_INVALID_PTX, 218), and the example failed.

The example now also checks the module's own .target before loading:

• an arch-specific target (a) needs that exact compute capability;
• a family target (f) needs the same major and at least its minor;
• a plain target needs at least its capability.

Otherwise the example takes the existing PTX-only path, which assembles the module with ptxas. The gate on the GPU generation (can_execute_tcgen05_ptx) is unchanged.

Changes

• cuda-oxide/crates/rustc-codegen-cuda/examples/tcgen05/src/main.rs: before loading the embedded module, compare its .target with the GPU's compute capability, and take the PTX-only path when the module cannot run there.

Testing

scripts/smoketest.sh -o 'tcgen05|gemm_sol':

Example	Before	After
gemm_sol	PASS (executed)	PASS (executed)
gemm_sol_final	PASS (executed)	PASS (executed)
tcgen05	FAIL (exit=1: CUDA_ERROR_INVALID_PTX)	PASS (PTX compiled; ptxas sm_100a ok)
tcgen05_matmul	PASS (executed)	PASS (executed)

scripts/smoketest.sh -o '^redux_f32$|generated_intrinsics_blackwell': 2/2 PASS.

GPU: B300 (cc 10.3, x86_64) | driver 610.57.04 | CUDA 13.3 | rustc 1.99.1 | cuda-rust fc747196

Not run on B200 (cc 10.0). There the module target matches the GPU, so the example executes as before.

• just -f cuda-oxide/Justfile check passes -- not run for this change; the change is confined to the example, which is its own workspace
• cuda-oxide/scripts/smoketest.sh -o 'tcgen05|gemm_sol' (above)
• New example added (if applicable) -- not applicable

Checklist

• All commits signed off (git commit -s)
• SPDX headers on new source files -- no new files

[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.