One minor comment.
On Fri, May 16, 2014 at 02:52:52PM +0800, junyan.he at inbox.com wrote:
> From: Junyan He <junyan.he at linux.intel.com>
>
> Signed-off-by: Junyan He <junyan.he at linux.intel.com>
> ---
> kernels/image_1D_buffer.cl | 14 ++++++++
> utests/CMakeLists.txt | 1 +
> utests/image_1D_buffer.cpp | 81 ++++++++++++++++++++++++++++++++++++++++++++++
> 3 files changed, 96 insertions(+)
> create mode 100644 kernels/image_1D_buffer.cl
> create mode 100644 utests/image_1D_buffer.cpp
>
> diff --git a/kernels/image_1D_buffer.cl b/kernels/image_1D_buffer.cl
> new file mode 100644
> index 0000000..6e84dc4
> --- /dev/null
> +++ b/kernels/image_1D_buffer.cl
> @@ -0,0 +1,14 @@
> +__kernel void image_1D_buffer(image1d_buffer_t image1, image1d_t image2, sampler_t sampler, __global int *results)
> +{
> + 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;
> +}
> +
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 91e4ffb..e984048 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -178,6 +178,7 @@ set (utests_sources
> enqueue_copy_buf.cpp
> enqueue_copy_buf_unaligned.cpp
> enqueue_fill_buf.cpp
> + image_1D_buffer.cpp
> utest_assert.cpp
> utest.cpp
> utest_file_map.cpp
> diff --git a/utests/image_1D_buffer.cpp b/utests/image_1D_buffer.cpp
> new file mode 100644
> index 0000000..0045d91
> --- /dev/null
> +++ b/utests/image_1D_buffer.cpp
> @@ -0,0 +1,81 @@
> +#include <string.h>
> +#include "utest_helper.hpp"
> +
> +void image_1D_buffer(void)
> +{
> + size_t buffer_sz = 1024;
> + char *buf_content = (char *)malloc(buffer_sz * sizeof(char));
> + 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);
> +
> + cl_mem buff = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
> + buffer_sz, buf_content, &error);
> + OCL_ASSERT(error == CL_SUCCESS);
> +
> + 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);
> +
> +
> + 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);
> + OCL_ASSERT(error == CL_SUCCESS);
> +
> + // Create sampler to use
> + sampler = clCreateSampler(ctx, false, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
> + OCL_ASSERT(error == CL_SUCCESS);
> +
> + cl_mem result_buf = buf[0] = clCreateBuffer(ctx, 0, buffer_sz/sizeof(int32_t), NULL, &error);
> + OCL_ASSERT(error == CL_SUCCESS);
> +
> + 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);
> + 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_UNMAP_BUFFER(0);
> +
> + clReleaseSampler(sampler);
> + clReleaseMemObject(result_buf);
You should not release result_buf manually, as you put it to buf[0]
which is handled by the helper routines. No need to submit new version.
I will fix it, thanks.
> + clReleaseMemObject(image1);
> + clReleaseMemObject(image2);
> + clReleaseMemObject(buff);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(image_1D_buffer);
> --
> 1.8.3.2
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet