maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Yaxun (Sam) Liu 2020-05-04 17:27:41 -04:00
parent ad5fad0ac5
commit d75a6e93ae
2 changed files with 53 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
clang/lib/Sema/SemaCUDA.cpp
View File

@ -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) {

44
clang/test/SemaCUDA/union-init.cu Normal file
View File

@ -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.}}
}