From 200183781c93270c852eb9d6ca8eaa44f6b5adf0 Mon Sep 17 00:00:00 2001
From: srinivasyadav18 <ssinganaboin@nvidia.com>
Date: Thu, 6 Jun 2024 21:42:54 +0000
Subject: [PATCH 1/6] Make murmurhash3_x64_128 compatible with existing cuco
 data structures

---
 include/cuco/detail/probing_scheme_base.cuh | 30 +++++++++
 include/cuco/detail/probing_scheme_impl.inl | 19 ++++--
 include/cuco/detail/utils.cuh               | 19 +-----
 tests/CMakeLists.txt                        |  1 +
 tests/static_map/hash_test.cu               | 72 +++++++++++++++++++++
 5 files changed, 117 insertions(+), 24 deletions(-)
 create mode 100644 tests/static_map/hash_test.cu

diff --git a/include/cuco/detail/probing_scheme_base.cuh b/include/cuco/detail/probing_scheme_base.cuh
index a3d7c148a..9ce06da92 100644
--- a/include/cuco/detail/probing_scheme_base.cuh
+++ b/include/cuco/detail/probing_scheme_base.cuh
@@ -16,6 +16,8 @@
 
 #pragma once
 
+#include <cuco/detail/__config>
+
 #include <cstdint>
 
 namespace cuco {
@@ -30,6 +32,34 @@ namespace detail {
  */
 template <int32_t CGSize>
 class probing_scheme_base {
+ private:
+  template <typename SizeType, typename HashType>
+  __host__ __device__ constexpr SizeType sanitize_hash_positive(HashType hash) const noexcept
+  {
+    if constexpr (cuda::std::is_signed_v<SizeType>) {
+      return cuda::std::abs(static_cast<SizeType>(hash));
+    } else {
+      return static_cast<SizeType>(hash);
+    }
+  }
+
+ protected:
+  template <typename SizeType, typename HashType>
+  __host__ __device__ constexpr SizeType sanitize_hash(HashType hash) const noexcept
+  {
+    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));
+      return sanitize_hash_positive<SizeType>(static_cast<SizeType>(ret));
+    } else
+      return sanitize_hash_positive<SizeType>(hash);
+  }
+
  public:
   /**
    * @brief The size of the CUDA cooperative thread group.
diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
index 50d7c4dcc..33998168e 100644
--- a/include/cuco/detail/probing_scheme_impl.inl
+++ b/include/cuco/detail/probing_scheme_impl.inl
@@ -107,7 +107,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)) % upper_bound,
+    probing_scheme_base_type::template sanitize_hash<size_type>(hash_(probe_key)) % upper_bound,
     1,  // step size is 1
     upper_bound};
 }
@@ -121,7 +121,10 @@ __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,
+    probing_scheme_base_type::template sanitize_hash<size_type>(
+      probing_scheme_base_type::template sanitize_hash<size_type>(hash_(probe_key)) +
+      g.thread_rank()) %
+      upper_bound,
     cg_size,
     upper_bound};
 }
@@ -148,9 +151,9 @@ __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)) % upper_bound,
+    probing_scheme_base_type::template sanitize_hash<size_type>(hash1_(probe_key)) % upper_bound,
     max(size_type{1},
-        cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) %
+        probing_scheme_base_type::template sanitize_hash<size_type>(hash2_(probe_key)) %
           upper_bound),  // step size in range [1, prime - 1]
     upper_bound};
 }
@@ -164,9 +167,13 @@ __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,
+    probing_scheme_base_type::template sanitize_hash<size_type>(
+      probing_scheme_base_type::template sanitize_hash<size_type>(hash1_(probe_key)) +
+      g.thread_rank()) %
+      upper_bound,
     static_cast<size_type>(
-      (cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) % (upper_bound / cg_size - 1) +
+      (probing_scheme_base_type::template sanitize_hash<size_type>(hash2_(probe_key)) %
+         (upper_bound / cg_size - 1) +
        1) *
       cg_size),
     upper_bound};  // TODO use fast_int operator
diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index 1cbe8fd26..f2aecc0ef 100644
--- a/include/cuco/detail/utils.cuh
+++ b/include/cuco/detail/utils.cuh
@@ -17,6 +17,7 @@
 
 #include <cuco/detail/bitwise_compare.cuh>
 
+#include <cuda/std/array>
 #include <cuda/std/bit>
 #include <cuda/std/cmath>
 #include <cuda/std/type_traits>
@@ -81,23 +82,5 @@ struct slot_is_filled {
   }
 };
 
-/**
- * @brief Converts a given hash value into a valid (positive) size type.
- *
- * @tparam SizeType The target type
- * @tparam HashType The input type
- *
- * @return Converted hash value
- */
-template <typename SizeType, typename HashType>
-__host__ __device__ constexpr SizeType sanitize_hash(HashType hash) noexcept
-{
-  if constexpr (cuda::std::is_signed_v<SizeType>) {
-    return cuda::std::abs(static_cast<SizeType>(hash));
-  } else {
-    return static_cast<SizeType>(hash);
-  }
-}
-
 }  // namespace detail
 }  // namespace cuco
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index a37f2d4e2..9d75d7a0e 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -76,6 +76,7 @@ ConfigureTest(STATIC_MAP_TEST
     static_map/custom_type_test.cu
     static_map/duplicate_keys_test.cu
     static_map/erase_test.cu
+    static_map/hash_test.cu
     static_map/heterogeneous_lookup_test.cu
     static_map/insert_and_find_test.cu
     static_map/insert_or_assign_test.cu
diff --git a/tests/static_map/hash_test.cu b/tests/static_map/hash_test.cu
new file mode 100644
index 000000000..c22eae998
--- /dev/null
+++ b/tests/static_map/hash_test.cu
@@ -0,0 +1,72 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <test_utils.hpp>
+
+#include <cuco/hash_functions.cuh>
+#include <cuco/static_map.cuh>
+
+#include <thrust/device_vector.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/transform_iterator.h>
+
+#include <catch2/catch_template_test_macros.hpp>
+
+using size_type = std::size_t;
+
+template <typename Key, typename Hash>
+void test_hash_function()
+{
+  using Value = int64_t;
+
+  constexpr size_type num_keys{400};
+
+  auto map = cuco::static_map<Key,
+                              Value,
+                              cuco::extent<size_type>,
+                              cuda::thread_scope_device,
+                              thrust::equal_to<Key>,
+                              cuco::linear_probing<1, Hash>,
+                              cuco::cuda_allocator<std::byte>,
+                              cuco::storage<2>>{
+    num_keys, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
+
+  auto keys_begin = thrust::counting_iterator<Key>(1);
+
+  auto pairs_begin = thrust::make_transform_iterator(
+    keys_begin, cuda::proclaim_return_type<cuco::pair<Key, Value>>([] __device__(auto i) {
+      return cuco::pair<Key, Value>(i, i);
+    }));
+
+  thrust::device_vector<bool> d_keys_exist(num_keys);
+
+  map.insert(pairs_begin, pairs_begin + num_keys);
+
+  REQUIRE(map.size() == num_keys);
+
+  map.contains(keys_begin, keys_begin + num_keys, d_keys_exist.begin());
+
+  REQUIRE(cuco::test::all_of(d_keys_exist.begin(), d_keys_exist.end(), thrust::identity{}));
+}
+
+TEMPLATE_TEST_CASE_SIG("static_map hash tests", "", ((typename Key)), (int32_t), (int64_t))
+{
+  test_hash_function<Key, cuco::murmurhash3_32<Key>>();
+  test_hash_function<Key, cuco::murmurhash3_x64_128<Key>>();
+  test_hash_function<Key, cuco::xxhash_32<Key>>();
+  test_hash_function<Key, cuco::xxhash_64<Key>>();
+}
\ No newline at end of file

From 56efe03f9412438f1a7d8ee89ccff4cf38816ef5 Mon Sep 17 00:00:00 2001
From: Srinivas Yadav Singanaboina <ssinganaboin@nvidia.com>
Date: Wed, 12 Jun 2024 20:50:51 +0000
Subject: [PATCH 2/6] move back sanitize_hash to utils.cuh

---
 include/cuco/detail/probing_scheme_base.cuh | 28 ------------
 include/cuco/detail/probing_scheme_impl.inl | 19 +++-----
 include/cuco/detail/utils.cuh               | 49 +++++++++++++++++++++
 3 files changed, 55 insertions(+), 41 deletions(-)

diff --git a/include/cuco/detail/probing_scheme_base.cuh b/include/cuco/detail/probing_scheme_base.cuh
index 9ce06da92..067f3f01a 100644
--- a/include/cuco/detail/probing_scheme_base.cuh
+++ b/include/cuco/detail/probing_scheme_base.cuh
@@ -32,34 +32,6 @@ namespace detail {
  */
 template <int32_t CGSize>
 class probing_scheme_base {
- private:
-  template <typename SizeType, typename HashType>
-  __host__ __device__ constexpr SizeType sanitize_hash_positive(HashType hash) const noexcept
-  {
-    if constexpr (cuda::std::is_signed_v<SizeType>) {
-      return cuda::std::abs(static_cast<SizeType>(hash));
-    } else {
-      return static_cast<SizeType>(hash);
-    }
-  }
-
- protected:
-  template <typename SizeType, typename HashType>
-  __host__ __device__ constexpr SizeType sanitize_hash(HashType hash) const noexcept
-  {
-    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));
-      return sanitize_hash_positive<SizeType>(static_cast<SizeType>(ret));
-    } else
-      return sanitize_hash_positive<SizeType>(hash);
-  }
-
  public:
   /**
    * @brief The size of the CUDA cooperative thread group.
diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
index 33998168e..3f4961a04 100644
--- a/include/cuco/detail/probing_scheme_impl.inl
+++ b/include/cuco/detail/probing_scheme_impl.inl
@@ -107,7 +107,7 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    probing_scheme_base_type::template sanitize_hash<size_type>(hash_(probe_key)) % upper_bound,
+    cuco::detail::sanitize_hash<size_type>(hash_(probe_key)) % upper_bound,
     1,  // step size is 1
     upper_bound};
 }
@@ -121,10 +121,7 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    probing_scheme_base_type::template sanitize_hash<size_type>(
-      probing_scheme_base_type::template sanitize_hash<size_type>(hash_(probe_key)) +
-      g.thread_rank()) %
-      upper_bound,
+    cuco::detail::sanitize_hash<size_type>(hash_(probe_key), g.thread_rank()) % upper_bound,
     cg_size,
     upper_bound};
 }
@@ -151,9 +148,9 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    probing_scheme_base_type::template sanitize_hash<size_type>(hash1_(probe_key)) % upper_bound,
+    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key)) % upper_bound,
     max(size_type{1},
-        probing_scheme_base_type::template sanitize_hash<size_type>(hash2_(probe_key)) %
+        cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) %
           upper_bound),  // step size in range [1, prime - 1]
     upper_bound};
 }
@@ -167,13 +164,9 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    probing_scheme_base_type::template sanitize_hash<size_type>(
-      probing_scheme_base_type::template sanitize_hash<size_type>(hash1_(probe_key)) +
-      g.thread_rank()) %
-      upper_bound,
+    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key), g.thread_rank()) % upper_bound,
     static_cast<size_type>(
-      (probing_scheme_base_type::template sanitize_hash<size_type>(hash2_(probe_key)) %
-         (upper_bound / cg_size - 1) +
+      (cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) % (upper_bound / cg_size - 1) +
        1) *
       cg_size),
     upper_bound};  // TODO use fast_int operator
diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index f2aecc0ef..5e0a0d6e0 100644
--- a/include/cuco/detail/utils.cuh
+++ b/include/cuco/detail/utils.cuh
@@ -82,5 +82,54 @@ struct slot_is_filled {
   }
 };
 
+template <typename SizeType, typename HashType>
+__host__ __device__ constexpr SizeType to_positive(HashType hash)
+{
+  if constexpr (cuda::std::is_signed_v<SizeType>) {
+    return cuda::std::abs(static_cast<SizeType>(hash));
+  } else {
+    return static_cast<SizeType>(hash);
+  }
+}
+
+/**
+ * @brief Converts a given hash value into a valid (positive) size type.
+ *
+ * @tparam SizeType The target type
+ * @tparam HashType The input type
+ *
+ * @return Converted hash value
+ */
+template <typename SizeType, typename HashType>
+__host__ __device__ constexpr SizeType sanitize_hash(HashType hash) noexcept
+{
+  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));
+    return to_positive<SizeType>(static_cast<SizeType>(ret));
+  } else {
+    return to_positive<SizeType>(hash);
+  }
+}
+
+/**
+ * @brief Converts a given hash value and cg_rank, into a valid (positive) size type.
+ *
+ * @tparam SizeType The target type
+ * @tparam HashType The input type
+ *
+ * @return Converted hash value
+ */
+template <typename SizeType, typename HashType>
+__host__ __device__ constexpr SizeType sanitize_hash(HashType hash, std::uint32_t cg_rank) noexcept
+{
+  return sanitize_hash<SizeType>(sanitize_hash<SizeType>(hash) + cg_rank);
+}
+
 }  // namespace detail
 }  // namespace cuco

From 3ae1afeb0e2b55d1b2ede3496c5de82a74541ce0 Mon Sep 17 00:00:00 2001
From: Srinivas Yadav Singanaboina <ssinganaboin@nvidia.com>
Date: Tue, 18 Jun 2024 04:25:13 +0000
Subject: [PATCH 3/6] fix sanitize_hash edge case

---
 include/cuco/detail/probing_scheme_impl.inl |  4 ++--
 include/cuco/detail/utils.cuh               | 15 ++++++++++++---
 2 files changed, 14 insertions(+), 5 deletions(-)

diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
index 3f4961a04..3a13cd25b 100644
--- a/include/cuco/detail/probing_scheme_impl.inl
+++ b/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) *
diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index 5e0a0d6e0..4324fd892 100644
--- a/include/cuco/detail/utils.cuh
+++ b/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

From c20316879e7f17b294f7aaee1ac5d240454f2e89 Mon Sep 17 00:00:00 2001
From: Srinivas Yadav Singanaboina <ssinganaboin@nvidia.com>
Date: Tue, 18 Jun 2024 17:36:16 +0000
Subject: [PATCH 4/6] minor styling changes with CG

---
 include/cuco/detail/probing_scheme_base.cuh | 2 --
 include/cuco/detail/probing_scheme_impl.inl | 6 ++++--
 include/cuco/detail/utils.cuh               | 6 +++---
 3 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/include/cuco/detail/probing_scheme_base.cuh b/include/cuco/detail/probing_scheme_base.cuh
index 067f3f01a..a3d7c148a 100644
--- a/include/cuco/detail/probing_scheme_base.cuh
+++ b/include/cuco/detail/probing_scheme_base.cuh
@@ -16,8 +16,6 @@
 
 #pragma once
 
-#include <cuco/detail/__config>
-
 #include <cstdint>
 
 namespace cuco {
diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
index 3a13cd25b..f004eee67 100644
--- a/include/cuco/detail/probing_scheme_impl.inl
+++ b/include/cuco/detail/probing_scheme_impl.inl
@@ -120,8 +120,9 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
   Extent upper_bound) const noexcept
 {
   using size_type = typename Extent::value_type;
+  using cg_type   = cooperative_groups::thread_block_tile<cg_size>;
   return detail::probing_iterator<Extent>{
-    cuco::detail::sanitize_hash<size_type>(hash_(probe_key), g) % upper_bound,
+    cuco::detail::sanitize_hash<cg_type, size_type>(g, hash_(probe_key)) % upper_bound,
     cg_size,
     upper_bound};
 }
@@ -163,8 +164,9 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
   Extent upper_bound) const noexcept
 {
   using size_type = typename Extent::value_type;
+  using cg_type   = cooperative_groups::thread_block_tile<cg_size>;
   return detail::probing_iterator<Extent>{
-    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key), g) % upper_bound,
+    cuco::detail::sanitize_hash<cg_type, size_type>(g, hash1_(probe_key)) % upper_bound,
     static_cast<size_type>(
       (cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) % (upper_bound / cg_size - 1) +
        1) *
diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index 4324fd892..5b210b90a 100644
--- a/include/cuco/detail/utils.cuh
+++ b/include/cuco/detail/utils.cuh
@@ -123,14 +123,14 @@ __host__ __device__ constexpr SizeType sanitize_hash(HashType hash) noexcept
 /**
  * @brief Converts a given hash value and cg_rank, into a valid (positive) size type.
  *
+ * @tparam CG Cooperative group type
  * @tparam SizeType The target type
  * @tparam HashType The input type
- * @tparam CG Cooperative group type
  *
  * @return Converted hash value
  */
-template <typename SizeType, typename HashType, typename CG>
-__host__ __device__ constexpr SizeType sanitize_hash(HashType hash, CG group) noexcept
+template <typename CG, typename SizeType, typename HashType>
+__device__ constexpr SizeType sanitize_hash(CG const& group, HashType hash) noexcept
 {
   auto const base_hash = sanitize_hash<SizeType>(hash);
   auto const max_size  = cuda::std::numeric_limits<SizeType>::max();

From 8feddcad4367bbd8549f89b0f7be76b18a366dc4 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Daniel=20J=C3=BCnger?= <djuenger@nvidia.com>
Date: Wed, 19 Jun 2024 04:12:46 +0200
Subject: [PATCH 5/6] Always use curly braces

Co-authored-by: Yunsong Wang <yunsongw@nvidia.com>
---
 include/cuco/detail/utils.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index 5b210b90a..36105c0b0 100644
--- a/include/cuco/detail/utils.cuh
+++ b/include/cuco/detail/utils.cuh
@@ -136,7 +136,7 @@ __device__ constexpr SizeType sanitize_hash(CG const& group, HashType hash) noex
   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);
+  if (base_hash > (max_size - cg_rank)) { return cg_rank - (max_size - base_hash); }
   return base_hash + cg_rank;
 }
 

From 6657e23bed4496b0870eb2393d86c8236d256fda Mon Sep 17 00:00:00 2001
From: Srinivas Yadav Singanaboina <ssinganaboin@nvidia.com>
Date: Wed, 19 Jun 2024 02:20:31 +0000
Subject: [PATCH 6/6] Re-order CG Type

---
 include/cuco/detail/probing_scheme_impl.inl | 6 ++----
 include/cuco/detail/utils.cuh               | 4 ++--
 2 files changed, 4 insertions(+), 6 deletions(-)

diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
index f004eee67..d0a67f87a 100644
--- a/include/cuco/detail/probing_scheme_impl.inl
+++ b/include/cuco/detail/probing_scheme_impl.inl
@@ -120,9 +120,8 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
   Extent upper_bound) const noexcept
 {
   using size_type = typename Extent::value_type;
-  using cg_type   = cooperative_groups::thread_block_tile<cg_size>;
   return detail::probing_iterator<Extent>{
-    cuco::detail::sanitize_hash<cg_type, size_type>(g, hash_(probe_key)) % upper_bound,
+    cuco::detail::sanitize_hash<size_type>(g, hash_(probe_key)) % upper_bound,
     cg_size,
     upper_bound};
 }
@@ -164,9 +163,8 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
   Extent upper_bound) const noexcept
 {
   using size_type = typename Extent::value_type;
-  using cg_type   = cooperative_groups::thread_block_tile<cg_size>;
   return detail::probing_iterator<Extent>{
-    cuco::detail::sanitize_hash<cg_type, size_type>(g, hash1_(probe_key)) % upper_bound,
+    cuco::detail::sanitize_hash<size_type>(g, hash1_(probe_key)) % upper_bound,
     static_cast<size_type>(
       (cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) % (upper_bound / cg_size - 1) +
        1) *
diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index 36105c0b0..21e4df759 100644
--- a/include/cuco/detail/utils.cuh
+++ b/include/cuco/detail/utils.cuh
@@ -123,13 +123,13 @@ __host__ __device__ constexpr SizeType sanitize_hash(HashType hash) noexcept
 /**
  * @brief Converts a given hash value and cg_rank, into a valid (positive) size type.
  *
- * @tparam CG Cooperative group type
  * @tparam SizeType The target type
+ * @tparam CG Cooperative group type
  * @tparam HashType The input type
  *
  * @return Converted hash value
  */
-template <typename CG, typename SizeType, typename HashType>
+template <typename SizeType, typename CG, typename HashType>
 __device__ constexpr SizeType sanitize_hash(CG const& group, HashType hash) noexcept
 {
   auto const base_hash = sanitize_hash<SizeType>(hash);