Test rapids-cmake cuco bump

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

cpp/include/cudf/reduction/bloom_filter.cuh

@@ -5,7 +5,7 @@
 
 #pragma once
 
-#include <cuco/bloom_filter_policies.cuh>
+#include <cuco/bloom_filter_policy.cuh>
 
 #include <cstdint>
 
@@ -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 <typename Hash>
+template <typename Key, typename Hash>
 using arrow_filter_policy =
-  cuco::parametric_filter_policy<Hash, std::uint32_t, 8, 8, 8, 1, 1, 8, false, false>;
+  cuco::bloom_filter_policy<Key, Hash, sizeof(std::uint32_t), 8, 8, 8, 1, 1, 8>;
 
 }  // namespace cudf

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<cuco::identity_hash<KeyType>>;
+using BloomFilterPolicy  = cudf::arrow_filter_policy<KeyType, cuco::identity_hash<KeyType>>;
 using BloomFilterRefType = cuco::bloom_filter_ref<KeyType,
                                                   cuco::extent<std::size_t>,
                                                   cuco::thread_scope_device,

cpp/libcudf_streaming/tests/streaming/test_bloom_filter.cu

@@ -25,7 +25,7 @@
 
 namespace {
 
-using policy_type = cudf::arrow_filter_policy<cuco::identity_hash<std::uint64_t>>;
+using policy_type = cudf::arrow_filter_policy<std::uint64_t, cuco::identity_hash<std::uint64_t>>;
 
 __global__ void block_index_kernel(std::uint32_t upper_hash,
                                    std::size_t num_blocks,

cpp/src/io/parquet/bloom_filter_reader.cu

@@ -45,15 +45,15 @@ 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`.
  *
  * @tparam Key The type of the values to generate a fingerprint for.
  */
 template <class Key>
-using arrow_filter_policy = cudf::arrow_filter_policy<cudf::hashing::detail::XXHash_64<Key>>;
+using arrow_filter_policy = cudf::arrow_filter_policy<Key, cudf::hashing::detail::XXHash_64<Key>>;
 
 /**
  * @brief Converts bloom filter membership results (for each column chunk) to a device column.

cpp/src/join/mark_join.cuh

@@ -10,6 +10,7 @@
 #include <cudf/detail/row_operator/hashing.cuh>
 #include <cudf/detail/row_operator/primitive_row_operators.cuh>
 #include <cudf/join/join.hpp>
+#include <cudf/reduction/bloom_filter.cuh>
 #include <cudf/table/table_view.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/memory_resource.hpp>
@@ -165,7 +166,8 @@ using storage_ref_type =
   cuco::bucket_storage_ref<mark_key_type, mark_join_bucket_size, cuco::extent<std::size_t>>;
 using right_key_type = cuco::pair<hash_value_type, rhs_index_type>;
 
-using bloom_filter_policy_type    = cuco::default_filter_policy<hash_value_type>;
+using bloom_filter_policy_type =
+  cudf::arrow_filter_policy<hash_value_type, cuco::xxhash_64<hash_value_type>>;
 using bloom_filter_allocator_type = rmm::mr::polymorphic_allocator<cuda::std::byte>;
 using bloom_filter_type           = cuco::bloom_filter<hash_value_type,
                                                        cuco::extent<std::size_t>,

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<cudf::hashing::detail::XXHash_64<key_type>>;
-  using word_type   = policy_type::word_type;
+  using key_type = StringType;
+  using policy_type =
+    cudf::arrow_filter_policy<key_type, cudf::hashing::detail::XXHash_64<key_type>>;
+  using word_type = policy_type::word_type;
 
   std::size_t constexpr num_filter_blocks = 4;
   auto stream                             = cudf::get_default_stream();