r221162 - Emit OpenCL local global variables without zeorinitializer

[Matt Arsenault]

Author: arsenm
Date: Mon Nov  3 10:51:53 2014
New Revision: 221162

URL: http://llvm.org/viewvc/llvm-project?rev=221162&view=rev
Log:
Emit OpenCL local global variables without zeorinitializer

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

Added:
    cfe/trunk/test/CodeGenOpenCL/local-initializer-undef.cl
Modified:
    cfe/trunk/lib/CodeGen/CGDecl.cpp
    cfe/trunk/test/CodeGenOpenCL/local.cl

Modified: cfe/trunk/lib/CodeGen/CGDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDecl.cpp?rev=221162&r1=221161&r2=221162&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDecl.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDecl.cpp Mon Nov  3 10:51:53 2014
@@ -189,10 +189,18 @@ llvm::Constant *CodeGenModule::getOrCrea
   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(Ty);
+  else
+    Init = llvm::UndefValue::get(LTy);
+
   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());

Added: cfe/trunk/test/CodeGenOpenCL/local-initializer-undef.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/local-initializer-undef.cl?rev=221162&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/local-initializer-undef.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/local-initializer-undef.cl Mon Nov  3 10:51:53 2014
@@ -0,0 +1,24 @@
+// 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;
+}

Modified: cfe/trunk/test/CodeGenOpenCL/local.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/local.cl?rev=221162&r1=221161&r2=221162&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/local.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/local.cl Mon Nov  3 10:51:53 2014
@@ -1,9 +1,11 @@
 // 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