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

[Song, Ruiling]

This patchset LGTM. This benchmark is important to compare buffer/image
performance on different generations.

Thanks!
Ruiling

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Meng Mengmeng
> Sent: Wednesday, November 4, 2015 4:17 PM
> To: beignet at lists.freedesktop.org
> Cc: Meng, Mengmeng
> Subject: [Beignet] [PATCH v2 1/2] add benckmark for copy data from buffer to
> buffer
> 
> 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 | 55
> +++++++++++++++++++++++++++++++++++++
>  kernels/bench_copy_buffer.cl        | 27 ++++++++++++++++++
>  3 files changed, 84 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..88983a7
> --- /dev/null
> +++ b/benchmark/benchmark_copy_buffer.cpp
> @@ -0,0 +1,55 @@
> +#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 w = 1920; \
> +  const size_t h = 1080; \
> +  const size_t sz = 4 * w * h; \
> + \
> +  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] = 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_bu
> ffer_ ##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..ed20352
> --- /dev/null
> +++ b/kernels/bench_copy_buffer.cl
> @@ -0,0 +1,27 @@
> +__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];
> +}
> +
> +__kernel void
> +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];
> +}
> +
> +__kernel void
> +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];
> +}
> +
> --
> 1.9.1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet