On Mon, 2017-01-16 at 11:02 -0800, [email protected] wrote: > From: Matt Arsenault <[email protected]>
Reviewed-by: Jan Vesely <[email protected]> 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]; > +}
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/piglit
