From 0f84b66650e965c30ac75deb2f8117411fd58f4c Mon Sep 17 00:00:00 2001 From: William Fan Date: Wed, 16 Sep 2026 09:27:04 -0400 Subject: [PATCH 1/6] Remove fast_int from extent and add unit tests --- include/cuco/detail/extent/extent.inl | 35 ++++++++++++++++++++++++--- tests/utility/extent_test.cu | 11 +++++++++ 2 files changed, 42 insertions(+), 4 deletions(-) diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl index 63faeaab2..e1a48cc56 100644 --- a/include/cuco/detail/extent/extent.inl +++ b/include/cuco/detail/extent/extent.inl @@ -110,12 +110,17 @@ struct valid_extent { }; template -struct valid_extent : cuco::utility::fast_int { - using value_type = - typename cuco::utility::fast_int::fast_int::value_type; ///< Extent value type +struct valid_extent { + using value_type = SizeType; ///< Extent value type + + __host__ __device__ constexpr value_type value() const noexcept { return value_; } + __host__ __device__ explicit constexpr operator value_type() const noexcept { return value(); } private: - using cuco::utility::fast_int::fast_int; + __host__ __device__ explicit constexpr valid_extent() noexcept : value_{} {} + __host__ __device__ explicit constexpr valid_extent(SizeType value) noexcept : value_{value} {} + + SizeType value_; // Friend declarations for all make_valid_extent overloads template @@ -135,6 +140,28 @@ struct valid_extent : cuco::utility::fast_int friend auto constexpr make_valid_extent(extent ext); + + // Operator overloads + template + friend __host__ __device__ constexpr value_type operator-(valid_extent const& lhs, + Rhs rhs) noexcept + { + return lhs.value() - rhs; + } + + template + friend __host__ __device__ constexpr value_type operator/(valid_extent const& lhs, + Rhs rhs) noexcept + { + return lhs.value() / rhs; + } + + template + friend __host__ __device__ constexpr value_type operator%(Lhs lhs, + valid_extent const& rhs) noexcept + { + return lhs % rhs.value(); + } }; // Primary implementation for fixed CGSize and BucketSize diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index 29edf1c3d..ab48c7552 100644 --- a/tests/utility/extent_test.cu +++ b/tests/utility/extent_test.cu @@ -54,6 +54,17 @@ TEMPLATE_TEST_CASE_SIG( REQUIRE(gold_reference == res.value()); } + SECTION("Dynamic valid extent supports arithmetic operators") + { + auto const size = cuco::extent{num}; + auto const res = cuco::make_valid_extent(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>; From c22feb8aa4cb3d5bd2e5e5ec663e2d9b15179282 Mon Sep 17 00:00:00 2001 From: William Fan Date: Wed, 16 Sep 2026 09:28:03 -0400 Subject: [PATCH 2/6] Reformatted --- tests/utility/extent_test.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index ab48c7552..924a40dae 100644 --- a/tests/utility/extent_test.cu +++ b/tests/utility/extent_test.cu @@ -64,7 +64,6 @@ TEMPLATE_TEST_CASE_SIG( REQUIRE((SizeType{5000} % res) == SizeType{5000} % gold_reference); } - SECTION("Invalid desired load factor throws exception") { using probing_scheme_type = cuco::linear_probing>; From 78e8703d224920cea6b7ca6443a8406445b4f494 Mon Sep 17 00:00:00 2001 From: William Fan Date: Wed, 16 Sep 2026 22:11:58 -0400 Subject: [PATCH 3/6] Restructure valid_extent operator overloads and remove stale imports and todo --- include/cuco/detail/extent/extent.inl | 63 ++++++------------- .../probing_scheme/probing_scheme_impl.inl | 2 +- tests/utility/extent_test.cu | 10 +++ 3 files changed, 31 insertions(+), 44 deletions(-) diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl index e1a48cc56..498c51835 100644 --- a/include/cuco/detail/extent/extent.inl +++ b/include/cuco/detail/extent/extent.inl @@ -10,7 +10,6 @@ #include #include #include -#include #include @@ -85,28 +84,6 @@ struct valid_extent { typename SizeType_, std::size_t N_> friend auto constexpr make_valid_extent(extent ext); - - // Operator overloads - template - friend __host__ __device__ constexpr value_type operator-(valid_extent const& lhs, - Rhs rhs) noexcept - { - return lhs.value() - rhs; - } - - template - friend __host__ __device__ constexpr value_type operator/(valid_extent const& lhs, - Rhs rhs) noexcept - { - return lhs.value() / rhs; - } - - template - friend __host__ __device__ constexpr value_type operator%(Lhs lhs, - valid_extent const& rhs) noexcept - { - return lhs % rhs.value(); - } }; template @@ -140,29 +117,29 @@ struct valid_extent { typename SizeType_, std::size_t N_> friend auto constexpr make_valid_extent(extent ext); +}; - // Operator overloads - template - friend __host__ __device__ constexpr value_type operator-(valid_extent const& lhs, - Rhs rhs) noexcept - { - return lhs.value() - rhs; - } +// Operator overloads +template +__host__ __device__ constexpr typename valid_extent::value_type operator-( + valid_extent const& lhs, Rhs rhs) noexcept +{ + return lhs.value() - rhs; +} - template - friend __host__ __device__ constexpr value_type operator/(valid_extent const& lhs, - Rhs rhs) noexcept - { - return lhs.value() / rhs; - } +template +__host__ __device__ constexpr typename valid_extent::value_type operator/( + valid_extent const& lhs, Rhs rhs) noexcept +{ + return lhs.value() / rhs; +} - template - friend __host__ __device__ constexpr value_type operator%(Lhs lhs, - valid_extent const& rhs) noexcept - { - return lhs % rhs.value(); - } -}; +template +__host__ __device__ constexpr typename valid_extent::value_type operator%( + Lhs lhs, valid_extent const& rhs) noexcept +{ + return lhs % rhs.value(); +} // Primary implementation for fixed CGSize and BucketSize template diff --git a/include/cuco/detail/probing_scheme/probing_scheme_impl.inl b/include/cuco/detail/probing_scheme/probing_scheme_impl.inl index 4ff6951bb..1e8757426 100644 --- a/include/cuco/detail/probing_scheme/probing_scheme_impl.inl +++ b/include/cuco/detail/probing_scheme/probing_scheme_impl.inl @@ -193,7 +193,7 @@ __host__ __device__ constexpr auto double_hashing::make_it hash2_(probe_key), static_cast(num_groups - 1)) + 1) * stride), - upper_bound}; // TODO use fast_int operator + upper_bound}; } template diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index 924a40dae..ba1058d20 100644 --- a/tests/utility/extent_test.cu +++ b/tests/utility/extent_test.cu @@ -54,6 +54,16 @@ TEMPLATE_TEST_CASE_SIG( REQUIRE(gold_reference == res.value()); } + SECTION("Static valid extent supports arithmetic operators") + { + auto constexpr size = cuco::extent{}; + auto constexpr res = cuco::make_valid_extent(size); + + REQUIRE((res - SizeType{10}) == gold_reference - 10); + REQUIRE((res / SizeType{2}) == gold_reference / 2); + REQUIRE((SizeType{5000} % res) == SizeType{5000} % gold_reference); + } + SECTION("Dynamic valid extent supports arithmetic operators") { auto const size = cuco::extent{num}; From bd6d09ca2c1656fc5f8a58b1bc30b9f64e55aadb Mon Sep 17 00:00:00 2001 From: William Fan Date: Wed, 16 Sep 2026 22:25:14 -0400 Subject: [PATCH 4/6] Add unit tests --- tests/utility/extent_test.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index ba1058d20..ee8a11646 100644 --- a/tests/utility/extent_test.cu +++ b/tests/utility/extent_test.cu @@ -59,9 +59,9 @@ TEMPLATE_TEST_CASE_SIG( auto constexpr size = cuco::extent{}; auto constexpr res = cuco::make_valid_extent(size); - REQUIRE((res - SizeType{10}) == gold_reference - 10); - REQUIRE((res / SizeType{2}) == gold_reference / 2); - REQUIRE((SizeType{5000} % res) == SizeType{5000} % gold_reference); + 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") From 72cf82a57dc933eadf8151c1211570cd49e4286a Mon Sep 17 00:00:00 2001 From: William Fan Date: Fri, 18 Sep 2026 22:45:45 -0400 Subject: [PATCH 5/6] Make valid_extent inherit from extent --- include/cuco/detail/extent/extent.inl | 50 +++++++-------------------- 1 file changed, 12 insertions(+), 38 deletions(-) diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl index 498c51835..8df8f9e12 100644 --- a/include/cuco/detail/extent/extent.inl +++ b/include/cuco/detail/extent/extent.inl @@ -56,48 +56,22 @@ constexpr std::uint64_t normalize_extent(SizeType size) } // namespace detail template -struct valid_extent { - using value_type = SizeType; ///< Extent value type +struct valid_extent : public extent { + using base_type = extent; - __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 - friend auto constexpr make_valid_extent(extent ext); - - template - friend auto constexpr make_valid_extent(extent ext); - - template