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 import itertools
  22 import os
  23 import random
  24 import textwrap
  25
  26 from modules import utils
  27 from genclbuiltins import MAX_VALUES, DATA_SIZES
  28
  29 TYPES = {
  30     'char': 'uchar',
  31     'uchar': 'uchar',
  32     'short': 'ushort',
  33     'ushort': 'ushort',
  34     'half': 'ushort',
  35     'int': 'uint',
  36     'uint': 'uint',
  37     'float': 'uint',
  38     'long': 'ulong',
  39     'ulong': 'ulong',
  40     'double': 'ulong'
  41 }
  42
  43 VEC_SIZES = ['2', '4', '8', '16']
  44 ELEMENTS = 8
  45
  46 DIR_NAME = os.path.join("cl", "builtin", "misc")
  47
  48
  49 def gen_array(size, m):
  50     return [random.randint(0, m) for i in range(size)]
  51
  52
  53 def permute(data, mask, ssize, dsize):
  54     return [data[(m % ssize) + ((i // dsize) * ssize)]
  55             for i, m in enumerate(mask)]
  56
  57
  58 def ext_req(type_name):
  59     if type_name[:6] == "double":
  60         return "require_device_extensions: cl_khr_fp64"
  61     if type_name[:4] == "half":
  62         return "require_device_extensions: cl_khr_fp16"
  63     return ""
  64
  65
  66 def print_config(f, type_name, utype_name):
  67     f.write(textwrap.dedent(("""\
  68     /*!
  69     [config]
  70     name: shuffle {type_name} {utype_name}
  71     dimensions: 1
  72     """ + ext_req(type_name))
  73     .format(type_name=type_name, utype_name=utype_name)))
  74
  75
  76 def begin_test(type_name, utype_name):
  77     fileName = os.path.join(DIR_NAME, 'builtin-shuffle-{}-{}.cl'.format(type_name, utype_name))
  78     print(fileName)
  79     f = open(fileName, 'w')
  80     print_config(f, type_name, utype_name)
  81     return f
  82
  83
  84 def main():
  85     random.seed(0)
  86     utils.safe_makedirs(DIR_NAME)
  87
  88     for t, ut in TYPES.items():
  89         f = begin_test(t, ut)
  90         for ss, ds in itertools.product(VEC_SIZES, VEC_SIZES):
  91             ssize = int(ss) * ELEMENTS
  92             dsize = int(ds) * ELEMENTS
  93             stype_name = t + ss
  94             dtype_name = t + ds
  95             utype_name = ut + ds
  96             data = gen_array(ssize, MAX_VALUES['ushort'])
  97             mask = gen_array(dsize, MAX_VALUES[ut])
  98             perm = permute(data, mask, int(ss), int(ds))
  99             f.write(textwrap.dedent("""
 100             [test]
 101             name: shuffle {stype_name} {utype_name}
 102             global_size: {elements} 0 0
 103             kernel_name: test_shuffle_{stype_name}_{utype_name}
 104             arg_out: 0 buffer {dtype_name}[{elements}] {perm}
 105             arg_in:  1 buffer {stype_name}[{elements}] {data}
 106             arg_in:  2 buffer {utype_name}[{elements}] {mask}
 107             """.format(stype_name=stype_name, utype_name=utype_name,
 108                        dtype_name=dtype_name, elements=ELEMENTS,
 109                        perm=' '.join([str(x) for x in perm]),
 110                        data=' '.join([str(x) for x in data]),
 111                        mask=' '.join([str(x) for x in mask]))))
 112
 113         f.write(textwrap.dedent("""!*/"""))
 114
 115         if t == "double":
 116             f.write(textwrap.dedent("""
 117             #pragma OPENCL EXTENSION cl_khr_fp64: enable
 118             """))
 119         if t == "half":
 120             f.write(textwrap.dedent("""
 121             #pragma OPENCL EXTENSION cl_khr_fp16: enable
 122             """))
 123
 124         for ss, ds in itertools.product(VEC_SIZES, VEC_SIZES):
 125             type_name = t + ss
 126             utype_name = ut + ds
 127             f.write(textwrap.dedent("""
 128             kernel void test_shuffle_{type_name}{ssize}_{utype_name}{dsize}(global {type_name}* out, global {type_name}* in, global {utype_name}* mask) {{
 129                 vstore{dsize}(shuffle(vload{ssize}(get_global_id(0), in), vload{dsize}(get_global_id(0), mask)), get_global_id(0), out);
 130             }}
 131             """.format(type_name=t, utype_name=ut, ssize=ss, dsize=ds)))
 132
 133         f.close()
 134
 135
 136 if __name__ == '__main__':
 137     main()