[Beignet] [PATCH] add cpu copy for 1Darray and 2darray related copy APIs.
xionghu.luo at intel.com
xionghu.luo at intel.com
Mon Jun 23 19:09:12 PDT 2014
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]};
+ 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
More information about the Beignet
mailing list