Merge branch 'fix-static-extent-issues' into migrate-shared-memory-code

NVIDIA · Dec 5, 2023 · bd57dc9 · bd57dc9
2 parents 6f4e79d + f27f088
commit bd57dc9
Show file tree

Hide file tree

Showing 5 changed files with 17 additions and 21 deletions.
diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl
@@ -31,8 +31,7 @@ template <typename SizeType, std::size_t N>
 struct window_extent {
   using value_type = SizeType;  ///< Extent value type
 
-  __host__ __device__ constexpr value_type value() const noexcept { return N; }
-  __host__ __device__ explicit constexpr operator value_type() const noexcept { return value(); }
+  __host__ __device__ explicit constexpr operator value_type() const noexcept { return N; }
 
  private:
   __host__ __device__ explicit constexpr window_extent() noexcept {}

diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
@@ -64,7 +64,7 @@ class probing_iterator {
   {
     // TODO: step_size_ can be a build time constant (e.g. linear probing)
     //  Worth passing another extent type?
-    curr_index_ = (curr_index_ + step_size_) % upper_bound_;
+    curr_index_ = (curr_index_ + step_size_) % static_cast<size_type>(upper_bound_);
     return *this;
   }
 
@@ -100,7 +100,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,
+    cuco::detail::sanitize_hash<size_type>(hash_(probe_key)) % static_cast<size_type>(upper_bound),
     1,  // step size is 1
     upper_bound};
 }
@@ -114,7 +114,8 @@ __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.thread_rank()) %
+      static_cast<size_type>(upper_bound),
     cg_size,
     upper_bound};
 }
@@ -133,10 +134,10 @@ __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,
+    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key)) % static_cast<size_type>(upper_bound),
     max(size_type{1},
         cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) %
-          upper_bound),  // step size in range [1, prime - 1]
+          static_cast<size_type>(upper_bound)),  // step size in range [1, prime - 1]
     upper_bound};
 }
 
@@ -149,9 +150,10 @@ __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.thread_rank()) %
+      static_cast<size_type>(upper_bound),
     static_cast<size_type>((cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) %
-                              (upper_bound.value() / cg_size - 1) +
+                              (static_cast<size_type>(upper_bound) / cg_size - 1) +
                             1) *
                            cg_size),
     upper_bound};  // TODO use fast_int operator

diff --git a/include/cuco/utility/fast_int.cuh b/include/cuco/utility/fast_int.cuh
@@ -51,13 +51,6 @@ struct fast_int {
     evaluate_magic_numbers();
   }
 
-  /**
-   * @brief Get the underlying integer value.
-   *
-   * @return Underlying value
-   */
-  __host__ __device__ constexpr value_type value() const noexcept { return value_; }
-
   /**
    * @brief Explicit conversion operator to the underlying value type.
    *
@@ -156,4 +149,4 @@ struct fast_int {
     return lhs - (lhs / rhs) * rhs.value_;
   }
 };
-}  // namespace cuco::utility
+}  // namespace cuco::utility
diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu
@@ -171,6 +171,8 @@ TEMPLATE_TEST_CASE_SIG(
   constexpr size_type gold_capacity = CGSize == 1 ? 422   // 211 x 1 x 2
                                                   : 412;  // 103 x 2 x 2
 
+  // XXX: test static extent on purpose, DO NOT CHANGE
+  using extent_type = cuco::experimental::extent<size_type, num_keys>;
   using probe =
     std::conditional_t<Probe == cuco::test::probe_sequence::linear_probing,
                        cuco::experimental::linear_probing<CGSize, cuco::murmurhash3_32<Key>>,
@@ -180,13 +182,13 @@ TEMPLATE_TEST_CASE_SIG(
 
   auto map = cuco::experimental::static_map<Key,
                                             Value,
-                                            cuco::experimental::extent<size_type>,
+                                            extent_type,
                                             cuda::thread_scope_device,
                                             thrust::equal_to<Key>,
                                             probe,
                                             cuco::cuda_allocator<std::byte>,
                                             cuco::experimental::storage<2>>{
-    num_keys, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
+    extent_type{}, cuco::empty_key<Key>{-1}, cuco::empty_value<Value>{-1}};
 
   REQUIRE(map.capacity() == gold_capacity);
 

diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu
@@ -44,13 +44,13 @@ TEMPLATE_TEST_CASE_SIG(
   {
     auto constexpr size = cuco::experimental::extent<SizeType, num>{};
     auto constexpr res  = cuco::experimental::make_window_extent<cg_size, window_size>(size);
-    STATIC_REQUIRE(gold_reference == res.value());
+    STATIC_REQUIRE(gold_reference == static_cast<SizeType>(res));
   }
 
   SECTION("Compute dynamic valid extent at run time.")
   {
     auto const size = cuco::experimental::extent<SizeType>{num};
     auto const res  = cuco::experimental::make_window_extent<cg_size, window_size>(size);
-    REQUIRE(gold_reference == res.value());
+    REQUIRE(gold_reference == static_cast<SizeType>(res));
   }
 }