Set the data format as 1920 * 1080 four channels and type as char,short and int.
Signed-off-by: Meng Mengmeng <mengmeng.m...@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; \ + \ + 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]; +} + -- 1.9.1 _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet