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

Submitted by Tom Stellard on April 7, 2015, 11:28 p.m.

Details

Message ID 1428449339-8235-1-git-send-email-thomas.stellard@amd.com
State New
Headers show

Not browsing as part of any series.

Commit Message

Tom Stellard April 7, 2015, 11:28 p.m.
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(-)

Patch hide | download patch | download mbox

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]):

Comments

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]):