Pushed.
> -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > Song, Ruiling > Sent: Wednesday, November 4, 2015 16:42 > To: Meng, Mengmeng; [email protected] > 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:[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_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 > > [email protected] > > http://lists.freedesktop.org/mailman/listinfo/beignet > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
