TopK

TopK#

Configuring the kernel#

Configuration for device level topk algorithm is not supported yet. This feature will be added in the future.

topk#

template<class Config = default_config, bool Descending = false, bool Ordered = false, bool Deterministic = false, bool Stable = false, class Decomposer = ::rocprim::identity_decomposer, class KeysInputIterator, class KeysOutputIterator, class SizeIn, class SizeOut>
inline hipError_t rocprim::topk(void *temporary_storage, size_t &storage_size, const KeysInputIterator keys_input, const KeysOutputIterator keys_output, const SizeIn size, const SizeOut K, Decomposer decomposer = {}, const hipStream_t stream = 0, const bool debug_synchronous = false)#

Find the largest or smallest K elements from an input array of keys.

Returns the K largest or smallest elements. These maybe be in any order.

A similar algorithm is nth_element. 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.

Example

In this example a device-level ascending top-k is performed on an array.

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 k;               // e.g., 3
int * input;            // e.g., [2, 3, 4, -1, -2, -3, 0, 5]
int * output;           // empty array of 3 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::topk(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size, k
);
// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);
// perform topk
rocprim::topk(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, input_size, k
);
// output should be [5, 4, 3] (order is not guaranteed)

Note

Ordered output is not supported yet. This feature will be added in the future.

Note

Deterministic output is not supported yet. This feature will be added in the future.

Note

Stable output is not supported yet. This feature will be added in the future.

Template Parameters:

  • Config – [optional] configuration of the primitive, must be default_config or topk_config.

  • Descending – [optional] determines the starting direction. If true, select the largest K elements, and vice versa.

  • Ordered – [optional] determines whether the output results are sorted by size.

  • Deterministic – [optional] to ensure that the results are exactly the same every time.

  • Stable – [optional] determines whether elements in the output are arranged according to their relative position in the input.

  • Decomposer – [optional] the type of the decomposer functor.

  • KeysInputIterator – [optional] 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 – [optional] random-access iterator type of the output range. Must meet the requirements of a C++ OutputIterator concept. It can be a simple pointer type.

  • SizeIn – [optional] integral type that represents the problem size.

  • SizeOut – [optional] integral type that counts the number of output elements.

Parameters:

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

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

  • keys_input – [in] pointer to the first element in the input range.

  • keys_output – [out] pointer to the first element in the output range.

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

  • K – [in] number of elements to be selected from the input.

  • decomposer – [in] decomposer functor that produces a tuple of references from the input key type.

  • 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 a successful top-k operation; otherwise a HIP runtime error of type hipError_t.

topk_pairs#

template<class Config = default_config, bool Descending = false, bool Ordered = false, bool Deterministic = false, bool Stable = false, class Decomposer = rocprim::identity_decomposer, class KeysInputIterator, class KeysOutputIterator, class ValuesInputIterator, class ValuesOutputIterator, class SizeIn, class SizeOut>
inline hipError_t rocprim::topk_pairs(void *temporary_storage, size_t &storage_size, const KeysInputIterator keys_input, const KeysOutputIterator keys_output, const ValuesInputIterator values_input, const ValuesOutputIterator values_output, const SizeIn size, const SizeOut K, const Decomposer decomposer = {}, const hipStream_t stream = 0, const bool debug_synchronous = false)#

Find the largest or smallest K elements from an input array of values based on their corresponding keys.

Returns the K largest or smallest (key, value) pairs. These maybe be in any order.

A similar algorithm is nth_element. 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.

Example

In this example a device-level ascending top-k is performed on an array.

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 k;               // e.g., 3
int * input_keys;       // e.g., [2, 3, 4, -1, -2, -3, 0, 5]
int * output_keys;      // empty array of 3 elements
int * input_vals;       // e.g., [0, 1, 2, 3, 4, 5, 6, 7]
int * output_vals;      // empty array of 3 elements

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::topk_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input_keys, output_keys, input_vals, output_vals, input_size, k
);
// allocate temporary storage
hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);
// perform topk_pairs
rocprim::topk_pairs(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input_keys, output_keys, input_vals, output_vals, input_size, k
);
// output_keys should be [5, 4, 3] (order is not guaranteed)
// output_vals should be [7, 2, 1] (order matches output_keys)

Note

Ordered output is not supported yet. This feature will be added in the future.

Note

Deterministic output is not supported yet. This feature will be added in the future.

Note

Stable output is not supported yet. This feature will be added in the future.

Template Parameters:

  • Config – [optional] configuration of the primitive, must be default_config or topk_config.

  • Descending – [optional] determines the starting direction. If true, select the largest K elements, and vice versa.

  • Ordered – [optional] determines whether the output results are sorted by size.

  • Deterministic – [optional] to ensure that the results are exactly the same every time.

  • Stable – [optional] determines whether elements in the output are arranged according to their relative position in the input.

  • Decomposer – [optional] the type of the decomposer functor.

  • KeysInputIterator – [optional] 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 – [optional] random-access iterator type of the output range. Must meet the requirements of a C++ OutputIterator concept. It can be a simple pointer type.

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

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

  • SizeIn – [optional] integral type that represents the problem size.

  • SizeOut – [optional] integral type that counts the number of output elements.

Parameters:

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

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

  • keys_input – [in] pointer to the first key in the input range.

  • keys_output – [out] pointer to the first key in the output range.

  • values_input – [in] pointer to the first value in the output range.

  • values_output – [out] pointer to the first value in the output range.

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

  • K – [in] number of elements to be selected from the input.

  • decomposer – [in] decomposer functor that produces a tuple of references from the input key type.

  • 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 a successful top-k operation; otherwise a HIP runtime error of type hipError_t.

Contents