From 3650bbeebc07f3d4cfd026bbb5088432bb1766e6 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Thu, 19 May 2016 20:13:53 +0000 Subject: [PATCH] [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 llvm-svn: 270108 --- clang/include/clang/Sema/Sema.h | 1 + clang/lib/Sema/SemaCUDA.cpp | 50 +++++++++++++- clang/lib/Sema/SemaDecl.cpp | 6 ++ clang/test/CodeGenCUDA/device-var-init.cu | 80 ++++++++++++++++++++++- clang/test/SemaCUDA/device-var-init.cu | 49 +++++++++++++- 5 files changed, 181 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 99d1379b2d03..c2b7cabced8f 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -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 //@{ diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index fee1ccf22b2c..7e201f4dae67 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -372,7 +372,7 @@ bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) { 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(CI->getInit())) @@ -384,6 +384,54 @@ bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) { 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; +} + // With -fcuda-host-device-constexpr, an unattributed constexpr function is // treated as implicitly __host__ __device__, unless: // * it is a variadic function (device-side variadic functions are not diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index d34d9b5688c3..9c442a61ce2e 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -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() ? diag::err_shared_var_init diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu index 8e29bddbd7c7..fa24dd95f921 100644 --- a/clang/test/CodeGenCUDA/device-var-init.cu +++ b/clang/test/CodeGenCUDA/device-var-init.cu @@ -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 __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. diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu index c717a3346330..ba63f6d8c0f5 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -58,6 +58,13 @@ __shared__ UC s_uc; __constant__ UC c_uc; // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__device__ UD d_ud; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__shared__ UD s_ud; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ UD c_ud; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} + __device__ ECI d_eci; // expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__device__ NED d_ned; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__shared__ NED s_ned; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NED c_ned; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} + __device__ NCV d_ncv; // expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__device__ VD d_vd; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__shared__ VD s_vd; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ VD c_vd; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} + __device__ NCF d_ncf; // expected-error@-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@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__shared__ T_B_NED s_t_b_ned; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ T_B_NED c_t_b_ned; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} + +__device__ T_F_NED d_t_f_ned; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__shared__ T_F_NED s_t_f_ned; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ T_F_NED c_t_f_ned; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} + +__device__ T_FA_NED d_t_fa_ned; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}} +__shared__ T_FA_NED s_t_fa_ned; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ T_FA_NED c_t_fa_ned; +// expected-error@-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@-1 {{initialization is not supported for __shared__ variables.}} static __shared__ UC s_uc; // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + static __shared__ NED s_ned; + // expected-error@-1 {{initialization is not supported for __shared__ variables.}} + static __device__ int ds; // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} static __constant__ int dc;