[Beignet] [PATCH] Add read_imagef benchmark for optimization.

[Yang, Rong R]

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> yan.wang at linux.intel.com
> Sent: Monday, September 5, 2016 14:52
> To: beignet at lists.freedesktop.org
> Cc: Yan Wang <yan.wang at linux.intel.com>
> Subject: [Beignet] [PATCH] Add read_imagef benchmark for optimization.
> 
> From: Yan Wang <yan.wang at linux.intel.com>
> 
> ---
>  benchmark/CMakeLists.txt                 |  1 +
>  benchmark/benchmark_read_image_float.cpp | 65
> ++++++++++++++++++++++++++++++++
>  kernels/compiler_read_image_float.cl     |  9 +++++
>  3 files changed, 75 insertions(+)
>  create mode 100644 benchmark/benchmark_read_image_float.cpp
>  create mode 100644 kernels/compiler_read_image_float.cl
> 
> diff --git a/benchmark/CMakeLists.txt b/benchmark/CMakeLists.txt index
> f9b246b..556275e 100644
> --- a/benchmark/CMakeLists.txt
> +++ b/benchmark/CMakeLists.txt
> @@ -15,6 +15,7 @@ set (benchmark_sources
>    benchmark_use_host_ptr_buffer.cpp
>    benchmark_read_buffer.cpp
>    benchmark_read_image.cpp
> +  benchmark_read_image_float.cpp
>    benchmark_copy_buffer_to_image.cpp
>    benchmark_copy_image_to_buffer.cpp
>    benchmark_copy_buffer.cpp
> diff --git a/benchmark/benchmark_read_image_float.cpp
> b/benchmark/benchmark_read_image_float.cpp
> new file mode 100644
> index 0000000..b0c2fb4
> --- /dev/null
> +++ b/benchmark/benchmark_read_image_float.cpp
> @@ -0,0 +1,65 @@
> +#include <string.h>
> +#include "utests/utest_helper.hpp"
> +#include <sys/time.h>
> +
> +#define NUM 400
> +
> +double benchmark_read_image_float(void) {
> +  struct timeval start,stop;
> +
> +  const size_t w = 128;
> +  const size_t h = 128;
> +  const size_t sz = w * h;
> +  cl_image_format format;
> +  cl_image_desc desc;
> +
> +  memset(&desc, 0x0, sizeof(cl_image_desc));  memset(&format, 0x0,
> + sizeof(cl_image_format));
> +
> +  // Setup kernel and images
> +  OCL_CREATE_KERNEL("compiler_read_image_float");
> +  buf_data[0] = (uint32_t*) malloc(sizeof(float) * sz);  for (uint32_t
> + i = 0; i < sz; ++i) {
> +    ((float*)buf_data[0])[i] = rand();
> +  }
> +
> +  format.image_channel_order = CL_R;
> +  format.image_channel_data_type = CL_FLOAT;  desc.image_type =
> + CL_MEM_OBJECT_IMAGE2D;  desc.image_width = w;  desc.image_height =
> h;
> + OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc,
> + buf_data[0]);  OCL_CREATE_BUFFER(buf[1], 0, sz * sizeof(float), NULL);
> +
> +  free(buf_data[0]);
> +  buf_data[0] = NULL;
> +
> +  // Run the kernel
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);  OCL_SET_ARG(1,
> + sizeof(cl_mem), &buf[1]);  OCL_SET_ARG(2, sizeof(cl_int), &w);
> + globals[0] = w;  globals[1] = h;  locals[0] = 16;  locals[1] = 16;
> +
> +  OCL_NDRANGE(2);
> +  OCL_FINISH();
> +
> +  gettimeofday(&start,0);
> +  for (size_t i=0; i<NUM; i++) {
> +    OCL_NDRANGE(2);
> +  }
> +  OCL_FINISH();
> +  gettimeofday(&stop,0);
> +
> +  free(buf_data[0]);
> +  buf_data[0] = NULL;
> +
> +  double elapsed = time_subtract(&stop, &start, 0);
> +
> +  return BANDWIDTH(sz * sizeof(float) * NUM, elapsed); }
> +
> +MAKE_BENCHMARK_FROM_FUNCTION(benchmark_read_image_float,
> "GB/S");
> diff --git a/kernels/compiler_read_image_float.cl
> b/kernels/compiler_read_image_float.cl
> new file mode 100644
> index 0000000..f581438
> --- /dev/null
> +++ b/kernels/compiler_read_image_float.cl
> @@ -0,0 +1,9 @@
> +__constant sampler_t sampler    = CLK_NORMALIZED_COORDS_FALSE |
> CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
Do you really need CLK_FILTER_LINEAR?
CLK_FILTER_LINEAR's image filter/cache's behavior is more complicated than CLK_FILTER_NEAREST. For benchmark, I think CLK_FILTER_NEAREST is better.

> +
> +__kernel void compiler_read_image_float(__read_only image2d_t src,
> +__global float* dst, int w) {
> +    int xi = get_global_id(0);
> +    int yi = get_global_id(1);
> +    float4 v = read_imagef(src, sampler, (float2)(xi, yi));
> +    *(dst + yi * w + xi) = v.x;
> +}
> --
> 1.9.1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet