generated_tests/genclbuiltins.py

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