1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
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)];
}
}
|
