[Beignet] [PATCH 2/2] utests: Add test case for box blur on image buffer

[Dag Lem]

This test demonstrates box blur on an image buffer, using an extremely
simple kernel.

Signed-off-by: Dag Lem <dag at nimrod.no>
---
 kernels/compiler_box_blur_image.cl | 18 +++++++++++++
 utests/CMakeLists.txt              |  1 +
 utests/compiler_box_blur_image.cpp | 52 ++++++++++++++++++++++++++++++++++++++
 utests/utest_helper.hpp            |  6 +++++
 4 files changed, 77 insertions(+)
 create mode 100644 kernels/compiler_box_blur_image.cl
 create mode 100644 utests/compiler_box_blur_image.cpp

diff --git a/kernels/compiler_box_blur_image.cl b/kernels/compiler_box_blur_image.cl
new file mode 100644
index 0000000..7bcbdeb
--- /dev/null
+++ b/kernels/compiler_box_blur_image.cl
@@ -0,0 +1,18 @@
+__kernel void compiler_box_blur_image(__read_only image2d_t src,
+                                      __write_only image2d_t dst)
+{
+  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                            CLK_ADDRESS_CLAMP_TO_EDGE |
+                            CLK_FILTER_NEAREST;
+  const int2 coord = (int2)(get_global_id(0), get_global_id(1));
+  int2 offset;
+  float4 sum = 0;
+
+  for (offset.y = -1; offset.y <= 1; offset.y++) {
+    for (offset.x = -1; offset.x <= 1; offset.x++) {
+      sum += read_imagef(src, sampler, coord + offset);
+    }
+  }
+
+  write_imagef(dst, coord, (1.0f/9.0f)*sum);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 63c873d..8700761 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -9,6 +9,7 @@ set (utests_sources
   compiler_mandelbrot.cpp
   compiler_mandelbrot_alternate.cpp
   compiler_box_blur_float.cpp
+  compiler_box_blur_image.cpp
   compiler_box_blur.cpp
   compiler_insert_to_constant.cpp
   compiler_argument_structure.cpp
diff --git a/utests/compiler_box_blur_image.cpp b/utests/compiler_box_blur_image.cpp
new file mode 100644
index 0000000..d94a97c
--- /dev/null
+++ b/utests/compiler_box_blur_image.cpp
@@ -0,0 +1,52 @@
+#include "utest_helper.hpp"
+
+static void compiler_box_blur_image()
+{
+  int w, h;
+  cl_image_format format = { };
+  cl_image_desc desc = { };
+  size_t origin[3] = { };
+  size_t region[3];
+  int *src, *dst;
+
+  OCL_CREATE_KERNEL("compiler_box_blur_image");
+
+  /* Load the picture */
+  src = cl_read_bmp("lenna128x128.bmp", &w, &h);
+
+  format.image_channel_order = CL_RGBA;
+  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_depth = 1;
+  desc.image_row_pitch = w*sizeof(uint32_t);
+
+  /* Run the kernel */
+  OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, src);
+  free(src);
+  desc.image_row_pitch = 0;
+  OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  globals[0] = w;
+  globals[1] = h;
+  locals[0] = 16;
+  locals[1] = 16;
+  OCL_NDRANGE(2);
+  dst = (int*)malloc(w*h*sizeof(uint32_t));
+  region[0] = w;
+  region[1] = h;
+  region[2] = 1;
+  OCL_READ_IMAGE(buf[1], origin, region, dst);
+
+  /* Save the image (for debug purpose) */
+  cl_write_bmp(dst, w, h, "compiler_box_blur_image.bmp");
+
+  /* Compare with the golden image */
+  OCL_CHECK_IMAGE(dst, w, h, "compiler_box_blur_ref.bmp");
+
+  free(dst);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_box_blur_image);
diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
index d882fc7..c28615e 100644
--- a/utests/utest_helper.hpp
+++ b/utests/utest_helper.hpp
@@ -88,6 +88,12 @@ extern EGLSurface  eglSurface;
 #define OCL_CREATE_IMAGE(IMAGE, FLAGS, FORMAT, DESC, DATA) \
     OCL_CALL2(clCreateImage, IMAGE, ctx, FLAGS, FORMAT, DESC, DATA)
 
+#define OCL_READ_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
+    OCL_CALL(clEnqueueReadImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
+
+#define OCL_WRITE_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
+    OCL_CALL(clEnqueueWriteImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
+
 #define OCL_CREATE_GL_IMAGE(IMAGE, FLAGS, TARGET, LEVEL, TEXTURE) \
     OCL_CALL2(clCreateFromGLTexture, IMAGE, ctx, FLAGS, TARGET, LEVEL, TEXTURE)
 
-- 
1.8.1.4