OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
dispatch_radix_sort.cuh File Reference

Go to the source code of this file.

Data Structures

union  cub::TempStorage
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >
 < Signed integer type for global offsets More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy200
 SM20. More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy300
 SM30. More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy350
 SM35. More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy500
 SM50. More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy600
 SM60 (GP100) More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy610
 SM61 (GP104) More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy620
 SM62 (Tegra, less RF) More...
 
struct  cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >::Policy700
 SM70 (GV100) More...
 
struct  cub::DispatchRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetT >
 < Signed integer type for global offsets More...
 
struct  cub::DispatchRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetT >::PassConfig< UpsweepKernelT, ScanKernelT, DownsweepKernelT >
 Pass configuration structure. More...
 
struct  cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >
 < Signed integer type for global offsets More...
 
struct  cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::PassConfig< SegmentedKernelT >
 PassConfig data structure. More...
 

Namespaces

 cub
 Optional outer namespace(s)
 

Typedefs

typedef AgentRadixSortUpsweep< typename If<(ALT_DIGIT_BITS), typename ChainedPolicyT::ActivePolicy::AltUpsweepPolicy, typename ChainedPolicyT::ActivePolicy::UpsweepPolicy >::Type, KeyT, OffsetT > cub::AgentRadixSortUpsweepT
 
typedef AgentRadixSortDownsweep< typename If<(ALT_DIGIT_BITS), typename ChainedPolicyT::ActivePolicy::AltDownsweepPolicy, typename ChainedPolicyT::ActivePolicy::DownsweepPolicy >::Type, IS_DESCENDING, KeyT, ValueT, OffsetT > cub::AgentRadixSortDownsweepT
 
typedef BlockRadixSort< KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, ValueT, ChainedPolicyT::ActivePolicy::SingleTilePolicy::RADIX_BITS,(ChainedPolicyT::ActivePolicy::SingleTilePolicy::RANK_ALGORITHM==RADIX_RANK_MEMOIZE), ChainedPolicyT::ActivePolicy::SingleTilePolicy::SCAN_ALGORITHM > cub::BlockRadixSortT
 
typedef BlockLoad< KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, ChainedPolicyT::ActivePolicy::SingleTilePolicy::LOAD_ALGORITHM > cub::BlockLoadKeys
 
typedef BlockLoad< ValueT, BLOCK_THREADS, ITEMS_PER_THREAD, ChainedPolicyT::ActivePolicy::SingleTilePolicy::LOAD_ALGORITHM > cub::BlockLoadValues
 
typedef Traits< KeyT >::UnsignedBits cub::UnsignedBitsT
 
typedef AgentRadixSortUpsweep< AgentRadixSortUpsweepPolicy< BLOCK_THREADS, ITEMS_PER_THREAD, SegmentedPolicyT::LOAD_MODIFIER, RADIX_BITS >, KeyT, OffsetT > cub::BlockUpsweepT
 
typedef BlockScan< OffsetT, BLOCK_THREADS > cub::DigitScanT
 
typedef AgentRadixSortDownsweep< SegmentedPolicyT, IS_DESCENDING, KeyT, ValueT, OffsetT > cub::BlockDownsweepT
 

Enumerations

enum  {
  BLOCK_THREADS = SegmentedPolicyT::BLOCK_THREADS, ITEMS_PER_THREAD = SegmentedPolicyT::ITEMS_PER_THREAD, RADIX_BITS = SegmentedPolicyT::RADIX_BITS, TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD,
  RADIX_DIGITS = 1 << RADIX_BITS, KEYS_ONLY = Equals<ValueT, NullType>::VALUE
}
 
enum  { cub::BINS_TRACKED_PER_THREAD = BlockDownsweepT::BINS_TRACKED_PER_THREAD }
 

Functions

template<typename ChainedPolicyT , bool ALT_DIGIT_BITS, bool IS_DESCENDING, typename KeyT , typename OffsetT >
 cub::__launch_bounds__ (int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltUpsweepPolicy::BLOCK_THREADS :ChainedPolicyT::ActivePolicy::UpsweepPolicy::BLOCK_THREADS)) __global__ void DeviceRadixSortUpsweepKernel(const KeyT *d_keys
 < Signed integer type for global offsets More...
 
even_share template cub::BlockInit< TILE_ITEMS, GRID_MAPPING_RAKE > ()
 
upsweep cub::ProcessRegion (even_share.block_offset, even_share.block_end)
 
 cub::CTA_SYNC ()
 
upsweep template cub::ExtractCounts< IS_DESCENDING > (d_spine, gridDim.x, blockIdx.x)
 
template<typename ChainedPolicyT , typename OffsetT >
 cub::__launch_bounds__ (int(ChainedPolicyT::ActivePolicy::ScanPolicy::BLOCK_THREADS), 1) __global__ void RadixSortScanBinsKernel(OffsetT *d_spine
 < Signed integer type for global offsets More...
 
AgentScanT cub::block_scan (temp_storage, d_spine, d_spine, cub::Sum(), OffsetT(0))
 
 cub::while (block_offset+AgentScanT::TILE_ITEMS<=num_counts)
 
template<typename ChainedPolicyT , bool ALT_DIGIT_BITS, bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetT >
 cub::__launch_bounds__ (int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltDownsweepPolicy::BLOCK_THREADS :ChainedPolicyT::ActivePolicy::DownsweepPolicy::BLOCK_THREADS)) __global__ void DeviceRadixSortDownsweepKernel(const KeyT *d_keys_in
 < Signed integer type for global offsets More...
 
 cub::AgentRadixSortDownsweepT (temp_storage, num_items, d_spine, d_keys_in, d_keys_out, d_values_in, d_values_out, current_bit, num_bits).ProcessRegion(even_share.block_offset
 
template<typename ChainedPolicyT , bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetT >
 cub::__launch_bounds__ (int(ChainedPolicyT::ActivePolicy::SingleTilePolicy::BLOCK_THREADS), 1) __global__ void DeviceRadixSortSingleTileKernel(const KeyT *d_keys_in
 < Signed integer type for global offsets More...
 
 cub::BlockLoadKeys (temp_storage.load_keys).Load(d_keys_in
 
 cub::if (!KEYS_ONLY)
 
 cub::BlockRadixSortT (temp_storage.sort).SortBlockedToStriped(keys
 
 cub::Int2Type< IS_DESCENDING > ()
 
 cub::Int2Type< KEYS_ONLY > ())
 
 cub::for (int ITEM=0;ITEM< ITEMS_PER_THREAD;++ITEM)
 
template<typename ChainedPolicyT , bool ALT_DIGIT_BITS, bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >
 cub::__launch_bounds__ (int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltSegmentedPolicy::BLOCK_THREADS :ChainedPolicyT::ActivePolicy::SegmentedPolicy::BLOCK_THREADS)) __global__ void DeviceSegmentedRadixSortKernel(const KeyT *d_keys_in
 < Signed integer type for global offsets More...
 
 cub::if (num_items<=0) return
 
BlockUpsweepT cub::upsweep (temp_storage.upsweep, d_keys_in, current_bit, pass_bits)
 
upsweep cub::ProcessRegion (segment_begin, segment_end)
 
upsweep cub::ExtractCounts (bin_count)
 
 cub::if (IS_DESCENDING)
 
 cub::DigitScanT (temp_storage.scan).ExclusiveSum(bin_count
 
BlockDownsweepT cub::downsweep (temp_storage.downsweep, bin_offset, num_items, d_keys_in, d_keys_out, d_values_in, d_values_out, current_bit, pass_bits)
 

Variables

OffsetT * cub::d_spine
 < [in] Input keys buffer More...
 
OffsetT cub::OffsetT
 [in] Total number of input data items
 
OffsetT int cub::current_bit
 [in] Bit position of current radix digit
 
OffsetT int int cub::num_bits
 [in] Number of bits of current radix digit
 
OffsetT int int GridEvenShare< OffsetT > cub::even_share
 < [in] Even-share descriptor for mapan equal number of tiles onto each thread block More...
 
AgentRadixSortUpsweepT cub::upsweep (temp_storage, d_keys, current_bit, num_bits)
 
int cub::num_counts
 < [in,out] Privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.) More...
 
int cub::block_offset = 0
 
BlockScanRunningPrefixOp< OffsetT, Sum > cub::prefix_op (0, Sum())
 
KeyT * cub::d_keys_out
 < [in] Input keys buffer More...
 
KeyT const ValueT * cub::d_values_in
 [in] Input values buffer More...
 
KeyT const ValueT ValueT * cub::d_values_out
 [in] Output values buffer
 
KeyT const ValueT ValueT OffsetT OffsetT cub::num_items = segment_end - segment_begin
 [in] Total number of input data items More...
 
even_share cub::block_end
 
KeyT const ValueT ValueT OffsetT int int cub::end_bit
 < [in] The past-the-end (most-significant) bit index needed for key comparison More...
 
KeyT cub::keys [ITEMS_PER_THREAD]
 
ValueT cub::values [ITEMS_PER_THREAD]
 
UnsignedBitsT cub::default_key_bits = (IS_DESCENDING) ? Traits<KeyT>::LOWEST_KEY : Traits<KeyT>::MAX_KEY
 
KeyT cub::default_key = reinterpret_cast<KeyT&>(default_key_bits)
 
KeyT const ValueT ValueT OffsetIteratorT cub::d_begin_offsets
 [in] Pointer to the sequence of beginning offsets of length num_segments, such that d_begin_offsets[i] is the first element of the ith data segment in d_keys_* and d_values_*
 
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT cub::d_end_offsets
 [in] Pointer to the sequence of ending offsets of length num_segments, such that d_end_offsets[i]-1 is the last element of the ith data segment in d_keys_* and d_values_*. If d_end_offsets[i]-1 <= d_begin_offsets[i], the ith is considered empty.
 
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT cub::int
 [in] The number of segments that comprise the sorting data
 
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT int int cub::pass_bits
 < [in] Number of bits of current radix digit More...
 
OffsetT cub::segment_begin = d_begin_offsets[blockIdx.x]
 
OffsetT cub::segment_end = d_end_offsets[blockIdx.x]
 
OffsetT cub::bin_count [BINS_TRACKED_PER_THREAD]
 
OffsetT cub::bin_offset [BINS_TRACKED_PER_THREAD]
 

Detailed Description

cub::DeviceRadixSort provides device-wide, parallel operations for computing a radix sort across a sequence of data items residing within device-accessible memory.

Definition in file dispatch_radix_sort.cuh.