Clean up cuco::pair (#319)

Contributes to #110 Depends on #314 This PR: - deprecates `cuco::pair_type` alias - fixes issues with `cuco::make_pair` - separates `pair` declarations and implementation details
NVIDIA · Jun 27, 2023 · 88ff1e4 · 88ff1e4
1 parent c0cce86
commit 88ff1e4
Show file tree

Hide file tree

Showing 40 changed files with 567 additions and 224 deletions.
diff --git a/benchmarks/hash_table/dynamic_map/contains_bench.cu b/benchmarks/hash_table/dynamic_map/contains_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_contains(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const initial_size  = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE);

diff --git a/benchmarks/hash_table/dynamic_map/erase_bench.cu b/benchmarks/hash_table/dynamic_map/erase_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_erase(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const initial_size  = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE);

diff --git a/benchmarks/hash_table/dynamic_map/find_bench.cu b/benchmarks/hash_table/dynamic_map/find_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_find(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const initial_size  = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE);

diff --git a/benchmarks/hash_table/dynamic_map/insert_bench.cu b/benchmarks/hash_table/dynamic_map/insert_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> dynamic_map_insert(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys     = state.get_int64_or_default("NumInputs", defaults::N);
   auto const initial_size = state.get_int64_or_default("InitSize", defaults::INITIAL_SIZE);

diff --git a/benchmarks/hash_table/static_map/contains_bench.cu b/benchmarks/hash_table/static_map/contains_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_contains(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);

diff --git a/benchmarks/hash_table/static_map/erase_bench.cu b/benchmarks/hash_table/static_map/erase_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_erase(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);

diff --git a/benchmarks/hash_table/static_map/find_bench.cu b/benchmarks/hash_table/static_map/find_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_find(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);

diff --git a/benchmarks/hash_table/static_map/insert_bench.cu b/benchmarks/hash_table/static_map/insert_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_map_insert(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys  = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);

diff --git a/benchmarks/hash_table/static_multimap/count_bench.cu b/benchmarks/hash_table/static_multimap/count_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_count(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);

diff --git a/benchmarks/hash_table/static_multimap/insert_bench.cu b/benchmarks/hash_table/static_multimap/insert_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_insert(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys  = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);

diff --git a/benchmarks/hash_table/static_multimap/query_bench.cu b/benchmarks/hash_table/static_multimap/query_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_query(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);
@@ -99,4 +99,4 @@ NVBENCH_BENCH_TYPES(static_multimap_query,
   .set_name("static_multimap_query_uniform_multiplicity")
   .set_type_axes_names({"Key", "Value", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
-  .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE);
+  .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE);
diff --git a/benchmarks/hash_table/static_multimap/retrieve_bench.cu b/benchmarks/hash_table/static_multimap/retrieve_bench.cu
@@ -35,7 +35,7 @@ template <typename Key, typename Value, typename Dist>
 std::enable_if_t<(sizeof(Key) == sizeof(Value)), void> static_multimap_retrieve(
   nvbench::state& state, nvbench::type_list<Key, Value, Dist>)
 {
-  using pair_type = cuco::pair_type<Key, Value>;
+  using pair_type = cuco::pair<Key, Value>;
 
   auto const num_keys      = state.get_int64_or_default("NumInputs", defaults::N);
   auto const occupancy     = state.get_float64_or_default("Occupancy", defaults::OCCUPANCY);
@@ -98,4 +98,4 @@ NVBENCH_BENCH_TYPES(static_multimap_retrieve,
   .set_name("static_multimap_retrieve_uniform_multiplicity")
   .set_type_axes_names({"Key", "Value", "Distribution"})
   .set_max_noise(defaults::MAX_NOISE)
-  .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE);
+  .add_int64_axis("Multiplicity", defaults::MULTIPLICITY_RANGE);
diff --git a/examples/static_multimap/host_bulk_example.cu b/examples/static_multimap/host_bulk_example.cu
@@ -60,7 +60,7 @@ int main(void)
   // The `_outer` suffix indicates that the occurrence of a non-match is 1.
   auto const output_size = map.count_outer(keys_to_find.begin(), keys_to_find.end());
 
-  thrust::device_vector<cuco::pair_type<key_type, value_type>> d_results(output_size);
+  thrust::device_vector<cuco::pair<key_type, value_type>> d_results(output_size);
 
   // Finds all keys {0, 1, 2, ...} and stores associated key/value pairs into `d_results`
   // If a key `keys_to_find[i]` doesn't exist, `d_results[i].second == empty_value_sentinel`

diff --git a/include/cuco/detail/dynamic_map.inl b/include/cuco/detail/dynamic_map.inl
@@ -146,7 +146,7 @@ void dynamic_map<Key, Value, Scope, Allocator>::insert(
       auto const n         = std::min(capacity_remaining, num_to_insert);
       auto const grid_size = (tile_size * n + stride * block_size - 1) / (stride * block_size);
 
-      detail::insert<block_size, tile_size, cuco::pair_type<key_type, mapped_type>>
+      detail::insert<block_size, tile_size, cuco::pair<key_type, mapped_type>>
         <<<grid_size, block_size, 0, stream>>>(first,
                                                first + n,
                                                submap_views_.data().get(),

diff --git a/include/cuco/detail/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing_ref_impl.cuh
@@ -17,7 +17,7 @@
 #pragma once
 
 #include <cuco/detail/equal_wrapper.cuh>
-#include <cuco/detail/pair.cuh>
+#include <cuco/pair.cuh>
 
 #include <thrust/distance.h>
 #include <thrust/pair.h>

diff --git a/include/cuco/detail/pair.inl b/include/cuco/detail/pair.inl
@@ -0,0 +1,51 @@
+/*
+ * Copyright (c) 2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <type_traits>
+#include <utility>
+
+namespace cuco {
+
+template <typename First, typename Second>
+__host__ __device__ constexpr pair<First, Second>::pair(First const& f, Second const& s)
+  : first{f}, second{s}
+{
+}
+
+template <typename First, typename Second>
+template <typename F, typename S>
+__host__ __device__ constexpr pair<First, Second>::pair(pair<F, S> const& p)
+  : first{p.first}, second{p.second}
+{
+}
+
+template <typename F, typename S>
+__host__ __device__ constexpr pair<std::decay_t<F>, std::decay_t<S>> make_pair(F&& f,
+                                                                               S&& s) noexcept
+{
+  return pair<std::decay_t<F>, std::decay_t<S>>(std::forward<F>(f), std::forward<S>(s));
+}
+
+template <class T1, class T2, class U1, class U2>
+__host__ __device__ constexpr bool operator==(cuco::pair<T1, T2> const& lhs,
+                                              cuco::pair<U1, U2> const& rhs) noexcept
+{
+  return lhs.first == rhs.first and lhs.second == rhs.second;
+}
+
+}  // namespace cuco
diff --git a/include/cuco/detail/probe_sequence_impl.cuh b/include/cuco/detail/probe_sequence_impl.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,7 +16,8 @@
 
 #pragma once
 
-#include <cuco/detail/pair.cuh>
+#include <cuco/detail/utils.cuh>
+#include <cuco/pair.cuh>
 
 #include <cuda/std/atomic>
 
@@ -71,13 +72,13 @@ template <typename Key,
           uint32_t CGSize>
 class probe_sequence_impl_base {
  protected:
-  using value_type         = cuco::pair_type<Key, Value>;       ///< Type of key/value pairs
+  using value_type         = cuco::pair<Key, Value>;            ///< Type of key/value pairs
   using key_type           = Key;                               ///< Key type
   using mapped_type        = Value;                             ///< Type of mapped values
   using atomic_key_type    = cuda::atomic<key_type, Scope>;     ///< Type of atomic keys
   using atomic_mapped_type = cuda::atomic<mapped_type, Scope>;  ///< Type of atomic mapped values
   /// Pair type of atomic key and atomic mapped value
-  using pair_atomic_type = cuco::pair_type<atomic_key_type, atomic_mapped_type>;
+  using pair_atomic_type = cuco::pair<atomic_key_type, atomic_mapped_type>;
   /// Type of the forward iterator to `pair_atomic_type`
   using iterator = pair_atomic_type*;
   /// Type of the forward iterator to `const pair_atomic_type`

diff --git a/include/cuco/detail/static_multimap/kernels.cuh b/include/cuco/detail/static_multimap/kernels.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -15,7 +15,7 @@
  */
 #pragma once
 
-#include <cuco/detail/pair.cuh>
+#include <cuco/pair.cuh>
 
 #include <thrust/type_traits/is_contiguous_iterator.h>
 

diff --git a/include/cuco/detail/traits.hpp b/include/cuco/detail/traits.hpp
@@ -0,0 +1,59 @@
+/*
+ * Copyright (c) 2023, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ */
+
+#pragma once
+
+#include <thrust/device_reference.h>
+#include <thrust/tuple.h>
+
+#include <cuda/std/tuple>
+#include <cuda/std/type_traits>
+
+namespace cuco::detail {
+
+template <typename T, typename = void>
+struct is_std_pair_like : cuda::std::false_type {
+};
+
+template <typename T>
+struct is_std_pair_like<T,
+                        cuda::std::void_t<decltype(cuda::std::get<0>(cuda::std::declval<T>())),
+                                          decltype(cuda::std::get<1>(cuda::std::declval<T>()))>>
+  : cuda::std::conditional_t<cuda::std::tuple_size<T>::value == 2,
+                             cuda::std::true_type,
+                             cuda::std::false_type> {
+};
+
+template <typename T, typename = void>
+struct is_thrust_pair_like_impl : cuda::std::false_type {
+};
+
+template <typename T>
+struct is_thrust_pair_like_impl<
+  T,
+  cuda::std::void_t<decltype(thrust::get<0>(cuda::std::declval<T>())),
+                    decltype(thrust::get<1>(cuda::std::declval<T>()))>>
+  : cuda::std::conditional_t<thrust::tuple_size<T>::value == 2,
+                             cuda::std::true_type,
+                             cuda::std::false_type> {
+};
+
+template <typename T>
+struct is_thrust_pair_like
+  : is_thrust_pair_like_impl<cuda::std::remove_reference_t<decltype(thrust::raw_reference_cast(
+      cuda::std::declval<T>()))>> {
+};
+
+}  // namespace cuco::detail