[Piglit] [PATCH 2/2] cl: Add tests for fdiv with neg/abs inputs
arsenm2 at gmail.com
arsenm2 at gmail.com
Mon Jan 16 19:02:27 UTC 2017
From: Matt Arsenault <arsenm2 at gmail.com>
---
tests/cl/program/execute/fdiv-modifiers-f32.cl | 257 ++++++++++++++++++++++++
tests/cl/program/execute/fdiv-modifiers-f64.cl | 262 +++++++++++++++++++++++++
2 files changed, 519 insertions(+)
create mode 100644 tests/cl/program/execute/fdiv-modifiers-f32.cl
create mode 100644 tests/cl/program/execute/fdiv-modifiers-f64.cl
diff --git a/tests/cl/program/execute/fdiv-modifiers-f32.cl b/tests/cl/program/execute/fdiv-modifiers-f32.cl
new file mode 100644
index 000000000..ce74dc5e2
--- /dev/null
+++ b/tests/cl/program/execute/fdiv-modifiers-f32.cl
@@ -0,0 +1,257 @@
+/*!
+[config]
+name: fdiv with neg or abs inputs
+clc_version_min: 10
+build_options: -cl-denorms-are-zero
+
+dimensions: 1
+global_size: 12 0 0
+
+## Division ##
+
+[test]
+name: fdiv -x, y
+kernel_name: fdiv_neg_x_y_f32
+
+arg_out: 0 buffer float[12] \
+ -1.0 -1.0 1.0 1.0 \
+ -2.0 2.0 2.0 -2.0 \
+ -0.5 0.5 0.5 -0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv x, -y
+kernel_name: fdiv_x_neg_y_f32
+
+arg_out: 0 buffer float[12] \
+ -1.0 -1.0 1.0 1.0 \
+ -2.0 2.0 2.0 -2.0 \
+ -0.5 0.5 0.5 -0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv -x, -y
+kernel_name: fdiv_neg_x_neg_y_f32
+
+arg_out: 0 buffer float[12] \
+ 1.0 1.0 -1.0 -1.0 \
+ 2.0 -2.0 -2.0 2.0 \
+ 0.5 -0.5 -0.5 0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv |x|, y
+kernel_name: fdiv_abs_x_y_f32
+
+arg_out: 0 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 0.5 -0.5 0.5 -0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv x, |y|
+kernel_name: fdiv_x_abs_y_f32
+
+arg_out: 0 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 2.0 2.0 -2.0 -2.0 \
+ 0.5 0.5 -0.5 -0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv |x|, |y|
+kernel_name: fdiv_abs_x_abs_y_f32
+
+arg_out: 0 buffer float[12] \
+ 1.0 1.0 1.0 1.0 \
+ 2.0 2.0 2.0 2.0 \
+ 0.5 0.5 0.5 0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv -|x|, y
+kernel_name: fdiv_neg_abs_x_y_f32
+
+arg_out: 0 buffer float[12] \
+ -1.0 1.0 1.0 -1.0 \
+ -2.0 2.0 -2.0 2.0 \
+ -0.5 0.5 -0.5 0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv x, -|y|
+kernel_name: fdiv_x_neg_abs_y_f32
+
+arg_out: 0 buffer float[12] \
+ -1.0 1.0 -1.0 1.0 \
+ -2.0 -2.0 2.0 2.0 \
+ -0.5 -0.5 0.5 0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv -|x|, -|y|
+kernel_name: fdiv_neg_abs_x_neg_abs_y_f32
+
+arg_out: 0 buffer float[12] \
+ 1.0 1.0 1.0 1.0 \
+ 2.0 2.0 2.0 2.0 \
+ 0.5 0.5 0.5 0.5
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer float[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv 4.0, y
+kernel_name: fdiv_4_y_f32
+
+arg_out: 0 buffer float[12] \
+ 4.0 -4.0 8.0 -8.0 \
+ 1.0 1.0 -1.0 -1.0 \
+ 2.0 -2.0 nan 0.0
+
+arg_in: 1 buffer float[12] \
+ 1.0 -1.0 0.5 -0.5 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 -2.0 0.0 inf
+
+!*/
+
+kernel void fdiv_neg_x_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -in0[id] / in1[id];
+}
+
+kernel void fdiv_x_neg_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = in0[id] / -in1[id];
+}
+
+kernel void fdiv_neg_x_neg_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -in0[id] / -in1[id];
+}
+
+kernel void fdiv_abs_x_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = fabs(in0[id]) / in1[id];
+}
+
+kernel void fdiv_x_abs_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = in0[id] / fabs(in1[id]);
+}
+
+kernel void fdiv_abs_x_abs_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = fabs(in0[id]) / fabs(in1[id]);
+}
+
+kernel void fdiv_neg_abs_x_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -fabs(in0[id]) / in1[id];
+}
+
+kernel void fdiv_x_neg_abs_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = in0[id] / -fabs(in1[id]);
+}
+
+kernel void fdiv_neg_abs_x_neg_abs_y_f32(global float* out, global float* in0, global float* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -fabs(in0[id]) / -fabs(in1[id]);
+}
+
+kernel void fdiv_4_y_f32(global float* out, global float* in0)
+{
+ int id = get_global_id(0);
+ out[id] = 4.0f / in0[id];
+}
diff --git a/tests/cl/program/execute/fdiv-modifiers-f64.cl b/tests/cl/program/execute/fdiv-modifiers-f64.cl
new file mode 100644
index 000000000..57d21769a
--- /dev/null
+++ b/tests/cl/program/execute/fdiv-modifiers-f64.cl
@@ -0,0 +1,262 @@
+/*!
+[config]
+name: fdiv with neg or abs inputs
+clc_version_min: 10
+require_device_extensions: cl_khr_fp64
+
+dimensions: 1
+global_size: 12 0 0
+
+## Division ##
+
+
+[test]
+name: fdiv -x, y
+kernel_name: fdiv_neg_x_y_f64
+
+arg_out: 0 buffer double[12] \
+ -1.0 -1.0 1.0 1.0 \
+ -2.0 2.0 2.0 -2.0 \
+ -0.5 0.5 0.5 -0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv x, -y
+kernel_name: fdiv_x_neg_y_f64
+
+arg_out: 0 buffer double[12] \
+ -1.0 -1.0 1.0 1.0 \
+ -2.0 2.0 2.0 -2.0 \
+ -0.5 0.5 0.5 -0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv -x, -y
+kernel_name: fdiv_neg_x_neg_y_f64
+
+arg_out: 0 buffer double[12] \
+ 1.0 1.0 -1.0 -1.0 \
+ 2.0 -2.0 -2.0 2.0 \
+ 0.5 -0.5 -0.5 0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv |x|, y
+kernel_name: fdiv_abs_x_y_f64
+
+arg_out: 0 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 0.5 -0.5 0.5 -0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv x, |y|
+kernel_name: fdiv_x_abs_y_f64
+
+arg_out: 0 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 2.0 2.0 -2.0 -2.0 \
+ 0.5 0.5 -0.5 -0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv |x|, |y|
+kernel_name: fdiv_abs_x_abs_y_f64
+
+arg_out: 0 buffer double[12] \
+ 1.0 1.0 1.0 1.0 \
+ 2.0 2.0 2.0 2.0 \
+ 0.5 0.5 0.5 0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv -|x|, y
+kernel_name: fdiv_neg_abs_x_y_f64
+
+arg_out: 0 buffer double[12] \
+ -1.0 1.0 1.0 -1.0 \
+ -2.0 2.0 -2.0 2.0 \
+ -0.5 0.5 -0.5 0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv x, -|y|
+kernel_name: fdiv_x_neg_abs_y_f64
+
+arg_out: 0 buffer double[12] \
+ -1.0 1.0 -1.0 1.0 \
+ -2.0 -2.0 2.0 2.0 \
+ -0.5 -0.5 0.5 0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+[test]
+name: fdiv -|x|, -|y|
+kernel_name: fdiv_neg_abs_x_neg_abs_y_f64
+
+arg_out: 0 buffer double[12] \
+ 1.0 1.0 1.0 1.0 \
+ 2.0 2.0 2.0 2.0 \
+ 0.5 0.5 0.5 0.5
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 1.0 -1.0 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 2.0 -2.0 -2.0
+
+arg_in: 2 buffer double[12] \
+ 1.0 -1.0 -1.0 1.0 \
+ 2.0 -2.0 2.0 -2.0 \
+ 4.0 -4.0 4.0 -4.0
+
+
+[test]
+name: fdiv 4.0, y
+kernel_name: fdiv_4_y_f64
+
+arg_out: 0 buffer double[12] \
+ 4.0 -4.0 8.0 -8.0 \
+ 1.0 1.0 -1.0 -1.0 \
+ 2.0 -2.0 nan 0.0
+
+arg_in: 1 buffer double[12] \
+ 1.0 -1.0 0.5 -0.5 \
+ 4.0 4.0 -4.0 -4.0 \
+ 2.0 -2.0 0.0 inf
+
+!*/
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+
+kernel void fdiv_neg_x_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -in0[id] / in1[id];
+}
+
+kernel void fdiv_x_neg_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = in0[id] / -in1[id];
+}
+
+kernel void fdiv_neg_x_neg_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -in0[id] / -in1[id];
+}
+
+kernel void fdiv_abs_x_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = fabs(in0[id]) / in1[id];
+}
+
+kernel void fdiv_x_abs_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = in0[id] / fabs(in1[id]);
+}
+
+kernel void fdiv_abs_x_abs_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = fabs(in0[id]) / fabs(in1[id]);
+}
+
+kernel void fdiv_neg_abs_x_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -fabs(in0[id]) / in1[id];
+}
+
+kernel void fdiv_x_neg_abs_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = in0[id] / -fabs(in1[id]);
+}
+
+kernel void fdiv_neg_abs_x_neg_abs_y_f64(global double* out, global double* in0, global double* in1)
+{
+ int id = get_global_id(0);
+ out[id] = -fabs(in0[id]) / -fabs(in1[id]);
+}
+
+kernel void fdiv_4_y_f64(global double* out, global double* in0)
+{
+ int id = get_global_id(0);
+ out[id] = 4.0 / in0[id];
+}
--
2.11.0
More information about the Piglit
mailing list