From 31554f8f6e842e778dddcb8f5a248759d1f38149 Mon Sep 17 00:00:00 2001 From: Georgy Evtushenko Date: Mon, 18 Sep 2023 23:58:22 +0000 Subject: [PATCH 1/2] Test signed offset types in segmented reduce --- .../catch2_test_device_segmented_reduce.cu | 31 ++++++++++--------- ..._test_device_segmented_reduce_iterators.cu | 12 ++++--- ...device_segmented_reduce_iterators_64bit.cu | 11 +++---- 3 files changed, 28 insertions(+), 26 deletions(-) diff --git a/cub/test/catch2_test_device_segmented_reduce.cu b/cub/test/catch2_test_device_segmented_reduce.cu index 416cad26d398..a21ca1767157 100644 --- a/cub/test/catch2_test_device_segmented_reduce.cu +++ b/cub/test/catch2_test_device_segmented_reduce.cu @@ -79,14 +79,17 @@ type_pair // clang-format on #endif +using offsets = c2h::type_list; + CUB_TEST("Device reduce works with all device interfaces", "[segmented][reduce][device]", - full_type_list) + full_type_list, + offsets) { - using params = params_t; - using item_t = typename params::item_t; - using output_t = typename params::output_t; - using offset_t = int32_t; + using type_pair_t = typename c2h::get<0, TestType>; + using input_t = typename type_pair_t::input_t; + using output_t = typename type_pair_t::output_t; + using offset_t = typename c2h::get<1, TestType>; constexpr int min_items = 1; constexpr int max_items = 1000000; @@ -116,7 +119,7 @@ CUB_TEST("Device reduce works with all device interfaces", auto d_offsets_it = thrust::raw_pointer_cast(segment_offsets.data()); // Generate input data - thrust::device_vector in_items(num_items); + thrust::device_vector in_items(num_items); c2h::gen(CUB_SEED(2), in_items); auto d_in_it = thrust::raw_pointer_cast(in_items.data()); @@ -128,7 +131,7 @@ CUB_TEST("Device reduce works with all device interfaces", auto reduction_op = unwrap_op(reference_extended_fp(d_in_it), op_t{}); // Prepare verification data - using accum_t = cub::detail::accumulator_t; + using accum_t = cub::detail::accumulator_t; thrust::host_vector expected_result(num_segments); compute_segmented_problem_reference(in_items, segment_offsets, @@ -158,7 +161,7 @@ CUB_TEST("Device reduce works with all device interfaces", SECTION("sum") { using op_t = cub::Sum; - using accum_t = cub::detail::accumulator_t; + using accum_t = cub::detail::accumulator_t; // Prepare verification data thrust::host_vector expected_result(num_segments); @@ -187,7 +190,7 @@ CUB_TEST("Device reduce works with all device interfaces", compute_segmented_problem_reference(in_items, segment_offsets, op_t{}, - cub::NumericTraits::Max(), + cub::NumericTraits::Max(), expected_result.begin()); // Run test @@ -212,7 +215,7 @@ CUB_TEST("Device reduce works with all device interfaces", compute_segmented_problem_reference(in_items, segment_offsets, op_t{}, - cub::NumericTraits::Lowest(), + cub::NumericTraits::Lowest(), expected_result.begin()); // Run test @@ -230,14 +233,14 @@ CUB_TEST("Device reduce works with all device interfaces", SECTION("argmax") { - using result_t = cub::KeyValuePair; + using result_t = cub::KeyValuePair; // Prepare verification data thrust::host_vector expected_result(num_segments); compute_segmented_argmax_reference(in_items, segment_offsets, expected_result.begin()); // Run test - using result_t = cub::KeyValuePair; + using result_t = cub::KeyValuePair; thrust::device_vector out_result(num_segments); device_segmented_arg_max(d_in_it, thrust::raw_pointer_cast(out_result.data()), @@ -251,10 +254,10 @@ CUB_TEST("Device reduce works with all device interfaces", SECTION("argmin") { - using result_t = cub::KeyValuePair; + using result_t = cub::KeyValuePair; // Prepare verification data - thrust::host_vector host_items(in_items); + thrust::host_vector host_items(in_items); thrust::host_vector expected_result(num_segments); compute_segmented_argmin_reference(in_items, segment_offsets, expected_result.begin()); diff --git a/cub/test/catch2_test_device_segmented_reduce_iterators.cu b/cub/test/catch2_test_device_segmented_reduce_iterators.cu index 4da9bf8048f8..10c87471a17a 100644 --- a/cub/test/catch2_test_device_segmented_reduce_iterators.cu +++ b/cub/test/catch2_test_device_segmented_reduce_iterators.cu @@ -52,15 +52,17 @@ DECLARE_CDP_WRAPPER(cub::DeviceSegmentedReduce::Sum, device_segmented_sum); // List of types to test using custom_t = c2h::custom_type_t; using iterator_type_list = c2h::type_list, type_pair>; +using offsets = c2h::type_list; CUB_TEST("Device segmented reduce works with fancy input iterators", "[reduce][device]", - iterator_type_list) + iterator_type_list, + offsets) { - using params = params_t; - using item_t = typename params::item_t; - using output_t = typename params::output_t; - using offset_t = uint32_t; + using type_pair_t = typename c2h::get<0, TestType>; + using item_t = typename type_pair_t::input_t; + using output_t = typename type_pair_t::output_t; + using offset_t = typename c2h::get<1, TestType>; constexpr int min_items = 1; constexpr int max_items = 1000000; diff --git a/cub/test/catch2_test_device_segmented_reduce_iterators_64bit.cu b/cub/test/catch2_test_device_segmented_reduce_iterators_64bit.cu index 3cc7ccc02912..f39a0d50112e 100644 --- a/cub/test/catch2_test_device_segmented_reduce_iterators_64bit.cu +++ b/cub/test/catch2_test_device_segmented_reduce_iterators_64bit.cu @@ -45,17 +45,13 @@ DECLARE_CDP_WRAPPER(cub::DeviceSegmentedReduce::Sum, device_segmented_sum); // %PARAM% TEST_CDP cdp 0:1 // List of types to test -using offset_type_list = c2h::type_list; -using num_segments_type_list = c2h:: - type_list; +using offsets = c2h::type_list; CUB_TEST("Device segmented reduce works with fancy input iterators and 64-bit offsets", "[reduce][device]", - offset_type_list, - num_segments_type_list) + offsets) { using offset_t = typename c2h::get<0, TestType>; - using num_segments_t = typename c2h::get<1, TestType>; using op_t = cub::Sum; constexpr offset_t offset_zero = 0; @@ -63,7 +59,8 @@ CUB_TEST("Device segmented reduce works with fancy input iterators and 64-bit of constexpr offset_t iterator_value = 2; constexpr offset_t min_items_per_segment = offset_one << 31; constexpr offset_t max_items_per_segment = offset_one << 33; - constexpr num_segments_t num_segments = 2; + + constexpr int num_segments = 2; // generate individual segment lengths and store cumulative sum in segment_offsets const offset_t num_items_in_first_segment = From 46dc4c8b44ccad4e722c833e87612e840f7e94a2 Mon Sep 17 00:00:00 2001 From: Georgy Evtushenko Date: Tue, 19 Sep 2023 20:03:53 +0000 Subject: [PATCH 2/2] Return old offset type to segmented arg reductions --- cub/cub/device/device_segmented_reduce.cuh | 8 ++++++-- cub/test/catch2_test_device_segmented_reduce.cu | 5 ++--- 2 files changed, 8 insertions(+), 5 deletions(-) diff --git a/cub/cub/device/device_segmented_reduce.cuh b/cub/cub/device/device_segmented_reduce.cuh index c205c5508fe8..a776ac0f9e7c 100644 --- a/cub/cub/device/device_segmented_reduce.cuh +++ b/cub/cub/device/device_segmented_reduce.cuh @@ -731,7 +731,9 @@ struct DeviceSegmentedReduce cudaStream_t stream = 0) { // Integer type for global offsets - using OffsetT = detail::common_iterator_value_t; + // Using common iterator value type is a breaking change, see: + // https://github.com/NVIDIA/cccl/pull/414#discussion_r1330632615 + using OffsetT = int; // detail::common_iterator_value_t; // The input type using InputValueT = cub::detail::value_t; @@ -1097,7 +1099,9 @@ struct DeviceSegmentedReduce cudaStream_t stream = 0) { // Integer type for global offsets - using OffsetT = detail::common_iterator_value_t; + // Using common iterator value type is a breaking change, see: + // https://github.com/NVIDIA/cccl/pull/414#discussion_r1330632615 + using OffsetT = int; // detail::common_iterator_value_t; // The input type using InputValueT = cub::detail::value_t; diff --git a/cub/test/catch2_test_device_segmented_reduce.cu b/cub/test/catch2_test_device_segmented_reduce.cu index a21ca1767157..f9dca2dc8d46 100644 --- a/cub/test/catch2_test_device_segmented_reduce.cu +++ b/cub/test/catch2_test_device_segmented_reduce.cu @@ -233,14 +233,13 @@ CUB_TEST("Device reduce works with all device interfaces", SECTION("argmax") { - using result_t = cub::KeyValuePair; + using result_t = cub::KeyValuePair; // Prepare verification data thrust::host_vector expected_result(num_segments); compute_segmented_argmax_reference(in_items, segment_offsets, expected_result.begin()); // Run test - using result_t = cub::KeyValuePair; thrust::device_vector out_result(num_segments); device_segmented_arg_max(d_in_it, thrust::raw_pointer_cast(out_result.data()), @@ -254,7 +253,7 @@ CUB_TEST("Device reduce works with all device interfaces", SECTION("argmin") { - using result_t = cub::KeyValuePair; + using result_t = cub::KeyValuePair; // Prepare verification data thrust::host_vector host_items(in_items);