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

[Yang, Rong R]

Pushed.

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Song, Ruiling
> Sent: Wednesday, November 4, 2015 16:42
> To: Meng, Mengmeng; beignet at lists.freedesktop.org
> Cc: Meng, Mengmeng
> Subject: Re: [Beignet] [PATCH v2 1/2] add benckmark for copy data from
> buffer to buffer
> 
> 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_cop
> y_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
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet