Reimplement `cudf::merge` for nested types without using comparators by divyegala · Pull Request #14250 · NVIDIA/cudf

divyegala · 2023-10-04T20:48:56Z

Description

This PR also uses new experimental comparators for non-nested types by introducing a new device constructor for cudf::experimental::row::lexicographic::device_row_comparator. In the case of non-nested types, preprocessing can be skipped so comparators can be created on the fly. This solution helps us avoid creating 3 comparator types because thrust::merge can call the operator with indices from either side of the table.

Furthermore, the PR reworks cudf/detail/merge.cuh by removing any CUDA headers/components to expose a true detail API of the form cudf/detail/merge.hpp.

Benchmark comparison for non-nested types

Compilation time increases from ~6 mins to ~7 mins.

Checklist

I am familiar with the Contributing Guidelines.
New or existing tests cover these changes.
The documentation is up to date with these changes.

divyegala · 2023-10-04T20:49:49Z

Benchmark                                          Time             CPU      Time Old      Time New       CPU Old       CPU New
-------------------------------------------------------------------------------------------------------------------------------
Merge/pow2tables/2/manual_time                  +0.0034         +0.0051             0             0             0             0
Merge/pow2tables/4/manual_time                  +0.0150         +0.0136             1             1             1             1
Merge/pow2tables/8/manual_time                  +0.0045         +0.0042             2             2             2             2
Merge/pow2tables/16/manual_time                 +0.0108         +0.0108             4             4             4             4
Merge/pow2tables/32/manual_time                 +0.0128         +0.0127             9            10            10            10
Merge/pow2tables/64/manual_time                 +0.0123         +0.0123            22            23            22            23
Merge/pow2tables/128/manual_time                +0.0147         +0.0147            52            53            52            53
OVERALL_GEOMEAN                                 +0.0105         +0.0105             0             0             0             0

copy-pr-bot · 2023-10-16T19:47:01Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

divyegala · 2023-10-16T19:47:13Z

/ok to test

divyegala · 2023-10-16T19:54:36Z

/ok to test

bdice

This is great. I have a couple early suggestions on the draft.

divyegala · 2023-10-16T22:25:20Z

/ok to test

divyegala · 2023-10-16T23:44:07Z

/ok to test

divyegala · 2023-10-17T17:15:28Z

/ok to test

divyegala · 2023-10-20T17:34:46Z

/ok to test

bdice

Awesome work. I have a few comments.

Co-authored-by: Bradley Dice <bdice@bradleydice.com>

divyegala · 2023-10-26T00:41:59Z

/ok to test

bdice

Approving. I had one comment about whether a comment is really a TODO or not, but otherwise LGTM.

divyegala · 2023-10-26T22:28:33Z

/ok to test

divyegala · 2023-10-26T22:30:17Z

/ok to test

mythrocks

This is looking very good. I have a couple of trivial nitpicks.

I'm not done. I've just hit the meat of the change. I'll resume the review shortly.

mythrocks · 2023-10-27T18:08:54Z

+   * `null_order::BEFORE` for all columns.
+   * @param comparator Physical element relational comparison functor.
+   */
+  template <bool nested_disable = not has_nested_columns, CUDF_ENABLE_IF(nested_disable)>


Interesting phrasing. I'll remember this for when I need it.

Is nested_disable used elsewhere?
If not, would CUDF_ENABLE_IF(not has_nested_columns) be less readable?

I can change the variable name if you'd like, but I cannot use the class-template has_nested_columns directly here.

Let me explain why, and my apologies if you already knew this but I misunderstood what you are suggesting.

If I did:

template <CUDF_ENABLE_IF(not has_nested_columns)>

then we'd not be able to trigger SFINAE. This is because has_nested_columns is a class-template. So, when you write device_comparator<true> d_comp, the substitution for the enable-if on the constructor will trigger a failure as enable_if::type will be ill-formed.

As a result of which, we want to defer the SFINAE at the member function level so that it behaves the way it is intended to: as an overload.

mythrocks · 2023-10-27T21:35:39Z

+{
+  size_type const left_size  = left_table.num_rows();
+  size_type const right_size = right_table.num_rows();
+  size_type const total_size = left_size + right_size;


Another impertinent question: Do we need to protect against the case where the total_size overflows?

Good point! I added a check here https://github.com/rapidsai/cudf/pull/14250/files#diff-bea7124eddd4e287917be21bd0a0b7cb946816ad08c2a77943177a54569d87d0R641-R647

mythrocks

LGTM. A couple of minor questions.
Edit: Thanks, that was an interesting read. I've learnt a thing or two from this change.

divyegala · 2023-10-27T22:08:14Z

/ok to test

divyegala · 2023-10-28T00:31:21Z

/merge

bdice · 2023-10-30T14:16:58Z

+  CUDF_EXPECTS(std::accumulate(tables_to_merge.cbegin(),
+                               tables_to_merge.cend(),
+                               cudf::size_type{0},
+                               [](auto const& running_sum, auto const& tbl) {
+                                 return running_sum + tbl.num_rows();
+                               }) <= std::numeric_limits<cudf::size_type>::max(),
+               "Total number of merged rows exceeds row limit");


@divyegala (cc: @mythrocks) I don't think this actually protects against overflow. By the definition of overflow (if it were even defined behavior for signed types, which it's not...), the sum will always be less than the max limit of size_type when it rolls over. You need to use a larger type and accumulate in std::size_t precision. Can you file a follow-up PR with this change?

Suggested change

CUDF_EXPECTS(std::accumulate(tables_to_merge.cbegin(),

tables_to_merge.cend(),

cudf::size_type{0},

[](auto const& running_sum, auto const& tbl) {

return running_sum + tbl.num_rows();

}) <= std::numeric_limits<cudf::size_type>::max(),

"Total number of merged rows exceeds row limit");

CUDF_EXPECTS(std::accumulate(tables_to_merge.cbegin(),

tables_to_merge.cend(),

size_t{0},

[](auto const& running_sum, auto const& tbl) {

return running_sum + static_cast<size_t>(tbl.num_rows());

}) <= static_cast<size_t>(std::numeric_limits<cudf::size_type>::max()),

"Total number of merged rows exceeds row limit");

Examples of this pattern:

https://github.com/rapidsai/cudf/blob/2a923dfff84a8ce3422a126e9d0035282a0bea2f/cpp/src/strings/copying/concatenate.cu#L221

https://github.com/rapidsai/cudf/blob/2a923dfff84a8ce3422a126e9d0035282a0bea2f/cpp/src/copying/scatter.cu#L324

https://github.com/rapidsai/cudf/blob/2a923dfff84a8ce3422a126e9d0035282a0bea2f/cpp/src/scalar/scalar.cpp#L70

Thanks for pointing that out @bdice! I'll raise a PR

@bdice

#14250 added a check to ensure `cudf::merge` throws when the total number of merged rows exceed `cudf::size_type` limit, however @bdice pointed out that the check was not correct because the accumulation was still occurring in `cudf::size_type`. This PR computes the accumulation in `std::size_t`. Authors: - Divye Gala (https://github.com/divyegala) Approvers: - Bradley Dice (https://github.com/bdice) - MithunR (https://github.com/mythrocks) - Nghia Truong (https://github.com/ttnghia) URL: #14345

passing tests

b0dc5d3

divyegala added improvement Improvement / enhancement to an existing function non-breaking Non-breaking change labels Oct 4, 2023

github-actions Bot added the libcudf Affects libcudf (C++/CUDA) code. label Oct 4, 2023

divyegala added 3 commits October 4, 2023 13:57

use new algorithm only for nested types

3343ddb

auto const everything

2426c0b

experimental comparators for non-nested case

264b112

Merge branch 'branch-23.12' into merge-new-algo

9d9fe7c

divyegala changed the title ~~Reimplement cudf::merge without using comparators~~ Reimplement cudf::merge for nested types without using comparators Oct 16, 2023

bdice reviewed Oct 16, 2023

View reviewed changes

Comment thread cpp/src/merge/merge.cu

Comment thread cpp/src/merge/merge.cu Outdated

add device constructor for experimental comparator

61dec01

divyegala added 4 commits October 16, 2023 16:39

fix regression, address initial review

5329634

Merge branch 'branch-23.12' into merge-new-algo

66964b6

default construction unnecessary

377c9cd

add noexcept back

29d51c0

Merge branch 'branch-23.12' into merge-new-algo

c4036e0

divyegala added 2 commits October 20, 2023 10:29

passing tests, benchmarks, and detail API rework

b62fe89

delete accidentally included files

d3a2425

github-actions Bot added the CMake CMake build issue label Oct 20, 2023

Merge branch 'branch-23.12' into merge-new-algo

b799e5f

divyegala marked this pull request as ready for review October 20, 2023 17:34

divyegala requested a review from a team as a code owner October 20, 2023 17:34

bdice reviewed Oct 24, 2023

View reviewed changes

Comment thread cpp/benchmarks/sort/rank_lists.cpp

Comment thread cpp/benchmarks/sort/rank_structs.cpp

Comment thread cpp/include/cudf/merge.hpp Outdated

Comment thread cpp/src/merge/merge.cu Outdated

Comment thread cpp/src/merge/merge.cu Outdated

divyegala and others added 3 commits October 25, 2023 19:35

Update cpp/include/cudf/merge.hpp

6bcff40

Co-authored-by: Bradley Dice <bdice@bradleydice.com>

header includes according to dev guide

7774caf

address review

29a8428

divyegala requested a review from bdice October 26, 2023 00:18

Merge branch 'branch-23.12' into merge-new-algo

88eea0b

bdice approved these changes Oct 26, 2023

View reviewed changes

Comment thread cpp/src/merge/merge.cu Outdated

Merge branch 'branch-23.12' into merge-new-algo

611ace2

GregoryKimball mentioned this pull request Oct 26, 2023

[FEA] Implement full support for nested types #11844

Closed

no column checks in comparator functor needed

e5a6d5d

mythrocks reviewed Oct 27, 2023

View reviewed changes

Comment thread cpp/src/merge/merge.cu Outdated

mythrocks reviewed Oct 27, 2023

View reviewed changes

mythrocks approved these changes Oct 27, 2023

View reviewed changes

divyegala added 2 commits October 27, 2023 15:06

address review feedback

d84391e

Merge remote-tracking branch 'upstream/branch-23.12' into merge-new-algo

5ca5d42

rapids-bot Bot merged commit 2bc454a into NVIDIA:branch-23.12 Oct 28, 2023

bdice reviewed Oct 30, 2023

View reviewed changes

divyegala mentioned this pull request Oct 30, 2023

Fix overflow check in cudf::merge #14345

Merged

3 tasks

GregoryKimball mentioned this pull request Nov 16, 2023

[FEA] Add support for shallow lists in cudf::merge #13514

Closed

This was referenced Nov 20, 2023

Update cudf::merge to experimental row comparators #13347

Closed

[FEA] Support for non key nested types in cudf::merge #8050

Closed

This was referenced Dec 22, 2025

[FEA] Enable efficient Table.merge for nested types in out-of-core sort NVIDIA/cudf-spark#14046

Open

[FEA] Enable efficient merge sort (Table.merge) for nested types NVIDIA/cudf-spark#14059

Closed

Uh oh!

Conversation

divyegala commented Oct 4, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

divyegala commented Oct 4, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

copy-pr-bot Bot commented Oct 16, 2023

Uh oh!

divyegala commented Oct 16, 2023

Uh oh!

divyegala commented Oct 16, 2023

Uh oh!

bdice left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

divyegala commented Oct 16, 2023

Uh oh!

divyegala commented Oct 16, 2023

Uh oh!

divyegala commented Oct 17, 2023

Uh oh!

divyegala commented Oct 20, 2023

Uh oh!

bdice left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

divyegala commented Oct 26, 2023

Uh oh!

bdice left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

divyegala commented Oct 26, 2023

Uh oh!

divyegala commented Oct 26, 2023

Uh oh!

mythrocks left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

mythrocks Oct 27, 2023

Choose a reason for hiding this comment

Uh oh!

divyegala Oct 27, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

mythrocks Oct 27, 2023

Choose a reason for hiding this comment

Uh oh!

divyegala Oct 27, 2023

Choose a reason for hiding this comment

Uh oh!

mythrocks left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

divyegala commented Oct 27, 2023

Uh oh!

divyegala commented Oct 28, 2023

Uh oh!

bdice Oct 30, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

divyegala Oct 30, 2023

divyegala commented Oct 4, 2023 •

edited

Loading

divyegala commented Oct 4, 2023 •

edited

Loading

divyegala Oct 27, 2023 •

edited

Loading

mythrocks left a comment •

edited

Loading

bdice Oct 30, 2023 •

edited

Loading