blob: 488eb9eb03cf7a5f0f84052a30c18192775741c3 [file] [log] [blame]
Austin Schuhae856ca2023-11-18 14:04:00 -08001diff --git a/include/cuco/detail/pair/pair.inl b/include/cuco/detail/pair/pair.inl
2index 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 {
60diff --git a/include/cuco/detail/static_map.inl b/include/cuco/detail/static_map.inl
61index 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
82diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl
83index 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,
108diff --git a/include/cuco/pair.cuh b/include/cuco/pair.cuh
109index 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
122diff --git a/include/cuco/static_map_ref.cuh b/include/cuco/static_map_ref.cuh
123index 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