Set the data format as 1920 * 1080 four channels(RGBA) and type as char,short and int.
Signed-off-by: Meng Mengmeng <[email protected]> --- benchmark/CMakeLists.txt | 3 +- benchmark/benchmark_copy_image.cpp | 70 ++++++++++++++++++++++++++++++++++++++ kernels/bench_copy_image.cl | 15 ++++++++ 3 files changed, 87 insertions(+), 1 deletion(-) create mode 100644 benchmark/benchmark_copy_image.cpp create mode 100644 kernels/bench_copy_image.cl diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt index 03a56f2..dd33829 100644 --- a/benchmark/CMakeLists.txt +++ b/benchmark/CMakeLists.txt @@ -17,7 +17,8 @@ set (benchmark_sources benchmark_read_image.cpp benchmark_copy_buffer_to_image.cpp benchmark_copy_image_to_buffer.cpp - benchmark_copy_buffer.cpp) + benchmark_copy_buffer.cpp + benchmark_copy_image.cpp) SET(CMAKE_CXX_FLAGS "-DBUILD_BENCHMARK ${CMAKE_CXX_FLAGS}") diff --git a/benchmark/benchmark_copy_image.cpp b/benchmark/benchmark_copy_image.cpp new file mode 100644 index 0000000..92dffc9 --- /dev/null +++ b/benchmark/benchmark_copy_image.cpp @@ -0,0 +1,70 @@ +#include <string.h> +#include "utests/utest_helper.hpp" +#include <sys/time.h> + +#define BENCH_COPY_IMAGE(T, M, Q) \ +double benchmark_copy_image_ ##T(void) \ +{ \ + struct timeval start,stop; \ +\ + const size_t w = 1920; \ + const size_t h = 1080; \ + const size_t sz = 4 * w * h; \ + cl_image_format format; \ + cl_image_desc desc; \ +\ + memset(&desc, 0x0, sizeof(cl_image_desc)); \ + memset(&format, 0x0, sizeof(cl_image_format)); \ +\ + OCL_CREATE_KERNEL("bench_copy_image"); \ + buf_data[0] = (uint32_t*) malloc(sizeof(M) * sz); \ + for (uint32_t i = 0; i < sz; ++i) { \ + ((M*)buf_data[0])[i] = rand(); \ + } \ +\ + format.image_channel_order = CL_RGBA; \ + format.image_channel_data_type = Q; \ + desc.image_type = CL_MEM_OBJECT_IMAGE2D; \ + desc.image_width = w; \ + desc.image_height = h; \ + desc.image_row_pitch = desc.image_width * sizeof(M) * 4; \ + OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, buf_data[0]); \ +\ + desc.image_row_pitch = 0; \ + OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL); \ +\ + free(buf_data[0]); \ + buf_data[0] = NULL; \ +\ + 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_image_ ##T,true); + +BENCH_COPY_IMAGE(uchar,unsigned char,CL_UNSIGNED_INT8) +BENCH_COPY_IMAGE(ushort,unsigned short,CL_UNSIGNED_INT16) +BENCH_COPY_IMAGE(uint,unsigned int,CL_UNSIGNED_INT32) diff --git a/kernels/bench_copy_image.cl b/kernels/bench_copy_image.cl new file mode 100644 index 0000000..fb6d4f3 --- /dev/null +++ b/kernels/bench_copy_image.cl @@ -0,0 +1,15 @@ +__kernel void +bench_copy_image(__read_only image2d_t src, __write_only image2d_t dst) +{ + uint4 color = 0; + int2 coord; + int x = (int)get_global_id(0); + int y = (int)get_global_id(1); + + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_CLAMP| CLK_FILTER_NEAREST; + + coord.x = x; + coord.y = y; + color=read_imageui(src, sampler, coord); + write_imageui(dst, coord, color); +} -- 1.9.1 _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
