[Beignet] [PATCH] Add a test case that trigger vstore3 bug.

[Weng, Chuanbo]

Hi Zhigang,
	I've updated to the latest commit.
	The main difference between our environment is the llvm version.
When I switch to llvm-3.3, this case pass. It fails when using llvm-3.2.

-----Original Message-----
From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com] 
Sent: Thursday, August 29, 2013 2:42 PM
To: Weng, Chuanbo
Cc: beignet at lists.freedesktop.org; Zou, Nanhai; Gong, Zhigang
Subject: Re: [Beignet] [PATCH] Add a test case that trigger vstore3 bug.

Hi chuanbo,

I just tried this case on my machine. It just passed as below:

compiler_vstore3_no_workgroup:
  compiler_vstore3_no_workgroup()    [SUCCESS]

Did you use the latest master branch to do the test?

On Tue, Aug 27, 2013 at 10:47:37PM +0800, Chuanbo Weng wrote:
> When there is no workgroup, the vstore3 will store incorrect value to 
> buffer.
> 
> Signed-off-by: Chuanbo Weng <chuanbo.weng at intel.com>
> ---
>  kernels/compiler_vstore3_no_workgroup.cl |   11 +++++++++
>  utests/CMakeLists.txt                    |    1 +
>  utests/compiler_vstore3_no_workgroup.cpp |   39 ++++++++++++++++++++++++++++++
>  3 files changed, 51 insertions(+)
>  create mode 100644 kernels/compiler_vstore3_no_workgroup.cl
>  create mode 100644 utests/compiler_vstore3_no_workgroup.cpp
> 
> diff --git a/kernels/compiler_vstore3_no_workgroup.cl 
> b/kernels/compiler_vstore3_no_workgroup.cl
> new file mode 100644
> index 0000000..aa641ae
> --- /dev/null
> +++ b/kernels/compiler_vstore3_no_workgroup.cl
> @@ -0,0 +1,11 @@
> +__kernel
> +void compiler_vstore3_no_workgroup(__global uint *vectorArray){
> +  int col = get_global_id(0);
> +
> +  uint3 val3;
> +  val3.x = col*3 + 0;
> +  val3.y = col*3 + 1;
> +  val3.z = col*3 + 2;
> +  size_t thumb_offset = col;
> +  vstore3(val3, thumb_offset, vectorArray); }
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 
> 97b7519..d9cb024 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -135,6 +135,7 @@ set (utests_sources
>    compiler_long_asr.cpp
>    compiler_long_mult.cpp
>    compiler_long_cmp.cpp
> +  compiler_vstore3_no_workgroup.cpp
>    utest_assert.cpp
>    utest.cpp
>    utest_file_map.cpp
> diff --git a/utests/compiler_vstore3_no_workgroup.cpp 
> b/utests/compiler_vstore3_no_workgroup.cpp
> new file mode 100644
> index 0000000..96491cd
> --- /dev/null
> +++ b/utests/compiler_vstore3_no_workgroup.cpp
> @@ -0,0 +1,39 @@
> +#include "utest_helper.hpp"
> +
> +static void cpu(int global_id, unsigned int *dst) {
> +  for(unsigned int j = 0; j < 3; j++){
> +    *(dst + global_id*3 + j) = global_id*3 + j;
> +  }
> +}
> +
> +void compiler_vstore3_no_workgroup(void){
> +
> +  const size_t n = 16;
> +  unsigned int cpu_dst[16*3];
> +	
> +  // Setup kernel and buffers
> +  OCL_CREATE_KERNEL("compiler_vstore3_no_workgroup");
> +  OCL_CREATE_BUFFER(buf[0], 0, n*3*sizeof(unsigned int), NULL);  
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);  globals[0] = n;
> +
> +  // Run the kernel on GPU
> +  OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, 1, NULL, globals, 
> + NULL, 0, NULL, NULL);
> +
> +  // Run on CPU
> +  for (int32_t i = 0; i < (int32_t) n; ++i)
> +    cpu(i, cpu_dst);
> +
> +  // Compare
> +  OCL_MAP_BUFFER(0);
> +  unsigned int *px = (unsigned int *)buf_data[0];  for (int32_t i = 
> + 0; i < (int32_t) n; ++i)
> +    for (int32_t j = 0; j < 3; ++j){
> +      OCL_ASSERT(*(px+i*3+j) == *(cpu_dst+i*3+j));
> +    }
> +  OCL_UNMAP_BUFFER(0);
> +
> +}
> +
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_vstore3_no_workgroup)
> --
> 1.7.9.5
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet