diff --git a/rocprim/include/rocprim/device/device_adjacent_difference.hpp b/rocprim/include/rocprim/device/device_adjacent_difference.hpp
index f20ccc6cb..917bf4328 100644
--- a/rocprim/include/rocprim/device/device_adjacent_difference.hpp
+++ b/rocprim/include/rocprim/device/device_adjacent_difference.hpp
@@ -245,7 +245,7 @@ hipError_t adjacent_difference_impl(void* const          temporary_storage,
 /// \endcode
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// `adjacent_difference_config`.
+/// `adjacent_difference_config` or a class derived from it.
 /// \tparam InputIt - [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.
 /// \tparam OutputIt - [inferred] random-access iterator type of the output range. Must meet the
@@ -341,7 +341,7 @@ hipError_t adjacent_difference(void* const          temporary_storage,
 /// \endcode
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// `adjacent_difference_config`.
+/// `adjacent_difference_config` or a class derived from it.
 /// \tparam InputIt - [inferred] random-access iterator type of the value range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam BinaryFunction - [inferred] binary operation function object that will be applied to
@@ -394,7 +394,7 @@ hipError_t adjacent_difference_inplace(void* const          temporary_storage,
 /// \endcode
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// `adjacent_difference_config`.
+/// `adjacent_difference_config` or a class derived from it.
 /// \tparam InputIt - [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.
 /// \tparam OutputIt - [inferred] random-access iterator type of the output range. Must meet the
@@ -490,7 +490,7 @@ hipError_t adjacent_difference_right(void* const          temporary_storage,
 /// \endcode
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// `adjacent_difference_config`.
+/// `adjacent_difference_config` or a class derived from it.
 /// \tparam InputIt - [inferred] random-access iterator type of the value range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam BinaryFunction - [inferred] binary operation function object that will be applied to
diff --git a/rocprim/include/rocprim/device/device_binary_search.hpp b/rocprim/include/rocprim/device/device_binary_search.hpp
index 38cac71fe..53914438d 100644
--- a/rocprim/include/rocprim/device/device_binary_search.hpp
+++ b/rocprim/include/rocprim/device/device_binary_search.hpp
@@ -114,8 +114,8 @@ struct is_default_or_has_tag<default_config, Tag>
 /// in the semi-open interval <tt>[lower_output[i], upper_output[i])</tt> of \p haystack, in case of
 /// being present at all.
 ///
-/// @tparam Config - [optional] Configuration of the primitive. It can be \p lower_bound_config or
-/// a class with the same members. Default is \p default_config.
+/// @tparam Config - [optional] Configuration of the primitive. It has to be \p lower_bound_config or
+/// a class derived from it. Default is \p default_config.
 /// @tparam HaystackIterator - [inferred] Random-access iterator type of the search range. Must meet
 /// the requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// @tparam NeedlesIterator - [inferred] Random-access iterator type of the input range. Must meet
@@ -251,7 +251,7 @@ hipError_t lower_bound(void * temporary_storage,
 /// being present at all.
 ///
 /// @tparam Config - [optional] Configuration of the primitive. It can be \p upper_bound_config or
-/// a class with the same members. Default is \p default_config.
+/// a class derived from it. Default is \p default_config.
 /// @tparam HaystackIterator - [inferred] Random-access iterator type of the search range. Must meet
 /// the requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// @tparam NeedlesIterator - [inferred] Random-access iterator type of the input range. Must meet
@@ -381,7 +381,7 @@ hipError_t upper_bound(void * temporary_storage,
 /// is written to \p storage_size and the function returns without performing the search operation.
 ///
 /// @tparam Config - [optional] Configuration of the primitive. It can be \p binary_search_config or
-/// a class with the same members. Default is \p default_config.
+/// a class derived from it. Default is \p default_config.
 /// @tparam HaystackIterator - [inferred] Random-access iterator type of the search range. Must meet
 /// the requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// @tparam NeedlesIterator - [inferred] Random-access iterator type of the input range. Must meet
diff --git a/rocprim/include/rocprim/device/device_histogram.hpp b/rocprim/include/rocprim/device/device_histogram.hpp
index 3dd56abb3..0e8ff3700 100644
--- a/rocprim/include/rocprim/device/device_histogram.hpp
+++ b/rocprim/include/rocprim/device/device_histogram.hpp
@@ -449,7 +449,7 @@ inline hipError_t histogram_range_impl(void*          temporary_storage,
 /// * Returns the required size of \p temporary_storage in \p storage_size
 /// if \p temporary_storage in a null pointer.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -552,7 +552,7 @@ inline hipError_t histogram_even(void*          temporary_storage,
 /// * Returns the required size of \p temporary_storage in \p storage_size
 /// if \p temporary_storage in a null pointer.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -663,7 +663,7 @@ inline hipError_t histogram_even(void*          temporary_storage,
 ///
 /// \tparam Channels - number of channels interleaved in the input samples.
 /// \tparam ActiveChannels - number of channels being used for computing histograms.
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -775,7 +775,7 @@ inline hipError_t multi_histogram_even(void*          temporary_storage,
 ///
 /// \tparam Channels - number of channels interleaved in the input samples.
 /// \tparam ActiveChannels - number of channels being used for computing histograms.
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -882,7 +882,7 @@ inline hipError_t multi_histogram_even(void*          temporary_storage,
 /// * Returns the required size of \p temporary_storage in \p storage_size
 /// if \p temporary_storage in a null pointer.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -979,7 +979,7 @@ inline hipError_t histogram_range(void*          temporary_storage,
 /// * Returns the required size of \p temporary_storage in \p storage_size
 /// if \p temporary_storage in a null pointer.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -1085,7 +1085,7 @@ inline hipError_t histogram_range(void*          temporary_storage,
 ///
 /// \tparam Channels - number of channels interleaved in the input samples.
 /// \tparam ActiveChannels - number of channels being used for computing histograms.
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
@@ -1192,7 +1192,7 @@ inline hipError_t multi_histogram_range(void*          temporary_storage,
 ///
 /// \tparam Channels - number of channels interleaved in the input samples.
 /// \tparam ActiveChannels - number of channels being used for computing histograms.
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p histogram_config or a class derived from it.
 /// \tparam SampleIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam Counter - integer type for histogram bin counters.
diff --git a/rocprim/include/rocprim/device/device_merge.hpp b/rocprim/include/rocprim/device/device_merge.hpp
index bba58a845..30ef30e31 100644
--- a/rocprim/include/rocprim/device/device_merge.hpp
+++ b/rocprim/include/rocprim/device/device_merge.hpp
@@ -216,7 +216,7 @@ hipError_t merge_impl(void * temporary_storage,
 /// if \p temporary_storage in a null pointer.
 /// * Accepts custom compare_functions for merging across the device.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p merge_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p merge_config or a class derived from it.
 /// \tparam 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.
 /// \tparam InputIterator2 - random-access iterator type of the second input range. Must meet the
@@ -320,7 +320,7 @@ hipError_t merge(void * temporary_storage,
 /// if \p temporary_storage in a null pointer.
 /// * Accepts custom compare_functions for merging across the device.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p merge_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p merge_config or a class derived from it.
 /// \tparam 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.
 /// \tparam KeysInputIterator2 - random-access iterator type of the second keys input range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_partition.hpp b/rocprim/include/rocprim/device/device_partition.hpp
index ae8c9f1e4..3b89a3ed1 100644
--- a/rocprim/include/rocprim/device/device_partition.hpp
+++ b/rocprim/include/rocprim/device/device_partition.hpp
@@ -382,7 +382,7 @@ hipError_t partition_impl(void * temporary_storage,
 /// * Range specified by \p selected_count_output must have at least 1 element.
 /// * Relative order is preserved.
 ///
-/// \tparam Config - [optional] configuration of the primitive. If has to be an instance of \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. If has to be an instance of \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be a simple
 /// pointer type.
 /// \tparam SelectedOutputIterator - random-access iterator type of the selected output range. It
@@ -528,7 +528,7 @@ inline hipError_t partition_two_way(void*                       temporary_storag
 /// * Values of \p flag range should be implicitly convertible to `bool` type.
 /// * The relative order of elements in both output ranges matches the input range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. If has to be an instance of \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. If has to be an instance of \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam FlagIterator - random-access iterator type of the flag range. It can be
@@ -662,7 +662,7 @@ inline hipError_t partition_two_way(void*                       temporary_storag
 /// * Relative order is preserved for the elements for which the corresponding values from \p flags
 /// are \p true. Other elements are copied in reverse order.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam FlagIterator - random-access iterator type of the flag range. It can be
@@ -783,7 +783,7 @@ hipError_t partition(void * temporary_storage,
 /// * Relative order is preserved for the elements for which the \p predicate returns \p true. Other
 /// elements are copied in reverse order.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. It can be
@@ -927,7 +927,7 @@ hipError_t partition(void * temporary_storage,
 /// minus the number of elements written to \p output_first_part minus the number of elements written
 /// to \p output_second_part.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam FirstOutputIterator - random-access iterator type of the first output range. It can be
diff --git a/rocprim/include/rocprim/device/device_radix_sort.hpp b/rocprim/include/rocprim/device/device_radix_sort.hpp
index 1ff1af2c0..c4676b264 100644
--- a/rocprim/include/rocprim/device/device_radix_sort.hpp
+++ b/rocprim/include/rocprim/device/device_radix_sort.hpp
@@ -732,7 +732,7 @@ inline hipError_t
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -838,7 +838,7 @@ hipError_t radix_sort_keys(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -945,7 +945,7 @@ hipError_t radix_sort_keys_desc(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -1070,7 +1070,7 @@ hipError_t radix_sort_pairs(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -1195,7 +1195,7 @@ hipError_t radix_sort_pairs_desc(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam Size - integral type that represents the problem size.
 ///
@@ -1307,7 +1307,7 @@ hipError_t radix_sort_keys(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam Size - integral type that represents the problem size.
 ///
@@ -1419,7 +1419,7 @@ hipError_t radix_sort_keys_desc(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam Value - value type.
 /// \tparam Size - integral type that represents the problem size.
@@ -1546,7 +1546,7 @@ hipError_t radix_sort_pairs(void * temporary_storage,
 /// can be improved by setting \p begin_bit and \p end_bit, for example if all keys are in range
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam Value - value type.
 /// \tparam Size - integral type that represents the problem size.
diff --git a/rocprim/include/rocprim/device/device_reduce.hpp b/rocprim/include/rocprim/device/device_reduce.hpp
index 167f6d93b..649e1c583 100644
--- a/rocprim/include/rocprim/device/device_reduce.hpp
+++ b/rocprim/include/rocprim/device/device_reduce.hpp
@@ -256,7 +256,7 @@ hipError_t reduce_impl(void * temporary_storage,
 /// * By default, the input type is used for accumulation. A custom type
 /// can be specified using <tt>rocprim::transform_iterator</tt>, see the example below.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p reduce_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p reduce_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
@@ -403,7 +403,7 @@ hipError_t reduce(void * temporary_storage,
 /// * By default, the input type is used for accumulation. A custom type
 /// can be specified using <tt>rocprim::transform_iterator</tt>, see the example below.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p reduce_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p reduce_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_reduce_by_key.hpp b/rocprim/include/rocprim/device/device_reduce_by_key.hpp
index d57c111bd..315f75668 100644
--- a/rocprim/include/rocprim/device/device_reduce_by_key.hpp
+++ b/rocprim/include/rocprim/device/device_reduce_by_key.hpp
@@ -397,7 +397,7 @@ hipError_t reduce_by_key_impl(void*                     temporary_storage,
 /// * Ranges specified by \p unique_output and \p aggregates_output must have at least
 /// <tt>*unique_count_output</tt> (i.e. the number of unique keys) elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be `reduce_by_key_config`.
+/// \tparam Config - [optional] configuration of the primitive. It has to be `reduce_by_key_config` or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam ValuesInputIterator - random-access iterator type of the input range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_run_length_encode.hpp b/rocprim/include/rocprim/device/device_run_length_encode.hpp
index 563180067..109561cb1 100644
--- a/rocprim/include/rocprim/device/device_run_length_encode.hpp
+++ b/rocprim/include/rocprim/device/device_run_length_encode.hpp
@@ -76,7 +76,7 @@ namespace detail
 /// * Ranges specified by \p unique_output and \p counts_output must have at least
 /// <tt>*runs_count_output</tt> (i.e. the number of runs) elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p run_length_encode_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p run_length_encode_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam UniqueOutputIterator - random-access iterator type of the output range. Must meet the
@@ -191,7 +191,7 @@ hipError_t run_length_encode(void * temporary_storage,
 /// * Ranges specified by \p offsets_output and \p counts_output must have at least
 /// <tt>*runs_count_output</tt> (i.e. the number of non-trivial runs) elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p run_length_encode_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p run_length_encode_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OffsetsOutputIterator - random-access iterator type of the output range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_scan.hpp b/rocprim/include/rocprim/device/device_scan.hpp
index 7a32b7999..2cb648dad 100644
--- a/rocprim/include/rocprim/device/device_scan.hpp
+++ b/rocprim/include/rocprim/device/device_scan.hpp
@@ -457,7 +457,7 @@ inline auto scan_impl(void*               temporary_storage,
 /// * By default, the input type is used for accumulation. A custom type
 /// can be specified using <tt>rocprim::transform_iterator</tt>, see the example below.
 ///
-/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_config.
+/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
@@ -586,7 +586,7 @@ inline hipError_t inclusive_scan(void*             temporary_storage,
 /// if \p temporary_storage in a null pointer.
 /// * Ranges specified by \p input and \p output must have at least \p size elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_config.
+/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_scan_by_key.hpp b/rocprim/include/rocprim/device/device_scan_by_key.hpp
index 6699bf10a..8d8c4fecb 100644
--- a/rocprim/include/rocprim/device/device_scan_by_key.hpp
+++ b/rocprim/include/rocprim/device/device_scan_by_key.hpp
@@ -320,7 +320,7 @@ inline hipError_t scan_by_key_impl(void* const           temporary_storage,
 /// * Ranges specified by \p keys_input, \p values_input, and \p values_output must have
 /// at least \p size elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_by_key_config.
+/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_by_key_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam ValuesInputIterator - random-access iterator type of the input range. It can be
@@ -443,7 +443,7 @@ inline hipError_t inclusive_scan_by_key(void* const                temporary_sto
 /// * Ranges specified by \p keys_input, \p values_input, and \p values_output must have
 /// at least \p size elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_by_key_config.
+/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_by_key_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam ValuesInputIterator - random-access iterator type of the input range. It can be
diff --git a/rocprim/include/rocprim/device/device_segmented_radix_sort.hpp b/rocprim/include/rocprim/device/device_segmented_radix_sort.hpp
index 38fd3bb17..576789f06 100644
--- a/rocprim/include/rocprim/device/device_segmented_radix_sort.hpp
+++ b/rocprim/include/rocprim/device/device_segmented_radix_sort.hpp
@@ -590,7 +590,7 @@ hipError_t segmented_radix_sort_impl(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -712,7 +712,7 @@ hipError_t segmented_radix_sort_keys(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -835,7 +835,7 @@ hipError_t segmented_radix_sort_keys_desc(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -976,7 +976,7 @@ hipError_t segmented_radix_sort_pairs(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam KeysInputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam KeysOutputIterator - random-access iterator type of the output range. Must meet the
@@ -1117,7 +1117,7 @@ hipError_t segmented_radix_sort_pairs_desc(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam OffsetIterator - random-access iterator type of segment offsets. Must meet the
 /// requirements of a C++ OutputIterator concept. It can be a simple pointer type.
@@ -1245,7 +1245,7 @@ hipError_t segmented_radix_sort_keys(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam OffsetIterator - random-access iterator type of segment offsets. Must meet the
 /// requirements of a C++ OutputIterator concept. It can be a simple pointer type.
@@ -1373,7 +1373,7 @@ hipError_t segmented_radix_sort_keys_desc(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam Value - value type.
 /// \tparam OffsetIterator - random-access iterator type of segment offsets. Must meet the
@@ -1516,7 +1516,7 @@ hipError_t segmented_radix_sort_pairs(void * temporary_storage,
 /// [100, 10000], <tt>begin_bit = 0</tt> and <tt>end_bit = 14</tt> will cover the whole range.
 ///
 /// \tparam Config - [optional] configuration of the primitive. It has to be
-/// \p segmented_radix_sort_config.
+/// \p segmented_radix_sort_config or a class derived from it.
 /// \tparam Key - key type. Must be an integral type or a floating-point type.
 /// \tparam Value - value type.
 /// \tparam OffsetIterator - random-access iterator type of segment offsets. Must meet the
diff --git a/rocprim/include/rocprim/device/device_segmented_reduce.hpp b/rocprim/include/rocprim/device/device_segmented_reduce.hpp
index 4377a48b1..424b291ec 100644
--- a/rocprim/include/rocprim/device/device_segmented_reduce.hpp
+++ b/rocprim/include/rocprim/device/device_segmented_reduce.hpp
@@ -161,7 +161,7 @@ hipError_t segmented_reduce_impl(void * temporary_storage,
 /// <tt>segments + 1</tt> elements: <tt>offsets</tt> for \p begin_offsets and
 /// <tt>offsets + 1</tt> for \p end_offsets.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p reduce_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p reduce_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_segmented_scan.hpp b/rocprim/include/rocprim/device/device_segmented_scan.hpp
index 719905862..4cad6e0f4 100644
--- a/rocprim/include/rocprim/device/device_segmented_scan.hpp
+++ b/rocprim/include/rocprim/device/device_segmented_scan.hpp
@@ -166,7 +166,7 @@ hipError_t segmented_scan_impl(void * temporary_storage,
 /// <tt>segments + 1</tt> elements: <tt>offsets</tt> for \p begin_offsets and
 /// <tt>offsets + 1</tt> for \p end_offsets.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ RandomAccessIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
@@ -280,7 +280,7 @@ hipError_t segmented_inclusive_scan(void * temporary_storage,
 /// <tt>segments + 1</tt> elements: <tt>offsets</tt> for \p begin_offsets and
 /// <tt>offsets + 1</tt> for \p end_offsets.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ RandomAccessIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
@@ -397,7 +397,7 @@ hipError_t segmented_exclusive_scan(void * temporary_storage,
 /// * Ranges specified by \p input, \p output, and \p flags must have at least \p size elements.
 /// * \p value_type of \p HeadFlagIterator iterator should be convertible to \p bool type.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ RandomAccessIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
@@ -509,7 +509,7 @@ hipError_t segmented_inclusive_scan(void * temporary_storage,
 /// * Ranges specified by \p input, \p output, and \p flags must have at least \p size elements.
 /// * \p value_type of \p HeadFlagIterator iterator should be convertible to \p bool type.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p scan_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ RandomAccessIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
diff --git a/rocprim/include/rocprim/device/device_select.hpp b/rocprim/include/rocprim/device/device_select.hpp
index 43507c8af..68ce3bbb5 100644
--- a/rocprim/include/rocprim/device/device_select.hpp
+++ b/rocprim/include/rocprim/device/device_select.hpp
@@ -58,7 +58,7 @@ namespace detail
 /// * Range specified by \p selected_count_output must have at least 1 element.
 /// * Values of \p flag range should be implicitly convertible to `bool` type.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam FlagIterator - random-access iterator type of the flag range. It can be
@@ -181,7 +181,7 @@ hipError_t select(void * temporary_storage,
 /// values can be copied into it.
 /// * Range specified by \p selected_count_output must have at least 1 element.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p select_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. It can be
 /// a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. It can be
diff --git a/rocprim/include/rocprim/device/device_transform.hpp b/rocprim/include/rocprim/device/device_transform.hpp
index dddc92298..683b10b5a 100644
--- a/rocprim/include/rocprim/device/device_transform.hpp
+++ b/rocprim/include/rocprim/device/device_transform.hpp
@@ -82,7 +82,7 @@ ROCPRIM_KERNEL
 /// \par Overview
 /// * Ranges specified by \p input and \p output must have at least \p size elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p transform_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p transform_config or a class derived from it.
 /// \tparam InputIterator - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
@@ -207,7 +207,7 @@ inline hipError_t transform(InputIterator     input,
 /// \par Overview
 /// * Ranges specified by \p input1, \p input2, and \p output must have at least \p size elements.
 ///
-/// \tparam Config - [optional] configuration of the primitive. It has to be \p transform_config.
+/// \tparam Config - [optional] configuration of the primitive. It has to be \p transform_config or a class derived from it.
 /// \tparam InputIterator1 - random-access iterator type of the input range. Must meet the
 /// requirements of a C++ InputIterator concept. It can be a simple pointer type.
 /// \tparam InputIterator2 - random-access iterator type of the input range. Must meet the