Replace cub::ArrayWrapper by cuda::std::array and deprecate it
#1764
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
bernhardmgruber
force-pushed
the
arraywrapper
branch
2 times, most recently
from
8feecc0
to
85a7610
Compare
🟨 CI Results [ Failed: 8 | Passed: 190 | Total: 198 ]
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
gevtushenko
reviewed
May 21, 2024
There was a problem hiding this comment.
It's exciting to see CUB types being replaced by standard utilities!
This change seems to lead to many codegen differences (many kernels are just BPT.TRAP 0x1 now). Before getting this PR out of draft, could you make sure that the resulting SASS matches one without your modifications? That'd save us a lot of time on benchmarking.
cub/cub/util_type.cuh
Outdated
| @@ -885,7 +885,7 @@ struct KeyValuePair<K, V, false, true> | |||
| * \brief A wrapper for passing simple static arrays as kernel parameters | |||
| */ | |||
| template <typename T, int COUNT> | |||
| struct ArrayWrapper | |||
| struct CUB_DEPRECATED ArrayWrapper | |||
There was a problem hiding this comment.
important: we tend to document deprecation in a way that orient users on what to use instead, especially given that this wrapper is used by clients:
* @deprecated [Since 2.5.0] The `cub::ArrayWrapper` is deprecated.
* Use `cuda::std::array` instead.
@miscco if I remember correctly, there were some issues with @deprecated on Sphinx end, is it still the case?
@jrhemstad if I remember correctly, we tried to postpone all deprecations until a major release to avoid source compatibility issues on minor CTK updates. Is it still the recommended way or we can relax this requirement given the small ArrayWrapper user base?
There was a problem hiding this comment.
we tried to postpone all deprecations until a major release to avoid source compatibility issues on minor CTK updates. Is it still the recommended way or we can relax this requirement given the small ArrayWrapper user base?
Adding new deprecation warnings in a minor release is fair game as far as I'm concerned. We should just make sure to include in the emitted warning to tell people how to turn the warning off.
There was a problem hiding this comment.
@gevtushenko yes currently we cannot build any \deprecated and @deprecated doxygen comment
There was a problem hiding this comment.
@miscco what does this mean for this PR now? Can I add a @deprecated bla bla comment now, or not?
There was a problem hiding this comment.
It will break our docs soon, so I believe you should do deprecated bla bla so that it is there but is not picked up by doxygen
That way when we fix this it will be easy to bring in
miscco
reviewed
May 22, 2024
bernhardmgruber
commented
May 22, 2024
bernhardmgruber
force-pushed
the
arraywrapper
branch
from
85a7610
to
65d716a
Compare
🟨 CI Results: Pass: 95%/198 | Total Time: 4d 04h | Avg Time: 30m 30s | Hits: 18%/112452
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
bernhardmgruber
force-pushed
the
arraywrapper
branch
from
65d716a
to
ca7e1dd
Compare
🟨 CI Results: Pass: 95%/198 | Total Time: 2d 22h | Avg Time: 21m 19s | Hits: 64%/112452
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
bernhardmgruber
force-pushed
the
arraywrapper
branch
from
ca7e1dd
to
c901d5d
Compare
bernhardmgruber
force-pushed
the
arraywrapper
branch
from
c901d5d
to
5345335
Compare
miscco
approved these changes
May 28, 2024
🟩 CI Results: Pass: 100%/198 | Total Time: 4d 09h | Avg Time: 32m 05s | Hits: 22%/132439
|
| # | Runner |
|---|---|
| 154 | linux-amd64-cpu16 |
| 16 | linux-arm64-cpu16 |
| 16 | linux-amd64-gpu-v100-latest-1 |
| 12 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| Thrust | |
| CUDA Experimental |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| libcu++ | |
| +/- | CUB |
| +/- | Thrust |
| CUDA Experimental |
|
I'm a bit late to the party, but I think there should be a better way than to use |
I fully agree. But I don't see where you would need to use this combination as a user of CUB. Could you please provide an example where you are forced to do that? You can do this as an issue if you like to! |
|
@bernhardmgruber Basically all the CUB block algorithms take C-style arrays in their API. At the moment I use my own array wrapper because I did not know about CUB's built-in one and I figured that I can't use As one can't return C-style arrays from functions, I think the cleanest solution would be to let CUB's API take |
So I conclude that this PR is no breaking change for you then?
Ah, I understand where you are getting at. You have a feature request for CUB APIs to take |
It isn't for me personally but I posted here because it is breaking for people that do use |
1 task
AFAIK, CUB's block algorithms take references to native arrays, so users can use any type to supply those. I agree that
I see IIUC you are asking about the combination of both use cases: Using the same type to pass arrays by value to a kernel and then using the same type+instance again to pass a native array inside the kernel to a CUB block algortihm? This use case is indeed affected. |
|
@bernhardmgruber Ok, I was not aware that passing arrays to kernels was the main use case of |
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.
The type
cub::ArrayWrapperis only used by CUB's histogram algorithms and can be replaced bycuda::std::array, which offers a superset of the functionality.The generated SASS is identical before and after this PR, except for the mangling of function names.