fix sanitize_hash edge case

Author: srinivasyadav18

include/cuco/detail/probing_scheme_impl.inl

@@ -121,7 +121,7 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    cuco::detail::sanitize_hash<size_type>(hash_(probe_key), g.thread_rank()) % upper_bound,
+    cuco::detail::sanitize_hash<size_type>(hash_(probe_key), g) % upper_bound,
     cg_size,
     upper_bound};
 }
@@ -164,7 +164,7 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key), g.thread_rank()) % upper_bound,
+    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key), g) % upper_bound,
     static_cast<size_type>(
       (cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) % (upper_bound / cg_size - 1) +
        1) *

include/cuco/detail/utils.cuh

@@ -20,9 +20,12 @@
 #include <cuda/std/array>
 #include <cuda/std/bit>
 #include <cuda/std/cmath>
+#include <cuda/std/limits>
 #include <cuda/std/type_traits>
 #include <thrust/tuple.h>
 
+#include <cstddef>
+
 namespace cuco {
 namespace detail {
 
@@ -122,13 +125,19 @@ __host__ __device__ constexpr SizeType sanitize_hash(HashType hash) noexcept
  *
  * @tparam SizeType The target type
  * @tparam HashType The input type
+ * @tparam CG Cooperative group type
  *
  * @return Converted hash value
  */
-template <typename SizeType, typename HashType>
-__host__ __device__ constexpr SizeType sanitize_hash(HashType hash, std::uint32_t cg_rank) noexcept
+template <typename SizeType, typename HashType, typename CG>
+__host__ __device__ constexpr SizeType sanitize_hash(HashType hash, CG group) noexcept
 {
-  return sanitize_hash<SizeType>(sanitize_hash<SizeType>(hash) + cg_rank);
+  auto const base_hash = sanitize_hash<SizeType>(hash);
+  auto const max_size  = cuda::std::numeric_limits<SizeType>::max();
+  auto const cg_rank   = static_cast<SizeType>(group.thread_rank());
+
+  if (base_hash > (max_size - cg_rank)) return cg_rank - (max_size - base_hash);
+  return base_hash + cg_rank;
 }
 
 }  // namespace detail