cub::AgentRadixSortHistogramPolicy#

template<int _BLOCK_THREADS, int _ITEMS_PER_THREAD, int NOMINAL_4B_NUM_PARTS, typename ComputeT, int _RADIX_BITS>
struct AgentRadixSortHistogramPolicy#

Public Types

enum [anonymous]#

Values:

enumerator BLOCK_THREADS = _BLOCK_THREADS#
enumerator ITEMS_PER_THREAD = _ITEMS_PER_THREAD#
enumerator NUM_PARTS = cuda::std::max(1, NOMINAL_4B_NUM_PARTS * 4 / cuda::std::max(int{sizeof(ComputeT)}, 4))#

NUM_PARTS is the number of private histograms (parts) each histogram is split into.

Each warp lane is assigned to a specific part based on the lane ID. However, lanes with the same ID in different warp use the same private histogram. This arrangement helps reduce the degree of conflicts in atomic operations.

enumerator RADIX_BITS = _RADIX_BITS#