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

Submitted by Meng, Mengmeng on Nov. 2, 2015, 6:29 a.m.

Details

Message ID 1446445761-11101-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. 2, 2015, 6:29 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 | 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

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

Comments

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

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

> Meng Mengmeng

> Sent: Monday, November 2, 2015 2:29 PM

> To: beignet@lists.freedesktop.org

> Cc: Meng, Mengmeng

> Subject: [Beignet] [PATCH 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 | 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; \


Please use same work-group setting as copy-image. So we can compare them.
  globals[0] = w; \
  globals[1] = h; \
  locals[0] = 16; \
  locals[1] = 4; \
Thanks your comments, I'll send v2.

-----Original Message-----
From: Song, Ruiling 

Sent: Wednesday, November 04, 2015 3:17 PM
To: Meng, Mengmeng; beignet@lists.freedesktop.org
Cc: Meng, Mengmeng
Subject: RE: [Beignet] [PATCH 1/2] add benckmark for copy data from buffer to buffer



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

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

> Of Meng Mengmeng

> Sent: Monday, November 2, 2015 2:29 PM

> To: beignet@lists.freedesktop.org

> Cc: Meng, Mengmeng

> Subject: [Beignet] [PATCH 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 | 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; \


Please use same work-group setting as copy-image. So we can compare them.
  globals[0] = w; \
  globals[1] = h; \
  locals[0] = 16; \
  locals[1] = 4; \