LGTM. BTW: Does the function GenRegister::h2 also have this issue?
-----Original Message----- From: [email protected] [mailto:[email protected]] On Behalf Of Ruiling Song Sent: Wednesday, November 06, 2013 1:40 PM To: [email protected] Cc: Song, Ruiling Subject: [Beignet] [PATCH] GBE: fix a 64bit scalar register issue. For scalar register, should use stride 0. also change the unit test to hit the point. Signed-off-by: Ruiling Song <[email protected]> --- backend/src/backend/gen_register.hpp | 3 ++- kernels/compiler_long.cl | 7 ++++--- utests/compiler_long.cpp | 2 ++ 3 files changed, 8 insertions(+), 4 deletions(-) diff --git a/backend/src/backend/gen_register.hpp b/backend/src/backend/gen_register.hpp index 66bc288..2f80c72 100644 --- a/backend/src/backend/gen_register.hpp +++ b/backend/src/backend/gen_register.hpp @@ -277,7 +277,8 @@ namespace gbe GBE_ASSERT(isint64()); GenRegister r = *this; r.type = type == GEN_TYPE_UL ? GEN_TYPE_UD : GEN_TYPE_D; - r.hstride = GEN_HORIZONTAL_STRIDE_2; + if(hstride != GEN_HORIZONTAL_STRIDE_0) + r.hstride = GEN_HORIZONTAL_STRIDE_2; r.vstride = GEN_VERTICAL_STRIDE_16; 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
