[libcudacxx] Experimental try_cancel exposure #3671
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
miscco
requested changes
Feb 4, 2025
libcudacxx/include/cuda/try_cancel
Outdated
Comment on lines
145
to
149
| _LIBCUDACXX_BEGIN_NAMESPACE_CUDA | ||
|
|
||
| namespace experimental { | ||
|
|
||
| namespace __detail { |
There was a problem hiding this comment.
This should rather be put into cudax rather than libcu++
There was a problem hiding this comment.
cudax is not shipped with the CTK but experimental exposure of HW features should be available in the CTK early to allow programmers to use the HW early on without having to familiarize themselves with all the details.
libcudacxx/include/cuda/try_cancel
Outdated
| namespace __detail { | ||
|
|
||
| struct __empty_t { | ||
| __device__ void operator()(dim3); |
There was a problem hiding this comment.
Important: We need to at least define it
Suggested change
| __device__ void operator()(dim3); | |
| _CCCL_DEVICE _CCCL_HIDE_FROM_ABI void operator()(dim3) const noexcept {}; |
There was a problem hiding this comment.
I wanted this to trigger an error if a CUDA C++ program uses it.
Would defining it support that?
I think otherwise I can define it, and then issue a static_assert below.
Co-authored-by: Michael Schellenberger Costa <[email protected]>
Co-authored-by: Michael Schellenberger Costa <[email protected]>
|
pre-commit.ci autofix |
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
This PR provides experimental low-level exposure in libcudacxx for PTX
try_cancel.While experimental, it'd be very beneficial for this exposure to be available in the CTK, hence adding it to libcu++.
We should aim to stabilize it during the CUDA 13.x cycle.
This low-level wrapper delivers reasonable functionality to end-users while hiding the gnarly synchronization and pipelining required and enabling us to deliver bug-fixes and performance improvements in the future. The functionality is exposed in a portable way, and backported to all SM's that CUDA supports, but the focus is on sm_100+ (we can deliver better performance in sm_70+ later).
This PR intentionally only covers
for_each_cancelled_block.Once that API is done,
for_each_cancelled_clusterwould be an analogous API to add.Future potential extensions worth exploring:
int NStages = 1later tht enables users to do, e.g., double-buffering.Checklist