[Beignet] [PATCH v2 4/4] Add a benchmark which test do 3*3 median filter in image.

[Meng Mengmeng]

It's basic image test for uchar, ushort and uint.

v2:
convert uint to float before do median filter and
use intermediate variable in if loop.

Signed-off-by: Meng Mengmeng <mengmeng.meng at intel.com>
---
 benchmark/benchmark_copy_image.cpp | 17 ++++++++++-------
 kernels/bench_copy_image.cl        | 37 +++++++++++++++++++++++++++++++++++++
 2 files changed, 47 insertions(+), 7 deletions(-)

diff --git a/benchmark/benchmark_copy_image.cpp b/benchmark/benchmark_copy_image.cpp
index 40cb647..1d35188 100644
--- a/benchmark/benchmark_copy_image.cpp
+++ b/benchmark/benchmark_copy_image.cpp
@@ -2,8 +2,8 @@
 #include "utests/utest_helper.hpp"
 #include <sys/time.h>
 
-#define BENCH_COPY_IMAGE(T, M, Q) \
-double benchmark_copy_image_ ##T(void) \
+#define BENCH_COPY_IMAGE(J, T, K, M, Q) \
+double benchmark_ ##J ##_image_ ##T(void) \
 { \
   struct timeval start,stop; \
 \
@@ -16,7 +16,7 @@ double benchmark_copy_image_ ##T(void) \
   memset(&desc, 0x0, sizeof(cl_image_desc)); \
   memset(&format, 0x0, sizeof(cl_image_format)); \
 \
-  OCL_CREATE_KERNEL("bench_copy_image"); \
+  OCL_CREATE_KERNEL_FROM_FILE("bench_copy_image",K); \
   buf_data[0] = (uint32_t*) malloc(sizeof(M) * sz); \
   for (uint32_t i = 0; i < sz; ++i) { \
     ((M*)buf_data[0])[i] = rand(); \
@@ -63,8 +63,11 @@ double benchmark_copy_image_ ##T(void) \
   return (double)(1000 / (elapsed * 1e-3)); \
 } \
 \
-MAKE_BENCHMARK_FROM_FUNCTION_KEEP_PROGRAM(benchmark_copy_image_ ##T, true, "FPS");
+MAKE_BENCHMARK_FROM_FUNCTION_KEEP_PROGRAM(benchmark_ ##J ##_image_ ##T, true, "FPS");
 
-BENCH_COPY_IMAGE(uchar,unsigned char,CL_UNSIGNED_INT8)
-BENCH_COPY_IMAGE(ushort,unsigned short,CL_UNSIGNED_INT16)
-BENCH_COPY_IMAGE(uint,unsigned int,CL_UNSIGNED_INT32)
+BENCH_COPY_IMAGE(copy,uchar, "bench_copy_image", unsigned char, CL_UNSIGNED_INT8)
+BENCH_COPY_IMAGE(copy,ushort, "bench_copy_image", unsigned short, CL_UNSIGNED_INT16)
+BENCH_COPY_IMAGE(copy,uint, "bench_copy_image", unsigned int,CL_UNSIGNED_INT32)
+BENCH_COPY_IMAGE(filter,uchar, "bench_filter_image", unsigned char,CL_UNSIGNED_INT8)
+BENCH_COPY_IMAGE(filter,ushort, "bench_filter_image", unsigned short,CL_UNSIGNED_INT16)
+BENCH_COPY_IMAGE(filter,uint, "bench_filter_image", unsigned int,CL_UNSIGNED_INT32)
diff --git a/kernels/bench_copy_image.cl b/kernels/bench_copy_image.cl
index fb6d4f3..e6548f3 100644
--- a/kernels/bench_copy_image.cl
+++ b/kernels/bench_copy_image.cl
@@ -1,3 +1,4 @@
+const constant float filter_flag = 0.111111f;
 __kernel void
 bench_copy_image(__read_only image2d_t src, __write_only image2d_t dst)
 {
@@ -13,3 +14,39 @@ bench_copy_image(__read_only image2d_t src, __write_only image2d_t dst)
   color=read_imageui(src, sampler, coord);
   write_imageui(dst, coord, color);
 }
+
+__kernel void
+bench_filter_image(__read_only image2d_t src, __write_only image2d_t dst)
+{
+  float4 color = 0;
+  int2 coord_00, coord_01, coord_02, coord_10, coord_11, coord_12, coord_20, coord_21, coord_22;
+  int x = (int)get_global_id(0);
+  int y = (int)get_global_id(1);
+  int x_sz = (int)get_global_size(0);
+  int y_sz = (int)get_global_size(1);
+
+  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_CLAMP| CLK_FILTER_NEAREST;
+
+  int x0 = x - 1; int x1 = x + 1;
+  int y0 = y - 1; int y1 = y + 1 ;
+  int x_left = (x > 0)?x0:x; int x_right = (x > x_sz - 2)?x:x1;
+  int y_top = (y > 0)?y0:y; int y_bottom = (y > y_sz - 2)?y:y1;
+
+  coord_00.x = x_left;  coord_00.y = y_top;
+  coord_01.x = x; coord_01.y = y_top;
+  coord_02.x = x_right; coord_02.y = y_top;
+
+  coord_10.x = x_left; coord_10.y = y;
+  coord_11.x = x; coord_11.y = y;
+  coord_12.x = x_right; coord_12.y = y;
+
+  coord_20.x = x_left; coord_20.y = y_bottom;
+  coord_21.x = x; coord_21.y = y_bottom;
+  coord_22.x = x_right; coord_22.y = y_bottom;
+
+  color = convert_float4(read_imageui(src, sampler, coord_00)) + convert_float4(read_imageui(src, sampler, coord_01)) + convert_float4(read_imageui(src, sampler, coord_02))
+        + convert_float4(read_imageui(src, sampler, coord_10)) + convert_float4(read_imageui(src, sampler, coord_11)) + convert_float4(read_imageui(src, sampler, coord_12))
+        + convert_float4(read_imageui(src, sampler, coord_20)) + convert_float4(read_imageui(src, sampler, coord_21)) + convert_float4(read_imageui(src, sampler, coord_22));
+
+  write_imageui(dst, coord_11, convert_uint4(color * filter_flag));
+}
-- 
1.9.1