From: Ruiling Song <ruiling.song@intel.com>
Date: Fri, 18 Oct 2013 07:11:29 +0000 (+0800)
Subject: GBE: Handle all-zero constant.
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=8d0f73f2d8a9099acffb36b9c582f3ee6162aa9e;p=contrib%2Fbeignet.git

GBE: Handle all-zero constant.

Also refine Undef value support.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: "Yang, Rong R" <rong.r.yang@intel.com>
---

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index ea34675..968283a 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -577,12 +577,16 @@ namespace gbe
 
     GBE_ASSERT(c);
     if(isa<UndefValue>(c)) {
-      uint32_t n = c->getNumOperands();
-      Type * opTy = type->getArrayElementType();
-      uint32_t size = opTy->getIntegerBitWidth()/ 8;
-      offset += size*n;
+      uint32_t size = getTypeByteSize(unit, type);
+      offset += size;
+      return;
+    } else if(isa<ConstantAggregateZero>(c)) {
+      uint32_t size = getTypeByteSize(unit, type);
+      memset((char*)mem+offset, 0, size);
+      offset += size;
       return;
     }
+
     switch(id) {
       case Type::TypeID::StructTyID:
         {
diff --git a/kernels/compiler_global_constant.cl b/kernels/compiler_global_constant.cl
index 71fe86c..53e24b3 100644
--- a/kernels/compiler_global_constant.cl
+++ b/kernels/compiler_global_constant.cl
@@ -19,9 +19,16 @@ struct Test2 {
   char a0;
   int a1;
 };
+struct Test3 {
+  int a0;
+  int a1;
+};
+struct Test4 {
+  float a0;
+  float a1;
+};
 
 constant struct Person james= {{"james"}, (int3)(1, 2, 3)};
-
 constant struct Test1 t0 = {1, 2};
 constant struct Test2 t1 = {1, 2};
 
@@ -29,6 +36,10 @@ constant int3 c[3] = {(int3)(0, 1, 2), (int3)(3, 4, 5), (int3)(6,7,8) };
 constant char4 d[3] = {(char4)(0, 1, 2, 3), (char4)(4, 5, 6, 7), (char4)(8, 9, 10, 11)};
 
 constant struct Person members[3] = {{{"abc"}, (int3)(1, 2, 3)}, { {"defg"}, (int3)(4,5,6)}, { {"hijk"}, (int3)(7,8,9)} };
+constant struct Test3 zero_struct = {0, 0};
+constant int3 zero_vec = {0,0,0};
+constant int zero_arr[3] = {0,0,0};
+constant float zero_flt[3] = {0.0f, 0.0f, 0.0f};
 
 __kernel void
 compiler_global_constant(__global int *dst, int e, int r)
@@ -36,7 +47,7 @@ compiler_global_constant(__global int *dst, int e, int r)
   int id = (int)get_global_id(0);
 
   int4 x = a + b;
-  dst[id] = m[id%3] * n * o[2] + e + r *x.y * a.x;
+  dst[id] = m[id%3] * n * o[2] + e + r *x.y * a.x + zero_struct.a0 + zero_vec.x + zero_arr[1] + (int)zero_flt[2];
 }
 // array of vectors
 __kernel void