Emit OpenCL local global variables without zeorinitializer

[arsenm]

Summary

Local variables are not initialized, and every target has
been (incorrectly) ignoring the unnecessary request for
zero initialization.

[joey]

LGTM! (Apart from those tiny changes)

lib/CodeGen/CGDecl.cpp

196

Use 'Ty' (initialised on line 179)

198

Use 'LTy' (initialised on line 189)

[arsenm]

r221162

lib/CodeGen/CGDecl.cpp

llvm::Constant *CodeGenModule::getOrCreateStaticVarDecl(
   if (D.hasAttr<AsmLabelAttr>())
     Name = getMangledName(&D);
   else
     Name = getStaticDeclName(*this, D);
 
   llvm::Type *LTy = getTypes().ConvertTypeForMem(Ty);
   unsigned AddrSpace =
       GetGlobalVarAddressSpace(&D, getContext().getTargetAddressSpace(Ty));
+
+  // Local address space cannot have an initializer.
+  llvm::Constant *Init = nullptr;
+  if (Ty.getAddressSpace() != LangAS::opencl_local)
+    Init = EmitNullConstant(D.getType());
+  else
+    Init = llvm::UndefValue::get(getTypes().ConvertTypeForMem(D.getType()));
+
   llvm::GlobalVariable *GV =
     new llvm::GlobalVariable(getModule(), LTy,
                              Ty.isConstant(getContext()), Linkage,
-                             EmitNullConstant(D.getType()), Name, nullptr,
+                             Init, Name, nullptr,
                              llvm::GlobalVariable::NotThreadLocal,
                              AddrSpace);
   GV->setAlignment(getContext().getDeclAlign(&D).getQuantity());
   setGlobalVisibility(GV, &D);
 
   if (D.getTLSKind())
     setTLSMode(GV, D);
 

test/CodeGenOpenCL/local-initializer-undef.cl

+// RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+typedef struct Foo {
+    int x;
+    float y;
+    float z;
+} Foo;
+
+// CHECK-DAG: @test.lds_int = internal addrspace(2) global i32 undef
+// CHECK-DAG: @test.lds_int_arr = internal addrspace(2) global [128 x i32] undef
+// CHECK-DAG: @test.lds_struct = internal addrspace(2) global %struct.Foo undef
+// CHECK-DAG: @test.lds_struct_arr = internal addrspace(2) global [64 x %struct.Foo] undef
+__kernel void test()
+{
+    __local int lds_int;
+    __local int lds_int_arr[128];
+    __local Foo lds_struct;
+    __local Foo lds_struct_arr[64];
+
+    lds_int = 1;
+    lds_int_arr[0] = 1;
+    lds_struct.x = 1;
+    lds_struct_arr[0].x = 1;
+}

test/CodeGenOpenCL/local.cl

 // RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck %s
 
+void func(local int*);
+
 __kernel void foo(void) {
-  // CHECK: @foo.i = internal unnamed_addr addrspace(2)
+  // CHECK: @foo.i = internal addrspace(2) global i32 undef
   __local int i;
-  ++i;
+  func(&i);
 }
 
 // CHECK-LABEL: define void @_Z3barPU7CLlocali
 __kernel void __attribute__((__overloadable__)) bar(local int *x) {
   *x = 5;
 }