The warp_sort class provides warp-wide methods for computing a parallel sort of items across thread warps. More...

#include <warp_sort.hpp>

Inheritance diagram for warp_sort< Key, WarpSize, Value >:

Collaboration diagram for warp_sort< Key, WarpSize, Value >:

Public Types
typedef base_type::storage_type	storage_type
	Struct used to allocate a temporary memory that is required for thread communication during operations provided by related parallel primitive. More...

Public Member Functions
template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &thread_key, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &thread_key, storage_type &storage, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type using temporary storage. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &, storage_type &, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type using temporary storage. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], storage_type &storage, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type using temporary storage. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], storage_type &, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort for any data type using temporary storage. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &thread_key, Value &thread_value, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &, Value &, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], Value(&thread_values)[ItemsPerThread], BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], Value(&thread_values)[ItemsPerThread], BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &thread_key, Value &thread_value, storage_type &storage, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type using temporary storage. More...

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key &, Value &, storage_type &, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type using temporary storage. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], Value(&thread_values)[ItemsPerThread], storage_type &storage, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize<=__AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type using temporary storage. More...

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>
ROCPRIM_DEVICE ROCPRIM_INLINE auto	sort (Key(&thread_keys)[ItemsPerThread], Value(&thread_values)[ItemsPerThread], storage_type &, BinaryFunction compare_function=BinaryFunction()) -> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void >::type
	Warp sort by key for any data type using temporary storage. More...

Detailed Description

template<class Key, unsigned int WarpSize = device_warp_size(), class Value = empty_type>
class warp_sort< Key, WarpSize, Value >

The warp_sort class provides warp-wide methods for computing a parallel sort of items across thread warps.

This class currently implements parallel bitonic sort, and only accepts warp sizes that are powers of two.

Template Parameters

Key	Data type for parameter Key
WarpSize	[optional] The number of threads in a warp
Value	[optional] Data type for parameter Value. By default, it's empty_type

Overview

WarpSize must be power of two.
WarpSize must be equal to or less than the size of hardware warp (see rocprim::device_warp_size()). If it is less, sort is performed separately within groups determined by WarpSize. For example, if WarpSize is 4, hardware warp is 64, sort will be performed in logical warps grouped like this: { {0, 1, 2, 3}, {4, 5, 6, 7 }, ..., {60, 61, 62, 63} } (thread is represented here by its id within hardware warp).
Accepts custom compare_functions for sorting across a warp.
Number of threads executing warp_sort's function must be a multiple of WarpSize.

Example:

Every thread within the warp uses the warp_sort class by first specializing the warp_sort type, and instantiating an object that will be used to invoke a member function.

__global__ void example_kernel(...)
{
    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    int value = input[i];
    rocprim::warp_sort<int, 64> wsort;
    wsort.sort(value);
    input[i] = value;
}

Below is a snippet demonstrating how to pass a custom compare function:

__device__ bool customCompare(const int& a, const int& b)
{
    return a < b;
}
...
__global__ void example_kernel(...)
{
    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    int value = input[i];
    rocprim::warp_sort<int, 64> wsort;
    wsort.sort(value, customCompare);
    input[i] = value;
}

Member Typedef Documentation

◆ storage_type

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

typedef base_type::storage_type warp_sort< Key, WarpSize, Value >::storage_type

Struct used to allocate a temporary memory that is required for thread communication during operations provided by related parallel primitive.

Depending on the implemention the operations exposed by parallel primitive may require a temporary storage for thread communication. The storage should be allocated using keywords shared. It can be aliased to an externally allocated memory, or be a part of a union with other storage types to increase shared memory reusability.

Member Function Documentation

◆ sort() [1/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	thread_key,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_key	- input/output to pass to other threads
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

◆ sort() [2/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type.

Invalid Warp Size

◆ sort() [3/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_keys	- input/output keys to pass to other threads
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

◆ sort() [4/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type.

Invalid Warp Size

◆ sort() [5/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	thread_key,
		storage_type &	storage,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type using temporary storage.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_key	- input/output to pass to other threads
storage	- temporary storage for inputs
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

Storage reusage: Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.: __global__ void example_kernel(...)
{
int value = ...;
using warp_sort_int = rp::warp_sort<int, 64>;
warp_sort_int wsort;
__shared__ typename warp_sort_int::storage_type storage;
wsort.sort(value, storage);
...
}

◆ sort() [6/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	,
		storage_type &	,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type using temporary storage.

Invalid Warp Size

◆ sort() [7/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		storage_type &	storage,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type using temporary storage.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_keys	- input/output keys to pass to other threads
storage	- temporary storage for inputs
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

Storage reusage: Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.: __global__ void example_kernel(...)
{
int value = ...;
using warp_sort_int = rp::warp_sort<int, 64>;
warp_sort_int wsort;
__shared__ typename warp_sort_int::storage_type storage;
wsort.sort(value, storage);
...
}

◆ sort() [8/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		storage_type &	,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort for any data type using temporary storage.

Invalid Warp Size

◆ sort() [9/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	thread_key,
		Value &	thread_value,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_key	- input/output key to pass to other threads
thread_value	- input/output value to pass to other threads
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

◆ sort() [10/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	,
		Value &	,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type.

Invalid Warp Size

◆ sort() [11/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		Value(&)	thread_values[ItemsPerThread],
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_keys	- input/output keys to pass to other threads
thread_values	- input/outputs values to pass to other threads
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

◆ sort() [12/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		Value(&)	thread_values[ItemsPerThread],
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type.

Invalid Warp Size

◆ sort() [13/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	thread_key,
		Value &	thread_value,
		storage_type &	storage,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type using temporary storage.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_key	- input/output key to pass to other threads
thread_value	- input/output value to pass to other threads
storage	- temporary storage for inputs
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

Storage reusage: Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.: __global__ void example_kernel(...)
{
int value = ...;
using warp_sort_int = rp::warp_sort<int, 64>;
warp_sort_int wsort;
__shared__ typename warp_sort_int::storage_type storage;
wsort.sort(key, value, storage);
...
}

◆ sort() [14/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key &	,
		Value &	,
		storage_type &	,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type using temporary storage.

Invalid Warp Size

◆ sort() [15/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		Value(&)	thread_values[ItemsPerThread],
		storage_type &	storage,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize <= __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type using temporary storage.

Template Parameters

BinaryFunction - type of binary function used for sort. Default type is rocprim::less<T>.

Parameters

thread_keys	- input/output keys to pass to other threads
thread_values	- input/output values to pass to other threads
storage	- temporary storage for inputs
compare_function	- binary operation function object that will be used for sort. The signature of the function should be equivalent to the following: `bool f(const T &a, const T &b);`. The signature does not need to have `const &`, but function object must not modify the objects passed to it.

Storage reusage: Synchronization barrier should be placed before storage is reused or repurposed: __syncthreads() or rocprim::syncthreads().

Example.: __global__ void example_kernel(...)
{
int value = ...;
using warp_sort_int = rp::warp_sort<int, 64>;
warp_sort_int wsort;
__shared__ typename warp_sort_int::storage_type storage;
wsort.sort(key, value, storage);
...
}

◆ sort() [16/16]

template<class Key , unsigned int WarpSize = device_warp_size(), class Value = empty_type>

template<unsigned int ItemsPerThread, class BinaryFunction = ::rocprim::less<Key>, unsigned int FunctionWarpSize = WarpSize>

ROCPRIM_DEVICE ROCPRIM_INLINE auto warp_sort< Key, WarpSize, Value >::sort	(	Key(&)	thread_keys[ItemsPerThread],
		Value(&)	thread_values[ItemsPerThread],
		storage_type &	,
		BinaryFunction	compare_function = `BinaryFunction()`
	)		-> typename std::enable_if<(FunctionWarpSize > __AMDGCN_WAVEFRONT_SIZE), void>::type

inline

Warp sort by key for any data type using temporary storage.

Invalid Warp Size

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

rocprim/include/rocprim/warp/warp_sort.hpp

Public Types

Public Member Functions

Detailed Description

template<class Key, unsigned int WarpSize = device_warp_size(), class Value = empty_type> class warp_sort< Key, WarpSize, Value >

Member Typedef Documentation

◆ storage_type

Member Function Documentation

◆ sort() [1/16]

◆ sort() [2/16]

◆ sort() [3/16]

◆ sort() [4/16]

◆ sort() [5/16]

◆ sort() [6/16]

◆ sort() [7/16]

◆ sort() [8/16]

◆ sort() [9/16]

◆ sort() [10/16]

◆ sort() [11/16]

◆ sort() [12/16]

◆ sort() [13/16]

◆ sort() [14/16]

◆ sort() [15/16]

◆ sort() [16/16]

template<class Key, unsigned int WarpSize = device_warp_size(), class Value = empty_type>
class warp_sort< Key, WarpSize, Value >