[Piglit] [PATCH] cl: Add tests for local memory atomic_inc and atomic_add builtins

[Aaron Watry]

Reviewed-by: Aaron Watry <awatry at gmail.com>

On Tue, Aug 20, 2013 at 6:19 PM, Tom Stellard <tom at stellard.net> wrote:
> From: Tom Stellard <thomas.stellard at amd.com>
>
> ---
>  tests/all_cl.tests                                 |  1 +
>  .../execute/builtin/atomic/atomic_add-local.cl     | 61 ++++++++++++++++++++++
>  .../execute/builtin/atomic/atomic_inc-local.cl     | 58 ++++++++++++++++++++
>  3 files changed, 120 insertions(+)
>  create mode 100644 tests/cl/program/execute/builtin/atomic/atomic_add-local.cl
>  create mode 100644 tests/cl/program/execute/builtin/atomic/atomic_inc-local.cl
>
> diff --git a/tests/all_cl.tests b/tests/all_cl.tests
> index efb04e4..d0c279d 100644
> --- a/tests/all_cl.tests
> +++ b/tests/all_cl.tests
> @@ -105,6 +105,7 @@ program["Execute"] = program_execute
>  add_program_test_dir(program_build, 'tests/cl/program/build')
>  add_program_test_dir(program_build_fail, 'tests/cl/program/build/fail')
>  add_program_test_dir(program_execute, 'tests/cl/program/execute')
> +add_program_test_dir(program_execute, 'tests/cl/program/execute/builtin/atomic')
>
>  #Run generated built-in tests
>  program_execute_builtin = Group()
> diff --git a/tests/cl/program/execute/builtin/atomic/atomic_add-local.cl b/tests/cl/program/execute/builtin/atomic/atomic_add-local.cl
> new file mode 100644
> index 0000000..3a30005
> --- /dev/null
> +++ b/tests/cl/program/execute/builtin/atomic/atomic_add-local.cl
> @@ -0,0 +1,61 @@
> +/*!
> +[config]
> +name: atomic_add 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] -4 1
> +arg_in:  1 buffer int[1] NULL
> +arg_in:  2 int           -4
> +arg_in:  3 int           5
> +
> +[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] 4 9
> +arg_in:  1 buffer uint[1] NULL
> +arg_in:  2 uint           4
> +arg_in:  3 uint           5
> +
> +[test]
> +name: threads
> +kernel_name: threads_int
> +dimensions: 1
> +global_size: 8 0 0
> +local_size:  8 0 0
> +arg_out: 0 buffer int[1] 28
> +arg_in:  1 buffer int[1] NULL
> +
> +!*/
> +
> +#define SIMPLE_TEST(TYPE) \
> +kernel void simple_##TYPE(global TYPE *out, local TYPE *mem, TYPE initial, TYPE value) { \
> +       *mem = initial; \
> +       TYPE a = atomic_add(mem, value); \
> +       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); \
> +       TYPE id = get_local_id(0); \
> +       atomic_add(mem, id); \
> +       barrier(CLK_LOCAL_MEM_FENCE); \
> +       *out = *mem; \
> +}
> +
> +SIMPLE_TEST(int)
> +SIMPLE_TEST(uint)
> +
> +THREADS_TEST(int)
> +THREADS_TEST(uint)
> diff --git a/tests/cl/program/execute/builtin/atomic/atomic_inc-local.cl b/tests/cl/program/execute/builtin/atomic/atomic_inc-local.cl
> new file mode 100644
> index 0000000..ba7e7a0
> --- /dev/null
> +++ b/tests/cl/program/execute/builtin/atomic/atomic_inc-local.cl
> @@ -0,0 +1,58 @@
> +/*!
> +[config]
> +name: atomic_inc 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] -2 -1
> +arg_in:  1 buffer int[1] NULL
> +arg_in:  2 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
> +
> +[test]
> +name: threads
> +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:  1 buffer int[1] NULL
> +
> +!*/
> +
> +#define SIMPLE_TEST(TYPE) \
> +kernel void simple_##TYPE(global TYPE *out, local TYPE *mem, TYPE initial) { \
> +       *mem = initial; \
> +       TYPE a = atomic_inc(mem); \
> +       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_inc(mem); \
> +       barrier(CLK_LOCAL_MEM_FENCE); \
> +       *out = *mem; \
> +}
> +
> +SIMPLE_TEST(int)
> +SIMPLE_TEST(uint)
> +
> +THREADS_TEST(int)
> +THREADS_TEST(uint)
> --
> 1.7.11.4
>
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/piglit