Use cuda::proclaim_return_type on device lambdas #1662

hyperbolic2346 · 2023-12-19T04:16:01Z

This PR makes spark-rapids-jni compatible with Thrust 2 by adding cuda::proclaim_return_type. This is following suit with cudf, who recently did this work with rapidsai/cudf#14577

closes #1639

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

hyperbolic2346 · 2023-12-19T04:25:31Z

build

ttnghia · 2023-12-19T04:33:15Z

src/main/cpp/src/parse_uri.cu

@@ -164,11 +166,11 @@ bool __device__ validate_ipv6(string_view s)
  int address_char_count{0};
  bool address_has_hex{false};

-  auto const leading_double_colon = [&]() {
+  auto const leading_double_colon = cuda::proclaim_return_type<bool>([&]() {


This is not device lambda so don't need to change here.

Wait, I just realized that this lambda is unused. The only place that refers to it is

spark-rapids-jni/src/main/cpp/src/parse_uri.cu

Line 212 in 6784570

// if (colon_count == max_colons && !leading_double_colon) { return false; }

which is commented out.

Yeah. I wrote it all to match the RFC, but Spark/Java don't follow it exactly and I need to match them. I documented the differences by commenting out the original path. I don't know the take on this and if I should remove it or not.

If you commented out the caller then you should also comment out the function too, since it is now unused.

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

ttnghia · 2023-12-19T18:10:44Z

FYI, CCCL 2 has been merged into cudf (rapidsai/cudf#14576). You can test locally this PR with the latest cudf commit to make sure this compiles successfully.

mythrocks · 2023-12-19T19:36:30Z

For the record, we'll also need to change CMakeLists.txt, to fetch CCCL 2:

diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt
index 18c0cd1..a905297 100644
--- a/src/main/cpp/CMakeLists.txt
+++ b/src/main/cpp/CMakeLists.txt
@@ -94,11 +94,8 @@ include(cmake/Modules/ConfigureCUDA.cmake) # set other CUDA compilation flags
 # ##################################################################################################
 # * dependencies ----------------------------------------------------------------------------------
 
-# find libcu++
-include(${rapids-cmake-dir}/cpm/libcudacxx.cmake)
-
-# find thrust/cub
-include(${CUDF_DIR}/cpp/cmake/thirdparty/get_thrust.cmake)
+# find CCCL
+include(${CUDF_DIR}/cpp/cmake/thirdparty/get_cccl.cmake)
 
 # JNI
 find_package(JNI REQUIRED)

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

mythrocks · 2023-12-19T21:24:33Z

I think the changes in this PR should be good to go.

They now include the changes I had independently. (i.e. changes to CMakeLists.txt, map_utils.cu, and row_conversion.cu.)
(minimal.diff.txt, attached for reference.)

The changes in this PR are more comprehensive, if not yet strictly required.

I do see that we have trouble running tests/ROW_CONVERSION successfully with CCCL2. But that might well be a pre-existing problem.

ttnghia · 2023-12-19T21:33:55Z

Yes the issue I was trying to solve (#1567) now shows up---my previous fix doesn't work anymore. We now need to fix it some other way.

mythrocks · 2023-12-19T21:41:29Z

My conservative vote would be not to change the thirdparty/cudf submodule commit hash to go to CCCL-2, for now.

But we need to, there might be value in commenting out RowToColumnTests.ManyStrings to allow the build to run through with CCCL-2. We can then chase down the inclusive_scan problem that row_conversion.cu runs into (#1579).

abellina · 2023-12-19T23:22:07Z

My conservative vote would be not to change the thirdparty/cudf submodule commit hash to go to CCCL-2, for now.

I am confused @mythrocks, why do you want to hold?

mythrocks · 2023-12-20T00:02:05Z

I'm 👍 on this change. We might need to rebase it for the latest in parse_uri.cu.

mythrocks · 2023-12-20T00:05:30Z

I am confused @mythrocks, why do you want to hold?

@abellina: Don't mind me. I'm irrationally apprehensive about disruptive changes near the holidays.

hyperbolic2346 · 2023-12-20T00:53:29Z

I feel the same way. I don't see the value in checking in something and then going on holiday. So many are out until the new year already. If we run across some major issue, we may not be around to fix or triage it. I'd prefer to leave things running until January than check this in and have unknown issues crop up during that reduced productivity phase. I'm all for fast-following cudf, but I don't see that making this change to CCCL 2.2 now vs in 2 weeks makes a huge difference and it comes with a potential for extended disruptions.

This ship may have already sailed though since even with a pinned cudf, we are pulling in the latest rmm and our builds are broken.

…claim_return_type

hyperbolic2346 · 2023-12-20T00:57:01Z

I'm 👍 on this change. We might need to rebase it for the latest in parse_uri.cu.

Done

hyperbolic2346 · 2023-12-20T00:57:11Z

build

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

ttnghia · 2023-12-20T01:47:51Z

We need to wait until tomorrow so the changes in cudf can propagate here to build.

bdice

All the changes thus far look good to me. One comment on size_type vs. int32_t.

src/main/cpp/src/row_conversion.cu

mythrocks

LGTM, save for what @bdice suggested regarding size_type vs int32_t.

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

src/main/cpp/CMakeLists.txt

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

hyperbolic2346 · 2023-12-20T22:32:08Z

build

ttnghia · 2023-12-21T00:39:16Z

Depends on #1668.

Signed-off-by: Nghia Truong <nghiat@nvidia.com>

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

Signed-off-by: Nghia Truong <nghiat@nvidia.com>

# Conflicts: # src/main/cpp/src/row_conversion.cu

ttnghia · 2023-12-21T06:18:10Z

build

bdice

The CCCL 2.2.0 C++ and CMake changes look good to me.

This reverts commit 763406c. # Conflicts: # src/main/cpp/benchmarks/row_conversion.cpp # src/main/cpp/src/bloom_filter.cu # src/main/cpp/src/map_utils.cu # src/main/cpp/src/xxhash64.cu # thirdparty/cudf Signed-off-by: Nghia Truong <nghiat@nvidia.com>

* Revert "Use cuda::proclaim_return_type on device lambdas (#1662)" This reverts commit 763406c. # Conflicts: # src/main/cpp/benchmarks/row_conversion.cpp # src/main/cpp/src/bloom_filter.cu # src/main/cpp/src/map_utils.cu # src/main/cpp/src/xxhash64.cu # thirdparty/cudf Signed-off-by: Nghia Truong <nghiat@nvidia.com> * Remove redundant header Signed-off-by: Nghia Truong <nghiat@nvidia.com> * Update copyright year Signed-off-by: Nghia Truong <nghiat@nvidia.com> --------- Signed-off-by: Nghia Truong <nghiat@nvidia.com>

adding proclaim_return_type to device lambdas

9076ad5

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

hyperbolic2346 added the tech debt label Dec 19, 2023

hyperbolic2346 requested review from ttnghia and nvdbaranec December 19, 2023 04:16

hyperbolic2346 self-assigned this Dec 19, 2023

clang-format

6784570

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

ttnghia reviewed Dec 19, 2023

View reviewed changes

No cuda::proclaim_return_type on non-device lambda

8fb716c

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

Adding Mithun's changes for CCCL 2

44d3720

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

mythrocks mentioned this pull request Dec 19, 2023

[BUG] Build is broken, possibly with CCCL 2.2 changes upstream #1665

Closed

Merge remote-tracking branch 'upstream/branch-24.02' into mwilson/pro…

32f65e2

…claim_return_type

linting

7d12f40

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

bdice reviewed Dec 20, 2023

View reviewed changes

src/main/cpp/src/row_conversion.cu Outdated Show resolved Hide resolved

mythrocks previously approved these changes Dec 20, 2023

View reviewed changes

updating return type

4691bfb

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

hyperbolic2346 dismissed mythrocks’s stale review via 4691bfb December 20, 2023 18:40

bdice reviewed Dec 20, 2023

View reviewed changes

src/main/cpp/CMakeLists.txt Show resolved Hide resolved

Update src/main/cpp/CMakeLists.txt

11a872e

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

ttnghia and others added 5 commits December 20, 2023 16:39

Update jni

a87f3e4

Signed-off-by: Nghia Truong <nghiat@nvidia.com>

Apply suggestions from code review

3208e7d

Fix styles

8c6487b

Signed-off-by: Nghia Truong <nghiat@nvidia.com>

linting

380c218

Signed-off-by: Mike Wilson <knobby@burntsheep.com>

Update submodule manually

d8b62d7

Signed-off-by: Nghia Truong <nghiat@nvidia.com>

ttnghia mentioned this pull request Dec 21, 2023

Update jni to remove row conversion code #1668

Closed

ttnghia added 2 commits December 20, 2023 21:51

Fix header

ae3561a

Signed-off-by: Nghia Truong <nghiat@nvidia.com>

Merge branch 'remove_row_conversion' into mwilson/proclaim_return_type

348299a

# Conflicts: # src/main/cpp/src/row_conversion.cu

ttnghia self-assigned this Dec 21, 2023

mythrocks approved these changes Dec 21, 2023

View reviewed changes

bdice approved these changes Dec 21, 2023

View reviewed changes

ttnghia merged commit 763406c into NVIDIA:branch-24.02 Dec 21, 2023
3 checks passed

hyperbolic2346 deleted the mwilson/proclaim_return_type branch December 22, 2023 04:26

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Use cuda::proclaim_return_type on device lambdas #1662

Use cuda::proclaim_return_type on device lambdas #1662

hyperbolic2346 commented Dec 19, 2023 •

edited

Loading

hyperbolic2346 commented Dec 19, 2023

ttnghia Dec 19, 2023

ttnghia Dec 19, 2023

hyperbolic2346 Dec 19, 2023

ttnghia Dec 19, 2023

ttnghia commented Dec 19, 2023 •

edited

Loading

mythrocks commented Dec 19, 2023

mythrocks commented Dec 19, 2023 •

edited

Loading

ttnghia commented Dec 19, 2023

mythrocks commented Dec 19, 2023

abellina commented Dec 19, 2023

mythrocks commented Dec 20, 2023

mythrocks commented Dec 20, 2023

hyperbolic2346 commented Dec 20, 2023

hyperbolic2346 commented Dec 20, 2023

hyperbolic2346 commented Dec 20, 2023

ttnghia commented Dec 20, 2023

bdice left a comment

mythrocks left a comment

hyperbolic2346 commented Dec 20, 2023

ttnghia commented Dec 21, 2023

ttnghia commented Dec 21, 2023

bdice left a comment

Use cuda::proclaim_return_type on device lambdas #1662

Use cuda::proclaim_return_type on device lambdas #1662

Conversation

hyperbolic2346 commented Dec 19, 2023 • edited Loading

hyperbolic2346 commented Dec 19, 2023

ttnghia Dec 19, 2023

Choose a reason for hiding this comment

ttnghia Dec 19, 2023

Choose a reason for hiding this comment

hyperbolic2346 Dec 19, 2023

Choose a reason for hiding this comment

ttnghia Dec 19, 2023

Choose a reason for hiding this comment

ttnghia commented Dec 19, 2023 • edited Loading

mythrocks commented Dec 19, 2023

mythrocks commented Dec 19, 2023 • edited Loading

ttnghia commented Dec 19, 2023

mythrocks commented Dec 19, 2023

abellina commented Dec 19, 2023

mythrocks commented Dec 20, 2023

mythrocks commented Dec 20, 2023

hyperbolic2346 commented Dec 20, 2023

hyperbolic2346 commented Dec 20, 2023

hyperbolic2346 commented Dec 20, 2023

ttnghia commented Dec 20, 2023

bdice left a comment

Choose a reason for hiding this comment

mythrocks left a comment

Choose a reason for hiding this comment

hyperbolic2346 commented Dec 20, 2023

ttnghia commented Dec 21, 2023

ttnghia commented Dec 21, 2023

bdice left a comment

Choose a reason for hiding this comment

hyperbolic2346 commented Dec 19, 2023 •

edited

Loading

ttnghia commented Dec 19, 2023 •

edited

Loading

mythrocks commented Dec 19, 2023 •

edited

Loading