[Piglit] [PATCH 03/10] cl: Add tests for different versions of fmin / fmax.
Jan Vesely
jan.vesely at rutgers.edu
Tue Dec 6 20:08:37 UTC 2016
On Tue, 2016-12-06 at 11:12 -0800, Matt Arsenault wrote:
> > On Dec 5, 2016, at 13:35, Jan Vesely <jan.vesely at rutgers.edu> wrote:
> >
> > On Mon, 2016-12-05 at 09:48 -0800, arsenm2 at gmail.com <mailto:arsenm2 at gmail.com> wrote:
> > > From: Matt Arsenault <arsenm2 at gmail.com>
> > >
> > > These do not use the normal simple format because the number
> > > of combinations that need to be tested is simply too large,
> > > especially when tests for min3/max3 are added.
> > >
> > > The unordered compare tests could be improved. Currently they truly
> > > test the unordered compare because of LLVM bug 21610, but
> > > ideally that would be fixed.
> > > ---
> > > tests/cl/program/CMakeLists.cl.txt | 1 +
> > > .../cl/program/execute/scalar-comparison-float.cl | 105 +++++
> > > tests/cl/program/float-min-max-kernels.cl | 492 +++++++++++++++++++++
> > > tests/cl/program/float-min-max.cpp | 475 ++++++++++++++++++++
> > > 4 files changed, 1073 insertions(+)
> > > create mode 100644 tests/cl/program/float-min-max-kernels.cl
> > > create mode 100644 tests/cl/program/float-min-max.cpp
> > >
> > > diff --git a/tests/cl/program/CMakeLists.cl.txt b/tests/cl/program/CMakeLists.cl.txt
> > > index c8d7307..5ef0f6b 100644
> > > --- a/tests/cl/program/CMakeLists.cl.txt
> > > +++ b/tests/cl/program/CMakeLists.cl.txt
> > > @@ -2,3 +2,4 @@ piglit_cl_add_program_test (tester program-tester.c)
> > > piglit_cl_add_program_test (max-work-item-sizes max-work-item-sizes.c)
> > > piglit_cl_add_program_test (bitcoin-phatk bitcoin-phatk.c)
> > > piglit_cl_add_program_test (predefined-macros predefined-macros.c)
> > > +piglit_cl_add_program_test (float-min-max float-min-max.cpp)
> > > diff --git a/tests/cl/program/execute/scalar-comparison-float.cl b/tests/cl/program/execute/scalar-comparison-float.cl
> > > index 4891fc5..598fae0 100644
> > > --- a/tests/cl/program/execute/scalar-comparison-float.cl
> > > +++ b/tests/cl/program/execute/scalar-comparison-float.cl
> > > @@ -148,6 +148,71 @@ arg_in: 1 float -3.5
> > > arg_in: 2 float 4.5
> > > arg_out: 0 buffer int[1] 1
> > >
> > > +
> > > +[test]
> > > +name: select_max_gt
> > > +kernel_name: select_max_gt
> > > +global_size: 24 0 0
> > > +
> > > +arg_out: 0 buffer float[24] \
> > > + 0.0 1.0 2.0 2.0 0.0 0.0 \
> > > + NAN NAN 1.0 NAN -1.0 NAN \
> > > + 0.0 0.0 97.0 INF INF INF \
> > > + NAN NAN INF NAN -INF INF
> > > +
> > > +arg_in: 1 buffer float[24] \
> > > + 0.0 1.0 1.0 2.0 0.0 -1.0 \
> > > + NAN 1.0 NAN -1.0 NAN 0.0 \
> > > + 0.0 -0.0 37.0 INF INF -INF \
> > > + -INF INF NAN -INF NAN 0.0
> > > +
> > > +arg_in: 2 buffer float[24] \
> > > + 0.0 1.0 2.0 1.0 -1.0 0.0 \
> > > + NAN NAN 1.0 NAN -1.0 NAN \
> > > + -0.0 0.0 97.0 INF -INF INF \
> > > + -INF NAN INF NAN -INF INF
> > > +
> > > +[test]
> > > +name: select_max_gte
> > > +kernel_name: select_max_gte
> > > +global_size: 15 0 0
> > > +
> > > +arg_out: 0 buffer float[15] \
> > > + 0.0 1.0 2.0 2.0 0.0 0.0 \
> > > + NAN NAN 1.0 NAN -1.0 NAN \
> > > + 0.0 0.0 97.0
> > > +
> > > +arg_in: 1 buffer float[15] \
> > > + 0.0 1.0 1.0 2.0 0.0 -1.0 \
> > > + NAN 1.0 NAN -1.0 NAN 0.0 \
> > > + 0.0 -0.0 37.0
> > > +
> > > +arg_in: 2 buffer float[15] \
> > > + 0.0 1.0 2.0 1.0 -1.0 0.0 \
> > > + NAN NAN 1.0 NAN -1.0 NAN \
> > > + -0.0 0.0 97.0
> > > +
> > > +[test]
> > > +name: select_min_gt
> > > +kernel_name: select_min_gt
> > > +global_size: 15 0 0
> > > +
> > > +arg_out: 0 buffer float[15] \
> > > + 0.0 1.0 1.0 1.0 -1.0 -1.0 \
> > > + NAN NAN NAN NAN NAN NAN \
> > > + 0.0 0.0 37.0
> > > +
> > > +arg_in: 1 buffer float[15] \
> > > + 0.0 1.0 1.0 2.0 0.0 -1.0 \
> > > + NAN 1.0 NAN -1.0 NAN 0.0 \
> > > + 0.0 -0.0 37.0
> > > +
> > > +arg_in: 2 buffer float[15] \
> > > + 0.0 1.0 2.0 1.0 -1.0 0.0 \
> > > + NAN NAN 1.0 NAN -1.0 NAN \
> > > + -0.0 0.0 97.0
> > > +
> > > +
> > > !*/
> > >
> > > kernel void eq(global int* out, float a, float b) {
> > > @@ -173,3 +238,43 @@ kernel void lt(global int* out, float a, float b) {
> > > kernel void lte(global int* out, float a, float b) {
> > > out[0] = a <= b;
> > > }
> > > +
> > > +kernel void select_max_gt(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] > b[id]) ? a[id] : b[id];
> > > +}
> > > +
> > > +kernel void select_max_gte(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] >= b[id]) ? a[id] : b[id];
> > > +}
> > > +
> > > +kernel void select_min_gt(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] > b[id]) ? b[id] : a[id];
> > > +}
> > > +
> > > +kernel void select_min_gte(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] >= b[id]) ? b[id] : a[id];
> > > +}
> > > +
> > > +kernel void select_min_lt(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] < b[id]) ? a[id] : b[id];
> > > +}
> > > +
> > > +kernel void select_max_lt(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] < b[id]) ? b[id] : a[id];
> > > +}
> > > +
> > > +kernel void select_min_lte(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] <= b[id]) ? a[id] : b[id];
> > > +}
> > > +
> > > +kernel void select_max_lte(global float* restrict out, global float* restrict a, global float* restrict b) {
> > > + int id = get_global_id(0);
> > > + out[id] = (a[id] <= b[id]) ? b[id] : a[id];
> > > +}
> > > diff --git a/tests/cl/program/float-min-max-kernels.cl b/tests/cl/program/float-min-max-kernels.cl
> > > new file mode 100644
> > > index 0000000..09f31d4
> > > --- /dev/null
> > > +++ b/tests/cl/program/float-min-max-kernels.cl
> > > @@ -0,0 +1,492 @@
> > > +#pragma OPENCL EXTENSION cl_khr_fp64 : enable
> > > +
> > > +kernel void select_max_gt_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a > b) ? a : b;
> > > +}
> > > +
> > > +kernel void select_max_ge_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a >= b) ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_gt_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a > b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_ge_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a >= b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_lt_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a < b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_le_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a <= b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_lt_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a < b) ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_le_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a <= b) ? a : b;
> > > +}
> > > +
> > > +kernel void test_fmin_f32(global float* restrict out,
> > > + constant float* in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = fmin(a, b);
> > > +}
> > > +
> > > +kernel void test_fmax_f32(global float* restrict out,
> > > + constant float* in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + out[n * idx + idy] = fmax(a, b);
> > > +}
> > > +
> > > +// FIXME: It is a canonicalization bug that an unordered comparison is
> > > +// emitted for this if the intermediate cmp variable is used.
> > > +kernel void select_max_ugt_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a <= b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +kernel void select_max_uge_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a < b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_ugt_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a <= b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_uge_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a < b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_ult_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a >= b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_ule_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a > b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_ult_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a >= b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_ule_f32(global float* restrict out,
> > > + constant float* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + float a = in[idx];
> > > + float b = in[idy];
> > > +
> > > + bool cmp = !(a > b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +#if cl_khr_fp64
> > > +kernel void select_max_gt_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a > b) ? a : b;
> > > +}
> > > +
> > > +kernel void select_max_ge_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a >= b) ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_gt_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a > b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_ge_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a >= b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_lt_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a < b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_le_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a <= b) ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_lt_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a < b) ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_le_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = (a <= b) ? a : b;
> > > +}
> > > +
> > > +kernel void test_fmin_f64(global double* restrict out,
> > > + constant double* in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = fmin(a, b);
> > > +}
> > > +
> > > +kernel void test_fmax_f64(global double* restrict out,
> > > + constant double* in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + out[n * idx + idy] = fmax(a, b);
> > > +}
> > > +
> > > +// FIXME: It is a canonicalization bug that an unordered comparison is
> > > +// emitted for this if the intermediate cmp variable is used.
> > > +kernel void select_max_ugt_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a <= b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +kernel void select_max_uge_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a < b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_ugt_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a <= b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_uge_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a < b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_ult_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a >= b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_max_ule_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a > b);
> > > + out[n * idx + idy] = cmp ? b : a;
> > > +}
> > > +
> > > +kernel void select_min_ult_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a >= b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +kernel void select_min_ule_f64(global double* restrict out,
> > > + constant double* restrict in,
> > > + int n)
> > > +{
> > > + int idx = get_global_id(0);
> > > + int idy = get_global_id(1);
> > > +
> > > + double a = in[idx];
> > > + double b = in[idy];
> > > +
> > > + bool cmp = !(a > b);
> > > + out[n * idx + idy] = cmp ? a : b;
> > > +}
> > > +
> > > +#endif
> > > diff --git a/tests/cl/program/float-min-max.cpp b/tests/cl/program/float-min-max.cpp
> > > new file mode 100644
> > > index 0000000..296b446
> > > --- /dev/null
> > > +++ b/tests/cl/program/float-min-max.cpp
> > > @@ -0,0 +1,475 @@
> > > +
> > > +extern "C" {
> > > +#include "piglit-framework-cl-program.h"
> > > +}
> > > +
> > > +PIGLIT_CL_PROGRAM_TEST_CONFIG_BEGIN
> > > +
> > > + config.name = "Run kernels which will use select min / max instructions";
> > > +
> > > + config.run_per_device = true;
> > > +
> > > + config.program_source_file = "float-min-max-kernels.cl";
> > > + config.kernel_name = NULL; // We have many kernels.
> > > +
> > > +PIGLIT_CL_PROGRAM_TEST_CONFIG_END
> > > +
> > > +
> > > +template <typename Real>
> > > +struct TestFunction
> > > +{
> > > + typedef Real (*MinMaxFunc)(Real, Real);
> > > +
> > > + const char* kernel_name;
> > > + MinMaxFunc ref_func;
> > > +};
> > > +
> > > +static const size_t n_cases = 32;
> > > +static const size_t n_denormal_cases = 8;
> > > +static const size_t n_tests = 18;
> > > +
> > > +template <typename Real>
> > > +class FMinFMaxTest
> > > +{
> > > +public:
> > > + typedef typename TestFunction<Real>::MinMaxFunc MinMaxFunc;
> > > +
> > > + static const Real cases[n_cases];
> > > + static const TestFunction<Real> test_minmax_fns[n_tests];
> > > +
> > > +
> > > + FMinFMaxTest() { }
> > > + ~FMinFMaxTest() { }
> > > +
> > > + static cl_mem create_input_buffer(const piglit_cl_program_test_env* env,
> > > + bool denormals);
> > > +
> > > + static bool verify_results(MinMaxFunc func,
> > > + const Real* results,
> > > + bool test_denormals);
> > > + static bool run_minmax_test(const piglit_cl_program_test_env* env,
> > > + const TestFunction<Real>* test_fn,
> > > + cl_mem input,
> > > + bool test_denormals);
> > > +
> > > + static piglit_result run_tests(const piglit_cl_program_test_env* env,
> > > + cl_device_fp_config fp_config);
> > > +
> > > +};
> > > +
> > > +template <typename Real>
> > > +static Real select_max_gt(Real a, Real b)
> > > +{
> > > + return (a > b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_ge(Real a, Real b)
> > > +{
> > > + return (a >= b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_gt(Real a, Real b)
> > > +{
> > > + return (a > b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_ge(Real a, Real b)
> > > +{
> > > + return (a >= b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_lt(Real a, Real b)
> > > +{
> > > + return (a < b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_le(Real a, Real b)
> > > +{
> > > + return (a <= b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_lt(Real a, Real b)
> > > +{
> > > + return (a < b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_le(Real a, Real b)
> > > +{
> > > + return (a <= b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_ugt(Real a, Real b)
> > > +{
> > > + return !(a <= b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_uge(Real a, Real b)
> > > +{
> > > + return !(a < b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_ugt(Real a, Real b)
> > > +{
> > > + return !(a <= b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_uge(Real a, Real b)
> > > +{
> > > + return !(a < b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_ult(Real a, Real b)
> > > +{
> > > + return !(a >= b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_max_ule(Real a, Real b)
> > > +{
> > > + return !(a > b) ? b : a;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_ult(Real a, Real b)
> > > +{
> > > + return !(a >= b) ? a : b;
> > > +}
> > > +
> > > +template <typename Real>
> > > +static Real select_min_ule(Real a, Real b)
> > > +{
> > > + return !(a > b) ? a : b;
> > > +}
> > > +
> > > +#define TYPE_SUFFIX "_f32"
> > > +#define TYPE_NAME float
> > > +#define TEST_FN(name) { #name TYPE_SUFFIX, name<TYPE_NAME> }
> > > +
> > > +template <>
> > > +const TestFunction<float> FMinFMaxTest<float>::test_minmax_fns[n_tests] = {
> > > + TEST_FN(select_max_gt),
> > > + TEST_FN(select_max_ge),
> > > + TEST_FN(select_min_gt),
> > > + TEST_FN(select_min_ge),
> > > +
> > > + TEST_FN(select_max_lt),
> > > + TEST_FN(select_max_le),
> > > + TEST_FN(select_min_lt),
> > > + TEST_FN(select_min_le),
> > > +
> > > + TEST_FN(select_max_ugt),
> > > + TEST_FN(select_max_uge),
> > > + TEST_FN(select_min_ugt),
> > > + TEST_FN(select_min_uge),
> > > +
> > > + TEST_FN(select_max_ult),
> > > + TEST_FN(select_max_ule),
> > > + TEST_FN(select_min_ult),
> > > + TEST_FN(select_min_ule),
> > > +
> > > + { "test_fmin_f32", fminf },
> > > + { "test_fmax_f32", fmaxf }
> > > +};
> > > +
> > > +#undef TYPE_SUFFIX
> > > +#undef TYPE_NAME
> > > +#define TYPE_SUFFIX "_f64"
> > > +#define TYPE_NAME double
> > > +
> > > +template <>
> > > +const TestFunction<double> FMinFMaxTest<double>::test_minmax_fns[n_tests] = {
> > > + TEST_FN(select_max_gt),
> > > + TEST_FN(select_max_ge),
> > > + TEST_FN(select_min_gt),
> > > + TEST_FN(select_min_ge),
> > > +
> > > + TEST_FN(select_max_lt),
> > > + TEST_FN(select_max_le),
> > > + TEST_FN(select_min_lt),
> > > + TEST_FN(select_min_le),
> > > +
> > > + TEST_FN(select_max_ugt),
> > > + TEST_FN(select_max_uge),
> > > + TEST_FN(select_min_ugt),
> > > + TEST_FN(select_min_uge),
> > > +
> > > + TEST_FN(select_max_ult),
> > > + TEST_FN(select_max_ule),
> > > + TEST_FN(select_min_ult),
> > > + TEST_FN(select_min_ule),
> > > +
> > > + { "test_fmin_f64", fmin },
> > > + { "test_fmax_f64", fmax }
> > > +};
> > > +
> > > +#undef TYPE_SUFFIX
> > > +#undef TYPE_NAME
> > > +
> > > +template <>
> > > +const float FMinFMaxTest<float>::cases[n_cases] = {
> > > + 0.0f,
> > > + -0.0f,
> > > +
> > > + 0.5f,
> > > + -0.5f,
> > > +
> > > + -1.0f,
> > > + 1.0f,
> > > +
> > > + -2.0f,
> > > + 2.0f,
> > > +
> > > + 3.0f,
> > > + -3.0f,
> > > +
> > > + 4.0f,
> > > + -4.0f,
> > > +
> > > + 12345.0,
> > > +
> > > + 0x1p-126f, // Minimum normal number
> > > + -0x1p-126f,
> > > +
> > > + 0x1p-126f, // Min float
> > > + -0x1p-126f,
> > > +
> > > + 0x1.fffffep+127f, // Max float
> > > + -0x1.fffffep+127f,
> > > +
> > > + 0x1p-23f, // Epsilon
> > > + -0x1p-23f,
> > > +
> > > + INFINITY,
> > > + -INFINITY,
> > > + NAN,
> > > +
> > > + // Denormals.
> > > + 0x1p-149f, // Denorm min
> > > + -0x1p-149f,
> > > +
> > > + 0x1p-148f, // Denorm min * 2.0
> > > + -0x1p-148f,
> > > +
> > > + 0x1.fffffcp-127f, // Max denormal
> > > + -0x1.fffffcp-127f,
> > > +
> > > + 0x1.fffffcp-128f, // Max denormal / 2.0
> > > + -0x1.fffffcp-128f
> > > +};
> > > +
> > > +template <>
> > > +const double FMinFMaxTest<double>::cases[n_cases] = {
> > > + 0.0,
> > > + -0.0,
> > > +
> > > + 0.5,
> > > + -0.5,
> > > +
> > > + -1.0,
> > > + 1.0,
> > > +
> > > + -2.0,
> > > + 2.0,
> > > +
> > > + 3.0,
> > > + -3.0,
> > > +
> > > + 4.0,
> > > + -4.0,
> > > +
> > > + 12345.0,
> > > +
> > > + 0x1.fffffffffffffp+1023, // Maximum double
> > > + -0x1.fffffffffffffp+1023,
> > > +
> > > + 0x1p-1022, // Minimum normal number
> > > + -0x1p-1022,
> > > +
> > > + 0x1p-52, // Epsilon
> > > + -0x1p-52,
> > > +
> > > + INFINITY,
> > > + -INFINITY,
> > > + NAN,
> > > +
> > > + // Denormals.
> > > + 0x0.0000000000001p-1022, // Denorm min
> > > + -0x0.0000000000001p-1022,
> > > +
> > > + 0x0.0000000000002p-1022, // Denorm min * 2.0
> > > + -0x0.0000000000002p-1022,
> > > +
> > > + 0x0.fffffffffffffp-1022, // Max denormal
> > > + -0x0.fffffffffffffp-1022,
> > > +
> > > + 0x0.8p-1022, // Max denormal / 2.0
> > > + -0x0.8p-1022
> > > +};
> > > +
> > > +
> > > +
> > > +template <typename Real>
> > > +cl_mem FMinFMaxTest<Real>::create_input_buffer(const piglit_cl_program_test_env* env,
> > > + bool denormals)
> > > +{
> > > + const size_t n = denormals ? n_cases : (n_cases - n_denormal_cases);
> > > + cl_mem buf = piglit_cl_create_buffer(env->context, CL_MEM_READ_ONLY, n * n * sizeof(Real));
> > > + if (!buf)
> > > + return NULL;
> > > +
> > > + for (size_t i = 0; i < n; ++i)
> > > + {
> > > + if (!piglit_cl_write_buffer(env->context->command_queues[0], buf, i * n * sizeof(Real), n * sizeof(Real), cases))
> > > + {
> > > + // Leaking buf
> > > + return NULL;
> > > + }
> > > + }
> > > +
> > > + return buf;
> > > +}
> > > +
> > > +template <typename Real>
> > > +bool FMinFMaxTest<Real>::verify_results(MinMaxFunc func,
> > > + const Real* results,
> > > + bool test_denormals)
> > > +{
> > > + bool failed = false;
> > > +
> > > + const size_t n = test_denormals ? n_cases : (n_cases - n_denormal_cases);
> > > +
> > > + for (size_t i = 0; i < n; ++i)
> > > + {
> > > + for (size_t j = 0; j < n; ++j)
> > > + {
> > > + Real x = cases[i];
> > > + Real y = cases[j];
> > > +
> > > + Real ref = func(x, y);
> > > + Real result = results[n * i + j];
> > > +
> > > + failed |= !piglit_cl_probe_floating(result, ref, 0);
> > > + }
> > > + }
> > > +
> > > + return failed;
> > > +}
> > > +
> > > +template <typename Real>
> > > +bool FMinFMaxTest<Real>::run_minmax_test(const piglit_cl_program_test_env* env,
> > > + const TestFunction<Real>* test_fn,
> > > + cl_mem input,
> > > + bool test_denormals)
> > > +{
> > > + const size_t n = test_denormals ? n_cases : (n_cases - n_denormal_cases);
> > > + const cl_int n_i = (cl_int) n;
> > > + const size_t global_size[2] = { n, n };
> > > + printf("Create kernel '%s'\n", test_fn->kernel_name);
> > > + cl_kernel kernel = piglit_cl_create_kernel(env->program,
> > > + test_fn->kernel_name);
> > > + if (!kernel)
> > > + {
> > > + return true;
> > > + }
> > > +
> > > + Real* ptr_out = new Real[n * n]();
> > > + if (!ptr_out)
> > > + {
> > > + return true;
> > > + }
> > > +
> > > + cl_mem mem_out = piglit_cl_create_buffer(env->context, CL_MEM_WRITE_ONLY,
> > > + n * n * sizeof(Real));
> > > + if (!mem_out)
> > > + {
> > > + delete[] ptr_out;
> > > + return true;
> > > + }
> > > +
> > > + piglit_cl_set_kernel_buffer_arg(kernel, 0, &mem_out);
> > > + piglit_cl_set_kernel_buffer_arg(kernel, 1, &input);
> > > + piglit_cl_set_kernel_arg(kernel, 2, sizeof(cl_int), &n_i);
> > > +
> > > + piglit_cl_execute_ND_range_kernel(env->context->command_queues[0],
> > > + kernel,
> > > + 2,
> > > + NULL,
> > > + global_size,
> > > + NULL);
> > > +
> > > + bool failed = !piglit_cl_read_buffer(env->context->command_queues[0], mem_out, 0,
> > > + n * n * sizeof(Real), ptr_out);
> > > + if (!failed)
> > > + failed = verify_results(test_fn->ref_func, ptr_out, test_denormals);
> > > +
> > > + delete[] ptr_out;
> > > +
> > > + // Leaking mem_out
> > > + return failed;
> > > +}
> > > +
> > > +template <typename Real>
> > > +piglit_result FMinFMaxTest<Real>::run_tests(const piglit_cl_program_test_env* env,
> > > + cl_device_fp_config fp_config)
> > > +{
> > > + bool failed = false;
> > > +
> > > + bool test_denormals = (fp_config & CL_FP_DENORM) != 0;
> > > +
> > > + cl_mem input = create_input_buffer(env, test_denormals);
> > > + if (!input)
> > > + return PIGLIT_FAIL;
> > > +
> > > + for (size_t i = 0; i < n_tests; ++i)
> > > + {
> > > + if (run_minmax_test(env, &test_minmax_fns[i], input, test_denormals))
> > > + {
> > > + fprintf(stderr, "Failure testing kernel %s\n", test_minmax_fns[i].kernel_name);
> > > + return PIGLIT_FAIL;
> > > + }
> > > + }
> > > +
> > > + return failed ? PIGLIT_FAIL : PIGLIT_PASS;
> > > +}
> > > +
> > > +piglit_result
> > > +piglit_cl_test(const int argc,
> > > + const char** argv,
> > > + const piglit_cl_program_test_config* config,
> > > + const piglit_cl_program_test_env* env)
> > > +{
> > > + cl_device_fp_config float_fp_config = 0;
> > > + cl_device_fp_config double_fp_config = 0;
> > > +
> > > + clGetDeviceInfo(env->device_id, CL_DEVICE_SINGLE_FP_CONFIG,
> > > + sizeof(cl_device_fp_config), &float_fp_config, NULL);
> > > +
> > > + clGetDeviceInfo(env->device_id, CL_DEVICE_DOUBLE_FP_CONFIG,
> > > + sizeof(cl_device_fp_config), &double_fp_config, NULL);
> > > +
> > > + piglit_result result = FMinFMaxTest<float>().run_tests(env,
> > > + float_fp_config);
> > > + if (result != PIGLIT_PASS)
> > > + return result;
> > > +
> > > + if (double_fp_config != 0) {
> > > + piglit_result result = FMinFMaxTest<double>().run_tests(env,
> > > + double_fp_config);
> > > + if (result != PIGLIT_PASS)
> > > + return result;
> > > + }
> > > +
> > > + return PIGLIT_PASS;
> > > +}
> >
> > why are these tests created at runtime instead of build time? the c++
> > code would be better as python generator.
> >
> > Jan
>
>
> The number of combinations started growing, and repeating the same
> list of inputs for the tests and manually adding the output results
> grew tiresome. I didn’t see other precedent for other tests to
> generate the test outputs.
that's what generated tests are for.
generated_tests/gen_cl_{,v}store_tests.py use the approach of having
external include files for CL kernels (like
generated_tests/cl/store/store-kernels-global.inc),and python code that
generates test inputs.
it also produces tests that run all the cases irrespective of
successes/failures.
I'm not sure if program-tester supports cl_device_info values. I think
it would be nicer to add cl_device_info support to program tester,
rather than having only this dedicated test. I'm OK with the current
approach if the test is fixed to use subtests and always produces the
same number of results.
regards,
Jan
>
> -Matt
>
-------------- 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/20161206/5d352c8d/attachment-0001.sig>
More information about the Piglit
mailing list