Index: lib/CodeGen/CGDeclCXX.cpp =================================================================== --- lib/CodeGen/CGDeclCXX.cpp +++ lib/CodeGen/CGDeclCXX.cpp @@ -300,10 +300,97 @@ PtrArray->setComdat(C); } +static bool hasNonEmptyDefaultConstructors(CodeGenModule &CGM, + const CXXRecordDecl *RD) { + for (auto C : RD->ctors()) + if (C->isDefaultConstructor() && !CGM.isEmptyCudaConstructor(C)) + return true; + return false; +} + +bool CodeGenModule::isEmptyCudaConstructor(const CXXConstructorDecl *CD) { + // (E.2.3.1, CUDA 7.5) A constructor for a class type is considered + // empty at a point in the translation unit, if it is either a + // trivial constructor + if (CD->isTrivial()) + return true; + + // ... or it satisfies all of the following conditions: + // The constructor function has been defined. + if (!CD->isDefined()) + return false; + + // The constructor function has no parameters, + if (CD->getNumParams() != 0) + return false; + + // the initializer list is empty + for (const CXXCtorInitializer *CI: CD->inits()) + if (CI->isAnyMemberInitializer() && CI->isWritten()) + return false; + + // and the function body is an empty compound statement. + // That does not always work. + if (!CD->hasTrivialBody()) + return false; + + const CXXRecordDecl *RD = CD->getParent(); + // Its class has no virtual functions + for (auto Method: RD->methods()) + if (Method->isVirtual()) + return false; + + // .. and no virtual base classes. + if (RD->getNumVBases() != 0) + return false; + + // The default constructors of all base classes of its class can be + // considered empty. + for (auto &Base : RD->bases()) + if (hasNonEmptyDefaultConstructors(*this, + Base.getType()->getAsCXXRecordDecl())) + return false; + + // For all the nonstatic data members of its class that are of class type + // (or array thereof), the default constructors can be considered empty. + for (const auto *I : RD->decls()) + if (const FieldDecl *V = dyn_cast(I)) { + QualType T = V->getType(); + + if (const ArrayType *Ty = dyn_cast(T)) + while ((Ty = dyn_cast(T))) + T = Ty->getElementType(); + + if (const CXXRecordDecl *R = T->getAsCXXRecordDecl()) + if (hasNonEmptyDefaultConstructors(*this, R)) + return false; + } + + return true; +} + void CodeGenModule::EmitCXXGlobalVarDeclInitFunc(const VarDecl *D, llvm::GlobalVariable *Addr, bool PerformInit) { + + // According to E.2.3.1 in CUDA-7.5 Programming guide: + // __device__, __constant__ and __shared__ variables defined in + // namespace scope, that are of class type, cannot have a non-empty + // constructor... + if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice && + (D->hasAttr() || D->hasAttr() || + D->hasAttr())) { + if (const Expr *InitExpr = D->getAnyInitializer()) { + const CXXConstructExpr *CE = dyn_cast(InitExpr); + if (CE == nullptr || !isEmptyCudaConstructor(CE->getConstructor())) + Error(D->getLocation(), "dynamic initialization is not supported for " + "__device__, __constant__ and __shared__ " + "variables."); + } + return; + } + // Check if we've already initialized this decl. auto I = DelayedCXXInitPosition.find(D); if (I != DelayedCXXInitPosition.end() && I->second == ~0U) Index: lib/CodeGen/CodeGenModule.h =================================================================== --- lib/CodeGen/CodeGenModule.h +++ lib/CodeGen/CodeGenModule.h @@ -1119,6 +1119,10 @@ /// \breif Get the declaration of std::terminate for the platform. llvm::Constant *getTerminateFn(); + /// Returns whether given CXXConstructorDecl is an empty constructor + /// allowed by CUDA (E2.2.1, CUDA 7.5). + bool isEmptyCudaConstructor(const CXXConstructorDecl *CD); + private: llvm::Constant * GetOrCreateLLVMFunction(StringRef MangledName, llvm::Type *Ty, GlobalDecl D, Index: lib/CodeGen/CodeGenModule.cpp =================================================================== --- lib/CodeGen/CodeGenModule.cpp +++ lib/CodeGen/CodeGenModule.cpp @@ -1344,6 +1344,11 @@ if (LangOpts.OpenMP && LangOpts.OpenMPUseTLS && getContext().getTargetInfo().isTLSSupported() && isa(Global)) return false; + // Delay codegen for device-side CUDA variables. We need to have all + // constructor definitions available before we can determine whether + // we can skip them or produce an error. + if (LangOpts.CUDA && LangOpts.CUDAIsDevice && isa(Global)) + return false; return true; } @@ -2197,9 +2202,9 @@ && D->hasAttr()) { if (InitExpr) { const auto *C = dyn_cast(InitExpr); - if (C == nullptr || !C->getConstructor()->hasTrivialBody()) - Error(D->getLocation(), - "__shared__ variable cannot have an initialization."); + if (C == nullptr || !isEmptyCudaConstructor(C->getConstructor())) + Error(D->getLocation(), "initialization is not supported for " + "__shared__ variables."); } Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy)); } else if (!InitExpr) { Index: test/CodeGenCUDA/device-var-init.cu =================================================================== --- /dev/null +++ test/CodeGenCUDA/device-var-init.cu @@ -0,0 +1,371 @@ +// REQUIRES: nvptx-registered-target + +// Make sure we don't allow dynamic initialization for device +// variables, but accept empty constructors allowed by CUDA. + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \ +// RUN: | FileCheck %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \ +// RUN: -DERROR_CASE -verify -o /dev/null %s + +#include "Inputs/cuda.h" + +// no-constructor +struct NC { + int nc; +}; + +// empty constructor +struct EC { + int ec; + __device__ EC() {} +}; + +// empty constructor w/ initializer list +struct ECI { + int eci; + __device__ ECI() : eci(1) {} +}; + +// non-empty constructor +struct NEC { + int nec; + __device__ NEC() { nec = 1; } +}; + +// no-constructor, virtual method +struct NCV { + virtual void vm() {} +}; + +// no-constructor, no-constructor base class +struct NC_B_NC : NC { + int nc_b_nc; +}; + +// no-constructor, empty-constructor base class +struct NC_B_EC : EC { + int nc_b_ec; +}; + +// no-constructor, base class w/ constructor+init list. +struct NC_B_ECI : ECI { +}; + +// no-constructor, non-empty-constructor base class +struct NC_B_NEC : NEC { + int nc_b_nec; +}; + +// no-constructor, base class w/ virtual method +struct NC_B_NCV : NCV { + int nc_b_ncv; +}; + +// empty constructor, no-constructor base class +struct EC_B_NC : NC { + __device__ EC_B_NC() {} +}; + +// empty constructor, empty-constructor base class +struct EC_B_EC : EC { + __device__ EC_B_EC() {} +}; + +// empty constructor, base class w/ constructor+init list. +struct EC_B_ECI : ECI { + __device__ EC_B_ECI() {} +}; + +// empty constructor, non-empty-constructor base class +struct EC_B_NEC : NEC { + __device__ EC_B_NEC() {} +}; + +// empty constructor, non-empty-constructor base class +struct EC_B_NCV : NCV { + __device__ EC_B_NCV() {} +}; + +// no-constructor, no-constructor virtual base class +struct NC_V_NC : virtual NC { +}; + +// no-constructor, empty constructor virtual base class +struct NC_V_EC : virtual EC { +}; + +// empty constructor, no-constructor virtual base class +struct EC_V_NC : virtual NC { + __device__ EC_V_NC() {} +}; + +// empty constructor, empty constructor virtual base class +struct EC_V_EC : virtual EC { + __device__ EC_V_EC() {} +}; + +// no-constructor, no-constructor field +struct NC_F_NC { + NC nc_f_nc; +}; + +// no-constructor, empty-constructor field +struct NC_F_EC{ + EC nc_f_ec; +}; + +// no-constructor, empty-constructor+initializer field +struct NC_F_ECI{ + ECI nc_f_ec; +}; + +// no-constructor, non-empty-constructor field +struct NC_F_NEC { + NEC nc_f_nec; +}; + +// no-constructor, field w/ virtual method +struct NC_F_NCV { + NCV nc_f_ncv; +}; + +// no-constructor, no-constructor field +struct NC_FA_NC { + NC nc_fa_nc[2]; +}; + +// no-constructor, empty-constructor field +struct NC_FA_EC{ + EC nc_fa_ec[2]; +}; + +// no-constructor, non-empty-constructor field +struct NC_FA_NEC { + NEC nc_fa_nec[2]; +}; + +// no-constructor, field w/ virtual method +struct NC_FA_NCV { + NCV nc_fa_ncv[2]; +}; + +// No constructor, no initializer +__device__ NC nc_d; +// CHECK: @nc_d = addrspace(1) externally_initialized global %struct.NC zeroinitializer, +__shared__ NC nc_s; +// CHECK: @nc_s = addrspace(3) global %struct.NC undef +__constant__ NC nc_c; +// CHECK: @nc_c = addrspace(4) externally_initialized global %struct.NC zeroinitializer, + +// No constructor, initializer +__device__ NC nc_di = {1}; +// CHECK: @nc_di = addrspace(1) externally_initialized global %struct.NC { i32 1 } +#ifdef ERROR_CASE +__shared__ NC nc_si = {2}; // expected-error {{initialization is not supported for __shared__ variables.}} +#endif +__constant__ NC nc_ci = {3}; +// CHECK: @nc_ci = addrspace(4) externally_initialized global %struct.NC { i32 3 } + +// Empty constructor. +__device__ EC ec_d; +// CHECK: @ec_d = addrspace(1) externally_initialized global %struct.EC zeroinitializer +__shared__ EC ec_s; +// CHECK: @ec_s = addrspace(3) global %struct.EC undef +__constant__ EC ec_c; +// CHECK: @ec_c = addrspace(4) externally_initialized global %struct.EC zeroinitializer + +#ifdef ERROR_CASE +__device__ ECI deci; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ ECI seci; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ ECI ceci; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NEC dnec; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NEC snec; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NEC cnec; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NCV dncv; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NCV sncv; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NCV cncv; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +#endif + +// Make sure we apply initializer checks to base classes + +__device__ NC_B_NC nc_b_nc_d; +// CHECK: @nc_b_nc_d = addrspace(1) externally_initialized global %struct.NC_B_NC zeroinitializer, +__shared__ NC_B_NC nc_b_nc_s; +// CHECK: @nc_b_nc_s = addrspace(3) global %struct.NC_B_NC undef +__constant__ NC_B_NC nc_b_nc_c; +// CHECK: @nc_b_nc_c = addrspace(4) externally_initialized global %struct.NC_B_NC zeroinitializer + +__device__ NC_B_EC nc_b_ec_d; +// CHECK: @nc_b_ec_d = addrspace(1) externally_initialized global %struct.NC_B_EC zeroinitializer, +__shared__ NC_B_EC nc_b_ec_s; +// CHECK: @nc_b_ec_s = addrspace(3) global %struct.NC_B_EC undef +__constant__ NC_B_EC nc_b_ec_c; +// CHECK: @nc_b_ec_c = addrspace(4) externally_initialized global %struct.NC_B_EC zeroinitializer + +__device__ EC_B_NC ec_b_nc_d; +// CHECK: @ec_b_nc_d = addrspace(1) externally_initialized global %struct.EC_B_NC zeroinitializer, +__shared__ EC_B_NC ec_b_nc_s; +// CHECK: @ec_b_nc_s = addrspace(3) global %struct.EC_B_NC undef +__constant__ EC_B_NC ec_b_nc_c; +// CHECK: @ec_b_nc_c = addrspace(4) externally_initialized global %struct.EC_B_NC zeroinitializer + +__device__ EC_B_EC ec_b_ec_d; +// CHECK: @ec_b_ec_d = addrspace(1) externally_initialized global %struct.EC_B_EC zeroinitializer, +__shared__ EC_B_EC ec_b_ec_s; +// CHECK: @ec_b_ec_s = addrspace(3) global %struct.EC_B_EC undef +__constant__ EC_B_EC ec_b_ec_c; +// CHECK: @ec_b_ec_c = addrspace(4) externally_initialized global %struct.EC_B_EC zeroinitializer + +#ifdef ERROR_CASE +__device__ NC_B_ECI nc_b_eci_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_B_ECI nc_b_eci_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_B_ECI nc_b_eci_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_B_NEC nc_b_nec_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_B_NEC nc_b_nec_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_B_NEC nc_b_nec_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_B_NCV nc_b_ncv_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_B_NCV nc_b_ncv_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_B_NCV nc_b_ncv_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_F_NEC nc_f_nec_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_F_NEC nc_f_nec_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_F_NEC nc_f_nec_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_F_NCV nc_f_ncv_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_F_NCV nc_f_ncv_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_F_NCV nc_f_ncv_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_FA_NEC nc_fa_nec_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_FA_NEC nc_fa_nec_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_FA_NEC nc_fa_nec_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_FA_NCV nc_fa_ncv_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_FA_NCV nc_fa_ncv_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_FA_NCV nc_fa_ncv_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ EC_B_NEC ec_b_nec_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ EC_B_NEC ec_b_nec_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ EC_B_NEC ec_b_nec_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ EC_B_NCV ec_b_ncv_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ EC_B_NCV ec_b_ncv_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ EC_B_NCV ec_b_ncv_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} + +__device__ NC_V_NC nc_v_nc_d; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +__shared__ NC_V_NC nc_v_nc_s; +// expected-error@-1 {{initialization is not supported for __shared__ variables.}} +__constant__ NC_V_NC nc_v_nc_c; +// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}} +#endif + +__device__ NC_F_NC nc_f_nc_d; +// CHECK: @nc_f_nc_d = addrspace(1) externally_initialized global %struct.NC_F_NC zeroinitializer, +__shared__ NC_F_NC nc_f_nc_s; +// CHECK: @nc_f_nc_s = addrspace(3) global %struct.NC_F_NC undef +__constant__ NC_F_NC nc_f_nc_c; +// CHECK: @nc_f_nc_c = addrspace(4) externally_initialized global %struct.NC_F_NC zeroinitializer + +__device__ NC_F_EC nc_f_ec_d; +// CHECK: @nc_f_ec_d = addrspace(1) externally_initialized global %struct.NC_F_EC zeroinitializer, +__shared__ NC_F_EC nc_f_ec_s; +// CHECK: @nc_f_ec_s = addrspace(3) global %struct.NC_F_EC undef +__constant__ NC_F_EC nc_f_ec_c; +// CHECK: @nc_f_ec_c = addrspace(4) externally_initialized global %struct.NC_F_EC zeroinitializer + +__device__ NC_FA_NC nc_fa_nc_d; +// CHECK: @nc_fa_nc_d = addrspace(1) externally_initialized global %struct.NC_FA_NC zeroinitializer, +__shared__ NC_FA_NC nc_fa_nc_s; +// CHECK: @nc_fa_nc_s = addrspace(3) global %struct.NC_FA_NC undef +__constant__ NC_FA_NC nc_fa_nc_c; +// CHECK: @nc_fa_nc_c = addrspace(4) externally_initialized global %struct.NC_FA_NC zeroinitializer + +// Note: Despite CUDA guide indicating that empty constructors are OK +// for "nonstatic data members of its class that are of class type (or +// array thereof)", nvcc throws an error for an array of records with +// empty constructors. Clang does accept them. +__device__ NC_FA_EC nc_fa_ec_d; +// CHECK: @nc_fa_ec_d = addrspace(1) externally_initialized global %struct.NC_FA_EC zeroinitializer, +__shared__ NC_FA_EC nc_fa_ec_s; +// CHECK: @nc_fa_ec_s = addrspace(3) global %struct.NC_FA_EC undef +__constant__ NC_FA_EC nc_fa_ec_c; +// CHECK: @nc_fa_ec_c = addrspace(4) externally_initialized global %struct.NC_FA_EC zeroinitializer + +// We should not emit global initializers for device-side variables. +// CHECK-NOT: @__cxx_global_var_init + +// Make sure that initialization restrictions do not apply to local +// variables. +__device__ void df() { + ECI eci; + NEC nec; + NCV ncv; + NC_B_ECI nc_b_eci; + NC_B_NEC nc_b_nec; + NC_B_NCV nc_b_ncv; + NC_F_ECI nc_f_eci; + NC_F_NEC nc_f_nec; + NC_F_NCV nc_f_ncv; + EC_B_NEC ec_b_nec; + EC_B_NCV ec_b_ncv; + NC_V_NC nc_v_nc; +} + +// CHECK: define void @_Z2dfv() +// CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci) +// CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec) +// CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv) #2 +// CHECK: call void @_ZN8NC_B_ECIC1Ev(%struct.NC_B_ECI* %nc_b_eci) +// CHECK: call void @_ZN8NC_B_NECC1Ev(%struct.NC_B_NEC* %nc_b_nec) +// CHECK: call void @_ZN8NC_B_NCVC1Ev(%struct.NC_B_NCV* %nc_b_ncv) +// CHECK: call void @_ZN8NC_F_ECIC1Ev(%struct.NC_F_ECI* %nc_f_eci) +// CHECK: call void @_ZN8NC_F_NECC1Ev(%struct.NC_F_NEC* %nc_f_nec) +// CHECK: call void @_ZN8NC_F_NCVC1Ev(%struct.NC_F_NCV* %nc_f_ncv) +// CHECK: call void @_ZN8EC_B_NECC1Ev(%struct.EC_B_NEC* %ec_b_nec) +// CHECK: call void @_ZN8EC_B_NCVC1Ev(%struct.EC_B_NCV* %ec_b_ncv) +// CHECK: call void @_ZN7NC_V_NCC1Ev(%struct.NC_V_NC* %nc_v_nc) #2 +// CHECK: ret void + +// We should not emit global init function. +// CHECK-NOT: @_GLOBAL__sub_I