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

[Song, Ruiling]

Good suggestion, I will modify h2() and send new version.

-----Original Message-----
From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com] 
Sent: Thursday, November 07, 2013 2:16 PM
To: Song, Ruiling
Cc: Yang, Rong R; beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH] GBE: fix a 64bit scalar register issue.

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