Skip to content

fix(mir-lower): preserve address spaces through 16-byte enum fillers - #1332

Open
YunusMutlu wants to merge 1 commit into
NVIDIA:mainfrom
YunusMutlu:fix/enum-filler-global-load
Open

YunusMutlu wants to merge 1 commit into
NVIDIA:mainfrom
YunusMutlu:fix/enum-filler-global-load

Conversation

@YunusMutlu

Copy link
Copy Markdown
Contributor

Summary

Preserve known global pointer address spaces when lowering niche-encoded enums with a layout-neutral 16-byte filler. Representing the filler as [2 x i64] instead of [16 x i8] allows aggregate scalarization to retain the pointer's global address space, restoring ld.global / st.global codegen for the affected vector load and store paths.

Fixes #1331

Changes

  • Represent layout-neutral 16-byte enum fillers as [2 x i64].
  • Keep the byte-array fallback when the required 8-byte alignment is not available.
  • Add enum layout coverage for 8-byte and 16-byte alignment cases.
  • Add vectorization regressions for both load and store paths through a wide Option payload.
  • Verify that the affected operations remain ld.global.v4.b32 and st.global.v4.b32.

Testing

  • just check passes (the local mirror of CI: fmt, clippy, tests, guards, docs)
  • cargo oxide run vectorization --arch sm_86 passes
  • New example added (if applicable)

Additional validation:

  • cargo test -p mir-lower
  • cargo clippy --workspace --all-targets -- -D warnings
  • (cd crates/rustc-codegen-cuda && cargo clippy --all-targets -- -D warnings)
  • git diff --check
  • Verified on sm_86 that the load changes from ld.v4.b32 / LD.E.128 to ld.global.v4.b32 / LDG.E.128
  • Verified on sm_86 that the store changes from st.v4.b32 / ST.E.128 to st.global.v4.b32 / STG.E.128

Checklist

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

@nihalpasham nihalpasham added the cuda-oxide SIMT programming model: rustc backend, cargo oxide, cuda-device, cuda-host, book label Oct 11, 2026
Signed-off-by: Yunus Mutlu <mutluyunus708@gmail.com>
@YunusMutlu
YunusMutlu force-pushed the fix/enum-filler-global-load branch from c626c73 to 55b5c39 Compare October 13, 2026 10:51
@copy-pr-bot

copy-pr-bot Bot commented Oct 13, 2026

Copy link
Copy Markdown

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.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

cuda-oxide SIMT programming model: rustc backend, cargo oxide, cuda-device, cuda-host, book

Projects

None yet

Development

Successfully merging this pull request may close these issues.

16-byte enum fillers can erase global pointer address spaces during lowering

3 participants