Skip to content

Commit

Permalink
Drop deprecated entities from CUB util_type (NVIDIA#3743)
Browse files Browse the repository at this point in the history
  • Loading branch information
@bernhardmgruber
bernhardmgruber authored Feb 9, 2025
1 parent d334eb4 commit 4a2c424
Show file tree
Hide file tree
Showing 15 changed files with 83 additions and 260 deletions.
2 changes: 1 addition & 1 deletion c2h/include/c2h/bfloat16.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -267,7 +267,7 @@ _LIBCUDACXX_END_NAMESPACE_STD
_CCCL_SUPPRESS_DEPRECATED_PUSH
template <>
struct CUB_NS_QUALIFIER::NumericTraits<bfloat16_t>
: CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, true, false, unsigned short, bfloat16_t>
: CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, unsigned short, bfloat16_t>
{};
_CCCL_SUPPRESS_DEPRECATED_POP

Expand Down
3 changes: 1 addition & 2 deletions c2h/include/c2h/half.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -361,8 +361,7 @@ _LIBCUDACXX_END_NAMESPACE_STD

_CCCL_SUPPRESS_DEPRECATED_PUSH
template <>
struct CUB_NS_QUALIFIER::NumericTraits<half_t>
: CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, true, false, unsigned short, half_t>
struct CUB_NS_QUALIFIER::NumericTraits<half_t> : CUB_NS_QUALIFIER::BaseTraits<FLOATING_POINT, unsigned short, half_t>
{};
_CCCL_SUPPRESS_DEPRECATED_POP

Expand Down
100 changes: 38 additions & 62 deletions c2h/include/c2h/test_util_vec.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -295,28 +295,22 @@ C2H_VEC_OVERLOAD(double, double)
/**
* Vector1 overloads
*/
# define C2H_VEC_1_TRAITS_OVERLOAD(T, BaseT) \
CUB_NAMESPACE_BEGIN \
template <> \
struct NumericTraits<T> \
{ \
static constexpr Category CATEGORY = NOT_A_NUMBER; \
enum \
{ \
PRIMITIVE = false, \
NULL_TYPE = false, \
}; \
static __host__ __device__ T Max() \
{ \
T retval = {NumericTraits<BaseT>::Max()}; \
return retval; \
} \
static __host__ __device__ T Lowest() \
{ \
T retval = {NumericTraits<BaseT>::Lowest()}; \
return retval; \
} \
}; \
# define C2H_VEC_1_TRAITS_OVERLOAD(T, BaseT) \
CUB_NAMESPACE_BEGIN \
template <> \
struct NumericTraits<T> \
{ \
static __host__ __device__ T Max() \
{ \
T retval = {NumericTraits<BaseT>::Max()}; \
return retval; \
} \
static __host__ __device__ T Lowest() \
{ \
T retval = {NumericTraits<BaseT>::Lowest()}; \
return retval; \
} \
}; \
CUB_NAMESPACE_END

/**
Expand All @@ -327,12 +321,6 @@ C2H_VEC_OVERLOAD(double, double)
template <> \
struct NumericTraits<T> \
{ \
static constexpr Category CATEGORY = NOT_A_NUMBER; \
enum \
{ \
PRIMITIVE = false, \
NULL_TYPE = false, \
}; \
static __host__ __device__ T Max() \
{ \
T retval = {NumericTraits<BaseT>::Max(), NumericTraits<BaseT>::Max()}; \
Expand All @@ -354,12 +342,6 @@ C2H_VEC_OVERLOAD(double, double)
template <> \
struct NumericTraits<T> \
{ \
static constexpr Category CATEGORY = NOT_A_NUMBER; \
enum \
{ \
PRIMITIVE = false, \
NULL_TYPE = false, \
}; \
static __host__ __device__ T Max() \
{ \
T retval = {NumericTraits<BaseT>::Max(), NumericTraits<BaseT>::Max(), NumericTraits<BaseT>::Max()}; \
Expand All @@ -376,34 +358,28 @@ C2H_VEC_OVERLOAD(double, double)
/**
* Vector4 overloads
*/
# define C2H_VEC_4_TRAITS_OVERLOAD(T, BaseT) \
CUB_NAMESPACE_BEGIN \
template <> \
struct NumericTraits<T> \
{ \
static constexpr Category CATEGORY = NOT_A_NUMBER; \
enum \
{ \
PRIMITIVE = false, \
NULL_TYPE = false, \
}; \
static __host__ __device__ T Max() \
{ \
T retval = {NumericTraits<BaseT>::Max(), \
NumericTraits<BaseT>::Max(), \
NumericTraits<BaseT>::Max(), \
NumericTraits<BaseT>::Max()}; \
return retval; \
} \
static __host__ __device__ T Lowest() \
{ \
T retval = {NumericTraits<BaseT>::Lowest(), \
NumericTraits<BaseT>::Lowest(), \
NumericTraits<BaseT>::Lowest(), \
NumericTraits<BaseT>::Lowest()}; \
return retval; \
} \
}; \
# define C2H_VEC_4_TRAITS_OVERLOAD(T, BaseT) \
CUB_NAMESPACE_BEGIN \
template <> \
struct NumericTraits<T> \
{ \
static __host__ __device__ T Max() \
{ \
T retval = {NumericTraits<BaseT>::Max(), \
NumericTraits<BaseT>::Max(), \
NumericTraits<BaseT>::Max(), \
NumericTraits<BaseT>::Max()}; \
return retval; \
} \
static __host__ __device__ T Lowest() \
{ \
T retval = {NumericTraits<BaseT>::Lowest(), \
NumericTraits<BaseT>::Lowest(), \
NumericTraits<BaseT>::Lowest(), \
NumericTraits<BaseT>::Lowest()}; \
return retval; \
} \
}; \
CUB_NAMESPACE_END

/**
Expand Down
10 changes: 0 additions & 10 deletions cub/cub/block/radix_rank_sort_operations.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -77,11 +77,6 @@ CUB_NAMESPACE_BEGIN
template <typename KeyT, bool IsFP = ::cuda::is_floating_point<KeyT>::value>
struct BaseDigitExtractor
{
// TODO(bgruber): sanity check, remove eventually
_CCCL_SUPPRESS_DEPRECATED_PUSH
static_assert(Traits<KeyT>::CATEGORY != FLOATING_POINT, "");
_CCCL_SUPPRESS_DEPRECATED_POP

using TraitsT = Traits<KeyT>;
using UnsignedBits = typename TraitsT::UnsignedBits;

Expand All @@ -94,11 +89,6 @@ struct BaseDigitExtractor
template <typename KeyT>
struct BaseDigitExtractor<KeyT, true>
{
// TODO(bgruber): sanity check, remove eventually
_CCCL_SUPPRESS_DEPRECATED_PUSH
static_assert(Traits<KeyT>::CATEGORY == FLOATING_POINT, "");
_CCCL_SUPPRESS_DEPRECATED_POP

using TraitsT = Traits<KeyT>;
using UnsignedBits = typename TraitsT::UnsignedBits;

Expand Down
4 changes: 2 additions & 2 deletions cub/cub/device/dispatch/tuning/tuning_run_length_encode.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -285,7 +285,7 @@ struct sm100_tuning<LengthT, KeyT, primitive_length::yes, primitive_key::yes, le
};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class LengthT>
// struct sm100_tuning<LengthT, __int128_t, primitive_length::yes, primitive_key::no, length_size::_4, key_size::_16>
// {
Expand Down Expand Up @@ -577,7 +577,7 @@ struct sm100_tuning<LengthT, double, primitive_length::yes, primitive_key::yes,
{};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class LengthT>
// struct sm100_tuning<LengthT, __int128_t, primitive_length::yes, primitive_key::no, length_size::_4, key_size::_16>
// {
Expand Down
2 changes: 1 addition & 1 deletion cub/cub/device/dispatch/tuning/tuning_scan.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -404,7 +404,7 @@ struct sm100_tuning<double, AccumT, OffsetT, op_type::plus, primitive_accum::yes
: sm90_tuning<double, primitive_op::yes, primitive_accum::yes, accum_size::_8>
{};

#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class OffsetT> struct sm100_tuning<__int128_t, OffsetT, op_type::plus, primitive_accum::no,
// offset_size::_8, accum_size::_16> : tuning<576, 21, 860, 630> {}; template <class OffsetT> struct
// sm100_tuning<__uint128_t, OffsetT, op_type::plus, primitive_accum::no, offset_size::_8, accum_size::_16>
Expand Down
4 changes: 2 additions & 2 deletions cub/cub/device/dispatch/tuning/tuning_select_if.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -667,7 +667,7 @@ struct sm100_tuning<Input,
// };

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <>
// struct sm100_tuning<__int128_t, flagged::no, keep_rejects::no, offset_size::_4, primitive::no, input_size::_16>
// {
Expand Down Expand Up @@ -837,7 +837,7 @@ struct sm100_tuning<Input,
};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <>
// struct sm100_tuning<__int128_t, flagged::yes, keep_rejects::no, offset_size::_4, primitive::no, input_size::_16>
// {
Expand Down
8 changes: 4 additions & 4 deletions cub/cub/device/dispatch/tuning/tuning_unique_by_key.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -576,7 +576,7 @@ struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::yes, key_si
};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class KeyT, class ValueT>
// struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::no, key_size::_1, val_size::_16>
// {
Expand Down Expand Up @@ -634,7 +634,7 @@ struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::yes, key_si
};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class KeyT, class ValueT>
// struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::no, key_size::_2, val_size::_16>
// {
Expand Down Expand Up @@ -695,7 +695,7 @@ struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::yes, key_si
};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class KeyT, class ValueT>
// struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::no, key_size::_4, val_size::_16>
// {
Expand Down Expand Up @@ -758,7 +758,7 @@ struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::yes, key_si
};

// TODO(gonidelis): Tune for I128.
#if CUB_IS_INT128_ENABLED
#if _CCCL_HAS_INT128()
// template <class KeyT, class ValueT>
// struct sm100_tuning<KeyT, ValueT, primitive_key::yes, primitive_val::no, key_size::_8, val_size::_16>
// {
Expand Down
Loading

0 comments on commit 4a2c424

Please sign in to comment.