Skip to content

Commit

Permalink
Fixes
Browse files Browse the repository at this point in the history
  • Loading branch information
@bernhardmgruber
bernhardmgruber committed Feb 7, 2025
1 parent da7adf6 commit d577ae2
Show file tree
Hide file tree
Showing 8 changed files with 11 additions and 164 deletions.
15 changes: 0 additions & 15 deletions c2h/include/c2h/bfloat16.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -274,21 +274,6 @@ struct CUB_NS_QUALIFIER::detail::unsigned_bits<bfloat16_t, void>
using type = unsigned short;
};

_CCCL_SUPPRESS_DEPRECATED_PUSH
template <>
struct CUB_NS_QUALIFIER::detail::FpLimits<bfloat16_t>
{
static __host__ __device__ __forceinline__ bfloat16_t Max()
{
return bfloat16_t::max();
}

static __host__ __device__ __forceinline__ bfloat16_t Lowest()
{
return bfloat16_t::lowest();
}
};

template <>
struct CUB_NS_QUALIFIER::detail::NumericTraits<bfloat16_t>
: CUB_NS_QUALIFIER::detail::BaseTraits<FLOATING_POINT, true, false, unsigned short, bfloat16_t>
Expand Down
15 changes: 0 additions & 15 deletions c2h/include/c2h/half.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -369,21 +369,6 @@ struct CUB_NS_QUALIFIER::detail::unsigned_bits<half_t, void>
using type = unsigned short;
};

_CCCL_SUPPRESS_DEPRECATED_PUSH
template <>
struct CUB_NS_QUALIFIER::detail::FpLimits<half_t>
{
static __host__ __device__ __forceinline__ half_t Max()
{
return (half_t::max)();
}

static __host__ __device__ __forceinline__ half_t Lowest()
{
return half_t::lowest();
}
};

template <>
struct CUB_NS_QUALIFIER::detail::NumericTraits<half_t>
: CUB_NS_QUALIFIER::detail::BaseTraits<FLOATING_POINT, true, false, unsigned short, half_t>
Expand Down
1 change: 0 additions & 1 deletion c2h/include/c2h/test_util_vec.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -340,7 +340,6 @@ C2H_VEC_TRAITS_OVERLOAD(ulong, unsigned long)
C2H_VEC_TRAITS_OVERLOAD(ulonglong, unsigned long long)
C2H_VEC_TRAITS_OVERLOAD(float, float)
C2H_VEC_TRAITS_OVERLOAD(double, double)
_CCCL_SUPPRESS_DEPRECATED_POP

# undef C2H_VEC_TRAITS_OVERLOAD
# undef C2H_VEC_TRAITS_OVERLOAD_IMPL
Expand Down
117 changes: 10 additions & 107 deletions cub/cub/util_type.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -791,37 +791,27 @@ struct BinaryOpHasIdxParam<T,
CCCL_DEPRECATED_BECAUSE("Use ::value instead") static constexpr bool HAS_PARAM = true;
};

/**
* \brief Basic type traits categories
*/
enum CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") Category
namespace detail
{
enum Category
{
NOT_A_NUMBER,
SIGNED_INTEGER,
UNSIGNED_INTEGER,
FLOATING_POINT
};

namespace detail
{
template <Category _CATEGORY, bool _PRIMITIVE, bool _NULL_TYPE, typename _UnsignedBits, typename T>
struct BaseTraits
{
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr Category CATEGORY = _CATEGORY;
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr bool PRIMITIVE = _PRIMITIVE;
CCCL_DEPRECATED static constexpr bool NULL_TYPE = _NULL_TYPE;
};
{};

template <typename _UnsignedBits, typename T>
struct BaseTraits<UNSIGNED_INTEGER, true, false, _UnsignedBits, T>
{
using UnsignedBits = _UnsignedBits;

CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr Category CATEGORY = UNSIGNED_INTEGER;
static constexpr UnsignedBits LOWEST_KEY = UnsignedBits(0);
static constexpr UnsignedBits MAX_KEY = UnsignedBits(-1);
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr bool PRIMITIVE = true;
static constexpr bool NULL_TYPE = false;
static constexpr UnsignedBits LOWEST_KEY = UnsignedBits(0);
static constexpr UnsignedBits MAX_KEY = UnsignedBits(-1);

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
{
Expand Down Expand Up @@ -855,12 +845,9 @@ struct BaseTraits<SIGNED_INTEGER, true, false, _UnsignedBits, T>
{
using UnsignedBits = _UnsignedBits;

CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr Category CATEGORY = SIGNED_INTEGER;
static constexpr UnsignedBits HIGH_BIT = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);
static constexpr UnsignedBits LOWEST_KEY = HIGH_BIT;
static constexpr UnsignedBits MAX_KEY = UnsignedBits(-1) ^ HIGH_BIT;
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr bool PRIMITIVE = true;
CCCL_DEPRECATED static constexpr bool NULL_TYPE = false;

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
{
Expand All @@ -885,31 +872,14 @@ struct BaseTraits<SIGNED_INTEGER, true, false, _UnsignedBits, T>
}
};

template <typename T>
struct FpLimits
{
static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
{
return ::cuda::std::numeric_limits<T>::max();
}

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
{
return ::cuda::std::numeric_limits<T>::lowest();
}
};

template <typename _UnsignedBits, typename T>
struct BaseTraits<FLOATING_POINT, true, false, _UnsignedBits, T>
{
using UnsignedBits = _UnsignedBits;

CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr Category CATEGORY = FLOATING_POINT;
static constexpr UnsignedBits HIGH_BIT = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);
static constexpr UnsignedBits LOWEST_KEY = UnsignedBits(-1);
static constexpr UnsignedBits MAX_KEY = UnsignedBits(-1) ^ HIGH_BIT;
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr bool PRIMITIVE = true;
CCCL_DEPRECATED static constexpr bool NULL_TYPE = false;

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
{
Expand All @@ -925,16 +895,12 @@ struct BaseTraits<FLOATING_POINT, true, false, _UnsignedBits, T>

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Max()
{
_CCCL_SUPPRESS_DEPRECATED_PUSH
return FpLimits<T>::Max();
_CCCL_SUPPRESS_DEPRECATED_POP
return ::cuda::std::numeric_limits<T>::max();
}

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T Lowest()
{
_CCCL_SUPPRESS_DEPRECATED_PUSH
return FpLimits<T>::Lowest();
_CCCL_SUPPRESS_DEPRECATED_POP
return ::cuda::std::numeric_limits<T>::lowest();
}
};

Expand Down Expand Up @@ -964,11 +930,8 @@ struct NumericTraits<__uint128_t>
using T = __uint128_t;
using UnsignedBits = __uint128_t;

CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr Category CATEGORY = UNSIGNED_INTEGER;
static constexpr UnsignedBits LOWEST_KEY = UnsignedBits(0);
static constexpr UnsignedBits MAX_KEY = UnsignedBits(-1);
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr bool PRIMITIVE = false;
CCCL_DEPRECATED static constexpr bool NULL_TYPE = false;

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
{
Expand Down Expand Up @@ -997,12 +960,9 @@ struct NumericTraits<__int128_t>
using T = __int128_t;
using UnsignedBits = __uint128_t;

CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr Category CATEGORY = SIGNED_INTEGER;
static constexpr UnsignedBits HIGH_BIT = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);
static constexpr UnsignedBits LOWEST_KEY = HIGH_BIT;
static constexpr UnsignedBits MAX_KEY = UnsignedBits(-1) ^ HIGH_BIT;
CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") static constexpr bool PRIMITIVE = false;
CCCL_DEPRECATED static constexpr bool NULL_TYPE = false;

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits TwiddleIn(UnsignedBits key)
{
Expand Down Expand Up @@ -1044,20 +1004,18 @@ template <> struct NumericTraits<__nv_fp8_e5m2> : BaseTraits<FLOATING_POINT, t

template <> struct NumericTraits<bool> : BaseTraits<UNSIGNED_INTEGER, true, false, typename UnitWord<bool>::VolatileWord, bool> {};
// clang-format on
_CCCL_SUPPRESS_DEPRECATED_POP

template <typename T>
struct Traits : NumericTraits<typename ::cuda::std::remove_cv<T>::type>
{};
} // namespace detail

using Category CCCL_DEPRECATED_BECAUSE("Use <cuda/std/type_traits> instead") = detail::Category;

template <Category _CATEGORY, bool _PRIMITIVE, bool _NULL_TYPE, typename _UnsignedBits, typename T>
using BaseTraits CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits and cuda::is_floating_point etc. instead") =
detail::BaseTraits<_CATEGORY, _PRIMITIVE, _NULL_TYPE, _UnsignedBits, T>;

template <typename T>
using FpLimits CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits instead") = detail::FpLimits<T>;

template <typename T>
using NumericTraits CCCL_DEPRECATED_BECAUSE("Use cuda::std::numeric_limits and cuda::is_floating_point etc. instead") =
detail::NumericTraits<T>;
Expand Down Expand Up @@ -1249,61 +1207,6 @@ struct key_traits<T, ::cuda::std::enable_if_t<::cuda::is_floating_point<T>::valu
}
};

template <>
struct twiddle<bool, void>
{
using UnsignedBits = unsigned_bits_t<bool>;

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits In(UnsignedBits key)
{
return key;
}

static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE UnsignedBits Out(UnsignedBits key)
{
return key;
}
};

// __uint128_t and __int128_t are not primitive
template <typename T>
using is_primitive = ::cuda::std::bool_constant<is_one_of<
T,
char,
signed char,
short,
int,
long,
long long,
unsigned char,
unsigned short,
unsigned int,
unsigned long,
unsigned long long,
bool,
float,
double
# if _CCCL_HAS_NVFP16()
,
__half
# endif // _CCCL_HAS_NVFP16()
# if _CCCL_HAS_NVBF16()
,
__nv_bfloat16
# endif // _CCCL_HAS_NVBF16()
# if _CCCL_HAS_NVFP8()
,
__nv_fp8_e4m3,
__nv_fp8_e5m2
# endif // _CCCL_HAS_NVFP8()
>()>;

# ifndef _CCCL_NO_VARIABLE_TEMPLATES
template <typename T>
inline constexpr bool is_primitive_v = is_primitive<T>::value;
# endif // !_CCCL_NO_VARIABLE_TEMPLATES
} // namespace detail

#endif // _CCCL_DOXYGEN_INVOKED

CUB_NAMESPACE_END
5 changes: 1 addition & 4 deletions cub/cub/warp/specializations/warp_reduce_shfl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -122,11 +122,8 @@ struct WarpReduceShfl
/// Whether the data type is a small (32b or less) integer for which we can use a single SHFL instruction per
/// exchange
IS_SMALL_UNSIGNED =
::cuda::std::is_integral<S>::value && ::cuda::std::is_unsigned<S>::value && (sizeof(S) <= sizeof(unsigned int)),
// TODO(bgruber): sanity check, remove later
old_IS_SMALL_UNSIGNED = (Traits<S>::CATEGORY == UNSIGNED_INTEGER) && (sizeof(S) <= sizeof(unsigned int))
::cuda::std::is_integral<S>::value && ::cuda::std::is_unsigned<S>::value && (sizeof(S) <= sizeof(unsigned int))
};
static_assert(IS_SMALL_UNSIGNED == old_IS_SMALL_UNSIGNED, "");
};

/// Shared memory storage layout type
Expand Down
4 changes: 0 additions & 4 deletions cub/cub/warp/specializations/warp_scan_shfl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -86,17 +86,13 @@ struct WarpScanShfl
template <typename S>
struct IntegerTraits
{
_CCCL_SUPPRESS_DEPRECATED_PUSH
enum
{
/// Whether the data type is a small (32b or less) integer for which we can use a single SFHL instruction per
/// exchange
IS_SMALL_UNSIGNED =
::cuda::std::is_integral<S>::value && ::cuda::std::is_unsigned<S>::value && (sizeof(S) <= sizeof(unsigned int)),
old_IS_SMALL_UNSIGNED = (Traits<S>::CATEGORY == UNSIGNED_INTEGER) && (sizeof(S) <= sizeof(unsigned int))
};
_CCCL_SUPPRESS_DEPRECATED_POP
static_assert(IS_SMALL_UNSIGNED == old_IS_SMALL_UNSIGNED, "");
};

/// Shared memory storage layout type
Expand Down
5 changes: 0 additions & 5 deletions cub/cub/warp/warp_scan.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -175,11 +175,6 @@ private:
/// Whether the data type is an integer (which has fully-associative addition)
IS_INTEGER = cuda::std::is_integral<T>::value
};
// TODO(bgruber): sanity check, remove eventually
_CCCL_SUPPRESS_DEPRECATED_PUSH
static_assert(IS_INTEGER == ((Traits<T>::CATEGORY == SIGNED_INTEGER) || (Traits<T>::CATEGORY == UNSIGNED_INTEGER)),
"");
_CCCL_SUPPRESS_DEPRECATED_POP

/// Internal specialization.
/// Use SHFL-based scan if LOGICAL_WARP_THREADS is a power-of-two
Expand Down
13 changes: 0 additions & 13 deletions cub/test/catch2_test_util_type.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -108,16 +108,3 @@ using types = c2h::type_list<
long double
#endif // _LIBCUDACXX_HAS_NO_LONG_DOUBLE
>;

C2H_TEST("Test FpLimits agrees with numeric_limits", "[util][type]", types)
{
using T = c2h::get<0, TestType>;
CAPTURE(c2h::type_name<T>());
_CCCL_SUPPRESS_DEPRECATED_PUSH
CHECK(cub::FpLimits<T>::Max() == cuda::std::numeric_limits<T>::max());
CHECK(cub::FpLimits<T>::Lowest() == cuda::std::numeric_limits<T>::lowest());

CHECK(cub::FpLimits<const T>::Max() == cuda::std::numeric_limits<const T>::max());
CHECK(cub::FpLimits<const T>::Lowest() == cuda::std::numeric_limits<const T>::lowest());
_CCCL_SUPPRESS_DEPRECATED_POP
}

0 comments on commit d577ae2

Please sign in to comment.