On Tue, 2014-09-09 at 17:11 -0500, Aaron Watry wrote: > v2: Properly test signed/unsigned values > > Signed-off-by: Aaron Watry <[email protected]> > --- > .../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 <[email protected]>
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/piglit
