utest: add new test for constant expression processing.

If we use 3-component vector in a union, it may introduce
some complex constant expression as below:

float bitcast (i32 trunc (i128 bitcast (<4 x i32> <i32 1065353216, i32 1073741824, i32 1077936128, i32 undef> to i128) to i32) to float).

To test the constant expression processing function.

Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
Reviewed-by: "Song, Ruiling" <ruiling.song@intel.com>

3 files changed, 59 insertions, 0 deletions

diff --git a/kernels/compiler_constant_expr.cl b/kernels/compiler_constant_expr.cl
new file mode 100644
index 0000000..d40cead
--- /dev/null
+++ b/kernels/compiler_constant_expr.cl
@@ -0,0 +1,23 @@
+float3 foo_pow3(float3 src0, float3 src1)
+{
+  union {
+    float3 f3;
+    float   farray[4];
+  } s0, s1, dst;
+  s0.f3 = src0;
+  s1.f3 = src1;
+  int i;
+  for(i = 0; i < 3; i++)
+    dst.farray[i] = pow(s0.farray[i], s1.farray[i]);
+  return dst.f3;
+}
+
+__kernel void
+compiler_constant_expr(__global float* src, __global float *dst)
+{
+  int gid = get_global_id(0);
+  float3 f3 = vload3(gid, src);
+  float3 cf3 = (float3)(1.f, 2.f, 3.f);
+  float3 result = foo_pow3(f3, cf3);
+  vstore3(result, gid, dst); 
+} 
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 561744d..2605abd 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -182,6 +182,7 @@ set (utests_sources
   enqueue_built_in_kernels.cpp
   image_1D_buffer.cpp
   compare_image_2d_and_1d_array.cpp
+  compiler_constant_expr.cpp
   utest_assert.cpp
   utest.cpp
   utest_file_map.cpp
diff --git a/utests/compiler_constant_expr.cpp b/utests/compiler_constant_expr.cpp
new file mode 100644
index 0000000..8bed03b
--- /dev/null
+++ b/utests/compiler_constant_expr.cpp
@@ -0,0 +1,35 @@
+#include "utest_helper.hpp"
+#include <math.h>
+
+static void compiler_constant_expr(void)
+{
+  const size_t n = 48;
+
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_constant_expr");
+  buf_data[0] = (uint32_t*) malloc(sizeof(float) * n);
+  for (uint32_t i = 0; i < n; ++i) ((float*)buf_data[0])[i] = i;
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(float), buf_data[0]);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
+  free(buf_data[0]);
+  buf_data[0] = NULL;
+
+  // Run the kernel
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = 16;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < n; ++i) {
+    float expect = pow(((float*)buf_data[0])[i], (i % 3) + 1);
+    float err = fabs(((float*)buf_data[1])[i] - expect);
+    OCL_ASSERT(err <= 100 * cl_FLT_ULP(expect));
+  }
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_constant_expr);
+