Search N#

Configuring the kernel#

template<unsigned int BlockSize, unsigned int ItemsPerThread>
struct search_n_config : public rocprim::detail::search_n_config_params#

Configuration of device-level search_n.

Template Parameters:
  • BlockSize – number of threads in a block.

  • ItemsPerThread – number of items processed by each thread.

search_n#

template<class Config = default_config, class InputIterator, class OutputIterator, class BinaryPredicate = rocprim::equal_to<typename std::iterator_traits<InputIterator>::value_type>>
inline hipError_t rocprim::search_n(void *temporary_storage, size_t &storage_size, InputIterator input, OutputIterator output, const size_t size, const size_t count, const typename std::iterator_traits<InputIterator>::value_type *value, const BinaryPredicate binary_predicate = BinaryPredicate(), const hipStream_t stream = static_cast<hipStream_t>(0), const bool debug_synchronous = false)#

Searches for the first occurrence of a sequence of count elements all equal to value.

The equality of the elements of the sequence and the given value is determined according to a given comparison function. If found, the index of the first item of the found sequence in the input is returned. Otherwise, returns the size of the input.

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 search across the device.

  • Streams in graph capture mode are supported

Template Parameters:
  • Config – [optional] configuration of the primitive. It must be default_config or search_n_config.

  • InputIterator – [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.

  • OutputIterator – [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.

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

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

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

  • input[in] iterator to the input range.

  • output[out] iterator to the output range. The output is one element.

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

  • count[in] number of elements in the sequence. Must be less or equal than size, otherwise hipErrorInvalidValue will be returned.

  • value[in] value of the elements to search for.

  • binary_predicate[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++ requirements of BinaryPredicate. The default value is BinaryPredicate().

  • 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 search; otherwise a HIP runtime error of type hipError_t.