[Piglit] [PATCH 1/2] cl: Add generated tests for fract

[Tom Stellard]

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)
+
+
 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
+    },
     '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]):
-- 
2.0.4