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

yxsamliu · Makarand Maydeo · commit 876f99a2dc86 · 2023-11-21T00:14:11.000-05:00
Make trivial ctor/dtor implicitly host device functions so that they can be used to initialize file-scope device variables to match nvcc behavior. Fixes: llvm#72261 Fixes: SWDEV-432412 cherry-pick of: llvm#72394 [CUDA][HIP] ignore implicit host/device attr for override When deciding whether a previous function declaration is an overload or override, implicit host/device attrs should not be considered. This fixes the failure for the following code: `template <typename T> class C { explicit C() {}; }; template <> C<int>::C() {}; ` The issue was introduced by llvm#72394 sine the template specialization is treated as overload due to implicit host/device attrs are considered for overload/override differentiation. cherry-pick of llvm#72815 Change-Id: Ie896cc0e4d5d82d5af48e08a996a3b392285d9ee
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
@@ -13193,6 +13193,10 @@ 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);
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
@@ -730,6 +730,22 @@ 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.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
@@ -15884,6 +15884,9 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
   if (FD && !FD->isDeleted())
     checkTypeSupport(FD->getType(), FD->getLocation(), FD);
 
+  if (LangOpts.CUDA)
+    maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FD);
+
   return dcl;
 }
 
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
@@ -1404,8 +1404,10 @@ 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),
-                         OldTarget = IdentifyCUDATarget(Old);
+      CUDAFunctionTarget NewTarget = IdentifyCUDATarget(
+                             New, isa<CXXConstructorDecl>(New)),
+                         OldTarget = IdentifyCUDATarget(
+                             Old, isa<CXXConstructorDecl>(New));
       if (NewTarget != CFT_InvalidTarget) {
         assert((OldTarget != CFT_InvalidTarget) &&
                "Unexpected invalid target.");
diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -12,7 +12,7 @@ extern "C" void host_fn() {}
 struct Dummy {};
 
 struct S {
-  S() {}
+  S() { static int nontrivial_ctor = 1; }
   // expected-note@-1 2 {{'S' declared here}}
   ~S() { host_fn(); }
   // expected-note@-1 {{'~S' declared here}}
diff --git a/clang/test/SemaCUDA/default-ctor.cu b/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; // expected-error{{no matching constructor for initialization of 'Out'}}
+  Out out;
   OutD outd;
   OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}}
   OutHD outhd;
diff --git a/clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/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() {}
+  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
 };
 
 struct B1_with_device_ctor {
diff --git a/clang/test/SemaCUDA/implicit-member-target-collision.cu b/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() {}
+  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
 };
 
 struct B1_with_device_ctor {
diff --git a/clang/test/SemaCUDA/implicit-member-target-inherited.cu b/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() {}
+  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
 };
 // expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
 // expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
@@ -39,6 +39,7 @@ 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;
@@ -83,7 +84,7 @@ void hostfoo3() {
 // Test 4: infer inherited default ctor from a field, not a base
 
 struct A4_with_host_ctor {
-  A4_with_host_ctor() {}
+  A4_with_host_ctor() { static int nontrivial_ctor = 1; }
 };
 
 struct B4_with_inherited_host_ctor : A4_with_host_ctor{
diff --git a/clang/test/SemaCUDA/implicit-member-target.cu b/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() {}
+  A1_with_host_ctor() { static int nontrivial_ctor = 1; }
 };
 
 // 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() {}
+  A4_with_host_ctor() { static int nontrivial_ctor = 1; }
 };
 
 struct B4_with_implicit_default_ctor {
diff --git a/clang/test/SemaCUDA/trivial-ctor-dtor.cu b/clang/test/SemaCUDA/trivial-ctor-dtor.cu
@@ -0,0 +1,56 @@
+// RUN: %clang_cc1 -isystem %S/Inputs  -fsyntax-only -verify %s
+// RUN: %clang_cc1 -isystem %S/Inputs -fcuda-is-device -fsyntax-only -verify %s
+
+#include <cuda.h>
+
+// Check trivial ctor/dtor
+struct A {
+  int x;
+  A() {}
+  ~A() {}
+};
+
+__device__ A a;
+
+// Check trivial ctor/dtor of template class
+template<typename T>
+struct TA {
+  T x;
+  TA() {}
+  ~TA() {}
+};
+
+__device__ TA<int> ta;
+
+// Check non-trivial ctor/dtor in parent template class
+template<typename T>
+struct TB {
+  T x;
+  TB() { static int nontrivial_ctor = 1; }
+  ~TB() {}
+};
+
+template<typename T>
+struct TC : TB<T> {
+  T x;
+  TC() {}
+  ~TC() {}
+};
+
+__device__ TC<int> tc; //expected-error {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}
+
+// Check trivial ctor specialization
+template <typename T>
+struct C { //expected-note {{candidate constructor (the implicit copy constructor) not viable}}
+           //expected-note@-1 {{candidate constructor (the implicit move constructor) not viable}}
+    explicit C() {};
+};
+
+template <> C<int>::C() {};
+__device__ C<int> ci_d;
+C<int> ci_h;
+
+// Check non-trivial ctor specialization
+template <> C<float>::C() { static int nontrivial_ctor = 1; } //expected-note {{candidate constructor not viable: call to __host__ function from __device__ function}}
+__device__ C<float> cf_d; //expected-error {{no matching constructor for initialization of 'C<float>'}}
+C<float> cf_h;

Original file line number	Diff line number	Diff line change
`@@ -15884,6 +15884,9 @@ Decl Sema::ActOnFinishFunctionBody(Decl dcl, Stmt *Body,`
`15884`	`15884`	`if (FD && !FD->isDeleted())`
`15885`	`15885`	`checkTypeSupport(FD->getType(), FD->getLocation(), FD);`
`15886`	`15886`
	`15887`	`+ if (LangOpts.CUDA)`
	`15888`	`+ maybeAddCUDAHostDeviceAttrsToTrivialCtorDtor(FD);`
	`15889`	`+`
`15887`	`15890`	`return dcl;`
`15888`	`15891`	`}`
`15889`	`15892`