r259592 - [CUDA] Do not allow dynamic initialization of global device side variables.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 2 14:29:49 PST 2016
Author: tra
Date: Tue Feb 2 16:29:48 2016
New Revision: 259592
URL: http://llvm.org/viewvc/llvm-project?rev=259592&view=rev
Log:
[CUDA] Do not allow dynamic initialization of global device side variables.
In general CUDA does not allow dynamic initialization of
global device-side variables. One exception is that CUDA allows
records with empty constructors as described in section E2.2.1 of
CUDA 7.5 Programming guide.
This patch applies initializer checks for all device-side variables.
Empty constructors are accepted, but no code is generated for them.
Differential Revision: http://reviews.llvm.org/D15305
Added:
cfe/trunk/test/CodeGenCUDA/device-var-init.cu
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/CodeGen/CGDeclCXX.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=259592&r1=259591&r2=259592&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Feb 2 16:29:48 2016
@@ -6436,6 +6436,11 @@ def err_variadic_device_fn : Error<
def err_va_arg_in_device : Error<
"CUDA device code does not support va_arg">;
def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">;
+def err_dynamic_var_init : Error<
+ "dynamic initialization is not supported for "
+ "__device__, __constant__, and __shared__ variables.">;
+def err_shared_var_init : Error<
+ "initialization is not supported for __shared__ variables.">;
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=259592&r1=259591&r2=259592&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Tue Feb 2 16:29:48 2016
@@ -8834,6 +8834,10 @@ public:
bool ConstRHS,
bool Diagnose);
+ /// \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);
+
/// \name Code completion
//@{
/// \brief Describes the context in which code completion occurs.
Modified: cfe/trunk/lib/CodeGen/CGDeclCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDeclCXX.cpp?rev=259592&r1=259591&r2=259592&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDeclCXX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDeclCXX.cpp Tue Feb 2 16:29:48 2016
@@ -304,6 +304,17 @@ 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. All
+ // the checks have been done in Sema by now. Whatever initializers
+ // are allowed are empty and we just need to ignore them here.
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+ (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+ D->hasAttr<CUDASharedAttr>()))
+ return;
+
// Check if we've already initialized this decl.
auto I = DelayedCXXInitPosition.find(D);
if (I != DelayedCXXInitPosition.end() && I->second == ~0U)
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=259592&r1=259591&r2=259592&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Tue Feb 2 16:29:48 2016
@@ -2334,18 +2334,13 @@ void CodeGenModule::EmitGlobalVarDefinit
const VarDecl *InitDecl;
const Expr *InitExpr = D->getAnyInitializer(InitDecl);
- // CUDA E.2.4.1 "__shared__ variables cannot have an initialization as part
- // of their declaration."
- if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice
- && 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.");
- }
+ // CUDA E.2.4.1 "__shared__ variables cannot have an initialization
+ // as part of their declaration." Sema has already checked for
+ // error cases, so we just need to set Init to UndefValue.
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+ D->hasAttr<CUDASharedAttr>())
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
- } else if (!InitExpr) {
+ else if (!InitExpr) {
// This is a tentative definition; tentative definitions are
// implicitly initialized with { 0 }.
//
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=259592&r1=259591&r2=259592&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Tue Feb 2 16:29:48 2016
@@ -14,6 +14,7 @@
#include "clang/Sema/Sema.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/Decl.h"
+#include "clang/AST/ExprCXX.h"
#include "clang/Lex/Preprocessor.h"
#include "clang/Sema/SemaDiagnostic.h"
#include "llvm/ADT/Optional.h"
@@ -419,3 +420,37 @@ bool Sema::inferCUDATargetForImplicitSpe
return false;
}
+
+bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
+ if (!CD->isDefined() && CD->isTemplateInstantiation())
+ InstantiateFunctionDefinition(Loc, CD->getFirstDecl());
+
+ // (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.
+ // The constructor function has no parameters,
+ // and the function body is an empty compound statement.
+ if (!(CD->hasTrivialBody() && CD->getNumParams() == 0))
+ return false;
+
+ // Its class has no virtual functions and no virtual base classes.
+ if (CD->getParent()->isDynamicClass())
+ return false;
+
+ // The only form of initializer allowed is an empty constructor.
+ // This will recursively checks all base classes and member initializers
+ if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
+ if (const CXXConstructExpr *CE =
+ dyn_cast<CXXConstructExpr>(CI->getInit()))
+ return isEmptyCudaConstructor(Loc, CE->getConstructor());
+ return false;
+ }))
+ 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=259592&r1=259591&r2=259592&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Feb 2 16:29:48 2016
@@ -10223,6 +10223,38 @@ Sema::FinalizeDeclaration(Decl *ThisDecl
}
}
+ // Perform check for initializers of device-side global variables.
+ // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
+ // 7.5). CUDA also allows constant initializers for __constant__ and
+ // __device__ variables.
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
+ const Expr *Init = VD->getInit();
+ const bool IsGlobal = VD->hasGlobalStorage() && !VD->isStaticLocal();
+ if (Init && IsGlobal &&
+ (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
+ VD->hasAttr<CUDASharedAttr>())) {
+ bool AllowedInit = false;
+ if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
+ AllowedInit =
+ isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+ // We'll allow constant initializers even if it's a non-empty
+ // constructor according to CUDA rules. This deviates from NVCC,
+ // but allows us to handle things like constexpr constructors.
+ if (!AllowedInit &&
+ (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+ AllowedInit = VD->getInit()->isConstantInitializer(
+ Context, VD->getType()->isReferenceType());
+
+ if (!AllowedInit) {
+ Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
+ ? diag::err_shared_var_init
+ : diag::err_dynamic_var_init)
+ << Init->getSourceRange();
+ VD->setInvalidDecl();
+ }
+ }
+ }
+
// Grab the dllimport or dllexport attribute off of the VarDecl.
const InheritableAttr *DLLAttr = getDLLAttr(VD);
Added: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/device-var-init.cu?rev=259592&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Tue Feb 2 16:29:48 2016
@@ -0,0 +1,393 @@
+// 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 -std=c++11 \
+// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
+// RUN: -emit-llvm -DERROR_CASE -verify -o /dev/null %s
+
+#ifdef __clang__
+#include "Inputs/cuda.h"
+#endif
+
+// Base classes with different initializer variants.
+
+// trivial constructor -- allowed
+struct T {
+ int t;
+};
+
+// empty constructor
+struct EC {
+ int ec;
+ __device__ EC() {} // -- allowed
+ __device__ EC(int) {} // -- not allowed
+};
+
+// empty templated constructor -- allowed with no arguments
+struct ETC {
+ template <typename... T> __device__ ETC(T...) {}
+};
+
+// undefined constructor -- not allowed
+struct UC {
+ int uc;
+ __device__ UC();
+};
+
+// empty constructor w/ initializer list -- not allowed
+struct ECI {
+ int eci;
+ __device__ ECI() : eci(1) {}
+};
+
+// non-empty constructor -- not allowed
+struct NEC {
+ int nec;
+ __device__ NEC() { nec = 1; }
+};
+
+// no-constructor, virtual method -- not allowed
+struct NCV {
+ int ncv;
+ __device__ virtual void vm() {}
+};
+
+// dynamic in-class field initializer -- not allowed
+__device__ int f();
+struct NCF {
+ int ncf = f();
+};
+
+// static in-class field initializer. NVCC does not allow it, but
+// clang generates static initializer for this, so we'll accept it.
+struct NCFS {
+ int ncfs = 3;
+};
+
+// undefined templated constructor -- not allowed
+struct UTC {
+ template <typename... T> __device__ UTC(T...);
+};
+
+// non-empty templated constructor -- not allowed
+struct NETC {
+ int netc;
+ template <typename... T> __device__ NETC(T...) { netc = 1; }
+};
+
+__device__ int d_v;
+// CHECK: @d_v = addrspace(1) externally_initialized global i32 0,
+__shared__ int s_v;
+// CHECK: @s_v = addrspace(3) global i32 undef,
+__constant__ int c_v;
+// CHECK: addrspace(4) externally_initialized global i32 0,
+
+__device__ int d_v_i = 1;
+// CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1,
+#ifdef ERROR_CASE
+__shared__ int s_v_i = 1;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+#endif
+__constant__ int c_v_i = 1;
+// CHECK: @c_v_i = addrspace(4) externally_initialized global i32 1,
+
+#ifdef ERROR_CASE
+__device__ int d_v_f = f();
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ int s_v_f = f();
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ int c_v_f = f();
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+__device__ T d_t;
+// CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
+__shared__ T s_t;
+// CHECK: @s_t = addrspace(3) global %struct.T undef,
+__constant__ T c_t;
+// CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
+
+__device__ T d_t_i = {2};
+// CHECKL @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 },
+#ifdef ERROR_CASE
+__shared__ T s_t_i = {2};
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+#endif
+__constant__ T c_t_i = {2};
+// CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
+
+__device__ EC d_ec;
+// CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
+__shared__ EC s_ec;
+// CHECK: @s_ec = addrspace(3) global %struct.EC undef,
+__constant__ EC c_ec;
+// CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+
+#if ERROR_CASE
+__device__ EC d_ec_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ EC s_ec_i(3);
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ EC c_ec_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ EC d_ec_i2 = {3};
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ EC s_ec_i2 = {3};
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ EC c_ec_i2 = {3};
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+__device__ ETC d_etc;
+// CHETCK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
+__shared__ ETC s_etc;
+// CHETCK: @s_etc = addrspace(3) global %struct.ETC undef,
+__constant__ ETC c_etc;
+// CHETCK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
+
+#if ERROR_CASE
+__device__ ETC d_etc_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ ETC s_etc_i(3);
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ ETC c_etc_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ ETC d_etc_i2 = {3};
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ ETC s_etc_i2 = {3};
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ ETC c_etc_i2 = {3};
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ UC d_uc;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UC s_uc;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UC c_uc;
+// 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;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ ECI c_eci;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ NEC d_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NEC s_nec;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NEC c_nec;
+// 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;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NCV c_ncv;
+// 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;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NCF c_ncf;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+__device__ NCFS d_ncfs;
+// CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 }
+__constant__ NCFS c_ncfs;
+// CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
+
+#if ERROR_CASE
+__shared__ NCFS s_ncfs;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+
+__device__ UTC d_utc;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UTC s_utc;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UTC c_utc;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ UTC d_utc_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UTC s_utc_i(3);
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UTC c_utc_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ NETC d_netc;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NETC s_netc;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NETC c_netc;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ NETC d_netc_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NETC s_netc_i(3);
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NETC c_netc_i(3);
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+// Regular base class -- allowed
+struct T_B_T : T {};
+__device__ T_B_T d_t_b_t;
+// CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
+__shared__ T_B_T s_t_b_t;
+// CHECK: @s_t_b_t = addrspace(3) global %struct.T_B_T undef,
+__constant__ T_B_T c_t_b_t;
+// CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
+
+// Incapsulated object of allowed class -- allowed
+struct T_F_T {
+ T t;
+};
+__device__ T_F_T d_t_f_t;
+// CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
+__shared__ T_F_T s_t_f_t;
+// CHECK: @s_t_f_t = addrspace(3) global %struct.T_F_T undef,
+__constant__ T_F_T c_t_f_t;
+// CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
+
+// array of allowed objects -- allowed
+struct T_FA_T {
+ T t[2];
+};
+__device__ T_FA_T d_t_fa_t;
+// CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
+__shared__ T_FA_T s_t_fa_t;
+// CHECK: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef,
+__constant__ T_FA_T c_t_fa_t;
+// CHECK: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer,
+
+
+// Calling empty base class initializer is OK
+struct EC_I_EC : EC {
+ __device__ EC_I_EC() : EC() {}
+};
+__device__ EC_I_EC d_ec_i_ec;
+// CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
+__shared__ EC_I_EC s_ec_i_ec;
+// CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
+__constant__ EC_I_EC c_ec_i_ec;
+// CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
+
+// .. though passing arguments is not allowed.
+struct EC_I_EC1 : EC {
+ __device__ EC_I_EC1() : EC(1) {}
+};
+#if ERROR_CASE
+__device__ EC_I_EC1 d_ec_i_ec1;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ EC_I_EC1 s_ec_i_ec1;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ EC_I_EC1 c_ec_i_ec1;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+// Virtual base class -- not allowed
+struct T_V_T : virtual T {};
+#if ERROR_CASE
+__device__ T_V_T d_t_v_t;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_V_T s_t_v_t;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_V_T c_t_v_t;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+// Make sure that we don't allow if we inherit or incapsulate
+// something with disallowed initializer.
+
+// Inherited from or incapsulated class with non-empty constructor --
+// not allowed
+struct T_B_NEC : NEC {};
+struct T_F_NEC {
+ NEC nec;
+};
+struct T_FA_NEC {
+ NEC nec[2];
+};
+
+#if ERROR_CASE
+__device__ T_B_NEC d_t_b_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_B_NEC s_t_b_nec;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_B_NEC c_t_b_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_F_NEC d_t_f_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_F_NEC s_t_f_nec;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_F_NEC c_t_f_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_FA_NEC d_t_fa_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_FA_NEC s_t_fa_nec;
+// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_FA_NEC c_t_fa_nec;
+// expected-error at -1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+#endif
+
+// 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() {
+ T t;
+ EC ec;
+ ETC etc;
+ UC uc;
+ ECI eci;
+ NEC nec;
+ NCV ncv;
+ NCF ncf;
+ NCFS ncfs;
+ UTC utc;
+ NETC netc;
+ T_B_T t_b_t;
+ T_F_T t_f_t;
+ T_FA_T t_fa_t;
+ EC_I_EC ec_i_ec;
+ EC_I_EC1 ec_i_ec1;
+ T_V_T t_v_t;
+ T_B_NEC t_b_nec;
+ T_F_NEC t_f_nec;
+ T_FA_NEC t_fa_nec;
+ static __shared__ UC s_uc;
+}
+
+// CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec)
+// CHECK: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
+// CHECK: call void @_ZN2UCC1Ev(%struct.UC* %uc)
+// CHECK: call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
+// CHECK: call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+// CHECK: call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
+// CHECK: call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+// CHECK: call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+// CHECK: call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
+// CHECK: call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
+// CHECK: call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
+// CHECK: call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
+// CHECK: call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t) #3
+// CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
+// CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
+// CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+// CHECK: call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*))
+// CHECK: ret void
+
+// We should not emit global init function.
+// CHECK-NOT: @_GLOBAL__sub_I
More information about the cfe-commits
mailing list