[Beignet] [PATCH] add cpu copy for 1Darray and 2darray related copy APIs.
Zhigang Gong
zhigang.gong at linux.intel.com
Tue Jun 24 06:22:34 PDT 2014
I double checked the spec and found the spec doesn't define the error code
when try to copy a non-1D image to a 1D image. The CL_IMAGE_FORMAT_NOT_SUPPORTED
is not for this purpose. IMO, the most closing error code should be
CL_IMAGE_FORMAT_MISMATCH.
On Tue, Jun 24, 2014 at 09:40:15AM +0000, Luo, Xionghu wrote:
> Image1d could only be copied to image1d right now, no need for other option, shall I return CL_IMAGE_FORMAT_NOT_SUPPORTED for all else branches?
>
> Luo Xionghu
> Best Regards
>
>
> -----Original Message-----
> From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com]
> Sent: Tuesday, June 24, 2014 3:51 PM
> To: Luo, Xionghu
> Cc: beignet at lists.freedesktop.org
> Subject: Re: [Beignet] [PATCH] add cpu copy for 1Darray and 2darray related copy APIs.
>
> On Tue, Jun 24, 2014 at 10:09:12AM +0800, xionghu.luo at intel.com wrote:
> > From: Luo <xionghu.luo at intel.com>
> >
> > detail cases: 1Darray, 2Darray, 2Darrayto2D, 2Darrayto3D, 2Dto2Darray, 3Dto2Darray.
> >
> > 1d used gpu copy.
> >
> > Signed-off-by: Luo <xionghu.luo at intel.com>
> > ---
> > src/CMakeLists.txt | 4 +-
> > src/cl_context.h | 1 +
> > src/cl_mem.c | 73 +++++++++++++++++++++++++-
> > src/cl_mem.h | 4 ++
> > src/kernels/cl_internal_copy_image_1d_to_1d.cl | 19 +++++++
> > 5 files changed, 97 insertions(+), 4 deletions(-) create mode 100644
> > src/kernels/cl_internal_copy_image_1d_to_1d.cl
> >
> > diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index
> > 8651af6..82b6df0 100644
> > --- a/src/CMakeLists.txt
> > +++ b/src/CMakeLists.txt
> > @@ -41,8 +41,8 @@ set (KERNEL_STR_FILES) set (KERNEL_NAMES
> > cl_internal_copy_buf_align4
> > cl_internal_copy_buf_align16 cl_internal_copy_buf_unalign_same_offset
> > cl_internal_copy_buf_unalign_dst_offset
> > cl_internal_copy_buf_unalign_src_offset
> > -cl_internal_copy_buf_rect cl_internal_copy_image_2d_to_2d
> > cl_internal_copy_image_3d_to_2d -cl_internal_copy_image_2d_to_3d
> > cl_internal_copy_image_3d_to_3d
> > +cl_internal_copy_buf_rect cl_internal_copy_image_1d_to_1d
> > +cl_internal_copy_image_2d_to_2d cl_internal_copy_image_3d_to_2d
> > +cl_internal_copy_image_2d_to_3d cl_internal_copy_image_3d_to_3d
> > cl_internal_copy_image_2d_to_buffer
> > cl_internal_copy_image_3d_to_buffer
> > cl_internal_copy_buffer_to_image_2d
> > cl_internal_copy_buffer_to_image_3d
> > cl_internal_fill_buf_align8 cl_internal_fill_buf_align4 diff --git
> > a/src/cl_context.h b/src/cl_context.h index cba0a0a..74e31c7 100644
> > --- a/src/cl_context.h
> > +++ b/src/cl_context.h
> > @@ -47,6 +47,7 @@ enum _cl_internal_ker_type {
> > CL_ENQUEUE_COPY_BUFFER_UNALIGN_DST_OFFSET,
> > CL_ENQUEUE_COPY_BUFFER_UNALIGN_SRC_OFFSET,
> > CL_ENQUEUE_COPY_BUFFER_RECT,
> > + CL_ENQUEUE_COPY_IMAGE_1D_TO_1D, //copy image 1d to image 1d
> > CL_ENQUEUE_COPY_IMAGE_2D_TO_2D, //copy image 2d to image 2d
> > CL_ENQUEUE_COPY_IMAGE_3D_TO_2D, //copy image 3d to image 2d
> > CL_ENQUEUE_COPY_IMAGE_2D_TO_3D, //copy image 2d to image 3d
> > diff --git a/src/cl_mem.c b/src/cl_mem.c index e0c4ec9..8bb7215 100644
> > --- a/src/cl_mem.c
> > +++ b/src/cl_mem.c
> > @@ -542,6 +542,38 @@ cl_mem_copy_image_region(const size_t *origin, const size_t *region,
> > }
> > }
> >
> > +void
> > +cl_mem_copy_image_to_image(const size_t *dst_origin,const size_t *src_origin, const size_t *region,
> > + const struct _cl_mem_image *dst_image,
> > +const struct _cl_mem_image *src_image) {
> > + //printf("origin:%u,%u,%u to %u,%u,%u\n",
> > +src_origin[0],src_origin[1], src_origin[2],
> > +dst_origin[0],dst_origin[1], dst_origin[2]);
> > + //printf("region:%u,%u,%u \n", region[0],region[1], region[2]);
> > + //printf("pitch:%u,%u to %u,%u\n", src_image->row_pitch,
> > +src_image->slice_pitch,dst_image->row_pitch, dst_image->slice_pitch);
> > +
> > + char* dst= cl_mem_map_auto((cl_mem)dst_image);
> > + char* src= cl_mem_map_auto((cl_mem)src_image);
> > + size_t dst_offset = dst_image->bpp * dst_origin[0] +
> > + dst_image->row_pitch * dst_origin[1] + dst_image->slice_pitch *
> > + dst_origin[2]; size_t src_offset = src_image->bpp * src_origin[0] +
> > + src_image->row_pitch * src_origin[1] + src_image->slice_pitch *
> > + src_origin[2]; dst= (char*)dst+ dst_offset; src= (char*)src+
> > + src_offset; cl_uint y, z; for (z = 0; z < region[2]; z++) {
> > + const char* src_ptr = src;
> > + char* dst_ptr = dst;
> > + for (y = 0; y < region[1]; y++) {
> > + memcpy(dst_ptr, src_ptr, src_image->bpp*region[0]);
> > + src_ptr += src_image->row_pitch;
> > + dst_ptr += dst_image->row_pitch;
> > + }
> > + src = (char*)src + src_image->slice_pitch;
> > + dst = (char*)dst + dst_image->slice_pitch; }
> > +
> > + cl_mem_unmap_auto((cl_mem)src_image);
> > + cl_mem_unmap_auto((cl_mem)dst_image);
> > +
> > +}
> > +
> > static void
> > cl_mem_copy_image(struct _cl_mem_image *image,
> > size_t row_pitch,
> > @@ -1377,7 +1409,16 @@ cl_mem_kernel_copy_image(cl_command_queue queue, struct _cl_mem_image* src_image
> > assert(src_image->base.ctx == dst_image->base.ctx);
> >
> > /* setup the kernel and run. */
> > - if(src_image->image_type == CL_MEM_OBJECT_IMAGE2D) {
> > + if(src_image->image_type == CL_MEM_OBJECT_IMAGE1D) {
> > + if(dst_image->image_type == CL_MEM_OBJECT_IMAGE1D) {
> > + extern char cl_internal_copy_image_1d_to_1d_str[];
> > + extern int cl_internal_copy_image_1d_to_1d_str_size;
> > +
> > + ker = cl_context_get_static_kernel_form_bin(queue->ctx, CL_ENQUEUE_COPY_IMAGE_1D_TO_1D,
> > + cl_internal_copy_image_1d_to_1d_str, (size_t)cl_internal_copy_image_1d_to_1d_str_size, NULL);
> > + }
> Did you forget the else branch here? What if the src image is Image 1D but the dst image is not?
> > +
> > + }else if(src_image->image_type == CL_MEM_OBJECT_IMAGE2D) {
> > if(dst_image->image_type == CL_MEM_OBJECT_IMAGE2D) {
> > extern char cl_internal_copy_image_2d_to_2d_str[];
> > extern int cl_internal_copy_image_2d_to_2d_str_size;
> > @@ -1390,8 +1431,33 @@ cl_mem_kernel_copy_image(cl_command_queue
> > queue, struct _cl_mem_image* src_image
> >
> > ker = cl_context_get_static_kernel_form_bin(queue->ctx, CL_ENQUEUE_COPY_IMAGE_2D_TO_3D,
> > cl_internal_copy_image_2d_to_3d_str,
> > (size_t)cl_internal_copy_image_2d_to_3d_str_size, NULL);
> > + }else if(dst_image->image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY) {
> > +
> > + cl_mem_copy_image_to_image(dst_origin, src_origin, region, dst_image, src_image);
> > + return CL_SUCCESS;
> > + }
> > + }else if(src_image->image_type == CL_MEM_OBJECT_IMAGE1D_ARRAY) {
> > + if(dst_image->image_type == CL_MEM_OBJECT_IMAGE1D_ARRAY) {
> > +
> > + const size_t src_origin_cpu[]={src_origin[0], src_origin[2], src_origin[1]};
> > + const size_t dst_origin_cpu[]={dst_origin[0], dst_origin[2], dst_origin[1]};
> > + const size_t region_cpu[]={region[0], region[2], region[1]};
> > + cl_mem_copy_image_to_image(dst_origin_cpu, src_origin_cpu, region_cpu, dst_image, src_image);
> > + return CL_SUCCESS;
> > + }
> > + }else if(src_image->image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY) {
> > + if(dst_image->image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY) {
> > +
> > + cl_mem_copy_image_to_image(dst_origin, src_origin, region, dst_image, src_image);
> > + return CL_SUCCESS;
> > + }else if(dst_image->image_type == CL_MEM_OBJECT_IMAGE2D) {
> > + cl_mem_copy_image_to_image(dst_origin, src_origin, region, dst_image, src_image);
> > + return CL_SUCCESS;
> > + }else if(dst_image->image_type == CL_MEM_OBJECT_IMAGE3D) {
> > + cl_mem_copy_image_to_image(dst_origin, src_origin, region, dst_image, src_image);
> > + return CL_SUCCESS;
> > }
> > - }else if(src_image->image_type == CL_MEM_OBJECT_IMAGE3D) {
> > + }else if(src_image->image_type == CL_MEM_OBJECT_IMAGE3D) {
> > if(dst_image->image_type == CL_MEM_OBJECT_IMAGE2D) {
> > extern char cl_internal_copy_image_3d_to_2d_str[];
> > extern int cl_internal_copy_image_3d_to_2d_str_size;
> > @@ -1404,6 +1470,9 @@ cl_mem_kernel_copy_image(cl_command_queue queue,
> > struct _cl_mem_image* src_image
> >
> > ker = cl_context_get_static_kernel_form_bin(queue->ctx, CL_ENQUEUE_COPY_IMAGE_3D_TO_3D,
> > cl_internal_copy_image_3d_to_3d_str,
> > (size_t)cl_internal_copy_image_3d_to_3d_str_size, NULL);
> > + }else if(dst_image->image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY) {
> > + cl_mem_copy_image_to_image(dst_origin, src_origin, region, dst_image, src_image);
> > + return CL_SUCCESS;
> > }
> > }
> >
> > diff --git a/src/cl_mem.h b/src/cl_mem.h index d589093..b8012a0 100644
> > --- a/src/cl_mem.h
> > +++ b/src/cl_mem.h
> > @@ -261,6 +261,10 @@ cl_mem_copy_image_region(const size_t *origin, const size_t *region,
> > const void *src, size_t src_row_pitch, size_t src_slice_pitch,
> > const struct _cl_mem_image *image);
> >
> > +void
> > +cl_mem_copy_image_to_image(const size_t *dst_origin,const size_t *src_origin, const size_t *region,
> > + const struct _cl_mem_image *dst_image,
> > +const struct _cl_mem_image *src_image);
> > +
> > extern cl_mem cl_mem_new_libva_buffer(cl_context ctx,
> > unsigned int bo_name,
> > cl_int *errcode); diff --git
> > a/src/kernels/cl_internal_copy_image_1d_to_1d.cl
> > b/src/kernels/cl_internal_copy_image_1d_to_1d.cl
> > new file mode 100644
> > index 0000000..dca82b2
> > --- /dev/null
> > +++ b/src/kernels/cl_internal_copy_image_1d_to_1d.cl
> > @@ -0,0 +1,19 @@
> > +kernel void __cl_copy_image_1d_to_1d(__read_only image1d_t src_image, __write_only image1d_t dst_image,
> > + unsigned int region0, unsigned int region1, unsigned int region2,
> > + unsigned int src_origin0, unsigned int src_origin1, unsigned int src_origin2,
> > + unsigned int dst_origin0, unsigned int
> > +dst_origin1, unsigned int dst_origin2) {
> > + int i = get_global_id(0);
> > + int j = get_global_id(1);
> > + int k = get_global_id(2);
> > + int4 color;
> > + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
> > +CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
> > + int src_coord;
> > + int dst_coord;
> > + if((i >= region0) || (j>= region1) || (k>=region2))
> > + return;
> > + src_coord = src_origin0 + i;
> > + dst_coord = dst_origin0 + i;
> > + color = read_imagei(src_image, sampler, src_coord);
> > + write_imagei(dst_image, dst_coord, color); }
> > --
> > 1.8.1.2
> >
> > _______________________________________________
> > Beignet mailing list
> > Beignet at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list