Skip to content

feat(wmma): add shared-u32 ldmatrix forms - #1236

Merged
nihalpasham merged 5 commits into
NVIDIA:mainfrom
lucifer1004:pr/1235-native-shared-address
Sep 6, 2026
Merged

nihalpasham merged 5 commits into
NVIDIA:mainfrom
lucifer1004:pr/1235-native-shared-address

Conversation

@lucifer1004

@lucifer1004 lucifer1004 commented Sep 5, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Add .shared::cta.u32 address forms for the classic ldmatrix APIs, so callers can convert a generic shared pointer once and keep byte-offset/swizzle arithmetic in the 32-bit PTX state-space address domain.

This is the first implementation slice of #1235. It does not add the counted register-MMA carrier described there.

Changes

  • add cvta_generic_to_shared_u32(*const u8) -> u32 alongside the existing u64 descriptor-address conversion
  • add x1/x2/x4 normal and transposed ldmatrix_*_shared_u32(u32) source forms
  • declare those source aliases in the ldmatrix catalog, so importer dispatch, collector recognition, ABI markers, references, and generated tests remain one projection
  • allow the existing nvvm.ldmatrix operation to carry either a compatible pointer or a .shared::cta.u32 state-space address
  • pass u32 directly to the libNVVM inline-PTX route; materialize an address-space-3 pointer only at the LLVM-NVPTX intrinsic boundary
  • extend the existing generated-ldmatrix GPU oracle and lowering tests

The large generated-file count is the repository's normal catalog-digest propagation; the semantic additions remain in the ldmatrix overlay/generator and handwritten device stubs.

Testing

  • just check passes (all formatting, Clippy, unit tests, and structural guards passed; the final cargo deny --locked check could not run because this environment has no cargo-deny subcommand)
  • cargo oxide run generated_ldmatrix --arch sm_120a
  • cargo oxide run generated_ldmatrix --emit-nvvm-ir --arch sm_120a
  • New example added (not applicable; extended the existing generated_ldmatrix oracle)

Both GPU runs passed all 128 x4 fragment checks and both 32-lane dynamic-index checks on sm_120 after the API was renamed to the explicit _shared_u32 spelling.

Additional targeted release tests passed for cuda-intrinsics-gen, cuda-device, dialect-nvvm, mir-importer, and mir-lower. Targeted Clippy passed with -D warnings, including the example in the CUDA 13.3 environment.

Checklist

  • All commits signed off (git commit -s)
  • SPDX headers on new source files (no new handwritten source file; generated files retain generated SPDX headers)

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
@lucifer1004 lucifer1004 changed the title feat(wmma): add native shared-address ldmatrix forms feat(wmma): add shared-u32 ldmatrix forms Sep 5, 2026
…dmatrix forms

intrinsics/catalog.json's SHA-256 moved to 3df52944 with this PR's new
cvta_generic_to_shared_u32 / ldmatrix_*_shared_u32 catalog entries; 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.
@nihalpasham nihalpasham added enhancement New feature or request intrinsics Device intrinsics and libdevice math mappings labels Sep 6, 2026
CI's per-example fmt loop reflows the three-line `unsafe { ... }` block at
src/main.rs:108 onto two lines. Whitespace only, straight output of
`cargo fmt --manifest-path .../generated_ldmatrix/Cargo.toml` under the
pinned nightly. Root `cargo fmt` never reaches the example workspaces,
which is how it slipped past.
…-address

# Conflicts:
#	crates/mir-importer/src/translator/terminator/intrinsics/memory.rs
@nihalpasham

Copy link
Copy Markdown
Collaborator

Reviewed and validated:

Layered three commits:

  • book catalog stamp to 3df52944; the six new _shared_u32 compatibility paths moved the hash
  • rustfmt on generated_ldmatrix/src/main.rs
  • merge of main after fix(mir-importer): preserve call destination evaluation order #1229 landed; it rewrote the intrinsic destination-write path the cvta lowering shares, resolved by keeping the prepare step and the widened result type, generated intrinsics verified current

Verified here: 128 fragments + 2x32 dynamic selections on sm_120, default and --arch sm_120a, full smoketest 226/226 on the reconciled head, fork CI green.

Merging.

@nihalpasham
nihalpasham merged commit a07f804 into NVIDIA:main Sep 6, 2026
53 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

enhancement New feature or request intrinsics Device intrinsics and libdevice math mappings

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants