diff --git a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu index 99d5c8ea0744..6245e4cad0b8 100644 --- a/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu +++ b/cpp/libcudf_streaming/src/detail/device_bloom_filter.cu @@ -58,7 +58,16 @@ using BloomFilterRefType = cuco::bloom_filter_ref, cuco::thread_scope_device, - cuco::arrow_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 74884d8c6b9e..000000000000 --- a/cpp/src/io/parquet/arrow_filter_policy.cuh +++ /dev/null @@ -1,118 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#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. - * - * Reference: - * https://github.com/apache/arrow/blob/be1dcdb96b030639c0b56955c4c62f9d6b03f473/cpp/src/parquet/bloom_filter.cc#L219-L230 - * - * @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_; -}; - -} // 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/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, 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;