Skip to content

fix(mir-importer): preserve call destination evaluation order - #1229

Merged
nihalpasham merged 4 commits into
NVIDIA:mainfrom
uurl:fix/projected-call-destination-ordering
Sep 6, 2026
Merged

nihalpasham merged 4 commits into
NVIDIA:mainfrom
uurl:fix/projected-call-destination-ordering

Conversation

@uurl

@uurl uurl commented Sep 4, 2026

Copy link
Copy Markdown
Contributor

Closes #1164

Summary

MIR call destinations must be evaluated before the callee executes. This matters for projected destinations such as:

RET.1[index]
(*ptr).field
array[index].field

where evaluating the destination after the call can observe state mutated by the callee and store the result at a different location.

This PR introduces a prepared destination-write path that materializes the destination before the producer executes, retains the typed storage information across the call or intrinsic, and performs the final store afterward.

What changed

  • Added a PreparedDestinationWrite abstraction for call results:

    • bare locals continue through ValueMap::store_local
    • ordinary projected destinations retain a precomputed address
    • canonical enum payload destinations retain their typed payload-store metadata and enclosing enum address
  • Ordinary function-item, closure, and other call paths now follow:

evaluate destination
    -> call callee
    -> write result through prepared destination

instead of evaluating the projected destination after the call.

  • Generated intrinsic importers now prepare their destination before emitting the result-producing NVVM operation and use the same prepared-result epilogue.

  • Regenerated the affected intrinsic importer sources from cuda-intrinsics-gen.

  • Migrated hand-written intrinsic result paths to the same destination-write abstraction, including unit-returning paths such as core::intrinsics::copy, typed_swap_nonoverlapping, and compiler-hint no-op intrinsics.

  • Removed the old result-store helpers and their independent single-projection grammar.

  • Preserved the existing canonical payload storage path, including storage coercions such as boolean payload representation and pointer storage representation.

  • Updated the intrinsic contributor documentation to reflect the new prepared-destination flow.

Correctness regression

The original ordering bug is observable when the callee mutates a local used by the destination projection.

For example, conceptually:

_1 = 0;
Call(RET[_1] = bump(&mut _1), ...)

where bump changes _1 to 2.

The destination must resolve to RET[0] before entering the callee. Resolving it afterward incorrectly stores into RET[2].

The regression fixture compares host and device behavior and now covers:

  • mutation of an index local by the callee
  • Field -> Index
  • Deref -> Field
  • Index -> Field
  • bare-local destination control
  • projected generated-intrinsic result destination
  • existing deref, field, index, float, function-item, and sincos cases

It also includes a negative importer test ensuring that if the general place-address walker cannot represent a destination, translation fails with an unsupported-construct error instead of emitting a semantically incorrect store.

Validation

cargo fmt --all -- --check
PASS

git diff --check
PASS

cargo run -p cuda-intrinsics-gen -- check
generated intrinsic outputs are current

cargo test -p mir-importer --all-targets
1145 passed; 0 failed

cargo clippy -p mir-importer --all-targets -- -D warnings
PASS

scripts/smoketest.sh --compile-only -o '^mir_projected_call_destination$'
PASS
ptxas sm_80 ok

Runtime host/device comparison:

host folded:   0x935299a71fef4287
device folded: 0x935299a71fef4287

PASS: host/device agree on twelve call-destination cases and the generated SREG projection

Signed-off-by: Raul Estrada <raulestradaa@gmail.com>
@uurl
uurl force-pushed the fix/projected-call-destination-ordering branch from 707b56a to f02781a Compare September 4, 2026 04:33
…ture

CI's per-example fmt loop wants the five-line signature on one line.
Whitespace only, straight output of `cargo fmt --manifest-path
.../mir_projected_call_destination/Cargo.toml` under the pinned nightly.
…nder test

The renderer now emits prepare_destination_write followed by
emit_prepared_result_and_goto; the test still asserted the removed
emit_store_result_and_goto and failed. Assert both routes, in order.
@nihalpasham nihalpasham added bug Something isn't working miscompile Produces incorrect PTX/IR (wrong generated code) codegen Device code-generation pipeline (Rust MIR to IR to PTX) labels Sep 6, 2026
@nihalpasham

Copy link
Copy Markdown
Collaborator

Reviewed and validated:

Layered three commits:

Verified here: 12/12 host/device cases on the 5090, full smoketest 226/226, fork CI green.

Merging.

@nihalpasham
nihalpasham merged commit c8614eb into NVIDIA:main Sep 6, 2026
53 checks passed
nihalpasham added a commit that referenced this pull request Sep 6, 2026
* feat(wmma): add native shared-address ldmatrix forms

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>

* refactor(wmma): name ldmatrix u32 carriers explicitly

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>

* book: refresh mlir-dialects.md catalog stamp for the new shared-u32 ldmatrix forms

intrinsics/catalog.json's SHA-256 moved to 3df52944 because six existing
ldmatrix entries gained `ldmatrix_*_shared_u32` compatibility paths (entry
count unchanged at 1025; the overlay digest moved with them); the book
page's stamp still read the pre-PR 20bc41c2. Op/module counts are unchanged
(576 ops / 42 modules), so only the stamp needed updating.

* generated_ldmatrix: rustfmt the shared-address conversion

CI's per-example fmt loop reflows the three-line `unsafe { ... }` block at
src/main.rs:108 onto two lines. Whitespace only.

* merge main (#1229 rewrote the intrinsic destination-write path this PR's
  cvta lowering shares; resolved by keeping the prepare step and the widened
  u32/u64 result type)

Co-authored-by: nihalpasham <nihalp@nvidia.com>
@uurl
uurl deleted the fix/projected-call-destination-ordering branch September 16, 2026 21:51
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

bug Something isn't working codegen Device code-generation pipeline (Rust MIR to IR to PTX) miscompile Produces incorrect PTX/IR (wrong generated code)

Projects

None yet

Development

Successfully merging this pull request may close these issues.

Support chained projected destinations for call and intrinsic results

2 participants