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

[Zhigang Gong]

It's better to never modify the hstride manually. Always modify it through h2 method.
And in h2 method, we need to check whether it is scalar, and only set to h2 for non
scalar.

Otherwise, we may have many chance to make mistake with this scalar/non-scalar stride issue.

This patch is good for me. Will push latter. Thanks.

On Thu, Nov 07, 2013 at 06:05:23AM +0000, Song, Ruiling wrote:
> Yes, h2 seems like same issue, but I see it only works on address register. Suppose it will not be scalar register, right?
> I also prefer fix h2().
> Anybody else have comments? Thanks!
> 
> -----Original Message-----
> From: Yang, Rong R 
> Sent: Thursday, November 07, 2013 1:45 PM
> To: Song, Ruiling; beignet at lists.freedesktop.org
> Cc: Song, Ruiling
> Subject: RE: [Beignet] [PATCH] GBE: fix a 64bit scalar register issue.
> 
> LGTM.
> 
> BTW: Does the function GenRegister::h2 also have this issue?
> 
> -----Original Message-----
> From: beignet-bounces at lists.freedesktop.org [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of Ruiling Song
> Sent: Wednesday, November 06, 2013 1:40 PM
> To: beignet at lists.freedesktop.org
> 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 <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
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet