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.g...@linux.intel.com] Sent: Tuesday, June 24, 2014 3:51 PM To: Luo, Xionghu Cc: beignet@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....@intel.com wrote: > From: Luo <xionghu....@intel.com> > > detail cases: 1Darray, 2Darray, 2Darrayto2D, 2Darrayto3D, 2Dto2Darray, > 3Dto2Darray. > > 1d used gpu copy. > > Signed-off-by: Luo <xionghu....@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@lists.freedesktop.org > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet
