template<class Config = default_config, class InputIterator, class OutputIterator, class UniqueCountOutputIterator, class EqualityOp = ::rocprim::equal_to<typename std::iterator_traits<InputIterator>::value_type>>
inline hipError_t rocprim::unique(void *temporary_storage, size_t &storage_size, InputIterator input, OutputIterator output, UniqueCountOutputIterator unique_count_output, const size_t size, EqualityOp equality_op = EqualityOp(), const hipStream_t stream = 0, const bool debug_synchronous = false)#

Device-level parallel unique primitive.

From given input range unique primitive eliminates all but the first element from every consecutive group of equivalent elements and copies them into output.


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

  • Range specified by input must have at least size elements.

  • Range specified by output must have at least so many elements, that all selected values can be copied into it.

  • Range specified by unique_count_output must have at least 1 element.

  • By default InputIterator::value_type’s equality operator is used to check if elements are equivalent.


In this example a device-level unique operation is performed on an array of integer values.

#include <rocprim/rocprim.hpp>

// Prepare input and output (declare pointers, allocate device memory etc.)
size_t input_size;     // e.g., 8
int * input;           // e.g., [1, 4, 2, 4, 4, 7, 7, 7]
int * output;          // empty array of 8 elements
size_t * output_count; // empty array of 1 element

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, output_count,

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

// perform unique operation
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, output, output_count,
// output: [1, 4, 2, 4, 7]
// output_count: 5

Template Parameters:
  • InputIterator – - random-access iterator type of the input range. It can be a simple pointer type.

  • OutputIterator – - random-access iterator type of the output range. It can be a simple pointer type.

  • UniqueCountOutputIterator – - random-access iterator type of the unique_count_output value used to return number of unique values. It can be a simple pointer type.

  • EqualityOp – - type of an binary operator used to compare values for equality.

  • 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 unique operation.

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

  • input[in] - iterator to the first element in the range to select values from.

  • output[out] - iterator to the first element in the output range.

  • unique_count_output[out] - iterator to the total number of selected values (length of output).

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

  • equality_op[in] - [optional] binary function object used to compare input values for equality. The signature of the function should be equivalent to the following: bool equal_to(const T &a, const T &b);. The signature does not need to have const &, but function object must not modify the object passed to it.

  • stream[in] - [optional] HIP stream object. The default is 0 (default stream).

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



doxygenfunction: Cannot find function “rocprim::unique_by_key” in doxygen xml output for project “rocPRIM Documentation” from directory: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprim/checkouts/docs-5.1.3/docs/.doxygen/docBin/xml