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

[Zhigang Gong]

Pushed, thanks.

On Fri, Oct 18, 2013 at 03:11:29PM +0800, Ruiling Song wrote:
> 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