Skip to content

[TIR][CUDA] Avoid cp.async for noncontiguous vector copies - #20911

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

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

Conversation

@0z5a

@0z5a 0z5a commented Nov 4, 2026

Copy link
Copy Markdown

InjectPTXAsyncCopy currently strips the base from vector ramps without checking their strides. A 16-byte cp.async then copies contiguous bytes even if the source or destination is a gather/scatter.

Require both ramp strides to be provably one before injecting the instruction, including predicated copies and merged dynamic shared memory. Otherwise keep the original tensor store.

Fixes #20910.

Testing

  • On GH200, the original implementation returned [0, 1, 2, 3] for a source ramp selecting [0, 2, 4, 6]; the standalone change produced the expected result.
  • The focused test file passed all 7 cases on GH200: 4 IR cases and 3 numerical GPU cases for non-unit source and destination strides. These used the TVM change alone with the pinned TIRx version.
  • ruff check, ruff format --check, and clang-format passed for the changed files.

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 treats strided vector accesses as contiguous

1 participant