On Mon, 2017-10-02 at 10:32 -0700, Matt Arsenault wrote: > ping > > > On Sep 19, 2017, at 19:25, Matt Arsenault <[email protected]> wrote: > > > > These aren't great since they need more test > > values, and generating half results is kind > > of a pain from any other tool. The perfect > > values used don't really stress the conversions, > > but this as at least enough to make sure the > > encoding is correct.
what is this test supposed to test? generic fp16 encoding
/functionality? if so, why is it better than vstore_half/vload_half
tests?
does it test fp16 version of clamp? why is it better than generated
clamp tests?
Jan
> > ---
> > tests/cl/program/execute/mad-mix.cl | 283
> > ++++++++++++++++++++++++++++++++++++
> > 1 file changed, 283 insertions(+)
> > create mode 100644 tests/cl/program/execute/mad-mix.cl
> >
> > diff --git a/tests/cl/program/execute/mad-mix.cl
> > b/tests/cl/program/execute/mad-mix.cl
> > new file mode 100644
> > index 000000000..dd7a5a516
> > --- /dev/null
> > +++ b/tests/cl/program/execute/mad-mix.cl
> > @@ -0,0 +1,283 @@
> > +/*!
> > +
> > +[config]
> > +name: f32 mad with conversion from f16
> > +clc_version_min: 10
> > +build_options: -cl-denorms-are-zero
> > +require_device_extensions: cl_khr_fp16
> > +
> > +dimensions: 1
> > +
> > +[test]
> > +name: mad mix f32 f16lo f16lo f16lo
> > +kernel_name: mad_mix_f32_f16lo_f16lo_f16lo
> > +global_size: 4 0 0
> > +
> > +arg_out: 0 buffer float[4] \
> > + 0.0 1.0 1.0 -1.0 \
redundant backslash. I don't see the point of this weird alignment. it
devirges from other test and if anything it makes it less readable.
> > +
> > +arg_in: 1 buffer half[4] \
> > + 0.0 1.0 0.0 -1.0
> > +
> > +arg_in: 2 buffer half[4] \
> > + 0.0 1.0 1.0 1.0 \
also here
> > +
> > +arg_in: 3 buffer half[4] \
> > + 0.0 0.0 1.0 0.0
> > +
> > +
> > +[test]
> > +name: mad mix f32 fneg(f16lo) f16lo f16lo
> > +kernel_name: mad_mix_f32_negf16lo_f16lo_f16lo
> > +global_size: 4 0 0
> > +
> > +arg_out: 0 buffer float[4] \
> > + 0.0 -1.0 1.0 1.0 \
and here
> > +
> > +arg_in: 1 buffer half[4] \
> > + 0.0 1.0 0.0 -1.0
> > +
> > +arg_in: 2 buffer half[4] \
> > + 0.0 1.0 1.0 1.0 \
and here
> > +
> > +arg_in: 3 buffer half[4] \
> > + 0.0 0.0 1.0 0.0
> > +
> > +
> > +[test]
> > +name: mad mix f32 f16lo f16lo f16hi
> > +kernel_name: mad_mix_f32_f16lo_f16lo_f16hi
> > +global_size: 4 0 0
> > +
> > +arg_out: 0 buffer float[4] \
> > + 0.0 \
> > + 1.0 \
> > + 1.0 \
> > + -1.0
> > +
> > +arg_in: 1 buffer half[4] \
> > + 0.0 \
> > + 1.0 \
> > + 0.0 \
> > + -1.0
> > +
> > +arg_in: 2 buffer half[4] \
> > + 0.0 \
> > + 1.0 \
> > + 1.0 \
> > + 1.0
> > +
> > +arg_in: 3 buffer half2[4] \
> > + 1000.0 0.0 \
> > + 1000.0 0.0 \
> > + 1000.0 1.0 \
> > + 1000.0 0.0
> > +
> > +
> > +[test]
> > +name: mad mix f32 f16lo f16lo neg(f16hi)
> > +kernel_name: mad_mix_f32_f16lo_f16lo_negf16hi
> > +global_size: 5 0 0
> > +
> > +arg_out: 0 buffer float[5] \
> > + 0.0 \
> > + 1.0 \
> > + -1.0 \
> > + -1.0 \
> > + 0.0
> > +
> > +arg_in: 1 buffer half[5] \
> > + 0.0 \
> > + 1.0 \
> > + 0.0 \
> > + -1.0 \
> > + 2.0
> > +
> > +arg_in: 2 buffer half[5] \
> > + 0.0 \
> > + 1.0 \
> > + 1.0 \
> > + 1.0 \
> > + 2.0
> > +
> > +arg_in: 3 buffer half2[5] \
> > + 1000.0 0.0 \
> > + 1000.0 0.0 \
> > + 1000.0 1.0 \
> > + 1000.0 0.0 \
> > + 1000.0 4.0
> > +
> > +
> > +[test]
> > +name: mad mix f16lo fneg(f16lo) f16lo f16lo
> > +kernel_name: mad_mix_f16lo_negf16lo_f16lo_f16lo
> > +global_size: 4 0 0
> > +
> > +arg_out: 0 buffer half[4] \
> > + 0.0 -1.0 1.0 1.0 \
and here
> > +
> > +arg_in: 1 buffer half[4] \
> > + 0.0 1.0 0.0 -1.0
> > +
> > +arg_in: 2 buffer half[4] \
> > + 0.0 1.0 1.0 1.0 \
and here
> > +
> > +arg_in: 3 buffer half[4] \
> > + 0.0 0.0 1.0 0.0
> > +
> > +
> > +[test]
> > +name: mad mix f16hi fneg(f16lo) f16lo f16lo
> > +kernel_name: mad_mix_f16hi_negf16lo_f16lo_f16lo
> > +global_size: 4 0 0
> > +
> > +arg_out: 0 buffer half2[4] \
> > + 2.0 0.0 \
> > + 2.0 -1.0 \
> > + 2.0 1.0 \
> > + 2.0 1.0
> > +
> > +arg_in: 1 buffer half[4] \
> > + 0.0 1.0 0.0 -1.0
> > +
> > +arg_in: 2 buffer half[4] \
> > + 0.0 1.0 1.0 1.0 \
here
> > +
> > +arg_in: 3 buffer half[4] \
> > + 0.0 0.0 1.0 0.0
> > +
> > +
> > +
> > +[test]
> > +name: mad mix f32 f16lo f16lo f16lo with clamp
> > +kernel_name: mad_mix_f32_f16lo_f16lo_f16lo_clamp
> > +global_size: 5 0 0
> > +
> > +arg_out: 0 buffer float[5] \
> > + 0.0 1.0 0.0 0.75 \
> > + 1.0
> > +
> > +arg_in: 1 buffer half[5] \
> > + 0.0 2.0 -2.0 0.5 \
> > + 0.5
> > +
> > +arg_in: 2 buffer half[5] \
> > + 0.0 1.0 1.0 0.5 \
> > + 1.0
> > +
> > +arg_in: 3 buffer half[5] \
> > + 0.0 1.0 1.0 0.5 \
> > + 0.5
> > +
> > +
> > +[test]
> > +name: mad mix f16lo f16lo f16lo f16lo with clamp
> > +kernel_name: mad_mix_f16lo_f16lo_f16lo_f16lo_clamp
> > +global_size: 5 0 0
> > +
> > +arg_out: 0 buffer half[5] \
> > + 0.0 1.0 0.0 0.75 \
> > + 1.0
> > +
> > +arg_in: 1 buffer half[5] \
> > + 0.0 2.0 -2.0 0.5 \
> > + 0.5
> > +
> > +arg_in: 2 buffer half[5] \
> > + 0.0 1.0 1.0 0.5 \
> > + 1.0
> > +
> > +arg_in: 3 buffer half[5] \
> > + 0.0 1.0 1.0 0.5 \
> > + 0.5
> > +
> > +
> > +[test]
> > +name: mad mix f16hi f16lo f16lo f16lo with clamp
> > +kernel_name: mad_mix_f16hi_f16lo_f16lo_f16lo_clamp
> > +global_size: 5 0 0
> > +
> > +arg_out: 0 buffer half2[5] \
> > + 2.0 0.0 \
> > + 2.0 1.0 \
> > + 2.0 0.0 \
> > + 2.0 0.75 \
> > + 2.0 1.0
> > +
> > +arg_in: 1 buffer half[5] \
> > + 0.0 2.0 -2.0 0.5 \
> > + 0.5
> > +
> > +arg_in: 2 buffer half[5] \
> > + 0.0 1.0 1.0 0.5 \
> > + 1.0
> > +
> > +arg_in: 3 buffer half[5] \
> > + 0.0 1.0 1.0 0.5 \
> > + 0.5
> > +
> > +
> > +!*/
> > +
> > +#pragma OPENCL EXTENSION cl_khr_fp16 : enable
> > +
> > +kernel void mad_mix_f32_f16lo_f16lo_f16lo(global float* out, global half*
> > in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + out[id] = (float)in0[id] * (float)in1[id] + (float)in2[id];
> > +}
> > +
> > +kernel void mad_mix_f32_negf16lo_f16lo_f16lo(global float* out, global
> > half* in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + out[id] = (float)-in0[id] * (float)in1[id] + (float)in2[id];
> > +}
> > +
> > +kernel void mad_mix_f32_f16lo_f16lo_f16hi(global float* out, global half*
> > in0, global half* in1, volatile global half2* in2)
> > +{
> > + int id = get_global_id(0);
> > + out[id] = (float)in0[id] * (float)in1[id] + (float)in2[id].y;
> > +}
> > +
> > +kernel void mad_mix_f32_f16lo_f16lo_negf16hi(global float* out, global
> > half* in0, global half* in1, volatile global half2* in2)
> > +{
> > + int id = get_global_id(0);
> > + out[id] = (float)in0[id] * (float)in1[id] + (float)-in2[id].y;;
> > +}
> > +
> > +kernel void mad_mix_f16lo_negf16lo_f16lo_f16lo(global half* out, global
> > half* in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + float mad = (float)-in0[id] * (float)in1[id] + (float)in2[id];
> > + out[id] = (half)mad;
> > +}
> > +
> > +kernel void mad_mix_f16hi_negf16lo_f16lo_f16lo(volatile global half2* out,
> > global half* in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + float mad = (float)-in0[id] * (float)in1[id] + (float)in2[id];
> > + half2 result = { 2.0h, (half)mad };
> > + out[id] = result;
> > +}
> > +
> > +kernel void mad_mix_f32_f16lo_f16lo_f16lo_clamp(global float* out, global
> > half* in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + float mad = (float)in0[id] * (float)in1[id] + (float)in2[id];
> > + out[id] = clamp(mad, 0.0f, 1.0f);
> > +}
> > +
> > +kernel void mad_mix_f16lo_f16lo_f16lo_f16lo_clamp(global half* out, global
> > half* in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + float mad = (float)in0[id] * (float)in1[id] + (float)in2[id];
> > + out[id] = clamp((half)mad, 0.0h, 1.0h);
> > +}
> > +
> > +kernel void mad_mix_f16hi_f16lo_f16lo_f16lo_clamp(volatile global half2*
> > out, global half* in0, global half* in1, global half* in2)
> > +{
> > + int id = get_global_id(0);
> > + float mad = (float)in0[id] * (float)in1[id] + (float)in2[id];
> > + half2 result = { 2.0h, clamp((half)mad, 0.0h, 1.0h) };
> > + out[id] = result;
> > +}
Jan
> > --
> > 2.11.0
> >
>
> _______________________________________________
> Piglit mailing list
> [email protected]
> https://lists.freedesktop.org/mailman/listinfo/piglit
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/piglit
