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

[Meng Mengmeng]

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

Signed-off-by: Meng Mengmeng <mengmeng.meng at intel.com>
---
 benchmark/benchmark_copy_buffer.cpp | 15 ++++----
 kernels/bench_copy_buffer.cl        | 69 +++++++++++++++++++++++++++++++++++--
 2 files changed, 75 insertions(+), 9 deletions(-)

diff --git a/benchmark/benchmark_copy_buffer.cpp b/benchmark/benchmark_copy_buffer.cpp
index 93bf1b6..6cf3023 100644
--- a/benchmark/benchmark_copy_buffer.cpp
+++ b/benchmark/benchmark_copy_buffer.cpp
@@ -1,8 +1,8 @@
 #include "utests/utest_helper.hpp"
 #include <sys/time.h>
 
-#define BENCH_COPY_BUFFER(T, K, M) \
-double benchmark_copy_buffer_ ##T(void) \
+#define BENCH_COPY_BUFFER(J, T, K, M) \
+double benchmark_ ##J ##_buffer_ ##T(void) \
 { \
   struct timeval start,stop; \
  \
@@ -48,8 +48,11 @@ double benchmark_copy_buffer_ ##T(void) \
   return (double)(1000 / (elapsed * 1e-3)); \
 } \
  \
-MAKE_BENCHMARK_FROM_FUNCTION_KEEP_PROGRAM(benchmark_copy_buffer_ ##T, true, "FPS");
+MAKE_BENCHMARK_FROM_FUNCTION_KEEP_PROGRAM(benchmark_ ##J ##_buffer_ ##T, true, "FPS");
 
-BENCH_COPY_BUFFER(uchar,"bench_copy_buffer_uchar",unsigned char)
-BENCH_COPY_BUFFER(ushort,"bench_copy_buffer_ushort",unsigned short)
-BENCH_COPY_BUFFER(uint,"bench_copy_buffer_uint",unsigned int)
+BENCH_COPY_BUFFER(copy, uchar, "bench_copy_buffer_uchar", unsigned char)
+BENCH_COPY_BUFFER(copy, ushort, "bench_copy_buffer_ushort", unsigned short)
+BENCH_COPY_BUFFER(copy, uint, "bench_copy_buffer_uint", unsigned int)
+BENCH_COPY_BUFFER(filter, uchar, "bench_filter_buffer_uchar", unsigned char)
+BENCH_COPY_BUFFER(filter, ushort, "bench_filter_buffer_ushort", unsigned short)
+BENCH_COPY_BUFFER(filter, uint, "bench_filter_buffer_uint", unsigned int)
diff --git a/kernels/bench_copy_buffer.cl b/kernels/bench_copy_buffer.cl
index ed20352..8d8afd8 100644
--- a/kernels/bench_copy_buffer.cl
+++ b/kernels/bench_copy_buffer.cl
@@ -1,10 +1,11 @@
+const constant float filter_flag = 0.111111f;
 __kernel void
 bench_copy_buffer_uchar(__global uchar4* src, __global uchar4* dst)
 {
   int x = (int)get_global_id(0);
   int y = (int)get_global_id(1);
   int x_sz = (int)get_global_size(0);
-  dst[y * x_sz + x] =  src[y * x_sz + x];
+  dst[y * x_sz + x] = src[y * x_sz + x];
 }
 
 __kernel void
@@ -13,7 +14,7 @@ bench_copy_buffer_ushort(__global ushort4* src, __global ushort4* dst)
   int x = (int)get_global_id(0);
   int y = (int)get_global_id(1);
   int x_sz = (int)get_global_size(0);
-  dst[y * x_sz + x] =  src[y * x_sz + x];
+  dst[y * x_sz + x] = src[y * x_sz + x];
 }
 
 __kernel void
@@ -22,6 +23,68 @@ bench_copy_buffer_uint(__global uint4* src, __global uint4* dst)
   int x = (int)get_global_id(0);
   int y = (int)get_global_id(1);
   int x_sz = (int)get_global_size(0);
-  dst[y * x_sz + x] =  src[y * x_sz + x];
+  dst[y * x_sz + x] = src[y * x_sz + x];
 }
 
+__kernel void
+bench_filter_buffer_uchar(__global uchar4* src, __global uchar4* dst)
+{
+  float4 result;
+  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);
+
+  int x0 = x - 1; int x1 = x + 1;
+  int y0 = y - 1; int y1 = y + 1 ;
+  int x_left = (x0 > 0)?x0:x; int x_right = (x1 > x_sz - 1)?x:x1;
+  int y_top = (y0 > 0)?y0:y; int y_bottom = (y1 > y_sz - 1)?y:y1;
+
+  result = convert_float4(src[y_top * x_sz + x_left]) + convert_float4(src[y_top * x_sz + x]) + convert_float4(src[y_top * x_sz + x_right])
+         + convert_float4(src[y * x_sz + x_left]) + convert_float4(src[y * x_sz + x]) + convert_float4(src[y * x_sz + x_right])
+         + convert_float4(src[y_bottom * x_sz + x_left]) + convert_float4(src[y_bottom * x_sz + x]) + convert_float4(src[y_bottom * x_sz +x_right]);
+
+  dst[y * x_sz + x] = convert_uchar4(result * filter_flag);
+}
+
+__kernel void
+bench_filter_buffer_ushort(__global ushort4* src, __global ushort4* dst)
+{
+  float4 result;
+  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);
+
+  int x0 = x - 1; int x1 = x + 1;
+  int y0 = y - 1; int y1 = y + 1 ;
+  int x_left = (x0 > 0)?x0:x; int x_right = (x1 > x_sz - 1)?x:x1;
+  int y_top = (y0 > 0)?y0:y; int y_bottom = (y1 > y_sz - 1)?y:y1;
+
+  result = convert_float4(src[y_top * x_sz + x_left]) + convert_float4(src[y_top * x_sz + x]) + convert_float4(src[y_top * x_sz + x_right])
+         + convert_float4(src[y * x_sz + x_left]) + convert_float4(src[y * x_sz + x]) + convert_float4(src[y * x_sz + x_right])
+         + convert_float4(src[y_bottom * x_sz + x_left]) + convert_float4(src[y_bottom * x_sz + x]) + convert_float4(src[y_bottom * x_sz +x_right]);
+
+  dst[y * x_sz + x] = convert_ushort4(result * filter_flag);
+}
+
+__kernel void
+bench_filter_buffer_uint(__global uint4* src, __global uint4* dst)
+{
+  float4 result;
+  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);
+
+  int x0 = x - 1; int x1 = x + 1;
+  int y0 = y - 1; int y1 = y + 1 ;
+  int x_left = (x0 > 0)?x0:x; int x_right = (x1 > x_sz - 1)?x:x1;
+  int y_top = (y0 > 0)?y0:y; int y_bottom = (y1 > y_sz - 1)?y:y1;
+
+  result = convert_float4(src[y_top * x_sz + x_left]) + convert_float4(src[y_top * x_sz + x]) + convert_float4(src[y_top * x_sz + x_right])
+         + convert_float4(src[y * x_sz + x_left]) + convert_float4(src[y * x_sz + x]) + convert_float4(src[y * x_sz + x_right])
+         + convert_float4(src[y_bottom * x_sz + x_left]) + convert_float4(src[y_bottom * x_sz + x]) + convert_float4(src[y_bottom * x_sz +x_right]);
+
+  dst[y * x_sz + x] = convert_uint4(result * filter_flag);
+}
-- 
1.9.1