[Beignet] [PATCH 3/4] add basic utest for block_motion_estimate_intel

[Chuanbo Weng]

From: Guo Yejun <yejun.guo at intel.com>

if the CL device does not support this builtin kernel,
the test returns PASS.

Signed-off-by: Guo Yejun <yejun.guo at intel.com>
---
 utests/CMakeLists.txt                              |   1 +
 .../builtin_kernel_block_motion_estimate_intel.cpp | 109 +++++++++++++++++++++
 utests/utest_helper.hpp                            |   1 +
 3 files changed, 111 insertions(+)
 create mode 100644 utests/builtin_kernel_block_motion_estimate_intel.cpp

diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index e7a9e26..f5abb0c 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -203,6 +203,7 @@ set (utests_sources
   test_printf.cpp
   enqueue_fill_buf.cpp
   builtin_kernel_max_global_size.cpp
+  builtin_kernel_block_motion_estimate_intel.cpp
   image_1D_buffer.cpp
   compare_image_2d_and_1d_array.cpp
   compiler_fill_image_1d_array.cpp
diff --git a/utests/builtin_kernel_block_motion_estimate_intel.cpp b/utests/builtin_kernel_block_motion_estimate_intel.cpp
new file mode 100644
index 0000000..12bcb7d
--- /dev/null
+++ b/utests/builtin_kernel_block_motion_estimate_intel.cpp
@@ -0,0 +1,109 @@
+#include "utest_helper.hpp"
+#include <string.h>
+
+void builtin_kernel_block_motion_estimate_intel(void)
+{
+  char* built_in_kernel_names;
+  size_t built_in_kernels_size;
+  cl_int err = CL_SUCCESS;
+  size_t ret_sz;
+
+  OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, 0, 0, &built_in_kernels_size);
+  built_in_kernel_names = (char* )malloc(built_in_kernels_size * sizeof(char) );
+  OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, built_in_kernels_size, (void*)built_in_kernel_names, &ret_sz);
+  OCL_ASSERT(ret_sz == built_in_kernels_size);
+
+  if (strstr(built_in_kernel_names, "block_motion_estimate_intel") == NULL)
+  {
+        free(built_in_kernel_names);
+        return;
+  }
+
+  cl_program built_in_prog = clCreateProgramWithBuiltInKernels(ctx, 1, &device, built_in_kernel_names, &err);
+  OCL_ASSERT(built_in_prog != NULL);
+  kernel = clCreateKernel(built_in_prog, "block_motion_estimate_intel",  &err);
+  OCL_ASSERT(kernel != NULL);
+
+  cl_motion_estimation_desc_intel vmedesc = {CL_ME_MB_TYPE_16x16_INTEL,       //0x0
+                                          CL_ME_SUBPIXEL_MODE_INTEGER_INTEL,  //0x0
+                                          CL_ME_SAD_ADJUST_MODE_NONE_INTEL,   //0x0
+                                          CL_ME_SEARCH_PATH_RADIUS_16_12_INTEL //0x5
+                                          };
+  cl_accelerator_intel accel = clCreateAcceleratorINTEL(ctx, CL_ACCELERATOR_TYPE_MOTION_ESTIMATION_INTEL,sizeof(cl_motion_estimation_desc_intel), &vmedesc, &err);
+  OCL_ASSERT(accel != NULL);
+
+  const size_t w = 71; //80
+  const size_t h = 41; //48
+
+  cl_image_format format;
+  cl_image_desc desc;
+
+  memset(&desc, 0x0, sizeof(cl_image_desc));
+  memset(&format, 0x0, sizeof(cl_image_format));
+
+  uint8_t* image_data1 = (uint8_t *)malloc(w * h);    //src
+  uint8_t* image_data2 = (uint8_t *)malloc(w * h);    //ref
+  for (size_t j = 0; j < h; j++) {
+    for (size_t i = 0; i < w; i++) {
+      if (i >= 32 && i <= 47 && j >= 16 && j <= 31)
+        image_data2[w * j + i] = image_data1[w * j + i] = 100;
+      else
+        image_data2[w * j + i] = image_data1[w * j + i] = 0;
+    }
+  }
+
+  format.image_channel_order = CL_R;
+  format.image_channel_data_type = CL_UNORM_INT8;
+  desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+  desc.image_width = w;
+  desc.image_height = h;
+  desc.image_row_pitch = 0;
+  OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data1);        //src
+  OCL_CREATE_IMAGE(buf[1], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data2);        //ref
+
+  const size_t mv = (80/16) * (48/16);
+  OCL_CREATE_BUFFER(buf[2], 0, mv * sizeof(int) * 4, NULL);
+
+  OCL_SET_ARG(0, sizeof(cl_accelerator_intel), &accel);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(2, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(3, sizeof(cl_mem), NULL);
+  OCL_SET_ARG(4, sizeof(cl_mem), &buf[2]);
+  OCL_SET_ARG(5, sizeof(cl_mem), NULL);
+
+  globals[0] = w;
+  globals[1] = h;
+  OCL_CALL(clEnqueueNDRangeKernel, queue, kernel, 2, NULL, globals, NULL, 0, NULL, NULL);
+
+  OCL_MAP_BUFFER(2);
+  short expected[] = {-64, -48,
+                    -64, -48,
+                    -64, -48,
+                    -64, -48,
+                    -64, -48,
+                    -64, -48,
+                    -64, -48,
+                    0, 0,
+                    0, -48,
+                    -64, -48,
+                    -64, -48,
+                    -64, -48,
+                    -64, -48,
+                    0, -48,
+                    -64, -48};
+  short* res = (short*)buf_data[2];
+  for (uint32_t j = 0; j < mv; ++j) {
+    OCL_ASSERT(res[j * 2 + 0] == expected[j * 2 + 0]);
+    OCL_ASSERT(res[j * 2 + 1] == expected[j * 2 + 1]);
+  }
+  OCL_UNMAP_BUFFER(2);
+
+  clReleaseAcceleratorINTEL(accel);
+  clReleaseKernel(kernel);
+  clReleaseProgram(built_in_prog);
+  free(built_in_kernel_names);
+  free(image_data1);
+  free(image_data2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(builtin_kernel_block_motion_estimate_intel);
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index 3b17606..d8f8a50 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -26,6 +26,7 @@
 #define __UTEST_HELPER_HPP__
 
 #include "CL/cl.h"
+#include "CL/cl_ext.h"
 #include "CL/cl_intel.h"
 #include "utest.hpp"
 #include "utest_assert.hpp"
-- 
1.9.1