[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

llvm-svn: 259592

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 2f059c6..f176344 100644 (file)
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -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 "

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 745dca8..75b3e0a 100644 (file)
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -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.

diff --git a/clang/lib/CodeGen/CGDeclCXX.cpp b/clang/lib/CodeGen/CGDeclCXX.cpp
index adba731..98d904b 100644 (file)
--- a/clang/lib/CodeGen/CGDeclCXX.cpp
+++ b/clang/lib/CodeGen/CGDeclCXX.cpp
@@ -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)

diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index de580f2..4d7f627 100644 (file)
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2334,18 +2334,13 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   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 }.
     //

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 568c765..84fccd5 100644 (file)
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -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::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
 
   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;
+}

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 5db2c37..ffd71e3 100644 (file)
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -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);
 

diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu
new file mode 100644 (file)
index 0000000..774c590
--- /dev/null
+++ b/clang/test/CodeGenCUDA/device-var-init.cu
@@ -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@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ int s_v_f = f();
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ int c_v_f = f();
+// expected-error@-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@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ EC s_ec_i(3);
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ EC c_ec_i(3);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ EC d_ec_i2 = {3};
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ EC s_ec_i2 = {3};
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ EC c_ec_i2 = {3};
+// expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ ETC s_etc_i(3);
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ ETC c_etc_i(3);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ ETC d_etc_i2 = {3};
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ ETC s_etc_i2 = {3};
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ ETC c_etc_i2 = {3};
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ UC d_uc;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UC s_uc;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UC c_uc;
+// 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;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ ECI c_eci;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ NEC d_nec;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NEC s_nec;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NEC c_nec;
+// 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;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NCV c_ncv;
+// 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;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NCF c_ncf;
+// expected-error@-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@-1 {{initialization is not supported for __shared__ variables.}}
+
+__device__ UTC d_utc;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UTC s_utc;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UTC c_utc;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ UTC d_utc_i(3);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UTC s_utc_i(3);
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UTC c_utc_i(3);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ NETC d_netc;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NETC s_netc;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NETC c_netc;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ NETC d_netc_i(3);
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NETC s_netc_i(3);
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NETC c_netc_i(3);
+// expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ EC_I_EC1 s_ec_i_ec1;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ EC_I_EC1 c_ec_i_ec1;
+// expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_V_T s_t_v_t;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_V_T c_t_v_t;
+// expected-error@-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@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_B_NEC s_t_b_nec;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_B_NEC c_t_b_nec;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_F_NEC d_t_f_nec;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_F_NEC s_t_f_nec;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_F_NEC c_t_f_nec;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_FA_NEC d_t_fa_nec;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_FA_NEC s_t_fa_nec;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_FA_NEC c_t_fa_nec;
+// expected-error@-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