On Wed, Sep 5, 2018 at 10:16 AM Matt Arsenault <arse...@gmail.com> wrote:
> ping
>
merged last week as a9752f23ef3531c6c44fc60cfff9811862fd594a.
Jan
>
> > On Aug 22, 2018, at 15:41, Matt Arsenault <arse...@gmail.com> wrote:
> >
> > Also fixes apparently missing coverage for special
> > input arguments not passed in registers.
> > ---
> > tests/cl/program/execute/calls-workitem-id.cl | 136 ++++++++++++++++++
> > 1 file changed, 136 insertions(+)
> >
> > diff --git a/tests/cl/program/execute/calls-workitem-id.cl
> b/tests/cl/program/execute/calls-workitem-id.cl
> > index 7edfad7e9..b42c85959 100644
> > --- a/tests/cl/program/execute/calls-workitem-id.cl
> > +++ b/tests/cl/program/execute/calls-workitem-id.cl
> > @@ -38,6 +38,56 @@ arg_out: 2 buffer uint[64] \
> > 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 \
> > 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
> >
> > +[test]
> > +name: Callee function stack passed get_local_id
> > +kernel_name: kernel_call_too_many_argument_regs_get_local_id_012
> > +dimensions: 3
> > +global_size: 8 4 2
> > +local_size: 8 4 2
> > +
> > +arg_out: 0 buffer uint[64] \
> > + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 \
> > + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 \
> > + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7 \
> > + 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7
> > +
> > +arg_out: 1 buffer uint[64] \
> > + 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 \
> > + 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3 \
> > + 0 0 0 0 0 0 0 0 1 1 1 1 1 1 1 1 \
> > + 2 2 2 2 2 2 2 2 3 3 3 3 3 3 3 3
> > +
> > +arg_out: 2 buffer uint[64] \
> > + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 \
> > + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 \
> > + 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 \
> > + 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
> > +
> > +[test]
> > +name: Callee function stack passed get_local_id with byval
> > +kernel_name: kernel_call_too_many_argument_regs_byval_get_local_id_012
> > +dimensions: 3
> > +global_size: 8 4 2
> > +local_size: 8 4 2
> > +
> > +arg_out: 0 buffer uint[64] \
> > + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 \
> > + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 \
> > + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52 \
> > + 45 46 47 48 49 50 51 52 45 46 47 48 49 50 51 52
> > +
> > +arg_out: 1 buffer uint[64] \
> > + 47 47 47 47 47 47 47 47 48 48 48 48 48 48 48 48 \
> > + 49 49 49 49 49 49 49 49 50 50 50 50 50 50 50 50 \
> > + 47 47 47 47 47 47 47 47 48 48 48 48 48 48 48 48 \
> > + 49 49 49 49 49 49 49 49 50 50 50 50 50 50 50 50
> > +
> > +arg_out: 2 buffer uint[64] \
> > + 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 \
> > + 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 \
> > + 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 \
> > + 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51 51
> > +
> > !*/
> >
> > #define NOINLINE __attribute__((noinline))
> > @@ -75,3 +125,89 @@ kernel void
> kernel_call_pass_get_global_id_012(global uint *out0,
> > {
> > func_get_global_id_012(out0, out1, out2);
> > }
> > +
> > +// On amdgcn, this will require the workitem IDs be passed as values
> > +// on the stack after the arguments.
> > +NOINLINE
> > +uint3 too_many_argument_regs_get_local_id_012(
> > + int arg0, int arg1, int arg2, int arg3,
> > + int arg4, int arg5, int arg6, int arg7,
> > + int arg8, int arg9, int arg10, int arg11,
> > + int arg12, int arg13, int arg14, int arg15,
> > + int arg16, int arg17, int arg18, int arg19,
> > + int arg20, int arg21, int arg22, int arg23,
> > + int arg24, int arg25, int arg26, int arg27,
> > + int arg28, int arg29, int arg30, int arg31)
> > +{
> > + uint3 id;
> > + id.x = get_local_id(0);
> > + id.y = get_local_id(1);
> > + id.z = get_local_id(2);
> > + return id;
> > +}
> > +
> > +kernel void kernel_call_too_many_argument_regs_get_local_id_012(global
> uint* out0, global uint* out1, global uint* out2)
> > +{
> > + uint id0 = get_global_id(0);
> > + uint id1 = get_global_id(1);
> > + uint id2 = get_global_id(2);
> > + uint flat_id = (id2 * get_global_size(1) + id1) *
> get_global_size(0) + id0;
> > +
> > + uint3 result = too_many_argument_regs_get_local_id_012(
> > + 1234, 999, 42, 5555, 8888, 9009, 777, 4242,
> > + 202020, 6359, 8344, 1443, 552323, 33424, 666, 98765,
> > + 2222, 232556, 57777, 934121, 94991, 1337, 0xdead, 0xbeef,
> > + 0x5555, 0x3333, 0x666, 0x4141, 0x1234, 0x8888, 0xaaaa,
> 0xbbbb);
> > +
> > + out0[flat_id] = result.x;
> > + out1[flat_id] = result.y;
> > + out2[flat_id] = result.z;
> > +}
> > +
> > +
> > +typedef struct ByValStruct {
> > + long array[9];
> > +} ByValStruct;
> > +
> > +// Same as previous, with an additional byval passed argument.
> > +NOINLINE
> > +uint3 too_many_argument_regs_byval_get_local_id_012(
> > + ByValStruct byval_arg,
> > + int arg0, int arg1, int arg2, int arg3,
> > + int arg4, int arg5, int arg6, int arg7,
> > + int arg8, int arg9, int arg10, int arg11,
> > + int arg12, int arg13, int arg14, int arg15,
> > + int arg16, int arg17, int arg18, int arg19,
> > + int arg20, int arg21, int arg22, int arg23,
> > + int arg24, int arg25, int arg26, int arg27,
> > + int arg28, int arg29, int arg30, int arg31)
> > +{
> > + uint3 id;
> > + id.x = get_local_id(0) + byval_arg.array[3]; // + 42 + 3
> > + id.y = get_local_id(1) + byval_arg.array[5]; // + 42 + 5
> > + id.z = get_local_id(2) + byval_arg.array[8]; // + 42 + 8
> > + return id;
> > +}
> > +
> > +kernel void
> kernel_call_too_many_argument_regs_byval_get_local_id_012(global uint*
> out0, global uint* out1, global uint* out2)
> > +{
> > + uint id0 = get_global_id(0);
> > + uint id1 = get_global_id(1);
> > + uint id2 = get_global_id(2);
> > + uint flat_id = (id2 * get_global_size(1) + id1) *
> get_global_size(0) + id0;
> > +
> > + ByValStruct byval;
> > + for (int i = 0; i < 9; ++i)
> > + byval.array[i] = 42 + i;
> > +
> > + uint3 result = too_many_argument_regs_byval_get_local_id_012(
> > + byval,
> > + 1234, 999, 42, 5555, 8888, 9009, 777, 4242,
> > + 202020, 6359, 8344, 1443, 552323, 33424, 666, 98765,
> > + 2222, 232556, 57777, 934121, 94991, 1337, 0xdead, 0xbeef,
> > + 0x5555, 0x3333, 0x666, 0x4141, 0x1234, 0x8888, 0xaaaa,
> 0xbbbb);
> > +
> > + out0[flat_id] = result.x;
> > + out1[flat_id] = result.y;
> > + out2[flat_id] = result.z;
> > +}
> > --
> > 2.17.1
> >
>
>
_______________________________________________
Piglit mailing list
Piglit@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/piglit
