From b5dcc7c620f7805d73fc850d056664ef5da739b2 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Fri, 14 Aug 2026 22:51:30 +0000 Subject: [PATCH 1/7] Test rapids-cmake cuco bump --- cmake/rapids_config.cmake | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/cmake/rapids_config.cmake b/cmake/rapids_config.cmake index c9b1967245ac..18da8a799972 100644 --- a/cmake/rapids_config.cmake +++ b/cmake/rapids_config.cmake @@ -1,6 +1,6 @@ # ============================================================================= # cmake-format: off -# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 # cmake-format: on # ============================================================================= @@ -34,6 +34,10 @@ endif() if(NOT rapids-cmake-branch) set(rapids-cmake-branch "${RAPIDS_BRANCH}") endif() + +# Test the cuco bump from PointKernel/rapids-cmake +set(rapids-cmake-repo "PointKernel/rapids-cmake") +set(rapids-cmake-branch "bump-cuco-4b26118") include("${CMAKE_CURRENT_LIST_DIR}/RAPIDS.cmake") # Don't use sccache-dist for CMake's compiler tests From 5ed54f5457691e99b4fd247c12004d5406bf2d9b Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Sat, 15 Aug 2026 01:46:33 +0000 Subject: [PATCH 2/7] Migrate to cuco Bloom filter policy --- cpp/include/cudf/reduction/bloom_filter.cuh | 9 +++++---- cpp/libcudf_streaming/src/detail/device_bloom_filter.cu | 2 +- .../tests/streaming/test_bloom_filter.cu | 2 +- cpp/src/io/parquet/bloom_filter_reader.cu | 4 ++-- cpp/src/join/mark_join.cuh | 2 +- cpp/tests/io/parquet_bloom_filter_test.cu | 7 ++++--- 6 files changed, 14 insertions(+), 12 deletions(-) diff --git a/cpp/include/cudf/reduction/bloom_filter.cuh b/cpp/include/cudf/reduction/bloom_filter.cuh index c8c7b3d56df8..97d70d9b79fb 100644 --- a/cpp/include/cudf/reduction/bloom_filter.cuh +++ b/cpp/include/cudf/reduction/bloom_filter.cuh @@ -5,7 +5,7 @@ #pragma once -#include +#include #include @@ -14,14 +14,15 @@ namespace cudf { /** * @brief Policy describing the Apache Arrow Block-Split Bloom Filter layout. * - * Uses cuco's `parametric_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x + * Uses cuco's `bloom_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x * `uint32_t`), 8 fingerprint bits per key, fully horizontal add (Theta=8), and fully vertical * contains (Phi=8). This layout is bit-compatible with Apache Arrow. * + * @tparam Key The key type to generate a fingerprint for. * @tparam Hash The hash function used to generate a hash for each key. */ -template +template using arrow_filter_policy = - cuco::parametric_filter_policy; + cuco::bloom_filter_policy; } // namespace cudf diff --git a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu index ad3a2217ba13..d6ad8a116174 100644 --- a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu +++ b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu @@ -55,7 +55,7 @@ namespace cudf_streaming::detail { namespace { using KeyType = std::uint64_t; -using BloomFilterPolicy = cudf::arrow_filter_policy>; +using BloomFilterPolicy = cudf::arrow_filter_policy>; using BloomFilterRefType = cuco::bloom_filter_ref, cuco::thread_scope_device, diff --git a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu index b8fc5b2e978a..906ed416438f 100644 --- a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu +++ b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu @@ -25,7 +25,7 @@ namespace { -using policy_type = cudf::arrow_filter_policy>; +using policy_type = cudf::arrow_filter_policy>; __global__ void block_index_kernel(std::uint32_t upper_hash, std::size_t num_blocks, diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 3418c8ca9ada..a7605f69c6a3 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -45,7 +45,7 @@ namespace { * `XXHash_64` (so that `cudf::string_view` and other cudf types are hashed by content, matching the * Apache Parquet/Arrow bloom filter specification). * - * Uses cuco's `parametric_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x + * Uses cuco's `bloom_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x * `uint32_t`), 8 fingerprint bits per key, fully horizontal add (Theta=8) and fully vertical * contains (Phi=8). This layout is bit-compatible with Apache Arrow, as verified by cuCollections * `tests/bloom_filter/arrow_compat_test.cu`. @@ -53,7 +53,7 @@ namespace { * @tparam Key The type of the values to generate a fingerprint for. */ template -using arrow_filter_policy = cudf::arrow_filter_policy>; +using arrow_filter_policy = cudf::arrow_filter_policy>; /** * @brief Converts bloom filter membership results (for each column chunk) to a device column. diff --git a/cpp/src/join/mark_join.cuh b/cpp/src/join/mark_join.cuh index 6a3179f3b493..2a26946b9eb2 100644 --- a/cpp/src/join/mark_join.cuh +++ b/cpp/src/join/mark_join.cuh @@ -165,7 +165,7 @@ using storage_ref_type = cuco::bucket_storage_ref>; using right_key_type = cuco::pair; -using bloom_filter_policy_type = cuco::default_filter_policy; +using bloom_filter_policy_type = cuco::bloom_filter_policy; using bloom_filter_allocator_type = rmm::mr::polymorphic_allocator; using bloom_filter_type = cuco::bloom_filter, diff --git a/cpp/tests/io/parquet_bloom_filter_test.cu b/cpp/tests/io/parquet_bloom_filter_test.cu index 7fbb21714eff..adef120d0d6d 100644 --- a/cpp/tests/io/parquet_bloom_filter_test.cu +++ b/cpp/tests/io/parquet_bloom_filter_test.cu @@ -26,9 +26,10 @@ class ParquetBloomFilterTest : public cudf::test::BaseFixture {}; TEST_F(ParquetBloomFilterTest, TestStrings) { - using key_type = StringType; - using policy_type = cudf::arrow_filter_policy>; - using word_type = policy_type::word_type; + using key_type = StringType; + using policy_type = + cudf::arrow_filter_policy>; + using word_type = policy_type::word_type; std::size_t constexpr num_filter_blocks = 4; auto stream = cudf::get_default_stream(); From 783fd9c50e9b542f6a5e76894fe6e184feb26121 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 17 Aug 2026 16:04:14 +0000 Subject: [PATCH 3/7] Use parametric Bloom filter policy --- cpp/include/cudf/reduction/bloom_filter.cuh | 7 +++---- .../src/detail/device_bloom_filter.cu | 2 +- .../tests/streaming/test_bloom_filter.cu | 2 +- cpp/src/io/parquet/bloom_filter_reader.cu | 4 ++-- cpp/src/join/mark_join.cuh | 11 ++++++++++- cpp/tests/io/parquet_bloom_filter_test.cu | 7 +++---- 6 files changed, 20 insertions(+), 13 deletions(-) diff --git a/cpp/include/cudf/reduction/bloom_filter.cuh b/cpp/include/cudf/reduction/bloom_filter.cuh index 97d70d9b79fb..16e932bd223e 100644 --- a/cpp/include/cudf/reduction/bloom_filter.cuh +++ b/cpp/include/cudf/reduction/bloom_filter.cuh @@ -14,15 +14,14 @@ namespace cudf { /** * @brief Policy describing the Apache Arrow Block-Split Bloom Filter layout. * - * Uses cuco's `bloom_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x + * Uses cuco's `parametric_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x * `uint32_t`), 8 fingerprint bits per key, fully horizontal add (Theta=8), and fully vertical * contains (Phi=8). This layout is bit-compatible with Apache Arrow. * - * @tparam Key The key type to generate a fingerprint for. * @tparam Hash The hash function used to generate a hash for each key. */ -template +template using arrow_filter_policy = - cuco::bloom_filter_policy; + cuco::parametric_filter_policy; } // namespace cudf diff --git a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu index d6ad8a116174..ad3a2217ba13 100644 --- a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu +++ b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu @@ -55,7 +55,7 @@ namespace cudf_streaming::detail { namespace { using KeyType = std::uint64_t; -using BloomFilterPolicy = cudf::arrow_filter_policy>; +using BloomFilterPolicy = cudf::arrow_filter_policy>; using BloomFilterRefType = cuco::bloom_filter_ref, cuco::thread_scope_device, diff --git a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu index 906ed416438f..b8fc5b2e978a 100644 --- a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu +++ b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu @@ -25,7 +25,7 @@ namespace { -using policy_type = cudf::arrow_filter_policy>; +using policy_type = cudf::arrow_filter_policy>; __global__ void block_index_kernel(std::uint32_t upper_hash, std::size_t num_blocks, diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index a7605f69c6a3..3418c8ca9ada 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -45,7 +45,7 @@ namespace { * `XXHash_64` (so that `cudf::string_view` and other cudf types are hashed by content, matching the * Apache Parquet/Arrow bloom filter specification). * - * Uses cuco's `bloom_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x + * Uses cuco's `parametric_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x * `uint32_t`), 8 fingerprint bits per key, fully horizontal add (Theta=8) and fully vertical * contains (Phi=8). This layout is bit-compatible with Apache Arrow, as verified by cuCollections * `tests/bloom_filter/arrow_compat_test.cu`. @@ -53,7 +53,7 @@ namespace { * @tparam Key The type of the values to generate a fingerprint for. */ template -using arrow_filter_policy = cudf::arrow_filter_policy>; +using arrow_filter_policy = cudf::arrow_filter_policy>; /** * @brief Converts bloom filter membership results (for each column chunk) to a device column. diff --git a/cpp/src/join/mark_join.cuh b/cpp/src/join/mark_join.cuh index 2a26946b9eb2..887b30ff7822 100644 --- a/cpp/src/join/mark_join.cuh +++ b/cpp/src/join/mark_join.cuh @@ -165,7 +165,16 @@ using storage_ref_type = cuco::bucket_storage_ref>; using right_key_type = cuco::pair; -using bloom_filter_policy_type = cuco::bloom_filter_policy; +using bloom_filter_policy_type = cuco::parametric_filter_policy, + std::uint32_t, + 8, + 8, + 8, + 1, + 1, + 8, + false, + false>; using bloom_filter_allocator_type = rmm::mr::polymorphic_allocator; using bloom_filter_type = cuco::bloom_filter, diff --git a/cpp/tests/io/parquet_bloom_filter_test.cu b/cpp/tests/io/parquet_bloom_filter_test.cu index adef120d0d6d..7fbb21714eff 100644 --- a/cpp/tests/io/parquet_bloom_filter_test.cu +++ b/cpp/tests/io/parquet_bloom_filter_test.cu @@ -26,10 +26,9 @@ class ParquetBloomFilterTest : public cudf::test::BaseFixture {}; TEST_F(ParquetBloomFilterTest, TestStrings) { - using key_type = StringType; - using policy_type = - cudf::arrow_filter_policy>; - using word_type = policy_type::word_type; + using key_type = StringType; + using policy_type = cudf::arrow_filter_policy>; + using word_type = policy_type::word_type; std::size_t constexpr num_filter_blocks = 4; auto stream = cudf::get_default_stream(); From 929fa8dcb52c24a7745fe09e611dd0c05538a219 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 17 Aug 2026 16:05:16 +0000 Subject: [PATCH 4/7] Remove temporary rapids-cmake override --- cmake/rapids_config.cmake | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/cmake/rapids_config.cmake b/cmake/rapids_config.cmake index 18da8a799972..c9b1967245ac 100644 --- a/cmake/rapids_config.cmake +++ b/cmake/rapids_config.cmake @@ -1,6 +1,6 @@ # ============================================================================= # cmake-format: off -# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. # SPDX-License-Identifier: Apache-2.0 # cmake-format: on # ============================================================================= @@ -34,10 +34,6 @@ endif() if(NOT rapids-cmake-branch) set(rapids-cmake-branch "${RAPIDS_BRANCH}") endif() - -# Test the cuco bump from PointKernel/rapids-cmake -set(rapids-cmake-repo "PointKernel/rapids-cmake") -set(rapids-cmake-branch "bump-cuco-4b26118") include("${CMAKE_CURRENT_LIST_DIR}/RAPIDS.cmake") # Don't use sccache-dist for CMake's compiler tests From 30e53f6712cb2a7bca0bdb19339937f7903170d3 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 17 Aug 2026 16:09:36 +0000 Subject: [PATCH 5/7] Reuse Arrow Bloom filter policy in mark join --- cpp/src/join/mark_join.cuh | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/cpp/src/join/mark_join.cuh b/cpp/src/join/mark_join.cuh index be9e10b344cd..f8c6e93d1dfd 100644 --- a/cpp/src/join/mark_join.cuh +++ b/cpp/src/join/mark_join.cuh @@ -10,6 +10,7 @@ #include #include #include +#include #include #include #include @@ -165,16 +166,7 @@ using storage_ref_type = cuco::bucket_storage_ref>; using right_key_type = cuco::pair; -using bloom_filter_policy_type = cuco::parametric_filter_policy, - std::uint32_t, - 8, - 8, - 8, - 1, - 1, - 8, - false, - false>; +using bloom_filter_policy_type = cudf::arrow_filter_policy>; using bloom_filter_allocator_type = rmm::mr::polymorphic_allocator; using bloom_filter_type = cuco::bloom_filter, From a2c364444736852c14f9425a746d9cc224444d00 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 17 Aug 2026 16:29:02 +0000 Subject: [PATCH 6/7] Adopt new cuco Bloom filter policy API --- cpp/include/cudf/reduction/bloom_filter.cuh | 7 ++++--- cpp/libcudf_streaming/src/detail/device_bloom_filter.cu | 2 +- cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu | 2 +- cpp/src/io/parquet/bloom_filter_reader.cu | 4 ++-- cpp/src/join/mark_join.cuh | 3 ++- cpp/tests/io/parquet_bloom_filter_test.cu | 7 ++++--- 6 files changed, 14 insertions(+), 11 deletions(-) diff --git a/cpp/include/cudf/reduction/bloom_filter.cuh b/cpp/include/cudf/reduction/bloom_filter.cuh index 16e932bd223e..97d70d9b79fb 100644 --- a/cpp/include/cudf/reduction/bloom_filter.cuh +++ b/cpp/include/cudf/reduction/bloom_filter.cuh @@ -14,14 +14,15 @@ namespace cudf { /** * @brief Policy describing the Apache Arrow Block-Split Bloom Filter layout. * - * Uses cuco's `parametric_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x + * Uses cuco's `bloom_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x * `uint32_t`), 8 fingerprint bits per key, fully horizontal add (Theta=8), and fully vertical * contains (Phi=8). This layout is bit-compatible with Apache Arrow. * + * @tparam Key The key type to generate a fingerprint for. * @tparam Hash The hash function used to generate a hash for each key. */ -template +template using arrow_filter_policy = - cuco::parametric_filter_policy; + cuco::bloom_filter_policy; } // namespace cudf diff --git a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu index ad3a2217ba13..d6ad8a116174 100644 --- a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu +++ b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu @@ -55,7 +55,7 @@ namespace cudf_streaming::detail { namespace { using KeyType = std::uint64_t; -using BloomFilterPolicy = cudf::arrow_filter_policy>; +using BloomFilterPolicy = cudf::arrow_filter_policy>; using BloomFilterRefType = cuco::bloom_filter_ref, cuco::thread_scope_device, diff --git a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu index b8fc5b2e978a..906ed416438f 100644 --- a/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu +++ b/cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu @@ -25,7 +25,7 @@ namespace { -using policy_type = cudf::arrow_filter_policy>; +using policy_type = cudf::arrow_filter_policy>; __global__ void block_index_kernel(std::uint32_t upper_hash, std::size_t num_blocks, diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 3418c8ca9ada..a7605f69c6a3 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -45,7 +45,7 @@ namespace { * `XXHash_64` (so that `cudf::string_view` and other cudf types are hashed by content, matching the * Apache Parquet/Arrow bloom filter specification). * - * Uses cuco's `parametric_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x + * Uses cuco's `bloom_filter_policy` with the Apache Arrow layout: 256-bit blocks (8 x * `uint32_t`), 8 fingerprint bits per key, fully horizontal add (Theta=8) and fully vertical * contains (Phi=8). This layout is bit-compatible with Apache Arrow, as verified by cuCollections * `tests/bloom_filter/arrow_compat_test.cu`. @@ -53,7 +53,7 @@ namespace { * @tparam Key The type of the values to generate a fingerprint for. */ template -using arrow_filter_policy = cudf::arrow_filter_policy>; +using arrow_filter_policy = cudf::arrow_filter_policy>; /** * @brief Converts bloom filter membership results (for each column chunk) to a device column. diff --git a/cpp/src/join/mark_join.cuh b/cpp/src/join/mark_join.cuh index f8c6e93d1dfd..df2965bf5599 100644 --- a/cpp/src/join/mark_join.cuh +++ b/cpp/src/join/mark_join.cuh @@ -166,7 +166,8 @@ using storage_ref_type = cuco::bucket_storage_ref>; using right_key_type = cuco::pair; -using bloom_filter_policy_type = cudf::arrow_filter_policy>; +using bloom_filter_policy_type = + cudf::arrow_filter_policy>; using bloom_filter_allocator_type = rmm::mr::polymorphic_allocator; using bloom_filter_type = cuco::bloom_filter, diff --git a/cpp/tests/io/parquet_bloom_filter_test.cu b/cpp/tests/io/parquet_bloom_filter_test.cu index 7fbb21714eff..adef120d0d6d 100644 --- a/cpp/tests/io/parquet_bloom_filter_test.cu +++ b/cpp/tests/io/parquet_bloom_filter_test.cu @@ -26,9 +26,10 @@ class ParquetBloomFilterTest : public cudf::test::BaseFixture {}; TEST_F(ParquetBloomFilterTest, TestStrings) { - using key_type = StringType; - using policy_type = cudf::arrow_filter_policy>; - using word_type = policy_type::word_type; + using key_type = StringType; + using policy_type = + cudf::arrow_filter_policy>; + using word_type = policy_type::word_type; std::size_t constexpr num_filter_blocks = 4; auto stream = cudf::get_default_stream(); From 602bc9d35826b0ecf35cbd9eb04be897684c11d3 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 17 Aug 2026 16:31:13 +0000 Subject: [PATCH 7/7] Test latest cuco Bloom filter policy --- cmake/rapids_config.cmake | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/cmake/rapids_config.cmake b/cmake/rapids_config.cmake index c9b1967245ac..18da8a799972 100644 --- a/cmake/rapids_config.cmake +++ b/cmake/rapids_config.cmake @@ -1,6 +1,6 @@ # ============================================================================= # cmake-format: off -# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. +# SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 # cmake-format: on # ============================================================================= @@ -34,6 +34,10 @@ endif() if(NOT rapids-cmake-branch) set(rapids-cmake-branch "${RAPIDS_BRANCH}") endif() + +# Test the cuco bump from PointKernel/rapids-cmake +set(rapids-cmake-repo "PointKernel/rapids-cmake") +set(rapids-cmake-branch "bump-cuco-4b26118") include("${CMAKE_CURRENT_LIST_DIR}/RAPIDS.cmake") # Don't use sccache-dist for CMake's compiler tests