[Piglit] [PATCH] cl: Add tests for mad mix
Jan Vesely
jan.vesely at rutgers.edu
Thu Oct 5 15:09:01 UTC 2017
On Mon, 2017-10-02 at 10:32 -0700, Matt Arsenault wrote:
> ping
>
> > On Sep 19, 2017, at 19:25, Matt Arsenault <arsenm2 at gmail.com> 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
> Piglit at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/piglit
-------------- 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/20171005/c28141ff/attachment.sig>
More information about the Piglit
mailing list