cub::DispatchSegmentedSort

template<SortOrder Order, typename KeyT, typename ValueT, typename OffsetT, typename BeginOffsetIteratorT, typename EndOffsetIteratorT, typename PolicyHub = detail::segmented_sort::policy_hub<KeyT, ValueT>>
struct DispatchSegmentedSort

Public Types

using local_segment_index_t = detail::segmented_sort::local_segment_index_t

using global_segment_offset_t = detail::segmented_sort::global_segment_offset_t

Public Functions

inline DispatchSegmentedSort(

void *d_temp_storage,

size_t &temp_storage_bytes,

DoubleBuffer<KeyT> &d_keys,

DoubleBuffer<ValueT> &d_values,

::cuda::std::int64_t num_items,

global_segment_offset_t num_segments,

BeginOffsetIteratorT d_begin_offsets,

EndOffsetIteratorT d_end_offsets,

bool is_overwrite_okay,

cudaStream_t stream,

)

template<typename ActivePolicyT>
inline cudaError_t Invoke()

Public Members

void *d_temp_storage

Device-accessible allocation of temporary storage.

When nullptr, the required allocation size is written to temp_storage_bytes and no work is done.

size_t &temp_storage_bytes

Reference to size in bytes of d_temp_storage allocation.

DoubleBuffer<KeyT> &d_keys

Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys.

DoubleBuffer<ValueT> &d_values

Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values.

::cuda::std::int64_t num_items

Number of items to sort.

global_segment_offset_t num_segments

The number of segments that comprise the sorting data.

BeginOffsetIteratorT d_begin_offsets

Random-access input iterator to the sequence of beginning offsets of length num_segments, such that d_begin_offsets[i] is the first element of the i^th data segment in d_keys_* and d_values_*

EndOffsetIteratorT d_end_offsets

Random-access input iterator to the sequence of ending offsets of length num_segments, such that d_end_offsets[i]-1 is the last element of the i^th data segment in d_keys_* and d_values_*.

If d_end_offsets[i]-1 <= d_begin_offsets[i], the i^th is considered empty.

bool is_overwrite_okay

Whether is okay to overwrite source buffers.

cudaStream_t stream

CUDA stream to launch kernels within.

Public Static Functions

static inline cudaError_t Dispatch(

void *d_temp_storage,

size_t &temp_storage_bytes,

DoubleBuffer<KeyT> &d_keys,

DoubleBuffer<ValueT> &d_values,

::cuda::std::int64_t num_items,

global_segment_offset_t num_segments,

BeginOffsetIteratorT d_begin_offsets,

EndOffsetIteratorT d_end_offsets,

bool is_overwrite_okay,

cudaStream_t stream,

)

Public Static Attributes

static constexpr int KEYS_ONLY = ::cuda::std::is_same_v<ValueT, NullType>

static constexpr size_t num_selected_groups = 2

struct LargeSegmentsSelectorT

Public Functions

inline LargeSegmentsSelectorT(

OffsetT value,

BeginOffsetIteratorT d_offset_begin,

EndOffsetIteratorT d_offset_end,

)

inline bool operator()(local_segment_index_t segment_id) const

Public Members

OffsetT value = {}

BeginOffsetIteratorT d_offset_begin = {}

EndOffsetIteratorT d_offset_end = {}

global_segment_offset_t base_segment_offset = {}

struct SmallSegmentsSelectorT

Public Functions

inline SmallSegmentsSelectorT(

OffsetT value,

BeginOffsetIteratorT d_offset_begin,

EndOffsetIteratorT d_offset_end,

)

inline bool operator()(local_segment_index_t segment_id) const

Public Members

OffsetT value = {}

BeginOffsetIteratorT d_offset_begin = {}

EndOffsetIteratorT d_offset_end = {}

global_segment_offset_t base_segment_offset = {}