On Mon, 2016-12-05 at 09:48 -0800, arse...@gmail.com wrote: > From: Matt Arsenault <arse...@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
signature.asc
Description: This is a digitally signed message part
_______________________________________________ Piglit mailing list Piglit@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/piglit