[Piglit] [PATCH] cl: Add tests for some cases that were broken with function calls

[Jan Vesely]

On Thu, 2019-05-30 at 08:40 -0400, Matt Arsenault wrote:
> Ping
> 
> > On May 23, 2019, at 7:59 PM, arsenm2 at gmail.com wrote:
> > 
> > From: Matt Arsenault <arsenm2 at gmail.com>
> > 
> > ---
> > .../program/execute/call-clobbers-amdgcn.cl   | 102 ++++++++++++++++++
> > 1 file changed, 102 insertions(+)
> > 
> > diff --git a/tests/cl/program/execute/call-clobbers-amdgcn.cl b/tests/cl/program/execute/call-clobbers-amdgcn.cl
> > index 18e657ce3..b0a1f8c70 100644
> > --- a/tests/cl/program/execute/call-clobbers-amdgcn.cl
> > +++ b/tests/cl/program/execute/call-clobbers-amdgcn.cl
> > @@ -19,6 +19,49 @@ dimensions: 1
> > global_size: 1 0 0
> > arg_out: 0 buffer int[1] 0xabcd1234
> > 
> > +[test]
> > +name: Conditional call
> > +kernel_name: conditional_call
> > +dimensions: 1
> > +local_size: 64 0 0
> > +global_size: 64 0 0
> > +arg_out: 0 buffer int[64] \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234
> > +
> > +[test]
> > +name: Conditional call partial dispatch
> > +kernel_name: conditional_call
> > +dimensions: 1
> > +local_size: 16 0 0
> > +global_size: 16 0 0
> > +arg_out: 0 buffer int[16] \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 \
> > +  0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234 0xabcd1234
> > +
> > +
> > +[test]
> > +name: Skip call no lanes
> > +kernel_name: skip_call_no_lanes
> > +dimensions: 1
> > +local_size: 64 0 0
> > +global_size: 64 0 0
> > +arg_out: 0 buffer int[64] \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123 \
> > +  123 123 123 123 123 123 123 123
> > +
> > !*/
> > 
> > #ifndef __AMDGCN__
> > @@ -65,3 +108,62 @@ kernel void call_clobber_v40(__global int* ret)
> >                   : "v40");
> >     *ret = tmp;
> > }
> > +
> > +__attribute__((noinline))
> > +void spill_sgpr_to_csr_vgpr()
> > +{
> > +    __asm volatile(
> > +        "s_nop 1" :::
> > +        "v0","v1","v2","v3","v4","v5","v6","v7",
> > +        "v8","v9","v10","v11","v12","v13","v14","v15",
> > +        "v16","v17","v18","v19","v20","v21","v22","v23",
> > +        "v24","v25","v26","v27","v28","v29","v30","v31",
> > +
> > +        "s0","s1","s2","s3","s4","s5","s6","s7",
> > +        "s8","s9","s10","s11","s12","s13","s14","s15",
> > +        "s16","s17","s18","s19","s20","s21","s22","s23",
> > +        "s24","s25","s26","s27","s28","s29","s30","s31",
> > +        "s32", "s33", "s34", "s35", "s36", "s37", "s38");
> > +}
> > +
> > +// A CSR VGPR needs to be spilled/restored in the prolog/epilog, but
> > +// all lanes need to be made active to avoid clobbering lanes that did
> > +// not enter the call.
> > +kernel void conditional_call(global int* ret)
> > +{
> > +    __asm volatile("v_mov_b32 v32, 0xabcd1234" : : : "v32");
> > +
> > +    int id = get_local_id(0);
> > +    if (id == 0)
> > +    {
> > +        spill_sgpr_to_csr_vgpr();
> > +    }
> > +
> > +    int tmp;
> > +    __asm volatile("v_mov_b32 %0, v32"
> > +                   : "=v"(tmp)
> > +                   :
> > +                   : "v32");
> > +    ret[id] = tmp;
> > +}
> > +
> > +__attribute__((noinline))
> > +void hang_if_all_inactive()
> > +{
> > +    __builtin_amdgcn_s_sendmsghalt(0, 0);
> > +}
> > +
> > +// If all lanes could be dynamically false, the call must not be taken
> > +// in case a side effecting scalar op is called inside.
> > +kernel void skip_call_no_lanes(global int* ret)
> > +{
> > +    int divergent_false;
> > +    __asm volatile("v_mov_b32 %0, 0" : "=v"(divergent_false));
> > +
> > +    if (divergent_false)
> > +    {
> > +        hang_if_all_inactive();

this looks like it will hang the GPU on test failure, which is a no-
go.

Jan

> > +    }
> > +
> > +    ret[get_global_id(0)] = 123;
> > +}
> > -- 
> > 2.17.1
> > 
> 
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/piglit

-- 
Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20190601/b2e2d444/attachment.sig>