From 0c16a977487c81dfca3b22d9934ea890f3f055e7 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 30 Jun 2026 20:17:16 +0000 Subject: [PATCH 1/6] Test rapids-cmake PR #1044 --- cmake/rapids_config.cmake | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/cmake/rapids_config.cmake b/cmake/rapids_config.cmake index c9b1967245ac..7e65b0009b42 100644 --- a/cmake/rapids_config.cmake +++ b/cmake/rapids_config.cmake @@ -34,6 +34,10 @@ endif() if(NOT rapids-cmake-branch) set(rapids-cmake-branch "${RAPIDS_BRANCH}") endif() + +# Test rapids-cmake PR #1044 (bump cuco for the new bloom filter) +set(rapids-cmake-repo "PointKernel/rapids-cmake") +set(rapids-cmake-branch "cuco-new-filter") include("${CMAKE_CURRENT_LIST_DIR}/RAPIDS.cmake") # Don't use sccache-dist for CMake's compiler tests From 1438c3612fb5a974b59b830d827b5d9929b9713e Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 30 Jun 2026 23:08:23 +0000 Subject: [PATCH 2/6] Adapt arrow_filter_policy and mark_join to new cuco bloom filter policy --- cpp/src/io/parquet/arrow_filter_policy.cuh | 118 ++++----------------- cpp/src/join/mark_join.cuh | 5 +- 2 files changed, 25 insertions(+), 98 deletions(-) diff --git a/cpp/src/io/parquet/arrow_filter_policy.cuh b/cpp/src/io/parquet/arrow_filter_policy.cuh index 74884d8c6b9e..ab6e93715d7c 100644 --- a/cpp/src/io/parquet/arrow_filter_policy.cuh +++ b/cpp/src/io/parquet/arrow_filter_policy.cuh @@ -1,23 +1,38 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once #include -#include -#include -#include +#include #include namespace cudf::io::parquet::detail { /** - * @brief A policy that defines how Arrow Block-Split Bloom Filter generates and stores a key's - * fingerprint. + * @brief Hasher adapter that exposes the `argument_type` member required by cuco's + * `parametric_filter_policy`. + * + * @tparam Key The type of the values to generate a fingerprint for. + */ +template +struct arrow_hasher : cudf::hashing::detail::XXHash_64 { + using argument_type = Key; + using cudf::hashing::detail::XXHash_64::XXHash_64; +}; + +/** + * @brief A policy that defines how the Apache Arrow Block-Split Bloom Filter generates and stores a + * key's fingerprint. + * + * Implemented in terms of 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 is bit-compatible with Apache Arrow, as verified by cuCollections + * `tests/bloom_filter/arrow_compat_test.cu`. * * Reference: * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 @@ -25,94 +40,7 @@ namespace cudf::io::parquet::detail { * @tparam Key The type of the values to generate a fingerprint for. */ template -class arrow_filter_policy { - public: - using hasher = - cudf::hashing::detail::XXHash_64; ///< 64-bit XXHash hasher for Arrow bloom filter policy - using word_type = std::uint32_t; ///< uint32_t for Arrow bloom filter policy - using key_type = Key; ///< Hash function input type - using hash_result_type = std::uint64_t; ///< hash function output type - - static constexpr std::uint32_t bits_set_per_block = 8; ///< bits set per Arrow filter block - static constexpr std::uint32_t words_per_block = 8; ///< words per Arrow filter block - - static constexpr std::uint32_t bytes_per_filter_block = - 32; ///< Number of bytes in one Arrow filter block - static constexpr std::uint32_t max_arrow_filter_bytes = - 128 * 1024 * 1024; ///< Max bytes in Arrow bloom filter - static constexpr std::uint32_t max_filter_blocks = - (max_arrow_filter_bytes / - bytes_per_filter_block); ///< Max sub-filter blocks allowed in Arrow bloom filter - - /** - * @brief Constructs the `arrow_filter_policy` object. - * - * @note The number of filter blocks with Arrow policy must be in the - * range of [1, 4194304]. If the bloom filter is constructed with a larger - * number of blocks, only the first 4194304 (128MB) blocks will be used. - * - * @param hash Hash function used to generate a key's fingerprint - */ - CUDF_HOST_DEVICE constexpr arrow_filter_policy(hasher hash = {}) : hash_{hash} {} - - /** - * @brief Generates the hash value for a given key. - * - * @param key The key to hash - * - * @return The hash value of the key - */ - __device__ constexpr hash_result_type hash(key_type const& key) const { return hash_(key); } - - /** - * @brief Determines the filter block a key is added into. - * - * @note The number of filter blocks with Arrow policy must be in the - * range of [1, 4194304]. Passing a larger `num_blocks` will still - * upperbound the number of blocks used to the mentioned range. - * - * @tparam Extent Size type that is used to determine the number of blocks in the filter - * - * @param hash Hash value of the key - * @param num_blocks Number of block in the filter - * - * @return The block index for the given key's hash value - */ - template - __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const - { - constexpr auto hash_bits = cuda::std::numeric_limits::digits; - auto const max_blocks = cuda::std::min(num_blocks, max_filter_blocks); - // Make sure we are only contained within the `max_filter_blocks` blocks - return static_cast(((hash >> hash_bits) * max_blocks) >> hash_bits) % max_blocks; - } - - /** - * @brief Determines the fingerprint pattern for a word/segment within the filter block for a - * given key's hash value. - * - * @param hash Hash value of the key - * @param word_index Target word/segment within the filter block - * - * @return The bit pattern for the word/segment in the filter block - */ - __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const - { - constexpr std::uint32_t salts[words_per_block] = {0x47b6137bU, - 0x44974d91U, - 0x8824ad5bU, - 0xa2b7289dU, - 0x705495c7U, - 0x2df1424bU, - 0x9efc4947U, - 0x5c6bfb31U}; - word_type const key = static_cast(hash); - auto const salt = salts[word_index]; - return word_type{1} << ((key * salt) >> 27); - } - - private: - hasher hash_; -}; +using arrow_filter_policy = + cuco::parametric_filter_policy, std::uint32_t, 8, 8, 8, 1, 1, 8, false, false>; } // namespace cudf::io::parquet::detail diff --git a/cpp/src/join/mark_join.cuh b/cpp/src/join/mark_join.cuh index ece70cf026a5..8f65fc517601 100644 --- a/cpp/src/join/mark_join.cuh +++ b/cpp/src/join/mark_join.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -165,8 +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, hash_value_type, 2U>; +using bloom_filter_policy_type = cuco::default_filter_policy; using bloom_filter_allocator_type = rmm::mr::polymorphic_allocator; using bloom_filter_type = cuco::bloom_filter, From 189ac849b02651e5bccc293b48cb9ba5b738053f Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 30 Jun 2026 23:17:49 +0000 Subject: [PATCH 3/6] Update copyright on rapids_config.cmake --- cmake/rapids_config.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cmake/rapids_config.cmake b/cmake/rapids_config.cmake index 7e65b0009b42..45cdf91a2be2 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 # ============================================================================= From 763fb86b3fc4d70a8834627175cd4298908fc4ae Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 1 Jul 2026 00:20:17 +0000 Subject: [PATCH 4/6] Share arrow_bloom_filter_policy between parquet and streaming --- .../utilities/arrow_bloom_filter_policy.cuh | 29 +++++++++++++++++++ .../src/detail/device_bloom_filter.cu | 3 +- cpp/src/io/parquet/arrow_filter_policy.cuh | 21 ++++---------- 3 files changed, 36 insertions(+), 17 deletions(-) create mode 100644 cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh diff --git a/cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh b/cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh new file mode 100644 index 000000000000..32a1d6402003 --- /dev/null +++ b/cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh @@ -0,0 +1,29 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include + +#include + +namespace cudf::detail { + +/** + * @brief Apache Arrow Block-Split Bloom Filter policy, parameterized on the fingerprint hasher. + * + * 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`. + * + * @tparam Hash Fingerprint hasher type. Must expose an `argument_type` member (e.g. + * `cuco::identity_hash`, or an adapter over `cudf::hashing::detail::XXHash_64`). + */ +template +using arrow_bloom_filter_policy = + cuco::parametric_filter_policy; + +} // namespace cudf::detail diff --git a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu index 99d5c8ea0744..494279122936 100644 --- a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu +++ b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu @@ -28,6 +28,7 @@ #pragma GCC diagnostic pop #endif +#include #include #include #include @@ -58,7 +59,7 @@ using BloomFilterRefType = cuco::bloom_filter_ref, cuco::thread_scope_device, - cuco::arrow_filter_policy>; + cudf::detail::arrow_bloom_filter_policy>>; using StorageType = BloomFilterRefType::filter_block_type; } // namespace diff --git a/cpp/src/io/parquet/arrow_filter_policy.cuh b/cpp/src/io/parquet/arrow_filter_policy.cuh index ab6e93715d7c..566e640fe9bd 100644 --- a/cpp/src/io/parquet/arrow_filter_policy.cuh +++ b/cpp/src/io/parquet/arrow_filter_policy.cuh @@ -5,12 +5,9 @@ #pragma once +#include #include -#include - -#include - namespace cudf::io::parquet::detail { /** @@ -26,21 +23,13 @@ struct arrow_hasher : cudf::hashing::detail::XXHash_64 { }; /** - * @brief A policy that defines how the Apache Arrow Block-Split Bloom Filter generates and stores a - * key's fingerprint. - * - * Implemented in terms of 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 is bit-compatible with Apache Arrow, as verified by cuCollections - * `tests/bloom_filter/arrow_compat_test.cu`. - * - * Reference: - * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 + * @brief Policy describing the Apache Arrow Block-Split Bloom Filter, hashing keys with cudf's + * `XXHash_64` (so that `cudf::string_view` and other cudf types are hashed by content, matching the + * Apache Parquet/Arrow bloom filter specification). * * @tparam Key The type of the values to generate a fingerprint for. */ template -using arrow_filter_policy = - cuco::parametric_filter_policy, std::uint32_t, 8, 8, 8, 1, 1, 8, false, false>; +using arrow_filter_policy = cudf::detail::arrow_bloom_filter_policy>; } // namespace cudf::io::parquet::detail From a731eeaa29097be2abe84c1a4fdf72bb1402a820 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 2 Jul 2026 19:13:30 +0000 Subject: [PATCH 5/6] Updates --- .../utilities/arrow_bloom_filter_policy.cuh | 29 --------------- .../src/detail/device_bloom_filter.cu | 12 +++++-- cpp/src/io/parquet/arrow_filter_policy.cuh | 35 ------------------- cpp/src/io/parquet/bloom_filter_reader.cu | 28 +++++++++++++-- cpp/tests/io/parquet_bloom_filter_test.cu | 22 +++++++++--- 5 files changed, 53 insertions(+), 73 deletions(-) delete mode 100644 cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh delete mode 100644 cpp/src/io/parquet/arrow_filter_policy.cuh diff --git a/cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh b/cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh deleted file mode 100644 index 32a1d6402003..000000000000 --- a/cpp/include/cudf/detail/utilities/arrow_bloom_filter_policy.cuh +++ /dev/null @@ -1,29 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include - -#include - -namespace cudf::detail { - -/** - * @brief Apache Arrow Block-Split Bloom Filter policy, parameterized on the fingerprint hasher. - * - * 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`. - * - * @tparam Hash Fingerprint hasher type. Must expose an `argument_type` member (e.g. - * `cuco::identity_hash`, or an adapter over `cudf::hashing::detail::XXHash_64`). - */ -template -using arrow_bloom_filter_policy = - cuco::parametric_filter_policy; - -} // namespace cudf::detail diff --git a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu index 494279122936..6245e4cad0b8 100644 --- a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu +++ b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu @@ -28,7 +28,6 @@ #pragma GCC diagnostic pop #endif -#include #include #include #include @@ -59,7 +58,16 @@ using BloomFilterRefType = cuco::bloom_filter_ref, cuco::thread_scope_device, - cudf::detail::arrow_bloom_filter_policy>>; + cuco::parametric_filter_policy, + std::uint32_t, + 8, + 8, + 8, + 1, + 1, + 8, + false, + false>>; using StorageType = BloomFilterRefType::filter_block_type; } // namespace diff --git a/cpp/src/io/parquet/arrow_filter_policy.cuh b/cpp/src/io/parquet/arrow_filter_policy.cuh deleted file mode 100644 index 566e640fe9bd..000000000000 --- a/cpp/src/io/parquet/arrow_filter_policy.cuh +++ /dev/null @@ -1,35 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include -#include - -namespace cudf::io::parquet::detail { - -/** - * @brief Hasher adapter that exposes the `argument_type` member required by cuco's - * `parametric_filter_policy`. - * - * @tparam Key The type of the values to generate a fingerprint for. - */ -template -struct arrow_hasher : cudf::hashing::detail::XXHash_64 { - using argument_type = Key; - using cudf::hashing::detail::XXHash_64::XXHash_64; -}; - -/** - * @brief Policy describing the Apache Arrow Block-Split Bloom Filter, hashing keys with cudf's - * `XXHash_64` (so that `cudf::string_view` and other cudf types are hashed by content, matching the - * Apache Parquet/Arrow bloom filter specification). - * - * @tparam Key The type of the values to generate a fingerprint for. - */ -template -using arrow_filter_policy = cudf::detail::arrow_bloom_filter_policy>; - -} // namespace cudf::io::parquet::detail diff --git a/cpp/src/io/parquet/bloom_filter_reader.cu b/cpp/src/io/parquet/bloom_filter_reader.cu index 2160b7179cc7..e0591f2d2115 100644 --- a/cpp/src/io/parquet/bloom_filter_reader.cu +++ b/cpp/src/io/parquet/bloom_filter_reader.cu @@ -1,9 +1,8 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ -#include "arrow_filter_policy.cuh" #include "compact_protocol_reader.hpp" #include "expression_transform_helpers.hpp" #include "io/utilities/time_utils.hpp" @@ -25,6 +24,7 @@ #include #include +#include #include #include #include @@ -36,6 +36,30 @@ namespace cudf::io::parquet::detail { namespace { +/** + * @brief Policy describing the Apache Arrow Block-Split Bloom Filter, hashing keys with cudf's + * `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 + * `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`. + * + * @tparam Key The type of the values to generate a fingerprint for. + */ +template +using arrow_filter_policy = cuco::parametric_filter_policy, + std::uint32_t, + 8, + 8, + 8, + 1, + 1, + 8, + false, + false>; + /** * @brief Converts bloom filter membership results (for each column chunk) to a device column. * diff --git a/cpp/tests/io/parquet_bloom_filter_test.cu b/cpp/tests/io/parquet_bloom_filter_test.cu index 7bb46eaaac32..4de68ff3c991 100644 --- a/cpp/tests/io/parquet_bloom_filter_test.cu +++ b/cpp/tests/io/parquet_bloom_filter_test.cu @@ -1,10 +1,8 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ -#include "src/io/parquet/arrow_filter_policy.cuh" - #include #include #include @@ -18,6 +16,9 @@ #include #include +#include + +#include using StringType = cudf::string_view; @@ -25,8 +26,19 @@ class ParquetBloomFilterTest : public cudf::test::BaseFixture {}; TEST_F(ParquetBloomFilterTest, TestStrings) { - using key_type = StringType; - using policy_type = cudf::io::parquet::detail::arrow_filter_policy; + using key_type = StringType; + // Apache Arrow Block-Split Bloom Filter layout, hashing keys with cudf's `XXHash_64` (matching + // `cudf::io::parquet::detail::arrow_filter_policy`). + using policy_type = cuco::parametric_filter_policy, + std::uint32_t, + 8, + 8, + 8, + 1, + 1, + 8, + false, + false>; using word_type = policy_type::word_type; std::size_t constexpr num_filter_blocks = 4; From f74183cab81d9ead66a44eced0a396878d18f8d8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 6 Jul 2026 17:30:34 +0000 Subject: [PATCH 6/6] Revert cmake changes --- 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 45cdf91a2be2..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 rapids-cmake PR #1044 (bump cuco for the new bloom filter) -set(rapids-cmake-repo "PointKernel/rapids-cmake") -set(rapids-cmake-branch "cuco-new-filter") include("${CMAKE_CURRENT_LIST_DIR}/RAPIDS.cmake") # Don't use sccache-dist for CMake's compiler tests