diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl
new file mode 100644
index 000000000..9d79bc907
--- /dev/null
+++ b/include/cuco/detail/extent/extent.inl
@@ -0,0 +1,126 @@
+/*
+ * Copyright (c) 2023, 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.
+ */
+
+#pragma once
+
+#include <cuco/detail/error.hpp>
+#include <cuco/detail/prime.hpp>  // TODO move to detail/extent/
+#include <cuco/detail/utils.hpp>
+#include <cuco/utility/fast_int.cuh>
+
+#include <type_traits>
+
+namespace cuco {
+namespace experimental {
+
+template <int32_t CGSize, int32_t WindowSize, typename SizeType, std::size_t N = dynamic_extent>
+struct window_extent {
+  using value_type = SizeType;  ///< Extent value type
+
+  static auto constexpr cg_size     = CGSize;
+  static auto constexpr window_size = WindowSize;
+
+  __host__ __device__ constexpr value_type value() const noexcept { return N; }
+  __host__ __device__ explicit constexpr operator value_type() const noexcept { return value(); }
+
+ private:
+  __host__ __device__ explicit constexpr window_extent() noexcept {}
+  __host__ __device__ explicit constexpr window_extent(SizeType) noexcept {}
+
+  template <int32_t CGSize_, int32_t WindowSize_, typename SizeType_, std::size_t N_>
+  friend auto constexpr make_window_extent(extent<SizeType_, N_> ext);
+};
+
+template <int32_t CGSize, int32_t WindowSize, typename SizeType>
+struct window_extent<CGSize, WindowSize, SizeType, dynamic_extent>
+  : cuco::utility::fast_int<SizeType> {
+  using value_type =
+    typename cuco::utility::fast_int<SizeType>::fast_int::value_type;  ///< Extent value type
+
+  static auto constexpr cg_size     = CGSize;
+  static auto constexpr window_size = WindowSize;
+
+ private:
+  using cuco::utility::fast_int<SizeType>::fast_int;
+
+  template <int32_t CGSize_, int32_t WindowSize_, typename SizeType_, std::size_t N_>
+  friend auto constexpr make_window_extent(extent<SizeType_, N_> ext);
+};
+
+template <typename Container, typename SizeType, std::size_t N>
+[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext)
+{
+  return make_window_extent<Container::cg_size, Container::window_size>(ext);
+}
+
+template <typename Container>
+[[nodiscard]] std::size_t constexpr make_window_extent(std::size_t size)
+{
+  return make_window_extent<Container::cg_size, Container::window_size>(size);
+}
+
+template <int32_t CGSize, int32_t WindowSize, typename SizeType, std::size_t N>
+[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext)
+{
+  auto constexpr max_prime = cuco::detail::primes.back();
+  auto constexpr max_value =
+    (static_cast<uint64_t>(std::numeric_limits<SizeType>::max()) < max_prime)
+      ? std::numeric_limits<SizeType>::max()
+      : static_cast<SizeType>(max_prime);
+  auto const size = SDIV(ext, CGSize * WindowSize);
+  if (size <= 0 or size > max_value) { CUCO_FAIL("Invalid input extent"); }
+
+  if constexpr (N == dynamic_extent) {
+    return window_extent<CGSize, WindowSize, SizeType>{static_cast<SizeType>(
+      *cuco::detail::lower_bound(
+        cuco::detail::primes.begin(), cuco::detail::primes.end(), static_cast<uint64_t>(size)) *
+      CGSize)};
+  }
+  if constexpr (N != dynamic_extent) {
+    return window_extent<CGSize,
+                         WindowSize,
+                         SizeType,
+                         static_cast<std::size_t>(
+                           *cuco::detail::lower_bound(cuco::detail::primes.begin(),
+                                                      cuco::detail::primes.end(),
+                                                      static_cast<uint64_t>(size)) *
+                           CGSize)>{};
+  }
+}
+
+template <int32_t CGSize, int32_t WindowSize>
+[[nodiscard]] std::size_t constexpr make_window_extent(std::size_t size)
+{
+  return static_cast<std::size_t>(make_window_extent<CGSize, WindowSize>(extent{size}));
+}
+
+namespace detail {
+
+template <typename...>
+struct is_window_extent : std::false_type {
+};
+
+template <int32_t CGSize, int32_t WindowSize, typename SizeType, std::size_t N>
+struct is_window_extent<window_extent<CGSize, WindowSize, SizeType, N>> : std::true_type {
+};
+
+template <typename T>
+inline constexpr bool is_window_extent_v = is_window_extent<T>::value;
+
+}  // namespace detail
+
+}  // namespace experimental
+}  // namespace cuco
\ No newline at end of file
diff --git a/include/cuco/detail/open_addressing_impl.cuh b/include/cuco/detail/open_addressing_impl.cuh
index 2c8688869..dbf169dca 100644
--- a/include/cuco/detail/open_addressing_impl.cuh
+++ b/include/cuco/detail/open_addressing_impl.cuh
@@ -88,7 +88,7 @@ class open_addressing_impl {
   using key_type   = Key;    ///< Key type
   using value_type = Value;  ///< The storage value type, NOT payload type
   /// Extent type
-  using extent_type = decltype(make_valid_extent<cg_size, window_size>(std::declval<Extent>()));
+  using extent_type = decltype(make_window_extent<open_addressing_impl>(std::declval<Extent>()));
   using size_type   = typename extent_type::value_type;  ///< Size type
   using key_equal   = KeyEqual;                          ///< Key equality comparator type
   using storage_type =
@@ -103,7 +103,7 @@ class open_addressing_impl {
    * capacity, sentinel values and CUDA stream.
    *
    * @note The actual capacity depends on the given `capacity`, the probing scheme, CG size, and the
-   * window size and it's computed via `make_valid_extent` factory. Insert operations will not
+   * window size and it is computed via the `make_window_extent` factory. Insert operations will not
    * automatically grow the container. Attempting to insert more unique keys than the capacity of
    * the container results in undefined behavior.
    * @note Any `*_sentinel`s are reserved and behavior is undefined when attempting to insert
@@ -130,7 +130,7 @@ class open_addressing_impl {
       empty_slot_sentinel_{empty_slot_sentinel},
       predicate_{pred},
       probing_scheme_{probing_scheme},
-      storage_{make_valid_extent<cg_size, window_size>(capacity), alloc}
+      storage_{make_window_extent<open_addressing_impl>(capacity), alloc}
   {
     this->clear_async(stream);
   }
diff --git a/include/cuco/detail/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing_ref_impl.cuh
index 0e9cfdbe4..99187cc51 100644
--- a/include/cuco/detail/open_addressing_ref_impl.cuh
+++ b/include/cuco/detail/open_addressing_ref_impl.cuh
@@ -17,6 +17,7 @@
 #pragma once
 
 #include <cuco/detail/equal_wrapper.cuh>
+#include <cuco/extent.cuh>
 #include <cuco/pair.cuh>
 
 #include <thrust/distance.h>
@@ -62,6 +63,13 @@ class open_addressing_ref_impl {
                       ProbingScheme>,
     "ProbingScheme must inherit from cuco::detail::probing_scheme_base");
 
+  static_assert(is_window_extent_v<typename StorageRef::extent_type>,
+                "Extent is not a valid cuco::window_extent");
+  static_assert(ProbingScheme::cg_size == StorageRef::extent_type::cg_size,
+                "Extent has incompatible CG size");
+  static_assert(StorageRef::window_size == StorageRef::extent_type::window_size,
+                "Extent has incompatible window size");
+
  public:
   using key_type            = Key;                                     ///< Key type
   using probing_scheme_type = ProbingScheme;                           ///< Type of probing scheme
@@ -138,7 +146,7 @@ class open_addressing_ref_impl {
                          Predicate const& predicate) noexcept
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
-    auto probing_iter = probing_scheme_(key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(key, storage_ref_.window_extent());
 
     while (true) {
       auto const window_slots = storage_ref_[*probing_iter];
@@ -180,7 +188,7 @@ class open_addressing_ref_impl {
                          value_type const& value,
                          Predicate const& predicate) noexcept
   {
-    auto probing_iter = probing_scheme_(group, key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(group, key, storage_ref_.window_extent());
 
     while (true) {
       auto const window_slots = storage_ref_[*probing_iter];
@@ -244,7 +252,7 @@ class open_addressing_ref_impl {
                                                           Predicate const& predicate) noexcept
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
-    auto probing_iter = probing_scheme_(key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(key, storage_ref_.window_extent());
 
     while (true) {
       auto const window_slots = storage_ref_[*probing_iter];
@@ -301,7 +309,7 @@ class open_addressing_ref_impl {
     value_type const& value,
     Predicate const& predicate) noexcept
   {
-    auto probing_iter = probing_scheme_(group, key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(group, key, storage_ref_.window_extent());
 
     while (true) {
       auto const window_slots = storage_ref_[*probing_iter];
@@ -375,7 +383,7 @@ class open_addressing_ref_impl {
                                          Predicate const& predicate) const noexcept
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
-    auto probing_iter = probing_scheme_(key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(key, storage_ref_.window_extent());
 
     while (true) {
       // TODO atomic_ref::load if insert operator is present
@@ -413,7 +421,7 @@ class open_addressing_ref_impl {
     ProbeKey const& key,
     Predicate const& predicate) const noexcept
   {
-    auto probing_iter = probing_scheme_(group, key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(group, key, storage_ref_.window_extent());
 
     while (true) {
       auto const window_slots = storage_ref_[*probing_iter];
@@ -455,7 +463,7 @@ class open_addressing_ref_impl {
                                                Predicate const& predicate) const noexcept
   {
     static_assert(cg_size == 1, "Non-CG operation is incompatible with the current probing scheme");
-    auto probing_iter = probing_scheme_(key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(key, storage_ref_.window_extent());
 
     while (true) {
       // TODO atomic_ref::load if insert operator is present
@@ -497,7 +505,7 @@ class open_addressing_ref_impl {
        ProbeKey const& key,
        Predicate const& predicate) const noexcept
   {
-    auto probing_iter = probing_scheme_(group, key, storage_ref_.num_windows());
+    auto probing_iter = probing_scheme_(group, key, storage_ref_.window_extent());
 
     while (true) {
       auto const window_slots = storage_ref_[*probing_iter];
diff --git a/include/cuco/detail/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme_impl.inl
index 4b617a133..3090d026e 100644
--- a/include/cuco/detail/probing_scheme_impl.inl
+++ b/include/cuco/detail/probing_scheme_impl.inl
@@ -16,6 +16,8 @@
 
 #pragma once
 
+#include <cuco/detail/utils.cuh>
+
 namespace cuco {
 namespace experimental {
 namespace detail {
@@ -97,9 +99,10 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
   ProbeKey const& probe_key, Extent upper_bound) const noexcept
 {
   using size_type = typename Extent::value_type;
-  return detail::probing_iterator<Extent>{static_cast<size_type>(hash_(probe_key) % upper_bound),
-                                          1,  // step size is 1
-                                          upper_bound};
+  return detail::probing_iterator<Extent>{
+    cuco::detail::sanitize_hash<size_type>(hash_(probe_key)) % upper_bound,
+    1,  // step size is 1
+    upper_bound};
 }
 
 template <int32_t CGSize, typename Hash>
@@ -111,7 +114,7 @@ __host__ __device__ constexpr auto linear_probing<CGSize, Hash>::operator()(
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    static_cast<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};
 }
@@ -130,9 +133,10 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    static_cast<size_type>(hash1_(probe_key) % upper_bound),
-    static_cast<size_type>(hash2_(probe_key) % (upper_bound - 1) +
-                           1),  // step size in range [1, prime - 1]
+    cuco::detail::sanitize_hash<size_type>(hash1_(probe_key)) % upper_bound,
+    max(size_type{1},
+        cuco::detail::sanitize_hash<size_type>(hash2_(probe_key)) %
+          upper_bound),  // step size in range [1, prime - 1]
     upper_bound};
 }
 
@@ -145,9 +149,12 @@ __host__ __device__ constexpr auto double_hashing<CGSize, Hash1, Hash2>::operato
 {
   using size_type = typename Extent::value_type;
   return detail::probing_iterator<Extent>{
-    static_cast<size_type>((hash1_(probe_key) + g.thread_rank()) % upper_bound),
-    static_cast<size_type>((hash2_(probe_key) % (upper_bound / cg_size - 1) + 1) * cg_size),
-    upper_bound};
+    cuco::detail::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.value() / cg_size - 1) +
+                            1) *
+                           cg_size),
+    upper_bound};  // TODO use fast_int operator
 }
 }  // namespace experimental
 }  // namespace cuco
diff --git a/include/cuco/detail/storage/aow_storage.cuh b/include/cuco/detail/storage/aow_storage.cuh
index ac86508de..1f01dad37 100644
--- a/include/cuco/detail/storage/aow_storage.cuh
+++ b/include/cuco/detail/storage/aow_storage.cuh
@@ -67,7 +67,7 @@ class aow_storage_base : public storage_base<Extent> {
    *
    * @return The total number of slot windows
    */
-  [[nodiscard]] __host__ __device__ constexpr extent_type num_windows() const noexcept
+  [[nodiscard]] __host__ __device__ constexpr size_type num_windows() const noexcept
   {
     return storage_base<Extent>::capacity();
   }
@@ -77,9 +77,19 @@ class aow_storage_base : public storage_base<Extent> {
    *
    * @return The total number of slots
    */
-  [[nodiscard]] __host__ __device__ constexpr auto capacity() const noexcept
+  [[nodiscard]] __host__ __device__ constexpr size_type capacity() const noexcept
   {
-    return storage_base<Extent>::capacity().template multiply<window_size>();
+    return storage_base<Extent>::capacity() * window_size;
+  }
+
+  /**
+   * @brief Gets the window extent of the current storage.
+   *
+   * @return The window extent.
+   */
+  [[nodiscard]] __host__ __device__ constexpr extent_type window_extent() const noexcept
+  {
+    return storage_base<Extent>::extent();
   }
 };
 
@@ -278,7 +288,7 @@ class aow_storage : public aow_storage_base<WindowSize, T, Extent> {
    * @brief Constructor of AoW storage.
    *
    * @note The input `size` should be exclusively determined by the return value of
-   * `make_valid_extent` since it depends on the requested low-bound value, the probing scheme, and
+   * `make_window_extent` since it depends on the requested low-bound value, the probing scheme, and
    * the storage.
    *
    * @param size Number of windows to (de)allocate
@@ -325,7 +335,7 @@ class aow_storage : public aow_storage_base<WindowSize, T, Extent> {
    */
   [[nodiscard]] constexpr ref_type ref() const noexcept
   {
-    return ref_type{this->num_windows(), this->data()};
+    return ref_type{this->window_extent(), this->data()};
   }
 
   /**
diff --git a/include/cuco/detail/storage/counter_storage.cuh b/include/cuco/detail/storage/counter_storage.cuh
index 12c963530..bb36b15e2 100644
--- a/include/cuco/detail/storage/counter_storage.cuh
+++ b/include/cuco/detail/storage/counter_storage.cuh
@@ -38,7 +38,7 @@ namespace detail {
 template <typename SizeType, cuda::thread_scope Scope, typename Allocator>
 class counter_storage : public storage_base<cuco::experimental::extent<SizeType, 1>> {
  public:
-  using storage_base<cuco::experimental::extent<SizeType, 1>>::capacity_;  ///< Storage size
+  using storage_base<cuco::experimental::extent<SizeType, 1>>::capacity;  ///< Storage size
 
   using size_type      = SizeType;                        ///< Size type
   using value_type     = cuda::atomic<size_type, Scope>;  ///< Type of the counter
@@ -56,8 +56,8 @@ class counter_storage : public storage_base<cuco::experimental::extent<SizeType,
     : storage_base<cuco::experimental::extent<SizeType, 1>>{cuco::experimental::extent<size_type,
                                                                                        1>{}},
       allocator_{allocator},
-      counter_deleter_{capacity_, allocator_},
-      counter_{allocator_.allocate(capacity_), counter_deleter_}
+      counter_deleter_{this->capacity(), allocator_},
+      counter_{allocator_.allocate(this->capacity()), counter_deleter_}
   {
   }
 
diff --git a/include/cuco/detail/storage/storage_base.cuh b/include/cuco/detail/storage/storage_base.cuh
index ada0726db..15ec30472 100644
--- a/include/cuco/detail/storage/storage_base.cuh
+++ b/include/cuco/detail/storage/storage_base.cuh
@@ -71,20 +71,30 @@ class storage_base {
    *
    * @param size Number of elements to (de)allocate
    */
-  explicit constexpr storage_base(Extent size) : capacity_{size} {}
+  explicit constexpr storage_base(Extent size) : extent_{size} {}
 
   /**
    * @brief Gets the total number of elements in the current storage.
    *
    * @return The total number of elements
    */
-  [[nodiscard]] __host__ __device__ constexpr extent_type capacity() const noexcept
+  [[nodiscard]] __host__ __device__ constexpr size_type capacity() const noexcept
   {
-    return capacity_;
+    return static_cast<size_type>(extent_);
+  }
+
+  /**
+   * @brief Gets the extent of the current storage.
+   *
+   * @return The extent.
+   */
+  [[nodiscard]] __host__ __device__ constexpr extent_type extent() const noexcept
+  {
+    return extent_;
   }
 
  protected:
-  extent_type capacity_;  ///< Total number of elements
+  extent_type extent_;  ///< Total number of elements
 };
 
 }  // namespace detail
diff --git a/include/cuco/detail/utils.cuh b/include/cuco/detail/utils.cuh
index fdded70f5..793854694 100644
--- a/include/cuco/detail/utils.cuh
+++ b/include/cuco/detail/utils.cuh
@@ -20,6 +20,8 @@
 #include <thrust/tuple.h>
 
 #include <cuda/std/bit>
+#include <cuda/std/cmath>
+#include <cuda/std/type_traits>
 
 namespace cuco {
 namespace detail {
@@ -104,6 +106,24 @@ struct strong_type {
   T value;  ///< Underlying value
 };
 
+/**
+ * @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);
+  }
+}
+
 /**
  * @brief Gives value to use as alignment for a pair type that is at least the
  * size of the sum of the size of the first type and second type, or 16,
diff --git a/include/cuco/extent.cuh b/include/cuco/extent.cuh
index b825188ed..a08f77a39 100644
--- a/include/cuco/extent.cuh
+++ b/include/cuco/extent.cuh
@@ -16,8 +16,6 @@
 
 #pragma once
 
-#include <cuco/detail/prime.hpp>
-
 #include <cstddef>
 #include <cstdint>
 
@@ -46,19 +44,6 @@ struct extent {
    * @return Extent size
    */
   __host__ __device__ constexpr operator value_type() const noexcept { return N; }
-
-  /**
-   * @brief Multiplies the current extent with the given `Value`.
-   *
-   * @tparam Value The input value to multiply with
-   *
-   * @return Resulting static extent
-   */
-  template <int32_t Value>
-  __host__ __device__ constexpr auto multiply() const noexcept
-  {
-    return extent<value_type, N * Value>{};
-  }
 };
 
 /**
@@ -84,25 +69,67 @@ struct extent<SizeType, dynamic_extent> {
    */
   __host__ __device__ constexpr operator value_type() const noexcept { return value_; }
 
-  /**
-   * @brief Multiplies the current extent with the given `Value`.
-   *
-   * @tparam Value The input value to multiply with
-   *
-   * @return Resulting extent
-   */
-  template <int32_t Value>
-  __host__ __device__ constexpr auto multiply() const noexcept
-  {
-    return extent<value_type>{Value * value_};
-  }
-
  private:
   value_type value_;  ///< Extent value
 };
 
 /**
- * @brief Computes valid extent based on given parameters.
+ * @brief Window extent strong type.
+ *
+ * @note This type is used internally and can only be constructed using the `make_window_extent'
+ * factory method.
+ *
+ * @tparam SizeType Size type
+ * @tparam N Extent
+ *
+ */
+template <int32_t CGSize, int32_t WindowSize, typename SizeType, std::size_t N>
+struct window_extent;
+
+/**
+ * @brief Computes a valid window extent/capacity for a given container type.
+ *
+ * @note The actual capacity of a container (map/set) should be exclusively determined by the return
+ * value of this utility since the output depends on the requested low-bound size, the probing
+ * scheme, and the storage. This utility is used internally during container constructions while for
+ * container ref constructions, it would be users' responsibility to use this function to determine
+ * the capacity ctor argument for the container.
+ *
+ * @tparam Container Container type to compute the extent for
+ * @tparam SizeType Size type
+ * @tparam N Extent
+ *
+ * @param ext The input extent
+ *
+ * @throw If the input extent is invalid
+ *
+ * @return Resulting valid `window extent`
+ */
+template <typename Container, typename SizeType, std::size_t N>
+[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext);
+
+/**
+ * @brief Computes a valid capacity for a given container type.
+ *
+ * @note The actual capacity of a container (map/set) should be exclusively determined by the return
+ * value of this utility since the output depends on the requested low-bound size, the probing
+ * scheme, and the storage. This utility is used internally during container constructions while for
+ * container ref constructions, it would be users' responsibility to use this function to determine
+ * the capacity ctor argument for the container.
+ *
+ * @tparam Container Container type to compute the extent for
+ *
+ * @param size The input size
+ *
+ * @throw If the input size is invalid
+ *
+ * @return Resulting valid extent as `std::size_t`
+ */
+template <typename Container>
+[[nodiscard]] std::size_t constexpr make_window_extent(std::size_t size);
+
+/**
+ * @brief Computes valid window extent based on given parameters.
  *
  * @note The actual capacity of a container (map/set) should be exclusively determined by the return
  * value of this utility since the output depends on the requested low-bound size, the probing
@@ -115,35 +142,37 @@ struct extent<SizeType, dynamic_extent> {
  * @tparam SizeType Size type
  * @tparam N Extent
  *
+ * @param ext The input extent
+ *
  * @throw If the input extent is invalid
  *
  * @return Resulting valid extent
  */
 template <int32_t CGSize, int32_t WindowSize, typename SizeType, std::size_t N>
-[[nodiscard]] auto constexpr make_valid_extent(extent<SizeType, N> ext)
-{
-  auto constexpr max_prime = cuco::detail::primes.back();
-  auto constexpr max_value =
-    (static_cast<uint64_t>(std::numeric_limits<SizeType>::max()) < max_prime)
-      ? std::numeric_limits<SizeType>::max()
-      : static_cast<SizeType>(max_prime);
-  auto const size = SDIV(ext, CGSize * WindowSize);
-  if (size <= 0 or size > max_value) { CUCO_FAIL("Invalid input extent"); }
-
-  if constexpr (N == dynamic_extent) {
-    return extent<SizeType>{static_cast<SizeType>(
-      *cuco::detail::lower_bound(
-        cuco::detail::primes.begin(), cuco::detail::primes.end(), static_cast<uint64_t>(size)) *
-      CGSize)};
-  }
-  if constexpr (N != dynamic_extent) {
-    return extent<SizeType,
-                  static_cast<std::size_t>(*cuco::detail::lower_bound(cuco::detail::primes.begin(),
-                                                                      cuco::detail::primes.end(),
-                                                                      static_cast<uint64_t>(size)) *
-                                           CGSize)>{};
-  }
-}
+[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext);
+
+/**
+ * @brief Computes valid window extent/capacity based on given parameters.
+ *
+ * @note The actual capacity of a container (map/set) should be exclusively determined by the return
+ * value of this utility since the output depends on the requested low-bound size, the probing
+ * scheme, and the storage. This utility is used internally during container constructions while for
+ * container ref constructions, it would be users' responsibility to use this function to determine
+ * the capacity ctor argument for the container.
+ *
+ * @tparam CGSize Number of elements handled per CG
+ * @tparam WindowSize Number of elements handled per Window
+ *
+ * @param size The input size
+ *
+ * @throw If the input size is invalid
+ *
+ * @return Resulting valid extent as `std::size_t`
+ */
+template <int32_t CGSize, int32_t WindowSize>
+[[nodiscard]] std::size_t constexpr make_window_extent(std::size_t size);
 
 }  // namespace experimental
 }  // namespace cuco
+
+#include <cuco/detail/extent/extent.inl>
diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh
index a14325c84..38f3b92c9 100644
--- a/include/cuco/static_map.cuh
+++ b/include/cuco/static_map.cuh
@@ -158,7 +158,7 @@ class static_map {
    * and CUDA stream.
    *
    * The actual map capacity depends on the given `capacity`, the probing scheme, CG size, and the
-   * window size and it's computed via `make_valid_extent` factory. Insert operations will not
+   * window size and it is computed via the `make_window_extent` factory. Insert operations will not
    * automatically grow the map. Attempting to insert more unique keys than the capacity of the map
    * results in undefined behavior.
    *
diff --git a/include/cuco/static_set.cuh b/include/cuco/static_set.cuh
index 9aa67a498..0cb558a65 100644
--- a/include/cuco/static_set.cuh
+++ b/include/cuco/static_set.cuh
@@ -134,7 +134,7 @@ class static_set {
    * and CUDA stream.
    *
    * The actual set capacity depends on the given `capacity`, the probing scheme, CG size, and the
-   * window size and it's computed via `make_valid_extent` factory. Insert operations will not
+   * window size and it is computed via the `make_window_extent` factory. Insert operations will not
    * automatically grow the set. Attempting to insert more unique keys than the capacity of the map
    * results in undefined behavior.
    *
diff --git a/include/cuco/utility/fast_int.cuh b/include/cuco/utility/fast_int.cuh
index 5cd2998f6..6616e2c5c 100644
--- a/include/cuco/utility/fast_int.cuh
+++ b/include/cuco/utility/fast_int.cuh
@@ -51,6 +51,13 @@ 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.
    *
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index ebc37e39b..a72ff52d5 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -58,7 +58,6 @@ ConfigureTest(UTILITY_TEST
 ###################################################################################################
 # - static_set tests ------------------------------------------------------------------------------
 ConfigureTest(STATIC_SET_TEST
-    static_set/capacity_test.cu
     static_set/heterogeneous_lookup_test.cu
     static_set/insert_and_find_test.cu
     static_set/large_input_test.cu
diff --git a/tests/static_set/capacity_test.cu b/tests/static_set/capacity_test.cu
deleted file mode 100644
index e144325d5..000000000
--- a/tests/static_set/capacity_test.cu
+++ /dev/null
@@ -1,102 +0,0 @@
-/*
- * Copyright (c) 2023, 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 <cuco/static_set.cuh>
-
-#include <catch2/catch_test_macros.hpp>
-
-TEST_CASE("Static set capacity", "")
-{
-  constexpr std::size_t num_keys{400};
-  using Key        = int32_t;
-  using ProbeT     = cuco::experimental::double_hashing<1, cuco::default_hash_function<Key>>;
-  using Equal      = thrust::equal_to<Key>;
-  using AllocatorT = cuco::cuda_allocator<std::byte>;
-  using StorageT   = cuco::experimental::aow_storage<2>;
-
-  SECTION("Static extent must be evaluated at compile time.")
-  {
-    auto constexpr gold_capacity = 422;  // 211 x 2
-
-    using extent_type = cuco::experimental::extent<std::size_t, num_keys>;
-    cuco::experimental::
-      static_set<Key, extent_type, cuda::thread_scope_device, Equal, ProbeT, AllocatorT, StorageT>
-        set{extent_type{}, cuco::empty_key<Key>{-1}};
-    auto const capacity = set.capacity();
-    STATIC_REQUIRE(capacity == gold_capacity);
-
-    auto ref                = set.ref(cuco::experimental::insert);
-    auto const ref_capacity = ref.capacity();
-    STATIC_REQUIRE(ref_capacity == gold_capacity);
-  }
-
-  SECTION("Dynamic extent is evaluated at run time.")
-  {
-    auto constexpr gold_capacity = 422;  // 211 x 2
-
-    using extent_type = cuco::experimental::extent<std::size_t>;
-    cuco::experimental::
-      static_set<Key, extent_type, cuda::thread_scope_device, Equal, ProbeT, AllocatorT, StorageT>
-        set{num_keys, cuco::empty_key<Key>{-1}};
-    auto const capacity = set.capacity();
-    REQUIRE(capacity == gold_capacity);
-
-    auto ref                = set.ref(cuco::experimental::insert);
-    auto const ref_capacity = ref.capacity();
-    REQUIRE(ref_capacity == gold_capacity);
-  }
-
-  SECTION("Static extent must be evaluated at compile time.")
-  {
-    auto constexpr gold_capacity = 412;  // 103 x 2 x 2
-
-    using extent_type = cuco::experimental::extent<std::size_t, num_keys>;
-    using probe       = cuco::experimental::linear_probing<2, cuco::default_hash_function<Key>>;
-    auto set          = cuco::experimental::
-      static_set<Key, extent_type, cuda::thread_scope_device, Equal, probe, AllocatorT, StorageT>{
-        extent_type{}, cuco::empty_key<Key>{-1}};
-
-    REQUIRE(set.capacity() == gold_capacity);
-
-    auto const capacity = set.capacity();
-    STATIC_REQUIRE(capacity == gold_capacity);
-
-    auto ref                = set.ref(cuco::experimental::insert);
-    auto const ref_capacity = ref.capacity();
-    STATIC_REQUIRE(ref_capacity == gold_capacity);
-  }
-
-  SECTION("Dynamic extent is evaluated at run time.")
-  {
-    auto constexpr gold_capacity = 412;  // 103 x 2 x 2
-
-    using probe = cuco::experimental::linear_probing<2, cuco::default_hash_function<Key>>;
-    auto set    = cuco::experimental::static_set<Key,
-                                              cuco::experimental::extent<std::size_t>,
-                                              cuda::thread_scope_device,
-                                              Equal,
-                                              probe,
-                                              AllocatorT,
-                                              StorageT>{num_keys, cuco::empty_key<Key>{-1}};
-
-    auto const capacity = set.capacity();
-    REQUIRE(capacity == gold_capacity);
-
-    auto ref                = set.ref(cuco::experimental::insert);
-    auto const ref_capacity = ref.capacity();
-    REQUIRE(ref_capacity == gold_capacity);
-  }
-}
diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu
index 2623a8ae5..d44e20368 100644
--- a/tests/utility/extent_test.cu
+++ b/tests/utility/extent_test.cu
@@ -43,14 +43,14 @@ TEMPLATE_TEST_CASE_SIG(
   SECTION("Compute static valid extent at compile time.")
   {
     auto constexpr size = cuco::experimental::extent<SizeType, num>{};
-    auto constexpr res  = cuco::experimental::make_valid_extent<cg_size, window_size>(size);
-    STATIC_REQUIRE(gold_reference == res);
+    auto constexpr res  = cuco::experimental::make_window_extent<cg_size, window_size>(size);
+    STATIC_REQUIRE(gold_reference == res.value());
   }
 
   SECTION("Compute dynamic valid extent at run time.")
   {
     auto const size = cuco::experimental::extent<SizeType>{num};
-    auto const res  = cuco::experimental::make_valid_extent<cg_size, window_size>(size);
-    REQUIRE(gold_reference == res);
+    auto const res  = cuco::experimental::make_window_extent<cg_size, window_size>(size);
+    REQUIRE(gold_reference == res.value());
   }
 }