[Beignet] [PATCH 2/4] utests: refine image 1d buffer test case.
Yang, Rong R
rong.r.yang at intel.com
Wed Aug 26 23:51:11 PDT 2015
The patchset is pushed, with mirror changes.
GBE: a potential bug in instruction scheduling.
+ || node->insn.opcode == SEL_OP_ELSE
utests: refine image 1d buffer test case.
- int offset = x;
> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Zhigang Gong
> Sent: Thursday, August 13, 2015 10:24
> To: beignet at lists.freedesktop.org
> Cc: Gong, Zhigang
> Subject: [Beignet] [PATCH 2/4] utests: refine image 1d buffer test case.
>
> We need to test large image 1d buffer read and write testing.
>
> Signed-off-by: Zhigang Gong <zhigang.gong at intel.com>
> ---
> kernels/image_1D_buffer.cl | 11 ++----- utests/image_1D_buffer.cpp | 73
> ++++++++++++++++++----------------------------
> 2 files changed, 32 insertions(+), 52 deletions(-)
>
> diff --git a/kernels/image_1D_buffer.cl b/kernels/image_1D_buffer.cl index
> e8e0a86..4f982e6 100644
> --- a/kernels/image_1D_buffer.cl
> +++ b/kernels/image_1D_buffer.cl
> @@ -1,13 +1,8 @@
> -__kernel void image_1D_buffer(image1d_buffer_t image1, image1d_t
> image2, sampler_t sampler, __global int *results)
> +__kernel void image_1D_buffer(image1d_buffer_t image1,
> image1d_buffer_t
> +image2)
> {
> int x = get_global_id(0);
> int offset = x;
>
> - int4 col = read_imagei(image1, x);
> - int4 test = (col != read_imagei(image2, sampler, x));
> -
> - if (test.x || test.y || test.z || test.w)
> - results[offset] = 0;
> - else
> - results[offset] = 1;
> + uint4 color = read_imageui(image1, x);
> + write_imageui(image2, x, color);
> }
> diff --git a/utests/image_1D_buffer.cpp b/utests/image_1D_buffer.cpp
> index d8d761f..e2cfcde 100644
> --- a/utests/image_1D_buffer.cpp
> +++ b/utests/image_1D_buffer.cpp
> @@ -3,78 +3,63 @@
>
> void image_1D_buffer(void)
> {
> - size_t buffer_sz = 1024;
> - char *buf_content = (char *)malloc(buffer_sz * sizeof(char));
> + size_t buffer_sz = 65536;
> + char *buf_content = (char *)malloc(buffer_sz * sizeof(int));
> int error;
> cl_image_desc image_desc;
> cl_image_format image_format;
> - cl_sampler sampler;
> - cl_mem image1, image2;
> cl_mem ret_mem = NULL;
>
> OCL_CREATE_KERNEL("image_1D_buffer");
>
> for (int32_t i = 0; i < (int32_t)buffer_sz; ++i)
> - buf_content[i] = (rand() & 127);
> + buf_content[i] = (rand() & 0xFFFFFFFF);
>
> - cl_mem buff = clCreateBuffer(ctx, CL_MEM_READ_ONLY |
> CL_MEM_COPY_HOST_PTR,
> - buffer_sz, buf_content, &error);
> - OCL_ASSERT(error == CL_SUCCESS);
> + OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_ONLY |
> CL_MEM_COPY_HOST_PTR,
> + buffer_sz * sizeof(int), buf_content); OCL_CREATE_BUFFER(buf[1],
> + CL_MEM_READ_WRITE, buffer_sz * sizeof(int), NULL);
>
> memset(&image_desc, 0x0, sizeof(cl_image_desc));
> memset(&image_format, 0x0, sizeof(cl_image_format));
>
> image_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
> - image_desc.image_row_pitch = buffer_sz;
> - image_desc.image_width = buffer_sz / sizeof(uint32_t); //assume rgba32
> - image_desc.buffer = buff;
> -
> - image_format.image_channel_order = CL_RGBA;
> - image_format.image_channel_data_type = CL_UNSIGNED_INT8;
> -
> - image1 = clCreateImage(ctx, CL_MEM_READ_ONLY, &image_format,
> - &image_desc, NULL, &error );
> - OCL_ASSERT(error == CL_SUCCESS);
> -
> - error = clGetImageInfo(image1, CL_IMAGE_BUFFER, sizeof(ret_mem),
> &ret_mem, NULL);
> - OCL_ASSERT(error == CL_SUCCESS);
> - OCL_ASSERT(ret_mem == buff);
> + image_desc.image_row_pitch = buffer_sz * sizeof(int);
> + image_desc.image_width = buffer_sz; //assume r32 image_desc.buffer =
> + buf[0];
>
> + image_format.image_channel_order = CL_R;
> + image_format.image_channel_data_type = CL_UNSIGNED_INT32;
>
> - memset(&image_desc, 0x0, sizeof(cl_image_desc));
> - image_desc.image_type = CL_MEM_OBJECT_IMAGE1D;
> - image_desc.image_width = buffer_sz / sizeof(uint32_t);
> - image2 = clCreateImage(ctx,
> CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
> - &image_format, &image_desc, buf_content, &error);
> + // Create the source image1d_buffer.
> + OCL_CREATE_IMAGE(buf[2], CL_MEM_READ_ONLY, &image_format,
> + &image_desc, NULL); error = clGetImageInfo(buf[2], CL_IMAGE_BUFFER,
> + sizeof(ret_mem), &ret_mem, NULL);
> OCL_ASSERT(error == CL_SUCCESS);
> + OCL_ASSERT(ret_mem == buf[0]);
>
> - // Create sampler to use
> - sampler = clCreateSampler(ctx, false, CL_ADDRESS_NONE,
> CL_FILTER_NEAREST, &error );
> + // Create the destination image1d_buffer.
> + image_desc.buffer = buf[1];
> + OCL_CREATE_IMAGE(buf[3], CL_MEM_READ_ONLY, &image_format,
> + &image_desc, NULL); error = clGetImageInfo(buf[3], CL_IMAGE_BUFFER,
> + sizeof(ret_mem), &ret_mem, NULL);
> OCL_ASSERT(error == CL_SUCCESS);
> + OCL_ASSERT(ret_mem == buf[1]);
>
> - cl_mem result_buf = buf[0] = clCreateBuffer(ctx, 0, buffer_sz, NULL,
> &error);
> - OCL_ASSERT(error == CL_SUCCESS);
> + OCL_SET_ARG(0, sizeof(cl_mem), &buf[2]); OCL_SET_ARG(1,
> + sizeof(cl_mem), &buf[3]);
>
> - OCL_SET_ARG(0, sizeof(cl_mem), &image1);
> - OCL_SET_ARG(1, sizeof(cl_mem), &image2);
> - OCL_SET_ARG(2, sizeof(sampler), &sampler);
> - OCL_SET_ARG(3, sizeof(cl_mem), &result_buf);
> -
> - globals[0] = buffer_sz/sizeof(int32_t);
> + globals[0] = buffer_sz;
> locals[0] = 16;
>
> OCL_NDRANGE(1);
>
> /* Now check the result. */
> OCL_MAP_BUFFER(0);
> - for (uint32_t i = 0; i < buffer_sz/sizeof(int32_t); i++)
> - OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 1);
> + OCL_MAP_BUFFER(1);
> + for (uint32_t i = 0; i < buffer_sz; i++) {
> + if (((uint32_t*)buf_data[1])[i] != ((uint32_t*)buf_data[0])[i])
> + printf("i %d expected %x got %x \n", i, ((uint32_t*)buf_data[0])[i],
> ((uint32_t*)buf_data[1])[i]);
> + OCL_ASSERT(((uint32_t*)buf_data[1])[i] ==
> + ((uint32_t*)buf_data[0])[i]); }
> OCL_UNMAP_BUFFER(0);
> -
> - clReleaseSampler(sampler);
> - clReleaseMemObject(image1);
> - clReleaseMemObject(image2);
> - clReleaseMemObject(buff);
> + OCL_UNMAP_BUFFER(1);
> }
>
> MAKE_UTEST_FROM_FUNCTION(image_1D_buffer);
> --
> 1.9.1
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list