Custom type traits in rocPRIM

Contents

Custom type traits in rocPRIM#

This interface is designed to enable users to provide additional type trait information to rocPRIM, facilitating better compatibility with custom types.

Accurately describing custom types is important for performance optimization and computational correctness.

Custom types that implement arithmetic operators can behave like built-in arithmetic types but might still be interpreted by rocPRIM algorithms as generic struct or class types.

The rocPRIM type traits interface lets users add custom trait information for their types, improving compatibility between these types and rocPRIM algorithms.

This interface is similar to operator overloading.

Traits should be implemented as required by specific algorithms. Some traits can’t be defined if they can be inferred from others.

Interface#

template<class T>
struct define#

Overview

This template struct provides an interface for downstream libraries to implement type traits for their custom types. Users can utilize this template struct to define traits for these types. Users should only implement traits as required by specific algorithms, and some traits cannot be defined if they can be inferred from others. This API is not static because of ODR.

Example

The example below demonstrates how to implement traits for a custom floating-point type.

// Your type definition
struct custom_float_type
{};
// Implement the traits
template<>
struct rocprim::traits::define<custom_float_type>
{
    using is_arithmetic = rocprim::traits::is_arithmetic::values<true>;
    using number_format = rocprim::traits::number_format::values<traits::number_format::kind::floating_point_type>;
    using float_bit_mask = rocprim::traits::float_bit_mask::values<uint32_t, 10, 10, 10>;
};
The example below demonstrates how to implement traits for a custom integral type.
// Your type definition
struct custom_int_type
{};
// Implement the traits
template<>
struct rocprim::traits::define<custom_int_type>
{
    using is_arithmetic = rocprim::traits::is_arithmetic::values<true>;
    using number_format = rocprim::traits::number_format::values<traits::number_format::kind::integral_type>;
    using integral_sign = rocprim::traits::integral_sign::values<traits::integral_sign::kind::signed_type>;
};

Template Parameters:

T – The type for which you want to define traits.

template<class T>
struct get#

Overview

This template struct is designed to allow rocPRIM algorithms to retrieve trait information from C++ build-in arithmetic types, rocPRIM types, and custom types. This API is not static because of ODR.

  • All member functions are compiled only when invoked.

  • Different algorithms require different traits.

Example

The following code demonstrates how to retrieve the traits of type T.

// Get the trait in a template parameter
template<class T, std::enable_if<rocprim::traits::get<T>().is_integral()>::type* = nullptr>
void get_traits_in_template_parameter(){}
// Get the trait in a function body
template<class T>
void get_traits_in_function_body(){
    constexpr auto input_traits = rocprim::traits::get<InputType>();
    // Then you can use the member functinos
    constexpr bool is_arithmetic = input_traits.is_arithmetic();
}

Template Parameters:

T – The type from which you want to retrieve the traits.

Public Functions

inline constexpr bool is_arithmetic() const#

Get the value of trait is_arithmetic.

Returns:

true if std::is_arithmetic_v<T> is true, or if type T is a rocPRIM arithmetic type, or if the is_arithmetic trait has been defined as true; otherwise, returns false.

inline constexpr bool is_fundamental() const#

Get trait is_fundamental.

Returns:

true if T is a fundamental type (that is, rocPRIM arithmetic type, void, or nullptr_t); otherwise, returns false.

inline constexpr bool is_build_in() const#

Check if the type is a build_in type, this function is different from is_fundamental, because, by implementing traits, downstream code can “hack” into rocprim to let a type be arithmetic, and by following the rules of std::is_fundamental, rocprim::is_fundamental returns a union set of std::is_fundamental and rocprim::is_arithmetic. So, to check wether a type is a build-in type, please use this function.

Returns:

true if T is a build_in type (that is, char, unsigned char, short, unsigned short, int unsigned int, long long, unsigned long long, rocprim::int128_t, rocprim::uint128_t, rocprim::half, float, double);

inline constexpr bool is_compound() const#

If T is fundamental type, then returns false.

Returns:

false if T is a fundamental type (that is, rocPRIM arithmetic type, void, or nullptr_t); otherwise, returns true.

inline constexpr bool is_floating_point() const#

To check if T is floating-point type.

Warning

You cannot call this function when is_arithmetic() returns false; doing so will result in a compile-time error.

inline constexpr bool is_integral() const#

To check if T is integral type.

Warning

You cannot call this function when is_arithmetic() returns false; doing so will result in a compile-time error.

inline constexpr bool is_signed() const#

To check if T is signed integral type.

Warning

You cannot call this function when is_integral() returns false; doing so will result in a compile-time error.

inline constexpr bool is_unsigned() const#

To check if T is unsigned integral type.

Warning

You cannot call this function when is_integral() returns false; doing so will result in a compile-time error.

inline constexpr bool is_scalar() const#

Get trait is_scalar.

Returns:

true if std::is_scalar_v<T> is true, or if type T is a rocPRIM arithmetic type, or if the is_scalar trait has been defined as true; otherwise, returns false.

inline constexpr auto float_bit_mask() const#

Get trait float_bit_mask.

Warning

You cannot call this function when is_floating_point() returns false; doing so will result in a compile-time error.

Returns:

A constexpr instance of the specialization of rocprim::traits::float_bit_mask::values as provided in the traits definition of type T. If the float_bit_mask trait is not defined, it returns the rocprim::detail::float_bit_mask values, provided a specialization of rocprim::detail::float_bit_mask<T> exists.

template<bool Descending = false>
inline constexpr auto radix_key_codec() const#

Get trait radix_key_codec.

Returns:

A constexpr instance of the specialization of rocprim::traits::radix_key_codec::codec as provided in the traits definition of type T.

Available traits#

struct is_arithmetic#

Definability

  • Undefinable: For types with predefined traits.

  • Optional: For other types.

How to define
using is_arithmetic = rocprim::traits::is_arithmetic::values<true>;

How to use
rocprim::traits::get<InputType>().is_arithmetic();

template<bool Val>
struct values#

Value of this trait.

Public Static Attributes

static constexpr auto value = Val#

This indicates if the InputType is arithmetic.

struct is_scalar#

Arithmetic types, pointers, member pointers, and null pointers are considered scalar types.

Definability

  • Undefinable: For types with predefined traits.

  • Optional: For other types. If both is_arithmetic and is_scalar are defined, their values must be consistent; otherwise, a compile-time error will occur.

How to define
using is_scalar = rocprim::traits::is_scalar::values<true>;

How to use
rocprim::traits::get<InputType>().is_scalar();

template<bool Val>
struct values#

Value of this trait.

Public Static Attributes

static constexpr auto value = Val#

This indicates if the InputType is scalar.

struct number_format#

Definability

  • Undefinable: For types with predefined traits and non-arithmetic types.

  • Required: If you define is_arithmetic as true, you must also define this trait; otherwise, a compile-time error will occur.

How to define
using number_format = rocprim::traits::number_format::values<number_format::kind::integral_type>;

How to use
rocprim::traits::get<InputType>().is_integral();
rocprim::traits::get<InputType>().is_floating_point();

Public Types

enum class kind#

The kind enum that indecates the values avaliable for this trait.

Values:

enumerator unknown_type#
enumerator floating_point_type#
enumerator integral_type#
template<kind Val>
struct values#

Value of this trait.

Public Static Attributes

static constexpr auto value = Val#

This indicates if the InputType is floating_point_type or integral_type or unknown_type.

struct integral_sign#

Definability

  • Undefinable: For types with predefined traits, non-arithmetic types and floating-point types.

  • Required: If you define number_format as number_format::kind::floating_point_type, you must also define this trait; otherwise, a compile-time error will occur.

How to define
using integral_sign = rocprim::traits::integral_sign::values<traits::integral_sign::kind::signed_type>;

How to use
rocprim::traits::get<InputType>().is_signed();
rocprim::traits::get<InputType>().is_unsigned();

Public Types

enum class kind#

The kind enum that indecates the values avaliable for this trait.

Values:

enumerator unknown_type#
enumerator signed_type#
enumerator unsigned_type#
template<kind Val>
struct values#

Value of this trait.

Public Static Attributes

static constexpr auto value = Val#

This indicates if the InputType is signed_type or unsigned_type or unknown_type.

struct float_bit_mask#

Definability

  • Undefinable: For types with predefined traits, non-arithmetic types and integral types.

  • Required: If you define number_format as number_format::kind::unknown_type, you must also define this trait; otherwise, a compile-time error will occur.

How to define
using float_bit_mask = rocprim::traits::float_bit_mask::values<int,1,1,1>;

How to use
rocprim::traits::get<InputType>().float_bit_mask();

template<class BitType, BitType SignBit, BitType Exponent, BitType Mantissa>
struct values#

Value of this trait.

Public Static Attributes

static constexpr BitType sign_bit = SignBit#

Trait sign_bit for the InputType.

static constexpr BitType exponent = Exponent#

Trait exponent for the InputType.

static constexpr BitType mantissa = Mantissa#

Trait mantissa for the InputType.

Warning

doxygenstruct: Cannot find class “rocprim::traits::is_fundamental” in doxygen xml output for project “rocPRIM” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprim/checkouts/develop/projects/rocprim/docs/doxygen/xml

struct radix_key_codec#

Definability

  • Undefinable: For all types.

Overview This trait is automatically generated.

How to use
constexpr auto codec = rocprim::traits::get<InputType>().radix_key_codec();
using codec_t = decltype(codec);

Public Types

template<class T>
using radix_key_fundamental = typename has_bit_key_type<codec_base<T>>::result#

Determines whether the type is fundamental for radix_key.

Public Static Functions

template<class Key, bool Descending>
static inline constexpr auto get()#

The getter of this trait.

Template Parameters:

Key – type of the radix key

Returns:

The specialization of rocprim::traits::radix_key_codec::codec.

template<class Key, bool Descending = false, bool is_fundamental = radix_key_fundamental<Key>::value>
class codec : protected codec_base<Key>#

codec_base wrapper for fundamental radix key types

Public Types

using bit_key_type = typename base_type::bit_key_type#

Type of the encoded key.

Public Static Functions

template<class Decomposer = ::rocprim::identity_decomposer>
__host__ __device__ static inline bit_key_type encode(Key key, Decomposer decomposer = {})#

Encodes a key of type Key into bit_key_type.

Template Parameters:

Decomposer – Being Key a fundamental type, Decomposer should be identity_decomposer. This is also the type by default.

Parameters:
  • key[in] Key to encode.

  • decomposer[in] [optional] Decomposer functor.

Returns:

A bit_key_type encoded key.

template<class Decomposer = ::rocprim::identity_decomposer>
__host__ __device__ static inline void encode_inplace(Key &key, Decomposer decomposer = {})#

Encodes in-place a key of type Key.

Template Parameters:

Decomposer – Being Key a fundamental type, Decomposer should be identity_decomposer. This is also the type by default.

Parameters:
  • key[inout] Key to encode.

  • decomposer[in] [optional] Decomposer functor.

template<class Decomposer = ::rocprim::identity_decomposer>
__host__ __device__ static inline Key decode(bit_key_type bit_key, Decomposer decomposer = {})#

Decodes an encoded key of type bit_key_type back into Key.

Template Parameters:

Decomposer – Being Key a fundamental type, Decomposer should be identity_decomposer. This is also the type by default.

Parameters:
  • bit_key[in] Key to decode.

  • decomposer[in] [optional] Decomposer functor.

Returns:

A Key decoded key.

template<class Decomposer = ::rocprim::identity_decomposer>
__host__ __device__ static inline void decode_inplace(Key &key, Decomposer decomposer = {})#

Decodes in-place an encoded key of type Key.

Template Parameters:

Decomposer – Being Key a fundamental type, Decomposer should be identity_decomposer. This is also the type by default.

Parameters:
  • key[inout] Key to decode.

  • decomposer[in] [optional] Decomposer functor.

__host__ __device__ static inline unsigned int extract_digit(bit_key_type bit_key, unsigned int start, unsigned int radix_bits)#

Extracts the specified bits from a given encoded key.

Parameters:
  • bit_key[in] Encoded key.

  • start[in] Start bit of the sequence of bits to extract.

  • radix_bits[in] How many bits to extract.

Returns:

Requested bits from the key.

template<class Decomposer = ::rocprim::identity_decomposer>
__host__ __device__ static inline unsigned int extract_digit(Key key, unsigned int start, unsigned int radix_bits, Decomposer decomposer = {})#

Extracts the specified bits from a given in-place encoded key.

Template Parameters:

Decomposer – Being Key a fundamental type, Decomposer should be identity_decomposer. This is also the type by default.

Parameters:
  • key[in] Key.

  • start[in] Start bit of the sequence of bits to extract.

  • radix_bits[in] How many bits to extract.

  • decomposer[in] [optional] Decomposer functor.

Returns:

Requested bits from the key.

template<class Decomposer = ::rocprim::identity_decomposer>
__host__ __device__ static inline Key get_out_of_bounds_key(Decomposer decomposer = {})#

Gives the default value for out-of-bound keys of type Key.

Template Parameters:

Decomposer – Being Key a fundamental type, Decomposer should be identity_decomposer. This is also the type by default.

Parameters:

decomposer[in] [optional] Decomposer functor.

Returns:

Out-of-bound keys’ value.

template<class Key, bool Descending>
class codec<Key, Descending, false>#

Specialization of class codec for non-fundamental radix key types.

Public Types

using bit_key_type = Key#

The key in this case is a custom type, so bit_key_type cannot be the type of the encoded key because it depends on the decomposer used. It is thus set as the type Key.

Public Static Functions

template<class Decomposer>
__host__ __device__ static inline bit_key_type encode(Key key, Decomposer decomposer = {})#

Encodes a key of type Key into bit_key_type.

Template Parameters:

Decomposer – Decomposer functor type. Being Key a custom key type, the decomposer type must be other than the identity_decomposer.

Parameters:
  • key[in] Key to encode.

  • decomposer[in] [optional] Key is a custom key type, so a custom decomposer functor that returns a rocprim::tuple of references to fundamental types from a Key key is needed.

Returns:

A bit_key_type encoded key.

template<class Decomposer>
__host__ __device__ static inline void encode_inplace(Key &key, Decomposer decomposer = {})#

Encodes in-place a key of type Key.

Template Parameters:

Decomposer – Decomposer functor type. Being Key a custom key type, the decomposer type must be other than the identity_decomposer.

Parameters:
  • key[inout] Key to encode.

  • decomposer[in] [optional] Key is a custom key type, so a custom decomposer functor that returns a rocprim::tuple of references to fundamental types from a Key key is needed.

template<class Decomposer>
__host__ __device__ static inline Key decode(bit_key_type bit_key, Decomposer decomposer = {})#

Decodes an encoded key of type bit_key_type back into Key.

Template Parameters:

Decomposer – Decomposer functor type. Being Key a custom key type, the decomposer type must be other than the identity_decomposer.

Parameters:
  • bit_key[in] Key to decode.

  • decomposer[in] [optional] Key is a custom key type, so a custom decomposer functor that returns a rocprim::tuple of references to fundamental types from a Key key is needed.

Returns:

A Key decoded key.

template<class Decomposer>
__host__ __device__ static inline void decode_inplace(Key &key, Decomposer decomposer = {})#

Decodes in-place an encoded key of type Key.

Template Parameters:

Decomposer – Decomposer functor type. Being Key a custom key type, the decomposer type must be other than the identity_decomposer.

Parameters:
  • key[inout] Key to decode.

  • decomposer[in] [optional] Decomposer functor.

__host__ __device__ static inline unsigned int extract_digit(bit_key_type, unsigned int, unsigned int)#

Extracts the specified bits from a given encoded key.

Returns:

Requested bits from the key.

template<class Decomposer>
__host__ __device__ static inline unsigned int extract_digit(Key key, unsigned int start, unsigned int radix_bits, Decomposer decomposer)#

Extracts the specified bits from a given in-place encoded key.

Template Parameters:

Decomposer – Decomposer functor type. Being Key a custom key type, the decomposer type must be other than the identity_decomposer.

Parameters:
  • key[in] Key.

  • start[in] Start bit of the sequence of bits to extract.

  • radix_bits[in] How many bits to extract.

  • decomposer[in] Key is a custom key type, so a custom decomposer functor that returns a rocprim::tuple of references to fundamental types from a Key key is needed.

Returns:

Requested bits from the key.

template<class Decomposer>
__host__ __device__ static inline Key get_out_of_bounds_key(Decomposer decomposer)#

Gives the default value for out-of-bound keys of type Key.

Template Parameters:

Decomposer – Decomposer functor type. Being Key a custom key type, the decomposer type must be other than the identity_decomposer.

Parameters:

decomposer[in] Key is a custom key type, so a custom decomposer functor that returns a rocprim::tuple of references to fundamental types from a Key key is needed.

Returns:

Out-of-bound keys’ value.

Type traits wrappers#

template<class T>
struct is_floating_point : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_floating_point()>

An extension of std::is_floating_point that supports additional arithmetic types, including rocprim::half, rocprim::bfloat16, and any types with trait rocprim::traits::number_format::values<number_format::kind::floating_point_type> implemented.

template<class T>
struct is_integral : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_integral()>

An extension of std::is_integral that supports additional arithmetic types, including rocprim::int128_t, rocprim::uint128_t, and any types with trait rocprim::traits::number_format::values<number_format::kind::integral_type> implemented.

template<class T>
struct is_arithmetic : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_arithmetic()>

An extension of std::is_arithmetic that supports additional arithmetic types, including any types with trait rocprim::traits::is_arithmetic::values<true> implemented.

template<class T>
struct is_fundamental : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_fundamental()>

An extension of std::is_fundamental that supports additional arithmetic types, including any types with trait rocprim::traits::is_arithmetic::values<true> implemented.

template<class T>
struct is_unsigned : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_unsigned()>

An extension of std::is_unsigned that supports additional arithmetic types, including rocprim::uint128_t, and any types with trait rocprim::traits::integral_sign::values<integral_sign::kind::unsigned_type> implemented.

template<class T>
struct is_signed : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_signed()>

An extension of std::is_signed that supports additional arithmetic types, including rocprim::int128_t, and any types with trait rocprim::traits::integral_sign::values<integral_sign::kind::signed_type> implemented.

template<class T>
struct is_scalar : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_scalar()>

An extension of std::is_scalar that supports additional arithmetic types, including any types with trait rocprim::traits::is_scalar::values<true> implemented.

template<class T>
struct is_compound : public std::integral_constant<bool, ::rocprim::traits::get<T>().is_compound()>

An extension of std::is_scalar that supports additional non-arithmetic types.

Types with predefined traits#

template<>
struct define<float>

Public Types

using float_bit_mask = traits::float_bit_mask::values<uint32_t, 0x80000000, 0x7F800000, 0x007FFFFF>
template<>
struct define<double>

Public Types

using float_bit_mask = traits::float_bit_mask::values<uint64_t, 0x8000000000000000, 0x7FF0000000000000, 0x000FFFFFFFFFFFFF>
template<>
struct define<rocprim::bfloat16>

Public Types

using is_arithmetic = traits::is_arithmetic::values<true>
using number_format = traits::number_format::values<traits::number_format::kind::floating_point_type>
using float_bit_mask = traits::float_bit_mask::values<uint16_t, 0x8000, 0x7F80, 0x007F>
template<>
struct define<rocprim::half>

Public Types

using is_arithmetic = traits::is_arithmetic::values<true>
using number_format = traits::number_format::values<traits::number_format::kind::floating_point_type>
using float_bit_mask = traits::float_bit_mask::values<uint16_t, 0x8000, 0x7C00, 0x03FF>
template<>
struct define<rocprim::int128_t> : public std::conditional_t<std::is_arithmetic<rocprim::int128_t>::value, traits::define<void>, detail::define_int128_t>
template<>
struct define<rocprim::uint128_t> : public std::conditional_t<std::is_arithmetic<rocprim::uint128_t>::value, traits::define<void>, detail::define_uint128_t>