Austin Schuh | ae856ca | 2023-11-18 14:04:00 -0800 | [diff] [blame] | 1 | diff --git a/include/cuco/detail/pair/pair.inl b/include/cuco/detail/pair/pair.inl |
| 2 | index 3279a91..44be993 100644 |
| 3 | --- a/include/cuco/detail/pair/pair.inl |
| 4 | +++ b/include/cuco/detail/pair/pair.inl |
| 5 | @@ -51,7 +51,53 @@ __host__ __device__ constexpr bool operator==(cuco::pair<T1, T2> const& lhs, |
| 6 | } // namespace cuco |
| 7 | |
| 8 | namespace thrust { |
| 9 | -#include <cuco/detail/pair/tuple_helpers.inl> |
| 10 | +template <std::size_t I, typename T1, typename T2> |
| 11 | +__host__ __device__ constexpr auto get(cuco::pair<T1, T2>& p) -> |
| 12 | + typename tuple_element<I, cuco::pair<T1, T2>>::type& |
| 13 | +{ |
| 14 | + static_assert(I < 2); |
| 15 | + if constexpr (I == 0) { |
| 16 | + return p.first; |
| 17 | + } else { |
| 18 | + return p.second; |
| 19 | + } |
| 20 | +} |
| 21 | + |
| 22 | +template <std::size_t I, typename T1, typename T2> |
| 23 | +__host__ __device__ constexpr auto get(cuco::pair<T1, T2>&& p) -> |
| 24 | + typename tuple_element<I, cuco::pair<T1, T2>>::type&& |
| 25 | +{ |
| 26 | + static_assert(I < 2); |
| 27 | + if constexpr (I == 0) { |
| 28 | + return std::move(p.first); |
| 29 | + } else { |
| 30 | + return std::move(p.second); |
| 31 | + } |
| 32 | +} |
| 33 | + |
| 34 | +template <std::size_t I, typename T1, typename T2> |
| 35 | +__host__ __device__ constexpr auto get(cuco::pair<T1, T2> const& p) -> |
| 36 | + typename tuple_element<I, cuco::pair<T1, T2>>::type const& |
| 37 | +{ |
| 38 | + static_assert(I < 2); |
| 39 | + if constexpr (I == 0) { |
| 40 | + return p.first; |
| 41 | + } else { |
| 42 | + return p.second; |
| 43 | + } |
| 44 | +} |
| 45 | + |
| 46 | +template <std::size_t I, typename T1, typename T2> |
| 47 | +__host__ __device__ constexpr auto get(cuco::pair<T1, T2> const&& p) -> |
| 48 | + typename tuple_element<I, cuco::pair<T1, T2>>::type const&& |
| 49 | +{ |
| 50 | + static_assert(I < 2); |
| 51 | + if constexpr (I == 0) { |
| 52 | + return std::move(p.first); |
| 53 | + } else { |
| 54 | + return std::move(p.second); |
| 55 | + } |
| 56 | +} |
| 57 | } // namespace thrust |
| 58 | |
| 59 | namespace cuda::std { |
| 60 | diff --git a/include/cuco/detail/static_map.inl b/include/cuco/detail/static_map.inl |
| 61 | index 48799a7..336f6de 100644 |
| 62 | --- a/include/cuco/detail/static_map.inl |
| 63 | +++ b/include/cuco/detail/static_map.inl |
| 64 | @@ -443,7 +443,7 @@ __device__ |
| 65 | "insert_and_find is not supported for unpackable data on pre-Volta GPUs."); |
| 66 | #endif |
| 67 | |
| 68 | - auto current_slot{initial_slot(insert_pair.first, hash)}; |
| 69 | + auto current_slot{this->initial_slot(insert_pair.first, hash)}; |
| 70 | |
| 71 | while (true) { |
| 72 | key_type const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed); |
| 73 | @@ -514,7 +514,7 @@ __device__ |
| 74 | |
| 75 | // if we couldn't insert the key, but it wasn't a duplicate, then there must |
| 76 | // have been some other key there, so we keep looking for a slot |
| 77 | - current_slot = next_slot(current_slot); |
| 78 | + current_slot = this->next_slot(current_slot); |
| 79 | } |
| 80 | } |
| 81 | |
| 82 | diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl |
| 83 | index f27f21e..e90948f 100644 |
| 84 | --- a/include/cuco/detail/static_map/static_map_ref.inl |
| 85 | +++ b/include/cuco/detail/static_map/static_map_ref.inl |
| 86 | @@ -141,21 +141,6 @@ static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...> |
| 87 | return impl_.empty_value_sentinel(); |
| 88 | } |
| 89 | |
| 90 | -template <typename Key, |
| 91 | - typename T, |
| 92 | - cuda::thread_scope Scope, |
| 93 | - typename KeyEqual, |
| 94 | - typename ProbingScheme, |
| 95 | - typename StorageRef, |
| 96 | - typename... Operators> |
| 97 | -template <typename... NewOperators> |
| 98 | -auto static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::with( |
| 99 | - NewOperators...) && noexcept |
| 100 | -{ |
| 101 | - return static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, NewOperators...>( |
| 102 | - std::move(*this)); |
| 103 | -} |
| 104 | - |
| 105 | namespace detail { |
| 106 | |
| 107 | template <typename Key, |
| 108 | diff --git a/include/cuco/pair.cuh b/include/cuco/pair.cuh |
| 109 | index d28cae5..1caaa24 100644 |
| 110 | --- a/include/cuco/pair.cuh |
| 111 | +++ b/include/cuco/pair.cuh |
| 112 | @@ -87,8 +87,7 @@ struct alignas(detail::pair_alignment<First, Second>()) pair { |
| 113 | */ |
| 114 | template <typename T, std::enable_if_t<detail::is_std_pair_like<T>::value>* = nullptr> |
| 115 | __host__ __device__ constexpr pair(T const& p) |
| 116 | - : pair{cuda::std::get<0>(thrust::raw_reference_cast(p)), |
| 117 | - cuda::std::get<1>(thrust::raw_reference_cast(p))} |
| 118 | + : pair{std::get<0>(thrust::raw_reference_cast(p)), std::get<1>(thrust::raw_reference_cast(p))} |
| 119 | { |
| 120 | } |
| 121 | |
| 122 | diff --git a/include/cuco/static_map_ref.cuh b/include/cuco/static_map_ref.cuh |
| 123 | index 88e40f8..7cf1d74 100644 |
| 124 | --- a/include/cuco/static_map_ref.cuh |
| 125 | +++ b/include/cuco/static_map_ref.cuh |
| 126 | @@ -174,7 +174,11 @@ class static_map_ref |
| 127 | * @return `*this` with `NewOperators...` |
| 128 | */ |
| 129 | template <typename... NewOperators> |
| 130 | - [[nodiscard]] __host__ __device__ auto with(NewOperators... ops) && noexcept; |
| 131 | + [[nodiscard]] __host__ __device__ auto with(NewOperators... ops) && noexcept |
| 132 | + { |
| 133 | + return static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, NewOperators...>( |
| 134 | + std::move(*this)); |
| 135 | + } |
| 136 | |
| 137 | private: |
| 138 | impl_type impl_; ///< Static map ref implementation |