Add utest case for movforphi's undef case.

This case will trigger MovForPhi to handle a undef vector
element.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Reviewed-by: Homer Hsing <homer.xing@intel.com>

diff --git a/kernels/test_movforphi_undef.cl b/kernels/test_movforphi_undef.cl
new file mode 100644 (file)
index 0000000..035c02a
--- /dev/null
+++ b/kernels/test_movforphi_undef.cl
@@ -0,0 +1,18 @@
+__kernel void
+test_movforphi_undef(__read_only image2d_t src, __write_only image2d_t dst, sampler_t sampler)
+{
+  int2 coord, dstCoord;
+  int4 color;
+  int x = get_global_id(0);
+  int y = get_global_id(1);
+  dstCoord.x = x;
+  dstCoord.y = y;
+  coord.y = y;
+  for(int j = -8; j < 2; j++)
+  {
+    coord.x = j + x;
+    color = read_imagei(src, sampler, coord);
+    if (j == 1 + x)
+      write_imagei(dst, dstCoord, color);
+  }
+}

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 50e0baa..46d4c86 100644 (file)
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -63,6 +63,7 @@ ADD_LIBRARY(utests SHARED
   compiler_local_memory_two_ptr.cpp
   compiler_local_memory_barrier.cpp
   compiler_local_memory_barrier_wg64.cpp
+  compiler_movforphi_undef.cpp
   utest_assert.cpp
   utest.cpp
   utest_file_map.cpp

diff --git a/utests/compiler_movforphi_undef.cpp b/utests/compiler_movforphi_undef.cpp
new file mode 100644 (file)
index 0000000..f2ee009
--- /dev/null
+++ b/utests/compiler_movforphi_undef.cpp
@@ -0,0 +1,51 @@
+#include "utest_helper.hpp"
+
+static void compiler_movforphi_undef(void)
+{
+  const size_t w = 16;
+  const size_t h = 16;
+  cl_image_format format;
+  cl_sampler sampler;
+
+  format.image_channel_order = CL_RGBA;
+  format.image_channel_data_type = CL_UNSIGNED_INT8;
+
+  // Setup kernel and images
+  OCL_CREATE_KERNEL("test_movforphi_undef");
+  buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * w * h);
+  for (uint32_t j = 0; j < h; ++j)
+    for (uint32_t i = 0; i < w; i++)
+      ((uint32_t*)buf_data[0])[j * w + i] = j * w + i;
+
+  OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, w, h, w * sizeof(uint32_t), buf_data[0]);
+  OCL_CREATE_IMAGE(buf[1], 0, &format, w, h, w * sizeof(uint32_t), NULL);
+  OCL_CREATE_SAMPLER(sampler, CL_ADDRESS_REPEAT, CL_FILTER_NEAREST);
+  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]);
+  OCL_SET_ARG(2, sizeof(sampler), &sampler);
+  globals[0] = w;
+  globals[1] = h;
+  locals[0] = 16;
+  locals[1] = 16;
+  OCL_NDRANGE(2);
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  // Just compare the initial 2 data is enough for this case, as the initial 2 data must in the first
+  // tile box and we can just get the correct coords.
+  for (uint32_t j = 0; j < 1; ++j)
+    for (uint32_t i = 0; i < 3; i++)
+    {
+      if (i < w - 1)
+       OCL_ASSERT(((uint32_t*)buf_data[0])[j * w + i + 1] == ((uint32_t*)buf_data[1])[j * w + i]);
+    }
+  OCL_UNMAP_BUFFER(0);
+  OCL_UNMAP_BUFFER(1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_movforphi_undef);