[PATCH] Emit OpenCL local global variables without zeorinitializer

[Matt Arsenault]

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

http://reviews.llvm.org/D6077

Files:
  lib/CodeGen/CGDecl.cpp
  test/CodeGenOpenCL/local-initializer-undef.cl
  test/CodeGenOpenCL/local.cl

Index: lib/CodeGen/CGDecl.cpp
===================================================================
--- lib/CodeGen/CGDecl.cpp
+++ lib/CodeGen/CGDecl.cpp
@@ -189,10 +189,18 @@
   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());
Index: test/CodeGenOpenCL/local-initializer-undef.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/local-initializer-undef.cl
@@ -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;
+}
Index: test/CodeGenOpenCL/local.cl
===================================================================
--- test/CodeGenOpenCL/local.cl
+++ test/CodeGenOpenCL/local.cl
@@ -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

_______________________________________________
cfe-commits mailing list
cfe-commits@cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits