rocPRIM
Classes | Public Types | Public Member Functions | Static Public Attributes | List of all members
detail::onesweep_histograms_helper< KeyType, BlockSize, ItemsPerThread, RadixBits, Descending > Struct Template Reference

Classes

struct  storage_type
 

Public Types

using counter_type = uint32_t
 
using key_codec = radix_key_codec< KeyType, Descending >
 
using bit_key_type = typename key_codec::bit_key_type
 

Public Member Functions

ROCPRIM_DEVICE ROCPRIM_INLINE counter_type & get_counter (const unsigned stripe_index, const unsigned int place, const unsigned int digit, storage_type &storage)
 
ROCPRIM_DEVICE ROCPRIM_INLINE void clear_histogram (const unsigned int flat_id, storage_type &storage)
 
template<bool IsFull>
ROCPRIM_DEVICE void count_digits_at_place (const unsigned int flat_id, const unsigned int stripe, const bit_key_type(&bit_keys)[ItemsPerThread], const unsigned int place, const unsigned int start_bit, const unsigned int current_radix_bits, const unsigned int valid_count, storage_type &storage)
 
template<bool IsFull, class KeysInputIterator , class Offset >
ROCPRIM_DEVICE void count_digits (KeysInputIterator keys_input, Offset *global_digit_counts, const unsigned int valid_count, const unsigned int begin_bit, const unsigned int end_bit, storage_type &storage)
 

Static Public Attributes

static constexpr unsigned int radix_size = 1u << RadixBits
 
static constexpr unsigned int max_digit_places = ::rocprim::detail::ceiling_div(sizeof(KeyType) * 8, RadixBits)
 
static constexpr unsigned int items_per_block = BlockSize * ItemsPerThread
 
static constexpr unsigned int digits_per_thread = ::rocprim::detail::ceiling_div(radix_size, BlockSize)
 
static constexpr unsigned int atomic_stripes = 4
 
static constexpr unsigned int histogram_counters = radix_size * max_digit_places * atomic_stripes
 

The documentation for this struct was generated from the following file: