Use multiply-high range reduction for probing

include/cuco/detail/hash_functions/identity_hash.cuh

@@ -13,11 +13,7 @@ namespace cuco::detail {
 /**
  * @brief An Identity hash function to hash the given argument on host and device
  *
- * @note `identity_hash` is perfect if `hash_table_capacity >= |input set|`
- *
- * @note `identity_hash` is only intended to be used perfectly.
- *
- * @note Perfect hashes are deterministic, and thus do not need seeds.
+ * @note Identity hashes are deterministic, and thus do not need seeds.
  *
  * @tparam Key The type of the values to hash
  */

include/cuco/detail/probing_scheme/probing_scheme_impl.inl

@@ -102,7 +102,7 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::make_iterator(
   using size_type             = typename Extent::value_type;
   size_type const num_buckets = upper_bound / BucketSize;
   size_type const init =
-    cuco::detail::sanitize_hash<size_type>(hash_(probe_key), num_buckets) * BucketSize;
+    cuco::detail::hash_to_index<size_type>(hash_(probe_key), num_buckets) * BucketSize;
   return detail::probing_iterator<Extent>{init, static_cast<size_type>(BucketSize), upper_bound};
 }
 
@@ -117,7 +117,7 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::make_iterator(
   size_type constexpr stride = cg_size * BucketSize;
   size_type const num_groups = upper_bound / stride;
   size_type const init =
-    cuco::detail::sanitize_hash<size_type>(hash_(probe_key), num_groups) * stride +
+    cuco::detail::hash_to_index<size_type>(hash_(probe_key), num_groups) * stride +
     g.thread_rank() * BucketSize;
   return detail::probing_iterator<Extent>{init, stride, upper_bound};
 }
@@ -165,9 +165,9 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::make_it
   using size_type             = typename Extent::value_type;
   size_type const num_buckets = upper_bound / BucketSize;
   return detail::probing_iterator<Extent>{
-    static_cast<size_type>(cuco::detail::sanitize_hash<size_type>(hash1_(probe_key), num_buckets) *
+    static_cast<size_type>(cuco::detail::hash_to_index<size_type>(hash1_(probe_key), num_buckets) *
                            BucketSize),
-    static_cast<size_type>((cuco::detail::sanitize_hash<size_type>(
+    static_cast<size_type>((cuco::detail::hash_to_index<size_type>(
                               hash2_(probe_key), static_cast<size_type>(num_buckets - 1)) +
                             1) *
                            BucketSize),  // step size in range [1, prime - 1]
@@ -186,10 +186,10 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::make_it
   size_type const num_groups = upper_bound / stride;
 
   return detail::probing_iterator<Extent>{
-    static_cast<size_type>(cuco::detail::sanitize_hash<size_type>(hash1_(probe_key), num_groups) *
+    static_cast<size_type>(cuco::detail::hash_to_index<size_type>(hash1_(probe_key), num_groups) *
                              stride +
                            g.thread_rank() * BucketSize),
-    static_cast<size_type>((cuco::detail::sanitize_hash<size_type>(
+    static_cast<size_type>((cuco::detail::hash_to_index<size_type>(
                               hash2_(probe_key), static_cast<size_type>(num_groups - 1)) +
                             1) *
                            stride),

include/cuco/detail/utils.cuh

@@ -7,6 +7,7 @@
 
 #include <cuco/detail/__config>
 
+#include <cuda/cmath>
 #include <cuda/std/array>
 #include <cuda/std/cmath>
 #include <cuda/std/cstdint>
@@ -26,55 +27,61 @@ __device__ __forceinline__ cuda::std::int32_t count_least_significant_bits(cuda:
   return __popc(x & (1 << n) - 1);
 }
 
-template <typename SizeType, typename HashType>
-__host__ __device__ constexpr cuda::std::make_unsigned_t<SizeType> to_positive(HashType hash)
-{
-  using unsigned_size_type = cuda::std::make_unsigned_t<SizeType>;
-  auto const value         = static_cast<unsigned_size_type>(hash);
-
-  if constexpr (cuda::std::is_signed_v<SizeType>) {
-    auto constexpr max =
-      static_cast<unsigned_size_type>(cuda::std::numeric_limits<SizeType>::max());
-    return value > max ? unsigned_size_type{0} - value : value;
-  } else {
-    return value;
-  }
-}
-
 /**
- * @brief Converts a hash value into a valid index for the given modulus.
+ * @brief Maps an unsigned hash bit pattern into `[0, upper_bound)` using multiply-high.
  *
- * @note Hash values wider than `SizeType` are narrowed before reduction, preserving the existing
- * low-bit mapping policy.
+ * @note Hash bits should be uniformly distributed across the hash type's width. This is range
+ * reduction, not a remainder operation. A 128-bit hash contributes its high 64 bits (the second
+ * element of a two-word hash), which are then reduced like a 64-bit hash.
  *
- * @tparam SizeType The target type
- * @tparam HashType The input type
+ * @tparam SizeType Integral index type
+ * @tparam HashType Integral hash type of at most 128 bits, or a two-word 128-bit hash
  *
  * @param hash The hash value
- * @param modulus Exclusive upper bound for the returned index
+ * @param upper_bound Positive exclusive upper bound for the returned index
  *
- * @return An index in `[0, modulus)`
+ * @return An index in `[0, upper_bound)`
  */
 template <typename SizeType, typename HashType>
-__host__ __device__ constexpr SizeType sanitize_hash(HashType hash, SizeType modulus) noexcept
+__host__ __device__ constexpr SizeType hash_to_index(HashType hash, SizeType upper_bound) noexcept
 {
-  using unsigned_size_type = cuda::std::make_unsigned_t<SizeType>;
-
-  unsigned_size_type magnitude;
+  using unsigned_size_type   = cuda::std::make_unsigned_t<SizeType>;
+  constexpr auto hash_bytes  = sizeof(HashType);
+  constexpr auto index_bytes = sizeof(SizeType);
   if constexpr (cuda::std::is_same_v<HashType, cuda::std::array<std::uint64_t, 2>>) {
-#if !defined(CUCO_HAS_INT128)
-    static_assert(false,
-                  "CUCO_HAS_INT128 undefined. Need unsigned __int128 type when sanitizing "
-                  "cuda::std::array<std::uint64_t, 2>");
-#endif
-    unsigned __int128 ret{};
-    memcpy(&ret, &hash, sizeof(unsigned __int128));
-    magnitude = to_positive<SizeType>(ret);
+    return hash_to_index<SizeType>(hash[1], upper_bound);
+  } else if constexpr (hash_bytes > 8) {
+    static_assert(cuda::std::is_integral_v<HashType> and hash_bytes == 16);
+    using unsigned_hash_type = cuda::std::make_unsigned_t<HashType>;
+    return hash_to_index<SizeType>(
+      static_cast<std::uint64_t>(static_cast<unsigned_hash_type>(hash) >> 64), upper_bound);
+  } else if constexpr (hash_bytes == 4 and index_bytes == 8) {
+    auto const bits  = static_cast<cuda::std::uint32_t>(hash);
+    auto const bound = static_cast<cuda::std::uint64_t>(upper_bound);
+    // Decompose the 96-bit product without introducing a full 64-bit high multiply.
+    return static_cast<SizeType>(cuda::mul_hi(bits, static_cast<cuda::std::uint32_t>(bound)) +
+                                 static_cast<cuda::std::uint64_t>(bits) *
+                                   static_cast<cuda::std::uint32_t>(bound >> 32));
+  } else if constexpr (hash_bytes == 8 and index_bytes == 4) {
+    auto const bits  = static_cast<cuda::std::uint64_t>(hash);
+    auto const bound = static_cast<cuda::std::uint32_t>(upper_bound);
+    // Only the upper 32 bits of the 96-bit product are needed.
+    auto const middle =
+      static_cast<cuda::std::uint64_t>(static_cast<cuda::std::uint32_t>(bits >> 32)) * bound +
+      cuda::mul_hi(static_cast<cuda::std::uint32_t>(bits), bound);
+    return static_cast<SizeType>(middle >> 32);
   } else {
-    magnitude = to_positive<SizeType>(hash);
+    using unsigned_hash_type = cuda::std::make_unsigned_t<HashType>;
+    using product_type =
+      cuda::std::conditional_t<(hash_bytes > index_bytes), unsigned_hash_type, unsigned_size_type>;
+    static_assert(sizeof(product_type) <= 8);
+    constexpr auto shift = cuda::std::numeric_limits<product_type>::digits -
+                           cuda::std::numeric_limits<unsigned_hash_type>::digits;
+    // Left-align a narrower hash so widening does not shrink the range of initial indices.
+    auto const bits = static_cast<product_type>(static_cast<unsigned_hash_type>(hash)) << shift;
+    return static_cast<SizeType>(
+      cuda::mul_hi(static_cast<product_type>(bits), static_cast<product_type>(upper_bound)));
   }
-
-  return static_cast<SizeType>(magnitude % static_cast<unsigned_size_type>(modulus));
 }
 
 }  // namespace detail

tests/CMakeLists.txt

@@ -45,6 +45,7 @@ ConfigureTest(UTILITY_TEST
     utility/storage_test.cu
     utility/fast_int_test.cu
     utility/hash_test.cu
+    utility/hash_to_index_test.cu
     utility/probing_scheme_test.cu)
 
 # catch_discover_tests executes each Catch2 test case in its own process. This matters for device

tests/utility/hash_to_index_test.cu

@@ -0,0 +1,163 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <test_utils.hpp>
+
+#include <cuco/detail/utils.cuh>
+
+#include <cuda/std/array>
+#include <cuda/std/limits>
+#include <cuda/std/type_traits>
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+
+#include <catch2/catch_template_test_macros.hpp>
+#include <catch2/catch_test_macros.hpp>
+
+#include <cstdint>
+#include <random>
+#include <vector>
+
+template <typename Hash, typename Size>
+__global__ void hash_to_indices(Hash const* hashes, Size const* bounds, Size* output, int count)
+{
+  auto const i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < count) { output[i] = cuco::detail::hash_to_index<Size>(hashes[i], bounds[i]); }
+}
+
+TEST_CASE("hash_to_index is constexpr and preserves unsigned hash bits", "")
+{
+  using cuco::detail::hash_to_index;
+  STATIC_REQUIRE(hash_to_index<std::int32_t>(std::uint32_t{0x80000000}, 10) == 5);
+  STATIC_REQUIRE(hash_to_index<std::int32_t>(std::int32_t{-1}, 10) == 9);
+  STATIC_REQUIRE(hash_to_index<std::int64_t>(std::uint32_t{0x80000000}, 10) == 5);
+  STATIC_REQUIRE(hash_to_index<std::uint64_t>(std::uint32_t{0x80000000}, 0x200000002ull) ==
+                 0x100000001ull);
+  STATIC_REQUIRE(hash_to_index<std::uint32_t>(std::uint64_t{0x8000000000000000ull}, 10) == 5);
+  STATIC_REQUIRE(hash_to_index<std::uint64_t>(std::uint64_t{0xffffffffffffffffull},
+                                              0xffffffffffffffffull) == 0xfffffffffffffffeull);
+  STATIC_REQUIRE(hash_to_index<std::uint32_t>(std::uint8_t{128}, 10) == 5);
+}
+
+#if defined(CUCO_HAS_INT128)
+TEMPLATE_TEST_CASE("hash_to_index uses the high word of scalar 128-bit hashes",
+                   "",
+                   std::int32_t,
+                   std::uint32_t,
+                   std::int64_t,
+                   std::uint64_t)
+{
+  constexpr auto hash =
+    (static_cast<unsigned __int128>(0x4000000080000000ull) << 64) | 0x123456789abcdef0ull;
+  constexpr TestType expected = 2;
+  STATIC_REQUIRE(cuco::detail::hash_to_index<TestType>(hash, 10) == expected);
+  STATIC_REQUIRE(cuco::detail::hash_to_index<TestType>(static_cast<__int128>(-1), 10) == 9);
+  // The low word does not contribute.
+  STATIC_REQUIRE(cuco::detail::hash_to_index<TestType>(
+                   static_cast<unsigned __int128>(0xffffffffffffffffull), 10) == 0);
+  thrust::device_vector<unsigned __int128> hashes(1, hash);
+  thrust::device_vector<TestType> bounds(1, 10);
+  thrust::device_vector<TestType> output(1);
+  hash_to_indices<<<1, 1>>>(thrust::raw_pointer_cast(hashes.data()),
+                            thrust::raw_pointer_cast(bounds.data()),
+                            thrust::raw_pointer_cast(output.data()),
+                            1);
+  CUCO_CUDA_TRY(cudaGetLastError());
+  REQUIRE(output[0] == expected);
+}
+
+TEMPLATE_TEST_CASE_SIG("hash_to_index matches a wide product on host and device",
+                       "",
+                       ((typename Hash, typename Size), Hash, Size),
+                       (std::uint32_t, std::int32_t),
+                       (std::uint32_t, std::uint32_t),
+                       (std::uint32_t, std::int64_t),
+                       (std::uint32_t, std::uint64_t),
+                       (std::uint64_t, std::int32_t),
+                       (std::uint64_t, std::uint32_t),
+                       (std::uint64_t, std::int64_t),
+                       (std::uint64_t, std::uint64_t),
+                       (std::int32_t, std::int64_t),
+                       (std::int64_t, std::int32_t))
+{
+  using unsigned_hash     = cuda::std::make_unsigned_t<Hash>;
+  constexpr auto bits     = cuda::std::numeric_limits<unsigned_hash>::digits;
+  constexpr auto max_hash = cuda::std::numeric_limits<unsigned_hash>::max();
+  constexpr auto max_size = cuda::std::numeric_limits<Size>::max();
+  std::vector<Hash> hashes;
+  std::vector<Size> bounds;
+  std::vector<Size> expected;
+  auto append = [&](unsigned_hash hash, Size bound) {
+    hashes.push_back(static_cast<Hash>(hash));
+    bounds.push_back(bound);
+    expected.push_back(static_cast<Size>((static_cast<unsigned __int128>(hash) * bound) >> bits));
+  };
+  for (auto hash : {unsigned_hash{0},
+                    unsigned_hash{1},
+                    unsigned_hash{max_hash / 2},
+                    unsigned_hash{max_hash / 2 + 1},
+                    unsigned_hash{max_hash - 1},
+                    max_hash}) {
+    for (auto bound : {Size{1},
+                       Size{2},
+                       Size{3},
+                       Size{10},
+                       Size{65537},
+                       Size{max_size / 2},
+                       Size{max_size - 1},
+                       max_size}) {
+      append(hash, bound);
+    }
+    if constexpr (sizeof(Size) == 8) {
+      for (auto bound : {Size{0xffffffffull}, Size{0x100000000ull}, Size{0x100000001ull}}) {
+        append(hash, bound);
+      }
+    }
+  }
+  std::mt19937_64 rng{24331};
+  for (int i = 0; i < 4096; ++i) {
+    append(static_cast<unsigned_hash>(rng()), static_cast<Size>(rng() % max_size + 1));
+  }
+  for (std::size_t i = 0; i < hashes.size(); ++i) {
+    REQUIRE(cuco::detail::hash_to_index<Size>(hashes[i], bounds[i]) == expected[i]);
+  }
+  thrust::device_vector<Hash> d_hashes(hashes);
+  thrust::device_vector<Size> d_bounds(bounds);
+  thrust::device_vector<Size> d_output(hashes.size());
+  hash_to_indices<<<(hashes.size() + 127) / 128, 128>>>(thrust::raw_pointer_cast(d_hashes.data()),
+                                                        thrust::raw_pointer_cast(d_bounds.data()),
+                                                        thrust::raw_pointer_cast(d_output.data()),
+                                                        static_cast<int>(hashes.size()));
+  CUCO_CUDA_TRY(cudaGetLastError());
+  thrust::host_vector<Size> output(d_output);
+  for (std::size_t i = 0; i < hashes.size(); ++i) {
+    REQUIRE(output[i] == expected[i]);
+  }
+}
+#endif
+
+TEMPLATE_TEST_CASE("hash_to_index uses the high word of two-word 128-bit hashes",
+                   "",
+                   std::int32_t,
+                   std::uint32_t,
+                   std::int64_t,
+                   std::uint64_t)
+{
+  using hash_type = cuda::std::array<std::uint64_t, 2>;
+  constexpr hash_type hash{0x123456789abcdef0ull, 0x4000000080000000ull};
+  constexpr TestType expected = 2;
+  STATIC_REQUIRE(cuco::detail::hash_to_index<TestType>(hash, 10) == expected);
+  STATIC_REQUIRE(cuco::detail::hash_to_index<TestType>(hash_type{0xffffffffffffffffull, 0}, 10) ==
+                 0);
+  thrust::device_vector<hash_type> hashes(1, hash);
+  thrust::device_vector<TestType> bounds(1, 10);
+  thrust::device_vector<TestType> output(1);
+  hash_to_indices<<<1, 1>>>(thrust::raw_pointer_cast(hashes.data()),
+                            thrust::raw_pointer_cast(bounds.data()),
+                            thrust::raw_pointer_cast(output.data()),
+                            1);
+  CUCO_CUDA_TRY(cudaGetLastError());
+  REQUIRE(output[0] == expected);
+}