Merge pull request #138 from PointKernel/add-pair-copy-ctor

Get rid of `std::move` when using `cuco::make_pair`
NVIDIA · Jan 28, 2022 · 6ec8b6d · 6ec8b6d
2 parents 2b83c76 + f382795
commit 6ec8b6d
Show file tree

Hide file tree

Showing 3 changed files with 22 additions and 23 deletions.
diff --git a/include/cuco/detail/pair.cuh b/include/cuco/detail/pair.cuh
@@ -132,7 +132,10 @@ union pair_converter {
   packed_type packed;
   pair_type pair;
 
-  __device__ pair_converter(pair_type _pair) : pair{_pair} {}
+  template <typename T>
+  __device__ pair_converter(T&& _pair) : pair{_pair}
+  {
+  }
 
   __device__ pair_converter(packed_type _packed) : packed{_packed} {}
 };
@@ -151,8 +154,7 @@ template <typename First, typename Second>
 struct alignas(detail::pair_alignment<First, Second>()) pair {
   using first_type  = First;
   using second_type = Second;
-  First first;
-  Second second;
+
   pair()            = default;
   ~pair()           = default;
   pair(pair const&) = default;
@@ -162,6 +164,11 @@ struct alignas(detail::pair_alignment<First, Second>()) pair {
 
   __host__ __device__ constexpr pair(First const& f, Second const& s) : first{f}, second{s} {}
 
+  template <typename F, typename S>
+  __host__ __device__ constexpr pair(pair<F, S> const& p) : first{p.first}, second{p.second}
+  {
+  }
+
   template <typename T, std::enable_if_t<detail::is_std_pair_like<T>::value>* = nullptr>
   __host__ __device__ constexpr pair(T const& t)
     : pair{std::get<0>(thrust::raw_reference_cast(t)), std::get<1>(thrust::raw_reference_cast(t))}
@@ -174,6 +181,9 @@ struct alignas(detail::pair_alignment<First, Second>()) pair {
            thrust::get<1>(thrust::raw_reference_cast(t))}
   {
   }
+
+  First first;
+  Second second;
 };
 
 template <typename K, typename V>

diff --git a/include/cuco/detail/static_map.inl b/include/cuco/detail/static_map.inl
@@ -168,7 +168,7 @@ static_map<Key, Value, Scope, Allocator>::device_mutable_view::packed_cas(
   auto expected_value = this->get_empty_value_sentinel();
 
   cuco::detail::pair_converter<value_type> expected_pair{
-    cuco::make_pair<Key, Value>(std::move(expected_key), std::move(expected_value))};
+    cuco::make_pair(expected_key, expected_value)};
   cuco::detail::pair_converter<value_type> new_pair{insert_pair};
 
   auto slot =

diff --git a/include/cuco/detail/static_multimap/device_view_impl.inl b/include/cuco/detail/static_multimap/device_view_impl.inl
@@ -236,7 +236,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_mutab
     auto expected_value = this->get_empty_value_sentinel();
 
     cuco::detail::pair_converter<value_type> expected_pair{
-      cuco::make_pair<Key, Value>(std::move(expected_key), std::move(expected_value))};
+      cuco::make_pair(expected_key, expected_value)};
     cuco::detail::pair_converter<value_type> new_pair{insert_pair};
 
     auto slot = reinterpret_cast<
@@ -927,26 +927,21 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
 
           if (first_equals) {
             auto const lane_offset = detail::count_least_significant_bits(first_exists, cg_lane_id);
-            Key key                = k;
-            output_buffer[output_idx + lane_offset] =
-              cuco::make_pair<Key, Value>(std::move(key), std::move(arr[0].second));
+            output_buffer[output_idx + lane_offset] = cuco::make_pair(k, arr[0].second);
           }
           if (second_equals) {
             auto const lane_offset =
               detail::count_least_significant_bits(second_exists, cg_lane_id);
-            Key key = k;
             output_buffer[output_idx + num_first_matches + lane_offset] =
-              cuco::make_pair<Key, Value>(std::move(key), std::move(arr[1].second));
+              cuco::make_pair(k, arr[1].second);
           }
         }
         if (probing_cg.any(first_slot_is_empty or second_slot_is_empty)) {
           running = false;
           if constexpr (is_outer) {
             if ((not found_match) && (cg_lane_id == 0)) {
               auto const output_idx     = atomicAdd(flushing_cg_counter, 1);
-              Key key                   = k;
-              output_buffer[output_idx] = cuco::make_pair<Key, Value>(
-                std::move(key), std::move(this->get_empty_value_sentinel()));
+              output_buffer[output_idx] = cuco::make_pair(k, this->get_empty_value_sentinel());
             }
           }
         }
@@ -1029,9 +1024,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
         if (equals) {
           // Each match computes its lane-level offset
           auto const lane_offset = detail::count_least_significant_bits(exists, lane_id);
-          Key key                = k;
-          output_buffer[output_idx + lane_offset] =
-            cuco::make_pair<Key, Value>(std::move(key), std::move(slot_contents.second));
+          output_buffer[output_idx + lane_offset] = cuco::make_pair(k, slot_contents.second);
         }
         if (0 == lane_id) { (*cg_counter) += num_matches; }
       }
@@ -1040,9 +1033,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
         if constexpr (is_outer) {
           if ((not found_match) && (lane_id == 0)) {
             output_idx                = (*cg_counter)++;
-            Key key                   = k;
-            output_buffer[output_idx] = cuco::make_pair<Key, Value>(
-              std::move(key), std::move(this->get_empty_value_sentinel()));
+            output_buffer[output_idx] = cuco::make_pair(k, this->get_empty_value_sentinel());
           }
         }
       }
@@ -1373,8 +1364,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
               auto const output_idx           = atomicAdd(flushing_cg_counter, 1);
               probe_output_buffer[output_idx] = pair;
               contained_output_buffer[output_idx] =
-                cuco::make_pair<Key, Value>(std::move(this->get_empty_key_sentinel()),
-                                            std::move(this->get_empty_value_sentinel()));
+                cuco::make_pair(this->get_empty_key_sentinel(), this->get_empty_value_sentinel());
             }
           }
         }
@@ -1484,8 +1474,7 @@ class static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::device_view_
             output_idx                      = (*cg_counter)++;
             probe_output_buffer[output_idx] = pair;
             contained_output_buffer[output_idx] =
-              cuco::make_pair<Key, Value>(std::move(this->get_empty_key_sentinel()),
-                                          std::move(this->get_empty_value_sentinel()));
+              cuco::make_pair(this->get_empty_key_sentinel(), this->get_empty_value_sentinel());
           }
         }
       }