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

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.

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.


https://reviews.llvm.org/D79367

Files:
  clang/lib/Sema/SemaCUDA.cpp
  clang/test/SemaCUDA/union-init.cu


Index: clang/test/SemaCUDA/union-init.cu
===================================================================
--- /dev/null
+++ 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.}}
+}
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -426,6 +426,10 @@
   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 @@
   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) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D79367.261931.patch
Type: text/x-patch
Size: 2361 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200504/f21e151a/attachment.bin>