blob: 488eb9eb03cf7a5f0f84052a30c18192775741c3 [file] [log] [blame]
diff --git a/include/cuco/detail/pair/pair.inl b/include/cuco/detail/pair/pair.inl
index 3279a91..44be993 100644
--- a/include/cuco/detail/pair/pair.inl
+++ b/include/cuco/detail/pair/pair.inl
@@ -51,7 +51,53 @@ __host__ __device__ constexpr bool operator==(cuco::pair<T1, T2> const& lhs,
} // namespace cuco
namespace thrust {
-#include <cuco/detail/pair/tuple_helpers.inl>
+template <std::size_t I, typename T1, typename T2>
+__host__ __device__ constexpr auto get(cuco::pair<T1, T2>& p) ->
+ typename tuple_element<I, cuco::pair<T1, T2>>::type&
+{
+ static_assert(I < 2);
+ if constexpr (I == 0) {
+ return p.first;
+ } else {
+ return p.second;
+ }
+}
+
+template <std::size_t I, typename T1, typename T2>
+__host__ __device__ constexpr auto get(cuco::pair<T1, T2>&& p) ->
+ typename tuple_element<I, cuco::pair<T1, T2>>::type&&
+{
+ static_assert(I < 2);
+ if constexpr (I == 0) {
+ return std::move(p.first);
+ } else {
+ return std::move(p.second);
+ }
+}
+
+template <std::size_t I, typename T1, typename T2>
+__host__ __device__ constexpr auto get(cuco::pair<T1, T2> const& p) ->
+ typename tuple_element<I, cuco::pair<T1, T2>>::type const&
+{
+ static_assert(I < 2);
+ if constexpr (I == 0) {
+ return p.first;
+ } else {
+ return p.second;
+ }
+}
+
+template <std::size_t I, typename T1, typename T2>
+__host__ __device__ constexpr auto get(cuco::pair<T1, T2> const&& p) ->
+ typename tuple_element<I, cuco::pair<T1, T2>>::type const&&
+{
+ static_assert(I < 2);
+ if constexpr (I == 0) {
+ return std::move(p.first);
+ } else {
+ return std::move(p.second);
+ }
+}
} // namespace thrust
namespace cuda::std {
diff --git a/include/cuco/detail/static_map.inl b/include/cuco/detail/static_map.inl
index 48799a7..336f6de 100644
--- a/include/cuco/detail/static_map.inl
+++ b/include/cuco/detail/static_map.inl
@@ -443,7 +443,7 @@ __device__
"insert_and_find is not supported for unpackable data on pre-Volta GPUs.");
#endif
- auto current_slot{initial_slot(insert_pair.first, hash)};
+ auto current_slot{this->initial_slot(insert_pair.first, hash)};
while (true) {
key_type const existing_key = current_slot->first.load(cuda::std::memory_order_relaxed);
@@ -514,7 +514,7 @@ __device__
// if we couldn't insert the key, but it wasn't a duplicate, then there must
// have been some other key there, so we keep looking for a slot
- current_slot = next_slot(current_slot);
+ current_slot = this->next_slot(current_slot);
}
}
diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl
index f27f21e..e90948f 100644
--- a/include/cuco/detail/static_map/static_map_ref.inl
+++ b/include/cuco/detail/static_map/static_map_ref.inl
@@ -141,21 +141,6 @@ static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>
return impl_.empty_value_sentinel();
}
-template <typename Key,
- typename T,
- cuda::thread_scope Scope,
- typename KeyEqual,
- typename ProbingScheme,
- typename StorageRef,
- typename... Operators>
-template <typename... NewOperators>
-auto static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::with(
- NewOperators...) && noexcept
-{
- return static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, NewOperators...>(
- std::move(*this));
-}
-
namespace detail {
template <typename Key,
diff --git a/include/cuco/pair.cuh b/include/cuco/pair.cuh
index d28cae5..1caaa24 100644
--- a/include/cuco/pair.cuh
+++ b/include/cuco/pair.cuh
@@ -87,8 +87,7 @@ struct alignas(detail::pair_alignment<First, Second>()) pair {
*/
template <typename T, std::enable_if_t<detail::is_std_pair_like<T>::value>* = nullptr>
__host__ __device__ constexpr pair(T const& p)
- : pair{cuda::std::get<0>(thrust::raw_reference_cast(p)),
- cuda::std::get<1>(thrust::raw_reference_cast(p))}
+ : pair{std::get<0>(thrust::raw_reference_cast(p)), std::get<1>(thrust::raw_reference_cast(p))}
{
}
diff --git a/include/cuco/static_map_ref.cuh b/include/cuco/static_map_ref.cuh
index 88e40f8..7cf1d74 100644
--- a/include/cuco/static_map_ref.cuh
+++ b/include/cuco/static_map_ref.cuh
@@ -174,7 +174,11 @@ class static_map_ref
* @return `*this` with `NewOperators...`
*/
template <typename... NewOperators>
- [[nodiscard]] __host__ __device__ auto with(NewOperators... ops) && noexcept;
+ [[nodiscard]] __host__ __device__ auto with(NewOperators... ops) && noexcept
+ {
+ return static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, NewOperators...>(
+ std::move(*this));
+ }
private:
impl_type impl_; ///< Static map ref implementation