hipCUB CCCL 3.0.3 updates

hipCUB CCCL 3.0.3 updates#

Changes have been made to hipCUB to achieve parity with CCCL 3.0.3, establish a hard dependency on libhipcxx, and align with the changes made in CUB.

hipCUB traits#

A new standard type traits system is now provided by libhipcxx and libcudacxx. Because of this the functionality and internal use of hipcub::Traits has been minimized. hipcub::Traits is now only used in hipCUB’s radix sort implementation for bit-twiddling.

hipcub::BaseTraits and hipcub::Traits can no longer be specialized for custom types. hipcub::FpLimits has been removed.

Classification of types should be done using the functionality in the hip/std/type_traits and hip/type_traits header files.

Floating-point limits should be obtained using hip::std::numeric_limits<T> instead of hipcub::FpLimits<T>.

hip::std::is_floating_point{_v} only recognizes C++ standard floating point types. hip::is_floating_point{_v} must be used to correctly classify extended types such as __half and hip_bfloat16.

Users can still specialize hipcub::NumericTraits for custom floating point types, inheriting from hipcub::BaseTraits and providing the necessary type information. The traits from libcu++ must also be specialized. For example, a custom floating point type my_half can be registered with hipCUB and libcu++ as follows:

template <>
inline constexpr bool ::hip::is_floating_point<my_half>::value = true;

template <>
class ::hip::std::numeric_limits<my_half> {
public:
  static constexpr bool is_specialized = true;
  static __host__ __device__ my_half max()    { return /* TODO */; }
  static __host__ __device__ my_half min()    { return /* TODO */; }
  static __host__ __device__ my_half lowest() { return /* TODO */; }
};

template <>
struct ::hipcub::NumericTraits<my_half> : ::hipcub::BaseTraits<FLOATING_POINT, true, uint16_t, my_half> {};

Removed macros#

Macro	Migration guidance
`HIPCUB_IF_CONSTEXPR`	Use `constexpr` instead.
`HIPCUB_IS_INT128_ENABLED`	No replacement.
`HIPCUB_MAX(a, b)`	Use `hip::std::max(a, b)` instead.
`HIPCUB_MIN(a, b)`	Use `hip::std::min(a, b)` instead.
`HIPCUB_QUOTIENT_CEILING(a, b)`	Use `hip::ceil_div(a, b)` instead.
`HIPCUB_QUOTIENT_FLOOR(a, b)`	Use plain integer division `a / b` instead.
`HIPCUB_ROUND_UP_NEAREST(a, b)`	Use `hip::round_up(a, b)` instead.
`LEGACY_PTX_ARCH`	No replacement.

Removed functions and classes#

API	Migration guidance
`hipcub::AliasTemporaries`	No replacement.
`hipcub::BAR`	No replacement.
`hipcub::BFI`	No replacement.
`hipcub::BaseTraits::CATEGORY`	Use the functionality available in `hip/std/type_traits` instead.
`hipcub::BaseTraits::nullptr_TYPE`	No replacement.
`hipcub::BaseTraits::PRIMITIVE`	Use the functionality available in `hip/std/type_traits` instead.
`hipcub::ConstantInputIterator`	Use `thrust::constant_iterator` instead.
`hipcub::CountingInputIterator`	Use `thrust::counting_iterator` instead.
`hipcub::CTA_SYNC`	Use `__syncthreads()` instead.
`hipcub::DeviceSpmv`	Use hipSPARSE instead.
`hipcub::Difference`	Use `hip::std::minus` instead.
`hipcub::DiscardOutputIterator`	Use `thrust::discard_iterator` instead.
`hipcub::DivideAndRoundUp`	Use `hip::round_up` instead.
`hipcub::Division`	Use `hip::std::divides` instead.
`hipcub::Equality`	Use `hip::std::equal_to` instead.
`hipcub::FpLimits<T>`	Use `hip::std::numeric_limits<T>` instead.
`hipcub::GridBarrier`	Use the APIs from cooperative groups instead.
`hipcub::GridBarrierLifetime`	Use the APIs from cooperative groups instead.
`hipcub::IADD3`	No replacement.
`hipcub::Inequality`	Use `hip::std::not_equal_to` instead.
`hipcub::Int2Type`	Use `hip::std::integral_constant` instead.
`hipcub::IterateThreadStore`	No replacement.
`hipcub::LaneId()`	Use `hip::ptx::get_sreg_laneid()` instead.
`hipcub::LaneMaskGe()`	Use `hip::ptx::get_sreg_lanemask_ge()` instead.
`hipcub::LaneMaskGt()`	Use `hip::ptx::get_sreg_lanemask_gt()` instead.
`hipcub::LaneMaskLe()`	Use `hip::ptx::get_sreg_lanemask_le()` instead.
`hipcub::LaneMaskLt()`	Use `hip::ptx::get_sreg_lanemask_lt()` instead.
`hipcub::max`	Use `hip::std::max` instead.
`hipcub::min`	Use `hip::std::min` instead.
`hipcub::PRMT`	Use `hip::ptx::prmt()` instead.
`hipcub::SHL_ADD`	No replacement.
`hipcub::SHR_ADD`	No replacement.
`hipcub::Sum`	Use `hip::std::plus` instead.
`hipcub::Swap(a, b)`	Use `hip::std::swap(a, b)` instead.
`hipcub::TransformInputIterator`	Use `thrust::transform_iterator` instead.
`hipcub::WARP_ALL(predicate, member_mask)`	Use `__all_sync()` instead.
`hipcub::WARP_ANY(predicate, member_mask)`	Use `__any_sync()` instead.
`hipcub::WARP_BALLOT(predicate, member_mask)`	Use `__ballot_sync()` instead.
`hipcub::WARP_SYNC`	Use `rocprim::wave_barrier()` instead.
`hipcub::WarpId()`	Use `hip::ptx::get_sreg_warpid()` instead.

Deprecations with planned removal#

hipcub::Traits<T>::Max() has been deprecated and will be removed in a later version. Users are encouraged to use hip::std::numeric_limits<T>::max() instead.

Other changes#

hipcub::DeviceReduce::{Arg}[Max|Min] now uses hip::std::numeric_limits<T>::[max|min]() instead of hipcub::Traits to determine the initial value.