[Beignet] [PATCH 3/3] add utest runtime_use_host_ptr_image

[Yang, Rong R]

The patchset LGTM, pushed.

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Guo Yejun
> Sent: Friday, September 25, 2015 4:28
> To: beignet at lists.freedesktop.org
> Cc: Guo, Yejun
> Subject: [Beignet] [PATCH 3/3] add utest runtime_use_host_ptr_image
> 
> Signed-off-by: Guo Yejun <yejun.guo at intel.com>
> ---
>  kernels/runtime_use_host_ptr_image.cl | 10 +++++
>  utests/CMakeLists.txt                 |  1 +
>  utests/runtime_use_host_ptr_image.cpp | 72
> +++++++++++++++++++++++++++++++++++
>  3 files changed, 83 insertions(+)
>  create mode 100644 kernels/runtime_use_host_ptr_image.cl
>  create mode 100644 utests/runtime_use_host_ptr_image.cpp
> 
> diff --git a/kernels/runtime_use_host_ptr_image.cl
> b/kernels/runtime_use_host_ptr_image.cl
> new file mode 100644
> index 0000000..7596ec8
> --- /dev/null
> +++ b/kernels/runtime_use_host_ptr_image.cl
> @@ -0,0 +1,10 @@
> +__kernel void
> +runtime_use_host_ptr_image(__read_only image2d_t src, __write_only
> +image2d_t dst) {
> +  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
> +CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
> +  int2 coord;
> +  coord.x = (int)get_global_id(0);
> +  coord.y = (int)get_global_id(1);
> +  float4 data = read_imagef(src, sampler, coord);
> +  write_imagef(dst, coord, data);
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index
> bfb902c..b26f4d0 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -213,6 +213,7 @@ set (utests_sources
>    vload_bench.cpp
>    runtime_use_host_ptr_buffer.cpp
>    runtime_alloc_host_ptr_buffer.cpp
> +  runtime_use_host_ptr_image.cpp
>    compiler_get_sub_group_size.cpp
>    compiler_get_sub_group_id.cpp
>    compiler_sub_group_shuffle.cpp)
> diff --git a/utests/runtime_use_host_ptr_image.cpp
> b/utests/runtime_use_host_ptr_image.cpp
> new file mode 100644
> index 0000000..2de9194
> --- /dev/null
> +++ b/utests/runtime_use_host_ptr_image.cpp
> @@ -0,0 +1,72 @@
> +#include "utest_helper.hpp"
> +#include <string.h>
> +
> +static void runtime_use_host_ptr_image(void) {
> +  const size_t w = 512;
> +  const size_t h = 512;
> +
> +  cl_image_format format;
> +  cl_image_desc desc;
> +
> +  memset(&desc, 0x0, sizeof(cl_image_desc));  memset(&format, 0x0,
> + sizeof(cl_image_format));
> +
> +  format.image_channel_order = CL_RGBA;
> format.image_channel_data_type
> + = CL_UNORM_INT8;  desc.image_type = CL_MEM_OBJECT_IMAGE2D;
> + desc.image_width = w;  desc.image_height = h;
> +
> +  //src image
> +  int ret = posix_memalign(&buf_data[0], 64, sizeof(uint32_t) * w * h);
> + OCL_ASSERT(ret == 0);  for (size_t i = 0; i < w*h; ++i)
> +    ((uint32_t*)buf_data[0])[i] = i;
> +
> +  OCL_CREATE_IMAGE(buf[0], CL_MEM_USE_HOST_PTR, &format, &desc,
> + buf_data[0]);
> +
> +  //dst image
> +  ret = posix_memalign(&buf_data[1], 64, sizeof(uint32_t) * w * h);
> + OCL_ASSERT(ret == 0);  for (size_t i = 0; i < w*h; ++i)
> +    ((uint32_t*)buf_data[1])[i] = 0;
> +
> +  OCL_CREATE_IMAGE(buf[1], CL_MEM_USE_HOST_PTR, &format, &desc,
> + buf_data[1]);
> +
> +  OCL_CREATE_KERNEL("runtime_use_host_ptr_image");
> +
> +  // Run the kernel
> +  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] = 16;  OCL_NDRANGE(2);
> +
> +  // Check result
> +  size_t origin[3];
> +  origin[0] = 0;
> +  origin[1] = 0;
> +  origin[2] = 0;
> +  size_t region[3];
> +  region[0] = w;
> +  region[1] = h;
> +  region[2] = 1;
> +  size_t pitch = 0;
> +  void* mapptr = (int*)clEnqueueMapImage(queue, buf[1], CL_TRUE,
> + CL_MAP_READ, origin, region, &pitch, NULL, 0, NULL, NULL, NULL);
> + OCL_ASSERT(mapptr == buf_data[1]);  for (uint32_t i = 0; i < w*h; ++i) {
> +    //printf("%d: src: 0x%x, dst: 0x%x\n", i, ((uint32_t*)buf_data[0])[i],
> ((uint32_t*)buf_data[1])[i]);
> +    OCL_ASSERT(((uint32_t*)buf_data[0])[i] ==
> + ((uint32_t*)buf_data[1])[i]);  }  clEnqueueUnmapMemObject(queue,
> + buf[1], mapptr, 0, NULL, NULL);
> +
> +  free(buf_data[0]);
> +  buf_data[0] = NULL;
> +  free(buf_data[1]);
> +  buf_data[1] = NULL;
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(runtime_use_host_ptr_image);
> --
> 1.9.1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet