On Monday 13 June 2016 11:23:49 Jan Vesely wrote: > Passes on CUDA, beignet, clover on kaveri, and intel CPU. > > Signed-off-by: Jan Vesely <[email protected]>
As spotted by git, there is an extra line added to the file atomic_cmpxchg- global-return.cl. The other test don't have empty line at the end. Otherwise the serie is Reviewed-by : Serge Martin <[email protected]> > --- > .../builtin/atomic/atomic_cmpxchg-global-return.cl | 74 > ++++++++++++++++++++++ .../builtin/atomic/atomic_cmpxchg-global.cl | > 69 ++++++++++++++++++++ 2 files changed, 143 insertions(+) > create mode 100644 > tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global-return.cl > create mode 100644 > tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global.cl > > diff --git > a/tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global-return.cl > b/tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global-return.cl > new file mode 100644 > index 0000000..2cc1031 > --- /dev/null > +++ > b/tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global-return.cl > @@ -0,0 +1,74 @@ > +/*! > +[config] > +name: atomic_cmpxchg global return > +clc_version_min: 11 > + > +[test] > +name: simple int > +kernel_name: simple_int > +dimensions: 1 > +global_size: 1 0 0 > +arg_out: 0 buffer int[2] 5 -4 > +arg_in: 0 buffer int[2] -4 -4 > +arg_in: 1 buffer int[2] -4 3 > +arg_in: 2 buffer int[2] 5 5 > +arg_out: 3 buffer int[2] -4 -4 > + > +[test] > +name: simple uint > +kernel_name: simple_uint > +dimensions: 1 > +global_size: 1 0 0 > +arg_out: 0 buffer uint[2] 5 4 > +arg_in: 0 buffer uint[2] 4 4 > +arg_in: 1 buffer uint[2] 4 3 > +arg_in: 2 buffer uint[2] 5 5 > +arg_out: 3 buffer uint[2] 4 4 > + > +[test] > +name: threads int > +kernel_name: threads_int > +dimensions: 1 > +global_size: 8 0 0 > +local_size: 8 0 0 > +arg_out: 0 buffer int[1] 8 > +arg_in: 0 buffer int[1] 0 > +arg_out: 1 buffer int[8] 0 1 2 3 4 5 6 7 > + > +[test] > +name: threads uint > +kernel_name: threads_uint > +dimensions: 1 > +global_size: 8 0 0 > +local_size: 8 0 0 > +arg_out: 0 buffer uint[1] 8 > +arg_in: 0 buffer uint[1] 0 > +arg_out: 1 buffer uint[8] 0 1 2 3 4 5 6 7 > + > +!*/ > + > +#define SIMPLE_TEST(TYPE) \ > +kernel void simple_##TYPE(global TYPE *initial, global TYPE *compare, > global TYPE *value, global TYPE *old) { \ + old[0] = > atomic_cmpxchg(initial, compare[0], value[0]); \ > + old[1] = atomic_cmpxchg(initial+1, compare[1], value[1]); \ > +} > + > +#define THREADS_TEST(TYPE) \ > +kernel void threads_##TYPE(global TYPE *out, global TYPE *old) { \ > + int i; \ > + barrier(CLK_GLOBAL_MEM_FENCE); \ > + TYPE id = get_global_id(0); \ > + for(i = 0; i < get_global_size(0); i++){ \ > + TYPE old_val = atomic_cmpxchg(out, id, id+1); \ > + if (old_val == id) /* success */ \ > + old[id] = old_val; \ > + barrier(CLK_GLOBAL_MEM_FENCE); \ > + } \ > +} > + > +SIMPLE_TEST(int) > +SIMPLE_TEST(uint) > + > +THREADS_TEST(int) > +THREADS_TEST(uint) > + > diff --git > a/tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global.cl > b/tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global.cl new file > mode 100644 > index 0000000..05cf401 > --- /dev/null > +++ b/tests/cl/program/execute/builtin/atomic/atomic_cmpxchg-global.cl > @@ -0,0 +1,69 @@ > +/*! > +[config] > +name: atomic_cmpxchg global > +clc_version_min: 11 > + > +[test] > +name: simple int > +kernel_name: simple_int > +dimensions: 1 > +global_size: 1 0 0 > +arg_out: 0 buffer int[2] 5 -4 > +arg_in: 0 buffer int[2] -4 -4 > +arg_in: 1 buffer int[2] -4 3 > +arg_in: 2 buffer int[2] 5 5 > + > +[test] > +name: simple uint > +kernel_name: simple_uint > +dimensions: 1 > +global_size: 1 0 0 > +arg_out: 0 buffer uint[2] 5 4 > +arg_in: 0 buffer uint[2] 4 4 > +arg_in: 1 buffer uint[2] 4 3 > +arg_in: 2 buffer uint[2] 5 5 > + > +[test] > +name: threads int > +kernel_name: threads_int > +dimensions: 1 > +global_size: 8 0 0 > +local_size: 8 0 0 > +arg_out: 0 buffer int[1] 8 > +arg_in: 0 buffer int[1] 0 > + > +[test] > +name: threads uint > +kernel_name: threads_uint > +dimensions: 1 > +global_size: 8 0 0 > +local_size: 8 0 0 > +arg_out: 0 buffer uint[1] 8 > +arg_in: 0 buffer uint[1] 0 > + > +!*/ > + > +#define SIMPLE_TEST(TYPE) \ > +kernel void simple_##TYPE(global TYPE *initial, global TYPE *compare, > global TYPE *value) { \ + atomic_cmpxchg(initial, compare[0], value[0]); \ > + atomic_cmpxchg(initial+1, compare[1], value[1]); \ > +} > + > +#define THREADS_TEST(TYPE) \ > +kernel void threads_##TYPE(global TYPE *out) { \ > + int i; \ > + barrier(CLK_GLOBAL_MEM_FENCE); \ > + TYPE id = get_global_id(0); \ > + for(i = 0; i < get_global_size(0); i++){ \ > + if (i == id){ \ > + atomic_cmpxchg(out, id, id+1); \ > + } \ > + barrier(CLK_GLOBAL_MEM_FENCE); \ > + } \ > +} > + > +SIMPLE_TEST(int) > +SIMPLE_TEST(uint) > + > +THREADS_TEST(int) > +THREADS_TEST(uint) _______________________________________________ Piglit mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/piglit
