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