[Piglit] [PATCH] cl: Add tests for mad mix
Matt Arsenault
arsenm2 at gmail.com
Thu Oct 5 19:36:43 UTC 2017
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.
v2: Fix backslash alignments
---
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..a5955361d
--- /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