Revert "[CUDA][HIP] make trivial ctor/dtor host device"

This reverts commit 20e8466.

Change-Id: Iacae3b6a2ed9a5a66c1956d494054a034e606a6c

Author: yxsamliu

clang/include/clang/Sema/Sema.h

@@ -13266,10 +13266,6 @@ class Sema final {
   void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD,
                                    const LookupResult &Previous);
 
-  /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to a
-  /// trivial cotr/dtor that does not have host and device attributes.
-  void maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FunctionDecl *FD);
-
   /// May add implicit CUDAConstantAttr attribute to VD, depending on VD
   /// and current compilation settings.
   void MaybeAddCUDAConstantAttr(VarDecl *VD);

clang/lib/Sema/SemaCUDA.cpp

@@ -730,22 +730,6 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
 }
 
-// If a trivial ctor/dtor has no host/device
-// attributes, make it implicitly host device function.
-void Sema::maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FunctionDecl *FD) {
-  bool IsTrivialCtor = false;
-  if (auto *CD = dyn_cast<CXXConstructorDecl>(FD))
-    IsTrivialCtor = isEmptyCudaConstructor(SourceLocation(), CD);
-  bool IsTrivialDtor = false;
-  if (auto *DD = dyn_cast<CXXDestructorDecl>(FD))
-    IsTrivialDtor = isEmptyCudaDestructor(SourceLocation(), DD);
-  if ((IsTrivialCtor || IsTrivialDtor) && !FD->hasAttr<CUDAHostAttr>() &&
-      !FD->hasAttr<CUDADeviceAttr>()) {
-    FD->addAttr(CUDAHostAttr::CreateImplicit(Context));
-    FD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
-  }
-}
-
 // TODO: `__constant__` memory may be a limited resource for certain targets.
 // A safeguard may be needed at the end of compilation pipeline if
 // `__constant__` memory usage goes beyond limit.

clang/lib/Sema/SemaDecl.cpp

@@ -15974,9 +15974,6 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
   if (FD && !FD->isDeleted())
     checkTypeSupport(FD->getType(), FD->getLocation(), FD);
 
-  if (LangOpts.CUDA)
-    maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FD);
-
   return dcl;
 }
 

clang/test/SemaCUDA/call-host-fn-from-device.cu

@@ -12,7 +12,7 @@ extern "C" void host_fn() {}
 struct Dummy {};
 
 struct S {
-  S() { static int nontrivial_ctor = 1; }
+  S() {}
   // expected-note@-1 2 {{'S' declared here}}
   ~S() { host_fn(); }
   // expected-note@-1 {{'~S' declared here}}

clang/test/SemaCUDA/default-ctor.cu

@@ -25,7 +25,7 @@ __device__ void fd() {
   InD ind;
   InH inh; // expected-error{{no matching constructor for initialization of 'InH'}}
   InHD inhd;
-  Out out;
+  Out out; // expected-error{{no matching constructor for initialization of 'Out'}}
   OutD outd;
   OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}}
   OutHD outhd;

clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu

@@ -6,7 +6,7 @@
 // Test 1: collision between two bases
 
 struct A1_with_host_ctor {
-  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+  A1_with_host_ctor() {}
 };
 
 struct B1_with_device_ctor {

clang/test/SemaCUDA/implicit-member-target-collision.cu

@@ -6,7 +6,7 @@
 // Test 1: collision between two bases
 
 struct A1_with_host_ctor {
-  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+  A1_with_host_ctor() {}
 };
 
 struct B1_with_device_ctor {

clang/test/SemaCUDA/implicit-member-target-inherited.cu

@@ -6,7 +6,7 @@
 // Test 1: infer inherited default ctor to be host.
 
 struct A1_with_host_ctor {
-  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+  A1_with_host_ctor() {}
 };
 // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
 // expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
@@ -83,7 +83,7 @@ void hostfoo3() {
 // Test 4: infer inherited default ctor from a field, not a base
 
 struct A4_with_host_ctor {
-  A4_with_host_ctor() { static int nontrivial_ctor = 1; }
+  A4_with_host_ctor() {}
 };
 
 struct B4_with_inherited_host_ctor : A4_with_host_ctor{

clang/test/SemaCUDA/implicit-member-target.cu

@@ -6,7 +6,7 @@
 // Test 1: infer default ctor to be host.
 
 struct A1_with_host_ctor {
-  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
+  A1_with_host_ctor() {}
 };
 
 // The implicit default constructor is inferred to be host because it only needs
@@ -75,7 +75,7 @@ void hostfoo3() {
 // Test 4: infer default ctor from a field, not a base
 
 struct A4_with_host_ctor {
-  A4_with_host_ctor() { static int nontrivial_ctor = 1; }
+  A4_with_host_ctor() {}
 };
 
 struct B4_with_implicit_default_ctor {