diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl index 3090d026e..abaf9e6ac 100644 --- a/include/cuco/detail/probing_scheme_impl.inl +++ b/include/cuco/detail/probing_scheme_impl.inl @@ -100,7 +100,7 @@ __host__ __device__ constexpr auto linear_probing::operator()( { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash_(probe_key)) % upper_bound, + cuco::detail::hash_to_index(hash_(probe_key)) % upper_bound, 1, // step size is 1 upper_bound}; } @@ -114,7 +114,7 @@ __host__ __device__ constexpr auto linear_probing::operator()( { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash_(probe_key) + g.thread_rank()) % upper_bound, + cuco::detail::hash_to_index(hash_(probe_key), g.thread_rank()) % upper_bound, cg_size, upper_bound}; } @@ -133,9 +133,9 @@ __host__ __device__ constexpr auto double_hashing::operato { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash1_(probe_key)) % upper_bound, + cuco::detail::hash_to_index(hash1_(probe_key)) % upper_bound, max(size_type{1}, - cuco::detail::sanitize_hash(hash2_(probe_key)) % + cuco::detail::hash_to_index(hash2_(probe_key)) % upper_bound), // step size in range [1, prime - 1] upper_bound}; } @@ -149,8 +149,8 @@ __host__ __device__ constexpr auto double_hashing::operato { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash1_(probe_key) + g.thread_rank()) % upper_bound, - static_cast((cuco::detail::sanitize_hash(hash2_(probe_key)) % + cuco::detail::hash_to_index(hash1_(probe_key), g.thread_rank()) % upper_bound, + static_cast((cuco::detail::hash_to_index(hash2_(probe_key)) % (upper_bound.value() / cg_size - 1) + 1) * cg_size), diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh index 22675d496..ca7160b04 100644 --- a/include/cuco/detail/utils.cuh +++ b/include/cuco/detail/utils.cuh @@ -112,15 +112,23 @@ struct strong_type { * @tparam SizeType The target type * @tparam HashType The input type * + * @param[in] hash The hash value + * @param[in] thread_rank Local thread offset during CG probing + * * @return Converted hash value */ template -__host__ __device__ constexpr SizeType sanitize_hash(HashType hash) noexcept +__host__ __device__ constexpr SizeType hash_to_index(HashType hash, + uint32_t thread_rank = 0) noexcept { - if constexpr (cuda::std::is_signed_v) { - return cuda::std::abs(static_cast(hash)); + if constexpr (std::is_unsigned_v) { + constexpr auto max_size = static_cast(cuda::std::numeric_limits::max()); + return static_cast((hash + thread_rank) & max_size); } else { - return static_cast(hash); + using larger_type = + cuda::std::conditional_t= sizeof(SizeType), HashType, SizeType>; + constexpr auto max_size = static_cast(cuda::std::numeric_limits::max()); + return static_cast((hash & max_size + thread_rank) & max_size); } }