[Piglit] [PATCH 2/2] cl: Add tests for calls with special inputs

Jan Vesely jan.vesely at rutgers.edu
Wed Sep 5 14:28:15 UTC 2018


On Wed, Sep 5, 2018 at 10:16 AM Matt Arsenault <arsenm2 at gmail.com> wrote:

> ping
>

merged last week as a9752f23ef3531c6c44fc60cfff9811862fd594a.

Jan


>
> > On Aug 22, 2018, at 15:41, Matt Arsenault <arsenm2 at 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
> >
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20180905/3cb66f93/attachment-0001.html>


More information about the Piglit mailing list