Ensure that we can use cuda::std::optional with types that are not …

…`__host__ __device__`
NVIDIA · Apr 24, 2024 · 7e0adb0 · 7e0adb0
1 parent 7333e0b
commit 7e0adb0
Show file tree

Hide file tree

Showing 39 changed files with 603 additions and 1 deletion.
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/optional b/libcudacxx/include/cuda/std/detail/libcxx/include/optional
@@ -268,6 +268,7 @@ struct __optional_destruct_base<_Tp, false>
   };
   bool __engaged_;
 
+  _CCCL_EXEC_CHECK_DISABLE
   _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 ~__optional_destruct_base()
   {
     if (__engaged_)
@@ -281,19 +282,22 @@ struct __optional_destruct_base<_Tp, false>
       , __engaged_(false)
   {}
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class... _Args>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __optional_destruct_base(in_place_t, _Args&&... __args)
       : __val_(_CUDA_VSTD::forward<_Args>(__args)...)
       , __engaged_(true)
   {}
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Fp, class... _Args>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr __optional_destruct_base(
     __optional_construct_from_invoke_tag, _Fp&& __f, _Args&&... __args)
       : __val_(_CUDA_VSTD::invoke(_CUDA_VSTD::forward<_Fp>(__f), _CUDA_VSTD::forward<_Args>(__args)...))
       , __engaged_(true)
   {}
 
+  _CCCL_EXEC_CHECK_DISABLE
   _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void reset() noexcept
   {
     if (__engaged_)
@@ -322,12 +326,14 @@ struct __optional_destruct_base<_Tp, true>
       , __engaged_(false)
   {}
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class... _Args>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr explicit __optional_destruct_base(in_place_t, _Args&&... __args)
       : __val_(_CUDA_VSTD::forward<_Args>(__args)...)
       , __engaged_(true)
   {}
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Fp, class... _Args>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr __optional_destruct_base(
     __optional_construct_from_invoke_tag, _Fp&& __f, _Args&&... __args)
@@ -373,6 +379,7 @@ struct __optional_storage_base : __optional_destruct_base<_Tp>
     return _CUDA_VSTD::move(this->__val_);
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class... _Args>
   _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 void __construct(_Args&&... __args)
   {
@@ -394,6 +401,7 @@ struct __optional_storage_base : __optional_destruct_base<_Tp>
     }
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _That>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr void __assign_from(_That&& __opt)
   {
@@ -695,6 +703,7 @@ public:
   constexpr optional& operator=(const optional&) = default;
   constexpr optional& operator=(optional&&)      = default;
 
+  _CCCL_EXEC_CHECK_DISABLE
   _LIBCUDACXX_TEMPLATE(class _Up = value_type)
   _LIBCUDACXX_REQUIRES(__opt_is_assignable_from_U<_Tp, _Up> _LIBCUDACXX_AND __opt_is_assignable<_Tp, _Up>)
   _LIBCUDACXX_INLINE_VISIBILITY constexpr optional& operator=(_Up&& __v)
@@ -744,6 +753,7 @@ public:
     return this->__get();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   _LIBCUDACXX_INLINE_VISIBILITY constexpr void swap(optional& __opt) noexcept(
     _LIBCUDACXX_TRAIT(is_nothrow_move_constructible, value_type) && _LIBCUDACXX_TRAIT(is_nothrow_swappable, value_type))
   {
@@ -852,6 +862,7 @@ public:
     return _CUDA_VSTD::move(this->__get());
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Up>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr value_type value_or(_Up&& __v) const&
   {
@@ -862,6 +873,7 @@ public:
     return this->has_value() ? this->__get() : static_cast<value_type>(_CUDA_VSTD::forward<_Up>(__v));
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Up>
   _LIBCUDACXX_INLINE_VISIBILITY constexpr value_type value_or(_Up&& __v) &&
   {
@@ -872,6 +884,7 @@ public:
     return this->has_value() ? _CUDA_VSTD::move(this->__get()) : static_cast<value_type>(_CUDA_VSTD::forward<_Up>(__v));
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto and_then(_Func&& __f) &
   {
@@ -885,6 +898,7 @@ public:
     return remove_cvref_t<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   and_then(_Func&& __f) const&
@@ -899,6 +913,7 @@ public:
     return remove_cvref_t<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   and_then(_Func&& __f) &&
@@ -913,6 +928,7 @@ public:
     return remove_cvref_t<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   and_then(_Func&& __f) const&&
@@ -927,6 +943,7 @@ public:
     return remove_cvref_t<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   transform(_Func&& __f) &
@@ -943,6 +960,7 @@ public:
     return optional<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   transform(_Func&& __f) const&
@@ -959,6 +977,7 @@ public:
     return optional<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   transform(_Func&& __f) &&
@@ -978,6 +997,7 @@ public:
     return optional<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   template <class _Func>
   _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_AVAILABILITY_THROW_BAD_OPTIONAL_ACCESS constexpr auto
   transform(_Func&& __f) const&&
@@ -997,6 +1017,7 @@ public:
     return optional<_Up>();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   _LIBCUDACXX_TEMPLATE(class _Func, class _Tp2 = _Tp)
   _LIBCUDACXX_REQUIRES(invocable<_Func> _LIBCUDACXX_AND _LIBCUDACXX_TRAIT(is_copy_constructible, _Tp2))
   _LIBCUDACXX_INLINE_VISIBILITY constexpr optional or_else(_Func&& __f) const&
@@ -1010,6 +1031,7 @@ public:
     return _CUDA_VSTD::forward<_Func>(__f)();
   }
 
+  _CCCL_EXEC_CHECK_DISABLE
   _LIBCUDACXX_TEMPLATE(class _Func, class _Tp2 = _Tp)
   _LIBCUDACXX_REQUIRES(invocable<_Func> _LIBCUDACXX_AND _LIBCUDACXX_TRAIT(is_move_constructible, _Tp2))
   _LIBCUDACXX_INLINE_VISIBILITY constexpr optional or_else(_Func&& __f) &&
@@ -1032,6 +1054,7 @@ _CCCL_HOST_DEVICE optional(_Tp) -> optional<_Tp>;
 #  endif
 
 // Comparisons between optionals
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() == declval<const _Up&>()), bool),
@@ -1049,6 +1072,7 @@ operator==(const optional<_Tp>& __x, const optional<_Up>& __y)
   return *__x == *__y;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() != declval<const _Up&>()), bool),
@@ -1066,6 +1090,7 @@ operator!=(const optional<_Tp>& __x, const optional<_Up>& __y)
   return *__x != *__y;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() < declval<const _Up&>()), bool),
@@ -1083,6 +1108,7 @@ operator<(const optional<_Tp>& __x, const optional<_Up>& __y)
   return *__x < *__y;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() > declval<const _Up&>()), bool),
@@ -1100,6 +1126,7 @@ operator>(const optional<_Tp>& __x, const optional<_Up>& __y)
   return *__x > *__y;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() <= declval<const _Up&>()), bool),
@@ -1117,6 +1144,7 @@ operator<=(const optional<_Tp>& __x, const optional<_Up>& __y)
   return *__x <= *__y;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() >= declval<const _Up&>()), bool),
@@ -1208,6 +1236,7 @@ _LIBCUDACXX_INLINE_VISIBILITY constexpr bool operator>=(nullopt_t, const optiona
 }
 
 // Comparisons with T
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() == declval<const _Up&>()), bool),
@@ -1217,6 +1246,7 @@ operator==(const optional<_Tp>& __x, const _Up& __v)
   return static_cast<bool>(__x) ? *__x == __v : false;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() == declval<const _Up&>()), bool),
@@ -1226,6 +1256,7 @@ operator==(const _Tp& __v, const optional<_Up>& __x)
   return static_cast<bool>(__x) ? __v == *__x : false;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() != declval<const _Up&>()), bool),
@@ -1235,6 +1266,7 @@ operator!=(const optional<_Tp>& __x, const _Up& __v)
   return static_cast<bool>(__x) ? *__x != __v : true;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() != declval<const _Up&>()), bool),
@@ -1244,6 +1276,7 @@ operator!=(const _Tp& __v, const optional<_Up>& __x)
   return static_cast<bool>(__x) ? __v != *__x : true;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() < declval<const _Up&>()), bool),
@@ -1253,6 +1286,7 @@ operator<(const optional<_Tp>& __x, const _Up& __v)
   return static_cast<bool>(__x) ? *__x < __v : true;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() < declval<const _Up&>()), bool),
@@ -1262,6 +1296,7 @@ operator<(const _Tp& __v, const optional<_Up>& __x)
   return static_cast<bool>(__x) ? __v < *__x : false;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() <= declval<const _Up&>()), bool),
@@ -1271,6 +1306,7 @@ operator<=(const optional<_Tp>& __x, const _Up& __v)
   return static_cast<bool>(__x) ? *__x <= __v : true;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() <= declval<const _Up&>()), bool),
@@ -1280,6 +1316,7 @@ operator<=(const _Tp& __v, const optional<_Up>& __x)
   return static_cast<bool>(__x) ? __v <= *__x : false;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() > declval<const _Up&>()), bool),
@@ -1289,6 +1326,7 @@ operator>(const optional<_Tp>& __x, const _Up& __v)
   return static_cast<bool>(__x) ? *__x > __v : false;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() > declval<const _Up&>()), bool),
@@ -1298,6 +1336,7 @@ operator>(const _Tp& __v, const optional<_Up>& __x)
   return static_cast<bool>(__x) ? __v > *__x : true;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() >= declval<const _Up&>()), bool),
@@ -1307,6 +1346,7 @@ operator>=(const optional<_Tp>& __x, const _Up& __v)
   return static_cast<bool>(__x) ? *__x >= __v : false;
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp, class _Up>
 _LIBCUDACXX_INLINE_VISIBILITY constexpr __enable_if_t<
   _LIBCUDACXX_TRAIT(is_convertible, decltype(declval<const _Tp&>() >= declval<const _Up&>()), bool),

diff --git a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.comp_with_t/equal.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.comp_with_t/equal.pass.cpp
@@ -34,6 +34,20 @@ __host__ __device__ constexpr bool operator==(const X& lhs, const X& rhs)
   return lhs.i_ == rhs.i_;
 }
 
+struct DeviceOnly
+{
+  __device__ friend constexpr bool operator==(const DeviceOnly& lhs, const DeviceOnly& rhs)
+  {
+    return true;
+  }
+};
+
+__global__ void ensure_exec_check_is_disabled()
+{
+  assert(optional<DeviceOnly>{} == DeviceOnly{});
+  assert(DeviceOnly{} == optional<DeviceOnly>{});
+}
+
 __host__ __device__ constexpr bool test()
 {
   {

diff --git a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.comp_with_t/greater.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.comp_with_t/greater.pass.cpp
@@ -34,6 +34,20 @@ __host__ __device__ constexpr bool operator>(const X& lhs, const X& rhs)
   return lhs.i_ > rhs.i_;
 }
 
+struct DeviceOnly
+{
+  __device__ friend constexpr bool operator>(const DeviceOnly& lhs, const DeviceOnly& rhs)
+  {
+    return true;
+  }
+};
+
+__global__ void ensure_exec_check_is_disabled()
+{
+  assert(optional<DeviceOnly>{} > DeviceOnly{});
+  assert(DeviceOnly{} > optional<DeviceOnly>{});
+}
+
 __host__ __device__ constexpr bool test()
 {
   {

diff --git a/...udacxx/test/libcudacxx/std/utilities/optional/optional.comp_with_t/greater_equal.pass.cpp b/...udacxx/test/libcudacxx/std/utilities/optional/optional.comp_with_t/greater_equal.pass.cpp
@@ -34,6 +34,20 @@ __host__ __device__ constexpr bool operator>=(const X& lhs, const X& rhs)
   return lhs.i_ >= rhs.i_;
 }
 
+struct DeviceOnly
+{
+  __device__ friend constexpr bool operator>=(const DeviceOnly& lhs, const DeviceOnly& rhs)
+  {
+    return true;
+  }
+};
+
+__global__ void ensure_exec_check_is_disabled()
+{
+  assert(optional<DeviceOnly>{} >= DeviceOnly{});
+  assert(DeviceOnly{} >= optional<DeviceOnly>{});
+}
+
 __host__ __device__ constexpr bool test()
 {
   {