[Beignet] [PATCH] Utest: Add a bitonic sort test for non-constant extractelement
Pan Xiuli
xiuli.pan at intel.com
Thu Dec 10 21:28:57 PST 2015
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
More information about the Beignet
mailing list