[Piglit] [PATCH] cl: Add tests for mad mix

Matt Arsenault arsenm2 at gmail.com
Mon Oct 2 17:32:24 UTC 2017


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.
> ---
> 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   \
> +
> +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 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   \
> +
> +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 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  \
> +
> +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 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    \
> +
> +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;
> +}
> -- 
> 2.11.0
> 



More information about the Piglit mailing list