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

Submitted by Meng, Mengmeng on Nov. 4, 2015, 8:16 a.m.

Details

Message ID 1446625018-5206-1-git-send-email-mengmeng.meng@intel.com
State New
Headers show
Series "Series without cover letter" ( rev: 1 ) in Beignet

Not browsing as part of any series.

Commit Message

Meng, Mengmeng Nov. 4, 2015, 8:16 a.m.
Set the data format as 1920 * 1080 four channels and type as char,short and int.

Signed-off-by: Meng Mengmeng <mengmeng.meng@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

Patch hide | download patch | download mbox

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_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..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];
+}
+

Comments

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

Thanks!
Ruiling

> -----Original Message-----

> From: Beignet [mailto:beignet-bounces@lists.freedesktop.org] On Behalf Of

> Meng Mengmeng

> Sent: Wednesday, November 4, 2015 4:17 PM

> To: beignet@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@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@lists.freedesktop.org

> http://lists.freedesktop.org/mailman/listinfo/beignet
Pushed.

> -----Original Message-----

> From: Beignet [mailto:beignet-bounces@lists.freedesktop.org] On Behalf Of

> Song, Ruiling

> Sent: Wednesday, November 4, 2015 16:42

> To: Meng, Mengmeng; beignet@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@lists.freedesktop.org] On Behalf

> > Of Meng Mengmeng

> > Sent: Wednesday, November 4, 2015 4:17 PM

> > To: beignet@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@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@lists.freedesktop.org

> > http://lists.freedesktop.org/mailman/listinfo/beignet

> _______________________________________________

> Beignet mailing list

> Beignet@lists.freedesktop.org

> http://lists.freedesktop.org/mailman/listinfo/beignet