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 <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: "Yang, Rong R" <rong.r.yang@intel.com>

diff --git a/backend/src/backend/gen_register.hpp b/backend/src/backend/gen_register.hpp
index 66bc288..47b0515 100644 (file)
--- 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 (file)
--- 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 (file)
--- 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;