[Piglit] [PATCH] CL: Add global memory read/write and barrier tests

[Tom Stellard]

On Mon, Oct 14, 2013 at 08:40:47PM -0500, Aaron Watry wrote:
> Based on local-memory tests.

Reviewed-by: Tom Stellard <thomas.stellard at amd.com>

> ---
>  tests/cl/program/execute/global-memory.cl | 83 +++++++++++++++++++++++++++++++
>  1 file changed, 83 insertions(+)
>  create mode 100644 tests/cl/program/execute/global-memory.cl
> 
> diff --git a/tests/cl/program/execute/global-memory.cl b/tests/cl/program/execute/global-memory.cl
> new file mode 100644
> index 0000000..7ad9e72
> --- /dev/null
> +++ b/tests/cl/program/execute/global-memory.cl
> @@ -0,0 +1,83 @@
> +/*!
> +[config]
> +name: global_memory
> +
> +[test]
> +name: Simple
> +kernel_name: simple
> +dimensions: 1
> +global_size: 1 0 0
> +local_size:  1 0 0
> +arg_out: 0 buffer int[2] -1 -1
> +arg_in: 1 buffer int[2] 0 0
> +
> +[test]
> +name: (16 x 1 x 1) (16 x 1 x 1)
> +kernel_name: global_memory_one_work_group
> +dimensions: 1
> +global_size: 16 0 0
> +local_size:  16 0 0
> +arg_out: 0 buffer int[16] 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0
> +arg_in: 1 buffer int[16] 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
> +
> +[test]
> +name: (64 x 1 x 1) (4 x 1 x 1)
> +kernel_name: global_memory_many_work_groups
> +dimensions: 1
> +global_size: 64 0 0
> +local_size:   4 0 0
> +arg_out: 0 buffer int[64] repeat 1 2 3 0
> +arg_in: 1 buffer int[64] 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 
> +
> +[test]
> +name: 2 global memory objects
> +kernel_name: global_memory_two_objects
> +dimensions: 1
> +global_size: 4 0 0
> +local_size:  4 0 0
> +arg_out: 0 buffer int[8] 3 2 1 0 6 4 2 0
> +arg_in: 1 buffer int[4] repeat 0
> +arg_in: 2 buffer int[4] repeat 0
> +
> +!*/
> +
> +kernel void simple(global int *out, global int *tmp_mem) {
> +	tmp_mem[0] = 0xffffffff;
> +	tmp_mem[1] = 0xffffffff;
> +	out[0] = tmp_mem[0];
> +	out[1] = tmp_mem[1];
> +}
> +
> +kernel void global_memory_one_work_group(global int *out, global int *tmp_mem) {
> +	int group_offset = get_group_id(0) * get_local_size(0);
> +	int index = get_local_id(0);
> +	int index2;
> +	tmp_mem[group_offset + index] = index;
> +	index2 = index + 1;
> +	index2 = (index2 % 16);
> +	barrier(CLK_GLOBAL_MEM_FENCE);
> +	out[index] = tmp_mem[group_offset + index2];
> +}
> +
> +kernel void global_memory_many_work_groups(global int *out, global int *tmp_mem) {
> +	int group_offset = get_group_id(0) * get_local_size(0);
> +	int index = get_local_id(0);
> +	int global_id = get_global_id(0);
> +	int index2;
> +	tmp_mem[group_offset + index] = index;
> +	index2 = index + 1;
> +	index2 = (index2 % 4);
> +	barrier(CLK_GLOBAL_MEM_FENCE);
> +	out[global_id] = tmp_mem[group_offset + index2];
> +}
> +
> +kernel void global_memory_two_objects(global int *out, global int *tmp_mem0, global int *tmp_mem1) {
> +	int group_offset = get_group_id(0) * get_local_size(0);
> +	int index = get_local_id(0);
> +	int index2 = 3 - index;
> +	tmp_mem0[group_offset + index] = index;
> +	tmp_mem1[group_offset + index] = index * 2;
> +	barrier(CLK_GLOBAL_MEM_FENCE);
> +	out[group_offset + index] = tmp_mem0[group_offset + index2];
> +	out[group_offset + index + 4] = tmp_mem1[group_offset + index2];
> +}
> -- 
> 1.8.4
> 
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/piglit