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 in storage_size if temporary_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 or adjacent_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 have const &, but the function object must not modify the object passed to it. The operator must meet the C++ named requirement BinaryPredicate. The default operation used is rocprim::equal_to<T>, where T is the type of the elements in the input range obtained with std::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 type hipError_t.