-
Notifications
You must be signed in to change notification settings - Fork 188
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Replace parts of cub::Traits by numeric_limits and deprecate those #3381
Comments
9 tasks
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 14, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 21, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 21, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 21, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 21, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 22, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 22, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 22, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 22, 2025
TODO: This PR contains a lot more changes at the moment to test the feasibility of this deprecation. Fixes: NVIDIA#3381
bernhardmgruber
changed the title
Deprecate
Deprecate and replace Jan 22, 2025
cub::Traits
cub::Traits
bernhardmgruber
changed the title
Deprecate and replace
Replace cub::Traits by numeric_limits and deprecate it
Jan 22, 2025
cub::Traits
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 29, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 30, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 30, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 30, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 30, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 30, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
changed the title
Replace cub::Traits by numeric_limits and deprecate it
Replace parts of cub::Traits by numeric_limits and deprecate those
Jan 30, 2025
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 31, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Jan 31, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
bernhardmgruber
added a commit
to bernhardmgruber/cccl
that referenced
this issue
Feb 4, 2025
* Consistently use ::cuda::std::numeric_limits in CUB Fixes: NVIDIA#3381
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
cub::Traits
and friends provide a mix of functionality that is also partially provided by the (CUDA) C++ standard library. There should not be multiple ways to query and define traits of types, so wherever a standard library facility exists, we should deprecate that part incub::Traits
.The initial goal was to get rid of
cub::Traits
by replacing as much as possible bycuda::std::numeric_limits
and internalizing the rest. The rest turned out to be some bit twiddling utilities mostly used by radix sort. However, after a while, I noticed that those utilities may actually be nice to use outside CUB and some code search results indicate that:cub::Traits
: https://github.com/search?type=code&q=cub%3A%3ATraits%3C (152 hits)cub::NumericTraits
: https://github.com/search?type=code&q=cub%3A%3ANumericTraits%3C (82 hits)cub::BaseTraits
: https://github.com/search?type=code&q=cub%3A%3ABaseTraits%3C (31 hits)cub::FpLimits
: https://github.com/search?type=code&q=cub%3A%3AFpLimits%3C (52 hits)So I considered re-exposing the twiddling functionality as a new trait to get the unsigned bits of a type and to twiddle-in/out. However, I later realized that
cub::Traits
serves even more uses:cuda/type_traits
andcuda/std/limits
)cub::Traits
with user defined types to enable radix sort for them.The last use case is the most problematic, since users rely on injecting custom data types into CUB and that CUB algorithms are going to ask
cub::Traits
about that type. So changing anycub::Traits<T>::Max()
tocuda::std::numeric_limits<T>::max()
, or changingcuda::Traits<T>::CATEGORY == FLOATING_POINT
tocuda::is_floating_point<T>
anywhere in CUB (or any piece pulled into CUB from Thrust or libcu++), is a breaking change. This puts us in a very bad corner, and is likely already broken today, since we also rely on traits and limits from libcu++ in all kinds of places.An example of such a specialization is pytorch registering their bfloat16 wrapper
c10::BFloat16
to CUB this way. We do a similar thing in our unit test to work around half/bfloat16 not having operators for host code before CTK 12.2.I think we should at least deprecate all parts of
cub::Traits
that are covered by libcu++ limits or traits. Users should get a deprecation warning when usingcub::Traits
to ask for a type trait or limit.Beyond that, I can think of:
cub::Traits
. This way, every user providing a specialization ofcub::Traits
sees a loud warning that they should specialize limits, traits etc. For the bit twiddling, we should introduce a new trait restricted to just that functionality or just drop the functionality completely.cub::Traits
limited to the bit twiddling functionality. User code specializing it and also adding information on limits etc. will continue to work, but CUB algorithms using/moving to libcu++ limits/traits will no longer regard the user provided information. That leaves us at risk of inconsistency when refactoring CUB algorithms.cub::Traits
into an alias to a template inside adetail
namespace. This way, user code using it to query information or do bit twiddling continues to work, but users would get an error when they would specialize it.Now I know that radix sort has a differente API (decomposer) as well to describe how a key type can be turned into bits, so maybe using
cub::Traits
to buy a new type into radix sort is actually a hack and we may want to prevent users from doing that anyway. In many instances it's also probably used to work around limitations of the built-in half and bfloat16 types.The text was updated successfully, but these errors were encountered: