[Beignet] [PATCH] GBE: Handle all-zero constant.

[Yang, Rong R]

LGTM, thanks.

-----Original Message-----
From: beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org [mailto:beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org] On Behalf Of Ruiling Song
Sent: Friday, October 18, 2013 3:11 PM
To: beignet at lists.freedesktop.org
Cc: Song, Ruiling
Subject: [Beignet] [PATCH] GBE: Handle all-zero constant.

Also refine Undef value support.

Signed-off-by: Ruiling Song <ruiling.song at intel.com>
---
 backend/src/llvm/llvm_gen_backend.cpp |   12 ++++++++----
 kernels/compiler_global_constant.cl   |   15 +++++++++++++--
 2 files changed, 21 insertions(+), 6 deletions(-)

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
--
1.7.9.5

_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet