Skip to content

Commit

Permalink
Fix merge error
Browse files Browse the repository at this point in the history
  • Loading branch information
sleeepyjack committed Nov 20, 2023
1 parent 614f7bf commit 9c93533
Showing 1 changed file with 28 additions and 27 deletions.
55 changes: 28 additions & 27 deletions include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -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 <typename T>
__device__ void wait_for_payload(T & slot, T const& sentinel) const noexcept
{
auto ref = cuda::atomic_ref<T, Scope>{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 <typename T>
__device__ void wait_for_payload(T& slot, T const& sentinel) const noexcept
{
auto ref = cuda::atomic_ref<T, Scope>{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<key_type, key_equal> 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<key_type, key_equal> predicate_; ///< Key equality binary callable
probing_scheme_type probing_scheme_; ///< Probing scheme
storage_ref_type storage_ref_; ///< Slot storage ref
};

} // namespace detail
} // namespace experimental
Expand Down

0 comments on commit 9c93533

Please sign in to comment.