[CUDA][HIP] Fix empty ctor/dtor check for union

[yaxunl]

Summary

union ctor does not call ctors of its data members. union dtor does not call dtors of its data members.
Also union does not have base class.

https://godbolt.org/z/8RxZeG

Currently when clang checks whether union has an empty ctor/dtor, it checks the ctors/dtors of its
data members. This causes incorrectly diagnose device side global variables and shared variables as
having non-empty ctors/dtors.

This patch fixes that.

[tra]

Nice! Thank you for the patch.

clang/lib/Sema/SemaCUDA.cpp

bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
   // and the function body is an empty compound statement.
   if (!(CD->hasTrivialBody() && CD->getNumParams() == 0))
     return false;
 
   // Its class has no virtual functions and no virtual base classes.
   if (CD->getParent()->isDynamicClass())
     return false;
 
+  // Union ctor does not call ctors of its data members.
+  if (CD->getParent()->isUnion())
+    return true;
+
   // The only form of initializer allowed is an empty constructor.
   // This will recursively check all base classes and member initializers
   if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
         if (const CXXConstructExpr *CE =
                 dyn_cast<CXXConstructExpr>(CI->getInit()))
           return isEmptyCudaConstructor(Loc, CE->getConstructor());
         return false;
       }))
if (!DD->hasTrivialBody())
     return false;
 
   const CXXRecordDecl *ClassDecl = DD->getParent();
 
   // Its class has no virtual functions and no virtual base classes.
   if (ClassDecl->isDynamicClass())
     return false;
 
+  // Union does not have base class and union dtor does not call dtors of its
+  // data members.
+  if (DD->getParent()->isUnion())
+    return true;
+
   // Only empty destructors are allowed. This will recursively check
   // destructors for all base classes...
   if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
         if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
           return isEmptyCudaDestructor(Loc, RD->getDestructor());
         return true;
       }))
     return false;

clang/test/SemaCUDA/union-init.cu

+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
+
+#include "Inputs/cuda.h"
+
+struct A {
+  int a;
+  __device__ A() { a = 1; }
+  __device__ ~A() { a = 2; }
+};
+
+// This can be a global var since ctor/dtors of data members are not called.
+union B {
+  A a;
+  __device__ B() {}
+  __device__ ~B() {}
+};
+
+// This cannot be a global var since it has a dynamic ctor.
+union C {
+  A a;
+  __device__ C() { a.a = 3; }
+  __device__ ~C() {}
+};
+
+// This cannot be a global var since it has a dynamic dtor.
+union D {
+  A a;
+  __device__ D() { }
+  __device__ ~D() { a.a = 4; }
+};
+
+__device__ B b;
+__device__ C c;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ D d;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ void foo() {
+  __shared__ B b;
+  __shared__ C c;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  __shared__ D d;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+}