Skip to content
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

Open
Tracked by #101
bernhardmgruber opened this issue Jan 14, 2025 · 0 comments · May be fixed by #3384
Open
Tracked by #101

Replace parts of cub::Traits by numeric_limits and deprecate those #3381

bernhardmgruber opened this issue Jan 14, 2025 · 0 comments · May be fixed by #3384

Comments

@bernhardmgruber
Copy link
Contributor

bernhardmgruber commented Jan 14, 2025

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 in cub::Traits.

The initial goal was to get rid of cub::Traits by replacing as much as possible by cuda::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:

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:

  • traits and limits (those could be replaced by libcu++ cuda/type_traits and cuda/std/limits)
  • bit twiddling utilities (mostly for radix sort)
  • specializing 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 any cub::Traits<T>::Max() to cuda::std::numeric_limits<T>::max(), or changing cuda::Traits<T>::CATEGORY == FLOATING_POINT to cuda::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 using cub::Traits to ask for a type trait or limit.

Beyond that, I can think of:

  1. Eager: deprecate and eventually remove cub::Traits . This way, every user providing a specialization of cub::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.
  2. Conservative: retaining 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.
  3. Half-way: turning cub::Traits into an alias to a template inside a detail 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.

@github-project-automation github-project-automation bot moved this to Todo in CCCL Jan 14, 2025
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
@cccl-authenticator-app cccl-authenticator-app bot moved this from Todo to In Progress in CCCL Jan 14, 2025
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 added a commit to bernhardmgruber/cccl that referenced this issue Jan 22, 2025
@bernhardmgruber bernhardmgruber changed the title Deprecate cub::Traits Deprecate and replace cub::Traits Jan 22, 2025
@bernhardmgruber bernhardmgruber changed the title Deprecate and replace cub::Traits Replace cub::Traits by numeric_limits and deprecate it Jan 22, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 22, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 22, 2025
@cccl-authenticator-app cccl-authenticator-app bot moved this from In Progress to In Review in CCCL Jan 22, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 23, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 23, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 24, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 26, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 27, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 27, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 28, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 28, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 28, 2025
bernhardmgruber added a commit to bernhardmgruber/cccl that referenced this issue Jan 29, 2025
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 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
@cccl-authenticator-app cccl-authenticator-app bot moved this from In Review to In Progress in CCCL Feb 4, 2025
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
@cccl-authenticator-app cccl-authenticator-app bot moved this from In Progress to In Review in CCCL Feb 5, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
Status: In Review
Development

Successfully merging a pull request may close this issue.

1 participant