[Beignet] [PATCH] utests: add utests for global imm optimized

[rander.wang]

	there are some global immediates in global var list of LLVM.
	these imm can be integrated in instructions. for this test,
        there are two global immediates:

        MOV(1)              %42<0>:UD	:	0x0:UD
        MOV(1)              %43<0>:UD	:	0x30:UD

        used by:
        ADD(16)             %49<1>:D	:	%42<0,1,0>:D	%48<8,8,1>:D
        ADD(16)             %54<1>:D	:	%43<0,1,0>:D	%53<8,8,1>:D

	it can be
        ADD(16)             %49<1>:D    :       %48<8,8,1>:D   0x0:UD
        ADD(16)             %54<1>:D    :       %53<8,8,1>:D   0x30:UD

Signed-off-by: rander.wang <rander.wang at intel.com>
---
 kernels/compiler_global_immediate_optimized.cl | 49 ++++++++++++++++++++++++++
 utests/CMakeLists.txt                          |  3 +-
 utests/compiler_global_immediate_optimized.cpp | 29 +++++++++++++++
 3 files changed, 80 insertions(+), 1 deletion(-)
 create mode 100644 kernels/compiler_global_immediate_optimized.cl
 create mode 100644 utests/compiler_global_immediate_optimized.cpp

diff --git a/kernels/compiler_global_immediate_optimized.cl b/kernels/compiler_global_immediate_optimized.cl
new file mode 100644
index 0000000..9dcd72f
--- /dev/null
+++ b/kernels/compiler_global_immediate_optimized.cl
@@ -0,0 +1,49 @@
+constant int m[3] = {71,72,73};
+const constant int n = 1;
+constant int o[3] = {3, 2, 1};
+
+constant int4 a= {1, 2, 3, 4};
+constant int4 b = {0, -1, -2, -3};
+
+struct Person {
+  char name[7];
+  int3 idNumber;
+};
+
+struct Test1 {
+  int a0;
+  char a1;
+};
+
+struct Test2 {
+  char a0;
+  int a1;
+};
+struct Test3 {
+  int a0;
+  int a1;
+};
+struct Test4 {
+  float a0;
+  float a1;
+};
+
+constant struct Person james= {{"james"}, (int3)(1, 2, 3)};
+constant struct Test1 t0 = {1, 2};
+constant struct Test2 t1 = {1, 2};
+
+constant int3 c[3] = {(int3)(0, 1, 2), (int3)(3, 4, 5), (int3)(6,7,8) };
+constant char4 d[3] = {(char4)(0, 1, 2, 3), (char4)(4, 5, 6, 7), (char4)(8, 9, 10, 11)};
+
+constant struct Person members[3] = {{{"abc"}, (int3)(1, 2, 3)}, { {"defg"}, (int3)(4,5,6)}, { {"hijk"}, (int3)(7,8,9)} };
+constant struct Test3 zero_struct = {0, 0};
+constant int3 zero_vec = {0,0,0};
+constant int zero_arr[3] = {0,0,0};
+constant float zero_flt[3] = {0.0f, 0.0f, 0.0f};
+
+__kernel void
+compiler_global_immediate_optimized(__global int *dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] = c[id%3].y + d[id%3].w;
+}
\ No newline at end of file
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index b84cdd6..cd061b2 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -303,7 +303,8 @@ set (utests_sources
   compiler_generic_pointer.cpp
   runtime_pipe_query.cpp
   compiler_pipe_builtin.cpp
-  compiler_device_enqueue.cpp)
+  compiler_device_enqueue.cpp
+  compiler_global_immediate_optimized)
 
 if (LLVM_VERSION_NODOT VERSION_GREATER 34)
   SET(utests_sources
diff --git a/utests/compiler_global_immediate_optimized.cpp b/utests/compiler_global_immediate_optimized.cpp
new file mode 100644
index 0000000..753ea24
--- /dev/null
+++ b/utests/compiler_global_immediate_optimized.cpp
@@ -0,0 +1,29 @@
+#include "utest_helper.hpp"
+
+void compiler_global_immediate_optimized(void)
+{
+  const size_t n = 32;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_global_immediate_optimized");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+
+  // Run the kernel
+  globals[0] = n;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  uint32_t data1[] = {1, 4, 7};
+  uint32_t data2[]= {3, 7, 11};
+
+  // Check results
+  OCL_MAP_BUFFER(0);
+  for (uint32_t i = 0; i < n; ++i)
+//    printf("%d result %d reference %d\n", i, ((uint32_t *)buf_data[0])[i], data1[i%3] + data2[i%3]);
+    OCL_ASSERT(((uint32_t *)buf_data[0])[i] == data1[i%3] + data2[i%3]);
+  OCL_UNMAP_BUFFER(0);
+}
+
+
+MAKE_UTEST_FROM_FUNCTION_KEEP_PROGRAM(compiler_global_immediate_optimized, true);
-- 
2.7.4