Skip to content

Improve cuda::memcpy_async codegen - #5996

Merged
bernhardmgruber merged 12 commits into
NVIDIA:mainfrom
bernhardmgruber:memcpy_async_hacks
Oct 27, 2025
Merged

bernhardmgruber merged 12 commits into
NVIDIA:mainfrom
bernhardmgruber:memcpy_async_hacks

Conversation

@bernhardmgruber

@bernhardmgruber bernhardmgruber commented Sep 23, 2025

Copy link
Copy Markdown
Contributor

This PR improves the codegen for cuda::memcpy_async by ensuring a uniform data path for the bulk copy instruction and updating the mbarrier's tx count in the same branch.

Fixes: #5995
(SASS of cuda::memcpy_async version is now identical to hand-written cuda::ptx)

@copy-pr-bot

copy-pr-bot Bot commented Sep 23, 2025

Copy link
Copy Markdown
Contributor

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@bernhardmgruber bernhardmgruber changed the title Hack cuda::memcpy_async to improve codegen Hack cuda::memcpy_async to improve codegen Sep 23, 2025
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Sep 23, 2025
Comment thread libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h Outdated
Comment on lines +62 to +65
(if (__elect_from_group(__g)) {
::cuda::ptx::cp_async_bulk(
::cuda::ptx::space_cluster, ::cuda::ptx::space_global, __dest, __src, __size, __bar_handle);
::cuda::ptx::mbarrier_arrive_expect_tx(ptx::sem_release, ptx::scope_cta, ptx::space_shared, __bar_handle, __size);

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We put the transaction count update into the same branch as the bulk copy. Improves codegen.

Comment thread libcudacxx/include/cuda/__memcpy_async/dispatch_memcpy_async.h Outdated
Comment thread libcudacxx/include/cuda/__barrier/barrier_block_scope.h Outdated
@bernhardmgruber

Copy link
Copy Markdown
Contributor Author

/ok to test 446535d

@github-actions

This comment has been minimized.

Comment thread libcudacxx/include/cuda/__barrier/barrier_block_scope.h Outdated
@bernhardmgruber

Copy link
Copy Markdown
Contributor Author

/ok to test cce3315

@github-actions

This comment has been minimized.

@bernhardmgruber
bernhardmgruber force-pushed the memcpy_async_hacks branch 3 times, most recently from 6cf386b to 5f4e3f9 Compare October 23, 2025 22:36
@bernhardmgruber bernhardmgruber changed the title Hack cuda::memcpy_async to improve codegen Improve cuda::memcpy_async codegen Oct 23, 2025
@bernhardmgruber
bernhardmgruber marked this pull request as ready for review October 23, 2025 23:29
@bernhardmgruber
bernhardmgruber requested review from a team as code owners October 23, 2025 23:29
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Oct 23, 2025
@github-actions

This comment has been minimized.

Comment thread libcudacxx/include/cuda/__memcpy_async/completion_mechanism.h
Comment thread docs/libcudacxx/extended_api/asynchronous_operations/memcpy_async.rst Outdated
Comment thread docs/libcudacxx/extended_api/asynchronous_operations/memcpy_async.rst Outdated
Comment thread libcudacxx/include/cuda/__memcpy_async/cp_async_bulk_shared_global.h Outdated
@bernhardmgruber
bernhardmgruber enabled auto-merge (squash) October 27, 2025 07:53
@github-actions

This comment has been minimized.

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 2h 37m: Pass: 100%/89 | Total: 18h 11m | Max: 1h 05m | Hits: 97%/212780

See results here.

@bernhardmgruber
bernhardmgruber merged commit 1d974fb into NVIDIA:main Oct 27, 2025
198 of 201 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Oct 27, 2025
@bernhardmgruber
bernhardmgruber deleted the memcpy_async_hacks branch October 27, 2025 12:18
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

cuda::memcpy_async with cuda::barrier implementation is inefficient on sm90+

2 participants