[Beignet] [PATCH V2] Fix some piglit constant buffer tests fail.

[Zhigang Gong]

On Tue, Jun 18, 2013 at 02:03:42PM +0800, Yang Rong wrote:
> If indirect move's source is scalrar reg, such as using cb[0] in kernel,
> should not unpack.
> 
> Also change test case compiler_function_constant0 to trigger this bug.
> 
> Signed-off-by: Yang Rong <rong.r.yang at intel.com>
> ---
>  backend/src/backend/gen_context.cpp    |    7 ++++++-
>  backend/src/backend/gen_register.hpp   |    3 ++-
>  kernels/compiler_function_constant0.cl |    2 +-
>  utests/compiler_function_constant0.cpp |    6 ++----
>  4 files changed, 11 insertions(+), 7 deletions(-)
> 
> diff --git a/backend/src/backend/gen_context.cpp b/backend/src/backend/gen_context.cpp
> index af651e7..d0c2ff8 100644
> --- a/backend/src/backend/gen_context.cpp
> +++ b/backend/src/backend/gen_context.cpp
> @@ -223,7 +223,12 @@ namespace gbe
>    }
>  
>    void GenContext::emitIndirectMoveInstruction(const SelectionInstruction &insn) {
> -    const GenRegister src = GenRegister::unpacked_uw(ra->genReg(insn.src(0)).nr, 0);
> +    GenRegister src = ra->genReg(insn.src(0));
> +    if(isScalarReg(src.reg()))
> +      src = GenRegister::retype(src, GEN_TYPE_UW);
> +    else
> +      src = GenRegister::unpacked_uw(src.nr, src.subnr);

Maybe fix the subnr number here is better. src.subnr/typeSize(GEN_TYPE_UW).
As one policy is that we only use byte as subnr unit at the physical register scope.
For any other macros, we use the element as subnr unit. You can check the other places
using subnr which should all comply with this policy.

For examples:

static INLINE GenRegister ub16(uint32_t file, uint32_t nr, uint32_t subnr)
static INLINE GenRegister uw1(uint32_t file, uint32_t nr, uint32_t subnr) 
...

The subnr are all using element as unit rather than byte. So I suggest you also use
element as unit for new function:

static INLINE GenRegister unpacked_uw(uint32_t nr, uint32_t subnr)

Any thoughts?

> +
>      const GenRegister dst = ra->genReg(insn.dst(0));
>      const GenRegister a0 = GenRegister::addr8(0);
>      uint32_t simdWidth = p->curr.execWidth;
> diff --git a/backend/src/backend/gen_register.hpp b/backend/src/backend/gen_register.hpp
> index d772b0d..eb33ae8 100644
> --- a/backend/src/backend/gen_register.hpp
> +++ b/backend/src/backend/gen_register.hpp
> @@ -726,9 +726,10 @@ namespace gbe
>      }
>  
>      static INLINE GenRegister unpacked_uw(uint32_t nr, uint32_t subnr) {
> +      GBE_ASSERT((subnr%typeSize(GEN_TYPE_UW)) == 0);
Then no need to add this assertion.
>        return GenRegister(GEN_GENERAL_REGISTER_FILE,
>                           nr,
> -                         subnr,
> +                         subnr / typeSize(GEN_TYPE_UW),  //will multiply again in GenRegister, keep unchanged
And don't need to change the subnr.
>                           GEN_TYPE_UW,
>                           GEN_VERTICAL_STRIDE_16,
>                           GEN_WIDTH_8,
> diff --git a/kernels/compiler_function_constant0.cl b/kernels/compiler_function_constant0.cl
> index f6efcef..363d84e 100644
> --- a/kernels/compiler_function_constant0.cl
> +++ b/kernels/compiler_function_constant0.cl
> @@ -2,5 +2,5 @@ __kernel void
>  compiler_function_constant0(__constant short *c0, __constant char *c1, __global int *dst, int value)
>  {
>    int id = (int)get_global_id(0);
> -  dst[id] = value + c0[id%69] + c1[15];
> +  dst[id] = value + c0[id%69] + c1[0];
>  }
> diff --git a/utests/compiler_function_constant0.cpp b/utests/compiler_function_constant0.cpp
> index de564f3..c0a8a9d 100644
> --- a/utests/compiler_function_constant0.cpp
> +++ b/utests/compiler_function_constant0.cpp
> @@ -8,7 +8,7 @@ void compiler_function_constant0(void)
>    // Setup kernel and buffers
>    OCL_CREATE_KERNEL("compiler_function_constant0");
>    OCL_CREATE_BUFFER(buf[0], 0, 75 * sizeof(short), NULL);
> -  OCL_CREATE_BUFFER(buf[1], 0, 256 * sizeof(char), NULL);
> +  OCL_CREATE_BUFFER(buf[1], 0, 1 * sizeof(char), NULL);
>    OCL_CREATE_BUFFER(buf[2], 0, n * sizeof(uint32_t), NULL);
>    OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
>    OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> @@ -21,9 +21,7 @@ void compiler_function_constant0(void)
>    OCL_UNMAP_BUFFER(0);
>  
>    OCL_MAP_BUFFER(1);
> -  for(uint32_t i = 0; i < 256; ++i)
> -    ((char *)buf_data[1])[i] = 10;
> -  ((char *)buf_data[1])[15] = 15;
> +  ((char *)buf_data[1])[0] = 15;
>    OCL_UNMAP_BUFFER(1);
>  
>    // Run the kernel
> -- 
> 1.7.10.4
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet