tra created this revision.
tra added reviewers: rsmith, jingyue, jpienaar.
tra added a subscriber: cfe-commits.
In general CUDA does not allow dynamic initialization of
global device-side variables except for records with empty constructors as
described in section [[
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory-qualifiers
| E.2.3.1 of
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 or a non-empty
> destructor.
> A constructor for a class type is considered empty at a point in the
> translation unit,
> if it is either a trivial constructor or it satisfies all of the following
> conditions:
> * The constructor function has been defined.
> * The constructor function has no parameters, the initializer list is empty
> and the function body is an empty compound statement.
> * Its class has no virtual functions and no virtual base classes.
> * The default constructors of all base classes of its class can be considered
> empty.
> * For all the nonstatic data members of its class that are of class type (or
> array thereof), the default constructors can be considered empty.
Clang is already enforcing no-initializers for __shared__ variables, but
currently allows dynamic initialization for __device__ and __constant__
variables.
This patch applies initializer checks for all device-side variables.
Empty constructors are accepted, but no code is generated for them.
http://reviews.llvm.org/D15305
Files:
lib/CodeGen/CGDeclCXX.cpp
lib/CodeGen/CodeGenModule.cpp
lib/CodeGen/CodeGenModule.h
test/CodeGenCUDA/device-var-init.cu
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
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<VarDecl>(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<VarDecl>(Global))
+ return false;
return true;
}
@@ -2197,9 +2202,9 @@
&& D->hasAttr<CUDASharedAttr>()) {
if (InitExpr) {
const auto *C = dyn_cast<CXXConstructExpr>(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: 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<FieldDecl>(I)) {
+ QualType T = V->getType();
+
+ if (const ArrayType *Ty = dyn_cast<ArrayType>(T))
+ while ((Ty = dyn_cast<ArrayType>(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<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+ D->hasAttr<CUDASharedAttr>())) {
+ if (const Expr *InitExpr = D->getAnyInitializer()) {
+ const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(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)
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
