Skip to content

[TIR][CUDA] Require aligned addresses for cp.async - #20913

Open
0z5a wants to merge 1 commit into
apache:mainfrom
0z5a:codex/async-copy-alignment
Open

0z5a wants to merge 1 commit into
apache:mainfrom
0z5a:codex/async-copy-alignment

Conversation

@0z5a

@0z5a 0z5a commented Nov 4, 2026

Copy link
Copy Markdown

InjectPTXAsyncCopy can lower a 16-byte copy to cp.async even when the source or destination address is not provably 16-byte aligned. Check declared buffer alignment and symbolic offsets before injecting the instruction; preserve the original store otherwise. The check accounts for the byte buffer used by merged dynamic shared memory.

Fixes #20912.

Testing

  • Six focused IR regression cases passed with this change alone, covering aligned addresses, unaligned source/destination row strides, and insufficient base alignment.
  • Existing aligned scalar, vector, and merged-shared-memory CPU lowering checks continued to inject the expected instructions.
  • On GH200, aligned and unaligned numerical cases passed with this change together with an experimental TIRx scalar-store fallback and the separate non-unit-stride fix. That is combined testing, not a standalone GPU result for this PR.

The current pinned TIRx may reject an unaligned vector store after the fallback rather than compile it; this PR prevents emitting an unsafe async copy, while a scalar-store fallback is a separate follow-up.

@0z5a
0z5a marked this pull request as ready for review November 5, 2026 22:00

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

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[TIR][CUDA] InjectPTXAsyncCopy may generate unaligned cp.async

1 participant