[Piglit] [PATCH] cl: Add local memory test

[Tom Stellard]

On Thu, Jun 13, 2013 at 06:38:34PM -0500, Aaron Watry wrote:
> It's probably noting somewhere that these tests will fail on CL
> implementations where the maximum local workgroup size is < 16 (e.g.
> CPU implementations that limit to 1 thread per core).
>

Good point.  I think this applies to other tests as well.  Maybe I'll
add a check for this in cl-program-tester and skip tests that use a
larger than supported workgroup size.

> Also, I'm not sure if it really matters, but instead of:
> index2 = index + 1;
> index2 = index2 == 4 ? 0 : index2
> 
> You could instead use:
> index2 = (index + 1) % 4
> 
> Which removes a dependency on flow control actually working :)
> 

I'll change this. I was trying to avoid integer division which is why I
didn't use the modulo operator, but modulo powers of two doesn't involve
division, so it should be fine.

-Tom

> Either way, with a comment about local workgroup max sizes added:
> Reviewed-by: Aaron Watry <awatry at gmail.com>
> 
> On Wed, Jun 12, 2013 at 7:39 PM, Tom Stellard <tom at stellard.net> wrote:
> > From: Tom Stellard <thomas.stellard at amd.com>
> >
> > ---
> >  tests/cl/program/execute/local-memory.cl | 81 ++++++++++++++++++++++++++++++++
> >  1 file changed, 81 insertions(+)
> >  create mode 100644 tests/cl/program/execute/local-memory.cl
> >
> > diff --git a/tests/cl/program/execute/local-memory.cl b/tests/cl/program/execute/local-memory.cl
> > new file mode 100644
> > index 0000000..3d2c55b
> > --- /dev/null
> > +++ b/tests/cl/program/execute/local-memory.cl
> > @@ -0,0 +1,81 @@
> > +/*!
> > +[config]
> > +name: local_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
> > +
> > +[test]
> > +name: (16 x 1 x 1) (16 x 1 x 1)
> > +kernel_name: local_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
> > +
> > +[test]
> > +name: (64 x 1 x 1) (4 x 1 x 1)
> > +kernel_name: local_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
> > +
> > +[test]
> > +name: 2 local memory objects
> > +kernel_name: local_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
> > +
> > +!*/
> > +
> > +kernel void simple(global int *out) {
> > +       local volatile int local_mem[2];
> > +       local_mem[0] = 0xffffffff;
> > +       local_mem[1] = 0xffffffff;
> > +       out[0] = local_mem[0];
> > +       out[1] = local_mem[1];
> > +}
> > +
> > +kernel void local_memory_one_work_group(global int *out) {
> > +       local int local_mem[16];
> > +       int index = get_local_id(0);
> > +       int index2;
> > +       local_mem[index] = index;
> > +       index2 = index + 1;
> > +       index2 = index2 == 16 ? 0 : index2;
> > +       barrier(CLK_LOCAL_MEM_FENCE);
> > +       out[index] = local_mem[index2];
> > +}
> > +
> > +kernel void local_memory_many_work_groups(global int *out) {
> > +       local int local_mem[4];
> > +       int index = get_local_id(0);
> > +       int global_id = get_global_id(0);
> > +       int index2;
> > +       local_mem[index] = index;
> > +       index2 = index + 1;
> > +       index2 = index2 == 4 ? 0 : index2;
> > +       barrier(CLK_LOCAL_MEM_FENCE);
> > +       out[global_id] = local_mem[index2];
> > +}
> > +
> > +kernel void local_memory_two_objects(global int *out) {
> > +       local int local_mem0[4];
> > +       local int local_mem1[4];
> > +
> > +       int index = get_local_id(0);
> > +       local_mem0[index] = index;
> > +       local_mem1[index] = index * 2;
> > +       int index2 = 3 - index;
> > +       barrier(CLK_LOCAL_MEM_FENCE);
> > +       out[index] = local_mem0[index2];
> > +       out[index + 4] = local_mem1[index2];
> > +}
> > --
> > 1.7.11.4
> >
> > _______________________________________________
> > Piglit mailing list
> > Piglit at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/piglit