/*!
[config]
name: Copying struct of three integers.
clc_version_min: 10

[test]
kernel_name: struct_field_copy
arg_in:  0 buffer uint[5] 42 2 3 4 5
arg_in:  1 buffer uint[1] 0
arg_out: 1 buffer uint[1] 1

[test]
kernel_name: struct_assign_copy
arg_in:  0 buffer uint[5] 42 2 3 4 5
arg_in:  1 buffer uint[1] 0
arg_out: 1 buffer uint[1] 1

[test]
kernel_name: struct_pass_copy
arg_in:  0 buffer uint[5] 42 2 3 4 5
arg_in:  1 buffer uint[1] 0
arg_out: 1 buffer uint[1] 1

[test]
kernel_name: struct_local_copy
arg_in:  0 buffer uint[5] 42 2 3 4 5
arg_in:  1 buffer uint[1] 0
arg_out: 1 buffer uint[1] 42
!*/

typedef struct {
	int f0, f1, f2;
} Block;

void process(__private const Block block, __private int *data) {
	if (block.f0 == 42) {
		*data = 1;
	}
}

#define CPY(src, dst) { \
	(dst).f0 = (src).f0; \
	(dst).f1 = (src).f1; \
	(dst).f2 = (src).f2; \
}

__kernel
void struct_field_copy(__global const Block *ip, __global int *out) {
	int tmp;
	Block block;

	CPY(ip[0], block);

	tmp = out[0];
	process(block, &tmp);
	*out = tmp;
}

__kernel
void struct_assign_copy(__global const Block *ip, __global int *out) {
	int tmp;
	Block block;

	block = ip[0];

	tmp = out[0];
	process(block, &tmp);
	*out = tmp;
}

__kernel
void struct_pass_copy(__global const Block *ip, __global int *out) {
	int tmp;

	tmp = out[0];
	process(ip[0], &tmp);
	*out = tmp;
}

__kernel
void struct_local_copy(__global const Block *ip, __global int *out) {
	Block block;

	block = ip[0];

	*out = block.f0;
}