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 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 9c83e7bfa855..df2965bf5599 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,7 +166,8 @@ 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 = + 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();