[Beignet] [PATCH 1/2] add benckmark for copy data from buffer to buffer

[Meng Mengmeng]

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

Signed-off-by: Meng Mengmeng <mengmeng.meng at intel.com>
---
 benchmark/CMakeLists.txt            |  3 ++-
 benchmark/benchmark_copy_buffer.cpp | 52 +++++++++++++++++++++++++++++++++++++
 kernels/bench_copy_buffer.cl        | 21 +++++++++++++++
 3 files changed, 75 insertions(+), 1 deletion(-)
 create mode 100644 benchmark/benchmark_copy_buffer.cpp
 create mode 100644 kernels/bench_copy_buffer.cl

diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt
index 3e43a21..03a56f2 100644
--- a/benchmark/CMakeLists.txt
+++ b/benchmark/CMakeLists.txt
@@ -16,7 +16,8 @@ set (benchmark_sources
   benchmark_read_buffer.cpp
   benchmark_read_image.cpp
   benchmark_copy_buffer_to_image.cpp
-  benchmark_copy_image_to_buffer.cpp)
+  benchmark_copy_image_to_buffer.cpp
+  benchmark_copy_buffer.cpp)
 
 
 SET(CMAKE_CXX_FLAGS "-DBUILD_BENCHMARK ${CMAKE_CXX_FLAGS}")
diff --git a/benchmark/benchmark_copy_buffer.cpp b/benchmark/benchmark_copy_buffer.cpp
new file mode 100644
index 0000000..68b9870
--- /dev/null
+++ b/benchmark/benchmark_copy_buffer.cpp
@@ -0,0 +1,52 @@
+#include "utests/utest_helper.hpp"
+#include <sys/time.h>
+
+#define BENCH_COPY_BUFFER(T, K, M) \
+double benchmark_copy_buffer_ ##T(void) \
+{ \
+  struct timeval start,stop; \
+ \
+  const size_t n = 1920 * 1080; \
+  const size_t sz = 4 * n ; \
+ \
+  OCL_CREATE_BUFFER(buf[0], 0, sz * sizeof(M), NULL); \
+  OCL_CREATE_BUFFER(buf[1], 0, sz * sizeof(M), NULL); \
+ \
+  OCL_CREATE_KERNEL_FROM_FILE("bench_copy_buffer",K); \
+ \
+  OCL_MAP_BUFFER(0); \
+  for (size_t i = 0; i < sz; i ++) { \
+    ((M *)(buf_data[0]))[i] = rand(); \
+  } \
+  OCL_UNMAP_BUFFER(0); \
+ \
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); \
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); \
+ \
+  globals[0] = 1920 * 1080; \
+  locals[0] = 64; \
+ \
+  gettimeofday(&start,0); \
+  for (size_t i=0; i<100; i++) { \
+    OCL_NDRANGE(1); \
+  } \
+  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_buffer_ ##T,true);
+
+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)
diff --git a/kernels/bench_copy_buffer.cl b/kernels/bench_copy_buffer.cl
new file mode 100644
index 0000000..e5bdb2e
--- /dev/null
+++ b/kernels/bench_copy_buffer.cl
@@ -0,0 +1,21 @@
+__kernel void
+bench_copy_buffer_uchar(__global uchar4* src, __global uchar4* dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] =  src[id];
+}
+
+__kernel void
+bench_copy_buffer_ushort(__global ushort4* src, __global ushort4* dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] =  src[id];
+}
+
+__kernel void
+bench_copy_buffer_uint(__global uint4* src, __global uint4* dst)
+{
+  int id = (int)get_global_id(0);
+  dst[id] =  src[id];
+}
+
-- 
1.9.1