Remove fast_int inheritance from dynamic valid extents

include/cuco/detail/extent/extent.inl

@@ -10,7 +10,6 @@
 #include <cuco/detail/utility/math.cuh>
 #include <cuco/probing_scheme.cuh>
 #include <cuco/storage.cuh>
-#include <cuco/utility/fast_int.cuh>
 
 #include <cuda/std/type_traits>
 
@@ -57,65 +56,22 @@ constexpr std::uint64_t normalize_extent(SizeType size)
 }  // namespace detail
 
 template <typename SizeType, std::size_t N>
-struct valid_extent {
-  using value_type = SizeType;  ///< Extent value type
+class valid_extent : public extent<SizeType, N> {
+  using base_type = extent<SizeType, N>;
 
-  __host__ __device__ constexpr value_type value() const noexcept { return N; }
-  __host__ __device__ explicit constexpr operator value_type() const noexcept { return value(); }
+ public:
+  using value_type = typename base_type::value_type;
 
- private:
-  __host__ __device__ explicit constexpr valid_extent() noexcept {}
-  __host__ __device__ explicit constexpr valid_extent(SizeType) noexcept {}
-
-  // Friend declarations for all make_valid_extent overloads
-  template <int32_t CGSize_, int32_t BucketSize_, typename SizeType_, std::size_t N_>
-  friend auto constexpr make_valid_extent(extent<SizeType_, N_> ext);
-
-  template <typename ProbingScheme, typename Storage, typename SizeType_, std::size_t N_>
-  friend auto constexpr make_valid_extent(extent<SizeType_, N_> ext);
-
-  template <template <typename> class ProbingScheme,
-            typename Storage,
-            typename SizeType_,
-            std::size_t N_>
-  friend auto constexpr make_valid_extent(extent<SizeType_, N_> ext);
-
-  template <template <typename, typename> class ProbingScheme,
-            typename Storage,
-            typename SizeType_,
-            std::size_t N_>
-  friend auto constexpr make_valid_extent(extent<SizeType_, N_> ext);
-
-  // Operator overloads
-  template <typename Rhs>
-  friend __host__ __device__ constexpr value_type operator-(valid_extent const& lhs,
-                                                            Rhs rhs) noexcept
+  __host__ __device__ constexpr value_type value() const noexcept
   {
-    return lhs.value() - rhs;
+    return base_type::operator value_type();
   }
 
-  template <typename Rhs>
-  friend __host__ __device__ constexpr value_type operator/(valid_extent const& lhs,
-                                                            Rhs rhs) noexcept
-  {
-    return lhs.value() / rhs;
-  }
-
-  template <typename Lhs>
-  friend __host__ __device__ constexpr value_type operator%(Lhs lhs,
-                                                            valid_extent const& rhs) noexcept
+ private:
+  __host__ __device__ explicit constexpr valid_extent(value_type value = {}) noexcept
+    : base_type{value}
   {
-    return lhs % rhs.value();
   }
-};
-
-template <typename SizeType>
-struct valid_extent<SizeType, dynamic_extent> : cuco::utility::fast_int<SizeType> {
-  using value_type =
-    typename cuco::utility::fast_int<SizeType>::fast_int::value_type;  ///< Extent value type
-
- private:
-  using cuco::utility::fast_int<SizeType>::fast_int;
 
   // Friend declarations for all make_valid_extent overloads
   template <int32_t CGSize_, int32_t BucketSize_, typename SizeType_, std::size_t N_>
@@ -137,6 +93,28 @@ struct valid_extent<SizeType, dynamic_extent> : cuco::utility::fast_int<SizeType
   friend auto constexpr make_valid_extent(extent<SizeType_, N_> ext);
 };
 
+// Operator overloads
+template <typename SizeType, std::size_t N, typename Rhs>
+__host__ __device__ constexpr typename valid_extent<SizeType, N>::value_type operator-(
+  valid_extent<SizeType, N> const& lhs, Rhs rhs) noexcept
+{
+  return lhs.value() - rhs;
+}
+
+template <typename SizeType, std::size_t N, typename Rhs>
+__host__ __device__ constexpr typename valid_extent<SizeType, N>::value_type operator/(
+  valid_extent<SizeType, N> const& lhs, Rhs rhs) noexcept
+{
+  return lhs.value() / rhs;
+}
+
+template <typename Lhs, typename SizeType, std::size_t N>
+__host__ __device__ constexpr typename valid_extent<SizeType, N>::value_type operator%(
+  Lhs lhs, valid_extent<SizeType, N> const& rhs) noexcept
+{
+  return lhs % rhs.value();
+}
+
 // Primary implementation for fixed CGSize and BucketSize
 template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
 [[nodiscard]] auto constexpr make_valid_extent(extent<SizeType, N> ext)

include/cuco/detail/probing_scheme/probing_scheme_impl.inl

@@ -193,7 +193,7 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::make_it
                               hash2_(probe_key), static_cast<size_type>(num_groups - 1)) +
                             1) *
                            stride),
-    upper_bound};  // TODO use fast_int operator
+    upper_bound};
 }
 
 template <int32_t CGSize, typename Hash1, typename Hash2>

tests/utility/extent_test.cu

@@ -54,6 +54,26 @@ TEMPLATE_TEST_CASE_SIG(
     REQUIRE(gold_reference == res.value());
   }
 
+  SECTION("Static valid extent supports arithmetic operators")
+  {
+    auto constexpr size = cuco::extent<SizeType, num>{};
+    auto constexpr res  = cuco::make_valid_extent<probing_t, storage_t>(size);
+
+    STATIC_REQUIRE((res - SizeType{10}) == gold_reference - 10);
+    STATIC_REQUIRE((res / SizeType{2}) == gold_reference / 2);
+    STATIC_REQUIRE((SizeType{5000} % res) == SizeType{5000} % gold_reference);
+  }
+
+  SECTION("Dynamic valid extent supports arithmetic operators")
+  {
+    auto const size = cuco::extent<SizeType>{num};
+    auto const res  = cuco::make_valid_extent<probing_t, storage_t>(size);
+
+    REQUIRE((res - SizeType{10}) == gold_reference - 10);
+    REQUIRE((res / SizeType{2}) == gold_reference / 2);
+    REQUIRE((SizeType{5000} % res) == SizeType{5000} % gold_reference);
+  }
+
   SECTION("Invalid desired load factor throws exception")
   {
     using probing_scheme_type = cuco::linear_probing<cg_size, cuco::default_hash_function<int>>;