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

[Zhigang Gong]

Pushed with ruiling and mine comments. Thanks.

On Wed, Jun 19, 2013 at 03:36:34PM +0800, Yang Rong wrote:
> If indirect move's source is scalrar reg, such as using cb[0] in kernel,
> should not unpack.
> 
> 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 ++++++-
>  kernels/compiler_function_constant0.cl |    2 +-
>  utests/compiler_function_constant0.cpp |    6 ++----
>  3 files changed, 9 insertions(+), 6 deletions(-)
> 
> diff --git a/backend/src/backend/gen_context.cpp b/backend/src/backend/gen_context.cpp
> index af651e7..70c5bcf 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 / typeSize(GEN_TYPE_UW));
> +
>      const GenRegister dst = ra->genReg(insn.dst(0));
>      const GenRegister a0 = GenRegister::addr8(0);
>      uint32_t simdWidth = p->curr.execWidth;
> 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