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

[Yang Rong]

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