[CUDA] Do not allow dynamic initialization of global device side variables.

[tra]

Summary

In general CUDA does not allow dynamic initialization of
global device-side variables except for records with empty constructors as described in section 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.

[tra]

lib/CodeGen/CGDeclCXX.cpp

329

@rsmith: is this a good way to find member initializer list items?

struct S {
    int a,b,c;
    S() : a(1),b(2),c(3) {}
};

I'm looking for a(),b(),c() which is what I think CUDA spec wants to check for, but CD->inits() appears to have other initializers on the list as well.

[tra]

Ping.
@rsmith -- Richard, can you take a look?

[tra]

ping.

[rsmith]

This should be checked and diagnosed in Sema, not in CodeGen.

lib/CodeGen/CGDeclCXX.cpp

323–324

What if the constructor is a C-style varargs function:

struct X { X(...) {} };

?

329

You shouldn't need to check isAnyMemberInitializer: if there's any written inits, the constructor violates the rules.

333

What doesn't always work?

347–367

Rather than checking these properties this way, I'd suggest you check the initialization expression in each CXXCtorInitializer only contains CXXConstructExprs for empty constructors (or any other whitelisted constructs). Your current approach will miss a couple of cases which the CUDA spec misses but presumably meant to exclude:

1. Default member initializers
   
   int f(); struct X { int n = f(); X() {} };

2. Cases where a constructor other than a default constructor is implicitly invoked
   
   struct A { template<typename ...T> A(T...); }; struct B : A { B() {} };

lib/CodeGen/CodeGenModule.cpp

1424–1428

According to the quoted specification, you're supposed to check whether the constructor can be considered empty at the point in the translation unit where the definition of the variable occurs, so I don't think you need to delay anything.

[tra]

Check all variable initializers and only allow 'empty constructors' as Richard has suggested.
Changed test structure so that we test for allowed/disallowed constructors separately from testing how we handle initialization of base classes or member fields.

[tra]

Richard, I've updated the patch as you've suggested -- it indeed simplifies things quite a bit and handles the corner cases you've mentioned.

lib/CodeGen/CGDeclCXX.cpp

323–324

CUDA does not support varargs on device side. nvcc fails with an error:

error: a "device" function cannot have ellipsis

That's another thing I'll need to fix (as a separate patch) as clang currently accepts varargs everywhere.

This patch will ignore number of arguments passed to varargs constructor, but the checks for empty body still do apply.

329

As it turns out, the rules don't apply to all written initializers. For instance, nvcc allows empty constructor on init list:

struct A {  __device__ A(){}; };
struct B {  __device__ B(){}; };

struct C : A {
  B b;
  __device__ C() : A(), b() {}
};

__device__ C c;

I've simplified the patch so that in only checks for constructor's 'emptiness', but disregards how that constructor gets to be executed.

333

It was leftover from early patch variant that didn't defer emitting global vars.
If I don't defer and need to emit a global var before constructor definition is available, hasTrivialBody() returns false and triggers diagnostics.

347–367

Nice. This has simplified the checks a lot.

lib/CodeGen/CodeGenModule.cpp

1424–1428

I guess it's a bug in their guide as nvcc accepts following code with constructor definition appearing *after* the variable:

struct S {  S(); };
__device__ S s;
S::S() {}

[rsmith]

I think you missed this from my previous review:

This should be checked and diagnosed in Sema, not in CodeGen.

lib/CodeGen/CGDeclCXX.cpp

333–337

You can check these conditions with RD->isDynamicClass().

[tra]

Moved initializer checks from CodeGen to Sema.
Added test cases for initializers of non-class variables.

[tra]

In D15305#327226, @rsmith wrote:

I think you missed this from my previous review:

This should be checked and diagnosed in Sema, not in CodeGen.

Done.

[rsmith]

lib/CodeGen/CGDeclCXX.cpp

312

areallowed -> are allowed

lib/CodeGen/CodeGenModule.cpp

2334

As this is a global variable, it should presumably still be statically zero-initialized.

[tra]

Typo fix.

lib/CodeGen/CodeGenModule.cpp

2334

There is no way to initialize shared variables. They are rough equivalent of local variables, only in this case CUDA allocates them per kernel invocation from a shared buffer with no guarantees regarding its contents.

They used to be zero-initialized by compiler, but that was intentionally changed to undef in r245786 / http://reviews.llvm.org/D12241

[jlebar]

tra asked me to check for coverage. Looks pretty good in that respect.

include/clang/Basic/DiagnosticSemaKinds.td

6419

Nit, but, since we're all language nerds here, suggest adding an Oxford comma.

lib/Sema/SemaCUDA.cpp

436

The test passes if I comment out this if statement. I'm not sure if that's expected; this may or may not be entirely covered below.

442

Tests pass if I comment out the isDefined check.

lib/Sema/SemaDecl.cpp

10183

We also allow constant initializers for constant and device variables.

Consider rephrasing this -- it sounds like this is a clang extension, but I just checked and it does not appear to be.

10186

Test passes if I comment out IsGlobal or CUDAIsDevice. (I'm not sure if you care to test the latter, but the former seems important.)

[tra]

Addressed Justin's comments.

[tra]

lib/Sema/SemaCUDA.cpp

436

According to CPP reference trivial constructor will pass all other checks below.

442

hasTrivialBody() would only return true if we have a body which only happens if function is defined. isDefined() is mostly for readability here.

lib/Sema/SemaDecl.cpp

10186

IsGlobal -- all test cases were using either global or local variables. I've added a static shared variable in the device function. Now IsGlobal check (or, rather !isStaticLocal() part of it) is required in order for the tests to succeed.

CUDAIsDevice is not triggered because all test cases are run with -fcuda-is-device.
It's hard to run host-side test with -verify here because I'd have to put #ifdef around every 'expected-error'

[jlebar]

jingyue/jpienaar/rsmith - friendly ping? Without this, -O0 builds don't work, because they emit empty global initializers that don't get optimized out.

[jpienaar]

@jlebar: We defer it to your and Richard's approval. Thanks

[rsmith]

Some minor things, but feel free to commit after addressing them.

I agree that we should figure out what to do about the zero/undef initialization separately.

lib/Sema/SemaCUDA.cpp

429–430

The function might still not be defined after this (if the template is not defined); you should presumably return false here in that case.

442

Please do remove the isDefined check here. Including it makes a reader wonder what case it's trying to handle.

455–457

Maybe reorder this before the CXXCtorInitializer check? It's a much cheaper test.

lib/Sema/SemaDecl.cpp

10191–10198

What should happen if the init is a constant initializer that is a CXXConstructExpr, but it uses a constructor that is not empty from CUDA's perspective? Such as:

struct X { constexpr X() { int n = 0; } };
__device__ X x;

I would assume this should be valid, but I think you'll reject it. Maybe change else if ( to if (!AllowedInit &&?

10196–10198

Might be clearer as

if (__device__ || __constant__)
  AllowedInit = isConstantInitializer(...)

[tra]

Addressed Richard's comments.
Relaxed restrictions a bit to allow constant initializers even those CUDA would not considered to be empty.
Updated test case accordingly.

[tra]

lib/Sema/SemaCUDA.cpp

429–430

I don't think it's needed. If it's still not definied, it will be caught by hasTrivialBody() check below.

lib/Sema/SemaDecl.cpp

10191–10198

NVCC produces an error (probably because it does not support c++14):
zz.cu(1): error: statement may not appear in a constexpr constructor

clang w/ this patch indeed considers it to be a non-empty initializer and produces an error.

I agree that allowing constant initializer is the right thing to do. Your example requires c++14, so there's no direct comparison with nvcc, but I think allowing it is indeed the right thing to do here.

include/clang/Basic/DiagnosticSemaKinds.td

 def err_global_call_not_config : Error<
   "call to global function %0 not configured">;
 def err_ref_bad_target : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
 def warn_host_calls_from_host_device : Warning<
   "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">,
   InGroup<CudaCompat>;
+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 "
   "%select{function|block|method|constructor}2; expected type from format "
   "string was %3">, InGroup<NonPODVarargs>, DefaultError;
 // The arguments to this diagnostic should match the warning above.
 def err_cannot_pass_objc_interface_to_vararg_format : Error<
   "cannot pass object with interface type %1 by value to variadic "

include/clang/Sema/Sema.h

   /// The result of this call is implicit CUDA target attribute(s) attached to
   /// the member declaration.
   bool inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
                                                CXXSpecialMember CSM,
                                                CXXMethodDecl *MemberDecl,
                                                bool ConstRHS,
                                                bool Diagnose);
 
+  bool isEmptyCudaConstructor(SourceLocation VarLoc,
+                              CXXConstructorDecl *CD);
   /// \name Code completion
   //@{
   /// \brief Describes the context in which code completion occurs.
   enum ParserCompletionContext {
     /// \brief Code completion occurs at top-level or namespace context.
     PCC_Namespace,
     /// \brief Code completion occurs within a class, struct, or union.
     PCC_Class,

lib/CodeGen/CGDeclCXX.cpp

void CodeGenModule::EmitPointerToInitFunc(const VarDecl *D,
   if (llvm::Comdat *C = GV->getComdat())
     PtrArray->setComdat(C);
 }
 
 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)
     return;
 
   llvm::FunctionType *FTy = llvm::FunctionType::get(VoidTy, false);
   SmallString<256> FnName;
   {
     llvm::raw_svector_ostream Out(FnName);
     getCXXABI().getMangleContext().mangleDynamicInitializer(D, Out);
   }
 
   // Create a variable initialization function.
   llvm::Function *Fn =
       CreateGlobalInitOrDestructFunction(FTy, FnName.str(),
                                          getTypes().arrangeNullaryFunction(),
                                          D->getLocation());
 
   auto *ISA = D->getAttr<InitSegAttr>();
   CodeGenFunction(*this).GenerateCXXGlobalVarDeclInitFunc(Fn, D, Addr,
                                                           PerformInit);
 
   llvm::GlobalVariable *COMDATKey =
       supportsCOMDAT() && D->isExternallyVisible() ? Addr : nullptr;
 
   if (D->getTLSKind()) {
     // FIXME: Should we support init_priority for thread_local?
     // FIXME: Ideally, initialization of instantiated thread_local static data
     // members of class templates should not trigger initialization of other
     // entities in the TU.
     // FIXME: We only need to register one __cxa_thread_atexit function for the
     // entire TU.
     CXXThreadLocalInits.push_back(Fn);
     CXXThreadLocalInitVars.push_back(D);
   } else if (PerformInit && ISA) {
     EmitPointerToInitFunc(D, Addr, Fn, ISA);
   } else if (auto *IPA = D->getAttr<InitPriorityAttr>()) {
     OrderGlobalInits Key(IPA->getPriority(), PrioritizedCXXGlobalInits.size());
     PrioritizedCXXGlobalInits.push_back(std::make_pair(Key, Fn));
   } else if (isTemplateInstantiation(D->getTemplateSpecializationKind())) {
     // C++ [basic.start.init]p2:
     //   Definitions of explicitly specialized class template static data
     //   members have ordered initialization. Other class template static data
     //   members (i.e., implicitly or explicitly instantiated specializations)
     //   have unordered initialization.
     //
     // As a consequence, we can put them into their own llvm.global_ctors entry.
     //
     // If the global is externally visible, put the initializer into a COMDAT
     // group with the global being initialized.  On most platforms, this is a
     // minor startup time optimization.  In the MS C++ ABI, there are no guard
     // variables, so this COMDAT key is required for correctness.
     AddGlobalCtor(Fn, 65535, COMDATKey);
   } else if (D->hasAttr<SelectAnyAttr>()) {
     // SelectAny globals will be comdat-folded. Put the initializer into a
     // COMDAT group associated with the global, so the initializers get folded
     // too.
     AddGlobalCtor(Fn, 65535, COMDATKey);

lib/CodeGen/CodeGenModule.cpp

if (FD->getTemplateSpecializationKind() == TSK_ImplicitInstantiation)
       // Implicit template instantiations may change linkage if they are later
       // explicitly instantiated, so they should not be emitted eagerly.
       return false;
   // If OpenMP is enabled and threadprivates must be generated like TLS, delay
   // codegen for global variables, because they may be marked as threadprivate.
   if (LangOpts.OpenMP && LangOpts.OpenMPUseTLS &&
       getContext().getTargetInfo().isTLSSupported() && isa<VarDecl>(Global))
     return false;
 
   return true;
 }
 
 ConstantAddress CodeGenModule::GetAddrOfUuidDescriptor(
     const CXXUuidofExpr* E) {
   // Sema has verified that IIDSource has a __declspec(uuid()), and that its
   // well-formed.
   StringRef Uuid = E->getUuidAsStringRef(Context);
   std::string Name = "_GUID_" + Uuid.lower();
   std::replace(Name.begin(), Name.end(), '-', '_');
 
   // Contains a 32-bit field.
void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   QualType ASTTy = D->getType();
   CXXRecordDecl *RD = ASTTy->getBaseElementTypeUnsafe()->getAsCXXRecordDecl();
   bool NeedsGlobalCtor = false;
   bool NeedsGlobalDtor = RD && !RD->hasTrivialDestructor();
 
   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 }.
     //
     // Note that tentative definitions are only emitted at the end of
     // a translation unit, so they should never have incomplete
     // type. In addition, EmitTentativeDefinition makes sure that we
     // never attempt to emit a tentative definition if a real one
     // exists. A use may still exists, however, so we still may need

lib/Sema/SemaCUDA.cpp

 //===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 /// \file
 /// \brief This file implements semantic analysis for CUDA constructs.
 ///
 //===----------------------------------------------------------------------===//
 
 #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"
 #include "llvm/ADT/SmallVector.h"
 using namespace clang;
 
 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                          MultiExprArg ExecConfig,
if (InferredTarget.hasValue()) {
     // If no target was inferred, mark this member as __host__ __device__;
     // it's the least restrictive option that can be invoked from any target.
     MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
     MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
   }
 
   return false;
 }
+
+bool Sema::isEmptyCudaConstructor(SourceLocation VarLoc,
+                                  CXXConstructorDecl *CD) {
+  if (!CD->isDefined() && CD->isTemplateInstantiation())
+    InstantiateFunctionDefinition(VarLoc, 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->isDefined() && CD->getNumParams() == 0 && CD->hasTrivialBody()))
+    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(VarLoc, CE->getConstructor());
+        return false;
+      }))
+    return false;
+
+  // Its class has no virtual functions and no virtual base classes.
+  if (CD->getParent()->isDynamicClass())
+    return false;
+
+  return true;
+}

lib/Sema/SemaDecl.cpp

if (FunctionDecl *FD =
       if (Attr *A = getDLLAttr(FD)) {
         auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext()));
         NewAttr->setInherited(true);
         VD->addAttr(NewAttr);
       }
     }
   }
 
+  // Perform check for initializers of device-side global variables.
+  // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
+  // 7.5).  We also allow constant initializers for __constant__ and
+  // __device__ variables.
+  const Expr *Init = VD->getInit();
+  const bool IsGlobal = VD->hasGlobalStorage() && !VD->isStaticLocal();
+  if (Init && IsGlobal && getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+      (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());
+    else if ((VD->hasAttr<CUDADeviceAttr>() ||
+              VD->hasAttr<CUDAConstantAttr>()) &&
+             VD->getInit()->isConstantInitializer(
+                 Context, VD->getType()->isReferenceType()))
+      AllowedInit = true;
+
+    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);
 
   // Imported static data members cannot be defined out-of-line.
   if (const auto *IA = dyn_cast_or_null<DLLImportAttr>(DLLAttr)) {
     if (VD->isStaticDataMember() && VD->isOutOfLine() &&
         VD->isThisDeclarationADefinition()) {
       // We allow definitions of dllimport class template static data members

test/CodeGenCUDA/device-var-init.cu

+// 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:     -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.
+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.}}
+
+__device__ NCFS d_ncfs;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __shared__ variables.}}
+__shared__ NCFS s_ncfs;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NCFS c_ncfs;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__ and __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;
+}
+
+// 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) #3
+// CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+// CHECK:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs) #3
+// 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: ret void
+
+// We should not emit global init function.
+// CHECK-NOT: @_GLOBAL__sub_I