Find first of

Find first of#

Configuring the kernel#

template<unsigned int BlockSize, unsigned int ItemsPerThread>
struct find_first_of_config : public rocprim::detail::find_first_of_config_params#

Configuration of device-level find_first_of.

Template Parameters:

  • BlockSize – number of threads in a block.

  • ItemsPerThread – number of items processed by each thread.

Subclassed by rocprim::detail::default_find_first_of_config< arch, value_type, enable >

find_first_of#

template<class Config = default_config, class InputIterator1, class InputIterator2, class OutputIterator, class BinaryFunction = ::rocprim::equal_to<typename std::iterator_traits<InputIterator1>::value_type>>
inline hipError_t rocprim::find_first_of(void *temporary_storage, size_t &storage_size, InputIterator1 input, InputIterator2 keys, OutputIterator output, size_t size, size_t keys_size, BinaryFunction compare_function = BinaryFunction(), hipStream_t stream = 0, bool debug_synchronous = false)#

Searches the range [input, input + size) for any of the elements in the range [keys, keys + keys_size).

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

  • Accepts custom compare_function.

Example

In this example a device-level find_first_of is performed where inputs and 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 size;                // e.g., 8
size_t keys_size;           // e.g., 2
unsigned int* input;        // e.g., [ 6, 3, 5, 4, 1, 8, 2, 7 ]
unsigned int* keys;         // e.g., [ 10, 5 ]
unsigned int* keys_output;  // 1 element

size_t temporary_storage_size_bytes;
void * temporary_storage_ptr = nullptr;
// Get required size of the temporary storage
rocprim::find_first_of(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, keys, output, size, keys_size
);

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

// perform find_first_of
rocprim::find_first_of(
    temporary_storage_ptr, temporary_storage_size_bytes,
    input, keys, output, size, keys_size
);
// output: [ 2 ]

Template Parameters:

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

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

  • InputIterator2 – [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 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 InputIterator1 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 range of elements to examine.

  • keys – [in] iterator to the range of elements to search for.

  • output – [out] iterator to the output range. output should be able to be written for 1 element. *output constains the position of the first element in the range [input, input + size) that is equal to an element from the range [keys, keys + keys_size).

  • size – [in] number of elements to examine.

  • keys_size – [in] number of elements to search for.

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

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

Contents