This patchset LGTM. This benchmark is important to compare buffer/image performance on different generations.
Thanks! Ruiling > -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > Meng Mengmeng > Sent: Wednesday, November 4, 2015 4:17 PM > To: [email protected] > 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 <[email protected]> > --- > 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 > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
