[Piglit] [PATCH 2/2] cl: Add tests for fdiv with neg/abs inputs
Jan Vesely
jan.vesely at rutgers.edu
Fri Jan 20 20:34:13 UTC 2017
On Mon, 2017-01-16 at 11:02 -0800, arsenm2 at gmail.com wrote:
> From: Matt Arsenault <arsenm2 at gmail.com>
Reviewed-by: Jan Vesely <jan.vesely at rutgers.edu>
Jan
>
> ---
> 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];
> +}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20170120/3e29f272/attachment.sig>
More information about the Piglit
mailing list