fix complex and removed user specialization

Author: fbusato

libcudacxx/include/cuda/__type_traits/is_trivially_copyable.h

@@ -42,7 +42,6 @@ _CCCL_BEGIN_NAMESPACE_CUDA
 template <typename _Tp, typename = void>
 inline constexpr bool __is_aggregate_trivially_copyable_v = false;
 
-//! Users are allowed to specialize this variable template for their own types
 template <typename _Tp>
 inline constexpr bool __is_trivially_copyable_v =
   ::cuda::std::is_trivially_copyable_v<_Tp> || ::cuda::std::__is_extended_floating_point_v<_Tp>
@@ -67,11 +66,17 @@ inline constexpr bool __is_trivially_copyable_v<::cuda::std::pair<_T1, _T2>> =
 template <typename... _Ts>
 inline constexpr bool __is_trivially_copyable_v<::cuda::std::tuple<_Ts...>> = (__is_trivially_copyable_v<_Ts> && ...);
 
-template <typename _Tp>
-inline constexpr bool __is_trivially_copyable_v<::cuda::std::complex<_Tp>> = __is_trivially_copyable_v<_Tp>;
+template <>
+inline constexpr bool __is_trivially_copyable_v<::cuda::std::complex<::__half>> = true;
 
-template <typename _Tp>
-inline constexpr bool __is_trivially_copyable_v<::cuda::complex<_Tp>> = __is_trivially_copyable_v<_Tp>;
+template <>
+inline constexpr bool __is_trivially_copyable_v<::cuda::std::complex<::__nv_bfloat16>> = true;
+
+template <>
+inline constexpr bool __is_trivially_copyable_v<::cuda::complex<::__half>> = true;
+
+template <>
+inline constexpr bool __is_trivially_copyable_v<::cuda::complex<::__nv_bfloat16>> = true;
 
 // if all the previous conditions fail, check if the type is an aggregate and all its members are trivially copyable
 template <typename _Tp>
@@ -88,7 +93,6 @@ inline constexpr bool
 template <typename _Tp>
 inline constexpr bool is_trivially_copyable_v = __is_trivially_copyable_v<::cuda::std::remove_const_t<_Tp>>;
 
-// defined as alias so users cannot specialize it (they should specialize the variable template instead)
 template <typename _Tp>
 using is_trivially_copyable = ::cuda::std::bool_constant<is_trivially_copyable_v<_Tp>>;
 

libcudacxx/test/libcudacxx/cuda/type_traits/is_trivially_copyable.aggr.pass.cpp

@@ -47,40 +47,6 @@ struct ArrayMember
   int values[2];
 };
 
-//----------------------------------------------------------------------------------------------------------------------
-// custom type
-
-struct CustomNonTrivialType
-{
-  int x;
-
-  CustomNonTrivialType() = default;
-  __host__ __device__ CustomNonTrivialType(const CustomNonTrivialType&) {}
-};
-
-template <>
-constexpr bool cuda::is_trivially_copyable_v<CustomNonTrivialType> = true;
-
-struct SingleMemberCustom
-{
-  CustomNonTrivialType x;
-};
-
-struct DerivedStructCustom : SingleMemberCustom
-{
-  CustomNonTrivialType y;
-};
-
-struct NestedStructCustom
-{
-  CustomNonTrivialType z;
-};
-
-struct ArrayMemberCustom
-{
-  CustomNonTrivialType values[2];
-};
-
 //----------------------------------------------------------------------------------------------------------------------
 // non trivially copyable type
 
@@ -89,9 +55,9 @@ struct NonTriviallyCopyable
   __host__ __device__ NonTriviallyCopyable(const NonTriviallyCopyable&) {};
 };
 
-struct RelaxedWithNonRelaxedMember
+struct AggregateWithNonTriviallyCopyableMember
 {
-  CustomNonTrivialType x;
+  int x;
   NonTriviallyCopyable y;
 };
 
@@ -101,10 +67,8 @@ __host__ __device__ void test()
   test_is_trivially_copyable<DerivedStruct>();
   test_is_trivially_copyable<NesterStruct>();
   test_is_trivially_copyable<ArrayMember>();
-  test_is_trivially_copyable<ArrayMember>();
-  test_is_trivially_copyable<DerivedStructCustom>();
-  test_is_trivially_copyable<NestedStructCustom>();
-  test_is_not_trivially_copyable<RelaxedWithNonRelaxedMember>();
+  test_is_not_trivially_copyable<NonTriviallyCopyable>();
+  test_is_not_trivially_copyable<AggregateWithNonTriviallyCopyableMember>();
 }
 
 int main(int, char**)