LGTM, thanks. -----Original Message----- From: [email protected] [mailto:[email protected]] On Behalf Of Ruiling Song Sent: Friday, October 18, 2013 3:11 PM To: [email protected] Cc: Song, Ruiling Subject: [Beignet] [PATCH] GBE: Handle all-zero constant.
Also refine Undef value support. Signed-off-by: Ruiling Song <[email protected]> --- 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 [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
