diff options
author | Matt Arsenault <arsenm2@gmail.com> | 2017-01-16 11:02:27 -0800 |
---|---|---|
committer | Jan Vesely <jan.vesely@rutgers.edu> | 2017-01-27 13:03:18 -0500 |
commit | 01eb7321b46dbd11d71a2ca07217ef387486db3a (patch) | |
tree | 586396586c5c9f4f97ba99cd61b84ae0ac7261df | |
parent | d28c2acf06d964696b01485cc252a6918bfa6936 (diff) |
cl: Add tests for fdiv with neg/abs inputs
Reviewed-by: Jan Vesely <jan.vesely@rutgers.edu>
-rw-r--r-- | tests/cl/program/execute/fdiv-modifiers-f32.cl | 257 | ||||
-rw-r--r-- | tests/cl/program/execute/fdiv-modifiers-f64.cl | 262 |
2 files changed, 519 insertions, 0 deletions
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]; +} |