[Beignet] [PATCH 2/2] utests: Add test case for box blur on image buffer
Zhigang Gong
zhigang.gong at linux.intel.com
Mon May 27 01:07:27 PDT 2013
Dag,
I gave a try of this test case, and I reproduce the failure. I also did
another test to apply this patch in addition on yout patch:
diff --git a/kernels/compiler_box_blur_image.cl b/kernels/compiler_box_blur_image.cl
index 7bcbdeb..d114cb4 100644
--- a/kernels/compiler_box_blur_image.cl
+++ b/kernels/compiler_box_blur_image.cl
@@ -1,16 +1,21 @@
__kernel void compiler_box_blur_image(__read_only image2d_t src,
__write_only image2d_t dst)
{
- const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+ const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;
const int2 coord = (int2)(get_global_id(0), get_global_id(1));
- int2 offset;
+ const float2 ncoord;
+ ncoord.x = coord.x / 128.0f;
+ ncoord.y = coord.y / 128.0f;
+ float2 offset;
float4 sum = 0;
for (offset.y = -1; offset.y <= 1; offset.y++) {
for (offset.x = -1; offset.x <= 1; offset.x++) {
- sum += read_imagef(src, sampler, coord + offset);
+ float2 scoord;
+ scoord = ncoord + offset / 128.0f;
+ sum += read_imagef(src, sampler, scoord);
}
}
which is to change to use normalized coords for the sampler. And it could pass
the test case. Currently, I'm not very sure whether this is a Gen platform limiation
or it's normal according to Open CL spec.
cc to Nanhai, do you have any comments here? Why only when we use normalized coords in the
sampler, we can get an expected sampling result (here, we are using clamp to edge.).
On Sat, May 25, 2013 at 10:14:17AM +0200, Dag Lem wrote:
> This test demonstrates box blur on an image buffer, using an extremely
> simple kernel.
>
> Signed-off-by: Dag Lem <dag at nimrod.no>
> ---
> kernels/compiler_box_blur_image.cl | 18 +++++++++++++
> utests/CMakeLists.txt | 1 +
> utests/compiler_box_blur_image.cpp | 52 ++++++++++++++++++++++++++++++++++++++
> utests/utest_helper.hpp | 6 +++++
> 4 files changed, 77 insertions(+)
> create mode 100644 kernels/compiler_box_blur_image.cl
> create mode 100644 utests/compiler_box_blur_image.cpp
>
> diff --git a/kernels/compiler_box_blur_image.cl b/kernels/compiler_box_blur_image.cl
> new file mode 100644
> index 0000000..7bcbdeb
> --- /dev/null
> +++ b/kernels/compiler_box_blur_image.cl
> @@ -0,0 +1,18 @@
> +__kernel void compiler_box_blur_image(__read_only image2d_t src,
> + __write_only image2d_t dst)
> +{
> + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
> + CLK_ADDRESS_CLAMP_TO_EDGE |
> + CLK_FILTER_NEAREST;
> + const int2 coord = (int2)(get_global_id(0), get_global_id(1));
> + int2 offset;
> + float4 sum = 0;
> +
> + for (offset.y = -1; offset.y <= 1; offset.y++) {
> + for (offset.x = -1; offset.x <= 1; offset.x++) {
> + sum += read_imagef(src, sampler, coord + offset);
> + }
> + }
> +
> + write_imagef(dst, coord, (1.0f/9.0f)*sum);
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 63c873d..8700761 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -9,6 +9,7 @@ set (utests_sources
> compiler_mandelbrot.cpp
> compiler_mandelbrot_alternate.cpp
> compiler_box_blur_float.cpp
> + compiler_box_blur_image.cpp
> compiler_box_blur.cpp
> compiler_insert_to_constant.cpp
> compiler_argument_structure.cpp
> diff --git a/utests/compiler_box_blur_image.cpp b/utests/compiler_box_blur_image.cpp
> new file mode 100644
> index 0000000..d94a97c
> --- /dev/null
> +++ b/utests/compiler_box_blur_image.cpp
> @@ -0,0 +1,52 @@
> +#include "utest_helper.hpp"
> +
> +static void compiler_box_blur_image()
> +{
> + int w, h;
> + cl_image_format format = { };
> + cl_image_desc desc = { };
> + size_t origin[3] = { };
> + size_t region[3];
> + int *src, *dst;
> +
> + OCL_CREATE_KERNEL("compiler_box_blur_image");
> +
> + /* Load the picture */
> + src = cl_read_bmp("lenna128x128.bmp", &w, &h);
> +
> + 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;
> + desc.image_depth = 1;
> + desc.image_row_pitch = w*sizeof(uint32_t);
> +
> + /* Run the kernel */
> + OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, src);
> + free(src);
> + desc.image_row_pitch = 0;
> + OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, 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] = 16;
> + OCL_NDRANGE(2);
> + dst = (int*)malloc(w*h*sizeof(uint32_t));
> + region[0] = w;
> + region[1] = h;
> + region[2] = 1;
> + OCL_READ_IMAGE(buf[1], origin, region, dst);
> +
> + /* Save the image (for debug purpose) */
> + cl_write_bmp(dst, w, h, "compiler_box_blur_image.bmp");
> +
> + /* Compare with the golden image */
> + OCL_CHECK_IMAGE(dst, w, h, "compiler_box_blur_ref.bmp");
> +
> + free(dst);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_box_blur_image);
> diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp
> index d882fc7..c28615e 100644
> --- a/utests/utest_helper.hpp
> +++ b/utests/utest_helper.hpp
> @@ -88,6 +88,12 @@ extern EGLSurface eglSurface;
> #define OCL_CREATE_IMAGE(IMAGE, FLAGS, FORMAT, DESC, DATA) \
> OCL_CALL2(clCreateImage, IMAGE, ctx, FLAGS, FORMAT, DESC, DATA)
>
> +#define OCL_READ_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
> + OCL_CALL(clEnqueueReadImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
> +
> +#define OCL_WRITE_IMAGE(IMAGE, ORIGIN, REGION, DATA) \
> + OCL_CALL(clEnqueueWriteImage, queue, IMAGE, CL_TRUE, ORIGIN, REGION, 0, 0, DATA, 0, NULL, NULL)
> +
> #define OCL_CREATE_GL_IMAGE(IMAGE, FLAGS, TARGET, LEVEL, TEXTURE) \
> OCL_CALL2(clCreateFromGLTexture, IMAGE, ctx, FLAGS, TARGET, LEVEL, TEXTURE)
>
> --
> 1.8.1.4
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list