Utest: Add a bitonic sort test for non-constant extractelement

Submitted by Pan Xiuli on Dec. 11, 2015, 5:28 a.m.

Details

Message ID 1449811737-32678-1-git-send-email-xiuli.pan@intel.com
State New
Headers show
Series "Utest: Add a bitonic sort test for non-constant extractelement" ( rev: 1 ) in Beignet

Not browsing as part of any series.

Commit Message

Pan Xiuli Dec. 11, 2015, 5:28 a.m.
This test case is for the new added non-constand index extractelement
path in llvm_scalarize pass.

Signed-off-by: Pan Xiuli <xiuli.pan@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

Patch hide | download patch | download mbox

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);

Comments

LGTM. Pushed, thanks.

> -----Original Message-----

> From: Beignet [mailto:beignet-bounces@lists.freedesktop.org] On Behalf Of

> Pan Xiuli

> Sent: Friday, December 11, 2015 13:29

> To: beignet@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@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@lists.freedesktop.org

> http://lists.freedesktop.org/mailman/listinfo/beignet