[Beignet] [PATCH] GBE: fix a 64bit scalar register issue.

[Ruiling Song]

For scalar register, should use stride 0.
also change the unit test to hit the point.

Signed-off-by: Ruiling Song <ruiling.song at intel.com>
---
 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