From bfa4dba15485b4192f8bbe13345e9658c97ecf76 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Picca=20Fr=C3=A9d=C3=A9ric-Emmanuel?= Date: Sat, 7 Oct 2017 07:59:01 +0200 Subject: New upstream version 0.6.0+dfsg --- silx/resources/opencl/array_utils.cl | 33 +++++++++++++++++++++++++++++++++ 1 file changed, 33 insertions(+) create mode 100644 silx/resources/opencl/array_utils.cl (limited to 'silx/resources/opencl/array_utils.cl') diff --git a/silx/resources/opencl/array_utils.cl b/silx/resources/opencl/array_utils.cl new file mode 100644 index 0000000..60677dc --- /dev/null +++ b/silx/resources/opencl/array_utils.cl @@ -0,0 +1,33 @@ +/** + * 2D Memcpy for float* arrays, + * replacing pyopencl "enqueue_copy" which does not return the expected result + * when dealing with rectangular buffers. + * ALL THE SIZES/OFFSETS ARE SPECIFIED IN PIXELS, NOT IN BYTES. + * In the (x, y) convention, x is the fast index (as in CUDA). + * + * :param dst: destination array + * :param src: source array + * :param dst_width: width of the dst array + * :param src_width: width of the src array + * :param dst_offset: tuple with the offset (x, y) in the dst array + * :param src_offset: tuple with the offset (x, y) in the src array + * :param transfer_shape: shape of the transfer array in the form (x, y) + * + */ +kernel void cpy2d( + global float* dst, + global float* src, + int dst_width, + int src_width, + int2 dst_offset, + int2 src_offset, + int2 transfer_shape) +{ + int gidx = get_global_id(0), + gidy = get_global_id(1); + if (gidx < transfer_shape.x && gidy < transfer_shape.y) + { + dst[(dst_offset.y + gidy)*dst_width + (dst_offset.x + gidx)] = src[(src_offset.y + gidy)*src_width + (src_offset.x + gidx)]; + } +} + -- cgit v1.2.3