Adjacent Find#
Configuring the kernel#
-
template<unsigned int BlockSize, unsigned int ItemsPerThread>
struct adjacent_find_config : public rocprim::detail::adjacent_find_config_params# Configuration of device-level adjacent_find.
- Template Parameters:
BlockSize – number of threads in a block.
ItemsPerThread – number of items processed by each thread.
Subclassed by rocprim::detail::default_adjacent_find_config< arch, input_type, enable >
adjacent_find#
-
template<typename Config = default_config, typename InputIterator, typename OutputIterator, typename BinaryPred = ::rocprim::equal_to<typename std::iterator_traits<InputIterator>::value_type>>
inline hipError_t rocprim::adjacent_find(void *const temporary_storage, std::size_t &storage_size, InputIterator input, OutputIterator output, const std::size_t size, BinaryPred op = BinaryPred{}, const hipStream_t stream = 0, const bool debug_synchronous = false)# Searches the input sequence for the first appearance of a consecutive pair of equal elements.
The returned index is either: the index within the input array of the first element of the first pair of consecutive equal elements found or the size of the input array if no such pair is found. Equivalent to the following code
if(size > 1) { for(std::size_t i = 0; i < size - 1 ; ++i) if (op(input[i], input[i + 1])) return i; } return size;
- Overview
The contents of the inputs are not altered by the function.
Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
is a null pointer.Accepts custom
op
.Streams in graph capture mode are supported.
- Example
In this example a device-level adjacent_find operation is performed on integer values.
#include <rocprim/rocprim.hpp> //or <rocprim/device/device_adjacent_find.hpp> // Custom boolean binary function auto equal_op = [](int a, int b) -> bool { return (a - b == 2); }; // Prepare input and output (declare pointers, allocate device memory etc.) std::size_t size; // e.g., 8 int* input; // e.g., [8, 7, 5, 4, 3, 2, 1, 0] std::size_t* output; // output index auto custom_op = equal_op{}; std::size_t temporary_storage_size_bytes; void* temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::adjacent_find( temporary_storage_ptr, temporary_storage_size_bytes, input, output, size, custom_op); // Allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // Perform adjacent find rocprim::adjacent_find( temporary_storage_ptr, temporary_storage_size_bytes, input, output, size, custom_op); // output: 1
- Template Parameters:
Config – [optional] Configuration of the primitive, must be
default_config
oradjacent_find_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 output index. Must meet the requirements of a C++ OutputIterator concept. It can be a simple pointer type.
BinaryPred – [inferred] Boolean binary operation function object that will be applied to consecutive items to check whether they are equal or not. 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 haveconst &
, but the function object must not modify the object passed to it. The operator must meet the C++ named requirementBinaryPredicate
. The default operation used isrocprim::equal_to<T>
, whereT
is the type of the elements in the input range obtained withstd::iterator_traits<InputIterator>::value_type
>.
- 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 the function returns without performing any device computation.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 index.
size – [in] Number of items in the input.
op – [in] [optional] The boolean binary operation to be used by the algorithm. Default is
rocprim::equal_to
specialized for the type of the input elements.stream – [in] [optional] HIP stream object. Default is
0
(the default stream).debug_synchronous – [in] [optional] If true, synchronization after every kernel launch is forced in order to check for errors and extra debugging info is printed to the standard output. Default value is
false
.
- Returns:
hipSuccess
(0) after a successful search, otherwise the HIP runtime error of typehipError_t
.