This patch looks good to me. -----Original Message----- From: [email protected] [mailto:[email protected]] On Behalf Of Ruiling Song Sent: Thursday, November 7, 2013 3:13 PM To: [email protected] Cc: Song, Ruiling Subject: [Beignet] [PATCH V2] GBE: fix a 64bit scalar register issue.
For scalar register, should use stride 0. also change the unit test to hit the point. v2: fix h2() Signed-off-by: Ruiling Song <[email protected]> --- backend/src/backend/gen_register.hpp | 6 +++--- kernels/compiler_long.cl | 7 ++++--- utests/compiler_long.cpp | 2 ++ 3 files changed, 9 insertions(+), 6 deletions(-) diff --git a/backend/src/backend/gen_register.hpp b/backend/src/backend/gen_register.hpp index 66bc288..47b0515 100644 --- a/backend/src/backend/gen_register.hpp +++ b/backend/src/backend/gen_register.hpp @@ -275,9 +275,8 @@ namespace gbe INLINE GenRegister bottom_half(void) const { GBE_ASSERT(isint64()); - GenRegister r = *this; + GenRegister r = h2(*this); r.type = type == GEN_TYPE_UL ? GEN_TYPE_UD : GEN_TYPE_D; - r.hstride = GEN_HORIZONTAL_STRIDE_2; r.vstride = GEN_VERTICAL_STRIDE_16; return r; } @@ -304,7 +303,8 @@ namespace gbe static INLINE GenRegister h2(GenRegister reg) { GenRegister r = reg; - r.hstride = GEN_HORIZONTAL_STRIDE_2; + if(r.hstride != GEN_HORIZONTAL_STRIDE_0) + r.hstride = GEN_HORIZONTAL_STRIDE_2; return r; } diff --git a/kernels/compiler_long.cl b/kernels/compiler_long.cl index 3087292..e69c5bf 100644 --- a/kernels/compiler_long.cl +++ b/kernels/compiler_long.cl @@ -1,7 +1,8 @@ -kernel void compiler_long(global long *src1, global long *src2, global long *dst) { +kernel void compiler_long(global long *src1, global long *src2, global +long *dst, long zero) { int i = get_global_id(0); + if(i < 5) - dst[i] = src1[i] + src2[i]; + dst[i] = src1[i] + src2[i] + src2[i]*zero; if(i > 5) - dst[i] = src1[i] - src2[i]; + dst[i] = src1[i] - src2[i] - zero; } diff --git a/utests/compiler_long.cpp b/utests/compiler_long.cpp index d7e1517..b525694 100644 --- a/utests/compiler_long.cpp +++ b/utests/compiler_long.cpp @@ -8,6 +8,7 @@ void compiler_long(void) const size_t n = 16; int64_t src1[n], src2[n]; + int64_t zero = 0; // Setup kernel and buffers OCL_CREATE_KERNEL("compiler_long"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int64_t), NULL); @@ -16,6 +17,7 @@ void compiler_long(void) OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]); + OCL_SET_ARG(3, sizeof(cl_long), &zero); globals[0] = n; locals[0] = 16; -- 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
