[Piglit] [PATCH 1/2] cl: Add atomic_max tests

[Jan Vesely]

On Tue, 2014-09-09 at 17:11 -0500, Aaron Watry wrote:
> v2: Properly test signed/unsigned values
> 
> Signed-off-by: Aaron Watry <awatry at gmail.com>
> ---
>  .../execute/builtin/atomic/atomic_max-local.cl     | 81 ++++++++++++++++++++++
>  1 file changed, 81 insertions(+)
>  create mode 100644 tests/cl/program/execute/builtin/atomic/atomic_max-local.cl
> 
> diff --git a/tests/cl/program/execute/builtin/atomic/atomic_max-local.cl b/tests/cl/program/execute/builtin/atomic/atomic_max-local.cl
> new file mode 100644
> index 0000000..68b513f
> --- /dev/null
> +++ b/tests/cl/program/execute/builtin/atomic/atomic_max-local.cl
> @@ -0,0 +1,81 @@
> +/*!
> +[config]
> +name: atomic_max local
> +clc_version_min: 11
> +
> +[test]
> +name: simple int
> +kernel_name: simple_int
> +dimensions: 1
> +global_size: 1 0 0
> +local_size: 1 0 0
> +arg_out: 0 buffer int[2] -1 2
> +arg_in:  1 buffer int[1] NULL
> +arg_in:  2 int           -1
> +arg_in:  3 int            2
> +
> +[test]
> +name: simple uint
> +kernel_name: simple_uint
> +dimensions: 1
> +global_size: 1 0 0
> +local_size:  1 0 0
> +arg_out: 0 buffer uint[2] 2 3
> +arg_in:  1 buffer uint[1] NULL
> +arg_in:  2 uint           2
> +arg_in:  3 uint           3
> +
> +[test]
> +name: simple uint 2
> +kernel_name: simple_uint
> +dimensions: 1
> +global_size: 1 0 0
> +local_size:  1 0 0
> +arg_out: 0 buffer uint[2] 3 4294967295
> +arg_in:  1 buffer uint[1] NULL
> +arg_in:  2 uint           3
> +arg_in:  3 uint           0xffffffff
> +
> +
> +[test]
> +name: threads

this name conflict with the name of the subtest below, same for the
second patch.

sorry for not noticing this earlier, only came to my attention after
piglit started complaining.

jan

> +kernel_name: threads_int
> +dimensions: 1
> +global_size: 8 0 0
> +local_size:  8 0 0
> +arg_out: 0 buffer int[1] 7
> +arg_in:  1 buffer int[1] NULL
> +
> +[test]
> +name: threads
> +kernel_name: threads_uint
> +dimensions: 1
> +global_size: 8 0 0
> +local_size:  8 0 0
> +arg_out: 0 buffer uint[1] 7
> +arg_in:  1 buffer uint[1] NULL
> +
> +!*/
> +
> +#define SIMPLE_TEST(TYPE) \
> +kernel void simple_##TYPE(global TYPE *out, local TYPE *mem, TYPE initial, TYPE other) { \
> +	*mem = initial; \
> +	TYPE a = atomic_max(mem, other); \
> +	out[0] = a; \
> +	out[1] = *mem; \
> +}
> +
> +#define THREADS_TEST(TYPE) \
> +kernel void threads_##TYPE(global TYPE *out, local TYPE *mem) { \
> +	*mem = 0; \
> +	barrier(CLK_LOCAL_MEM_FENCE); \
> +	atomic_max(mem, get_global_id(0)); \
> +	barrier(CLK_LOCAL_MEM_FENCE); \
> +	*out = *mem; \
> +}
> +
> +SIMPLE_TEST(int)
> +SIMPLE_TEST(uint)
> +
> +THREADS_TEST(int)
> +THREADS_TEST(uint)


-- 
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <http://lists.freedesktop.org/archives/piglit/attachments/20140925/2cea2ea4/attachment.sig>