[Piglit] [PATCH] cl: Add tests for some cases that were broken with function calls
Jan Vesely
jan.vesely at rutgers.edu
Sun Jun 2 18:36:11 UTC 2019
On Sun, 2019-06-02 at 09:49 -0400, Matt Arsenault wrote:
> > On Jun 1, 2019, at 10:57 PM, Jan Vesely <jan.vesely at rutgers.edu> wrote:
> >
> > 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
>
> Is there a way to specify a timeout? The alternatives require more ABI support
I think there's a timeout and it should be used by default. It also
defaults to None, so PiglitCLTest class would need to be modified to
pass something sensible by default.
Dylan would know more how to set this up.
The other problem is that even with killed process hung GPU usually
makes the machine unable to suspend or reboot on its own, which kills
remote testing.
I'd need to recheck if that's still the case with linux-5.1.
Jan
> _______________________________________________
> 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/20190602/f4c32b5d/attachment-0001.sig>
More information about the Piglit
mailing list