[Beignet] [PATCH] Implement cl_khr_image2d_from_buffer extension.
Zhigang Gong
zhigang.gong at linux.intel.com
Tue Dec 2 18:15:04 PST 2014
Thanks for the patch. But This implementation may hurt
performance. The major reason is that the image has some
special layout requirement and if we want to create an
image from eaxct the pitch * height buffer object, we
have to introduce an extra copy here for all cases.
This makes things even worse if the buffer is created
by CL_MEM_USE_HOST_PTR.
Before we find a way to eliminate the above overhead for this
extension, I prefer to not include it in beignet.
What do you think?
On Tue, Nov 25, 2014 at 07:07:13PM +0800, Yan Wang wrote:
> Implement cl_khr_image2d_from_buffer extension.
>
> ---
> CMakeLists.txt | 2 +
> kernels/image_2D_buffer.cl | 15 +++++
> src/CMakeLists.txt | 5 ++
> src/cl_api.c | 9 +++
> src/cl_device_id.c | 4 ++
> src/cl_device_id.h | 4 ++
> src/cl_extensions.c | 2 +-
> src/cl_gt_device.h | 4 ++
> src/cl_mem.c | 156 +++++++++++++++++++++++++++++++++++++++++++--
> utests/CMakeLists.txt | 6 ++
> utests/image_2D_buffer.cpp | 89 ++++++++++++++++++++++++++
> 11 files changed, 290 insertions(+), 6 deletions(-)
> create mode 100644 kernels/image_2D_buffer.cl
> create mode 100644 utests/image_2D_buffer.cpp
>
> diff --git a/CMakeLists.txt b/CMakeLists.txt
> index 49c8929..5ca7d90 100644
> --- a/CMakeLists.txt
> +++ b/CMakeLists.txt
> @@ -189,6 +189,8 @@ ELSE(OCLIcd_FOUND)
> MESSAGE(STATUS "Looking for OCL ICD header file - not found")
> ENDIF(OCLIcd_FOUND)
>
> +SET (OCL_IMAGE2D_BUFFER "true")
> +
> Find_Package(PythonInterp)
>
> ADD_SUBDIRECTORY(include)
> diff --git a/kernels/image_2D_buffer.cl b/kernels/image_2D_buffer.cl
> new file mode 100644
> index 0000000..6b9060c
> --- /dev/null
> +++ b/kernels/image_2D_buffer.cl
> @@ -0,0 +1,15 @@
> +__kernel void image_2D_buffer(image2d_t image1, image2d_t image2, sampler_t sampler, __global int *results)
> +{
> + int x = get_global_id(0);
> + int y = get_global_id(1);
> + int w = get_image_width(image1);
> + int offset = mad24(y, w, x);
> +
> + int4 pix = read_imagei(image1, (int2)(x, y));
> + int4 test = (pix != read_imagei(image2, sampler, (int2)(x, y)));
> +
> + if (test.x || test.y || test.z || test.w)
> + results[offset] = 0;
> + else
> + results[offset] = 1;
> +}
> diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
> index 7182bad..3ca5f1f 100644
> --- a/src/CMakeLists.txt
> +++ b/src/CMakeLists.txt
> @@ -114,6 +114,11 @@ SET(CMAKE_CXX_FLAGS "-DHAS_USERPTR ${CMAKE_CXX_FLAGS}")
> SET(CMAKE_C_FLAGS "-DHAS_USERPTR ${CMAKE_C_FLAGS}")
> endif (DRM_INTEL_USERPTR)
>
> +if (OCL_IMAGE2D_BUFFER)
> +SET(CMAKE_CXX_FLAGS "-DHAS_OCLImage2dBuffer ${CMAKE_CXX_FLAGS}")
> +SET(CMAKE_C_FLAGS "-DHAS_OCLImage2dBuffer ${CMAKE_C_FLAGS}")
> +endif (OCL_IMAGE2D_BUFFER)
> +
> set(GIT_SHA1 "git_sha1.h")
> add_custom_target(${GIT_SHA1} ALL
> COMMAND chmod +x ${CMAKE_CURRENT_SOURCE_DIR}/git_sha1.sh
> diff --git a/src/cl_api.c b/src/cl_api.c
> index 972c687..04095a2 100644
> --- a/src/cl_api.c
> +++ b/src/cl_api.c
> @@ -548,6 +548,14 @@ clCreateImage(cl_context context,
> err = CL_INVALID_IMAGE_DESCRIPTOR;
> goto error;
> }
> +#ifdef HAS_OCLImage2dBuffer
> + if ((image_desc->image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER &&
> + image_desc->image_type != CL_MEM_OBJECT_IMAGE2D) &&
> + image_desc->buffer) {
> + err = CL_INVALID_IMAGE_DESCRIPTOR;
> + goto error;
> + }
> +#else
> /* buffer refers to a valid buffer memory object if image_type is
> CL_MEM_OBJECT_IMAGE1D_BUFFER. Otherwise it must be NULL. */
> if (image_desc->image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER &&
> @@ -555,6 +563,7 @@ clCreateImage(cl_context context,
> err = CL_INVALID_IMAGE_DESCRIPTOR;
> goto error;
> }
> +#endif
> if (image_desc->num_mip_levels || image_desc->num_samples) {
> err = CL_INVALID_IMAGE_DESCRIPTOR;
> goto error;
> diff --git a/src/cl_device_id.c b/src/cl_device_id.c
> index 5ef0bde..c47d48c 100644
> --- a/src/cl_device_id.c
> +++ b/src/cl_device_id.c
> @@ -571,6 +571,10 @@ cl_get_device_info(cl_device_id device,
> DECL_FIELD(PARTITION_AFFINITY_DOMAIN, affinity_domain)
> DECL_FIELD(PARTITION_TYPE, partition_type)
> DECL_FIELD(REFERENCE_COUNT, device_reference_count)
> +#ifdef HAS_OCLImage2dBuffer
> + DECL_FIELD(IMAGE_PITCH_ALIGNMENT, image_pitch_alignment)
> + DECL_FIELD(IMAGE_BASE_ADDRESS_ALIGNMENT, image_base_address_alignment)
> +#endif
>
> case CL_DRIVER_VERSION:
> if (param_value_size_ret) {
> diff --git a/src/cl_device_id.h b/src/cl_device_id.h
> index ee6a8e6..8d8adac 100644
> --- a/src/cl_device_id.h
> +++ b/src/cl_device_id.h
> @@ -113,6 +113,10 @@ struct _cl_device_id {
> cl_device_affinity_domain affinity_domain;
> cl_device_partition_property partition_type[3];
> cl_uint device_reference_count;
> +#ifdef HAS_OCLImage2dBuffer
> + cl_uint image_pitch_alignment;
> + cl_uint image_base_address_alignment;
> +#endif
> };
>
> /* Get a device from the given platform */
> diff --git a/src/cl_extensions.c b/src/cl_extensions.c
> index d07a525..e31386f 100644
> --- a/src/cl_extensions.c
> +++ b/src/cl_extensions.c
> @@ -34,7 +34,7 @@ void check_opt1_extension(cl_extensions_t *extensions)
> {
> int id;
> for(id = OPT1_EXT_START_ID; id <= OPT1_EXT_END_ID; id++)
> - if (id == EXT_ID(khr_icd))
> + if (id == EXT_ID(khr_icd) || id == EXT_ID(khr_image2d_from_buffer))
> extensions->extensions[id].base.ext_enabled = 1;
> }
>
> diff --git a/src/cl_gt_device.h b/src/cl_gt_device.h
> index 37abfd2..f9c5ad4 100644
> --- a/src/cl_gt_device.h
> +++ b/src/cl_gt_device.h
> @@ -124,4 +124,8 @@ DECL_INFO_STRING(driver_version, LIBCL_DRIVER_VERSION_STRING)
> .affinity_domain = 0,
> .partition_type = {0},
> .device_reference_count = 1,
> +#ifdef HAS_OCLImage2dBuffer
> +.image_pitch_alignment = 4,
> +.image_base_address_alignment = 4,
> +#endif
>
> diff --git a/src/cl_mem.c b/src/cl_mem.c
> index 3323897..ab8efa8 100644
> --- a/src/cl_mem.c
> +++ b/src/cl_mem.c
> @@ -673,6 +673,7 @@ _cl_mem_new_image(cl_context ctx,
> size_t pitch,
> size_t slice_pitch,
> void *data,
> + int force_no_tiling,
> cl_int *errcode_ret)
> {
> cl_int err = CL_SUCCESS;
> @@ -736,7 +737,7 @@ _cl_mem_new_image(cl_context ctx,
> if (UNLIKELY(!data && pitch != 0)) DO_IMAGE_ERROR;
>
> /* Pick up tiling mode (we do only linear on SNB) */
> - if (cl_driver_get_ver(ctx->drv) != 6)
> + if (!force_no_tiling && cl_driver_get_ver(ctx->drv) != 6)
> tiling = cl_get_default_tiling(ctx->drv);
>
> depth = 1;
> @@ -935,7 +936,7 @@ _cl_mem_new_image_from_buffer(cl_context ctx,
> // We have to create a new image, and copy the buffer data to this new image.
> // And replace all the buffer object's reference to this image.
> image = _cl_mem_new_image(ctx, flags, image_format, image_desc->image_type,
> - mem_buffer->base.size / bpp, 0, 0, 0, 0, NULL, errcode_ret);
> + mem_buffer->base.size / bpp, 0, 0, 0, 0, NULL, 1, errcode_ret);
> if (image == NULL)
> return NULL;
> void *src = cl_mem_map(buffer, 0);
> @@ -953,7 +954,7 @@ _cl_mem_new_image_from_buffer(cl_context ctx,
>
> if (err != 0)
> goto error;
> -
> +
> // Now replace buffer's bo to this new bo, need to take care of sub buffer
> // case.
> cl_mem_replace_buffer(buffer, image->bo);
> @@ -974,6 +975,128 @@ error:
> return image;
> }
>
> +#ifdef HAS_OCLImage2dBuffer
> +static cl_mem
> +_cl_mem_new_image2d_from_buffer(cl_context ctx,
> + cl_mem_flags flags,
> + const cl_image_format* image_format,
> + const cl_image_desc *image_desc,
> + cl_int *errcode_ret)
> +{
> + cl_mem image = NULL;
> + cl_mem buffer = image_desc->buffer;
> + cl_int err = CL_SUCCESS;
> + *errcode_ret = err;
> + cl_mem_flags merged_flags;
> + uint32_t bpp;
> + uint32_t intel_fmt = INTEL_UNSUPPORTED_FORMAT;
> + size_t offset = 0;
> +
> + /* Get the size of each pixel */
> + if (UNLIKELY((err = cl_image_byte_per_pixel(image_format, &bpp)) != CL_SUCCESS))
> + goto error;
> +
> + /* Only a sub-set of the formats are supported */
> + intel_fmt = cl_image_get_intel_format(image_format);
> + if (UNLIKELY(intel_fmt == INTEL_UNSUPPORTED_FORMAT)) {
> + err = CL_INVALID_IMAGE_FORMAT_DESCRIPTOR;
> + goto error;
> + }
> +
> + if (!buffer) {
> + err = CL_INVALID_IMAGE_DESCRIPTOR;
> + goto error;
> + }
> +
> + if (flags & (CL_MEM_USE_HOST_PTR|CL_MEM_ALLOC_HOST_PTR|CL_MEM_COPY_HOST_PTR)) {
> + err = CL_INVALID_IMAGE_DESCRIPTOR;
> + goto error;
> + }
> +
> + /* access check. */
> + if ((buffer->flags & CL_MEM_WRITE_ONLY) &&
> + (flags & (CL_MEM_READ_WRITE|CL_MEM_READ_ONLY))) {
> + err = CL_INVALID_VALUE;
> + goto error;
> + }
> + if ((buffer->flags & CL_MEM_READ_ONLY) &&
> + (flags & (CL_MEM_READ_WRITE|CL_MEM_WRITE_ONLY))) {
> + err = CL_INVALID_VALUE;
> + goto error;
> + }
> + if ((buffer->flags & CL_MEM_HOST_WRITE_ONLY) &&
> + (flags & CL_MEM_HOST_READ_ONLY)) {
> + err = CL_INVALID_VALUE;
> + goto error;
> + }
> + if ((buffer->flags & CL_MEM_HOST_READ_ONLY) &&
> + (flags & CL_MEM_HOST_WRITE_ONLY)) {
> + err = CL_INVALID_VALUE;
> + goto error;
> + }
> + if ((buffer->flags & CL_MEM_HOST_NO_ACCESS) &&
> + (flags & (CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_WRITE_ONLY))) {
> + err = CL_INVALID_VALUE;
> + goto error;
> + }
> +
> + if (image_desc->image_width * image_desc->image_height * bpp > buffer->size) {
> + err = CL_INVALID_IMAGE_DESCRIPTOR;
> + goto error;
> + }
> +
> + merged_flags = buffer->flags;
> + if (flags & (CL_MEM_READ_WRITE|CL_MEM_READ_WRITE|CL_MEM_WRITE_ONLY)) {
> + merged_flags &= ~(CL_MEM_READ_WRITE|CL_MEM_READ_WRITE|CL_MEM_WRITE_ONLY);
> + merged_flags |= flags & (CL_MEM_READ_WRITE|CL_MEM_READ_WRITE|CL_MEM_WRITE_ONLY);
> + }
> + if (flags & (CL_MEM_HOST_WRITE_ONLY|CL_MEM_HOST_READ_ONLY|CL_MEM_HOST_NO_ACCESS)) {
> + merged_flags &= ~(CL_MEM_HOST_WRITE_ONLY|CL_MEM_HOST_READ_ONLY|CL_MEM_HOST_NO_ACCESS);
> + merged_flags |= flags & (CL_MEM_HOST_WRITE_ONLY|CL_MEM_HOST_READ_ONLY|CL_MEM_HOST_NO_ACCESS);
> + }
> + struct _cl_mem_buffer *mem_buffer = (struct _cl_mem_buffer*)buffer;
> + if (buffer->type == CL_MEM_SUBBUFFER_TYPE) {
> + offset = ((struct _cl_mem_buffer *)buffer)->sub_offset;
> + mem_buffer = mem_buffer->parent;
> + }
> + /* Get the size of each pixel */
> + if (UNLIKELY((err = cl_image_byte_per_pixel(image_format, &bpp)) != CL_SUCCESS))
> + goto error;
> +
> + // Per bspec, a image should has a at least 2 line vertical alignment,
> + // thus we can't simply attach a buffer to a 1d image surface which has the same size.
> + // We have to create a new image, and copy the buffer data to this new image.
> + // And replace all the buffer object's reference to this image.
> + image = _cl_mem_new_image(ctx, flags, image_format, image_desc->image_type,
> + image_desc->image_width, image_desc->image_height, 0, 0, 0, NULL, 1, errcode_ret);
> + if (image == NULL)
> + return NULL;
> +
> + if (err != 0)
> + goto error;
> +
> + // Now replace buffer's bo to this new bo, need to take care of sub buffer
> + // case.
> + cl_mem_replace_buffer(image, buffer->bo);
> + /* Now point to the right offset if buffer is a SUB_BUFFER. */
> + if (buffer->flags & CL_MEM_USE_HOST_PTR)
> + image->host_ptr = buffer->host_ptr + offset;
> + cl_mem_image(image)->offset = offset;
> + cl_mem_image(image)->w = image_desc->image_width;
> + cl_mem_image(image)->h = image_desc->image_height;
> + cl_mem_add_ref(buffer);
> + cl_mem_image(image)->buffer_1d = buffer;
> + return image;
> +
> +error:
> + if (image)
> + cl_mem_delete(image);
> + image = NULL;
> + *errcode_ret = err;
> + return image;
> +}
> +#endif
> +
> LOCAL cl_mem
> cl_mem_new_image(cl_context context,
> cl_mem_flags flags,
> @@ -983,19 +1106,37 @@ cl_mem_new_image(cl_context context,
> cl_int *errcode_ret)
> {
> switch (image_desc->image_type) {
> +#ifdef HAS_OCLImage2dBuffer
> + case CL_MEM_OBJECT_IMAGE1D:
> + case CL_MEM_OBJECT_IMAGE3D:
> + return _cl_mem_new_image(context, flags, image_format, image_desc->image_type,
> + image_desc->image_width, image_desc->image_height, image_desc->image_depth,
> + image_desc->image_row_pitch, image_desc->image_slice_pitch,
> + host_ptr, 0, errcode_ret);
> + case CL_MEM_OBJECT_IMAGE2D:
> + if (image_desc->buffer)
> + return _cl_mem_new_image2d_from_buffer(context, flags, image_format,
> + image_desc, errcode_ret);
> + else
> + return _cl_mem_new_image(context, flags, image_format, image_desc->image_type,
> + image_desc->image_width, image_desc->image_height, image_desc->image_depth,
> + image_desc->image_row_pitch, image_desc->image_slice_pitch,
> + host_ptr, 0, errcode_ret);
> +#else
> case CL_MEM_OBJECT_IMAGE1D:
> case CL_MEM_OBJECT_IMAGE2D:
> case CL_MEM_OBJECT_IMAGE3D:
> return _cl_mem_new_image(context, flags, image_format, image_desc->image_type,
> image_desc->image_width, image_desc->image_height, image_desc->image_depth,
> image_desc->image_row_pitch, image_desc->image_slice_pitch,
> - host_ptr, errcode_ret);
> + host_ptr, 0, errcode_ret);
> +#endif
> case CL_MEM_OBJECT_IMAGE1D_ARRAY:
> case CL_MEM_OBJECT_IMAGE2D_ARRAY:
> return _cl_mem_new_image(context, flags, image_format, image_desc->image_type,
> image_desc->image_width, image_desc->image_height, image_desc->image_array_size,
> image_desc->image_row_pitch, image_desc->image_slice_pitch,
> - host_ptr, errcode_ret);
> + host_ptr, 0, errcode_ret);
> case CL_MEM_OBJECT_IMAGE1D_BUFFER:
> return _cl_mem_new_image_from_buffer(context, flags, image_format,
> image_desc, errcode_ret);
> @@ -1024,7 +1165,12 @@ cl_mem_delete(cl_mem mem)
> /* iff we are a image, delete the 1d buffer if has. */
> if (IS_IMAGE(mem)) {
> if (cl_mem_image(mem)->buffer_1d) {
> +#ifdef HAS_OCLImage2dBuffer
> + assert(cl_mem_image(mem)->image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER ||
> + cl_mem_image(mem)->image_type == CL_MEM_OBJECT_IMAGE2D);
> +#else
> assert(cl_mem_image(mem)->image_type == CL_MEM_OBJECT_IMAGE1D_BUFFER);
> +#endif
> cl_mem_delete(cl_mem_image(mem)->buffer_1d);
> cl_mem_image(mem)->buffer_1d = NULL;
> }
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 8cc8b43..3736652 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -242,6 +242,12 @@ else()
> SET(UTESTS_REQUIRED_EGL_LIB "")
> endif()
>
> +if (OCL_IMAGE2D_BUFFER)
> +SET(utests_sources ${utests_sources} image_2D_buffer.cpp)
> +SET(CMAKE_CXX_FLAGS "-DHAS_OCLImage2dBuffer ${CMAKE_CXX_FLAGS} ${DEF_OCL_PCH_PCM_PATH}")
> +SET(CMAKE_C_FLAGS "-DHAS_OCLImage2dBuffer ${CMAKE_C_FLAGS} ${DEF_OCL_PCH_PCM_PATH}")
> +endif ()
> +
> if (COMPILER STREQUAL "CLANG")
> SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-tautological-compare")
> endif ()
> diff --git a/utests/image_2D_buffer.cpp b/utests/image_2D_buffer.cpp
> new file mode 100644
> index 0000000..e6e88d6
> --- /dev/null
> +++ b/utests/image_2D_buffer.cpp
> @@ -0,0 +1,89 @@
> +#include <string.h>
> +#include "utest_helper.hpp"
> +
> +#define TEST_SIZE 1024
> +
> +void image_2D_buffer(void)
> +{
> + size_t pix_w = TEST_SIZE;
> + size_t pix_h = TEST_SIZE;
> + size_t buffer_sz = pix_w * pix_h * sizeof(uint32_t);
> + 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_2D_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_IMAGE2D;
> + image_desc.image_row_pitch = pix_w * sizeof(uint32_t);
> + image_desc.image_width = pix_w;
> + image_desc.image_height = pix_h;
> + image_desc.image_array_size = 1;
> + 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_IMAGE2D;
> + image_desc.image_width = pix_w;
> + image_desc.image_height = pix_h;
> + 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, 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] = pix_w;
> + globals[1] = pix_h;
> + locals[0] = 16;
> + locals[1] = 16;
> +
> + OCL_NDRANGE(2);
> +
> + /* Now check the result. */
> + OCL_MAP_BUFFER(0);
> + for (uint32_t i = 0; i < buffer_sz/sizeof(uint32_t); i++)
> + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 1);
> + OCL_UNMAP_BUFFER(0);
> +
> + clReleaseSampler(sampler);
> + clReleaseMemObject(image1);
> + clReleaseMemObject(image2);
> + clReleaseMemObject(buff);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(image_2D_buffer);
> --
> 1.9.3
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list