diff --git a/include/cuco/detail/extent/extent.inl b/include/cuco/detail/extent/extent.inl index a7cd83dcd..007257ea9 100644 --- a/include/cuco/detail/extent/extent.inl +++ b/include/cuco/detail/extent/extent.inl @@ -31,8 +31,7 @@ template 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 index 3090d026e..2458315d5 100644 --- 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(upper_bound_); return *this; } @@ -100,7 +100,7 @@ __host__ __device__ constexpr auto linear_probing::operator()( { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash_(probe_key)) % upper_bound, + cuco::detail::sanitize_hash(hash_(probe_key)) % static_cast(upper_bound), 1, // step size is 1 upper_bound}; } @@ -114,7 +114,8 @@ __host__ __device__ constexpr auto linear_probing::operator()( { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash_(probe_key) + g.thread_rank()) % upper_bound, + cuco::detail::sanitize_hash(hash_(probe_key) + g.thread_rank()) % + static_cast(upper_bound), cg_size, upper_bound}; } @@ -133,10 +134,10 @@ __host__ __device__ constexpr auto double_hashing::operato { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash1_(probe_key)) % upper_bound, + cuco::detail::sanitize_hash(hash1_(probe_key)) % static_cast(upper_bound), max(size_type{1}, cuco::detail::sanitize_hash(hash2_(probe_key)) % - upper_bound), // step size in range [1, prime - 1] + static_cast(upper_bound)), // step size in range [1, prime - 1] upper_bound}; } @@ -149,9 +150,10 @@ __host__ __device__ constexpr auto double_hashing::operato { using size_type = typename Extent::value_type; return detail::probing_iterator{ - cuco::detail::sanitize_hash(hash1_(probe_key) + g.thread_rank()) % upper_bound, + cuco::detail::sanitize_hash(hash1_(probe_key) + g.thread_rank()) % + static_cast(upper_bound), static_cast((cuco::detail::sanitize_hash(hash2_(probe_key)) % - (upper_bound.value() / cg_size - 1) + + (static_cast(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 index 6616e2c5c..a9329a499 100644 --- 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 \ No newline at end of file +} // namespace cuco::utility diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu index 69fa69fb0..389453372 100644 --- 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; using probe = std::conditional_t>, @@ -180,13 +182,13 @@ TEMPLATE_TEST_CASE_SIG( auto map = cuco::experimental::static_map, + extent_type, cuda::thread_scope_device, thrust::equal_to, probe, cuco::cuda_allocator, cuco::experimental::storage<2>>{ - num_keys, cuco::empty_key{-1}, cuco::empty_value{-1}}; + extent_type{}, cuco::empty_key{-1}, cuco::empty_value{-1}}; REQUIRE(map.capacity() == gold_capacity); diff --git a/tests/utility/extent_test.cu b/tests/utility/extent_test.cu index d44e20368..d069f4dbd 100644 --- 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{}; auto constexpr res = cuco::experimental::make_window_extent(size); - STATIC_REQUIRE(gold_reference == res.value()); + STATIC_REQUIRE(gold_reference == static_cast(res)); } SECTION("Compute dynamic valid extent at run time.") { auto const size = cuco::experimental::extent{num}; auto const res = cuco::experimental::make_window_extent(size); - REQUIRE(gold_reference == res.value()); + REQUIRE(gold_reference == static_cast(res)); } }