generated_tests/gen_cl_shuffle2_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
  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(data1, data2, mask, ssize, dsize):
  54     ret = []
  55     for i, m in enumerate(mask):
  56         src = data1 if (m % (2 * ssize)) < ssize else data2
  57         ret.append(src[(m % ssize) + ((i // dsize) * ssize)])
  58     return ret
  59
  60
  61 def ext_req(type_name):
  62     if type_name[:6] == "double":
  63         return "require_device_extensions: cl_khr_fp64"
  64     if type_name[:4] == "half":
  65         return "require_device_extensions: cl_khr_fp16"
  66     return ""
  67
  68
  69 def print_config(f, type_name, utype_name):
  70     f.write(textwrap.dedent(("""\
  71     /*!
  72     [config]
  73     name: shuffle2 {type_name} {utype_name}
  74     dimensions: 1
  75     """ + ext_req(type_name))
  76     .format(type_name=type_name, utype_name=utype_name)))
  77
  78
  79 def begin_test(type_name, utype_name):
  80     fileName = os.path.join(DIR_NAME, 'builtin-shuffle2-{}-{}.cl'.format(type_name, utype_name))
  81     print(fileName)
  82     f = open(fileName, 'w')
  83     print_config(f, type_name, utype_name)
  84     return f
  85
  86
  87 def main():
  88     random.seed(0)
  89     utils.safe_makedirs(DIR_NAME)
  90
  91     for t, ut in TYPES.items():
  92         f = begin_test(t, ut)
  93         for ss, ds in itertools.product(VEC_SIZES, VEC_SIZES):
  94             ssize = int(ss) * ELEMENTS
  95             dsize = int(ds) * ELEMENTS
  96             stype_name = t + ss
  97             dtype_name = t + ds
  98             utype_name = ut + ds
  99             data1 = gen_array(ssize, MAX_VALUES['ushort'])
 100             data2 = gen_array(ssize, MAX_VALUES['ushort'])
 101             mask = gen_array(dsize, MAX_VALUES[ut])
 102             perm = permute(data1, data2, mask, int(ss), int(ds))
 103             f.write(textwrap.dedent("""
 104             [test]
 105             name: shuffle2 {stype_name} {utype_name}
 106             global_size: {elements} 0 0
 107             kernel_name: test_shuffle2_{stype_name}_{utype_name}
 108             arg_out: 0 buffer {dtype_name}[{elements}] {perm}
 109             arg_in:  1 buffer {stype_name}[{elements}] {data1}
 110             arg_in:  2 buffer {stype_name}[{elements}] {data2}
 111             arg_in:  3 buffer {utype_name}[{elements}] {mask}
 112             """.format(stype_name=stype_name, utype_name=utype_name,
 113                        dtype_name=dtype_name, elements=ELEMENTS,
 114                        perm=' '.join([str(x) for x in perm]),
 115                        data1=' '.join([str(x) for x in data1]),
 116                        data2=' '.join([str(x) for x in data2]),
 117                        mask=' '.join([str(x) for x in mask]))))
 118
 119         f.write(textwrap.dedent("""!*/"""))
 120
 121         if t == "double":
 122             f.write(textwrap.dedent("""
 123             #pragma OPENCL EXTENSION cl_khr_fp64: enable
 124             """))
 125         if t == "half":
 126             f.write(textwrap.dedent("""
 127             #pragma OPENCL EXTENSION cl_khr_fp16: enable
 128             """))
 129
 130         for ss, ds in itertools.product(VEC_SIZES, VEC_SIZES):
 131             type_name = t + ss
 132             utype_name = ut + ds
 133             f.write(textwrap.dedent("""
 134             kernel void test_shuffle2_{type_name}{ssize}_{utype_name}{dsize}(global {type_name}* out, global {type_name}* in1, global {type_name}* in2, global {utype_name}* mask) {{
 135                 vstore{dsize}(shuffle2(vload{ssize}(get_global_id(0), in1), vload{ssize}(get_global_id(0), in2), vload{dsize}(get_global_id(0), mask)), get_global_id(0), out);
 136             }}
 137             """.format(type_name=t, utype_name=ut, ssize=ss, dsize=ds)))
 138
 139         f.close()
 140
 141
 142 if __name__ == '__main__':
 143     main()