Merge#
Configuring the kernel#
-
template<unsigned int BlockSize, unsigned int ItemsPerThread>
using rocprim::merge_config = kernel_config<BlockSize, ItemsPerThread># Configuration of device-level merge primitives.
merge#
-
template<class Config = default_config, class InputIterator1, class InputIterator2, class OutputIterator, class BinaryFunction = ::rocprim::less<typename std::iterator_traits<InputIterator1>::value_type>>
inline hipError_t rocprim::merge(void *temporary_storage, size_t &storage_size, InputIterator1 input1, InputIterator2 input2, OutputIterator output, const size_t input1_size, const size_t input2_size, BinaryFunction compare_function = BinaryFunction(), const hipStream_t stream = 0, bool debug_synchronous = false)# Parallel merge primitive for device level.
merge
function performs a device-wide merge. Function merges two ordered sets of input values based on comparison function.- Overview
The contents of the inputs are not altered by the merging function.
Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.Accepts custom compare_functions for merging across the device.
- Example
In this example a device-level ascending merge is performed on an array of
int
values.#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) size_t input_size1; // e.g., 4 size_t input_size2; // e.g., 4 int * input1; // e.g., [0, 1, 2, 3] int * input2; // e.g., [0, 1, 2, 3] int * output; // empty array of 8 elements size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::merge( temporary_storage_ptr, temporary_storage_size_bytes, input1, input2, output, input_size1, input_size2 ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // perform merge rocprim::merge( temporary_storage_ptr, temporary_storage_size_bytes, input1, input2, output, input_size1, input_size2 ); // output: [0, 0, 1, 1, 2, 2, 3, 3]
- Template Parameters:
Config – - [optional] configuration of the primitive. It has to be
merge_config
or a class derived from it.InputIterator1 – - random-access iterator type of the first input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
InputIterator2 – - random-access iterator type of the second input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
OutputIterator – - random-access iterator type of the output range. Must meet the requirements of a C++ OutputIterator concept. It can be a simple pointer 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 function returns without performing the sort operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.input1 – [in] - iterator to the first element in the first range to merge.
input2 – [in] - iterator to the first element in the second range to merge.
output – [out] - iterator to the first element in the output range.
input1_size – [in] - number of element in the first input range.
input2_size – [in] - number of element in the second input range.
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 haveconst &
, but function object must not modify the objects passed to it. The default value isBinaryFunction()
.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 sort; otherwise a HIP runtime error of typehipError_t
.
-
template<class Config = default_config, class KeysInputIterator1, class KeysInputIterator2, class KeysOutputIterator, class ValuesInputIterator1, class ValuesInputIterator2, class ValuesOutputIterator, class BinaryFunction = ::rocprim::less<typename std::iterator_traits<KeysInputIterator1>::value_type>>
inline hipError_t rocprim::merge(void *temporary_storage, size_t &storage_size, KeysInputIterator1 keys_input1, KeysInputIterator2 keys_input2, KeysOutputIterator keys_output, ValuesInputIterator1 values_input1, ValuesInputIterator2 values_input2, ValuesOutputIterator values_output, const size_t input1_size, const size_t input2_size, BinaryFunction compare_function = BinaryFunction(), const hipStream_t stream = 0, bool debug_synchronous = false)# Parallel merge primitive for device level.
merge
function performs a device-wide merge of (key, value) pairs. Function merges two ordered sets of input keys and corresponding values based on key comparison function.- Overview
The contents of the inputs are not altered by the merging function.
Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.Accepts custom compare_functions for merging across the device.
- Example
In this example a device-level ascending merge is performed on an array of
int
values.#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) size_t input_size1; // e.g., 4 size_t input_size2; // e.g., 4 int * keys_input1; // e.g., [0, 1, 2, 3] int * keys_input2; // e.g., [0, 1, 2, 3] int * keys_output; // empty array of 8 elements int * values_input1; // e.g., [10, 11, 12, 13] int * values_input2; // e.g., [20, 21, 22, 23] int * values_output; // empty array of 8 elements size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::merge( temporary_storage_ptr, temporary_storage_size_bytes, keys_input1, keys_input2, keys_output, values_input1, values_input2, values_output, // input_size1, input_size2 ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // perform merge rocprim::merge( temporary_storage_ptr, temporary_storage_size_bytes, keys_input1, keys_input2, keys_output, values_input1, values_input2, values_output, // input_size1, input_size2 ); // keys_output: [0, 0, 1, 1, 2, 2, 3, 3] // values_output: [10, 20, 11, 21, 12, 22, 13, 23]
- Template Parameters:
Config – - [optional] configuration of the primitive. It has to be
merge_config
or a class derived from it.KeysInputIterator1 – - random-access iterator type of the first keys input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
KeysInputIterator2 – - random-access iterator type of the second keys input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
KeysOutputIterator – - random-access iterator type of the keys output range. Must meet the requirements of a C++ OutputIterator concept. It can be a simple pointer type.
ValuesInputIterator1 – - random-access iterator type of the first values input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
ValuesInputIterator2 – - random-access iterator type of the second values input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
ValuesOutputIterator – - random-access iterator type of the values output range. Must meet the requirements of a C++ OutputIterator concept. It can be a simple pointer 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 function returns without performing the sort operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.keys_input1 – [in] - iterator to the first key in the first range to merge.
keys_input2 – [in] - iterator to the first key in the second range to merge.
keys_output – [out] - iterator to the first key in the output range.
values_input1 – [in] - iterator to the first value in the first range to merge.
values_input2 – [in] - iterator to the first value in the second range to merge.
values_output – [out] - iterator to the first value in the output range.
input1_size – [in] - number of element in the first input range.
input2_size – [in] - number of element in the second input range.
compare_function – [in] - binary operation function object that will be used for key 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 haveconst &
, but function object must not modify the objects passed to it. The default value isBinaryFunction()
.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 sort; otherwise a HIP runtime error of typehipError_t
.