[Beignet] [PATCH] Utest: Add a bitonic sort test for non-constant extractelement

[Yang, Rong R]

LGTM. Pushed, thanks.

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Pan Xiuli
> Sent: Friday, December 11, 2015 13:29
> To: beignet at lists.freedesktop.org
> Cc: Pan, Xiuli
> Subject: [Beignet] [PATCH] Utest: Add a bitonic sort test for non-constant
> extractelement
> 
> This test case is for the new added non-constand index extractelement path
> in llvm_scalarize pass.
> 
> Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
> ---
>  kernels/compiler_bsort.cl | 47
> +++++++++++++++++++++++++++++++++++++++++++++++
>  utests/CMakeLists.txt     |  3 ++-
>  utests/compiler_bsort.cpp | 45
> +++++++++++++++++++++++++++++++++++++++++++++
>  3 files changed, 94 insertions(+), 1 deletion(-)  create mode 100644
> kernels/compiler_bsort.cl  create mode 100644 utests/compiler_bsort.cpp
> 
> diff --git a/kernels/compiler_bsort.cl b/kernels/compiler_bsort.cl new file
> mode 100644 index 0000000..db40da8
> --- /dev/null
> +++ b/kernels/compiler_bsort.cl
> @@ -0,0 +1,47 @@
> +#define UP 0
> +#define DOWN -1
> +
> +/* Sort elements in a vector */
> +#define SORT_VECTOR(input, dir)                                   \
> +   comp = input < shuffle(input, mask1) ^ dir;                    \
> +   input = shuffle(input, as_uint4(comp + add1));                 \
> +   comp = input < shuffle(input, mask2) ^ dir;                    \
> +   input = shuffle(input, as_uint4(comp * 2 + add2));             \
> +   comp = input < shuffle(input, mask3) ^ dir;                    \
> +   input = shuffle(input, as_uint4(comp + add3));                 \
> +
> +/* Sort elements between two vectors */
> +#define SWAP_VECTORS(input1, input2, dir)                         \
> +   temp = input1;                                                 \
> +   comp = (input1 < input2 ^ dir) * 4 + add4;                     \
> +   input1 = shuffle2(input1, input2, as_uint4(comp));             \
> +   input2 = shuffle2(input2, temp, as_uint4(comp));               \
> +
> +__kernel void compiler_bsort(__global float4 *data) {
> +
> +   float4 input1, input2, temp;
> +   int4 comp;
> +
> +   uint4 mask1 = (uint4)(1, 0, 3, 2);
> +   uint4 mask2 = (uint4)(2, 3, 0, 1);
> +   uint4 mask3 = (uint4)(3, 2, 1, 0);
> +
> +   int4 add1 = (int4)(1, 1, 3, 3);
> +   int4 add2 = (int4)(2, 3, 2, 3);
> +   int4 add3 = (int4)(1, 2, 2, 3);
> +   int4 add4 = (int4)(4, 5, 6, 7);
> +
> +   input1 = data[0];
> +   input2 = data[1];
> +
> +   SORT_VECTOR(input1, UP)
> +   SORT_VECTOR(input2, DOWN)
> +
> +   SWAP_VECTORS(input1, input2, UP)
> +
> +   SORT_VECTOR(input1, UP)
> +   SORT_VECTOR(input2, UP)
> +
> +   data[0] = input1;
> +   data[1] = input2;
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index
> b91c4ac..7c9822c 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -221,7 +221,8 @@ set (utests_sources
>    runtime_use_host_ptr_image.cpp
>    compiler_get_sub_group_size.cpp
>    compiler_get_sub_group_id.cpp
> -  compiler_sub_group_shuffle.cpp)
> +  compiler_sub_group_shuffle.cpp
> +  compiler_bsort.cpp)
> 
>  if (LLVM_VERSION_NODOT VERSION_GREATER 34)
>    SET(utests_sources
> diff --git a/utests/compiler_bsort.cpp b/utests/compiler_bsort.cpp new file
> mode 100644 index 0000000..3588f6b
> --- /dev/null
> +++ b/utests/compiler_bsort.cpp
> @@ -0,0 +1,45 @@
> +#include "utest_helper.hpp"
> +/*
> + * This test is for non-constant extractelement scalarize
> + * this bitonic sort test will use this path in
> + *
> + *  comp = input < shuffle(input, mask1) ^ dir;                    \
> + *  input = shuffle(input, as_uint4(comp + add1));                 \
> + *
> + * The origin buff is
> + * {3.0 5.0 4.0 6.0 0.0 7.0 2.0 1.0}
> + * and the expected result is
> + * {0.0 1.0 2.0 3.0 4.0 5.0 6.0 7.0}
> + */
> +void compiler_bsort(void)
> +{
> +  const int n = 8;
> +
> +  // Setup kernel and buffers
> +  OCL_CREATE_KERNEL("compiler_bsort");
> +  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);  globals[0] = 1;  locals[0] =
> + 1;
> +
> +  OCL_MAP_BUFFER(0);
> +  ((float *)(buf_data[0]))[0] = 3.0f;
> +  ((float *)(buf_data[0]))[1] = 5.0f;
> +  ((float *)(buf_data[0]))[2] = 4.0f;
> +  ((float *)(buf_data[0]))[3] = 6.0f;
> +  ((float *)(buf_data[0]))[4] = 0.0f;
> +  ((float *)(buf_data[0]))[5] = 7.0f;
> +  ((float *)(buf_data[0]))[6] = 2.0f;
> +  ((float *)(buf_data[0]))[7] = 1.0f;
> +  OCL_UNMAP_BUFFER(0);
> +
> +  OCL_NDRANGE(1);
> +
> +  OCL_MAP_BUFFER(0);
> +  for (int i = 0; i < n; i ++) {
> +    OCL_ASSERT(((float *)(buf_data[0]))[i] == (float)i);
> +  }
> +  OCL_UNMAP_BUFFER(0);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_bsort);
> --
> 2.1.4
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet