Nth Element

Nth Element#

Configuring the kernel#

template<unsigned int BlockSize, unsigned int ItemsPerThread, unsigned int StopRecursionSize, unsigned int NumberOfBuckets, block_radix_rank_algorithm RadixRankAlgorithm>
struct nth_element_config : public rocprim::detail::nth_element_config_params#

Configuration of device-level nth_element.

Template Parameters:

  • BlockSize – number of threads in a block.

  • ItemsPerThread – number of items processed by each thread.

  • StopRecursionSize – the size from where recursion is stopped to do a block sort

  • NumberOfBuckets – the number of buckets that are used in the algorithm

  • RadixRankAlgorithm – algorithm for radix rank

nth_element#

template<class Config = default_config, class KeysIterator, class BinaryFunction = ::rocprim::less<typename std::iterator_traits<KeysIterator>::value_type>>
inline hipError_t rocprim::nth_element(void *temporary_storage, size_t &storage_size, KeysIterator keys, size_t nth, size_t size, BinaryFunction compare_function = BinaryFunction(), hipStream_t stream = 0, bool debug_synchronous = false)#

Rearrange elements smaller than the n-th before and larger than n-th after the n-th element.

After nth_element the element in the position pointed to by nth is the element that would be in that position if the whole range were sorted. Additionally the other elements are rearranged such that for all values of i in [keys_output, keys_output + n) and all values of j in [keys_output + n, keys_output + size): comp(*i, *j) is false. Smaller elements than the n-th will be arranged before, and larger ones after the n-th element.

The topk algorithm is similar and selects the largest or smallest k elements from the input array. The main differences are:

  • topk returns arrays of size k, whereas nth_element returns an array matching the input size.

  • The element at index n is set to the element that would be at the n-th position in nth_element, while topk does not guarantee that.

  • Elements which are smaller than the n-th element will be placed in the front of the output, while topk ignores all elements which are smaller or larger than the k-th element (depending on descending or ascending).

  • In-place operation can be done by nth_element, while topk does not support that.

  • topk supports hipGraph, but nth_element does not.

Here are some tips for choosing between nth_element and topk:

  • If you need the n-th largest or smallest element to be placed at index n, please use nth_element.

  • If you want to care about both smaller elements and larger elements, please use nth_element.

  • If you want to use hipGraph, please use topk.

Overview

  • The contents of the inputs are not altered by the function.

  • Returns the required size of temporary_storage in storage_size if temporary_storage is a null pointer.

  • Accepts custom compare_functions for nth_element across the device.

  • Streams in graph capture mode are not supported

Stability

nth_element is not stable: it doesn’t necessarily preserve the relative ordering of equivalent keys. That is, given two keys a and b and a binary boolean operation op such that:

  • a precedes b in the input keys, and

  • op(a, b) and op(b, a) are both false, then it is not guaranteed that a will precede b as well in the output keys.

Example

In this example a device-level nth_element is performed where input keys are represented by an array of unsigned integers.

The full example is on GitHub.

#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
size_t nth;                 // e.g., 4
unsigned int * keys;  // e.g., [ 6, 3, 5, 4, 1, 8, 2, 7 ]

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, nth, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform nth_element
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys, nth, input_size
);
// possible keys:   [ 1, 3, 4, 2, 5, 8, 7, 6 ]

Template Parameters:

  • Config – [optional] configuration of the primitive. It has to be nth_element_config.

  • KeysIterator – [inferred] random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.

  • CompareFunction – [inferred] Type of binary function that accepts two arguments of the type KeysIterator and returns a value convertible to bool. Default type is rocprim::less<>.

Parameters:

  • temporary_storage – [in] pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to storage_size and function returns without performing the nth_element rearrangement.

  • storage_size – [inout] reference to a size (in bytes) of temporary_storage.

  • keys – [inout] iterator to the input range.

  • nth – [in] The index of the nth_element in the input range.

  • size – [in] number of element in the input range.

  • compare_function – [in] binary operation function object that will be used for comparison. 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. The comparator must meet the C++ named requirement Compare. The default value is BinaryFunction().

  • stream – [in] [optional] HIP stream object. Default is 0 (default stream).

  • debug_synchronous – [in] [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is false.

Returns:

hipSuccess (0) after successful rearrangement; otherwise a HIP runtime error of type hipError_t.

template<class Config = default_config, class KeysInputIterator, class KeysOutputIterator, class BinaryFunction = ::rocprim::less<typename std::iterator_traits<KeysInputIterator>::value_type>>
inline hipError_t rocprim::nth_element(void *temporary_storage, size_t &storage_size, KeysInputIterator keys_input, KeysOutputIterator keys_output, size_t nth, size_t size, BinaryFunction compare_function = BinaryFunction(), hipStream_t stream = 0, bool debug_synchronous = false)#

Rearrange elements smaller than the n-th before and larger than n-th after the n-th element.

After nth_element the element in the position pointed to by nth is the element that would be in that position if the whole range were sorted. Additionally the other elements are rearranged such that for all values of i in [keys_output, keys_output + n) and all values of j in [keys_output + n, keys_output + size): comp(*i, *j) is false. Smaller elements than the n-th will be arranged before, and larger ones after the n-th element.

The topk algorithm is similar and selects the largest or smallest k elements from the input array. The main differences are:

  • topk returns arrays of size k, whereas nth_element returns an array matching the input size.

  • The element at index n is set to the element that would be at the n-th position in nth_element, while topk does not guarantee that.

  • Elements which are smaller than the n-th element will be placed in the front of the output, while topk ignores all elements which are smaller or larger than the k-th element (depending on descending or ascending).

  • In-place operation can be done by nth_element, while topk does not support that.

  • topk supports hipGraph, but nth_element does not.

Here are some tips for choosing between nth_element and topk:

  • If you need the n-th largest or smallest element to be placed at index n, please use nth_element.

  • If you want to care about both smaller elements and larger elements, please use nth_element.

  • If you want to use hipGraph, please use topk.

Overview

  • The contents of the inputs are not altered by the function.

  • Returns the required size of temporary_storage in storage_size if temporary_storage is a null pointer.

  • Accepts custom compare_functions for nth_element across the device.

  • Streams in graph capture mode are not supported

Stability

nth_element is not stable: it doesn’t necessarily preserve the relative ordering of equivalent keys. That is, given two keys a and b and a binary boolean operation op such that:

  • a precedes b in the input keys, and

  • op(a, b) and op(b, a) are both false, then it is not guaranteed that a will precede b as well in the output keys.

Example

In this example a device-level nth_element is performed where input keys are represented by an array of unsigned integers.

#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;          // e.g., 8
size_t nth;                 // e.g., 4
unsigned int * keys_input;  // e.g., [ 6, 3, 5, 4, 1, 8, 2, 7 ]
unsigned int * keys_output; // empty array of 8 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, nth, input_size
);

// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);

// perform nth_element
rocprim::nth_element(
    temporary_storage_ptr, temporary_storage_size_bytes,
    keys_input, keys_output, nth, input_size
);
// possible keys_output:   [ 1, 3, 4, 2, 5, 8, 7, 6 ]

Template Parameters:

  • Config – [optional] configuration of the primitive. It has to be nth_element_config.

  • KeysInputIterator – [inferred] random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.

  • KeysOutputIterator – [inferred] random-access iterator type of the output range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.

  • CompareFunction – [inferred] Type of binary function that accepts two arguments of the type KeysIterator and returns a value convertible to bool. Default type is rocprim::less<>.

Parameters:

  • temporary_storage – [in] pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to storage_size and function returns without performing the nth_element rearrangement.

  • storage_size – [inout] reference to a size (in bytes) of temporary_storage.

  • keys_input – [in] iterator to the input range.

  • keys_output – [out] iterator to the output range. No overlap at all is allowed between keys_input and keys_output. keys_output should be able to be written and read from for size elements.

  • nth – [in] The index of the nth_element in the input range.

  • size – [in] number of element in the input range.

  • compare_function – [in] binary operation function object that will be used for comparison. 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. The comparator must meet the C++ named requirement Compare. The default value is BinaryFunction().

  • stream – [in] [optional] HIP stream object. Default is 0 (default stream).

  • debug_synchronous – [in] [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is false.

Returns:

hipSuccess (0) after successful rearrangement; otherwise a HIP runtime error of type hipError_t.

Contents