[0z5a]

Summary

cutile picks tileiras from CUDA_TOOLKIT_PATH (then CUDA_HOME, then the default install locations), but tileiras loads libNVVM and ptxas through CUDA_HOME / CUDA_PATH. When CUDA_TOOLKIT_PATH names one CUDA release and an inherited CUDA_HOME names another, the selected tileiras runs against the other release's libNVVM. Kernels with tcgen05 MMA then abort in LLVM's verifier (#1635). This PR runs a toolkit-resolved tileiras with CUDA_HOME set to the toolkit it came from and CUDA_PATH removed.

Refs #1635 (this is fix 1 from the issue; the docs note, fix 3, is not included).

Changes

• cutile-compiler/src/cuda_tile_runtime_utils.rs:
  • toolkit_root_for(tileiras) returns the toolkit root that the current environment resolves tileiras through, but only when the binary about to run is that resolved binary.
  • tileiras_command(tileiras, toolkit) sets CUDA_HOME=<root> and removes CUDA_PATH when a root is known.
  • run_tileiras_at uses it for the compile invocation.
  • A CUTILE_TILEIRAS_PATH override, a bare PATH lookup, and a binary from an earlier environment keep the inherited environment unchanged. For those, cutile does not know which toolkit the binary belongs to.
  • The --version and --list-versions probes do not load libNVVM and are unchanged. The fallback compile probes for assemblers without --list-versions are also unchanged.
• Tests:
  • tileiras_command_sets_cuda_home_only_for_a_toolkit_binary: inspects Command::get_envs.
  • toolkit_tileiras_compiles_against_its_own_toolkit: a fake toolkit's tileiras records the CUDA_HOME / CUDA_PATH it sees. With CUDA_TOOLKIT_PATH=<toolkit>, CUDA_HOME=<other> and CUDA_PATH=<other>, it gets <toolkit>|unset. Through a CUTILE_TILEIRAS_PATH override it gets the caller's <other>|<other>.

Testing

End to end, the gemm benchmark binary (--features smoke-test), with CUTILE_TILEIRAS_PATH unset and CUDA_TOOLKIT_PATH pointing at a 13.4 tile toolchain (tileiras, ptxas and libNVVM 13.4.92). Before = main, after = this branch:

B300 (sm_103), CUDA_HOME=<CUDA 13.3>
  before: exit 101  "Intrinsic has incorrect argument type! … tcgen05.mma …"
  after:  exit 0    6 timings
B300, CUDA_HOME=<CUDA 13.3> and CUDA_PATH=<CUDA 13.3>
  before: exit 101  (same abort)
  after:  exit 0
GH200 (sm_90), CUDA_HOME=<CUDA 13.2>
  before: exit 101  "failed to compile Tile IR program"
  after:  exit 0    6 timings

On GH200 (aarch64):

• cargo fmt --all --check: clean.

• cargo clippy -p cutile-compiler --all-targets -- -D warnings: clean.

• cargo clippy --workspace --exclude cuda-tile-rs --all-targets (the CI command): exit 0. No warnings in the changed file; the two warnings it prints are in other files and predate this change.

• cargo test -p cutile-compiler --lib: 149 passed (147 + 2 new).

• With only the call site reverted to the inherited environment, the new end-to-end test fails as expected:

  left:  ".../other|.../other\n"
  right: ".../toolkit|unset\n"
  

• just -f cuda-oxide/Justfile check passes (cutile-only change; ran the cutile fmt/clippy/tests above instead)

• the gemm benchmark passes on B300 and GH200 after the change (it failed before, as shown above)

• New example added (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.