[Beignet] [PATCH] add cpu copy for 1Darray and 2darray related copy APIs.
Zhigang Gong
zhigang.gong at linux.intel.com
Tue Jun 24 08:37:07 PDT 2014
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);
> + }
> +
> + }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]};
We also use array index as depth on cpu side. As it need to access each array by using
the slice pitch rather than the row pitch. I remove the above three lines don't use any _cpu
variables at below line.
> + 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
More information about the Beignet
mailing list