From 25ee915503d2947a32e89049e5a10cd2f4eb22a5 Mon Sep 17 00:00:00 2001 From: Tim Liu Date: Wed, 22 May 2024 20:56:50 +0800 Subject: [PATCH 001/159] Init version 24.08.0-SNAPSHOT Bump up JNI version to 24.08.0-SNAPSHOT Change submodule to rapidsai/cudf branch-24.08 Signed-off-by: Tim Liu --- .gitmodules | 2 +- CONTRIBUTING.md | 2 +- pom.xml | 2 +- src/main/cpp/CMakeLists.txt | 2 +- thirdparty/cudf | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/.gitmodules b/.gitmodules index 12b07c5b18..862e1ef3e6 100644 --- a/.gitmodules +++ b/.gitmodules @@ -1,4 +1,4 @@ [submodule "thirdparty/cudf"] path = thirdparty/cudf url = https://github.com/rapidsai/cudf.git - branch = branch-24.06 + branch = branch-24.08 diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 9f5c5be5c0..f4dd2940a2 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -160,7 +160,7 @@ $ ./build/build-in-docker install ... ``` Now cd to ~/repos/NVIDIA/spark-rapids and build with one of the options from -[spark-rapids instructions](https://github.com/NVIDIA/spark-rapids/blob/branch-24.06/CONTRIBUTING.md#building-from-source). +[spark-rapids instructions](https://github.com/NVIDIA/spark-rapids/blob/branch-24.08/CONTRIBUTING.md#building-from-source). ```bash $ ./build/buildall diff --git a/pom.xml b/pom.xml index 24daa4635e..9430eab648 100644 --- a/pom.xml +++ b/pom.xml @@ -21,7 +21,7 @@ com.nvidia spark-rapids-jni - 24.06.0-SNAPSHOT + 24.08.0-SNAPSHOT jar RAPIDS Accelerator JNI for Apache Spark diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt index 169067bfdd..d5d887739c 100644 --- a/src/main/cpp/CMakeLists.txt +++ b/src/main/cpp/CMakeLists.txt @@ -32,7 +32,7 @@ rapids_cuda_init_architectures(SPARK_RAPIDS_JNI) project( SPARK_RAPIDS_JNI - VERSION 24.06.00 + VERSION 24.08.00 LANGUAGES C CXX CUDA ) diff --git a/thirdparty/cudf b/thirdparty/cudf index 9a0612b3ad..968bea2a62 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 9a0612b3add9c76ea8cb45cc230b75b2474d91f7 +Subproject commit 968bea2a623acb22258f444c6be892d9fa61ea03 From db2e123146468edfa75cf42e82459364c0376b5c Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 23 May 2024 06:54:47 +0000 Subject: [PATCH 002/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 45dc595945..968bea2a62 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 45dc595945301f4076e66ec54a6e4de0b539cfb0 +Subproject commit 968bea2a623acb22258f444c6be892d9fa61ea03 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 245472ded1..df1ef9ff01 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -2f2a3c37af91d11347a8f64bc6e90197fc53eea2 +78da6e709bb893fe31587f419a2c80fa72cd66f0 From c81f57f2037ff8351dfe98cb1c6879463b12386e Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 24 May 2024 17:46:26 +0000 Subject: [PATCH 003/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 18 +++++++++--------- 3 files changed, 11 insertions(+), 11 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 8b5ff188e7..968bea2a62 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8b5ff188e79bb79ca0c2d581e94d3a91654a2d31 +Subproject commit 968bea2a623acb22258f444c6be892d9fa61ea03 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..df1ef9ff01 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +78da6e709bb893fe31587f419a2c80fa72cd66f0 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 662cbb8dc9..a0384f40f2 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -71,6 +71,14 @@ "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.06" }, + "NVTX3" : + { + "always_download" : true, + "git_shallow" : false, + "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", + "git_url" : "https://github.com/NVIDIA/NVTX.git", + "version" : "3.1.0" + }, "cuco" : { "always_download" : true, @@ -140,19 +148,11 @@ }, "version" : "3.0.6" }, - "nvtx3" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", - "git_url" : "https://github.com/NVIDIA/NVTX.git", - "version" : "3.1.0" - }, "rmm" : { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "46e153c18d17b07526d6ff2e04859fcbbd706879", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.06" }, From bac23db97b96fade260b6e9c633f283973f99e96 Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 24 May 2024 18:53:50 +0000 Subject: [PATCH 004/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 18 +++++++++--------- 3 files changed, 11 insertions(+), 11 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 8a405674a5..968bea2a62 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8a405674a5ba1554a0ced5d1f39f89fb424a768d +Subproject commit 968bea2a623acb22258f444c6be892d9fa61ea03 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..df1ef9ff01 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +78da6e709bb893fe31587f419a2c80fa72cd66f0 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 662cbb8dc9..a0384f40f2 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -71,6 +71,14 @@ "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.06" }, + "NVTX3" : + { + "always_download" : true, + "git_shallow" : false, + "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", + "git_url" : "https://github.com/NVIDIA/NVTX.git", + "version" : "3.1.0" + }, "cuco" : { "always_download" : true, @@ -140,19 +148,11 @@ }, "version" : "3.0.6" }, - "nvtx3" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", - "git_url" : "https://github.com/NVIDIA/NVTX.git", - "version" : "3.1.0" - }, "rmm" : { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "46e153c18d17b07526d6ff2e04859fcbbd706879", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.06" }, From 6b3a3027efce98f1305f23a8df0f84c9ff637c44 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 25 May 2024 05:11:20 +0800 Subject: [PATCH 005/159] [submodule-sync] bot-submodule-sync-branch-24.08 to branch-24.08 [skip ci] [bot] (#2071) * Update submodule cudf to a693e5b6f29a4c27ac845b6cccc17102b9e8ed46 Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> * Update submodule cudf to a693e5b6f29a4c27ac845b6cccc17102b9e8ed46 Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> * Update submodule cudf to 4a3315b55a89b2c92908eac8a6fd255a33843ba9 Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --------- Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 24 ++++++++++++------------ 3 files changed, 14 insertions(+), 14 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 968bea2a62..4a3315b55a 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 968bea2a623acb22258f444c6be892d9fa61ea03 +Subproject commit 4a3315b55a89b2c92908eac8a6fd255a33843ba9 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index df1ef9ff01..04673d4a09 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -78da6e709bb893fe31587f419a2c80fa72cd66f0 +9e882dd5f6dbf62407e30f3a67b208009341fcc3 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index a0384f40f2..71284874db 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -67,17 +67,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "261445be7993df57f624a3f4ee9fd15e7d26bb5e", + "git_tag" : "474ba658ff2a49cff7f7bc2e5e24be7866195aaf", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" - }, - "NVTX3" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", - "git_url" : "https://github.com/NVIDIA/NVTX.git", - "version" : "3.1.0" + "version" : "24.08" }, "cuco" : { @@ -148,13 +140,21 @@ }, "version" : "3.0.6" }, + "nvtx3" : + { + "always_download" : true, + "git_shallow" : false, + "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", + "git_url" : "https://github.com/NVIDIA/NVTX.git", + "version" : "3.1.0" + }, "rmm" : { "always_download" : true, "git_shallow" : false, - "git_tag" : "46e153c18d17b07526d6ff2e04859fcbbd706879", + "git_tag" : "f16de1f515fb7ae56bddb9999e7a35cdc3caab95", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From 01faa507176b1f81d859b35f90b40930f2e8bdae Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 25 May 2024 10:29:05 +0800 Subject: [PATCH 006/159] Update submodule cudf to 8458306ecbc17d3977a98e2e33752b678394f588 (#2077) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 4a3315b55a..8458306ecb 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 4a3315b55a89b2c92908eac8a6fd255a33843ba9 +Subproject commit 8458306ecbc17d3977a98e2e33752b678394f588 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 71284874db..4353cba0ea 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -152,7 +152,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "f16de1f515fb7ae56bddb9999e7a35cdc3caab95", + "git_tag" : "15c04666e20712f30882a903bfad79e208d8e852", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From d427ca7e86b3f6cc6775740527edcafe54275f00 Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 28 May 2024 14:35:17 +0000 Subject: [PATCH 007/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 8 ++++---- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 8a405674a5..8458306ecb 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8a405674a5ba1554a0ced5d1f39f89fb424a768d +Subproject commit 8458306ecbc17d3977a98e2e33752b678394f588 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..04673d4a09 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +9e882dd5f6dbf62407e30f3a67b208009341fcc3 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 662cbb8dc9..4353cba0ea 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -67,9 +67,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "261445be7993df57f624a3f4ee9fd15e7d26bb5e", + "git_tag" : "474ba658ff2a49cff7f7bc2e5e24be7866195aaf", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" + "version" : "24.08" }, "cuco" : { @@ -152,9 +152,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "15c04666e20712f30882a903bfad79e208d8e852", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From 8803a0edd498d64046b3199ae7c97020afb5197a Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 28 May 2024 17:01:09 +0000 Subject: [PATCH 008/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 28 ++++----------------------- 3 files changed, 6 insertions(+), 26 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 29429f7e4c..8458306ecb 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 29429f7e4c871758c0de930026347e6e3b0a5a9a +Subproject commit 8458306ecbc17d3977a98e2e33752b678394f588 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..04673d4a09 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +9e882dd5f6dbf62407e30f3a67b208009341fcc3 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 50d4c4c2f4..4353cba0ea 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -32,11 +32,6 @@ "fixed_in" : "", "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue." }, - { - "file" : "${current_json_dir}/revert_pr_211_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue." - }, { "file" : "cccl/kernel_pointer_hiding.diff", "fixed_in" : "2.4", @@ -47,30 +42,15 @@ "fixed_in" : "", "issue" : "Remove 64bit dispatching as not needed by libcudf and results in compiling twice as many kernels [https://github.com/rapidsai/cudf/pull/11437]" }, - { - "file" : "${current_json_dir}/thrust_disable_64bit_dispatching_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "Remove 64bit dispatching as not needed by libcudf and results in compiling twice as many kernels [https://github.com/rapidsai/cudf/pull/11437]" - }, { "file" : "${current_json_dir}/thrust_faster_sort_compile_times.diff", "fixed_in" : "", "issue" : "Improve Thrust sort compile times by not unrolling loops for inlined comparators [https://github.com/rapidsai/cudf/pull/10577]" }, - { - "file" : "${current_json_dir}/thrust_faster_sort_compile_times_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "Improve Thrust sort compile times by not unrolling loops for inlined comparators [https://github.com/rapidsai/cudf/pull/10577]" - }, { "file" : "${current_json_dir}/thrust_faster_scan_compile_times.diff", "fixed_in" : "", "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" - }, - { - "file" : "${current_json_dir}/thrust_faster_scan_compile_times_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" } ], "version" : "2.2.0" @@ -87,9 +67,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "1d87e5f7a19993e12726b610a9491d58748201b7", + "git_tag" : "474ba658ff2a49cff7f7bc2e5e24be7866195aaf", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" + "version" : "24.08" }, "cuco" : { @@ -172,9 +152,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "15c04666e20712f30882a903bfad79e208d8e852", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From 840938d81e8cf5dbc50a7bc4c7bdf4733aedb01c Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 29 May 2024 04:36:13 +0800 Subject: [PATCH 009/159] Update submodule cudf to 27220d6aba0728930f6e2922a7a27e9e503c2add (#2083) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 22 +++++++++++++++++++++- 2 files changed, 22 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 8458306ecb..27220d6aba 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8458306ecbc17d3977a98e2e33752b678394f588 +Subproject commit 27220d6aba0728930f6e2922a7a27e9e503c2add diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 4353cba0ea..900d93f9f8 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -32,6 +32,11 @@ "fixed_in" : "", "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue." }, + { + "file" : "${current_json_dir}/revert_pr_211_cccl_2.5.0.diff", + "fixed_in" : "", + "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue." + }, { "file" : "cccl/kernel_pointer_hiding.diff", "fixed_in" : "2.4", @@ -42,15 +47,30 @@ "fixed_in" : "", "issue" : "Remove 64bit dispatching as not needed by libcudf and results in compiling twice as many kernels [https://github.com/rapidsai/cudf/pull/11437]" }, + { + "file" : "${current_json_dir}/thrust_disable_64bit_dispatching_cccl_2.5.0.diff", + "fixed_in" : "", + "issue" : "Remove 64bit dispatching as not needed by libcudf and results in compiling twice as many kernels [https://github.com/rapidsai/cudf/pull/11437]" + }, { "file" : "${current_json_dir}/thrust_faster_sort_compile_times.diff", "fixed_in" : "", "issue" : "Improve Thrust sort compile times by not unrolling loops for inlined comparators [https://github.com/rapidsai/cudf/pull/10577]" }, + { + "file" : "${current_json_dir}/thrust_faster_sort_compile_times_cccl_2.5.0.diff", + "fixed_in" : "", + "issue" : "Improve Thrust sort compile times by not unrolling loops for inlined comparators [https://github.com/rapidsai/cudf/pull/10577]" + }, { "file" : "${current_json_dir}/thrust_faster_scan_compile_times.diff", "fixed_in" : "", "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" + }, + { + "file" : "${current_json_dir}/thrust_faster_scan_compile_times_cccl_2.5.0.diff", + "fixed_in" : "", + "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" } ], "version" : "2.2.0" @@ -67,7 +87,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "474ba658ff2a49cff7f7bc2e5e24be7866195aaf", + "git_tag" : "5e9f0c8970daa715e577ff5c078b390041027f01", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, From ffd6088179de2c01e204dc95ce97e1c8cf723df6 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 29 May 2024 16:31:14 +0800 Subject: [PATCH 010/159] Update submodule cudf to bdafa738cb7c0b4354efb22783ffd5d6edefebd6 (#2085) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 27220d6aba..bdafa738cb 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 27220d6aba0728930f6e2922a7a27e9e503c2add +Subproject commit bdafa738cb7c0b4354efb22783ffd5d6edefebd6 From d4fe5e133ff8e34043d5260b812436e8dc6112b3 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 29 May 2024 22:29:02 +0800 Subject: [PATCH 011/159] Update submodule cudf to ff981a4048a389b0e2582e94d3397a83096d16c9 (#2086) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index bdafa738cb..ff981a4048 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit bdafa738cb7c0b4354efb22783ffd5d6edefebd6 +Subproject commit ff981a4048a389b0e2582e94d3397a83096d16c9 From b3f68b6f9c72bc04cdca9737ae4664468e9d23b6 Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 29 May 2024 19:06:24 +0000 Subject: [PATCH 012/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 8 ++++---- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 29429f7e4c..ff981a4048 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 29429f7e4c871758c0de930026347e6e3b0a5a9a +Subproject commit ff981a4048a389b0e2582e94d3397a83096d16c9 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..04673d4a09 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +9e882dd5f6dbf62407e30f3a67b208009341fcc3 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 50d4c4c2f4..900d93f9f8 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,9 +87,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "1d87e5f7a19993e12726b610a9491d58748201b7", + "git_tag" : "5e9f0c8970daa715e577ff5c078b390041027f01", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" + "version" : "24.08" }, "cuco" : { @@ -172,9 +172,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "15c04666e20712f30882a903bfad79e208d8e852", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From b3433533a0e093860f8f29c4ae70e123bf0d0ec8 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 30 May 2024 05:11:29 +0800 Subject: [PATCH 013/159] Update submodule cudf to 12336da6ff3ae819635524127e65c0bfde0f3915 (#2091) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index ff981a4048..12336da6ff 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit ff981a4048a389b0e2582e94d3397a83096d16c9 +Subproject commit 12336da6ff3ae819635524127e65c0bfde0f3915 From 6f06a987e17ce7cf7c4992e51dea4e4cc180e5f8 Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 29 May 2024 22:00:54 +0000 Subject: [PATCH 014/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 8 ++++---- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 2b031e06a7..12336da6ff 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 2b031e06a7fe18eec462db445eea1c596b93a9f1 +Subproject commit 12336da6ff3ae819635524127e65c0bfde0f3915 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..04673d4a09 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +9e882dd5f6dbf62407e30f3a67b208009341fcc3 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 50d4c4c2f4..900d93f9f8 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,9 +87,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "1d87e5f7a19993e12726b610a9491d58748201b7", + "git_tag" : "5e9f0c8970daa715e577ff5c078b390041027f01", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" + "version" : "24.08" }, "cuco" : { @@ -172,9 +172,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "15c04666e20712f30882a903bfad79e208d8e852", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From ac7850ae01fcaa289308f6c95bdf85523dec17c7 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 30 May 2024 10:28:39 +0800 Subject: [PATCH 015/159] Update submodule cudf to 3a75f6db18c911d93727d12a0cf5abcdad22efda (#2094) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 12336da6ff..3a75f6db18 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 12336da6ff3ae819635524127e65c0bfde0f3915 +Subproject commit 3a75f6db18c911d93727d12a0cf5abcdad22efda From 8b621b0c98d66ca5287e154b7fbd4a0476f97873 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 30 May 2024 22:30:45 +0800 Subject: [PATCH 016/159] Update submodule cudf to 5ce95f05eeae469f4d46516b3cf6fe19902623f6 (#2097) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 3a75f6db18..5ce95f05ee 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 3a75f6db18c911d93727d12a0cf5abcdad22efda +Subproject commit 5ce95f05eeae469f4d46516b3cf6fe19902623f6 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 900d93f9f8..43549cba23 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -172,7 +172,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "15c04666e20712f30882a903bfad79e208d8e852", + "git_tag" : "805bcf1f6926c154d5ff00b0ed0e9c63f7437cc5", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From 7338e86ebdf7c7f2bcc573a7cf75672fbbf93b32 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 31 May 2024 05:24:38 +0800 Subject: [PATCH 017/159] Update submodule cudf to 789cbfdd69648fd7ec553922e64accb763ca3c57 (#2099) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 5ce95f05ee..789cbfdd69 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 5ce95f05eeae469f4d46516b3cf6fe19902623f6 +Subproject commit 789cbfdd69648fd7ec553922e64accb763ca3c57 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 43549cba23..380d866458 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,7 +87,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "5e9f0c8970daa715e577ff5c078b390041027f01", + "git_tag" : "2b3cbd067a689353bf522f9f1368fb886236b48e", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, From a5bc8c32328fe89f3079c0e62fe63063daeb73f7 Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 30 May 2024 23:09:56 +0000 Subject: [PATCH 018/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 8 ++++---- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 2b031e06a7..789cbfdd69 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 2b031e06a7fe18eec462db445eea1c596b93a9f1 +Subproject commit 789cbfdd69648fd7ec553922e64accb763ca3c57 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 20bbb44986..04673d4a09 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -41dc9623dbb8e5bdd2bccc22815efb9db6a49280 +9e882dd5f6dbf62407e30f3a67b208009341fcc3 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 50d4c4c2f4..380d866458 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,9 +87,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "1d87e5f7a19993e12726b610a9491d58748201b7", + "git_tag" : "2b3cbd067a689353bf522f9f1368fb886236b48e", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" + "version" : "24.08" }, "cuco" : { @@ -172,9 +172,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "dc1e17a03ed2dbc9329ccecc27922e414250f45a", + "git_tag" : "805bcf1f6926c154d5ff00b0ed0e9c63f7437cc5", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From bbcf5efad04a6e8313961399daab4d0ad44f3c39 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 31 May 2024 17:15:38 +0800 Subject: [PATCH 019/159] Update submodule cudf to 476db9fbb4a9969ea7406b916cead38990097fb9 (#2101) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 789cbfdd69..476db9fbb4 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 789cbfdd69648fd7ec553922e64accb763ca3c57 +Subproject commit 476db9fbb4a9969ea7406b916cead38990097fb9 From c7f388617eb0d3ff8f180a4ccc9968bf9554d0b9 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 31 May 2024 22:29:55 +0800 Subject: [PATCH 020/159] Update submodule cudf to dec0354b1ac2af981d4e8f13aceb45365838a1d8 (#2102) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 476db9fbb4..dec0354b1a 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 476db9fbb4a9969ea7406b916cead38990097fb9 +Subproject commit dec0354b1ac2af981d4e8f13aceb45365838a1d8 From ef8bf8e07eaca5438c88ed14eab80e6128fc4e3b Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 1 Jun 2024 04:31:21 +0800 Subject: [PATCH 021/159] Update submodule cudf to 7949a9cf6911066663e2245a4bb624e0f1847b06 (#2103) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index dec0354b1a..7949a9cf69 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit dec0354b1ac2af981d4e8f13aceb45365838a1d8 +Subproject commit 7949a9cf6911066663e2245a4bb624e0f1847b06 From 91f9582117506e669231487c905a611714864352 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 1 Jun 2024 10:29:38 +0800 Subject: [PATCH 022/159] Update submodule cudf to 1354abdb7a4f9eb58bfc6e359c49d0baabacb4e1 (#2104) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 7949a9cf69..1354abdb7a 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 7949a9cf6911066663e2245a4bb624e0f1847b06 +Subproject commit 1354abdb7a4f9eb58bfc6e359c49d0baabacb4e1 From 186005a32341b774203a3749ff08e97eed1f00e7 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 4 Jun 2024 05:15:51 +0800 Subject: [PATCH 023/159] Update submodule cudf to 4a17c451719a5d1e144b21703650bd323990e892 (#2106) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 1354abdb7a..4a17c45171 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 1354abdb7a4f9eb58bfc6e359c49d0baabacb4e1 +Subproject commit 4a17c451719a5d1e144b21703650bd323990e892 From 9396c92db13e525047f1295120bc6ae847bf2708 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 4 Jun 2024 11:12:10 +0800 Subject: [PATCH 024/159] Update submodule cudf to 382de32e8137a3a59a0800f46ef8a1de62b1a6e5 (#2107) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 4a17c45171..382de32e81 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 4a17c451719a5d1e144b21703650bd323990e892 +Subproject commit 382de32e8137a3a59a0800f46ef8a1de62b1a6e5 From 97619b8cbfe4f244669897a840243c85b03a7936 Mon Sep 17 00:00:00 2001 From: Peixin Date: Tue, 4 Jun 2024 11:53:26 +0800 Subject: [PATCH 025/159] Add new user to blossom-ci allow list (#2108) Signed-off-by: Peixin Li --- .github/workflows/blossom-ci.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/blossom-ci.yml b/.github/workflows/blossom-ci.yml index 33ccf50ea8..3d0f8486c0 100644 --- a/.github/workflows/blossom-ci.yml +++ b/.github/workflows/blossom-ci.yml @@ -66,6 +66,7 @@ jobs: binmahone,\ pmattione-nvidia,\ Feng-Jiang28,\ + pxLi,\ ', format('{0},', github.actor)) && github.event.comment.body == 'build' steps: - name: Check if comment is issued by authorized person From c0295023d59c7fd9c5b89db8d2b0b971e729ddd0 Mon Sep 17 00:00:00 2001 From: Peixin Date: Tue, 4 Jun 2024 12:53:25 +0800 Subject: [PATCH 026/159] Add default value for REF of premerge jenkinsfile to avoid bad overwritten (#2109) Signed-off-by: Peixin Li --- ci/Jenkinsfile.premerge | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ci/Jenkinsfile.premerge b/ci/Jenkinsfile.premerge index 0a00eb6f1b..fb7c3cd0de 100644 --- a/ci/Jenkinsfile.premerge +++ b/ci/Jenkinsfile.premerge @@ -57,7 +57,8 @@ pipeline { parameters { string(name: 'PARALLEL_LEVEL', defaultValue: '18', description: 'Parallel build cudf cpp with -DCPP_PARALLEL_LEVEL') - string(name: 'REF', defaultValue: '', + // Put a default value for REF to avoid error when running the pipeline manually + string(name: 'REF', defaultValue: 'main', description: 'Merged commit of specific PR') string(name: 'GITHUB_DATA', defaultValue: '', description: 'Json-formatted github data from upstream blossom-ci') From 5c4d9748ecf3c4f3575fabfb7462e840695dda63 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 4 Jun 2024 22:30:28 +0800 Subject: [PATCH 027/159] Update submodule cudf to eb460169786665b1624cb6c4f9b502b800810b37 (#2110) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 382de32e81..eb46016978 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 382de32e8137a3a59a0800f46ef8a1de62b1a6e5 +Subproject commit eb460169786665b1624cb6c4f9b502b800810b37 From 85b0b7dad5609c7aea837778feaf39755fb9d1cb Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 5 Jun 2024 04:35:08 +0800 Subject: [PATCH 028/159] Update submodule cudf to fe7412915a289e7a9469040ada1dcf74cda2c4d6 (#2111) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index eb46016978..fe7412915a 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit eb460169786665b1624cb6c4f9b502b800810b37 +Subproject commit fe7412915a289e7a9469040ada1dcf74cda2c4d6 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 380d866458..e4b4fcc43b 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -172,7 +172,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "805bcf1f6926c154d5ff00b0ed0e9c63f7437cc5", + "git_tag" : "8597c22a93b0950517c6c80b94f41f8e3313dd80", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From 729a485353d4dd9e880817d730d8a441d27b962f Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 5 Jun 2024 11:16:43 +0800 Subject: [PATCH 029/159] Update submodule cudf to 22ef0634f07f7b40d718e80bed176e88ac734ebe (#2112) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index fe7412915a..22ef0634f0 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit fe7412915a289e7a9469040ada1dcf74cda2c4d6 +Subproject commit 22ef0634f07f7b40d718e80bed176e88ac734ebe From 8404bedb07f14def40a7cb58860159614e681bb5 Mon Sep 17 00:00:00 2001 From: Feng Jiang <106386742+Feng-Jiang28@users.noreply.github.com> Date: Wed, 5 Jun 2024 13:11:24 +0800 Subject: [PATCH 030/159] Fix cast_string_to_float with trailing whitespaces for inf and nan string (#2063) * fixing trailing whitespaces issues when casting inf and nan string * removing debug comments * removing debug comments * remove comments Signed-off-by: fejiang * typo solved Signed-off-by: fejiang * removing comments * testcase for INIFINITY * remove cast related test cases together * shfl down logic changed * test cases for nan added * typo fixed * bug fix -nan should be treated as null --------- Signed-off-by: fejiang Co-authored-by: fejiang --- src/main/cpp/src/cast_string_to_float.cu | 39 +++++++++++---- .../spark/rapids/jni/CastStringsTest.java | 50 +++++++++++++++++++ 2 files changed, 78 insertions(+), 11 deletions(-) diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index e843d645ce..a1992be20a 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -102,7 +102,7 @@ class string_to_float { int sign = check_for_sign(); // check for leading nan - if (check_for_nan()) { + if (check_for_nan(sign)) { _out[_row] = NAN; compute_validity(_valid, _except); return; @@ -236,22 +236,31 @@ class string_to_float { // returns true if we encountered 'nan' // potentially changes: valid/except - __device__ bool check_for_nan() + __device__ bool check_for_nan(int const& sign) { auto const nan_mask = __ballot_sync(0xffffffff, (_warp_lane == 0 && (_c == 'N' || _c == 'n')) || (_warp_lane == 1 && (_c == 'A' || _c == 'a')) || (_warp_lane == 2 && (_c == 'N' || _c == 'n'))); if (nan_mask == 0x7) { - // if we start with 'nan', then even if we have other garbage character, this is a null row. - // - // if we're in ansi mode and this is not -precisely- nan, report that so that we can throw - // an exception later. - if (_len != 3) { - _valid = false; - _except = _len != 3; - } - return true; + // if we start with 'nan', then even if we have other garbage character(excluding + // whitespaces), this is a null row. but for e.g. : "nan " cases. spark will treat the as + // "nan", when the trailing characters are whitespaces, it is still a valid string. if we're + // in ansi mode and this is not -precisely- nan, report that so that we can throw an exception + // later. + + // move forward the current position by 3 + _bpos += 3; + _c = __shfl_down_sync(0xffffffff, _c, 3); + + // remove the trailing whitespaces, if there exits + remove_leading_whitespace(); + + // if we're at the end and sign is not '-', because Spark treats '-nan' as null + if (_bpos == _len && sign != -1) { return true; } + // if we reach out here, it means that we have other garbage character. + _valid = false; + _except = true; } return false; } @@ -299,11 +308,19 @@ class string_to_float { _bpos += 5; // if we're at the end if (_bpos == _len) { return true; } + _c = __shfl_down_sync(0xffffffff, _c, 5); } + // remove the remaining whitespace if exists + remove_leading_whitespace(); + + // if we're at the end + if (_bpos == _len) { return true; } + // if we reach here for any reason, it means we have "inf" or "infinity" at the start of the // string but also have additional characters, making this whole thing bogus/null _valid = false; + return true; } return false; diff --git a/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java b/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java index a9b1cfaa4d..254f8329e6 100644 --- a/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java +++ b/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java @@ -158,6 +158,56 @@ void castToFloatsTrimTest() { } } + @Test + void castToFloatNanTest(){ + Table.TestBuilder tb2 = new Table.TestBuilder(); + tb2.column("nan", "nan ", " nan ", "NAN", "nAn ", " NAn ", "Nan 0", "nan nan"); + + Table.TestBuilder tb = new Table.TestBuilder(); + tb.column(Float.NaN, Float.NaN, Float.NaN, Float.NaN, Float.NaN, Float.NaN, null, null); + + try (Table expected = tb.build()) { + List result = new ArrayList<>(); + try (Table origTable = tb2.build()) { + for (int i = 0; i < origTable.getNumberOfColumns(); i++) { + ColumnVector string_col = origTable.getColumn(i); + result.add(CastStrings.toFloat(string_col, false, expected.getColumn(i).getType())); + } + try (Table result_tbl = new Table(result.toArray(new ColumnVector[result.size()]))) { + AssertUtils.assertTablesAreEqual(expected, result_tbl); + } + } finally { + result.forEach(ColumnVector::close); + } + } + } + + @Test + void castToFloatsInfTest(){ + // The test data: Table.TestBuilder object with a column containing the string "inf" + Table.TestBuilder tb2 = new Table.TestBuilder(); + tb2.column("INFINITY ", "inf", "+inf ", " -INF ", "INFINITY AND BEYOND", "INF"); + + Table.TestBuilder tb = new Table.TestBuilder(); + tb.column(Float.POSITIVE_INFINITY, Float.POSITIVE_INFINITY, Float.POSITIVE_INFINITY, Float.NEGATIVE_INFINITY, null, Float.POSITIVE_INFINITY); + + try (Table expected = tb.build()) { + List result = new ArrayList<>(); + try (Table origTable = tb2.build()) { + for (int i = 0; i < origTable.getNumberOfColumns(); i++) { + ColumnVector string_col = origTable.getColumn(i); + result.add(CastStrings.toFloat(string_col, false, expected.getColumn(i).getType())); + } + System.out.println(result); + try (Table result_tbl = new Table(result.toArray(new ColumnVector[result.size()]))) { + AssertUtils.assertTablesAreEqual(expected, result_tbl); + } + } finally { + result.forEach(ColumnVector::close); + } + } + } + @Test void castToDecimalTest() { Table.TestBuilder tb = new Table.TestBuilder(); From 128dd8c00d06ce1559499e3416b540840f10f569 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 6 Jun 2024 05:18:07 +0800 Subject: [PATCH 031/159] Update submodule cudf to 20aa4442d27ca858796c7890ad0542dbaee542e1 (#2116) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 4 ++-- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 22ef0634f0..20aa4442d2 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 22ef0634f07f7b40d718e80bed176e88ac734ebe +Subproject commit 20aa4442d27ca858796c7890ad0542dbaee542e1 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 04673d4a09..a074172404 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -9e882dd5f6dbf62407e30f3a67b208009341fcc3 +d65526d7503aa4180b89f04b080e79c23c3a3ceb diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index e4b4fcc43b..9cbf2b0b97 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,7 +87,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "2b3cbd067a689353bf522f9f1368fb886236b48e", + "git_tag" : "43ce81964f3fb28dae95c3867f1905f26bc84e31", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, @@ -172,7 +172,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "8597c22a93b0950517c6c80b94f41f8e3313dd80", + "git_tag" : "203a3c099c2047436a7bc3865d3ab7baab61f7aa", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From 3ed476cc109fc1a66d1e31c9faf3c57427a57aea Mon Sep 17 00:00:00 2001 From: spark-rapids automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 6 Jun 2024 02:39:51 +0000 Subject: [PATCH 032/159] Auto-merge use branch-24.08 versions Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 8 ++++---- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 7c706cc400..20aa4442d2 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 7c706cc4004d5feaae92544b3b29a00c64f7ed86 +Subproject commit 20aa4442d27ca858796c7890ad0542dbaee542e1 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 356898bd1d..a074172404 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -3c754156fbe1b8ba0b6848d3e2cc1359d0d8918d +d65526d7503aa4180b89f04b080e79c23c3a3ceb diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 49c41c7c53..9cbf2b0b97 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,9 +87,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "5d4f4a8565c5cdff94f77832a775d73ec9e7513e", + "git_tag" : "43ce81964f3fb28dae95c3867f1905f26bc84e31", "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.06" + "version" : "24.08" }, "cuco" : { @@ -172,9 +172,9 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "f47ce3f0d46848cd9d5844d499bf150dd14d823a", + "git_tag" : "203a3c099c2047436a7bc3865d3ab7baab61f7aa", "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.06" + "version" : "24.08" }, "spdlog" : { From 5119cf68f3d2c3bca258a1cddba82236548f4f1c Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 6 Jun 2024 11:13:09 +0800 Subject: [PATCH 033/159] Update submodule cudf to d91380ef393e9156c34a078998041a6affca7923 (#2119) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 20aa4442d2..d91380ef39 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 20aa4442d27ca858796c7890ad0542dbaee542e1 +Subproject commit d91380ef393e9156c34a078998041a6affca7923 From acf1f29760ae65a7ba3270782790908c3868d0a6 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 6 Jun 2024 16:36:03 +0800 Subject: [PATCH 034/159] Update submodule cudf to 7fd6918f9f4bbfc499bc60a3532a464c357da4f4 (#2120) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index d91380ef39..7fd6918f9f 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit d91380ef393e9156c34a078998041a6affca7923 +Subproject commit 7fd6918f9f4bbfc499bc60a3532a464c357da4f4 From 85f132e0ed32c9ba75215ccfe8fe0d49ed223ada Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 6 Jun 2024 22:31:58 +0800 Subject: [PATCH 035/159] Update submodule cudf to 3b734ec2fd591f037fe1d8f8ce424c7049cb5a3e (#2122) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 7fd6918f9f..3b734ec2fd 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 7fd6918f9f4bbfc499bc60a3532a464c357da4f4 +Subproject commit 3b734ec2fd591f037fe1d8f8ce424c7049cb5a3e diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 9cbf2b0b97..9174f2f273 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -172,7 +172,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "203a3c099c2047436a7bc3865d3ab7baab61f7aa", + "git_tag" : "6902af9e17f3e9c41d23bd3b5855b2f4f51c83d6", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From 88c03c98b1f44aac9f3e3ff7d86d752661c51cb6 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 7 Jun 2024 05:16:09 +0800 Subject: [PATCH 036/159] Update submodule cudf to 3468fa1f5b9dfcf83a95bcb09fe5a4d8d3808620 (#2123) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 3b734ec2fd..3468fa1f5b 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 3b734ec2fd591f037fe1d8f8ce424c7049cb5a3e +Subproject commit 3468fa1f5b9dfcf83a95bcb09fe5a4d8d3808620 From acc5eba9fab7635e65fba5d49fae714244681a94 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 7 Jun 2024 10:31:54 +0800 Subject: [PATCH 037/159] Update submodule cudf to 451d12a2d8d69f63d2b9491286b8895ace6f87ba (#2126) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 3468fa1f5b..451d12a2d8 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 3468fa1f5b9dfcf83a95bcb09fe5a4d8d3808620 +Subproject commit 451d12a2d8d69f63d2b9491286b8895ace6f87ba From 9436ada1ad3f72a30b0c466ce14b600510e00ffb Mon Sep 17 00:00:00 2001 From: Feng Jiang <106386742+Feng-Jiang28@users.noreply.github.com> Date: Fri, 7 Jun 2024 14:27:51 +0800 Subject: [PATCH 038/159] [WIP] fix string to float using hand-picked values failed in CastOpSuite (#2121) * add logic for dealing with nan with sign * sign logic fixed * reformat Signed-off-by: fejiang --------- Signed-off-by: fejiang --- src/main/cpp/src/cast_string_to_float.cu | 24 ++++++++++++------- .../spark/rapids/jni/CastStringsTest.java | 4 ++-- 2 files changed, 17 insertions(+), 11 deletions(-) diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index a1992be20a..cd7de88110 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -112,7 +112,7 @@ class string_to_float { if (check_for_inf()) { if (_warp_lane == 0) { _out[_row] = - sign > 0 ? std::numeric_limits::infinity() : -std::numeric_limits::infinity(); + sign >= 0 ? std::numeric_limits::infinity() : -std::numeric_limits::infinity(); } compute_validity(_valid, _except); return; @@ -140,7 +140,9 @@ class string_to_float { _except = true; } - if (_warp_lane == 0) { _out[_row] = sign * static_cast(0); } + if (_warp_lane == 0) { + _out[_row] = sign >= 0 ? static_cast(0) : -static_cast(0); + } compute_validity(_valid, _except); return; } @@ -154,15 +156,15 @@ class string_to_float { // construct the final float value if (_warp_lane == 0) { // base value - double digitsf = sign * static_cast(digits); + double digitsf = sign >= 0 ? static_cast(digits) : -static_cast(digits); // exponent int exp_ten = exp_base + manual_exp; // final value if (exp_ten > std::numeric_limits::max_exponent10) { - _out[_row] = sign > 0 ? std::numeric_limits::infinity() - : -std::numeric_limits::infinity(); + _out[_row] = sign >= 0 ? std::numeric_limits::infinity() + : -std::numeric_limits::infinity(); } else { // make sure we don't produce a subnormal number. // - a normal number is one where the leading digit of the floating point rep is not zero. @@ -256,8 +258,8 @@ class string_to_float { // remove the trailing whitespaces, if there exits remove_leading_whitespace(); - // if we're at the end and sign is not '-', because Spark treats '-nan' as null - if (_bpos == _len && sign != -1) { return true; } + // if we're at the end and there is no sign, because Spark treats '-nan' and '+nan' as null. + if (_bpos == _len && sign == 0) { return true; } // if we reach out here, it means that we have other garbage character. _valid = false; _except = true; @@ -265,12 +267,16 @@ class string_to_float { return false; } - // returns 1 or -1 to indicate sign + // The `sign` variables is initialized to 0, indicating no sign. + // If a sign is detected, it sets `sign` to 1, indicating `+` sign. + // If `-` is then detected, it sets `sign` to -1. + // returns 1, 0, -1 to indicate signs. __device__ int check_for_sign() { auto const sign_mask = __ballot_sync(0xffffffff, _warp_lane == 0 && (_c == '+' || _c == '-')); - int sign = 1; + int sign = 0; if (sign_mask) { + sign = 1; // NOTE: warp lane 0 is the only thread that ever reads `sign`, so technically it would be // valid to just check if(c == '-'), but that would leave other threads with an incorrect // value. if this code ever changes, that could lead to hard-to-find bugs. diff --git a/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java b/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java index 254f8329e6..f784736819 100644 --- a/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java +++ b/src/test/java/com/nvidia/spark/rapids/jni/CastStringsTest.java @@ -161,10 +161,10 @@ void castToFloatsTrimTest() { @Test void castToFloatNanTest(){ Table.TestBuilder tb2 = new Table.TestBuilder(); - tb2.column("nan", "nan ", " nan ", "NAN", "nAn ", " NAn ", "Nan 0", "nan nan"); + tb2.column("nan", "nan ", " nan ", "NAN", "nAn ", " NAn ", "Nan 0", "+naN", "-nAn"); Table.TestBuilder tb = new Table.TestBuilder(); - tb.column(Float.NaN, Float.NaN, Float.NaN, Float.NaN, Float.NaN, Float.NaN, null, null); + tb.column(Float.NaN, Float.NaN, Float.NaN, Float.NaN, Float.NaN, Float.NaN, null, null, null); try (Table expected = tb.build()) { List result = new ArrayList<>(); From 36c0b7d3bbd2991a9058373b2016523bbc3d1745 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 7 Jun 2024 16:29:55 +0800 Subject: [PATCH 039/159] Update submodule cudf to 9bd16bb719e14ed1e0ee3edbd8c8417c03ac2f25 (#2127) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 451d12a2d8..9bd16bb719 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 451d12a2d8d69f63d2b9491286b8895ace6f87ba +Subproject commit 9bd16bb719e14ed1e0ee3edbd8c8417c03ac2f25 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 9174f2f273..0307699645 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -172,7 +172,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "6902af9e17f3e9c41d23bd3b5855b2f4f51c83d6", + "git_tag" : "49fcfeb771098c38b4fb83e969b2ae63f5dc66ab", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From f2b7d47d916cd46f723f86d74452a7fb37683f46 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 8 Jun 2024 04:31:45 +0800 Subject: [PATCH 040/159] Update submodule cudf to 39c5b86645dc61bf0c59d7bf733ca13872b46a44 (#2128) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 6 +++--- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 9bd16bb719..39c5b86645 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 9bd16bb719e14ed1e0ee3edbd8c8417c03ac2f25 +Subproject commit 39c5b86645dc61bf0c59d7bf733ca13872b46a44 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index a074172404..eeda3a55c1 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -d65526d7503aa4180b89f04b080e79c23c3a3ceb +e42a5c9c70092b7f379a4525acfd5ef4524c152e diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 0307699645..f3ab60c936 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -13,7 +13,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "36f379f29660761fe033a1306ca9dab6a88cb65c", + "git_tag" : "fde1cf79bde6744b6739636f502b5edfefe302e1", "git_url" : "https://github.com/NVIDIA/cccl.git", "patches" : [ @@ -73,7 +73,7 @@ "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" } ], - "version" : "2.2.0" + "version" : "2.5.0" }, "GTest" : { @@ -95,7 +95,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "2101cb31d0210b609cd02c88f9b538e10881d91d", + "git_tag" : "6923b3b1e526b3305de81f784fdca26736caf5ba", "git_url" : "https://github.com/NVIDIA/cuCollections.git", "version" : "0.0.1" }, From 81f63ad09b36dfeefe62c8e9430ce9dd67477f5b Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 8 Jun 2024 10:30:33 +0800 Subject: [PATCH 041/159] Update submodule cudf to 8e40fe7e6b01a399c3ea406a59d4cbcbc9bfce5c (#2131) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 39c5b86645..8e40fe7e6b 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 39c5b86645dc61bf0c59d7bf733ca13872b46a44 +Subproject commit 8e40fe7e6b01a399c3ea406a59d4cbcbc9bfce5c diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index eeda3a55c1..768177632f 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -e42a5c9c70092b7f379a4525acfd5ef4524c152e +5b0570870c86dd5f53ea00f0b46dd9e068f3f9f5 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index f3ab60c936..2f9e2e8d6e 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,7 +87,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "43ce81964f3fb28dae95c3867f1905f26bc84e31", + "git_tag" : "33e333b7409756926f2c9bba949d3aa70c9516ea", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, From e67cdd08113073cc9ca965ba83c59de6a86ea3bb Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Mon, 10 Jun 2024 16:28:38 +0800 Subject: [PATCH 042/159] Update submodule cudf to c02260f2fb1c162eabf0da0604cc6f08f2cc74ff (#2132) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 8e40fe7e6b..c02260f2fb 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8e40fe7e6b01a399c3ea406a59d4cbcbc9bfce5c +Subproject commit c02260f2fb1c162eabf0da0604cc6f08f2cc74ff From 302cfbe1176e2b5d1a1054ed80ff2b435b1ef35d Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 11 Jun 2024 04:32:10 +0800 Subject: [PATCH 043/159] Update submodule cudf to 58a15a84078c42b331ced4fd4384724d42328258 (#2134) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index c02260f2fb..58a15a8407 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit c02260f2fb1c162eabf0da0604cc6f08f2cc74ff +Subproject commit 58a15a84078c42b331ced4fd4384724d42328258 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 2f9e2e8d6e..f5dd484e59 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,7 +87,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "33e333b7409756926f2c9bba949d3aa70c9516ea", + "git_tag" : "64b51293874388852a9513a6996bba07c1523d26", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, @@ -172,7 +172,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "49fcfeb771098c38b4fb83e969b2ae63f5dc66ab", + "git_tag" : "dcca09818f1c5b5b1936f684b0ff9cca5683e63f", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From e56621332106852555e5aed26d65bcd5709f6274 Mon Sep 17 00:00:00 2001 From: Peixin Date: Tue, 11 Jun 2024 08:42:57 +0800 Subject: [PATCH 044/159] Update blossom-ci ACL to secure format (#2137) Signed-off-by: Peixin Li --- .github/workflows/blossom-ci.yml | 72 ++++++++++++++++---------------- 1 file changed, 37 insertions(+), 35 deletions(-) diff --git a/.github/workflows/blossom-ci.yml b/.github/workflows/blossom-ci.yml index 3d0f8486c0..686c655ca1 100644 --- a/.github/workflows/blossom-ci.yml +++ b/.github/workflows/blossom-ci.yml @@ -33,41 +33,43 @@ jobs: args: ${{ env.args }} # This job only runs for pull request comments - if: contains( '\ - abellina,\ - anfeng,\ - firestarman,\ - GaryShen2008,\ - jlowe,\ - mythrocks,\ - nartal1,\ - nvdbaranec,\ - NvTimLiu,\ - razajafri,\ - revans2,\ - rwlee,\ - sameerz,\ - tgravescs,\ - wbo4958,\ - wjxiz1992,\ - sperlingxx,\ - YanxuanLiu,\ - hyperbolic2346,\ - gerashegalov,\ - ttnghia,\ - nvliyuan,\ - res-life,\ - HaoYang670,\ - NVnavkumar,\ - yinqingh,\ - thirtiseven,\ - parthosa,\ - liurenjie1024,\ - binmahone,\ - pmattione-nvidia,\ - Feng-Jiang28,\ - pxLi,\ - ', format('{0},', github.actor)) && github.event.comment.body == 'build' + if: | + github.event.comment.body == 'build' && + ( + github.actor == 'abellina' || + github.actor == 'anfeng' || + github.actor == 'firestarman' || + github.actor == 'GaryShen2008' || + github.actor == 'jlowe' || + github.actor == 'mythrocks' || + github.actor == 'nartal1' || + github.actor == 'nvdbaranec' || + github.actor == 'NvTimLiu' || + github.actor == 'razajafri' || + github.actor == 'revans2' || + github.actor == 'rwlee' || + github.actor == 'sameerz' || + github.actor == 'tgravescs' || + github.actor == 'wbo4958' || + github.actor == 'wjxiz1992' || + github.actor == 'sperlingxx' || + github.actor == 'YanxuanLiu' || + github.actor == 'hyperbolic2346' || + github.actor == 'gerashegalov' || + github.actor == 'ttnghia' || + github.actor == 'nvliyuan' || + github.actor == 'res-life' || + github.actor == 'HaoYang670' || + github.actor == 'NVnavkumar' || + github.actor == 'yinqingh' || + github.actor == 'thirtiseven' || + github.actor == 'parthosa' || + github.actor == 'liurenjie1024' || + github.actor == 'binmahone' || + github.actor == 'pmattione-nvidia' || + github.actor == 'Feng-Jiang28' || + github.actor == 'pxLi' + ) steps: - name: Check if comment is issued by authorized person run: blossom-ci From 175266fb6b6b29ac69ae28b79058422220b11960 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 11 Jun 2024 16:31:38 +0800 Subject: [PATCH 045/159] [submodule-sync] bot-submodule-sync-branch-24.08 to branch-24.08 [skip ci] [bot] (#2139) * Update submodule cudf to 1bd210d76ab05c669aea230b9287b76a03328efa Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> * Update submodule cudf to ff1e4bb82ce4ab8ac54bc8715bf761a3700024bc Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --------- Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 58a15a8407..ff1e4bb82c 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 58a15a84078c42b331ced4fd4384724d42328258 +Subproject commit ff1e4bb82ce4ab8ac54bc8715bf761a3700024bc diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index f5dd484e59..1fa2686baa 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -87,7 +87,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "64b51293874388852a9513a6996bba07c1523d26", + "git_tag" : "88033f4182480a3244322c9bd680f92520f0d17b", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, From 27a1d57c626407fb0cb101e4f35b26b3b8daa8b0 Mon Sep 17 00:00:00 2001 From: Gera Shegalov Date: Tue, 11 Jun 2024 08:18:50 -0700 Subject: [PATCH 046/159] Pin rapids-cmake across all cmake invocations (#2136) Fixes #2135 Signed-off-by: Gera Shegalov --- pom.xml | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/pom.xml b/pom.xml index 1627ea94d1..2e91179a2c 100644 --- a/pom.xml +++ b/pom.xml @@ -429,6 +429,8 @@ + + + + Date: Wed, 12 Jun 2024 10:30:50 +0800 Subject: [PATCH 047/159] Update submodule cudf to f7ba6ab47ac994e6a1363119c01eee5dd6304181 (#2140) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 111 +---------------------------- 2 files changed, 2 insertions(+), 111 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index ff1e4bb82c..f7ba6ab47a 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit ff1e4bb82ce4ab8ac54bc8715bf761a3700024bc +Subproject commit f7ba6ab47ac994e6a1363119c01eee5dd6304181 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 1fa2686baa..ca0af67cac 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -1,14 +1,6 @@ { "packages" : { - "Arrow" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "7dd1d34074af176d9e861a360e135ae57b21cf96", - "git_url" : "https://github.com/apache/arrow.git", - "version" : "16.1.0" - }, "CCCL" : { "always_download" : true, @@ -83,114 +75,13 @@ "git_url" : "https://github.com/google/googletest.git", "version" : "1.13.0" }, - "KvikIO" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "88033f4182480a3244322c9bd680f92520f0d17b", - "git_url" : "https://github.com/rapidsai/kvikio.git", - "version" : "24.08" - }, - "cuco" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "6923b3b1e526b3305de81f784fdca26736caf5ba", - "git_url" : "https://github.com/NVIDIA/cuCollections.git", - "version" : "0.0.1" - }, - "dlpack" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "365b823cedb281cd0240ca601aba9b78771f91a3", - "git_url" : "https://github.com/dmlc/dlpack.git", - "version" : "0.8" - }, - "flatbuffers" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "595bf0007ab1929570c7671f091313c8fc20644e", - "git_url" : "https://github.com/google/flatbuffers.git", - "version" : "24.3.25" - }, - "fmt" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "f5e54359df4c26b6230fc61d38aa294581393084", - "git_url" : "https://github.com/fmtlib/fmt.git", - "patches" : - [ - { - "file" : "fmt/fix_10_1_1_version.diff", - "fixed_in" : "10.2.0", - "issue" : "fmt 10.1.1 produces a CMake package with version 10.1.0" - } - ], - "version" : "10.1.1" - }, - "jitify" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "e38b993f4cb3207745735c51d4f61cdaa735b7ac", - "git_url" : "https://github.com/rapidsai/jitify.git", - "version" : "2.0.0" - }, - "nanoarrow" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "11e73a8c85b45e3d49c8c541b4e1497a649fe03c", - "git_url" : "https://github.com/apache/arrow-nanoarrow.git", - "version" : "0.5.0" - }, - "nvcomp" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "v2.2.0", - "git_url" : "https://github.com/NVIDIA/nvcomp.git", - "proprietary_binary" : - { - "aarch64-linux" : "https://developer.download.nvidia.com/compute/nvcomp/${version}/local_installers/nvcomp_${version}_SBSA_${cuda-toolkit-version-major}.x.tgz", - "x86_64-linux" : "https://developer.download.nvidia.com/compute/nvcomp/${version}/local_installers/nvcomp_${version}_x86_64_${cuda-toolkit-version-major}.x.tgz" - }, - "version" : "3.0.6" - }, "nvtx3" : { "always_download" : true, "git_shallow" : false, - "git_tag" : "e170594ac7cf1dac584da473d4ca9301087090c1", + "git_tag" : "v${version}", "git_url" : "https://github.com/NVIDIA/NVTX.git", "version" : "3.1.0" - }, - "rmm" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "dcca09818f1c5b5b1936f684b0ff9cca5683e63f", - "git_url" : "https://github.com/rapidsai/rmm.git", - "version" : "24.08" - }, - "spdlog" : - { - "always_download" : true, - "git_shallow" : false, - "git_tag" : "7e635fca68d014934b4af8a1cf874f63989352b7", - "git_url" : "https://github.com/gabime/spdlog.git", - "patches" : - [ - { - "file" : "spdlog/nvcc_constexpr_fix.diff", - "fixed_in" : "1.13", - "issue" : "Fix constexpr mismatch between spdlog and fmt [https://github.com/gabime/spdlog/issues/2856]" - } - ], - "version" : "1.12.0" } } } \ No newline at end of file From 330d5bf6a79bbe1473ffa518a85e1d434e55d821 Mon Sep 17 00:00:00 2001 From: Gera Shegalov Date: Tue, 11 Jun 2024 23:06:44 -0700 Subject: [PATCH 048/159] Inline build-libcudf and simplify overall build (#2115) Inline build-libcudf and simplify overall build #2115 Additional changes - Deprecate libcudf.clean.skip. Default it to `false` - Move libcudf build to target/libcudf/cmake-build - Move jni build to target/jni/cmake-build - Reduce relative path hard-coding - Introduce BUILD_CUDF_TESTS/BENCHMARKS - Skip native tests when -DskipTests is passed Signed-off-by: Gera Shegalov --- CONTRIBUTING.md | 7 ---- build-libcudf.xml | 65 ---------------------------------- build/run-in-docker | 1 - pom.xml | 70 ++++++++++++++++++++++++++++--------- src/main/cpp/CMakeLists.txt | 43 +++++++++++++---------- 5 files changed, 79 insertions(+), 107 deletions(-) delete mode 100644 build-libcudf.xml diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index ae7a773610..e30f2da15d 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -64,12 +64,6 @@ so it should be invoked as one would invoke Maven, e.g.: `build/build-in-docker ### cudf Submodule and Build [RAPIDS cuDF](https://github.com/rapidsai/cudf) is being used as a submodule in this project. -Due to the lengthy build of libcudf, it is **not cleaned** during a normal Maven clean phase -unless built using `build/build-in-docker`. `build/build-in-docker` uses `ccache` by default -unless CCACHE_DISABLE=1 is set in the environment. - -`-Dlibcudf.clean.skip=false` can also be specified on the Maven command-line to force -libcudf to be cleaned during the Maven clean phase. Currently libcudf is only configured once and the build relies on cmake to re-configure as needed. This is because libcudf currently is rebuilding almost entirely when it is configured with the same @@ -93,7 +87,6 @@ to control aspects of the build: | `BUILD_BENCHMARKS` | Compile benchmarks | OFF | | `BUILD_FAULTINJ` | Compile fault injection | ON | | `libcudf.build.configure` | Force libcudf build to configure | false | -| `libcudf.clean.skip` | Whether to skip cleaning libcudf build | true | | `submodule.check.skip` | Whether to skip checking git submodules | false | diff --git a/build-libcudf.xml b/build-libcudf.xml deleted file mode 100644 index 765f50e8a0..0000000000 --- a/build-libcudf.xml +++ /dev/null @@ -1,65 +0,0 @@ - - - - - Configures and builds the libcudf library from the cudf submodule. - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - diff --git a/build/run-in-docker b/build/run-in-docker index 81152a1d9d..51e8a7af27 100755 --- a/build/run-in-docker +++ b/build/run-in-docker @@ -67,7 +67,6 @@ $DOCKER_CMD run $DOCKER_GPU_OPTS $DOCKER_RUN_EXTRA_ARGS -u $(id -u):$(id -g) --r -v "$LOCAL_CCACHE_DIR:$LOCAL_CCACHE_DIR:rw" \ -v "$LOCAL_MAVEN_REPO:$LOCAL_MAVEN_REPO:rw" \ --workdir "$REPODIR" \ - -e CCACHE_DISABLE \ -e CCACHE_DIR="$LOCAL_CCACHE_DIR" \ -e CMAKE_C_COMPILER_LAUNCHER="ccache" \ -e CMAKE_CXX_COMPILER_LAUNCHER="ccache" \ diff --git a/pom.xml b/pom.xml index 2e91179a2c..a1a22b3189 100644 --- a/pom.xml +++ b/pom.xml @@ -81,8 +81,11 @@ OFF OFF + Release OFF OFF + OFF + OFF ON ON false @@ -93,16 +96,16 @@ ${project.basedir}/thirdparty/cudf-pins/ 3.2.4 5.8.1 - ${cudf.path}/cpp/build + ${project.build.directory}/libcudf/cmake-build/ false - true + false ${project.build.directory}/libcudf-install pinned ${project.build.directory}/libcudfjni 1.8 1.8 2.25.0 - ${project.build.directory}/cmake-build + ${project.build.directory}/jni/cmake-build 1.10.0 UTF-8 1.7.30 @@ -322,11 +325,12 @@ failonerror="true" executable="cmake"> - + + ${skipTests} run @@ -392,18 +396,48 @@ build-libcudf validate - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + @@ -454,6 +488,10 @@ + + + + diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt index 8b165c300b..c78e9abc0f 100644 --- a/src/main/cpp/CMakeLists.txt +++ b/src/main/cpp/CMakeLists.txt @@ -14,7 +14,24 @@ cmake_minimum_required(VERSION 3.20.1 FATAL_ERROR) -include(../../../thirdparty/cudf/rapids_config.cmake) +set(CUDF_DIR + "$ENV{CUDF_ROOT}" + CACHE STRING "path to cudf repository" +) +set(CUDF_INSTALL_DIR + "$ENV{CUDF_INSTALL_DIR}" + CACHE STRING "path to libcudf install root" +) +set(CUDFJNI_BUILD_DIR + "$ENV{CUDFJNI_BUILD_DIR}" + CACHE STRING "path to libcudfjni build root" +) +set(CUDF_CPP_BUILD_DIR + "$ENV{CUDF_CPP_BUILD_DIR}" + CACHE STRING "path to libcudf build root" +) + +include("${CUDF_DIR}/rapids_config.cmake") include(rapids-cmake) include(rapids-cpm) include(rapids-cuda) @@ -78,18 +95,8 @@ set_property( ) message(VERBOSE "SPARK_RAPIDS_JNI: RMM_LOGGING_LEVEL = '${RMM_LOGGING_LEVEL}'.") -set(CUDF_INSTALL_DIR - "${PROJECT_BINARY_DIR}/../libcudf-install" - CACHE STRING "path to libcudf install root" -) - set(CMAKE_PREFIX_PATH ${CUDF_INSTALL_DIR}) -set(CUDF_DIR - "${PROJECT_BINARY_DIR}/../../thirdparty/cudf" - CACHE STRING "path to cudf repository" -) - # Set a default build type if none was specified rapids_cmake_build_type("Release") @@ -141,31 +148,31 @@ endif() # cudfjni find_library(CUDFJNI_LIB "libcudfjni.a" REQUIRED NO_DEFAULT_PATH - HINTS "${PROJECT_BINARY_DIR}/../libcudfjni" + HINTS "${CUDFJNI_BUILD_DIR}" ) # parquet find_library(PARQUET_LIB "libparquet.a" REQUIRED NO_DEFAULT_PATH - HINTS "${PROJECT_BINARY_DIR}/../libcudf-install/lib64" - HINTS "${PROJECT_BINARY_DIR}/../libcudf-install/lib" + HINTS "${CUDF_INSTALL_DIR}/lib64" + HINTS "${CUDF_INSTALL_DIR}/lib" ) # Internal parquet headers set (GENERATED_PARQUET_INCLUDE - "${CUDF_DIR}/cpp/build/_deps/arrow-src/cpp/src/" + "${CUDF_CPP_BUILD_DIR}/_deps/arrow-src/cpp/src/" CACHE STRING "generated parquet thrift headers" ) # thrift find_library(THRIFT_LIB "libthrift.a" REQUIRED NO_DEFAULT_PATH - HINTS "${CUDF_DIR}/cpp/build/_deps/arrow-build/thrift_ep-install/lib/" + HINTS "${CUDF_CPP_BUILD_DIR}/_deps/arrow-build/thrift_ep-install/lib/" ) set(CUDFJNI_INCLUDE_DIRS "${CUDF_DIR}/java/src/main/native/include" "${CUDF_DIR}/java/src/main/native/src" "${GENERATED_PARQUET_INCLUDE}" - "${CUDF_DIR}/cpp/build/_deps/arrow-build/thrift_ep-install/include/" + "${CUDF_CPP_BUILD_DIR}/_deps/arrow-build/thrift_ep-install/include/" ) # ################################################################################################## @@ -278,7 +285,7 @@ add_dependencies(cudfjnistub spark_rapids_jni) if(USE_GDS) include(${CUDF_DIR}/cpp/cmake/Modules/FindcuFile.cmake) find_library(CUFILEJNI_LIB "libcufilejni.a" REQUIRED NO_DEFAULT_PATH - HINTS "${SPARK_RAPIDS_JNI_BINARY_DIR}/../libcudfjni" + HINTS "${CUDFJNI_BUILD_DIR}" ) add_library(cufilejni SHARED src/emptyfile.cpp) set_target_properties( From 96c940d886e823d887faff7498ad54ae48d4b158 Mon Sep 17 00:00:00 2001 From: Jason Lowe Date: Wed, 12 Jun 2024 10:55:47 -0500 Subject: [PATCH 049/159] Profiler: disable collecting async allocations by default (#2105) Signed-off-by: Jason Lowe --- src/main/cpp/profiler/ProfilerJni.cpp | 36 ++++--- src/main/cpp/profiler/profiler_serializer.cpp | 20 +++- src/main/cpp/profiler/profiler_serializer.hpp | 4 +- .../com/nvidia/spark/rapids/jni/Profiler.java | 98 +++++++++++++++---- 4 files changed, 122 insertions(+), 36 deletions(-) diff --git a/src/main/cpp/profiler/ProfilerJni.cpp b/src/main/cpp/profiler/ProfilerJni.cpp index 1271b89d7b..87e01b2e92 100644 --- a/src/main/cpp/profiler/ProfilerJni.cpp +++ b/src/main/cpp/profiler/ProfilerJni.cpp @@ -184,7 +184,8 @@ struct free_buffer_tracker { void writer_thread_process(JavaVM* vm, jobject j_writer, size_t buffer_size, - size_t flush_threshold); + size_t flush_threshold, + bool async_alloc_capture); struct subscriber_state { CUpti_SubscriberHandle subscriber_handle; @@ -363,11 +364,16 @@ void setup_nvtx_env(JNIEnv* env, jstring j_lib_path) } // Main processing loop for the background writer thread -void writer_thread_process(JavaVM* vm, jobject j_writer, size_t buffer_size, size_t flush_threshold) +void writer_thread_process(JavaVM* vm, + jobject j_writer, + size_t buffer_size, + size_t flush_threshold, + bool async_alloc_capture) { try { JNIEnv* env = attach_to_jvm(vm); - profiler_serializer serializer(env, j_writer, buffer_size, flush_threshold); + profiler_serializer serializer( + env, j_writer, buffer_size, flush_threshold, async_alloc_capture); auto buffer = State->completed_buffers.get(); while (buffer) { serializer.process_cupti_buffer(buffer->data(), buffer->valid_size()); @@ -419,12 +425,14 @@ extern "C" { using namespace spark_rapids_jni::profiler; -JNIEXPORT void JNICALL Java_com_nvidia_spark_rapids_jni_Profiler_nativeInit(JNIEnv* env, - jclass, - jstring j_lib_path, - jobject j_writer, - jlong write_buffer_size, - jint flush_period_msec) +JNIEXPORT void JNICALL +Java_com_nvidia_spark_rapids_jni_Profiler_nativeInit(JNIEnv* env, + jclass, + jstring j_lib_path, + jobject j_writer, + jlong write_buffer_size, + jint flush_period_msec, + bool async_alloc_capture) { try { setup_nvtx_env(env, j_lib_path); @@ -432,9 +440,13 @@ JNIEXPORT void JNICALL Java_com_nvidia_spark_rapids_jni_Profiler_nativeInit(JNIE auto writer = static_cast(env->NewGlobalRef(j_writer)); if (!writer) { throw std::runtime_error("Unable to create a global reference to writer"); } State = new subscriber_state(writer, write_buffer_size); - State->writer_thread = std::thread( - writer_thread_process, get_jvm(env), writer, write_buffer_size, write_buffer_size); - auto rc = cuptiSubscribe(&State->subscriber_handle, callback_handler, nullptr); + State->writer_thread = std::thread(writer_thread_process, + get_jvm(env), + writer, + write_buffer_size, + write_buffer_size, + async_alloc_capture); + auto rc = cuptiSubscribe(&State->subscriber_handle, callback_handler, nullptr); check_cupti(rc, "Error initializing CUPTI"); rc = cuptiEnableCallback(1, State->subscriber_handle, diff --git a/src/main/cpp/profiler/profiler_serializer.cpp b/src/main/cpp/profiler/profiler_serializer.cpp index b47ff234ad..68da1baf4a 100644 --- a/src/main/cpp/profiler/profiler_serializer.cpp +++ b/src/main/cpp/profiler/profiler_serializer.cpp @@ -197,11 +197,13 @@ ShmemLimitConfig to_shmem_limit_config(CUpti_FuncShmemLimitConfig c) } // anonymous namespace -profiler_serializer::profiler_serializer(JNIEnv* env, - jobject writer, - size_t buffer_size, - size_t flush_threshold) - : env_(env), j_writer_(writer), flush_threshold_(flush_threshold), fbb_(buffer_size) +profiler_serializer::profiler_serializer( + JNIEnv* env, jobject writer, size_t buffer_size, size_t flush_threshold, bool capture_allocs) + : env_(env), + j_writer_(writer), + flush_threshold_(flush_threshold), + fbb_(buffer_size), + capture_allocs_(capture_allocs) { auto writer_class = env->GetObjectClass(writer); if (!writer_class) { throw std::runtime_error("Failed to locate class of data writer"); } @@ -332,6 +334,14 @@ void profiler_serializer::process_api_activity(CUpti_ActivityAPI const* r) case CUPTI_RUNTIME_TRACE_CBID_cudaGetLastError_v3020: case CUPTI_RUNTIME_TRACE_CBID_cudaPeekAtLastError_v3020: case CUPTI_RUNTIME_TRACE_CBID_cudaDeviceGetAttribute_v5000: return; + case CUPTI_RUNTIME_TRACE_CBID_cudaMallocAsync_v11020: + case CUPTI_RUNTIME_TRACE_CBID_cudaMallocAsync_ptsz_v11020: + case CUPTI_RUNTIME_TRACE_CBID_cudaMallocFromPoolAsync_v11020: + case CUPTI_RUNTIME_TRACE_CBID_cudaMallocFromPoolAsync_ptsz_v11020: + case CUPTI_RUNTIME_TRACE_CBID_cudaFreeAsync_v11020: + case CUPTI_RUNTIME_TRACE_CBID_cudaFreeAsync_ptsz_v11020: + if (capture_allocs_) { break; } + return; default: break; } } else { diff --git a/src/main/cpp/profiler/profiler_serializer.hpp b/src/main/cpp/profiler/profiler_serializer.hpp index 1feebf1b96..861cf9d1ab 100644 --- a/src/main/cpp/profiler/profiler_serializer.hpp +++ b/src/main/cpp/profiler/profiler_serializer.hpp @@ -29,7 +29,8 @@ namespace spark_rapids_jni::profiler { // Serializes profile data as flatbuffers struct profiler_serializer { - profiler_serializer(JNIEnv* env, jobject writer, size_t buffer_size, size_t flush_threshold); + profiler_serializer( + JNIEnv* env, jobject writer, size_t buffer_size, size_t flush_threshold, bool capture_allocs); void process_cupti_buffer(uint8_t* buffer, size_t valid_size); void flush(); @@ -51,6 +52,7 @@ struct profiler_serializer { jmethodID j_write_method_; jobject j_writer_; size_t flush_threshold_; + bool capture_allocs_; flatbuffers::FlatBufferBuilder fbb_; std::vector> api_offsets_; std::vector> device_offsets_; diff --git a/src/main/java/com/nvidia/spark/rapids/jni/Profiler.java b/src/main/java/com/nvidia/spark/rapids/jni/Profiler.java index 86d5b0edde..85e6b4a0a3 100644 --- a/src/main/java/com/nvidia/spark/rapids/jni/Profiler.java +++ b/src/main/java/com/nvidia/spark/rapids/jni/Profiler.java @@ -16,8 +16,6 @@ package com.nvidia.spark.rapids.jni; import ai.rapids.cudf.NativeDepsLoader; -import org.slf4j.Logger; -import org.slf4j.LoggerFactory; import java.io.File; import java.io.IOException; @@ -27,18 +25,32 @@ public class Profiler { private static final long DEFAULT_WRITE_BUFFER_SIZE = 1024 * 1024; private static final int DEFAULT_FLUSH_PERIOD_MILLIS = 0; + private static final boolean DEFAULT_ALLOC_ASYNC_CAPTURING = false; private static DataWriter writer = null; /** * Initialize the profiler in a standby state. The start method must be called after this * to start collecting profiling data. - * @param w data writer for writing profiling data + * @param config profiler configuration */ - public static void init(DataWriter w) { - init(w, DEFAULT_WRITE_BUFFER_SIZE, DEFAULT_FLUSH_PERIOD_MILLIS); + public static void init(DataWriter w, Config config) { + if (writer == null) { + File libPath; + try { + libPath = NativeDepsLoader.loadNativeDep("profilerjni", true); + } catch (IOException e) { + throw new RuntimeException("Error loading profiler library", e); + } + nativeInit(libPath.getAbsolutePath(), w, config.writeBufferSize, config.flushPeriodMillis, + config.allocAsyncCapturing); + writer = w; + } else { + throw new IllegalStateException("Already initialized"); + } } /** + * Deprecated. Use init(Config) instead. * Initialize the profiler in a standby state. The start method must be called after this * to start collecting profiling data. * @param w data writer for writing profiling data @@ -50,18 +62,11 @@ public static void init(DataWriter w) { * flushing. */ public static void init(DataWriter w, long writeBufferSize, int flushPeriodMillis) { - if (writer == null) { - File libPath; - try { - libPath = NativeDepsLoader.loadNativeDep("profilerjni", true); - } catch (IOException e) { - throw new RuntimeException("Error loading profiler library", e); - } - nativeInit(libPath.getAbsolutePath(), w, writeBufferSize, flushPeriodMillis); - writer = w; - } else { - throw new IllegalStateException("Already initialized"); - } + Config config = new Config.Builder() + .withWriteBufferSize(writeBufferSize) + .withFlushPeriodMillis(flushPeriodMillis) + .build(); + init(w, config); } /** @@ -105,7 +110,8 @@ public static void stop() { } private static native void nativeInit(String libPath, DataWriter writer, - long writeBufferSize, int flushPeriodMillis); + long writeBufferSize, int flushPeriodMillis, + boolean allocAsyncCapturing); private static native void nativeStart(); @@ -122,4 +128,60 @@ public interface DataWriter extends AutoCloseable { */ void write(ByteBuffer data); } + + /** Profiler configuration class. **/ + public static class Config { + private final long writeBufferSize; + private final int flushPeriodMillis; + private final boolean allocAsyncCapturing; + + Config(Builder builder) { + this.writeBufferSize = builder.writeBufferSize; + this.flushPeriodMillis = builder.flushPeriodMillis; + this.allocAsyncCapturing = builder.allocAsyncCapturing; + } + + /** Builder interface for profiler configuration. **/ + public static class Builder { + private long writeBufferSize = DEFAULT_WRITE_BUFFER_SIZE; + private int flushPeriodMillis = DEFAULT_FLUSH_PERIOD_MILLIS; + private boolean allocAsyncCapturing = DEFAULT_ALLOC_ASYNC_CAPTURING; + + /** + * Configure the size of the host memory buffer used for collecting profiling data. + * Recommended to be between 1 to 8 MB in size to balance callback overhead with + * latency. + * @param writeBufferSize size of buffer in bytes + */ + public Builder withWriteBufferSize(long writeBufferSize) { + this.writeBufferSize = writeBufferSize; + return this; + } + + /** + * Configure the time period to explicitly flush collected profiling data to the writer. + * @param flushPeriodMillis time period in milliseconds. A value <= 0 will disable explicit + * flushing. + */ + public Builder withFlushPeriodMillis(int flushPeriodMillis) { + this.flushPeriodMillis = flushPeriodMillis; + return this; + } + + /** + * Configure whether async allocation and free events are captured by the profiler. + * @param allocAsyncCapturing true if async allocation and free events should be captured, + * false otherwise. + */ + public Builder withAllocAsyncCapturing(boolean allocAsyncCapturing) { + this.allocAsyncCapturing = allocAsyncCapturing; + return this; + } + + /** Build a profiler configuration object. */ + public Config build() { + return new Config(this); + } + } + } } From 526cbd48c57c4af736c17fc7dcd2f09b16a7d514 Mon Sep 17 00:00:00 2001 From: Gera Shegalov Date: Thu, 13 Jun 2024 13:22:17 -0700 Subject: [PATCH 050/159] Replace hard-coded libcudf build path in submodule-sync (#2143) Fixes #2142 Signed-off-by: Gera Shegalov --- ci/submodule-sync.sh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ci/submodule-sync.sh b/ci/submodule-sync.sh index 1888696ba5..9794d40a0f 100755 --- a/ci/submodule-sync.sh +++ b/ci/submodule-sync.sh @@ -88,8 +88,9 @@ else echo "Test failed, will update the result" fi +LIBCUDF_BUILD_PATH=$(${MVN} help:evaluate -Dexpression=libcudf.build.path -q -DforceStdout) # Extract the rapids-cmake sha1 that we need to pin too -rapids_cmake_sha=$(git -C thirdparty/cudf/cpp/build/_deps/rapids-cmake-src/ rev-parse HEAD) +rapids_cmake_sha=$(git -C ${LIBCUDF_BUILD_PATH}/_deps/rapids-cmake-src/ rev-parse HEAD) echo "Update rapids-cmake pinned SHA1 to ${rapids_cmake_sha}" echo "${rapids_cmake_sha}" > thirdparty/cudf-pins/rapids-cmake.sha From a6889936899321e553080b3407c9cbe929f3e760 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 14 Jun 2024 05:28:04 +0800 Subject: [PATCH 051/159] Update submodule cudf to 3cb3df3255efaec4a5ebb6cb7606067f753e3554 (#2144) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index f7ba6ab47a..3cb3df3255 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit f7ba6ab47ac994e6a1363119c01eee5dd6304181 +Subproject commit 3cb3df3255efaec4a5ebb6cb7606067f753e3554 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 768177632f..eecc48514c 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -5b0570870c86dd5f53ea00f0b46dd9e068f3f9f5 +c4b5fd596335cbd91d20b79425656e2e9b71f376 From 41c1b52e3d084219976d3690d06df3dad33d9431 Mon Sep 17 00:00:00 2001 From: Nghia Truong <7416935+ttnghia@users.noreply.github.com> Date: Thu, 13 Jun 2024 20:01:41 -0700 Subject: [PATCH 052/159] Optimize `get_json_object` by calling the main kernel only once (#2129) * Implement the CPU code Signed-off-by: Nghia Truong * Reimplement kernel Signed-off-by: Nghia Truong * Fix the kernel caller Signed-off-by: Nghia Truong * Optimize validity computation Signed-off-by: Nghia Truong * Cleanup Signed-off-by: Nghia Truong * Cleanup Signed-off-by: Nghia Truong * Cleanup Signed-off-by: Nghia Truong * Add comment Signed-off-by: Nghia Truong * Turning kernel occupancy Signed-off-by: Nghia Truong * Cleanup Signed-off-by: Nghia Truong * Change padding for the scratch buffer Signed-off-by: Nghia Truong * Update docs Signed-off-by: Nghia Truong * Add test for overflow case Signed-off-by: Nghia Truong * Pad the output buffer using max row size Signed-off-by: Nghia Truong * Update test Signed-off-by: Nghia Truong * Change the padding ratio Signed-off-by: Nghia Truong --------- Signed-off-by: Nghia Truong --- src/main/cpp/src/get_json_object.cu | 250 +++++++++--------- .../spark/rapids/jni/GetJsonObjectTest.java | 25 +- 2 files changed, 151 insertions(+), 124 deletions(-) diff --git a/src/main/cpp/src/get_json_object.cu b/src/main/cpp/src/get_json_object.cu index 887b9887de..eaa68765ea 100644 --- a/src/main/cpp/src/get_json_object.cu +++ b/src/main/cpp/src/get_json_object.cu @@ -24,6 +24,7 @@ #include #include #include +#include #include #include #include @@ -37,7 +38,11 @@ #include #include +#include +#include +#include #include +#include #include namespace spark_rapids_jni { @@ -825,33 +830,21 @@ rmm::device_uvector construct_path_commands( * * * @param input The incoming json string - * @param input_len Size of the incoming json string - * @param path_commands_ptr The command buffer to be applied to the string. - * @param path_commands_size The command buffer size. - * @param out_buf Buffer user to store the results of the query - * (nullptr in the size computation step) - * @param out_buf_size Size of the output buffer - * @returns A pair containing the result code and the output buffer. + * @param path_commands The command buffer to be applied to the string + * @param out_buf Buffer user to store the string resulted from the query + * @returns A pair containing the result code and the output buffer */ -__device__ thrust::pair get_json_object_single( - char_range input, - cudf::device_span path_commands, - char* out_buf, - size_t out_buf_size) +__device__ thrust::pair get_json_object_single( + char_range input, cudf::device_span path_commands, char* out_buf) { json_parser j_parser(input); j_parser.next_token(); // JSON validation check if (json_token::ERROR == j_parser.get_current_token()) { return {false, 0}; } - // First pass: preprocess sizes. - // Second pass: writes output. - // The generator automatically determines which pass based on `out_buf`. - // If `out_buf_size` is zero, pass in `nullptr` to avoid generator writing trash output. - json_generator generator((out_buf_size == 0) ? nullptr : out_buf); + json_generator generator(out_buf); - bool const success = evaluate_path( - j_parser, generator, write_style::RAW, {path_commands.data(), path_commands.size()}); + bool const success = evaluate_path(j_parser, generator, write_style::RAW, path_commands); if (!success) { // generator may contain trash output, e.g.: generator writes some output, @@ -860,25 +853,23 @@ __device__ thrust::pair get_json_object_single( generator.set_output_len_zero(); } - return {success, generator.get_output_len()}; + return {success, static_cast(generator.get_output_len())}; } /** * @brief Kernel for running the JSONPath query. * - * This kernel operates in a 2-pass way. On the first pass it computes the - * output sizes. On the second pass, it fills in the provided output buffers - * (chars and validity). + * This kernel writes out the output strings and their lengths at the same time. If any output + * length exceed buffer size limit, a boolean flag will be turned on to inform to the caller. + * In such situation, another (larger) output buffer will be generated and the kernel is launched + * again. Otherwise, launching this kernel only once is sufficient to produce the desired output. * - * @param col Device view of the incoming string + * @param input The input JSON strings stored in a strings column + * @param offsets Offsets to the output locations in the output buffer * @param path_commands JSONPath command buffer - * @param d_sizes a buffer used to write the output sizes in the first pass, - * and is read back in on the second pass to compute offsets. - * @param output_offsets Buffer used to store the string offsets for the results - * of the query - * @param out_buf Buffer used to store the results of the query - * @param out_validity Output validity buffer - * @param out_valid_count Output count of # of valid bits + * @param out_stringviews The output array to store pointers to the output strings and their sizes + * @param out_buf Buffer used to store the strings resulted from the query + * @param has_out_of_bound Flag to indicate if any output string has length exceeds its buffer size */ template // We have 1 for the minBlocksPerMultiprocessor in the launch bounds to avoid spilling from @@ -889,61 +880,33 @@ template // the performance is really bad. This essentially tells NVCC to prefer using lots // of registers over spilling. __launch_bounds__(block_size, 1) CUDF_KERNEL - void get_json_object_kernel(cudf::column_device_view col, + void get_json_object_kernel(cudf::column_device_view input, + cudf::detail::input_offsetalator offsets, cudf::device_span path_commands, - cudf::size_type* d_sizes, - cudf::detail::input_offsetalator output_offsets, + thrust::pair* out_stringviews, char* out_buf, - cudf::bitmask_type* out_validity, - cudf::size_type* out_valid_count) + bool* has_out_of_bound) { - auto tid = cudf::detail::grid_1d::global_thread_id(); auto const stride = cudf::detail::grid_1d::grid_stride(); + for (auto tid = cudf::detail::grid_1d::global_thread_id(); tid < input.size(); tid += stride) { + char* const dst = out_buf + offsets[tid]; + bool is_valid = false; + cudf::size_type out_size = 0; - cudf::size_type warp_valid_count{0}; - - auto active_threads = __ballot_sync(0xffff'ffffu, tid < col.size()); - while (tid < col.size()) { - bool is_valid = false; - cudf::string_view const str = col.element(tid); + auto const str = input.element(tid); if (str.size_bytes() > 0) { - char* dst = out_buf != nullptr ? out_buf + output_offsets[tid] : nullptr; - size_t const dst_size = - out_buf != nullptr ? output_offsets[tid + 1] - output_offsets[tid] : 0; - - // process one single row - auto [result, output_size] = - get_json_object_single(str, {path_commands.data(), path_commands.size()}, dst, dst_size); - if (result) { is_valid = true; } - - // filled in only during the precompute step. during the compute step, the - // offsets are fed back in so we do -not- want to write them out - if (out_buf == nullptr) { d_sizes[tid] = static_cast(output_size); } - } else { - // valid JSON length is always greater than 0 - // if `str` size len is zero, output len is 0 and `is_valid` is false - if (out_buf == nullptr) { d_sizes[tid] = 0; } - } + auto const max_size = offsets[tid + 1] - offsets[tid]; - // validity filled in only during the output step - if (out_validity != nullptr) { - uint32_t mask = __ballot_sync(active_threads, is_valid); - // 0th lane of the warp writes the validity - if (!(tid % cudf::detail::warp_size)) { - out_validity[cudf::word_index(tid)] = mask; - warp_valid_count += __popc(mask); - } + // If `max_size == 0`, do not pass in the dst pointer to prevent writing garbage data. + thrust::tie(is_valid, out_size) = + get_json_object_single(str, path_commands, max_size != 0 ? dst : nullptr); + if (out_size > max_size) { *has_out_of_bound = true; } } - tid += stride; - active_threads = __ballot_sync(active_threads, tid < col.size()); - } - - // sum the valid counts across the whole block - if (out_valid_count != nullptr) { - cudf::size_type block_valid_count = - cudf::detail::single_lane_block_sum_reduce(warp_valid_count); - if (threadIdx.x == 0) { atomicAdd(out_valid_count, block_valid_count); } + // Write out `nullptr` in the output string_view to indicate that the output is a null. + // The situation `out_stringviews == nullptr` should only happen if the kernel is launched a + // second time due to out-of-bound write in the first launch. + if (out_stringviews) { out_stringviews[tid] = {is_valid ? dst : nullptr, out_size}; } } } @@ -953,64 +916,111 @@ std::unique_ptr get_json_object( rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - if (input.is_empty()) return cudf::make_empty_column(cudf::type_id::STRING); - if (instructions.size() > max_path_depth) { CUDF_FAIL("JSONPath query exceeds maximum depth"); } + if (input.is_empty()) { return cudf::make_empty_column(cudf::type_id::STRING); } - // get a string buffer to store all the names and convert to device std::string all_names; for (auto const& inst : instructions) { all_names += std::get<1>(inst); } - cudf::string_scalar all_names_scalar(all_names, true, stream); - // parse the json_path into a command buffer - auto path_commands = construct_path_commands( + auto const all_names_scalar = cudf::string_scalar(all_names, true, stream); + auto const path_commands = construct_path_commands( instructions, all_names_scalar, stream, rmm::mr::get_current_device_resource()); + auto const d_input_ptr = cudf::column_device_view::create(input.parent(), stream); + auto const in_offsets = cudf::detail::offsetalator_factory::make_input_iterator(input.offsets()); + + // A buffer to store the output strings without knowing their sizes. + // Since we do not know their sizes, we need to allocate the buffer a bit larger than the input + // size so that we will not write output strings into an out-of-bound position. + // Checking out-of-bound needs to be performed in the main kernel to make sure we will not have + // data corruption. + auto const scratch_size = [&] { + auto const max_row_size = thrust::transform_reduce( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(input.size()), + cuda::proclaim_return_type( + [in_offsets] __device__(auto const idx) { return in_offsets[idx + 1] - in_offsets[idx]; }), + int64_t{0}, + thrust::maximum{}); + + // Pad the scratch buffer by an additional size that is a multiple of max row size. + auto constexpr padding_rows = 10; + return input.chars_size(stream) + max_row_size * padding_rows; + }(); + auto output_scratch = rmm::device_uvector(scratch_size, stream); + auto out_stringviews = rmm::device_uvector>{ + static_cast(input.size()), stream}; + auto has_out_of_bound = rmm::device_scalar{false, stream}; + + constexpr int blocks_per_SM = 1; + constexpr int block_size = 256; + auto const num_blocks = [&] { + int device_id{}; + cudaDeviceProp props{}; + CUDF_CUDA_TRY(cudaGetDevice(&device_id)); + CUDF_CUDA_TRY(cudaGetDeviceProperties(&props, device_id)); + return props.multiProcessorCount * blocks_per_SM; + }(); - // compute output sizes - auto sizes = rmm::device_uvector( - input.size(), stream, rmm::mr::get_current_device_resource()); - auto d_offsets = cudf::detail::offsetalator_factory::make_input_iterator(input.offsets()); - - constexpr int block_size = 512; - cudf::detail::grid_1d const grid{input.size(), block_size}; - auto d_input_ptr = cudf::column_device_view::create(input.parent(), stream); - // preprocess sizes (returned in the offsets buffer) get_json_object_kernel - <<>>( - *d_input_ptr, path_commands, sizes.data(), d_offsets, nullptr, nullptr, nullptr); - - // convert sizes to offsets + <<>>(*d_input_ptr, + in_offsets, + path_commands, + out_stringviews.data(), + output_scratch.data(), + has_out_of_bound.data()); + + // If we didn't see any out-of-bound write, everything is good so far. + // Just gather the output strings and return. + if (!has_out_of_bound.value(stream)) { + return cudf::make_strings_column(out_stringviews, stream, mr); + } + // From here, we had out-of-bound write. Although this is very rare, it may still happen. + + // This scratch buffer is no longer needed. + output_scratch = rmm::device_uvector{0, stream}; + + // The string sizes computed in the previous kernel call will be used to allocate a new char + // buffer to store the output. + auto const size_it = cudf::detail::make_counting_transform_iterator( + 0, + cuda::proclaim_return_type( + [string_pairs = out_stringviews.data()] __device__(auto const idx) { + return string_pairs[idx].second; + })); auto [offsets, output_size] = - cudf::strings::detail::make_offsets_child_column(sizes.begin(), sizes.end(), stream, mr); - d_offsets = cudf::detail::offsetalator_factory::make_input_iterator(offsets->view()); + cudf::strings::detail::make_offsets_child_column(size_it, size_it + input.size(), stream, mr); - // allocate output string column - rmm::device_uvector chars(output_size, stream, mr); + // Also compute the null mask using the stored char pointers. + auto const validator = [] __device__(thrust::pair const item) { + return item.first != nullptr; + }; + auto [null_mask, null_count] = + cudf::detail::valid_if(out_stringviews.begin(), out_stringviews.end(), validator, stream, mr); - // potential optimization : if we know that all outputs are valid, we could - // skip creating the validity mask altogether - rmm::device_buffer validity = - cudf::detail::create_null_mask(input.size(), cudf::mask_state::UNINITIALIZED, stream, mr); + // No longer need it from here. Free up memory for now. + out_stringviews = rmm::device_uvector>{0, stream}; - // compute results - rmm::device_scalar d_valid_count{0, stream}; + auto chars = rmm::device_uvector(output_size, stream, mr); + auto const out_offsets = cudf::detail::offsetalator_factory::make_input_iterator(offsets->view()); + has_out_of_bound.set_value_to_zero_async(stream); get_json_object_kernel - <<>>( - *d_input_ptr, - path_commands, - sizes.data(), - d_offsets, - chars.data(), - static_cast(validity.data()), - d_valid_count.data()); - - return make_strings_column(input.size(), - std::move(offsets), - chars.release(), - input.size() - d_valid_count.value(stream), - std::move(validity)); + <<>>(*d_input_ptr, + out_offsets, + path_commands, + nullptr /*out_stringviews*/, + chars.data(), + has_out_of_bound.data()); + + // This kernel call should not see out-of-bound write. If it is still detected, there must be + // something wrong happened. + CUDF_EXPECTS(!has_out_of_bound.value(stream), + "Unexpected out-of-bound write in get_json_object kernel."); + + return cudf::make_strings_column( + input.size(), std::move(offsets), chars.release(), null_count, std::move(null_mask)); } } // namespace detail diff --git a/src/test/java/com/nvidia/spark/rapids/jni/GetJsonObjectTest.java b/src/test/java/com/nvidia/spark/rapids/jni/GetJsonObjectTest.java index bba6650d0f..1454a2d60b 100644 --- a/src/test/java/com/nvidia/spark/rapids/jni/GetJsonObjectTest.java +++ b/src/test/java/com/nvidia/spark/rapids/jni/GetJsonObjectTest.java @@ -31,9 +31,9 @@ void getJsonObjectTest() { namedPath("k") }; try (ColumnVector jsonCv = ColumnVector.fromStrings( "{\"k\": \"v\"}"); - ColumnVector expected = ColumnVector.fromStrings( - "v"); - ColumnVector actual = JSONUtils.getJsonObject(jsonCv, query)) { + ColumnVector expected = ColumnVector.fromStrings( + "v"); + ColumnVector actual = JSONUtils.getJsonObject(jsonCv, query)) { assertColumnsAreEqual(expected, actual); } } @@ -169,7 +169,7 @@ void getJsonObjectTest_Escape() { String JSON4 = "['a','b','\"C\"']"; // \\u4e2d\\u56FD is 中国 String JSON5 = "'\\u4e2d\\u56FD\\\"\\'\\\\\\/\\b\\f\\n\\r\\t\\b'"; - String JSON6 = "['\\u4e2d\\u56FD\\\"\\'\\\\\\/\\b\\f\\n\\r\\t\\b']"; + String JSON6 = "['\\u4e2d\\u56FD\\\"\\'\\\\\\/\\b\\f\\n\\r\\t\\b']"; String expectedStr1 = "{\"a\":\"A\"}"; String expectedStr2 = "{\"a\":\"A\\\"\"}"; @@ -600,6 +600,23 @@ void getJsonObjectTest_15() { } } + /** + * This test is when the JNI kernel is called twice. It happens when the output JSON strings + * have lengths that are larger than their corresponding input. + */ + @Test + void getJsonObjectTest_JNIKernelCalledTwice() { + // This is equivalent to the path '$'. + JSONUtils.PathInstructionJni[] query = new JSONUtils.PathInstructionJni[] {}; + try ( + ColumnVector input = ColumnVector.fromStrings("['\n']", "['\n\n\n\n\n\n\n\n\n\n']", + "", "", "", "", "", "", "", ""); + ColumnVector expected = ColumnVector.fromStrings("[\"\\n\"]", + "[\"\\n\\n\\n\\n\\n\\n\\n\\n\\n\\n\"]", null, null, null, null, null, null, null, null); + ColumnVector actual = JSONUtils.getJsonObject(input, query)) { + assertColumnsAreEqual(expected, actual); + } + } private JSONUtils.PathInstructionJni wildcardPath() { return new JSONUtils.PathInstructionJni(JSONUtils.PathInstructionType.WILDCARD, "", -1); From 6f6dd2cc3523773f41193e58fd483d96d3029c29 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 14 Jun 2024 11:14:19 +0800 Subject: [PATCH 053/159] Update submodule cudf to 31d909b0af9bcf9cf804ca1c3893ea71fbd5d765 (#2146) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 3cb3df3255..31d909b0af 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 3cb3df3255efaec4a5ebb6cb7606067f753e3554 +Subproject commit 31d909b0af9bcf9cf804ca1c3893ea71fbd5d765 From c7984e463dbcf816c1414fec8d6c305d9227e8cf Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 14 Jun 2024 16:32:19 +0800 Subject: [PATCH 054/159] Update submodule cudf to 34227d3cb687d465f1d4a5f12cbb37a47b97866e (#2147) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 31d909b0af..34227d3cb6 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 31d909b0af9bcf9cf804ca1c3893ea71fbd5d765 +Subproject commit 34227d3cb687d465f1d4a5f12cbb37a47b97866e From 93d82759f09a3f21f661f2e07beb04fd3e4474f5 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 15 Jun 2024 05:14:55 +0800 Subject: [PATCH 055/159] Update submodule cudf to f89cc07b50d3f89e7da8f98afb5fe8f9d9cf33c6 (#2148) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 35 ------------------------------ 2 files changed, 1 insertion(+), 36 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 34227d3cb6..f89cc07b50 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 34227d3cb687d465f1d4a5f12cbb37a47b97866e +Subproject commit f89cc07b50d3f89e7da8f98afb5fe8f9d9cf33c6 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index ca0af67cac..9e492e5efc 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -9,60 +9,25 @@ "git_url" : "https://github.com/NVIDIA/cccl.git", "patches" : [ - { - "file" : "cccl/bug_fixes.diff", - "fixed_in" : "2.3", - "issue" : "CCCL installs header-search.cmake files in nondeterministic order and has a typo in checking target creation that leads to duplicates" - }, - { - "file" : "cccl/hide_kernels.diff", - "fixed_in" : "2.3", - "issue" : "Mark all cub and thrust kernels with hidden visibility [https://github.com/nvidia/cccl/pulls/443]" - }, { "file" : "cccl/revert_pr_211.diff", "fixed_in" : "", "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue." }, - { - "file" : "${current_json_dir}/revert_pr_211_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "thrust::copy introduced a change in behavior that causes failures with cudaErrorInvalidValue." - }, - { - "file" : "cccl/kernel_pointer_hiding.diff", - "fixed_in" : "2.4", - "issue" : "Hide APIs that accept kernel pointers [https://github.com/NVIDIA/cccl/pull/1395]" - }, { "file" : "${current_json_dir}/thrust_disable_64bit_dispatching.diff", "fixed_in" : "", "issue" : "Remove 64bit dispatching as not needed by libcudf and results in compiling twice as many kernels [https://github.com/rapidsai/cudf/pull/11437]" }, - { - "file" : "${current_json_dir}/thrust_disable_64bit_dispatching_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "Remove 64bit dispatching as not needed by libcudf and results in compiling twice as many kernels [https://github.com/rapidsai/cudf/pull/11437]" - }, { "file" : "${current_json_dir}/thrust_faster_sort_compile_times.diff", "fixed_in" : "", "issue" : "Improve Thrust sort compile times by not unrolling loops for inlined comparators [https://github.com/rapidsai/cudf/pull/10577]" }, - { - "file" : "${current_json_dir}/thrust_faster_sort_compile_times_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "Improve Thrust sort compile times by not unrolling loops for inlined comparators [https://github.com/rapidsai/cudf/pull/10577]" - }, { "file" : "${current_json_dir}/thrust_faster_scan_compile_times.diff", "fixed_in" : "", "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" - }, - { - "file" : "${current_json_dir}/thrust_faster_scan_compile_times_cccl_2.5.0.diff", - "fixed_in" : "", - "issue" : "Improve Thrust scan compile times by reducing the number of kernels generated [https://github.com/rapidsai/cudf/pull/8183]" } ], "version" : "2.5.0" From 22ab2f54dc9d96a43a3679f4b0e5196987185a90 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 15 Jun 2024 10:30:17 +0800 Subject: [PATCH 056/159] Update submodule cudf to 74b382637e69d39df292c59938b5911d9ca3bdf9 (#2149) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index f89cc07b50..74b382637e 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit f89cc07b50d3f89e7da8f98afb5fe8f9d9cf33c6 +Subproject commit 74b382637e69d39df292c59938b5911d9ca3bdf9 From 0a444ffdc740542bf85ed0ba6ae8e3e3c1ee7efa Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Mon, 17 Jun 2024 20:36:54 +0800 Subject: [PATCH 057/159] Update submodule cudf to e9ebdea49d24f645a6ca5ff6d79e0525a114f5fc (#2150) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 74b382637e..e9ebdea49d 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 74b382637e69d39df292c59938b5911d9ca3bdf9 +Subproject commit e9ebdea49d24f645a6ca5ff6d79e0525a114f5fc From 673b8c14681e73cf68432b552f45d18e3cb5d751 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 18 Jun 2024 04:31:07 +0800 Subject: [PATCH 058/159] Update submodule cudf to 6ff4b4b27e7cc9b750146626531dca0a2a5307c4 (#2152) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index e9ebdea49d..6ff4b4b27e 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit e9ebdea49d24f645a6ca5ff6d79e0525a114f5fc +Subproject commit 6ff4b4b27e7cc9b750146626531dca0a2a5307c4 From 0ec3c13c196b69b18043af5f65532f37ed97f59b Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 18 Jun 2024 10:30:17 +0800 Subject: [PATCH 059/159] Update submodule cudf to 56e84425e84029b9b7c2ba07f0b8bbfd94846a40 (#2154) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 6ff4b4b27e..56e84425e8 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 6ff4b4b27e7cc9b750146626531dca0a2a5307c4 +Subproject commit 56e84425e84029b9b7c2ba07f0b8bbfd94846a40 From 9bded36e38e497a50e1a6850256777febfa92a12 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 18 Jun 2024 16:30:49 +0800 Subject: [PATCH 060/159] Update submodule cudf to 0bdf934f6017402b88c9c0fe798013203af8c39f (#2155) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 56e84425e8..0bdf934f60 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 56e84425e84029b9b7c2ba07f0b8bbfd94846a40 +Subproject commit 0bdf934f6017402b88c9c0fe798013203af8c39f From a7add46bc96079065e6165d0b1af1838a09ed307 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 18 Jun 2024 21:20:29 +0800 Subject: [PATCH 061/159] Update submodule cudf to dcc153b67c48909a7bd5fcecfd4ccc91844e55ec (#2156) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 0bdf934f60..dcc153b67c 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 0bdf934f6017402b88c9c0fe798013203af8c39f +Subproject commit dcc153b67c48909a7bd5fcecfd4ccc91844e55ec From 2185a5716826c69c942522fcd9be236f9ab6b06d Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 19 Jun 2024 05:15:01 +0800 Subject: [PATCH 062/159] Update submodule cudf to 9bc794aa355c8e4c42fbc611fe9d496c20a4db90 (#2157) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index dcc153b67c..9bc794aa35 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit dcc153b67c48909a7bd5fcecfd4ccc91844e55ec +Subproject commit 9bc794aa355c8e4c42fbc611fe9d496c20a4db90 From 3305f970c76e02f1aa94484bf84ca290a5f6e282 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 19 Jun 2024 10:30:24 +0800 Subject: [PATCH 063/159] Update submodule cudf to c83e5b3fdd7f9fe8a08c4f6874fbf847bba70c53 (#2158) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 9bc794aa35..c83e5b3fdd 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 9bc794aa355c8e4c42fbc611fe9d496c20a4db90 +Subproject commit c83e5b3fdd7f9fe8a08c4f6874fbf847bba70c53 From 2e9a294a769e2b64221e927807f82283aa345c69 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 20 Jun 2024 04:28:56 +0800 Subject: [PATCH 064/159] Update submodule cudf to f536e3017205be8b09f3dc2cfd448dc9c5a94d5d (#2159) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index c83e5b3fdd..f536e30172 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit c83e5b3fdd7f9fe8a08c4f6874fbf847bba70c53 +Subproject commit f536e3017205be8b09f3dc2cfd448dc9c5a94d5d From e8ca968d1cefb52fbb2b4aaf781a79b46f00c134 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Mon, 24 Jun 2024 22:29:32 +0800 Subject: [PATCH 065/159] Update submodule cudf to ac3c8dddda2fac2cb02c8a8ee58d827c00ddf867 (#2161) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index f536e30172..ac3c8dddda 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit f536e3017205be8b09f3dc2cfd448dc9c5a94d5d +Subproject commit ac3c8dddda2fac2cb02c8a8ee58d827c00ddf867 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 9e492e5efc..dbf3e547bf 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -5,7 +5,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "fde1cf79bde6744b6739636f502b5edfefe302e1", + "git_tag" : "e21d607157218540cd7c45461213fb96adf720b7", "git_url" : "https://github.com/NVIDIA/cccl.git", "patches" : [ From b9ed7b942493b04dd9cb476d100d6eeb2fb035c0 Mon Sep 17 00:00:00 2001 From: Gera Shegalov Date: Mon, 24 Jun 2024 08:35:02 -0700 Subject: [PATCH 066/159] Make run-in-docker usable in other repos (#2153) - Introduces WORKDIR to change the default local path available to Docker to point to another repo This PR makes it possible to write easy repro steps for cudf or any other path to be executed in the spark-rapids-jni Docker container Signed-off-by: Gera Shegalov --- CONTRIBUTING.md | 12 ++++++++++++ build/run-in-docker | 5 +++-- 2 files changed, 15 insertions(+), 2 deletions(-) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index e30f2da15d..7f83e2169b 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -61,6 +61,18 @@ the Docker container. The script passes all of its arguments onto the Maven command run inside the Docker container, so it should be invoked as one would invoke Maven, e.g.: `build/build-in-docker clean package` +#### Using spark-rapids-jni Docker Container with other Repos + +Spark RAPIDS project spans multiple repos. Some issues are discovered in +spark-rapids-jni but they need to be made easily reproducible in the cudf repo + +To this end export WORKDIR with the path pointing to a different repo + +``` +export WORKDIR=~/gits/rapidsai/cudf +~/gits/NVIDIA/spark-rapids-jni/build/run-in-docker head README.md +``` + ### cudf Submodule and Build [RAPIDS cuDF](https://github.com/rapidsai/cudf) is being used as a submodule in this project. diff --git a/build/run-in-docker b/build/run-in-docker index 51e8a7af27..2e3a732643 100755 --- a/build/run-in-docker +++ b/build/run-in-docker @@ -23,6 +23,7 @@ set -e # Base paths relative to this script's location SCRIPTDIR=$(cd $(dirname $0); pwd) REPODIR=$SCRIPTDIR/.. +WORKDIR=${WORKDIR:-$REPODIR} CUDA_VERSION=${CUDA_VERSION:-11.8.0} DOCKER_CMD=${DOCKER_CMD:-docker} @@ -63,10 +64,10 @@ $DOCKER_CMD run $DOCKER_GPU_OPTS $DOCKER_RUN_EXTRA_ARGS -u $(id -u):$(id -g) --r -v "/etc/passwd:/etc/passwd:ro" \ -v "/etc/shadow:/etc/shadow:ro" \ -v "/etc/sudoers.d:/etc/sudoers.d:ro" \ - -v "$REPODIR:$REPODIR:rw" \ + -v "$WORKDIR:$WORKDIR:rw" \ -v "$LOCAL_CCACHE_DIR:$LOCAL_CCACHE_DIR:rw" \ -v "$LOCAL_MAVEN_REPO:$LOCAL_MAVEN_REPO:rw" \ - --workdir "$REPODIR" \ + --workdir "$WORKDIR" \ -e CCACHE_DIR="$LOCAL_CCACHE_DIR" \ -e CMAKE_C_COMPILER_LAUNCHER="ccache" \ -e CMAKE_CXX_COMPILER_LAUNCHER="ccache" \ From b56a7eb8c59975f0814aa913bde904fb52e40e2c Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 25 Jun 2024 04:32:26 +0800 Subject: [PATCH 067/159] Update submodule cudf to bd76bf6b293b7f17a846df8392c18d92ced2b40f (#2167) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index ac3c8dddda..bd76bf6b29 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit ac3c8dddda2fac2cb02c8a8ee58d827c00ddf867 +Subproject commit bd76bf6b293b7f17a846df8392c18d92ced2b40f diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index eecc48514c..0dcb044f5f 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -c4b5fd596335cbd91d20b79425656e2e9b71f376 +5d47273ee4f64ec2710a6f3e7c9d9364813dcfda From 4fc888cbe6f47384664429aaa2163bbe40a21db9 Mon Sep 17 00:00:00 2001 From: Jason Lowe Date: Mon, 24 Jun 2024 16:48:47 -0500 Subject: [PATCH 068/159] Restore static CUDA runtime linking (#2165) Signed-off-by: Jason Lowe --- src/main/cpp/CMakeLists.txt | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt index c78e9abc0f..5e9dff7ac6 100644 --- a/src/main/cpp/CMakeLists.txt +++ b/src/main/cpp/CMakeLists.txt @@ -30,6 +30,8 @@ set(CUDF_CPP_BUILD_DIR "$ENV{CUDF_CPP_BUILD_DIR}" CACHE STRING "path to libcudf build root" ) +# libcudf's kvikio dependency requires this to be set when statically linking CUDA runtime +set(CUDA_STATIC_RUNTIME ON) include("${CUDF_DIR}/rapids_config.cmake") include(rapids-cmake) @@ -267,6 +269,7 @@ target_link_libraries( ${PARQUET_LIB} ${THRIFT_LIB} ) +rapids_cuda_set_runtime(spark_rapids_jni USE_STATIC ON) set_target_properties(spark_rapids_jni PROPERTIES LINK_LANGUAGE "CXX") # For backwards-compatibility with the cudf Java bindings and RAPIDS accelerated UDFs, # all of the code is built into libcudf.so that is statically linked to the CUDA runtime library. @@ -303,8 +306,8 @@ if(USE_GDS) -Wl,--no-whole-archive spark_rapids_jni ${cuFile_LIBRARIES} - CUDA::cudart_static ) + rapids_cuda_set_runtime(cufilejni USE_STATIC ON) endif() # ################################################################################################## From 0529efaacda521b5662ab0dd28018a4e158648ea Mon Sep 17 00:00:00 2001 From: Jason Lowe Date: Mon, 24 Jun 2024 16:49:19 -0500 Subject: [PATCH 069/159] Avoid re-generating thirdparty/cudf-pins/versions.json after libcudf build (#2166) Signed-off-by: Jason Lowe --- pom.xml | 4 +- thirdparty/cudf-pins/versions.json | 111 ++++++++++++++++++++++++++++- 2 files changed, 112 insertions(+), 3 deletions(-) diff --git a/pom.xml b/pom.xml index a1a22b3189..64641c21fb 100644 --- a/pom.xml +++ b/pom.xml @@ -463,7 +463,7 @@ - + - + Date: Tue, 25 Jun 2024 09:05:02 +0800 Subject: [PATCH 070/159] upgrade blossom ci actions version (#2160) Signed-off-by: YanxuanLiu --- .github/workflows/blossom-ci.yml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/.github/workflows/blossom-ci.yml b/.github/workflows/blossom-ci.yml index 686c655ca1..a17fbddd83 100644 --- a/.github/workflows/blossom-ci.yml +++ b/.github/workflows/blossom-ci.yml @@ -84,7 +84,7 @@ jobs: runs-on: ubuntu-latest steps: - name: Checkout code - uses: actions/checkout@v3 + uses: actions/checkout@v4 with: repository: ${{ fromJson(needs.Authorization.outputs.args).repo }} ref: ${{ fromJson(needs.Authorization.outputs.args).ref }} @@ -92,7 +92,7 @@ jobs: # repo specific steps - name: Setup java - uses: actions/setup-java@v3 + uses: actions/setup-java@v4 with: distribution: adopt java-version: 8 From 54bb39d131f9c47fe322db08c093191b682caaa3 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Tue, 25 Jun 2024 10:34:17 +0800 Subject: [PATCH 071/159] Update submodule cudf to b9a0b72773a3adf4ba9ae267911d8970f0db53b0 (#2170) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index bd76bf6b29..b9a0b72773 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit bd76bf6b293b7f17a846df8392c18d92ced2b40f +Subproject commit b9a0b72773a3adf4ba9ae267911d8970f0db53b0 From c484470174fcf5599187dc6db8b2188e3127fba1 Mon Sep 17 00:00:00 2001 From: Chong Gao Date: Tue, 25 Jun 2024 21:07:38 +0800 Subject: [PATCH 072/159] case when improvement: avoid copy_if_else (#2079) * case when improvement: avoid copy_if_else Signed-off-by: Chong Gao * update doc Signed-off-by: Chong Gao * Fix null handing bug * Refactor code: remove duplicated code * Address comments Signed-off-by: Chong Gao * Refactor * Update src/main/cpp/src/case_when.cu * Refactor --------- Signed-off-by: Chong Gao Co-authored-by: Chong Gao Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com> --- src/main/cpp/CMakeLists.txt | 2 + src/main/cpp/src/CaseWhenJni.cpp | 35 ++++++ src/main/cpp/src/case_when.cu | 102 ++++++++++++++++++ src/main/cpp/src/case_when.hpp | 53 +++++++++ .../com/nvidia/spark/rapids/jni/CaseWhen.java | 83 ++++++++++++++ .../nvidia/spark/rapids/jni/CaseWhenTest.java | 64 +++++++++++ 6 files changed, 339 insertions(+) create mode 100644 src/main/cpp/src/CaseWhenJni.cpp create mode 100644 src/main/cpp/src/case_when.cu create mode 100644 src/main/cpp/src/case_when.hpp create mode 100644 src/main/java/com/nvidia/spark/rapids/jni/CaseWhen.java create mode 100644 src/test/java/com/nvidia/spark/rapids/jni/CaseWhenTest.java diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt index 5e9dff7ac6..2de79a74c1 100644 --- a/src/main/cpp/CMakeLists.txt +++ b/src/main/cpp/CMakeLists.txt @@ -183,6 +183,7 @@ set(CUDFJNI_INCLUDE_DIRS add_library( spark_rapids_jni SHARED src/BloomFilterJni.cpp + src/CaseWhenJni.cpp src/CastStringJni.cpp src/DateTimeRebaseJni.cpp src/DecimalUtilsJni.cpp @@ -198,6 +199,7 @@ add_library( src/SparkResourceAdaptorJni.cpp src/ZOrderJni.cpp src/bloom_filter.cu + src/case_when.cu src/cast_decimal_to_string.cu src/format_float.cu src/cast_float_to_string.cu diff --git a/src/main/cpp/src/CaseWhenJni.cpp b/src/main/cpp/src/CaseWhenJni.cpp new file mode 100644 index 0000000000..2f99e85b4b --- /dev/null +++ b/src/main/cpp/src/CaseWhenJni.cpp @@ -0,0 +1,35 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "case_when.hpp" +#include "cudf_jni_apis.hpp" + +extern "C" { + +JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_CaseWhen_selectFirstTrueIndex( + JNIEnv* env, jclass, jlongArray bool_cols) +{ + JNI_NULL_CHECK(env, bool_cols, "array of column handles is null", 0); + try { + cudf::jni::auto_set_device(env); + cudf::jni::native_jpointerArray n_cudf_bool_columns(env, bool_cols); + auto bool_column_views = n_cudf_bool_columns.get_dereferenced(); + return cudf::jni::release_as_jlong( + spark_rapids_jni::select_first_true_index(cudf::table_view(bool_column_views))); + } + CATCH_STD(env, 0); +} +} diff --git a/src/main/cpp/src/case_when.cu b/src/main/cpp/src/case_when.cu new file mode 100644 index 0000000000..9857403898 --- /dev/null +++ b/src/main/cpp/src/case_when.cu @@ -0,0 +1,102 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "case_when.hpp" + +#include +#include +#include +#include + +#include + +namespace spark_rapids_jni { +namespace detail { +namespace { + +/** + * Select the column index for the first true in bool columns for the specified row + */ +struct select_first_true_fn { + // bool columns stores the results of executing `when` expressions + cudf::table_device_view const d_table; + + /** + * The number of bool columns is the size of case when branches. + * Note: reuturned index may be out of bound, valid bound is [0, col_num) + * When returning col_num index, it means final result is NULL value or ELSE value. + * + * e.g.: + * CASE WHEN 'a' THEN 'A' END + * The number of bool columns is 1 + * The number of scalars is 1 + * Max index is 1 which means using NULL(all when exprs are false). + * CASE WHEN 'a' THEN 'A' ELSE '_' END + * The number of bool columns is 1 + * The number of scalars is 2 + * Max index is also 1 which means using else value '_' + */ + __device__ cudf::size_type operator()(std::size_t row_idx) const + { + auto col_num = d_table.num_columns(); + for (auto col_idx = 0; col_idx < col_num; col_idx++) { + auto const& col = d_table.column(col_idx); + if (!col.is_null(row_idx) && col.element(row_idx)) { + // Predicate is true and not null + return col_idx; + } + } + return col_num; + } +}; + +} // anonymous namespace + +std::unique_ptr select_first_true_index(cudf::table_view const& when_bool_columns, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + // checks + auto const num_columns = when_bool_columns.num_columns(); + CUDF_EXPECTS(num_columns > 0, "At least one column must be specified"); + auto const row_count = when_bool_columns.num_rows(); + if (row_count == 0) { // empty begets empty + return cudf::make_empty_column(cudf::type_id::INT32); + } + // make output column + auto ret = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT32}, row_count, cudf::mask_state::UNALLOCATED, stream, mr); + + // select first true index + auto const d_table_ptr = cudf::table_device_view::create(when_bool_columns, stream); + thrust::transform(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(row_count), + ret->mutable_view().begin(), + select_first_true_fn{*d_table_ptr}); + return ret; +} + +} // namespace detail + +std::unique_ptr select_first_true_index(cudf::table_view const& when_bool_columns, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + return detail::select_first_true_index(when_bool_columns, stream, mr); +} + +} // namespace spark_rapids_jni diff --git a/src/main/cpp/src/case_when.hpp b/src/main/cpp/src/case_when.hpp new file mode 100644 index 0000000000..b7056c6c8f --- /dev/null +++ b/src/main/cpp/src/case_when.hpp @@ -0,0 +1,53 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include + +#include +#include + +#include + +namespace spark_rapids_jni { + +/** + * + * Select the column index for the first true in bool columns. + * For the row does not contain true, use end index(number of columns). + * + * e.g.: + * column 0 in table: true, false, false, false + * column 1 in table: false, true, false, false + * column 2 in table: false, false, true, false + * + * 1st row is: true, flase, false; first true index is 0 + * 2nd row is: false, true, false; first true index is 1 + * 3rd row is: false, flase, true; first true index is 2 + * 4th row is: false, false, false; do not find true, set index to the end index 3 + * + * output column: 0, 1, 2, 3 + * In the `case when` context, here 3 index means using NULL value. + * + */ +std::unique_ptr select_first_true_index( + cudf::table_view const& when_bool_columns, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); + +} // namespace spark_rapids_jni diff --git a/src/main/java/com/nvidia/spark/rapids/jni/CaseWhen.java b/src/main/java/com/nvidia/spark/rapids/jni/CaseWhen.java new file mode 100644 index 0000000000..dabefd0a98 --- /dev/null +++ b/src/main/java/com/nvidia/spark/rapids/jni/CaseWhen.java @@ -0,0 +1,83 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +package com.nvidia.spark.rapids.jni; + +import ai.rapids.cudf.*; + + +/** + * Exedute SQL `case when` semantic. + * If there are multiple branches and each branch uses scalar to generator value, + * then it's fast to use this class because it does not generate temp string columns. + * + * E.g.: + * SQL is: + * select + * case + * when bool_1_expr then "value_1" + * when bool_2_expr then "value_2" + * when bool_3_expr then "value_3" + * else "value_else" + * end + * from tab + * + * Execution steps: + * Execute bool exprs to get bool columns, e.g., gets: + * bool column 1: [true, false, false, false] // bool_1_expr result + * bool column 2: [false, true, false, flase] // bool_2_expr result + * bool column 3: [false, false, true, flase] // bool_3_expr result + * Execute `selectFirstTrueIndex` to get the column index for the first true in bool columns. + * Generate a column to store salars: "value_1", "value_2", "value_3", "value_else" + * Execute `Table.gather` to generate the final output column + * + */ +public class CaseWhen { + + /** + * + * Select the column index for the first true in bool columns. + * For the row does not contain true, use end index(number of columns). + * + * e.g.: + * column 0: true, false, false, false + * column 1: false, true, false, false + * column 2: false, false, true, false + * + * 1st row is: true, flase, false; first true index is 0 + * 2nd row is: false, true, false; first true index is 1 + * 3rd row is: false, flase, true; first true index is 2 + * 4th row is: false, false, false; do not find true, set index to the end index 3 + * + * output column: 0, 1, 2, 3 + * In the `case when` context, here 3 index means using NULL value. + * + */ + public static ColumnVector selectFirstTrueIndex(ColumnVector[] boolColumns) { + for (ColumnVector cv : boolColumns) { + assert(cv.getType().equals(DType.BOOL8)) : "Columns must be bools"; + } + + long[] boolHandles = new long[boolColumns.length]; + for (int i = 0; i < boolColumns.length; ++i) { + boolHandles[i] = boolColumns[i].getNativeView(); + } + + return new ColumnVector(selectFirstTrueIndex(boolHandles)); + } + + private static native long selectFirstTrueIndex(long[] boolHandles); +} diff --git a/src/test/java/com/nvidia/spark/rapids/jni/CaseWhenTest.java b/src/test/java/com/nvidia/spark/rapids/jni/CaseWhenTest.java new file mode 100644 index 0000000000..f0bff026ed --- /dev/null +++ b/src/test/java/com/nvidia/spark/rapids/jni/CaseWhenTest.java @@ -0,0 +1,64 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +package com.nvidia.spark.rapids.jni; + +import ai.rapids.cudf.*; + +import org.junit.jupiter.api.Test; + +import static ai.rapids.cudf.AssertUtils.assertColumnsAreEqual; + +public class CaseWhenTest { + + @Test + void selectIndexTest() { + try ( + ColumnVector b0 = ColumnVector.fromBooleans( + true, false, false, false); + ColumnVector b1 = ColumnVector.fromBooleans( + true, true, false, false); + ColumnVector b2 = ColumnVector.fromBooleans( + false, false, true, false); + ColumnVector b3 = ColumnVector.fromBooleans( + true, true, true, false); + ColumnVector expected = ColumnVector.fromInts(0, 1, 2, 4)) { + ColumnVector[] boolColumns = new ColumnVector[] { b0, b1, b2, b3 }; + try (ColumnVector actual = CaseWhen.selectFirstTrueIndex(boolColumns)) { + assertColumnsAreEqual(expected, actual); + } + } + } + + @Test + void selectIndexTestWithNull() { + try ( + ColumnVector b0 = ColumnVector.fromBoxedBooleans( + null, false, false, null, false); + ColumnVector b1 = ColumnVector.fromBoxedBooleans( + null, null, false, true, true); + ColumnVector b2 = ColumnVector.fromBoxedBooleans( + null, null, false, true, false); + ColumnVector b3 = ColumnVector.fromBoxedBooleans( + null, null, null, true, null); + ColumnVector expected = ColumnVector.fromInts(4, 4, 4, 1, 1)) { + ColumnVector[] boolColumns = new ColumnVector[] { b0, b1, b2, b3 }; + try (ColumnVector actual = CaseWhen.selectFirstTrueIndex(boolColumns)) { + assertColumnsAreEqual(expected, actual); + } + } + } +} From 30cf3a7da07c542793a5b38f15209c77b2cf2554 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Wed, 26 Jun 2024 04:35:15 +0800 Subject: [PATCH 073/159] Update submodule cudf to cdfb550f442e846623c721082128a095f02efff9 (#2172) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index b9a0b72773..cdfb550f44 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit b9a0b72773a3adf4ba9ae267911d8970f0db53b0 +Subproject commit cdfb550f442e846623c721082128a095f02efff9 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 11842d54e2..3ac8cb579e 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -52,7 +52,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "3cc6678e16f9c41833c9e7b2779e417bbc258181", + "git_tag" : "648cae352ea26637c433e1b4fb05641b0af55dbb", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, From ee3b1cf508041fb43e64b9e8a3fadfc7d5ca93cd Mon Sep 17 00:00:00 2001 From: Jihoon Son Date: Wed, 26 Jun 2024 08:45:37 -0700 Subject: [PATCH 074/159] All kernels in the JNI should have hidden visibility (#2168) * All kernels should have hidden visibility Signed-off-by: Jihoon Son --- src/main/cpp/faultinj/faultinj.cu | 4 +-- src/main/cpp/src/bloom_filter.cu | 8 +++--- src/main/cpp/src/cast_string.cu | 35 ++++++++++++------------ src/main/cpp/src/cast_string_to_float.cu | 16 +++++------ src/main/cpp/src/parse_uri.cu | 24 ++++++++-------- 5 files changed, 43 insertions(+), 44 deletions(-) diff --git a/src/main/cpp/faultinj/faultinj.cu b/src/main/cpp/faultinj/faultinj.cu index 13065a81ed..fcb4b3a12d 100644 --- a/src/main/cpp/faultinj/faultinj.cu +++ b/src/main/cpp/faultinj/faultinj.cu @@ -136,12 +136,12 @@ CUptiResult cuptiInitialize(void) return status; } -__global__ void faultInjectorKernelAssert(void) +__global__ static void faultInjectorKernelAssert(void) { assert(0 && "faultInjectorKernelAssert triggered"); } -__global__ void faultInjectorKernelTrap(void) { asm("trap;"); } +__global__ static void faultInjectorKernelTrap(void) { asm("trap;"); } boost::optional lookupConfig( boost::optional domainConfigs, diff --git a/src/main/cpp/src/bloom_filter.cu b/src/main/cpp/src/bloom_filter.cu index 5dfdd582ef..da4e3c5cb9 100644 --- a/src/main/cpp/src/bloom_filter.cu +++ b/src/main/cpp/src/bloom_filter.cu @@ -60,10 +60,10 @@ __device__ inline std::pair gpu_get_hash_ma } template -__global__ void gpu_bloom_filter_put(cudf::bitmask_type* const bloom_filter, - cudf::size_type bloom_filter_bits, - cudf::column_device_view input, - cudf::size_type num_hashes) +CUDF_KERNEL void gpu_bloom_filter_put(cudf::bitmask_type* const bloom_filter, + cudf::size_type bloom_filter_bits, + cudf::column_device_view input, + cudf::size_type num_hashes) { size_t const tid = threadIdx.x + blockIdx.x * blockDim.x; if (tid >= input.size()) { return; } diff --git a/src/main/cpp/src/cast_string.cu b/src/main/cpp/src/cast_string.cu index bfbbc3777d..4267daae37 100644 --- a/src/main/cpp/src/cast_string.cu +++ b/src/main/cpp/src/cast_string.cu @@ -156,14 +156,14 @@ process_value(bool first_value, T current_val, T const new_digit, bool adding) * @param ansi_mode true if ansi mode is required, which is more strict and throws */ template -void __global__ string_to_integer_kernel(T* out, - bitmask_type* validity, - const char* const chars, - size_type const* offsets, - bitmask_type const* incoming_null_mask, - size_type num_rows, - bool ansi_mode, - bool strip) +void CUDF_KERNEL string_to_integer_kernel(T* out, + bitmask_type* validity, + const char* const chars, + size_type const* offsets, + bitmask_type const* incoming_null_mask, + size_type num_rows, + bool ansi_mode, + bool strip) { auto const group = cooperative_groups::this_thread_block(); auto const warp = cooperative_groups::tiled_partition(group); @@ -386,18 +386,17 @@ __device__ thrust::optional> validate_and_exponent * @param scale scale of desired decimals * @param precision precision of desired decimals * @param ansi_mode true if ansi mode is required, which is more strict and throws - * @return __global__ */ template -__global__ void string_to_decimal_kernel(T* out, - bitmask_type* validity, - const char* const chars, - size_type const* offsets, - bitmask_type const* incoming_null_mask, - size_type num_rows, - int32_t scale, - int32_t precision, - bool strip) +CUDF_KERNEL void string_to_decimal_kernel(T* out, + bitmask_type* validity, + const char* const chars, + size_type const* offsets, + bitmask_type const* incoming_null_mask, + size_type num_rows, + int32_t scale, + int32_t precision, + bool strip) { auto const group = cooperative_groups::this_thread_block(); auto const warp = cooperative_groups::tiled_partition(group); diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index cd7de88110..c19a2a10fe 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -618,14 +618,14 @@ class string_to_float { }; template -__global__ void string_to_float_kernel(T* out, - bitmask_type* validity, - int32_t* ansi_except, - size_type* valid_count, - const char* const chars, - size_type const* offsets, - bitmask_type const* incoming_null_mask, - size_type const num_rows) +CUDF_KERNEL void string_to_float_kernel(T* out, + bitmask_type* validity, + int32_t* ansi_except, + size_type* valid_count, + const char* const chars, + size_type const* offsets, + bitmask_type const* incoming_null_mask, + size_type const num_rows) { size_type const tid = threadIdx.x + (blockDim.x * blockIdx.x); size_type const row = tid / 32; diff --git a/src/main/cpp/src/parse_uri.cu b/src/main/cpp/src/parse_uri.cu index 0e57366358..f0a78f4f52 100644 --- a/src/main/cpp/src/parse_uri.cu +++ b/src/main/cpp/src/parse_uri.cu @@ -770,13 +770,13 @@ uri_parts __device__ validate_uri(const char* str, * @param out_offsets Offsets to the start of the chunks * @param out_validity Bitmask of validity data, updated in function */ -__global__ void parse_uri_char_counter(column_device_view const in_strings, - URI_chunks chunk, - char const* const base_ptr, - size_type* const out_lengths, - size_type* const out_offsets, - bitmask_type* out_validity, - thrust::optional query_match) +CUDF_KERNEL void parse_uri_char_counter(column_device_view const in_strings, + URI_chunks chunk, + char const* const base_ptr, + size_type* const out_lengths, + size_type* const out_offsets, + bitmask_type* out_validity, + thrust::optional query_match) { // thread per row auto const tid = cudf::detail::grid_1d::global_thread_id(); @@ -850,11 +850,11 @@ __global__ void parse_uri_char_counter(column_device_view const in_strings, * @param offsets Offset value of each string associated with `out_chars` * @param out_chars Character buffer for the output string column */ -__global__ void parse_uri(column_device_view const in_strings, - char const* const base_ptr, - size_type const* const src_offsets, - size_type const* const offsets, - char* const out_chars) +CUDF_KERNEL void parse_uri(column_device_view const in_strings, + char const* const base_ptr, + size_type const* const src_offsets, + size_type const* const offsets, + char* const out_chars) { auto const tid = cudf::detail::grid_1d::global_thread_id(); From e0bf15efce0f211c3e8e0e3b0972c97412019ba1 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 27 Jun 2024 04:33:47 +0800 Subject: [PATCH 075/159] Update submodule cudf to e7cf69dc932636e9dfd32ea1bebefddc3f31d2f2 (#2175) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index cdfb550f44..e7cf69dc93 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit cdfb550f442e846623c721082128a095f02efff9 +Subproject commit e7cf69dc932636e9dfd32ea1bebefddc3f31d2f2 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 3ac8cb579e..8608894ec0 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -137,7 +137,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "1ab4920e50a07a97ec717db514ad9a3dccc9939a", + "git_tag" : "cc53275be319487d568f8d89bc14629efef67c87", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From afa28bcb613a3c29f5c0c272e263ca3d8c36a8b0 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 27 Jun 2024 10:30:00 +0800 Subject: [PATCH 076/159] Update submodule cudf to 6eac9207ca0804aeca64c83c533e16ad5963b0ba (#2176) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index e7cf69dc93..6eac9207ca 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit e7cf69dc932636e9dfd32ea1bebefddc3f31d2f2 +Subproject commit 6eac9207ca0804aeca64c83c533e16ad5963b0ba From 8410e8e07ae84a3bdf760721818229bf5f991c63 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 27 Jun 2024 16:32:49 +0800 Subject: [PATCH 077/159] Update submodule cudf to f267b1f068ec3e8fd49599fc28afa2fc0464118b (#2179) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 6eac9207ca..f267b1f068 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 6eac9207ca0804aeca64c83c533e16ad5963b0ba +Subproject commit f267b1f068ec3e8fd49599fc28afa2fc0464118b diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 8608894ec0..d2924dcb1c 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -52,7 +52,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "648cae352ea26637c433e1b4fb05641b0af55dbb", + "git_tag" : "157a29a8530494998642a9983aef27f06bd94b67", "git_url" : "https://github.com/rapidsai/kvikio.git", "version" : "24.08" }, @@ -137,7 +137,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "cc53275be319487d568f8d89bc14629efef67c87", + "git_tag" : "0f608ecfb16a5b82d5c7f2cfa59f5085b041c7d9", "git_url" : "https://github.com/rapidsai/rmm.git", "version" : "24.08" }, From 7b6e12b8b599bf521274254febd29b739a07ad22 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Thu, 27 Jun 2024 20:30:38 +0800 Subject: [PATCH 078/159] Update submodule cudf to fa8284ddb2de808573d5b21cc9e650578ddf6acc (#2180) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index f267b1f068..fa8284ddb2 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit f267b1f068ec3e8fd49599fc28afa2fc0464118b +Subproject commit fa8284ddb2de808573d5b21cc9e650578ddf6acc From fe3e432dee9375a0d5748cb4e3c941c79c704a7c Mon Sep 17 00:00:00 2001 From: Mike Wilson Date: Thu, 27 Jun 2024 14:34:13 -0400 Subject: [PATCH 079/159] Flipping MACRO ordering to match style (#2178) Signed-off-by: Mike Wilson --- src/main/cpp/src/cast_string.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/main/cpp/src/cast_string.cu b/src/main/cpp/src/cast_string.cu index 4267daae37..156dbeb7bf 100644 --- a/src/main/cpp/src/cast_string.cu +++ b/src/main/cpp/src/cast_string.cu @@ -156,7 +156,7 @@ process_value(bool first_value, T current_val, T const new_digit, bool adding) * @param ansi_mode true if ansi mode is required, which is more strict and throws */ template -void CUDF_KERNEL string_to_integer_kernel(T* out, +CUDF_KERNEL void string_to_integer_kernel(T* out, bitmask_type* validity, const char* const chars, size_type const* offsets, From 2fdb27cf4da5eae0513272bc7a8eb376f7d09bee Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 28 Jun 2024 04:31:49 +0800 Subject: [PATCH 080/159] Update submodule cudf to a71c249f9f320ecb61aa8135bbda300122e43491 (#2181) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index fa8284ddb2..a71c249f9f 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit fa8284ddb2de808573d5b21cc9e650578ddf6acc +Subproject commit a71c249f9f320ecb61aa8135bbda300122e43491 From ac06522b5e9a1b1aeb13c934d045e88bbe8c51ec Mon Sep 17 00:00:00 2001 From: Peixin Date: Fri, 28 Jun 2024 07:45:53 +0800 Subject: [PATCH 081/159] Improve cuda dependencies check during build (#2173) * Check cuda dependencies after build Signed-off-by: Peixin Li * udpate copyright Signed-off-by: Peixin Li * update script to check cudart only for all *.so * fix output --------- Signed-off-by: Peixin Li --- ci/check-cuda-dependencies.sh | 35 +++++++++++++++++++++++++++++++++++ ci/nightly-build.sh | 7 ++++++- ci/premerge-build.sh | 6 +++++- ci/submodule-sync.sh | 4 ++++ 4 files changed, 50 insertions(+), 2 deletions(-) create mode 100644 ci/check-cuda-dependencies.sh diff --git a/ci/check-cuda-dependencies.sh b/ci/check-cuda-dependencies.sh new file mode 100644 index 0000000000..9d988bedae --- /dev/null +++ b/ci/check-cuda-dependencies.sh @@ -0,0 +1,35 @@ +#!/bin/bash +# +# Copyright (c) 2024, NVIDIA CORPORATION. All rights reserved. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +# + +# common script to help check if packaged *.so files have dynamical link to CUDA Runtime + +set -exo pipefail + +jar_path=$1 +tmp_path=/tmp/"jni-$(date "+%Y%m%d%H%M%S")" +unzip -j "${jar_path}" "*64/Linux/*.so" -d "${tmp_path}" + +find "$tmp_path" -type f -name "*.so" | while read -r so_file; do + # Check if *.so file has a dynamic link to CUDA Runtime + if objdump -p "$so_file" | grep NEEDED | grep -qi cudart; then + echo "Dynamic link to CUDA Runtime found in $so_file..." + ldd "$so_file" + exit 1 + else + echo "No dynamic link to CUDA Runtime found in $so_file" + fi +done diff --git a/ci/nightly-build.sh b/ci/nightly-build.sh index 8a3c2dbacf..afa2b6a842 100755 --- a/ci/nightly-build.sh +++ b/ci/nightly-build.sh @@ -1,6 +1,6 @@ #!/bin/bash # -# Copyright (c) 2022-2023, NVIDIA CORPORATION. All rights reserved. +# Copyright (c) 2022-2024, NVIDIA CORPORATION. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -29,6 +29,7 @@ USE_GDS=${USE_GDS:-ON} USE_SANITIZER=${USE_SANITIZER:-ON} BUILD_FAULTINJ=${BUILD_FAULTINJ:-ON} ARM64=${ARM64:-false} +artifact_suffix="${CUDA_VER}" profiles="source-javadoc" if [ "${ARM64}" == "true" ]; then @@ -36,6 +37,7 @@ if [ "${ARM64}" == "true" ]; then USE_GDS="OFF" USE_SANITIZER="ON" BUILD_FAULTINJ="OFF" + artifact_suffix="${artifact_suffix}-arm64" fi ${MVN} clean package ${MVN_MIRROR} \ @@ -45,3 +47,6 @@ ${MVN} clean package ${MVN_MIRROR} \ -DUSE_GDS=${USE_GDS} -Dtest=*,!CuFileTest,!CudaFatalTest,!ColumnViewNonEmptyNullsTest \ -DBUILD_TESTS=ON -DBUILD_FAULTINJ=${BUILD_FAULTINJ} -Dcuda.version=$CUDA_VER \ -DUSE_SANITIZER=${USE_SANITIZER} + +build_name=$(${MVN} help:evaluate -Dexpression=project.build.finalName -q -DforceStdout) +. ci/check-cuda-dependencies.sh "target/${build_name}-${artifact_suffix}.jar" diff --git a/ci/premerge-build.sh b/ci/premerge-build.sh index e3adc10b3e..052dd7d6cc 100755 --- a/ci/premerge-build.sh +++ b/ci/premerge-build.sh @@ -1,6 +1,6 @@ #!/bin/bash # -# Copyright (c) 2022-2023, NVIDIA CORPORATION. All rights reserved. +# Copyright (c) 2022-2024, NVIDIA CORPORATION. All rights reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -28,3 +28,7 @@ ${MVN} verify ${MVN_MIRROR} \ -Dlibcudf.build.configure=true \ -DUSE_GDS=ON -Dtest=*,!CuFileTest,!CudaFatalTest,!ColumnViewNonEmptyNullsTest \ -DBUILD_TESTS=ON + +build_name=$(${MVN} help:evaluate -Dexpression=project.build.finalName -q -DforceStdout) +cuda_version=$(${MVN} help:evaluate -Dexpression=cuda.version -q -DforceStdout) +. ci/check-cuda-dependencies.sh "target/${build_name}-${cuda_version}.jar" diff --git a/ci/submodule-sync.sh b/ci/submodule-sync.sh index 9794d40a0f..a889d86eb0 100755 --- a/ci/submodule-sync.sh +++ b/ci/submodule-sync.sh @@ -88,6 +88,10 @@ else echo "Test failed, will update the result" fi +build_name=$(${MVN} help:evaluate -Dexpression=project.build.finalName -q -DforceStdout) +cuda_version=$(${MVN} help:evaluate -Dexpression=cuda.version -q -DforceStdout) +. ci/check-cuda-dependencies.sh "target/${build_name}-${cuda_version}.jar" + LIBCUDF_BUILD_PATH=$(${MVN} help:evaluate -Dexpression=libcudf.build.path -q -DforceStdout) # Extract the rapids-cmake sha1 that we need to pin too rapids_cmake_sha=$(git -C ${LIBCUDF_BUILD_PATH}/_deps/rapids-cmake-src/ rev-parse HEAD) From 7cddb42f1451b52513bdd1f96b28eaecefcdb1b9 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 28 Jun 2024 08:54:36 +0800 Subject: [PATCH 082/159] Update submodule cudf to c847b98291bd41f98ac417becf0c53293a392ce3 (#2182) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/rapids-cmake.sha | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index a71c249f9f..c847b98291 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit a71c249f9f320ecb61aa8135bbda300122e43491 +Subproject commit c847b98291bd41f98ac417becf0c53293a392ce3 diff --git a/thirdparty/cudf-pins/rapids-cmake.sha b/thirdparty/cudf-pins/rapids-cmake.sha index 0dcb044f5f..91edbb7134 100644 --- a/thirdparty/cudf-pins/rapids-cmake.sha +++ b/thirdparty/cudf-pins/rapids-cmake.sha @@ -1 +1 @@ -5d47273ee4f64ec2710a6f3e7c9d9364813dcfda +fa6d25b25ebfd0fcaf5e644ec7d12e0aaa9c7d6e From 83b24db9dfaf5dadbe4bda590a85ff8f91a1c0e2 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 28 Jun 2024 21:20:42 +0800 Subject: [PATCH 083/159] Update submodule cudf to 6b04fd3b704efdae7d39d09beba026fcbca5f996 (#2183) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index c847b98291..6b04fd3b70 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit c847b98291bd41f98ac417becf0c53293a392ce3 +Subproject commit 6b04fd3b704efdae7d39d09beba026fcbca5f996 From 25fa8125f4b6b9073f168bb02429baf23a01ca8b Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Fri, 28 Jun 2024 23:14:52 +0800 Subject: [PATCH 084/159] Update submodule cudf to 224ac5bad11465d0486af80e7935eac482269805 (#2184) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 6b04fd3b70..224ac5bad1 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 6b04fd3b704efdae7d39d09beba026fcbca5f996 +Subproject commit 224ac5bad11465d0486af80e7935eac482269805 From 8cb3be44d4948415656f031aec06eada8df39ca5 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 29 Jun 2024 04:31:18 +0800 Subject: [PATCH 085/159] Update submodule cudf to fb12d980342833a9d7092a19717eedad22328e6a (#2185) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 224ac5bad1..fb12d98034 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 224ac5bad11465d0486af80e7935eac482269805 +Subproject commit fb12d980342833a9d7092a19717eedad22328e6a From 1baa4c4f1f2c9de089981830ab0849e3cc18e383 Mon Sep 17 00:00:00 2001 From: Jenkins Automation <70000568+nvauto@users.noreply.github.com> Date: Sat, 29 Jun 2024 10:30:18 +0800 Subject: [PATCH 086/159] Update submodule cudf to 3c3edfef406288e164cc80ab82f9c64c0b88d0bd (#2186) Signed-off-by: spark-rapids automation <70000568+nvauto@users.noreply.github.com> --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index fb12d98034..3c3edfef40 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit fb12d980342833a9d7092a19717eedad22328e6a +Subproject commit 3c3edfef406288e164cc80ab82f9c64c0b88d0bd From 7aad93b4cc6be944fe559c923aacdc8033c94871 Mon Sep 17 00:00:00 2001 From: Liangcai Li Date: Sat, 29 Jun 2024 15:53:56 +0800 Subject: [PATCH 087/159] Add support for hive hash (#2171) This PR adds the support for Hive hash with the basic types. The supported types list is [bool, byte, short, int, long, string, date, timestamp, float, double]. --------- Signed-off-by: Firestarman Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com> Co-authored-by: MithunR --- src/main/cpp/CMakeLists.txt | 1 + src/main/cpp/src/HashJni.cpp | 17 +- src/main/cpp/src/hash.cuh | 39 --- src/main/cpp/src/hash.hpp | 75 ++++++ src/main/cpp/src/hive_hash.cu | 255 ++++++++++++++++++ src/main/cpp/tests/hash.cpp | 2 +- .../com/nvidia/spark/rapids/jni/Hash.java | 21 +- .../com/nvidia/spark/rapids/jni/HashTest.java | 130 ++++++++- 8 files changed, 494 insertions(+), 46 deletions(-) create mode 100644 src/main/cpp/src/hash.hpp create mode 100644 src/main/cpp/src/hive_hash.cu diff --git a/src/main/cpp/CMakeLists.txt b/src/main/cpp/CMakeLists.txt index 2de79a74c1..3c7181819c 100644 --- a/src/main/cpp/CMakeLists.txt +++ b/src/main/cpp/CMakeLists.txt @@ -217,6 +217,7 @@ add_library( src/timezones.cu src/utilities.cu src/xxhash64.cu + src/hive_hash.cu src/zorder.cu ) diff --git a/src/main/cpp/src/HashJni.cpp b/src/main/cpp/src/HashJni.cpp index 9e556cdd2d..c0adf38686 100644 --- a/src/main/cpp/src/HashJni.cpp +++ b/src/main/cpp/src/HashJni.cpp @@ -16,7 +16,7 @@ #include "cudf_jni_apis.hpp" #include "dtype_utils.hpp" -#include "hash.cuh" +#include "hash.hpp" #include "jni_utils.hpp" extern "C" { @@ -52,4 +52,19 @@ JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_Hash_xxhash64(JNIEnv* e } CATCH_STD(env, 0); } + +JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_Hash_hiveHash(JNIEnv* env, + jclass, + jlongArray column_handles) +{ + JNI_NULL_CHECK(env, column_handles, "array of column handles is null", 0); + + try { + cudf::jni::auto_set_device(env); + auto column_views = + cudf::jni::native_jpointerArray{env, column_handles}.get_dereferenced(); + return cudf::jni::release_as_jlong(spark_rapids_jni::hive_hash(cudf::table_view{column_views})); + } + CATCH_STD(env, 0); +} } diff --git a/src/main/cpp/src/hash.cuh b/src/main/cpp/src/hash.cuh index 8cf489a7e7..35969ca42a 100644 --- a/src/main/cpp/src/hash.cuh +++ b/src/main/cpp/src/hash.cuh @@ -19,15 +19,9 @@ #include #include -#include -#include - #include namespace spark_rapids_jni { - -constexpr int64_t DEFAULT_XXHASH64_SEED = 42; - /** * Normalization of floating point NaNs, passthrough for all other values. */ @@ -101,37 +95,4 @@ __device__ __inline__ std::pair<__int128_t, cudf::size_type> to_java_bigdecimal( return {big_endian_value, length}; } - -/** - * @brief Computes the murmur32 hash value of each row in the input set of columns. - * - * @param input The table of columns to hash - * @param seed Optional seed value to use for the hash function - * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the returned column's device memory - * - * @returns A column where each row is the hash of a column from the input. - */ -std::unique_ptr murmur_hash3_32( - cudf::table_view const& input, - uint32_t seed = 0, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); - -/** - * @brief Computes the xxhash64 hash value of each row in the input set of columns. - * - * @param input The table of columns to hash - * @param seed Optional seed value to use for the hash function - * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the returned column's device memory - * - * @returns A column where each row is the hash of a column from the input. - */ -std::unique_ptr xxhash64( - cudf::table_view const& input, - int64_t seed = DEFAULT_XXHASH64_SEED, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); - } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/hash.hpp b/src/main/cpp/src/hash.hpp new file mode 100644 index 0000000000..4021b9e75c --- /dev/null +++ b/src/main/cpp/src/hash.hpp @@ -0,0 +1,75 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include + +#include +#include + +namespace spark_rapids_jni { + +constexpr int64_t DEFAULT_XXHASH64_SEED = 42; + +/** + * @brief Computes the murmur32 hash value of each row in the input set of columns. + * + * @param input The table of columns to hash + * @param seed Optional seed value to use for the hash function + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * + * @returns A column where each row is the hash of a column from the input. + */ +std::unique_ptr murmur_hash3_32( + cudf::table_view const& input, + uint32_t seed = 0, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Computes the xxhash64 hash value of each row in the input set of columns. + * + * @param input The table of columns to hash + * @param seed Optional seed value to use for the hash function + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * + * @returns A column where each row is the hash of a column from the input. + */ +std::unique_ptr xxhash64( + cudf::table_view const& input, + int64_t seed = DEFAULT_XXHASH64_SEED, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); + +/** + * @brief Computes the Hive hash value of each row in the input set of columns. + * + * @param input The table of columns to hash + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * + * @returns A column where each row is the hash of a column from the input. + */ +std::unique_ptr hive_hash( + cudf::table_view const& input, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); + +} // namespace spark_rapids_jni diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu new file mode 100644 index 0000000000..85598565a9 --- /dev/null +++ b/src/main/cpp/src/hive_hash.cu @@ -0,0 +1,255 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "hash.cuh" + +#include +#include +#include +#include + +#include +#include + +#include +#include +#include + +namespace spark_rapids_jni { + +namespace { + +using hive_hash_value_t = int32_t; + +constexpr hive_hash_value_t HIVE_HASH_FACTOR = 31; +constexpr hive_hash_value_t HIVE_INIT_HASH = 0; + +hive_hash_value_t __device__ inline compute_int(int32_t key) { return key; } + +hive_hash_value_t __device__ inline compute_long(int64_t key) +{ + return (static_cast(key) >> 32) ^ key; +} + +hive_hash_value_t __device__ inline compute_bytes(int8_t const* data, cudf::size_type const len) +{ + hive_hash_value_t ret = HIVE_INIT_HASH; + for (auto i = 0; i < len; i++) { + ret = ret * HIVE_HASH_FACTOR + static_cast(data[i]); + } + return ret; +} + +template +struct hive_hash_function { + // 'seed' is not used in 'hive_hash_function', but required by 'element_hasher'. + constexpr hive_hash_function(uint32_t) {} + + [[nodiscard]] hive_hash_value_t __device__ inline operator()(Key const& key) const + { + CUDF_UNREACHABLE("Unsupported type for hive hash"); + } +}; // struct hive_hash_function + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()( + cudf::string_view const& key) const +{ + auto const data = reinterpret_cast(key.data()); + auto const len = key.size_bytes(); + return compute_bytes(data, len); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()(bool const& key) const +{ + return compute_int(static_cast(key)); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()(int8_t const& key) const +{ + return compute_int(static_cast(key)); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()( + int16_t const& key) const +{ + return compute_int(static_cast(key)); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()( + int32_t const& key) const +{ + return compute_int(key); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()( + int64_t const& key) const +{ + return compute_long(key); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()(float const& key) const +{ + auto normalized = spark_rapids_jni::normalize_nans(key); + auto* p_int = reinterpret_cast(&normalized); + return compute_int(*p_int); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()(double const& key) const +{ + auto normalized = spark_rapids_jni::normalize_nans(key); + auto* p_long = reinterpret_cast(&normalized); + return compute_long(*p_long); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()( + cudf::timestamp_D const& key) const +{ + auto* p_int = reinterpret_cast(&key); + return compute_int(*p_int); +} + +template <> +hive_hash_value_t __device__ inline hive_hash_function::operator()( + cudf::timestamp_us const& key) const +{ + auto time_as_long = *reinterpret_cast(&key); + constexpr int MICRO_PER_SEC = 1000000; + constexpr int NANO_PER_MICRO = 1000; + + int64_t ts = time_as_long / MICRO_PER_SEC; + int64_t tns = (time_as_long % MICRO_PER_SEC) * NANO_PER_MICRO; + + int64_t result = ts; + result <<= 30; + result |= tns; + + result = (static_cast(result) >> 32) ^ result; + return static_cast(result); +} + +/** + * @brief Computes the hash value of a row in the given table. + * + * This functor produces the same result as "HiveHash" in Spark for supported types. + * + * @tparam hash_function Hash functor to use for hashing elements. Must be hive_hash_function. + * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + */ +template