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@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-boun...@lists.freedesktop.org > [mailto:beignet-boun...@lists.freedesktop.org] On Behalf Of Ruiling Song > Sent: Wednesday, November 06, 2013 1:40 PM > To: beignet@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.s...@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@lists.freedesktop.org > http://lists.freedesktop.org/mailman/listinfo/beignet > _______________________________________________ > Beignet mailing list > Beignet@lists.freedesktop.org > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet
