[Beignet] [PATCH 3/3] Add stuct argument indirect load test.

[Zhigang Gong]

This patch LGTM, will push after the first 2 patches got merged.

Thanks,
Zhigang Gong.

On Mon, May 11, 2015 at 02:02:54PM +0800, Yang Rong wrote:
> 1. Enable compiler_argument_structure_indirect.
> 2. Add compiler_argument_structure_indirect, which has select address and load argument instruction.
> 
> Signed-off-by: Yang Rong <rong.r.yang at intel.com>
> ---
>  kernels/compiler_argument_structure_indirect.cl |  4 +--
>  kernels/compiler_argument_structure_select.cl   | 18 ++++++++++++
>  utests/CMakeLists.txt                           |  2 ++
>  utests/compiler_argument_structure_indirect.cpp |  7 +++--
>  utests/compiler_argument_structure_select.cpp   | 37 +++++++++++++++++++++++++
>  5 files changed, 63 insertions(+), 5 deletions(-)
>  create mode 100644 kernels/compiler_argument_structure_select.cl
>  create mode 100644 utests/compiler_argument_structure_select.cpp
> 
> diff --git a/kernels/compiler_argument_structure_indirect.cl b/kernels/compiler_argument_structure_indirect.cl
> index c4b062f..6fd7873 100644
> --- a/kernels/compiler_argument_structure_indirect.cl
> +++ b/kernels/compiler_argument_structure_indirect.cl
> @@ -1,7 +1,7 @@
> -struct hop { int x[16]; };
> +struct hop { int a, x[16]; };
>  
>  __kernel void
> -compiler_argument_structure(__global int *dst, struct hop h)
> +compiler_argument_structure_indirect(__global int *dst, struct hop h)
>  {
>    int id = (int)get_global_id(0);
>    dst[id] = h.x[get_local_id(0)];
> diff --git a/kernels/compiler_argument_structure_select.cl b/kernels/compiler_argument_structure_select.cl
> new file mode 100644
> index 0000000..295acf4
> --- /dev/null
> +++ b/kernels/compiler_argument_structure_select.cl
> @@ -0,0 +1,18 @@
> +typedef struct {
> +  int  offset;
> +  int  threshold0;
> +  int  threshold1;
> +}hop;
> +
> +__kernel void compiler_argument_structure_select(__global int *dst, hop h)
> +{
> +  int i = get_global_id (0);
> +  int threshold=0;
> +  if (i == 0)  {
> +    threshold = h.threshold0;
> +  } else {
> +    threshold = h.threshold1;
> +  }
> +  dst[i] = threshold;
> +}
> +
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index dcb3385..c9d96db 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -39,6 +39,8 @@ set (utests_sources
>    compiler_box_blur.cpp
>    compiler_insert_to_constant.cpp
>    compiler_argument_structure.cpp
> +  compiler_argument_structure_indirect.cpp
> +  compiler_argument_structure_select.cpp
>    compiler_arith_shift_right.cpp
>    compiler_mixed_pointer.cpp
>    compiler_array0.cpp
> diff --git a/utests/compiler_argument_structure_indirect.cpp b/utests/compiler_argument_structure_indirect.cpp
> index a4584d5..b54432e 100644
> --- a/utests/compiler_argument_structure_indirect.cpp
> +++ b/utests/compiler_argument_structure_indirect.cpp
> @@ -1,6 +1,6 @@
>  #include "utest_helper.hpp"
>  
> -struct hop { int x[16]; };
> +struct hop { int a, x[16]; };
>  
>  void compiler_argument_structure_indirect(void)
>  {
> @@ -21,8 +21,9 @@ void compiler_argument_structure_indirect(void)
>    OCL_MAP_BUFFER(0);
>  
>    // Check results
> -  for (uint32_t i = 0; i < n; ++i)
> -    OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 7);
> +  for (uint32_t i = 0; i < n; ++i ) {
> +    OCL_ASSERT(((uint32_t*)buf_data[0])[i] == (i%16));
> +  }
>  }
>  
>  MAKE_UTEST_FROM_FUNCTION(compiler_argument_structure_indirect);
> diff --git a/utests/compiler_argument_structure_select.cpp b/utests/compiler_argument_structure_select.cpp
> new file mode 100644
> index 0000000..b46e745
> --- /dev/null
> +++ b/utests/compiler_argument_structure_select.cpp
> @@ -0,0 +1,37 @@
> +#include "utest_helper.hpp"
> +
> +struct hop{
> +  int  offset;
> +  int  threshold0;
> +  int  threshold1;
> +};
> +
> +void compiler_argument_structure_select(void)
> +{
> +  const size_t n = 2048;
> +  hop h;
> +
> +  // Setup kernel and buffers
> +  OCL_CREATE_KERNEL("compiler_argument_structure_select");
> +  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> +  h.offset = 2;
> +  h.threshold0 = 5;
> +  h.threshold1 = 7;
> +  OCL_SET_ARG(1, sizeof(hop), &h);
> +
> +  // Run the kernel
> +  globals[0] = n;
> +  locals[0] = 16;
> +  OCL_NDRANGE(1);
> +  OCL_MAP_BUFFER(0);
> +
> +  // Check results
> +  OCL_ASSERT(((uint32_t*)buf_data[0])[0] == 5);
> +  for (uint32_t i = 1; i < n; ++i ) {
> +    OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 7);
> +  }
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_argument_structure_select);
> +
> -- 
> 1.8.3.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet