From code reviews

include/cuco/bloom_filter_policies.cuh

@@ -29,12 +29,12 @@ namespace cuco {
  * fingerprint.
  *
  * @tparam Key The type of the values to generate a fingerprint for.
- * @tparam Hash Hash function used to generate a key's fingerprint. By default, cuco::xxhash_64 will
- * be used.
+ * @tparam XXHash64 Custom (64 bit) XXHash hasher to generate a key's fingerprint.
+ * By default, cuco::xxhash_64 hasher will be used.
  *
  */
-template <class Key, class Hash = cuco::xxhash_64<Key>>
-using arrow_filter_policy = detail::arrow_filter_policy<Key, Hash>;
+template <class Key, class XXHash64 = cuco::xxhash_64<Key>>
+using arrow_filter_policy = detail::arrow_filter_policy<Key, XXHash64>;
 
 /**
  * @brief The default policy that defines how a Blocked Bloom Filter generates and stores a key's

include/cuco/detail/bloom_filter/arrow_filter_policy.cuh

@@ -78,11 +78,12 @@ namespace cuco::detail {
  * @endcode
  *
  * @tparam Key The type of the values to generate a fingerprint for.
+ * @tparam XXHash64 64-bit XXHash hasher implementation for fingerprint generation.
  */
-template <class Key, class Hash>
+template <class Key, class XXHash64>
 class arrow_filter_policy {
  public:
-  using hasher             = Hash;           ///< Hash function for Arrow bloom filter policy
+  using hasher             = XXHash64;       ///< 64-bit XXHash hasher for Arrow bloom filter policy
   using word_type          = std::uint32_t;  ///< uint32_t for Arrow bloom filter policy
   using hash_argument_type = typename hasher::argument_type;  ///< Hash function input type
   using hash_result_type   = decltype(std::declval<hasher>()(

include/cuco/detail/hash_functions/xxhash.cuh

@@ -100,17 +100,6 @@ struct XXHash_32 {
     }
   }
 
-  /**
-   * @brief Returns a hash value for its `span` like argument, as a value of type `result_type`.
-   *
-   * @param key The input argument to hash
-   * @return The resulting hash value for `span` like `key`
-   */
-  template <typename T = Key, typename = std::enable_if_t<is_span_like_v<T>>>
-  constexpr result_type __host__ __device__ operator()(Key const& key) const noexcept
-  {
-    return compute_hash(key.data(), key.size());
-  }
   /**
    * @brief Returns a hash value for its argument, as a value of type `result_type`.
    *

tests/bloom_filter/unique_sequence_test.cu

@@ -32,11 +32,15 @@
 
 using size_type = int32_t;
 
-template <typename Filter, typename Key>
-void test_unique_sequence(Filter& filter,
-                          thrust::device_vector<Key> const& keys,
-                          size_type num_keys)
+template <typename Filter>
+void test_unique_sequence(Filter& filter, size_type num_keys)
 {
+  using Key = typename Filter::key_type;
+
+  // Generate keys
+  thrust::device_vector<Key> keys(num_keys);
+  thrust::sequence(thrust::device, keys.begin(), keys.end());
+
   thrust::device_vector<bool> contained(num_keys, false);
 
   auto is_even =
@@ -100,11 +104,7 @@ TEMPLATE_TEST_CASE_SIG(
 
   auto filter = filter_type{1000, {}, {pattern_bits}};
 
-  // Generate keys
-  thrust::device_vector<Key> keys(num_keys);
-  thrust::sequence(thrust::device, keys.begin(), keys.end());
-
-  test_unique_sequence(filter, keys, num_keys);
+  test_unique_sequence(filter, num_keys);
 }
 
 TEMPLATE_TEST_CASE_SIG("Unique sequence with arrow policy",
@@ -120,27 +120,5 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence with arrow policy",
 
   auto filter = filter_type{1000};
 
-  // Generate keys
-  thrust::device_vector<Key> keys(num_keys);
-  thrust::sequence(thrust::device, keys.begin(), keys.end());
-
-  test_unique_sequence(filter, keys, num_keys);
-}
-
-TEMPLATE_TEST_CASE_SIG("Unique string sequence with arrow policy",
-                       "",
-                       ((class Key, class Policy), Key, Policy),
-                       (cuda::std::span<cuda::std::byte>,
-                        cuco::arrow_filter_policy<cuda::std::span<cuda::std::byte>>))
-{
-  using filter_type =
-    cuco::bloom_filter<Key, cuco::extent<size_t>, cuda::thread_scope_device, Policy>;
-  constexpr size_type num_keys{400};
-
-  auto filter = filter_type{1000};
-
-  // Generate keys (string spans) and the actual string data
-  auto [keys, data] = cuco::utility::generate_random_byte_sequences(num_keys, 20, 50);
-
-  test_unique_sequence(filter, keys, num_keys);
+  test_unique_sequence(filter, num_keys);
 }