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@-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.}}
+}
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) {
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@-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.}}
+}
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) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
