[rocm-libraries] ROCm/rocm-libraries#1036 (commit 539ba71)

[rocPRIM][Code Coverage][Cherry Pick] Increase code coverage
 for rocPRIM (#1036)

In this commit we have added test coverage, made some bug fixes and
added a wrapper function for a common pattern in the tests. This
`test_kernel_wrapper` function reduces the amount of code used.

In

ROCm/rocm-libraries@0318f42
we have added coverage of the `warp_scan_shuffle` by forcing dpp off in
cmake on the `warp_scan` test.

In

ROCm/rocm-libraries@30d54c1
we have added the missing test cases for the thread_algos in the thread
directory. (Except for thread_operators, these functions are mostly
duplicates probably needs to be cleanly moved/deprecated). We also found
a bug in `thread_reduce` with the tests and fixed this.

In

ROCm/rocm-libraries@48780ff
ROCm/rocm-libraries@37aa542
ROCm/rocm-libraries@05a9e12
ROCm/rocm-libraries@f5f1dbe
ROCm/rocm-libraries@343e275
ROCm/rocm-libraries@8007fe6
we have added unit tests for the files in the types directory.

In

ROCm/rocm-libraries@40ca3e2
we have introduced the `test_kernel_wrapper` and added the missing
`PartitionTwoWayFlag` test.

In

ROCm/rocm-libraries@7a4e7ba
unit tests for the rocprim::tuple type are added.

In

ROCm/rocm-libraries@29eac9a
we have added extra test coverage by introducing a new type that will go
into the untested path. We have also added this type to the benchmark to
show [this

specialization](ROCm/rocm-libraries@29eac9a#diff-22a70b2ad081732e222004ded43ce0db7145ff196f7b723289425dcb5b6c732dR228)
is still needed. We also changed the `std::is_integral` to
`rocprim::is_integral` to not include `(u)int128_t` for this
specialization, this does not impact performance. The specialization
enable_if was also slightly changed to make it clearer which path is
chosen (does not make a difference in actual executed code). Also custom
config and a iterator was added to the tests of merge_sort.

In

ROCm/rocm-libraries@f5614d5
test coverage was added for the device_scan_common.hpp file.

In

ROCm/rocm-libraries@1d2a40a
test coverage was increased by actually using `const fixed_array` and
including a `Level` type which goes into the base path of
sample_to_bin_even struct. We also added an iterator as a type to the
test. Also the new test wrapper was used.

In

ROCm/rocm-libraries@c518f42
test coverage was increased for the iterators, a lot of the operators
where missing. Also some cleaning up of the tests was done including the
wrapper when possible. There was also a bug found for the comperator in
`arg_index_iterator` and `texture_cache_iterator`, which was also fixed.
The `->` operator is not tested for `test_texture_cache`,
`arg_index_iterator`, `transform_iterator` and `zip_iterator`. They
currently do not seem to work, I am working on a possible fix but will
be added in a later PR.

Author: NguyenNhuDi

CHANGELOG.md

@@ -97,6 +97,8 @@ when the input or inital type was smaller than the output type.
 * Fixed an issue where `device_segmented_reduce` reported autotuning throughput being 5x lower than it was in reality.
 * Fixed device radix sort not returning the correct required temporary storage when a double buffer contains `nullptr`.
 * Fixed constness of equality operators (`==` and `!=`) in `rocprim::key_value_pair`.
+* Fixed an issue for the comparison operators in `arg_index_iterator` and `texture_cache_iterator`, where `<` and `>` comparators were swapped.
+* Fixed an issue for the `rocprim::thread_reduce` not working correctly with a prefix value.
 
 ### Known issues
 * When using `rocprim::deterministic_inclusive_scan_by_key` and `rocprim::deterministic_exclusive_scan_by_key` the intermediate values can change order on Navi3x

benchmark/benchmark_device_merge_sort.cpp

@@ -50,10 +50,16 @@ int main(int argc, char* argv[])
     CREATE_BENCHMARK(rocprim::int128_t)
     CREATE_BENCHMARK(rocprim::uint128_t)
 
-    using custom_float2          = common::custom_type<float, float>;
-    using custom_double2         = common::custom_type<double, double>;
-    using custom_int2            = common::custom_type<int, int>;
-    using custom_char_double     = common::custom_type<char, double>; // used by ssbk benchmark
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
+    using custom_double2_copy
+        = common::custom_type_copyable<double,
+                                       double>; // specific benchmark for specialization workaround
+    using custom_int2        = common::custom_type<int, int>;
+    using custom_char_double = common::custom_type<char, double>; // used by ssbk
+    using custom_char_double_copy
+        = common::custom_type_copyable<char,
+                                       double>; // specific benchmark for specialization workaround
     using custom_longlong_double = common::custom_type<long long, double>;
     using huge_float2_1024       = common::custom_huge_type<1024, float, float>;
     using huge_float2_2048       = common::custom_huge_type<2048, float, float>;
@@ -69,8 +75,10 @@ int main(int argc, char* argv[])
     CREATE_BENCHMARK(huge_float2_2048)
     CREATE_BENCHMARK(long long, custom_double2)
     CREATE_BENCHMARK(custom_double2, custom_double2)
+    CREATE_BENCHMARK(custom_double2, custom_double2_copy)
     CREATE_BENCHMARK(custom_int2, custom_double2)
     CREATE_BENCHMARK(custom_int2, custom_char_double)
+    CREATE_BENCHMARK(custom_int2, custom_char_double_copy)
     CREATE_BENCHMARK(custom_int2, custom_longlong_double)
     CREATE_BENCHMARK(rocprim::int128_t, rocprim::int128_t)
     CREATE_BENCHMARK(rocprim::uint128_t, rocprim::uint128_t)

benchmark/benchmark_device_merge_sort.hpp

@@ -140,7 +140,21 @@ struct device_merge_sort_benchmark : public benchmark_utils::autotune_interface
                                         seed.get_0());
 
         std::vector<value_type> values_input(size);
-        std::iota(values_input.begin(), values_input.end(), 0);
+
+        if constexpr(common::is_custom_type_copyable<value_type>::value)
+        {
+            for(size_t i = 0; i < size; i++)
+            {
+                value_type value;
+                value.x         = static_cast<typename value_type::first_type>(i);
+                value.y         = static_cast<typename value_type::second_type>(i);
+                values_input[i] = value;
+            }
+        }
+        else
+        {
+            std::iota(values_input.begin(), values_input.end(), 0);
+        }
 
         key_type* d_keys_input;
         key_type* d_keys_output;

benchmark/benchmark_utils.hpp

@@ -694,8 +694,8 @@ struct bench_naming
     static std::string format_name(std::string string)
     {
         format     format = get_format();
-        std::regex r(
-            "([A-z0-9]*):\\s*((?:common::custom_type<[A-z0-9,]*>)|[A-z:\\(\\)\\.<>\\s0-9]*)(\\}*)");
+        std::regex r("([A-z0-9]*):\\s*((?:common::custom_type<[A-z0-9,]*>)|(?:common::custom_type_"
+                     "copyable<[A-z0-9,]*>)|[A-z:\\(\\)\\.<>\\s0-9]*)(\\}*)");
         // First we perform some checks
         bool checks[4] = {false};
         for(std::sregex_iterator i = std::sregex_iterator(string.begin(), string.end(), r);
@@ -917,6 +917,16 @@ inline const char* Traits<rocprim::uint128_t>::name()
 {
     return "rocprim::uint128_t";
 }
+template<>
+inline const char* Traits<common::custom_type_copyable<char, double>>::name()
+{
+    return "common::custom_type_copyable<char,double>";
+}
+template<>
+inline const char* Traits<common::custom_type_copyable<double, double>>::name()
+{
+    return "common::custom_type_copyable<double,double>";
+}
 
 inline const char* get_block_scan_algorithm_name(rocprim::block_scan_algorithm alg)
 {

common/utils_custom_type.hpp

@@ -138,6 +138,85 @@ struct custom_type
     }
 };
 
+template<typename T, typename U = T, bool NonZero = false>
+struct custom_type_copyable
+{
+    using first_type  = T;
+    using second_type = U;
+
+    using value_type = std::conditional_t<std::is_same<T, U>::value, T, void>;
+
+    T x;
+    U y;
+
+    ROCPRIM_HOST_DEVICE constexpr inline custom_type_copyable()
+        : x(NonZero ? 12 : 0), y(NonZero ? 34 : 0)
+    {}
+
+    ROCPRIM_HOST_DEVICE inline custom_type_copyable(T x, U y) : x(x), y(y) {}
+
+    ROCPRIM_HOST_DEVICE inline custom_type_copyable(T xy) : x(xy), y(xy) {}
+
+    template<typename V, typename W>
+    ROCPRIM_HOST_DEVICE inline custom_type_copyable(
+        const custom_type_copyable<V, W, NonZero>& other)
+        : x(static_cast<T>(other.x)), y(static_cast<U>(other.y))
+    {}
+
+    ROCPRIM_HOST_DEVICE
+    inline bool
+        operator<(const custom_type_copyable& other) const
+    {
+        rocprim::less<T>     less_T;
+        rocprim::equal_to<T> equal_to_T;
+        rocprim::less<U>     less_U;
+        return (less_T(x, other.x) || (equal_to_T(x, other.x) && less_U(y, other.y)));
+    }
+
+    ROCPRIM_HOST_DEVICE
+    inline bool
+        operator>(const custom_type_copyable& other) const
+    {
+        rocprim::greater<T>  greater_T;
+        rocprim::equal_to<T> equal_to_T;
+        rocprim::greater<U>  greater_U;
+        return (greater_T(x, other.x) || (equal_to_T(x, other.x) && greater_U(y, other.y)));
+    }
+
+    ROCPRIM_HOST_DEVICE
+    inline bool
+        operator==(const custom_type_copyable& other) const
+    {
+        rocprim::equal_to<T> equal_to_T;
+        rocprim::equal_to<U> equal_to_U;
+        return (equal_to_T(x, other.x) && equal_to_U(y, other.y));
+    }
+
+    ROCPRIM_HOST_DEVICE
+    inline bool
+        operator!=(const custom_type_copyable& other) const
+    {
+        return !(*this == other);
+    }
+
+    friend inline std::ostream& operator<<(std::ostream& stream, const custom_type_copyable& value)
+    {
+        stream << "[" << value.x << "; " << value.y << "]";
+        return stream;
+    }
+};
+
+static_assert(std::is_trivially_copyable<custom_type_copyable<char, double>>::value,
+              "custom_type_copyable is not trivially copyable");
+
+template<typename T>
+struct is_custom_type_copyable : std::false_type
+{};
+
+template<typename T, typename U, bool NonZero>
+struct is_custom_type_copyable<common::custom_type_copyable<T, U, NonZero>> : std::true_type
+{};
+
 template<unsigned int Size, class T, typename U = T, bool NonZero = false>
 struct custom_huge_type : custom_type<T, U, NonZero>
 {
@@ -182,6 +261,10 @@ struct is_custom_type<common::custom_type<T, U, NonZero>> : std::true_type
 template<unsigned int Size, typename T, typename U, bool NonZero>
 struct is_custom_type<common::custom_huge_type<Size, T, U, NonZero>> : std::true_type
 {};
+
+template<typename T, typename U, bool NonZero>
+struct is_custom_type<common::custom_type_copyable<T, U, NonZero>> : std::true_type
+{};
 } // namespace common
 
 #endif // COMMON_UTILS_CUSTOM_TYPE_HPP_

rocprim/include/rocprim/config.hpp

@@ -150,11 +150,13 @@
 // Only defined when support is present, in contrast to ROCPRIM_DETAIL_USE_DPP, which should be
 // always defined
 #if defined(__HIP_DEVICE_COMPILE__) && defined(__AMDGCN__) \
-    && (!defined(__GFX6__) && !defined(__GFX7__))
+    && (!defined(__GFX6__) && !defined(__GFX7__))          \
+    && !(defined(ROCPRIM_TARGET_SPIRV) && ROCPRIM_TARGET_SPIRV == 1)
     #define ROCPRIM_DETAIL_HAS_DPP 1
 #endif
 
-#if !defined(ROCPRIM_DISABLE_DPP) && defined(ROCPRIM_DETAIL_HAS_DPP) && !ROCPRIM_TARGET_SPIRV
+#if(!defined(ROCPRIM_DISABLE_DPP) || ROCPRIM_DISABLE_DPP == 0) \
+    && (defined(ROCPRIM_DETAIL_HAS_DPP) && ROCPRIM_DETAIL_HAS_DPP == 1)
     #define ROCPRIM_DETAIL_USE_DPP 1
 #else
     #define ROCPRIM_DETAIL_USE_DPP 0

rocprim/include/rocprim/device/detail/device_histogram.hpp

@@ -306,8 +306,9 @@ ROCPRIM_DEVICE ROCPRIM_INLINE void
 }
 
 template<unsigned int BlockSize, unsigned int ActiveChannels, class Counter>
-ROCPRIM_DEVICE ROCPRIM_INLINE void init_histogram(fixed_array<Counter*, ActiveChannels> histogram,
-                                                  fixed_array<unsigned int, ActiveChannels> bins)
+ROCPRIM_DEVICE ROCPRIM_INLINE
+void init_histogram(fixed_array<Counter*, ActiveChannels>           histogram,
+                    const fixed_array<unsigned int, ActiveChannels> bins)
 {
     const unsigned int flat_id  = ::rocprim::detail::block_thread_id<0>();
     const unsigned int block_id = ::rocprim::detail::block_id<0>();
@@ -329,17 +330,17 @@ template<unsigned int BlockSize,
          class SampleIterator,
          class Counter,
          class SampleToBinOp>
-ROCPRIM_DEVICE ROCPRIM_INLINE void
-    histogram_shared(SampleIterator                             samples,
-                     unsigned int                               columns,
-                     unsigned int                               rows,
-                     unsigned int                               row_stride,
-                     unsigned int                               rows_per_block,
-                     unsigned int                               shared_histograms,
-                     fixed_array<Counter*, ActiveChannels>      histogram,
-                     fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
-                     fixed_array<unsigned int, ActiveChannels>  bins,
-                     unsigned int*                              block_histogram_start)
+ROCPRIM_DEVICE ROCPRIM_INLINE
+void histogram_shared(SampleIterator                                   samples,
+                      unsigned int                                     columns,
+                      unsigned int                                     rows,
+                      unsigned int                                     row_stride,
+                      unsigned int                                     rows_per_block,
+                      unsigned int                                     shared_histograms,
+                      fixed_array<Counter*, ActiveChannels>            histogram,
+                      const fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
+                      const fixed_array<unsigned int, ActiveChannels>  bins,
+                      unsigned int*                                    block_histogram_start)
 {
     using sample_type        = typename std::iterator_traits<SampleIterator>::value_type;
     using sample_vector_type = sample_vector<sample_type, Channels>;
@@ -456,13 +457,13 @@ template<unsigned int BlockSize,
          class SampleIterator,
          class Counter,
          class SampleToBinOp>
-ROCPRIM_DEVICE ROCPRIM_INLINE void
-    histogram_global(SampleIterator                             samples,
-                     unsigned int                               columns,
-                     unsigned int                               row_stride,
-                     fixed_array<Counter*, ActiveChannels>      histogram,
-                     fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
-                     fixed_array<unsigned int, ActiveChannels>  bins_bits)
+ROCPRIM_DEVICE ROCPRIM_INLINE
+void histogram_global(SampleIterator                                   samples,
+                      unsigned int                                     columns,
+                      unsigned int                                     row_stride,
+                      fixed_array<Counter*, ActiveChannels>            histogram,
+                      const fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
+                      const fixed_array<unsigned int, ActiveChannels>  bins_bits)
 {
     using sample_type        = typename std::iterator_traits<SampleIterator>::value_type;
     using sample_vector_type = sample_vector<sample_type, Channels>;

rocprim/include/rocprim/device/detail/device_merge_sort.hpp

@@ -230,7 +230,7 @@ struct block_permute_values_impl<Value,
                                  ItemsPerThread,
                                  std::enable_if_t<(std::is_trivially_copyable<Value>::value
                                                    && !rocprim::is_floating_point<Value>::value
-                                                   && !std::is_integral<Value>::value)>>
+                                                   && !rocprim::is_integral<Value>::value)>>
 {
     static constexpr unsigned int items_per_block = ItemsPerThread * BlockSize;
 

rocprim/include/rocprim/device/detail/device_merge_sort_mergepath.hpp

@@ -98,8 +98,9 @@ struct block_merge_impl<
     Value,
     BlockSize,
     ItemsPerThread,
-    std::enable_if_t<!std::is_trivially_copyable<Value>::value
-                     || rocprim::is_floating_point<Value>::value || std::is_integral<Value>::value>>
+    std::enable_if_t<
+        !std::is_trivially_copyable<Value>::value || rocprim::is_floating_point<Value>::value
+        || rocprim::is_integral<Value>::value || std::is_same<Value, ::rocprim::empty_type>::value>>
 {
 
     static constexpr bool         with_values = !std::is_same<Value, ::rocprim::empty_type>::value;
@@ -264,9 +265,9 @@ struct block_merge_impl<Key,
                         ItemsPerThread,
                         std::enable_if_t<std::is_trivially_copyable<Value>::value
                                          && !rocprim::is_floating_point<Value>::value
-                                         && !std::is_integral<Value>::value>>
+                                         && !rocprim::is_integral<Value>::value
+                                         && !std::is_same<Value, ::rocprim::empty_type>::value>>
 {
-
     static constexpr bool         with_values = !std::is_same<Value, ::rocprim::empty_type>::value;
     static constexpr unsigned int items_per_tile = BlockSize * ItemsPerThread;
 
@@ -301,7 +302,6 @@ struct block_merge_impl<Key,
                                                               const OffsetT*       merge_partitions,
                                                               storage_type&        storage)
         {
-
             auto& keys_shared   = storage.keys.get();
             auto& values_shared = storage.values.get();
 

rocprim/include/rocprim/device/device_histogram.hpp

@@ -44,8 +44,8 @@ namespace detail
 
 template<class Config, unsigned int ActiveChannels, class Counter>
 ROCPRIM_KERNEL ROCPRIM_LAUNCH_BOUNDS(device_params<Config>().histogram_config.block_size) void
-    init_histogram_kernel(fixed_array<Counter*, ActiveChannels>     histogram,
-                          fixed_array<unsigned int, ActiveChannels> bins)
+    init_histogram_kernel(fixed_array<Counter*, ActiveChannels>           histogram,
+                          const fixed_array<unsigned int, ActiveChannels> bins)
 {
     static constexpr histogram_config_params params = device_params<Config>();
 
@@ -59,15 +59,15 @@ template<class Config,
          class Counter,
          class SampleToBinOp>
 ROCPRIM_KERNEL ROCPRIM_LAUNCH_BOUNDS(device_params<Config>().histogram_config.block_size) void
-    histogram_shared_kernel(SampleIterator                             samples,
-                            unsigned int                               columns,
-                            unsigned int                               rows,
-                            unsigned int                               row_stride,
-                            unsigned int                               rows_per_block,
-                            unsigned int                               shared_histograms,
-                            fixed_array<Counter*, ActiveChannels>      histogram,
-                            fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
-                            fixed_array<unsigned int, ActiveChannels>  bins)
+    histogram_shared_kernel(SampleIterator                                   samples,
+                            unsigned int                                     columns,
+                            unsigned int                                     rows,
+                            unsigned int                                     row_stride,
+                            unsigned int                                     rows_per_block,
+                            unsigned int                                     shared_histograms,
+                            fixed_array<Counter*, ActiveChannels>            histogram,
+                            const fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
+                            const fixed_array<unsigned int, ActiveChannels>  bins)
 {
     static constexpr histogram_config_params params = device_params<Config>();
 
@@ -100,12 +100,12 @@ template<class Config,
          class Counter,
          class SampleToBinOp>
 ROCPRIM_KERNEL ROCPRIM_LAUNCH_BOUNDS(device_params<Config>().histogram_config.block_size) void
-    histogram_global_kernel(SampleIterator                             samples,
-                            unsigned int                               columns,
-                            unsigned int                               row_stride,
-                            fixed_array<Counter*, ActiveChannels>      histogram,
-                            fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
-                            fixed_array<unsigned int, ActiveChannels>  bins_bits)
+    histogram_global_kernel(SampleIterator                                   samples,
+                            unsigned int                                     columns,
+                            unsigned int                                     row_stride,
+                            fixed_array<Counter*, ActiveChannels>            histogram,
+                            const fixed_array<SampleToBinOp, ActiveChannels> sample_to_bin_op,
+                            const fixed_array<unsigned int, ActiveChannels>  bins_bits)
 {
     static constexpr histogram_config_params params = device_params<Config>();