[Backport branch/2.3.x] Add cuda::ptx::st_async
#1093
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
2 tasks
(cherry picked from commit 30ce5ab)
(cherry picked from commit dc1d934)
Because the size can be either 32 or 64 bit, this can catch a lot of errors. For instance: uint64_t * remote_buffer; uint64_t * remote_bar; cuda::ptx::st_async(remote_buffer, 1, remote_bar); would previously use the .b32 path because the `1` is an integer and determines the type resolution. Now, this will result in a compiler error. Resolution is to either (a) change the value type, or (b) change the buffer type. a) uint64_t * remote_buffer; cuda::ptx::st_async(remote_buffer, uint64_t(1), remote_bar); b) int32_t * remote_buffer; cuda::ptx::st_async(remote_buffer, 1, remote_bar); (cherry picked from commit 76044b7)
The type may be misleading on this one, so I added a note on alignment of the destination address. (cherry picked from commit 8230836)
jrhemstad
force-pushed
the
backport-1078-to-branch/2.3.x
branch
from
827d400
to
4a3f797
Compare
|
/ok to test |
jrhemstad
approved these changes
Nov 14, 2023
gevtushenko
pushed a commit
to gevtushenko/cccl
that referenced
this pull request
Dec 4, 2023
* Add st.async (cherry picked from commit 30ce5ab) * Fix usage example (cherry picked from commit dc1d934) * Use typed instead of void pointers Because the size can be either 32 or 64 bit, this can catch a lot of errors. For instance: uint64_t * remote_buffer; uint64_t * remote_bar; cuda::ptx::st_async(remote_buffer, 1, remote_bar); would previously use the .b32 path because the `1` is an integer and determines the type resolution. Now, this will result in a compiler error. Resolution is to either (a) change the value type, or (b) change the buffer type. a) uint64_t * remote_buffer; cuda::ptx::st_async(remote_buffer, uint64_t(1), remote_bar); b) int32_t * remote_buffer; cuda::ptx::st_async(remote_buffer, 1, remote_bar); (cherry picked from commit 76044b7) * Add note on alignment The type may be misleading on this one, so I added a note on alignment of the destination address. (cherry picked from commit 8230836) --------- Co-authored-by: Allard Hendriksen <ahendriksen@nvidia.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Description
Backport of #1078 to
branch/2.3.x.