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

[Artem Belevich via cfe-commits]

tra added inline comments.

================
Comment at: lib/CodeGen/CodeGenModule.cpp:2334
@@ -2339,1 +2333,3 @@
+      D->hasAttr<CUDASharedAttr>())
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
+  else if (!InitExpr) {
----------------
rsmith wrote:
> As this is a global variable, it should presumably still be statically zero-initialized.
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


http://reviews.llvm.org/D15305