[Piglit] [PATCH 1/2] cl: Add generated tests for fract
Jan Vesely
jan.vesely at rutgers.edu
Thu Apr 9 16:24:36 PDT 2015
On Tue, 2015-04-07 at 23:28 +0000, Tom Stellard wrote:
> This uses a custom function to generate the kernel code in order to handle the
> output parameters correctly.
> ---
> generated_tests/gen_cl_math_builtins.py | 75 +++++++++++++++++++++++++++++++++
> generated_tests/genclbuiltins.py | 18 +++++++-
> 2 files changed, 91 insertions(+), 2 deletions(-)
>
> diff --git a/generated_tests/gen_cl_math_builtins.py b/generated_tests/gen_cl_math_builtins.py
> index 30d0c43..34b364d 100644
> --- a/generated_tests/gen_cl_math_builtins.py
> +++ b/generated_tests/gen_cl_math_builtins.py
> @@ -30,6 +30,64 @@ from genclbuiltins import gen, NEGNAN
> from math import acos, acosh, asin, asinh, atan, atan2, atanh, cos, cosh, exp
> from math import fabs, fmod, log10, log1p, pi, pow, sin, sinh, sqrt, tan, tanh
>
> +
> +def gen_fract_kernels(functionDef, f):
> + dataType = 'float'
> + fnName = 'fract'
> + for vec_suffix in ['', '2', '4', '8', '16']:
> + if vec_suffix == '':
> + vec_count = 1
> + else:
> + vec_count = int(vec_suffix)
> + vec_type = '{}{}'.format(dataType, vec_suffix)
> +
> +
> + kernel ="""
> +
> +#if 1 == %(vec_count)s
> + #define STORE_RESULT(ptr, value) ptr[get_global_id(0)] = (value)
> + #define LOAD_ARG(arg) (arg[get_global_id(0)])
> +#else
> + #define STORE_RESULT(ptr, value) vstore%(vec_suffix)s((value), get_global_id(0), ptr)
> + #define LOAD_ARG(arg) (vload%(vec_suffix)s(get_global_id(0), arg))
> +#endif
> +
> +
> +
> +kernel void test_%(vec_count)s_%(fnName)s_%(dataType)s(
> + global %(dataType)s *outp_fract,
> + global %(dataType)s *outl_fract,
> + global %(dataType)s *outg_fract,
> + global %(dataType)s *outp_floor,
> + global %(dataType)s *outl_floor,
> + global %(dataType)s *outg_floor,
> + global %(dataType)s *in0) {
> +
> + %(vec_type)s floor_private;
> + local %(vec_type)s floor_local[1];
> +
> + global %(vec_type)s *outg_floor_vec = (global %(vec_type)s *)(outg_floor);
> +
> + // Test private address space
> + STORE_RESULT(outp_fract, fract(LOAD_ARG(in0), &floor_private));
> + STORE_RESULT(outp_floor, floor_private);
> +
> + // Test local address space
> + STORE_RESULT(outl_fract, fract(LOAD_ARG(in0), floor_local));
> + STORE_RESULT(outl_floor, floor_local[0]);
> +
> + // Test global address space
> + STORE_RESULT(outg_fract, fract(LOAD_ARG(in0), (outg_floor_vec + get_global_id(0))));
> +}
> +
> +#undef STORE_RESULT
> +#undef LOAD_ARG
> +
> +""" % locals()
> +
> + f.write(kernel)
I think it would be nicer to have this as separate file, similar to
store-kernels-*.inc. and have only "external_kernel" property.
especially if we want to add more multi-output functions (frexp, lgamma,
modf, remquo, sincos)
> +
> +
> CLC_VERSION_MIN = {
> 'acos' : 10,
> 'acosh' : 10,
> @@ -51,6 +109,7 @@ CLC_VERSION_MIN = {
> 'fabs' : 10,
> 'floor' : 10,
> 'fmod' : 10,
> + 'fract' : 10,
> 'ldexp' : 10,
> 'log10' : 10,
> 'log1p' : 10,
> @@ -258,6 +317,22 @@ tests = {
> ],
> 'tolerance' : 0
> },
> + 'fract' : {
> + 'arg_types': [F, F, F, F, F, F, F],
> + 'function_type': 'custom',
> + # For fract we have two outputs per address space.
> + 'values': [
> + [float("nan"), 0.0, 0.5, 0.0, float.fromhex('0x1.33333p-2'), float.fromhex('0x1.fffffep-1') ], #fract (private)
> + [float("nan"), 0.0, 0.5, 0.0, float.fromhex('0x1.33333p-2'), float.fromhex('0x1.fffffep-1') ], #fract (local)
> + [float("nan"), 0.0, 0.5, 0.0, float.fromhex('0x1.33333p-2'), float.fromhex('0x1.fffffep-1') ], #fract (global)
> + [float("nan"), float("inf"), 1.0, 2.0, -2.0, -1.0], #floor (private)
> + [float("nan"), float("inf"), 1.0, 2.0, -2.0, -1.0], #floor (local)
> + [float("nan"), float("inf"), 1.0, 2.0, -2.0, -1.0], #floor (global)
> + [float("nan"), float("inf"), 1.5, 2.0,float.fromhex('-0x1.b33334p+0'), float.fromhex('-0x1.000242p-24')] #src0
> + ],
> + 'num_out_args' : 6,
> + 'generate_kernels_fn' : gen_fract_kernels
we also need tolerance (1 for fract). I think we need a separate test
for corner values that have stricter precision requirement.
jan
> + },
> 'ldexp' : {
> 'arg_types': [F, F, I],
> 'function_type': 'tss',
> diff --git a/generated_tests/genclbuiltins.py b/generated_tests/genclbuiltins.py
> index cf95f9c..9a9607d 100644
> --- a/generated_tests/genclbuiltins.py
> +++ b/generated_tests/genclbuiltins.py
> @@ -186,6 +186,11 @@ def gen_kernel_3_arg_mixed_size_tts(f, fnName, inTypes, outType):
>
>
> def generate_kernels(f, dataType, fnName, fnDef):
> + # Functions that need special handling
> + if fnDef['function_type'] is 'custom':
> + fnDef['generate_kernels_fn'](fnDef, f)
> + return
> +
> argTypes = getArgTypes(dataType, fnDef['arg_types'])
>
> # For len(argTypes), remember that this includes the output arg
> @@ -301,6 +306,15 @@ def getArgTypes(baseType, argTypes):
> def isFloatType(t):
> return t not in U
>
> +def getNumOutArgs(functionDef):
> + if 'num_out_args' in functionDef:
> + return functionDef['num_out_args']
> + else:
> + return 1
> +
> +def isOutArg(functionDef, argIdx):
> + return argIdx < getNumOutArgs(functionDef)
> +
> # Print a test with all-vector inputs/outputs and/or mixed vector/scalar args
> def print_test(f, fnName, argType, functionDef, tests, numTests, vecSize, tss):
> # If the test allows mixed vector/scalar arguments, handle the case with
> @@ -334,7 +348,7 @@ def print_test(f, fnName, argType, functionDef, tests, numTests, vecSize, tss):
> for arg in range(0, argCount):
> argInOut = ''
> argVal = getStrVal(argType, tests[arg], (vecSize > 1))
> - if arg == 0:
> + if isOutArg(functionDef, arg):
> argInOut = 'arg_out: '
> else:
> argInOut = 'arg_in: '
> @@ -347,7 +361,7 @@ def print_test(f, fnName, argType, functionDef, tests, numTests, vecSize, tss):
> '[' + str(numTests * vecSize) + '] ' +
> ''.join(map(lambda x: (x + ' ') * vecSize, argVal.split()))
> )
> - if arg == 0:
> + if isOutArg(functionDef, arg) :
> f.write(' tolerance {0} '.format(tolerance))
> # Use ulp tolerance for float types
> if isFloatType(argTypes[arg]):
--
Jan Vesely <jan.vesely at rutgers.edu>
-------------- 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: <http://lists.freedesktop.org/archives/piglit/attachments/20150409/4a6a20e9/attachment.sig>
More information about the Piglit
mailing list