generated_tests/genclbuiltins.py

   1 from __future__ import print_function, division, absolute_import
   2 import os
   3
   4 import six
   5
   6 __all__ = ['gen', 'DATA_SIZES', 'MAX_VALUES', 'MAX', 'MIN', 'BMIN', 'BMAX',
   7            'SMIN', 'SMAX', 'UMIN', 'UMAX', 'TYPE', 'T', 'U', 'B']
   8
   9 _NUMERIC_TYPES = tuple(list(six.integer_types) + [float])
  10
  11
  12 DATA_SIZES = {
  13     'char': 8,
  14     'uchar': 8,
  15     'short': 16,
  16     'ushort': 16,
  17     'int': 32,
  18     'uint': 32,
  19     'long': 64,
  20     'ulong': 64
  21 }
  22
  23 # By default, just test what is part of the CL1.1 spec, leave vec3 for later
  24 # VEC_WIDTHS = [2, 3, 4, 8, 16]
  25 VEC_WIDTHS = [2, 4, 8, 16]
  26
  27 # ALL_WIDTHS = [1, 2, 3, 4, 8, 16]
  28 ALL_WIDTHS = [1, 2, 4, 8, 16]
  29
  30 MIN_VALUES = {
  31     'char': -128,
  32     'uchar': 0,
  33     'short': -32768,
  34     'ushort': 0,
  35     'int': -2147483648,
  36     'uint': 0,
  37     'long': -9223372036854775808,
  38     'ulong': 0
  39 }
  40
  41 MAX_VALUES = {
  42     'char': 127,
  43     'uchar': 255,
  44     'short': 32767,
  45     'ushort': 65535,
  46     'int': 2147483647,
  47     'uint': 4294967295,
  48     'long': 9223372036854775807,
  49     'ulong': 18446744073709551615
  50 }
  51
  52 # Define placeholders to reduce magic number usage
  53 MAX = 'MAX_VAL'
  54 MIN = 'MIN_VAL'
  55 BMIN = 'min_for_larger_type'
  56 BMAX = 'max_for_larger_type'
  57 SMIN = 'signed_min_for_type'
  58 SMAX = 'signed_max_for_type'
  59 UMIN = 'unsigned_min_for_type'
  60 UMAX = 'unsigned_max_for_type'
  61 TYPE = 'TYPE'
  62 SIZE = 'SIZE'
  63 TRUE = 'true_value_for_type' #1 for scalar, -1 for vector
  64 NEGNAN = 'Negative NAN as a string, because float("-nan") just produces nan'
  65
  66 # Identity type list
  67 T = {
  68     'char': 'char',
  69     'uchar': 'uchar',
  70     'short': 'short',
  71     'ushort': 'ushort',
  72     'int': 'int',
  73     'uint': 'uint',
  74     'long': 'long',
  75     'ulong': 'ulong'
  76 }
  77 # Signed type for each type
  78 SIGNED = {
  79     'char': 'char',
  80     'uchar': 'char',
  81     'short': 'short',
  82     'ushort': 'short',
  83     'int': 'int',
  84     'uint': 'int',
  85     'long': 'long',
  86     'ulong': 'long'
  87 }
  88 # Unsigned type for each source type
  89 U = {
  90     'char': 'uchar',
  91     'uchar': 'uchar',
  92     'short': 'ushort',
  93     'ushort': 'ushort',
  94     'int': 'uint',
  95     'uint': 'uint',
  96     'long': 'ulong',
  97     'ulong': 'ulong'
  98 }
  99 # Next larger type with same signedness
 100 B = {
 101     'char': 'short',
 102     'uchar': 'ushort',
 103     'short': 'int',
 104     'ushort': 'uint',
 105     'int': 'long',
 106     'uint': 'ulong',
 107 }
 108
 109
 110 # vecSizes has the layout [in0width, ..., inNwidth] where outType width is
 111 # assumed to match the width of the first input
 112 def gen_kernel(f, fnName, inTypes, outType, vecSizes, typePrefix):
 113     f.write('kernel void test_' + typePrefix + str(vecSizes[0]) + '_' + fnName
 114             + '_' + inTypes[0]+'(global '+outType+'* out')
 115     for arg in range(0, len(inTypes)):
 116         f.write(', global '+inTypes[arg]+'* in'+str(arg))
 117     f.write('){\n')
 118
 119     suffix = ';'
 120     if (vecSizes[0] == 1):
 121         f.write('  out[get_global_id(0)] = ')
 122     else:
 123         f.write('  vstore'+str(vecSizes[0])+'(')
 124         suffix = ', get_global_id(0), out)' + suffix
 125
 126     f.write(fnName+'(')
 127     suffix = ')' + suffix
 128
 129     for arg in range(0, len(inTypes)):
 130         if (arg > 0):
 131             f.write(', ')
 132         # if scalar, don't print vload/vstore
 133         if (vecSizes[arg] == 1):
 134             f.write('in'+str(arg)+'[get_global_id(0)]')
 135         else:
 136             f.write('vload'+str(vecSizes[arg])+'(get_global_id(0), in'+str(arg)+')')
 137
 138     f.write(suffix+'\n}\n\n')
 139
 140
 141
 142 def gen_kernel_1_arg(f, fnName, inType, outType):
 143     for vecSize in ALL_WIDTHS:
 144         gen_kernel(f, fnName, [inType], outType, [vecSize], '')
 145
 146
 147 #  2 argument kernel with input types that match their vector size
 148 def gen_kernel_2_arg_same_size(f, fnName, inTypes, outType):
 149     for vecSize in ALL_WIDTHS:
 150         gen_kernel(f, fnName, inTypes, outType, [vecSize, vecSize],
 151                    '')
 152
 153
 154 #  2 argument kernel with 1 vector and one scalar input argument
 155 def gen_kernel_2_arg_mixed_size(f, fnName, inTypes, outType):
 156     for vecSize in VEC_WIDTHS:
 157         gen_kernel(f, fnName, inTypes, outType, [vecSize, 1], 'tss_')
 158
 159
 160 #  2 argument kernel with 1 vector and one scalar input argument with multiple
 161 #    input data types
 162 def gen_kernel_2_arg_mixed_sign(f, fnName, inTypes, outType):
 163     for vecSize in ALL_WIDTHS:
 164         gen_kernel(f, fnName, inTypes, outType, [vecSize, vecSize],
 165                    '')
 166
 167
 168 #  3-argument built-in functions
 169
 170
 171 def gen_kernel_3_arg_same_type(f, fnName, inTypes, outType):
 172     for vecSize in ALL_WIDTHS:
 173         gen_kernel(f, fnName, inTypes, outType,
 174                    [vecSize, vecSize, vecSize], ''
 175         )
 176
 177 def gen_kernel_3_arg_mixed_size_tss(f, fnName, inTypes, outType):
 178     for vecSize in VEC_WIDTHS:
 179         gen_kernel(f, fnName, inTypes, outType,
 180                    [vecSize, 1, 1], 'tss_')
 181
 182 def gen_kernel_3_arg_mixed_size_tts(f, fnName, inTypes, outType):
 183     for vecSize in VEC_WIDTHS:
 184         gen_kernel(f, fnName, inTypes, outType,
 185                    [vecSize, vecSize, 1], 'tts_')
 186
 187
 188 def generate_kernels(f, dataType, fnName, fnDef):
 189     argTypes = getArgTypes(dataType, fnDef['arg_types'])
 190
 191     # For len(argTypes), remember that this includes the output arg
 192     if (len(argTypes) == 2):
 193         gen_kernel_1_arg(f, fnName, argTypes[1], argTypes[0])
 194         return
 195
 196     if (len(argTypes) == 3 and not fnName is 'upsample'):
 197         gen_kernel_2_arg_same_size(f, fnName,
 198                                 [argTypes[1], argTypes[2]], argTypes[0])
 199         if (fnDef['function_type'] is 'tss'):
 200             gen_kernel_2_arg_mixed_size(f, fnName,
 201                                 [argTypes[1], argTypes[2]], argTypes[0])
 202         return
 203
 204     if (len(argTypes) == 4):
 205         gen_kernel_3_arg_same_type(f, fnName,
 206                    [argTypes[1], argTypes[2], argTypes[3]], argTypes[0])
 207         if (fnDef['function_type'] is 'tss'):
 208             gen_kernel_3_arg_mixed_size_tss(f, fnName,
 209                    [argTypes[1], argTypes[2], argTypes[3]], argTypes[0])
 210         if (fnDef['function_type'] is 'tts'):
 211             gen_kernel_3_arg_mixed_size_tts(f, fnName,
 212                    [argTypes[1], argTypes[2], argTypes[3]], argTypes[0])
 213         return
 214
 215     if (fnName is 'upsample'):
 216         gen_kernel_2_arg_mixed_sign(f, fnName,
 217                                     [argTypes[1], argTypes[2]],
 218                                     argTypes[0])
 219         return
 220
 221 def getValue(type, val, isVector):
 222     # Check if val is a str, list, or value
 223     if (isinstance(val, str)):
 224         if (val == MIN):
 225             return MIN_VALUES[type]
 226         elif (val == MAX):
 227             return MAX_VALUES[type]
 228         elif (val == BMIN):
 229             return MIN_VALUES[B[type]]
 230         elif (val == BMAX):
 231             return MAX_VALUES[B[type]]
 232         elif (val == SMIN):
 233             return MIN_VALUES[SIGNED[type]]
 234         elif (val == SMAX):
 235             return MAX_VALUES[SIGNED[type]]
 236         elif (val == UMIN):
 237             return MIN_VALUES[U[type]]
 238         elif (val == UMAX):
 239             return MAX_VALUES[U[type]]
 240         elif (val == TYPE):
 241             return type
 242         elif (val == SIZE):
 243             return DATA_SIZES[type]
 244         elif (val == TRUE):
 245             if (isVector):
 246                 return -1
 247             else:
 248                 return 1
 249         elif (val == NEGNAN):
 250             return '-nan' #cl-program-tester translates this for us
 251         else:
 252             print('Unknown string value: ' + val + '\n')
 253     elif (isinstance(val, list)):
 254         # The list should be of the format: [op, arg1, ... argN] where op is a
 255         # Fn ref and arg[1-n] are either MIN/MAX or numbers (They could be
 256         # nested lists). The exception for arg1 is TYPE, which means to
 257         # substitute the data type
 258
 259         # Evaluate the value of the requested function and arguments
 260         # TODO: Change to varargs calls after unshifting the first list element
 261         if (callable(val[0])):
 262             if (len(val) == 2):
 263                 return (val[0])(getValue(type, val[1], isVector))
 264             elif (len(val) == 3):
 265                 return (val[0])(getValue(type, val[1], isVector), getValue(type, val[2], isVector))
 266             elif (len(val) == 4):
 267                 return (val[0])(getValue(type, val[1], isVector), getValue(type, val[2], isVector),
 268                             getValue(type, val[3], isVector))
 269             else:
 270                 return (val[0])(getValue(type, val[1], isVector), getValue(type, val[2], isVector),
 271                             getValue(type, val[3], isVector), getValue(type, val[4], isVector))
 272         else:
 273              return map(lambda x: getValue(type, x, isVector), val);
 274
 275     # At this point, we should have been passed a number
 276     if (isinstance(val, _NUMERIC_TYPES)):
 277         return val
 278
 279     print('Invalid value '+repr(val)+' encountered in getValue\n')
 280
 281
 282
 283 def getStrVal(type, val, isVector):
 284     return " ".join(map(str, getValue(type, val, isVector)))
 285
 286
 287 def getArgType(baseType, argType):
 288     # If the argType is a string, it's a literal data type... return it
 289     if (isinstance(argType, str)):
 290         return argType
 291     # otherwise it's a list to pick from
 292     return argType[baseType]
 293
 294
 295 def getArgTypes(baseType, argTypes):
 296     ret = []
 297     for argType in argTypes:
 298         ret.append(getArgType(baseType, argType))
 299     return ret
 300
 301 def isFloatType(t):
 302     return t not in U
 303
 304 # Print a test with all-vector inputs/outputs and/or mixed vector/scalar args
 305 def print_test(f, fnName, argType, functionDef, tests, numTests, vecSize, fntype):
 306     # If the test allows mixed vector/scalar arguments, handle the case with
 307     # only vector arguments through a recursive call.
 308     if (fntype is 'tss' or fntype is 'tts'):
 309         print_test(f, fnName, argType, functionDef, tests, numTests, vecSize,
 310                    'ttt')
 311
 312     # The tss && vecSize==1 case is handled in the non-tss case.
 313     if ((not fntype is 'ttt') and vecSize == 1):
 314         return
 315
 316     # If we're handling mixed vector/scalar input widths, the kernels have
 317     # different names than when the vector widths match
 318     tssStr = fntype + '_' if (not fntype is 'ttt') else ''
 319
 320     argTypes = getArgTypes(argType, functionDef['arg_types'])
 321     argCount = len(argTypes)
 322     tolerance = functionDef['tolerance'] if 'tolerance' in functionDef else 0
 323
 324     # Write the test header
 325     f.write('[test]\n' + 'name: ' + tssStr + fnName + ' ' + argType
 326             + str(vecSize) + '\n' + 'kernel_name: test_' + tssStr + str(vecSize)
 327             + '_' + fnName + '_' + argType + '\n' + 'global_size: '
 328             + str(numTests) + ' 0 0\n\n'
 329     )
 330
 331     # For each argument, write a line containing its type, index, and values
 332     for arg in range(0, argCount):
 333         argInOut = ''
 334         argVal = getStrVal(argType, tests[arg], (vecSize > 1))
 335         if arg == 0:
 336             argInOut = 'arg_out: '
 337         else:
 338             argInOut = 'arg_in: '
 339
 340         # The output argument and first tss argument are vectors, any that
 341         # follow are scalar. If !tss, then everything has a matching vector
 342         # width
 343         if (fntype is 'ttt' or (arg < 2 and fntype is 'tss') or (arg < 3 and fntype is 'tts')):
 344             f.write(argInOut + str(arg) + ' buffer ' + argTypes[arg] +
 345                     '[' + str(numTests * vecSize) + '] ' +
 346                     ''.join(map(lambda x: (x + ' ') * vecSize, argVal.split()))
 347             )
 348             if arg == 0:
 349                 f.write(' tolerance {0} '.format(tolerance))
 350                 # Use ulp tolerance for float types
 351                 if isFloatType(argTypes[arg]):
 352                     f.write('ulp')
 353             f.write('\n')
 354         else:
 355             argInOut = 'arg_in: '
 356             f.write(argInOut + str(arg) + ' buffer ' + argTypes[arg] + '[' +
 357                     str(numTests) + '] ' + argVal + '\n'
 358             )
 359
 360     # Blank line between tests for formatting reasons
 361     f.write('\n')
 362
 363
 364 def gen(types, minVersions, functions, testDefs, dirName):
 365     # Create the output directory if required
 366     if not os.path.exists(dirName):
 367         try:
 368             os.makedirs(dirName)
 369         except OSError as e:
 370             if e.errno == 17:  # file exists
 371                 pass
 372             raise
 373
 374     # Loop over all data types being tested. Create one output file per data
 375     # type
 376     for dataType in types:
 377         for fnName in functions:
 378             # Merge all of the generic/signed/unsigned/custom test definitions
 379             if (dataType, fnName) not in testDefs:
 380                 continue
 381             functionDef = testDefs[(dataType, fnName)]
 382
 383             # Check if the function actually exists for this data type
 384             if (not functionDef.keys()):
 385                 continue
 386
 387             clcVersionMin = minVersions[fnName]
 388
 389             fileName = 'builtin-' + dataType + '-' + fnName + '-' + \
 390                 str(float(clcVersionMin)/10)+'.generated.cl'
 391
 392             fileName = os.path.join(dirName, fileName)
 393
 394             f = open(fileName, 'w')
 395             print(fileName)
 396             # Write the file header
 397             f.write('/*!\n' +
 398                     '[config]\n' +
 399                     'name: Test '+dataType+' '+fnName+' built-in on CL 1.1\n' +
 400                     'clc_version_min: '+str(clcVersionMin)+'\n' +
 401                     'dimensions: 1\n'
 402             )
 403             if (dataType == 'double'):
 404                 f.write('require_device_extensions: cl_khr_fp64\n')
 405
 406             # Blank line  to provide separation between config header and tests
 407             f.write('\n')
 408
 409             # Write all tests for the built-in function
 410             tests = functionDef['values']
 411             argCount = len(functionDef['arg_types'])
 412             fnType = functionDef['function_type']
 413
 414             outputValues = tests[0]
 415             numTests = len(outputValues)
 416
 417             # Handle all available scalar/vector widths
 418             sizes = sorted(VEC_WIDTHS)
 419             sizes.insert(0, 1)  # Add 1-wide scalar to the vector widths
 420             for vecSize in sizes:
 421                 print_test(f, fnName, dataType, functionDef, tests,
 422                            numTests, vecSize, fnType)
 423
 424             # Terminate the header section
 425             f.write('!*/\n\n')
 426
 427             if (dataType == 'double'):
 428                 f.write('#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n\n')
 429
 430             # Generate the actual kernels
 431             generate_kernels(f, dataType, fnName, functionDef)
 432
 433         f.close()