r270108 - [CUDA] Do not allow non-empty destructors for global device-side variables.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Thu May 19 13:13:54 PDT 2016
Author: tra
Date: Thu May 19 15:13:53 2016
New Revision: 270108
URL: http://llvm.org/viewvc/llvm-project?rev=270108&view=rev
Log:
[CUDA] Do not allow non-empty destructors for global device-side variables.
According to Cuda Programming guide (v7.5, E2.3.1):
> __device__, __constant__ and __shared__ variables defined in namespace
> scope, that are of class type, cannot have a non-empty constructor or a
> non-empty destructor.
Clang already deals with device-side constructors (see D15305).
This patch enforces similar rules for destructors.
Differential Revision: http://reviews.llvm.org/D20140
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/CodeGenCUDA/device-var-init.cu
cfe/trunk/test/SemaCUDA/device-var-init.cu
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=270108&r1=270107&r2=270108&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Thu May 19 15:13:53 2016
@@ -9036,6 +9036,7 @@ public:
/// \return true if \p CD can be considered empty according to CUDA
/// (E.2.3.1 in CUDA 7.5 Programming guide).
bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
+ bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
/// \name Code completion
//@{
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=270108&r1=270107&r2=270108&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Thu May 19 15:13:53 2016
@@ -372,7 +372,7 @@ bool Sema::isEmptyCudaConstructor(Source
return false;
// The only form of initializer allowed is an empty constructor.
- // This will recursively checks all base classes and member initializers
+ // This will recursively check all base classes and member initializers
if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
if (const CXXConstructExpr *CE =
dyn_cast<CXXConstructExpr>(CI->getInit()))
@@ -381,6 +381,54 @@ bool Sema::isEmptyCudaConstructor(Source
}))
return false;
+ return true;
+}
+
+bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
+ // No destructor -> no problem.
+ if (!DD)
+ return true;
+
+ if (!DD->isDefined() && DD->isTemplateInstantiation())
+ InstantiateFunctionDefinition(Loc, DD->getFirstDecl());
+
+ // (E.2.3.1, CUDA 7.5) A destructor for a class type is considered
+ // empty at a point in the translation unit, if it is either a
+ // trivial constructor
+ if (DD->isTrivial())
+ return true;
+
+ // ... or it satisfies all of the following conditions:
+ // The destructor function has been defined.
+ // and the function body is an empty compound statement.
+ if (!DD->hasTrivialBody())
+ return false;
+
+ const CXXRecordDecl *ClassDecl = DD->getParent();
+
+ // Its class has no virtual functions and no virtual base classes.
+ if (ClassDecl->isDynamicClass())
+ return false;
+
+ // Only empty destructors are allowed. This will recursively check
+ // destructors for all base classes...
+ if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
+ if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
+ return isEmptyCudaDestructor(Loc, RD->getDestructor());
+ return true;
+ }))
+ return false;
+
+ // ... and member fields.
+ if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) {
+ if (CXXRecordDecl *RD = Field->getType()
+ ->getBaseElementTypeUnsafe()
+ ->getAsCXXRecordDecl())
+ return isEmptyCudaDestructor(Loc, RD->getDestructor());
+ return true;
+ }))
+ return false;
+
return true;
}
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=270108&r1=270107&r2=270108&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu May 19 15:13:53 2016
@@ -10442,6 +10442,12 @@ Sema::FinalizeDeclaration(Decl *ThisDecl
AllowedInit = VD->getInit()->isConstantInitializer(
Context, VD->getType()->isReferenceType());
+ // Also make sure that destructor, if there is one, is empty.
+ if (AllowedInit)
+ if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+ AllowedInit =
+ isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+
if (!AllowedInit) {
Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
? diag::err_shared_var_init
Modified: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/device-var-init.cu?rev=270108&r1=270107&r2=270108&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Thu May 19 15:13:53 2016
@@ -24,6 +24,16 @@ struct EC {
__device__ EC(int) {} // -- not allowed
};
+// empty destructor
+struct ED {
+ __device__ ~ED() {} // -- allowed
+};
+
+struct ECD {
+ __device__ ECD() {} // -- allowed
+ __device__ ~ECD() {} // -- allowed
+};
+
// empty templated constructor -- allowed with no arguments
struct ETC {
template <typename... T> __device__ ETC(T...) {}
@@ -35,6 +45,12 @@ struct UC {
__device__ UC();
};
+// undefined destructor -- not allowed
+struct UD {
+ int ud;
+ __device__ ~UD();
+};
+
// empty constructor w/ initializer list -- not allowed
struct ECI {
int eci;
@@ -47,12 +63,23 @@ struct NEC {
__device__ NEC() { nec = 1; }
};
+// non-empty destructor -- not allowed
+struct NED {
+ int ned;
+ __device__ ~NED() { ned = 1; }
+};
+
// no-constructor, virtual method -- not allowed
struct NCV {
int ncv;
__device__ virtual void vm() {}
};
+// virtual destructor -- not allowed.
+struct VD {
+ __device__ virtual ~VD() {}
+};
+
// dynamic in-class field initializer -- not allowed
__device__ int f();
struct NCF {
@@ -107,6 +134,20 @@ __shared__ EC s_ec;
__constant__ EC c_ec;
// CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+__device__ ED d_ed;
+// CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
+__shared__ ED s_ed;
+// CHECK: @s_ed = addrspace(3) global %struct.ED undef,
+__constant__ ED c_ed;
+// CHECK: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+
+__device__ ECD d_ecd;
+// CHECK: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
+__shared__ ECD s_ecd;
+// CHECK: @s_ecd = addrspace(3) global %struct.ECD undef,
+__constant__ ECD c_ecd;
+// CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+
__device__ ETC d_etc;
// CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
__shared__ ETC s_etc;
@@ -180,6 +221,17 @@ struct T_FA_NEC {
NEC nec[2];
};
+
+// Inherited from or incapsulated class with non-empty desstructor --
+// not allowed
+struct T_B_NED : NED {};
+struct T_F_NED {
+ NED ned;
+};
+struct T_FA_NED {
+ NED ned[2];
+};
+
// We should not emit global initializers for device-side variables.
// CHECK-NOT: @__cxx_global_var_init
@@ -190,16 +242,26 @@ __device__ void df() {
// CHECK-NOT: call
EC ec;
// CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec)
+ ED ed;
+ // CHECK-NOT: call
+ ECD ecd;
+ // CHECK: call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
ETC etc;
// CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
UC uc;
// CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc)
+ UD ud;
+ // CHECK-NOT: call
ECI eci;
// CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
NEC nec;
// CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+ NED ned;
+ // CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
NCV ncv;
// CHECK-NOT: call
+ VD vd;
+ // CHECK: call void @_ZN2VDC1Ev(%struct.VD* %vd)
NCF ncf;
// CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
NCFS ncfs;
@@ -226,6 +288,12 @@ __device__ void df() {
// CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
T_FA_NEC t_fa_nec;
// CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+ T_B_NED t_b_ned;
+ // CHECK-NOT: call
+ T_F_NED t_f_ned;
+ // CHECK-NOT: call
+ T_FA_NED t_fa_ned;
+ // CHECK-NOT: call
static __shared__ EC s_ec;
// CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
static __shared__ ETC s_etc;
@@ -234,9 +302,17 @@ __device__ void df() {
// anchor point separating constructors and destructors
df(); // CHECK: call void @_Z2dfv()
- // CHECK-NOT: call
+ // Verify that we only call non-empty destructors
+ // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned) #6
+ // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned) #6
+ // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned) #6
+ // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
+ // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
+ // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
+ // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
+ // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
- // CHECK: ret void
+ // CHECK-NEXT: ret void
}
// We should not emit global init function.
Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=270108&r1=270107&r2=270108&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu Thu May 19 15:13:53 2016
@@ -58,6 +58,13 @@ __shared__ UC s_uc;
__constant__ UC c_uc;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ UD d_ud;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UD s_ud;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UD c_ud;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
__device__ ECI d_eci;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
__shared__ ECI s_eci;
@@ -72,6 +79,13 @@ __shared__ NEC s_nec;
__constant__ NEC c_nec;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ NED d_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NED s_ned;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NED c_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
__device__ NCV d_ncv;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
__shared__ NCV s_ncv;
@@ -79,6 +93,13 @@ __shared__ NCV s_ncv;
__constant__ NCV c_ncv;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ VD d_vd;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ VD s_vd;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ VD c_vd;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
__device__ NCF d_ncf;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
__shared__ NCF s_ncf;
@@ -152,13 +173,37 @@ __shared__ T_FA_NEC s_t_fa_nec;
__constant__ T_FA_NEC c_t_fa_nec;
// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
-// Make sure that initialization restrictions do not apply to local
-// variables.
+__device__ T_B_NED d_t_b_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_B_NED s_t_b_ned;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_B_NED c_t_b_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_F_NED d_t_f_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_F_NED s_t_f_ned;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_F_NED c_t_f_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_FA_NED d_t_fa_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_FA_NED s_t_fa_ned;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_FA_NED c_t_fa_ned;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+// Verify that only __shared__ local variables may be static on device
+// side and that they are not allowed to be initialized.
__device__ void df_sema() {
static __shared__ NCFS s_ncfs;
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
static __shared__ UC s_uc;
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+ static __shared__ NED s_ned;
+ // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+
static __device__ int ds;
// expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
static __constant__ int dc;
More information about the cfe-commits
mailing list