Fix memmove optimization (#1937)

Copy the basic copy / move algorithms we use memmove when possible. However, we used the wrong constraint for that. We should always check for `is_trivially_copyable`
NVIDIA · Jul 4, 2024 · 7fba5aa · 7fba5aa
1 parent 166b97f
commit 7fba5aa
Show file tree

Hide file tree

Showing 8 changed files with 140 additions and 8 deletions.
diff --git a/libcudacxx/include/cuda/std/__algorithm/copy.h b/libcudacxx/include/cuda/std/__algorithm/copy.h
@@ -25,7 +25,7 @@
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_constant_evaluated.h>
 #include <cuda/std/__type_traits/is_same.h>
-#include <cuda/std/__type_traits/is_trivially_copy_assignable.h>
+#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/detail/libcxx/include/cstdlib>
 #include <cuda/std/detail/libcxx/include/cstring>
@@ -70,7 +70,7 @@ template <class _AlgPolicy,
           class _Tp,
           class _Up,
           __enable_if_t<_CCCL_TRAIT(is_same, __remove_const_t<_Tp>, _Up), int> = 0,
-          __enable_if_t<_CCCL_TRAIT(is_trivially_copy_assignable, _Up), int>   = 0>
+          __enable_if_t<_CCCL_TRAIT(is_trivially_copyable, _Up), int>          = 0>
 inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair<_Tp*, _Up*>
 __copy(_Tp* __first, _Tp* __last, _Up* __result)
 {

diff --git a/libcudacxx/include/cuda/std/__algorithm/copy_backward.h b/libcudacxx/include/cuda/std/__algorithm/copy_backward.h
@@ -24,7 +24,7 @@
 #include <cuda/std/__algorithm/unwrap_iter.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_same.h>
-#include <cuda/std/__type_traits/is_trivially_copy_assignable.h>
+#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/remove_const.h>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
@@ -43,7 +43,7 @@ __copy_backward(_BidirectionalIterator __first, _BidirectionalIterator __last, _
 template <class _Tp,
           class _Up,
           __enable_if_t<_CCCL_TRAIT(is_same, __remove_const_t<_Tp>, _Up), int> = 0,
-          __enable_if_t<_CCCL_TRAIT(is_trivially_copy_assignable, _Up), int>   = 0>
+          __enable_if_t<_CCCL_TRAIT(is_trivially_copyable, _Up), int>          = 0>
 inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 _Up*
 __copy_backward(_Tp* __first, _Tp* __last, _Up* __result)
 {

diff --git a/libcudacxx/include/cuda/std/__algorithm/move.h b/libcudacxx/include/cuda/std/__algorithm/move.h
@@ -25,7 +25,7 @@
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_copy_constructible.h>
 #include <cuda/std/__type_traits/is_same.h>
-#include <cuda/std/__type_traits/is_trivially_move_assignable.h>
+#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/__utility/pair.h>
 
@@ -47,7 +47,7 @@ template <class _AlgPolicy,
           class _Tp,
           class _Up,
           __enable_if_t<_CCCL_TRAIT(is_same, __remove_const_t<_Tp>, _Up), int> = 0,
-          __enable_if_t<_CCCL_TRAIT(is_trivially_move_assignable, _Up), int>   = 0>
+          __enable_if_t<_CCCL_TRAIT(is_trivially_copyable, _Up), int>          = 0>
 inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair<_Tp*, _Up*>
 __move(_Tp* __first, _Tp* __last, _Up* __result)
 {

diff --git a/libcudacxx/include/cuda/std/__algorithm/move_backward.h b/libcudacxx/include/cuda/std/__algorithm/move_backward.h
@@ -24,7 +24,7 @@
 #include <cuda/std/__algorithm/unwrap_iter.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_same.h>
-#include <cuda/std/__type_traits/is_trivially_move_assignable.h>
+#include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/__utility/pair.h>
 
@@ -46,7 +46,7 @@ template <class _AlgPolicy,
           class _Tp,
           class _Up,
           __enable_if_t<_CCCL_TRAIT(is_same, __remove_const_t<_Tp>, _Up), int> = 0,
-          __enable_if_t<_CCCL_TRAIT(is_trivially_move_assignable, _Up), int>   = 0>
+          __enable_if_t<_CCCL_TRAIT(is_trivially_copyable, _Up), int>          = 0>
 inline _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair<_Tp*, _Up*>
 __move_backward(_Tp* __first, _Tp* __last, _Up* __result)
 {

diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy.pass.cpp
@@ -72,6 +72,38 @@ struct NonTrivialCopy
   }
 };
 
+struct NonTrivialDestructor
+{
+  int data = 0;
+
+  NonTrivialDestructor() = default;
+
+  NonTrivialDestructor(NonTrivialDestructor&&) noexcept                 = default;
+  NonTrivialDestructor(const NonTrivialDestructor&) noexcept            = default;
+  NonTrivialDestructor& operator=(NonTrivialDestructor&&) noexcept      = default;
+  NonTrivialDestructor& operator=(const NonTrivialDestructor&) noexcept = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 ~NonTrivialDestructor() noexcept {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor(const int val) noexcept
+      : data(val)
+  {}
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor& operator=(const int val) noexcept
+  {
+    data = val;
+    return *this;
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool
+  operator==(const NonTrivialDestructor& lhs, const NonTrivialDestructor& rhs) noexcept
+  {
+    return lhs.data == rhs.data;
+  }
+  __host__ __device__ TEST_CONSTEXPR_CXX20 bool operator==(const int& other) const noexcept
+  {
+    return data == other;
+  }
+};
+
 template <class InIter, class OutIter>
 TEST_CONSTEXPR_CXX20 __host__ __device__ void test()
 {
@@ -149,6 +181,7 @@ TEST_CONSTEXPR_CXX20 __host__ __device__ bool test()
   test<const int*, int*>();
 
   test<const NonTrivialCopy*, NonTrivialCopy*>();
+  test<const NonTrivialDestructor*, NonTrivialDestructor*>();
 
   return true;
 }

diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_backward.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.copy/copy_backward.pass.cpp
@@ -72,6 +72,38 @@ struct NonTrivialCopy
   }
 };
 
+struct NonTrivialDestructor
+{
+  int data = 0;
+
+  NonTrivialDestructor() = default;
+
+  NonTrivialDestructor(NonTrivialDestructor&&) noexcept                 = default;
+  NonTrivialDestructor(const NonTrivialDestructor&) noexcept            = default;
+  NonTrivialDestructor& operator=(NonTrivialDestructor&&) noexcept      = default;
+  NonTrivialDestructor& operator=(const NonTrivialDestructor&) noexcept = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 ~NonTrivialDestructor() noexcept {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor(const int val) noexcept
+      : data(val)
+  {}
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor& operator=(const int val) noexcept
+  {
+    data = val;
+    return *this;
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool
+  operator==(const NonTrivialDestructor& lhs, const NonTrivialDestructor& rhs) noexcept
+  {
+    return lhs.data == rhs.data;
+  }
+  __host__ __device__ TEST_CONSTEXPR_CXX20 bool operator==(const int& other) const noexcept
+  {
+    return data == other;
+  }
+};
+
 template <class InIter, class OutIter>
 TEST_CONSTEXPR_CXX20 __host__ __device__ void test()
 {
@@ -138,6 +170,7 @@ TEST_CONSTEXPR_CXX20 __host__ __device__ bool test()
   test<const int*, int*>();
 
   test<const NonTrivialCopy*, NonTrivialCopy*>();
+  test<const NonTrivialDestructor*, NonTrivialDestructor*>();
 
   return true;
 }

diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move.pass.cpp
@@ -76,6 +76,38 @@ struct NonTrivialMove
   }
 };
 
+struct NonTrivialDestructor
+{
+  int data = 0;
+
+  NonTrivialDestructor() = default;
+
+  NonTrivialDestructor(NonTrivialDestructor&&) noexcept                 = default;
+  NonTrivialDestructor(const NonTrivialDestructor&) noexcept            = default;
+  NonTrivialDestructor& operator=(NonTrivialDestructor&&) noexcept      = default;
+  NonTrivialDestructor& operator=(const NonTrivialDestructor&) noexcept = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 ~NonTrivialDestructor() noexcept {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor(const int val) noexcept
+      : data(val)
+  {}
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor& operator=(const int val) noexcept
+  {
+    data = val;
+    return *this;
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool
+  operator==(const NonTrivialDestructor& lhs, const NonTrivialDestructor& rhs) noexcept
+  {
+    return lhs.data == rhs.data;
+  }
+  __host__ __device__ TEST_CONSTEXPR_CXX20 bool operator==(const int& other) const noexcept
+  {
+    return data == other;
+  }
+};
+
 template <class InIter, class OutIter>
 __host__ __device__ TEST_CONSTEXPR_CXX14 void test()
 {
@@ -175,6 +207,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test()
   test<int*, int*>();
 
   test<NonTrivialMove*, NonTrivialMove*>();
+  test<NonTrivialDestructor*, NonTrivialDestructor*>();
 
 #if defined(_LIBCUDACXX_HAS_MEMORY)
   test1<cpp17_input_iterator<cuda::std::unique_ptr<int>*>, cpp17_output_iterator<cuda::std::unique_ptr<int>*>>();

diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move_backward.pass.cpp b/libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.move/move_backward.pass.cpp
@@ -75,6 +75,38 @@ struct NonTrivialMove
   }
 };
 
+struct NonTrivialDestructor
+{
+  int data = 0;
+
+  NonTrivialDestructor() = default;
+
+  NonTrivialDestructor(NonTrivialDestructor&&) noexcept                 = default;
+  NonTrivialDestructor(const NonTrivialDestructor&) noexcept            = default;
+  NonTrivialDestructor& operator=(NonTrivialDestructor&&) noexcept      = default;
+  NonTrivialDestructor& operator=(const NonTrivialDestructor&) noexcept = default;
+  __host__ __device__ TEST_CONSTEXPR_CXX20 ~NonTrivialDestructor() noexcept {}
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor(const int val) noexcept
+      : data(val)
+  {}
+  __host__ __device__ TEST_CONSTEXPR_CXX20 NonTrivialDestructor& operator=(const int val) noexcept
+  {
+    data = val;
+    return *this;
+  }
+
+  __host__ __device__ TEST_CONSTEXPR_CXX20 friend bool
+  operator==(const NonTrivialDestructor& lhs, const NonTrivialDestructor& rhs) noexcept
+  {
+    return lhs.data == rhs.data;
+  }
+  __host__ __device__ TEST_CONSTEXPR_CXX20 bool operator==(const int& other) const noexcept
+  {
+    return data == other;
+  }
+};
+
 template <class InIter, class OutIter>
 __host__ __device__ TEST_CONSTEXPR_CXX14 void test()
 {
@@ -162,6 +194,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX14 bool test()
   test<int*, int*>();
 
   test<NonTrivialMove*, NonTrivialMove*>();
+  test<NonTrivialDestructor*, NonTrivialDestructor*>();
 
 #if defined(_LIBCUDACXX_HAS_MEMORY)
   test1<bidirectional_iterator<cuda::std::unique_ptr<int>*>, bidirectional_iterator<cuda::std::unique_ptr<int>*>>();