generated_tests/gen_cl_shuffle_builtins.py

   1 # coding=utf-8
   2 # Permission is hereby granted, free of charge, to any person obtaining a
   3 # copy of this software and associated documentation files (the "Software"),
   4 # to deal in the Software without restriction, including without limitation
   5 # the rights to use, copy, modify, merge, publish, distribute, sublicense,
   6 # and/or sell copies of the Software, and to permit persons to whom the
   7 # Software is furnished to do so, subject to the following conditions:
   8 #
   9 # The above copyright notice and this permission notice (including the next
  10 # paragraph) shall be included in all copies or substantial portions of the
  11 # Software.
  12 #
  13 # THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  14 # IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  15 # FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
  16 # THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  17 # LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
  18 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
  19 # SOFTWARE.
  20
  21 from __future__ import print_function, division, absolute_import
  22 import itertools
  23 import os
  24 import random
  25 import six
  26 import textwrap
  27
  28 from modules import utils
  29 from genclbuiltins import MAX_VALUES, DATA_SIZES
  30
  31 TYPES = {
  32     'char': 'uchar',
  33     'uchar': 'uchar',
  34     'short': 'ushort',
  35     'ushort': 'ushort',
  36     'half': 'ushort',
  37     'int': 'uint',
  38     'uint': 'uint',
  39     'float': 'uint',
  40     'long': 'ulong',
  41     'ulong': 'ulong',
  42     'double': 'ulong'
  43 }
  44
  45 VEC_SIZES = ['2', '4', '8', '16']
  46 ELEMENTS = 8
  47
  48 DIR_NAME = os.path.join("cl", "builtin", "misc")
  49
  50
  51 def gen_array(size, m):
  52     return [random.randint(0, m) for i in six.moves.range(size)]
  53
  54
  55 def permute(data, mask, ssize, dsize):
  56     return [data[(m % ssize) + ((i // dsize) * ssize)]
  57             for i, m in enumerate(mask)]
  58
  59
  60 def ext_req(type_name):
  61     if type_name[:6] == "double":
  62         return "require_device_extensions: cl_khr_fp64"
  63     if type_name[:4] == "half":
  64         return "require_device_extensions: cl_khr_fp16"
  65     return ""
  66
  67
  68 def print_config(f, type_name, utype_name):
  69     f.write(textwrap.dedent(("""\
  70     /*!
  71     [config]
  72     name: shuffle {type_name} {utype_name}
  73     dimensions: 1
  74     """ + ext_req(type_name))
  75     .format(type_name=type_name, utype_name=utype_name)))
  76
  77
  78 def begin_test(type_name, utype_name):
  79     fileName = os.path.join(DIR_NAME, 'builtin-shuffle-{}-{}.cl'.format(type_name, utype_name))
  80     print(fileName)
  81     f = open(fileName, 'w')
  82     print_config(f, type_name, utype_name)
  83     return f
  84
  85
  86 def main():
  87     random.seed(0)
  88     utils.safe_makedirs(DIR_NAME)
  89
  90     for t, ut in six.iteritems(TYPES):
  91         f = begin_test(t, ut)
  92         for ss, ds in itertools.product(VEC_SIZES, VEC_SIZES):
  93             ssize = int(ss) * ELEMENTS
  94             dsize = int(ds) * ELEMENTS
  95             stype_name = t + ss
  96             dtype_name = t + ds
  97             utype_name = ut + ds
  98             data = gen_array(ssize, MAX_VALUES['ushort'])
  99             mask = gen_array(dsize, MAX_VALUES[ut])
 100             perm = permute(data, mask, int(ss), int(ds))
 101             f.write(textwrap.dedent("""
 102             [test]
 103             name: shuffle {stype_name} {utype_name}
 104             global_size: {elements} 0 0
 105             kernel_name: test_shuffle_{stype_name}_{utype_name}
 106             arg_out: 0 buffer {dtype_name}[{elements}] {perm}
 107             arg_in:  1 buffer {stype_name}[{elements}] {data}
 108             arg_in:  2 buffer {utype_name}[{elements}] {mask}
 109             """.format(stype_name=stype_name, utype_name=utype_name,
 110                        dtype_name=dtype_name, elements=ELEMENTS,
 111                        perm=' '.join([str(x) for x in perm]),
 112                        data=' '.join([str(x) for x in data]),
 113                        mask=' '.join([str(x) for x in mask]))))
 114
 115         f.write(textwrap.dedent("""!*/"""))
 116
 117         if t == "double":
 118             f.write(textwrap.dedent("""
 119             #pragma OPENCL EXTENSION cl_khr_fp64: enable
 120             """))
 121         if t == "half":
 122             f.write(textwrap.dedent("""
 123             #pragma OPENCL EXTENSION cl_khr_fp16: enable
 124             """))
 125
 126         for ss, ds in itertools.product(VEC_SIZES, VEC_SIZES):
 127             type_name = t + ss
 128             utype_name = ut + ds
 129             f.write(textwrap.dedent("""
 130             kernel void test_shuffle_{type_name}{ssize}_{utype_name}{dsize}(global {type_name}* out, global {type_name}* in, global {utype_name}* mask) {{
 131                 vstore{dsize}(shuffle(vload{ssize}(get_global_id(0), in), vload{dsize}(get_global_id(0), mask)), get_global_id(0), out);
 132             }}
 133             """.format(type_name=t, utype_name=ut, ssize=ss, dsize=ds)))
 134
 135         f.close()
 136
 137
 138 if __name__ == '__main__':
 139     main()