Skip to content

Device reduce 64bit #1

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
jecs merged 2 commits into from
Sep 19, 2023
Merged

Device reduce 64bit #1

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
31554f8
Test signed offset types in segmented reduce
gevtushenko Sep 18, 2023
46dc4c8
Return old offset type to segmented arg reductions
gevtushenko Sep 19, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 6 additions & 2 deletions cub/cub/device/device_segmented_reduce.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -731,7 +731,9 @@ struct DeviceSegmentedReduce
cudaStream_t stream = 0)
{
// Integer type for global offsets
using OffsetT = detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
// 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<BeginOffsetIteratorT, EndOffsetIteratorT>;

// The input type
using InputValueT = cub::detail::value_t<InputIteratorT>;
Expand Down Expand Up @@ -1097,7 +1099,9 @@ struct DeviceSegmentedReduce
cudaStream_t stream = 0)
{
// Integer type for global offsets
using OffsetT = detail::common_iterator_value_t<BeginOffsetIteratorT, EndOffsetIteratorT>;
// 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<BeginOffsetIteratorT, EndOffsetIteratorT>;

// The input type
using InputValueT = cub::detail::value_t<InputIteratorT>;
Expand Down
26 changes: 14 additions & 12 deletions cub/test/catch2_test_device_segmented_reduce.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,14 +79,17 @@ type_pair<custom_t>
// clang-format on
#endif

using offsets = c2h::type_list<std::int32_t, std::uint32_t>;

CUB_TEST("Device reduce works with all device interfaces",
"[segmented][reduce][device]",
full_type_list)
full_type_list,
offsets)
{
using params = params_t<TestType>;
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;
Expand Down Expand Up @@ -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<item_t> in_items(num_items);
thrust::device_vector<input_t> in_items(num_items);
c2h::gen(CUB_SEED(2), in_items);
auto d_in_it = thrust::raw_pointer_cast(in_items.data());

Expand All @@ -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<op_t, output_t, item_t>;
using accum_t = cub::detail::accumulator_t<op_t, output_t, input_t>;
thrust::host_vector<output_t> expected_result(num_segments);
compute_segmented_problem_reference(in_items,
segment_offsets,
Expand Down Expand Up @@ -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<op_t, output_t, item_t>;
using accum_t = cub::detail::accumulator_t<op_t, output_t, input_t>;

// Prepare verification data
thrust::host_vector<output_t> expected_result(num_segments);
Expand Down Expand Up @@ -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<item_t>::Max(),
cub::NumericTraits<input_t>::Max(),
expected_result.begin());

// Run test
Expand All @@ -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<item_t>::Lowest(),
cub::NumericTraits<input_t>::Lowest(),
expected_result.begin());

// Run test
Expand All @@ -237,7 +240,6 @@ CUB_TEST("Device reduce works with all device interfaces",
compute_segmented_argmax_reference(in_items, segment_offsets, expected_result.begin());

// Run test
using result_t = cub::KeyValuePair<int, output_t>;
thrust::device_vector<result_t> out_result(num_segments);
device_segmented_arg_max(d_in_it,
thrust::raw_pointer_cast(out_result.data()),
Expand All @@ -254,7 +256,7 @@ CUB_TEST("Device reduce works with all device interfaces",
using result_t = cub::KeyValuePair<int, output_t>;

// Prepare verification data
thrust::host_vector<item_t> host_items(in_items);
thrust::host_vector<input_t> host_items(in_items);
thrust::host_vector<result_t> expected_result(num_segments);
compute_segmented_argmin_reference(in_items, segment_offsets, expected_result.begin());

Expand Down
12 changes: 7 additions & 5 deletions cub/test/catch2_test_device_segmented_reduce_iterators.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<c2h::accumulateable_t, c2h::equal_comparable_t>;
using iterator_type_list = c2h::type_list<type_pair<custom_t>, type_pair<std::int64_t>>;
using offsets = c2h::type_list<std::int32_t, std::uint32_t>;

CUB_TEST("Device segmented reduce works with fancy input iterators",
"[reduce][device]",
iterator_type_list)
iterator_type_list,
offsets)
{
using params = params_t<TestType>;
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;
Expand Down
11 changes: 4 additions & 7 deletions cub/test/catch2_test_device_segmented_reduce_iterators_64bit.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,25 +45,22 @@ 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<std::ptrdiff_t, std::size_t>;
using num_segments_type_list = c2h::
type_list<std::int16_t, std::int32_t, std::int64_t, std::uint16_t, std::uint32_t, std::uint64_t>;
using offsets = c2h::type_list<std::ptrdiff_t, std::size_t>;

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;
constexpr offset_t offset_one = 1;
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 =
Expand Down