From 9c935339bb1a7ac773ebd7a5223f80bc8cf16ca4 Mon Sep 17 00:00:00 2001 From: Daniel Juenger <2955913+sleeepyjack@users.noreply.github.com> Date: Mon, 20 Nov 2023 22:40:51 +0000 Subject: [PATCH] Fix merge error --- .../open_addressing_ref_impl.cuh | 55 ++++++++++--------- 1 file changed, 28 insertions(+), 27 deletions(-) diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh index 3640add3a..3f013210f 100644 --- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh @@ -1133,35 +1133,36 @@ class open_addressing_ref_impl { } else { return cas_dependent_write(address, expected, desired); } + } - /** - * @brief Waits until the slot payload has been updated - * - * @note The function will return once the slot payload is no longer equal to the sentinel - * value. - * - * @tparam T Map slot type - * - * @param slot The target slot to check payload with - * @param sentinel The slot sentinel value - */ - template - __device__ void wait_for_payload(T & slot, T const& sentinel) const noexcept - { - auto ref = cuda::atomic_ref{slot}; - T current; - // TODO exponential backoff strategy - do { - current = ref.load(cuda::std::memory_order_relaxed); - } while (cuco::detail::bitwise_compare(current, sentinel)); - } + /** + * @brief Waits until the slot payload has been updated + * + * @note The function will return once the slot payload is no longer equal to the sentinel + * value. + * + * @tparam T Map slot type + * + * @param slot The target slot to check payload with + * @param sentinel The slot sentinel value + */ + template + __device__ void wait_for_payload(T& slot, T const& sentinel) const noexcept + { + auto ref = cuda::atomic_ref{slot}; + T current; + // TODO exponential backoff strategy + do { + current = ref.load(cuda::std::memory_order_relaxed); + } while (cuco::detail::bitwise_compare(current, sentinel)); + } - // TODO: Clean up the sentinel handling since it's duplicated in ref and equal wrapper - value_type empty_slot_sentinel_; ///< Sentinel value indicating an empty slot - detail::equal_wrapper predicate_; ///< Key equality binary callable - probing_scheme_type probing_scheme_; ///< Probing scheme - storage_ref_type storage_ref_; ///< Slot storage ref - }; + // TODO: Clean up the sentinel handling since it's duplicated in ref and equal wrapper + value_type empty_slot_sentinel_; ///< Sentinel value indicating an empty slot + detail::equal_wrapper predicate_; ///< Key equality binary callable + probing_scheme_type probing_scheme_; ///< Probing scheme + storage_ref_type storage_ref_; ///< Slot storage ref +}; } // namespace detail } // namespace experimental