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

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2020-05-04T21:52:04-04:00
New Revision: d75a6e93ae99bfcd67219454307da56ebd155d45

URL: https://github.com/llvm/llvm-project/commit/d75a6e93ae99bfcd67219454307da56ebd155d45
DIFF: https://github.com/llvm/llvm-project/commit/d75a6e93ae99bfcd67219454307da56ebd155d45.diff

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

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.

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.

Differential Revision: https://reviews.llvm.org/D79367

Added: 
    clang/test/SemaCUDA/union-init.cu

Modified: 
    clang/lib/Sema/SemaCUDA.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index faab250e58ff..73d190891b0f 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -426,6 +426,10 @@ bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
   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) {
@@ -465,6 +469,11 @@ bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
   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) {

diff  --git a/clang/test/SemaCUDA/union-init.cu b/clang/test/SemaCUDA/union-init.cu
new file mode 100644
index 000000000000..a633975e3776
--- /dev/null
+++ b/clang/test/SemaCUDA/union-init.cu
@@ -0,0 +1,44 @@
+// 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 at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ D d;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ void foo() {
+  __shared__ B b;
+  __shared__ C c;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+  __shared__ D d;
+  // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+}