< Signed integer type for global offsets More...

Detailed Description

template<bool IS_DESCENDING, typename KeyT, typename ValueT, typename OffsetIteratorT, typename OffsetT>
struct cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >

< Signed integer type for global offsets

Utility class for dispatching the appropriately-tuned kernels for segmented device-wide radix sort

Definition at line 1307 of file dispatch_radix_sort.cuh.

Inheritance diagram for cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >:

Data Structures
struct	PassConfig
	PassConfig data structure. More...

Public Types
enum	{ KEYS_ONLY = (Equals<ValueT, NullType>::VALUE) }

Public Types inherited from cub::DeviceRadixSortPolicy< KeyT, ValueT, OffsetT >
enum	{ KEYS_ONLY = (Equals<ValueT, NullType>::VALUE) }

typedef If<(sizeof(ValueT)>4)&&(sizeof(KeyT)< sizeof(ValueT)), ValueT, KeyT >::Type	DominantT

typedef Policy700	MaxPolicy
	MaxPolicy.

Public Member Functions
CUB_RUNTIME_FUNCTION __forceinline__	DispatchSegmentedRadixSort (void *d_temp_storage, size_t &temp_storage_bytes, DoubleBuffer< KeyT > &d_keys, DoubleBuffer< ValueT > &d_values, OffsetT num_items, OffsetT num_segments, OffsetIteratorT d_begin_offsets, OffsetIteratorT d_end_offsets, int begin_bit, int end_bit, bool is_overwrite_okay, cudaStream_t stream, bool debug_synchronous, int ptx_version)
	Constructor.

template<typename PassConfigT >
CUB_RUNTIME_FUNCTION __forceinline__ cudaError_t	InvokePass (const KeyT d_keys_in, KeyT d_keys_out, const ValueT d_values_in, ValueT d_values_out, int &current_bit, PassConfigT &pass_config)
	Invoke a three-kernel sorting pass at the current bit.

template<typename ActivePolicyT , typename SegmentedKernelT >
CUB_RUNTIME_FUNCTION __forceinline__ cudaError_t	InvokePasses (SegmentedKernelT segmented_kernel, SegmentedKernelT alt_segmented_kernel)
	Invocation (run multiple digit passes)

template<typename ActivePolicyT >
CUB_RUNTIME_FUNCTION __forceinline__ cudaError_t	Invoke ()
	Invocation.

Static Public Member Functions
CUB_RUNTIME_FUNCTION static __forceinline__ cudaError_t	Dispatch (void *d_temp_storage, size_t &temp_storage_bytes, DoubleBuffer< KeyT > &d_keys, DoubleBuffer< ValueT > &d_values, int num_items, int num_segments, OffsetIteratorT d_begin_offsets, OffsetIteratorT d_end_offsets, int begin_bit, int end_bit, bool is_overwrite_okay, cudaStream_t stream, bool debug_synchronous)
	Internal dispatch routine.

Data Fields
void *	d_temp_storage
	[in] Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.

size_t &	temp_storage_bytes
	[in,out] Reference to size in bytes of `d_temp_storage` allocation

DoubleBuffer< KeyT > &	d_keys
	[in,out] 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
	[in,out] Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values

OffsetT	num_items
	[in] Number of items to sort

OffsetT	num_segments
	[in] The number of segments that comprise the sorting data

OffsetIteratorT	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 i^th data segment in `d_keys_` and `d_values_`

OffsetIteratorT	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 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.

int	begin_bit
	[in] The beginning (least-significant) bit index needed for key comparison

int	end_bit
	[in] The past-the-end (most-significant) bit index needed for key comparison

cudaStream_t	stream
	[in] CUDA stream to launch kernels within. Default is stream₀.

bool	debug_synchronous
	[in] Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is `false`.

int	ptx_version
	[in] PTX version

bool	is_overwrite_okay
	[in] Whether is okay to overwrite source buffers

Member Enumeration Documentation

◆ anonymous enum

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

anonymous enum

Definition at line 1314 of file dispatch_radix_sort.cuh.

Constructor & Destructor Documentation

◆ DispatchSegmentedRadixSort()

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

CUB_RUNTIME_FUNCTION __forceinline__ cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::DispatchSegmentedRadixSort	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		DoubleBuffer< KeyT > &	d_keys,
		DoubleBuffer< ValueT > &	d_values,
		OffsetT	num_items,
		OffsetT	num_segments,
		OffsetIteratorT	d_begin_offsets,
		OffsetIteratorT	d_end_offsets,
		int	begin_bit,
		int	end_bit,
		bool	is_overwrite_okay,
		cudaStream_t	stream,
		bool	debug_synchronous,
		int	ptx_version
	)

inline

Constructor.

Definition at line 1347 of file dispatch_radix_sort.cuh.

Member Function Documentation

◆ Dispatch()

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

CUB_RUNTIME_FUNCTION static __forceinline__ cudaError_t cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::Dispatch	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		DoubleBuffer< KeyT > &	d_keys,
		DoubleBuffer< ValueT > &	d_values,
		int	num_items,
		int	num_segments,
		OffsetIteratorT	d_begin_offsets,
		OffsetIteratorT	d_end_offsets,
		int	begin_bit,
		int	end_bit,
		bool	is_overwrite_okay,
		cudaStream_t	stream,
		bool	debug_synchronous
	)

inlinestatic

Internal dispatch routine.

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in,out]	d_keys	Double-buffer whose current buffer contains the unsorted input keys and, upon return, is updated to point to the sorted output keys
[in,out]	d_values	Double-buffer whose current buffer contains the unsorted input values and, upon return, is updated to point to the sorted output values
[in]	num_items	Number of items to sort
[in]	num_segments	The number of segments that comprise the sorting data
[in]	d_begin_offsets	Pointer 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_`
[in]	d_end_offsets	Pointer 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.
[in]	begin_bit	The beginning (least-significant) bit index needed for key comparison
[in]	end_bit	The past-the-end (most-significant) bit index needed for key comparison
[in]	is_overwrite_okay	Whether is okay to overwrite source buffers
[in]	stream	CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is `false`.

Definition at line 1575 of file dispatch_radix_sort.cuh.

◆ Invoke()

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

template<typename ActivePolicyT >

CUB_RUNTIME_FUNCTION __forceinline__ cudaError_t cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::Invoke ( )

inline

Invocation.

Definition at line 1557 of file dispatch_radix_sort.cuh.

◆ InvokePass()

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

template<typename PassConfigT >

CUB_RUNTIME_FUNCTION __forceinline__ cudaError_t cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::InvokePass	(	const KeyT *	d_keys_in,
		KeyT *	d_keys_out,
		const ValueT *	d_values_in,
		ValueT *	d_values_out,
		int &	current_bit,
		PassConfigT &	pass_config
	)

inline

Invoke a three-kernel sorting pass at the current bit.

Definition at line 1387 of file dispatch_radix_sort.cuh.

◆ InvokePasses()

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

template<typename ActivePolicyT , typename SegmentedKernelT >

CUB_RUNTIME_FUNCTION __forceinline__ cudaError_t cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::InvokePasses	(	SegmentedKernelT	segmented_kernel,
		SegmentedKernelT	alt_segmented_kernel
	)

inline

Invocation (run multiple digit passes)

< Function type of cub::DeviceSegmentedRadixSortKernel

Parameters

[in]	segmented_kernel	Kernel function pointer to parameterization of cub::DeviceSegmentedRadixSortKernel
[in]	alt_segmented_kernel	Alternate kernel function pointer to parameterization of cub::DeviceSegmentedRadixSortKernel

Definition at line 1455 of file dispatch_radix_sort.cuh.

Field Documentation

◆ begin_bit

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

int cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::begin_bit

[in] The beginning (least-significant) bit index needed for key comparison

Definition at line 1333 of file dispatch_radix_sort.cuh.

◆ d_begin_offsets

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

OffsetIteratorT cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::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 i^th data segment in d_keys_* and d_values_*

Definition at line 1331 of file dispatch_radix_sort.cuh.

◆ d_end_offsets

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

OffsetIteratorT cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::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 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.

Definition at line 1332 of file dispatch_radix_sort.cuh.

◆ d_keys

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

DoubleBuffer<KeyT>& cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::d_keys

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

Definition at line 1327 of file dispatch_radix_sort.cuh.

◆ d_temp_storage

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

void* cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::d_temp_storage

[in] Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to temp_storage_bytes and no work is done.

Definition at line 1325 of file dispatch_radix_sort.cuh.

◆ d_values

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

DoubleBuffer<ValueT>& cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::d_values

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

Definition at line 1328 of file dispatch_radix_sort.cuh.

◆ debug_synchronous

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

bool cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::debug_synchronous

[in] Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console. Default is false.

Definition at line 1336 of file dispatch_radix_sort.cuh.

◆ end_bit

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

int cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::end_bit

[in] The past-the-end (most-significant) bit index needed for key comparison

Definition at line 1334 of file dispatch_radix_sort.cuh.

◆ is_overwrite_okay

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

bool cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::is_overwrite_okay

[in] Whether is okay to overwrite source buffers

Definition at line 1338 of file dispatch_radix_sort.cuh.

◆ num_items

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

OffsetT cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::num_items

[in] Number of items to sort

Definition at line 1329 of file dispatch_radix_sort.cuh.

◆ num_segments

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

OffsetT cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::num_segments

[in] The number of segments that comprise the sorting data

Definition at line 1330 of file dispatch_radix_sort.cuh.

◆ ptx_version

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

int cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::ptx_version

[in] PTX version

Definition at line 1337 of file dispatch_radix_sort.cuh.

◆ stream

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

cudaStream_t cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::stream

[in] CUDA stream to launch kernels within. Default is stream₀.

Definition at line 1335 of file dispatch_radix_sort.cuh.

◆ temp_storage_bytes

template<bool IS_DESCENDING, typename KeyT , typename ValueT , typename OffsetIteratorT , typename OffsetT >

size_t& cub::DispatchSegmentedRadixSort< IS_DESCENDING, KeyT, ValueT, OffsetIteratorT, OffsetT >::temp_storage_bytes

[in,out] Reference to size in bytes of d_temp_storage allocation

Definition at line 1326 of file dispatch_radix_sort.cuh.

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

openfpm_data/src/util/cuda/cub_old/device/dispatch/dispatch_radix_sort.cuh

Detailed Description

Data Structures

Public Types

Public Member Functions

Static Public Member Functions

Data Fields

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ DispatchSegmentedRadixSort()

Member Function Documentation

◆ Dispatch()

◆ Invoke()

◆ InvokePass()

◆ InvokePasses()

Field Documentation

◆ begin_bit

◆ d_begin_offsets

◆ d_end_offsets

◆ d_keys

◆ d_temp_storage

◆ d_values

◆ debug_synchronous

◆ end_bit

◆ is_overwrite_okay

◆ num_items

◆ num_segments

◆ ptx_version

◆ stream

◆ temp_storage_bytes