[Beignet] [PATCH v2 2/2] add benckmark for copy data from image to image

[Meng Mengmeng]

Set the data format as 1920 * 1080 four channels(RGBA) and type as char,short and int.

Signed-off-by: Meng Mengmeng <mengmeng.meng at intel.com>
---
 benchmark/CMakeLists.txt           |  3 +-
 benchmark/benchmark_copy_image.cpp | 70 ++++++++++++++++++++++++++++++++++++++
 kernels/bench_copy_image.cl        | 15 ++++++++
 3 files changed, 87 insertions(+), 1 deletion(-)
 create mode 100644 benchmark/benchmark_copy_image.cpp
 create mode 100644 kernels/bench_copy_image.cl

diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt
index 03a56f2..dd33829 100644
--- a/benchmark/CMakeLists.txt
+++ b/benchmark/CMakeLists.txt
@@ -17,7 +17,8 @@ set (benchmark_sources
   benchmark_read_image.cpp
   benchmark_copy_buffer_to_image.cpp
   benchmark_copy_image_to_buffer.cpp
-  benchmark_copy_buffer.cpp)
+  benchmark_copy_buffer.cpp
+  benchmark_copy_image.cpp)
 
 
 SET(CMAKE_CXX_FLAGS "-DBUILD_BENCHMARK ${CMAKE_CXX_FLAGS}")
diff --git a/benchmark/benchmark_copy_image.cpp b/benchmark/benchmark_copy_image.cpp
new file mode 100644
index 0000000..92dffc9
--- /dev/null
+++ b/benchmark/benchmark_copy_image.cpp
@@ -0,0 +1,70 @@
+#include <string.h>
+#include "utests/utest_helper.hpp"
+#include <sys/time.h>
+
+#define BENCH_COPY_IMAGE(T, M, Q) \
+double benchmark_copy_image_ ##T(void) \
+{ \
+  struct timeval start,stop; \
+\
+  const size_t w = 1920; \
+  const size_t h = 1080; \
+  const size_t sz = 4 * w * h; \
+  cl_image_format format; \
+  cl_image_desc desc; \
+\
+  memset(&desc, 0x0, sizeof(cl_image_desc)); \
+  memset(&format, 0x0, sizeof(cl_image_format)); \
+\
+  OCL_CREATE_KERNEL("bench_copy_image"); \
+  buf_data[0] = (uint32_t*) malloc(sizeof(M) * sz); \
+  for (uint32_t i = 0; i < sz; ++i) { \
+    ((M*)buf_data[0])[i] = rand(); \
+  } \
+\
+  format.image_channel_order = CL_RGBA; \
+  format.image_channel_data_type = Q; \
+  desc.image_type = CL_MEM_OBJECT_IMAGE2D; \
+  desc.image_width = w; \
+  desc.image_height = h; \
+  desc.image_row_pitch = desc.image_width * sizeof(M) * 4; \
+  OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, buf_data[0]); \
+\
+  desc.image_row_pitch = 0; \
+  OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL); \
+\
+  free(buf_data[0]); \
+  buf_data[0] = 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] = 4; \
+\
+  gettimeofday(&start,0); \
+  for (size_t i=0; i<100; i++) { \
+    OCL_NDRANGE(2); \
+  } \
+  OCL_FINISH(); \
+\
+  OCL_MAP_BUFFER(1); \
+  OCL_UNMAP_BUFFER(1); \
+  gettimeofday(&stop,0); \
+\
+  clReleaseMemObject(buf[0]); \
+  free(buf_data[0]); \
+  buf_data[0] = NULL; \
+\
+  double elapsed = time_subtract(&stop, &start, 0); \
+\
+  return BANDWIDTH(sz * sizeof(M)*2 * 100, elapsed); \
+} \
+\
+MAKE_BENCHMARK_FROM_FUNCTION_KEEP_PROGRAM(benchmark_copy_image_ ##T,true);
+
+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)
diff --git a/kernels/bench_copy_image.cl b/kernels/bench_copy_image.cl
new file mode 100644
index 0000000..fb6d4f3
--- /dev/null
+++ b/kernels/bench_copy_image.cl
@@ -0,0 +1,15 @@
+__kernel void
+bench_copy_image(__read_only image2d_t src, __write_only image2d_t dst)
+{
+  uint4 color = 0;
+  int2 coord;
+  int x = (int)get_global_id(0);
+  int y = (int)get_global_id(1);
+
+  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_CLAMP| CLK_FILTER_NEAREST;
+
+  coord.x = x;
+  coord.y = y;
+  color=read_imageui(src, sampler, coord);
+  write_imageui(dst, coord, color);
+}
-- 
1.9.1