minor styling changes with CG

Author: srinivasyadav18

include/cuco/detail/probing_scheme_base.cuh

@@ -16,8 +16,6 @@
 
 #pragma once
 
-#include <cuco/detail/__config>
-
 #include <cstdint>
 
 namespace cuco {

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) *

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();