[CUDA][HIP] allow trivial ctor/dtor in device var init

Treat ctor/dtor in device var init as host device function
so that they can be used to initialize file-scope
device variables to match nvcc behavior. If they are non-trivial
they will be diagnosed.

We cannot add implicit host device attrs to non-trivial
ctor/dtor since determining whether they are non-trivial
needs to know whether they have a trivial body and all their
member and base classes' ctor/dtor have trivial body, which
is affected by where their bodies are defined or instantiated.

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

This reverts commit 876f99a.

Fixes: llvm#72261

Fixes: SWDEV-432412

Fixes: SWDEV-433956
Change-Id: I711db63a2166ce77dea06aad5d04cae10d96ce24

Author: yxsamliu

clang/include/clang/Sema/Sema.h

@@ -13193,10 +13193,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

@@ -225,6 +225,15 @@ Sema::CUDAFunctionPreference
 Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
                              const FunctionDecl *Callee) {
   assert(Callee && "Callee must be valid.");
+
+  // Treat ctor/dtor as host device function in device var initializer to allow
+  // trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor
+  // will be diagnosed by checkAllowedCUDAInitializer.
+  if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar &&
+      CurCUDATargetCtx.Target == CFT_Device &&
+      (isa<CXXConstructorDecl>(Callee) || isa<CXXDestructorDecl>(Callee)))
+    return CFP_HostDevice;
+
   CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
   CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
 
@@ -730,22 +739,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

@@ -15884,9 +15884,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/lib/Sema/SemaOverload.cpp

@@ -1404,10 +1404,8 @@ bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old,
     // Don't allow overloading of destructors.  (In theory we could, but it
     // would be a giant change to clang.)
     if (!isa<CXXDestructorDecl>(New)) {
-      CUDAFunctionTarget NewTarget = IdentifyCUDATarget(
-                             New, isa<CXXConstructorDecl>(New)),
-                         OldTarget = IdentifyCUDATarget(
-                             Old, isa<CXXConstructorDecl>(New));
+      CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New),
+                         OldTarget = IdentifyCUDATarget(Old);
       if (NewTarget != CFT_InvalidTarget) {
         assert((OldTarget != CFT_InvalidTarget) &&
                "Unexpected invalid target.");

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}}
@@ -39,7 +39,6 @@ struct A2_with_device_ctor {
 };
 // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
 // expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
-// expected-note@-4 {{candidate inherited constructor not viable: call to __device__ function from __host__ function}}
 
 struct B2_with_implicit_default_ctor : A2_with_device_ctor {
   using A2_with_device_ctor::A2_with_device_ctor;
@@ -84,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 {