diff options
author | Picca Frédéric-Emmanuel <picca@debian.org> | 2017-10-07 07:59:01 +0200 |
---|---|---|
committer | Picca Frédéric-Emmanuel <picca@debian.org> | 2017-10-07 07:59:01 +0200 |
commit | bfa4dba15485b4192f8bbe13345e9658c97ecf76 (patch) | |
tree | fb9c6e5860881fbde902f7cbdbd41dc4a3a9fb5d /silx/resources/opencl/array_utils.cl | |
parent | f7bdc2acff3c13a6d632c28c4569690ab106eed7 (diff) |
New upstream version 0.6.0+dfsg
Diffstat (limited to 'silx/resources/opencl/array_utils.cl')
-rw-r--r-- | silx/resources/opencl/array_utils.cl | 33 |
1 files changed, 33 insertions, 0 deletions
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)]; + } +} + |