[Piglit] [PATCH 2/2] cl: Add generator for shuffle2 builtins
Dylan Baker
dylan at pnwbakers.com
Fri Sep 8 18:32:31 UTC 2017
Quoting Jan Vesely (2017-09-01 12:30:52)
[snip]
> diff --git a/generated_tests/gen_cl_shuffle2_builtins.py b/generated_tests/gen_cl_shuffle2_builtins.py
> new file mode 100644
> index 000000000..677917210
> --- /dev/null
> +++ b/generated_tests/gen_cl_shuffle2_builtins.py
> @@ -0,0 +1,154 @@
> +# Copyright 2013 Advanced Micro Devices, Inc.
This also seems copy-n-pasted, probably on the other patch as well.
> +#
> +# Permission is hereby granted, free of charge, to any person obtaining a
> +# copy of this software and associated documentation files (the "Software"),
> +# to deal in the Software without restriction, including without limitation
> +# the rights to use, copy, modify, merge, publish, distribute, sublicense,
> +# and/or sell copies of the Software, and to permit persons to whom the
> +# Software is furnished to do so, subject to the following conditions:
> +#
> +# The above copyright notice and this permission notice (including the next
> +# paragraph) shall be included in all copies or substantial portions of the
> +# Software.
> +#
> +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
> +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
> +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
> +# THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
> +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
> +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
> +# SOFTWARE.
> +#
> +# Authors: Tom Stellard <thomas.stellard at amd.com>
same comment on this one as the other patch
> +#
> +#
> +
> +from __future__ import print_function, division, absolute_import
> +import os
> +import random
> +import textwrap
> +
> +from six.moves import range
> +
> +from modules import utils
> +from genclbuiltins import MAX_VALUES
> +
> +TYPES = {
> + 'char': 'uchar',
> + 'uchar': 'uchar',
> + 'short': 'ushort',
> + 'ushort': 'ushort',
> + 'half': 'ushort',
> + 'int': 'uint',
> + 'uint': 'uint',
> + 'float': 'uint',
> + 'long': 'ulong',
> + 'ulong': 'ulong',
> + 'double': 'ulong'
> +}
> +
> +VEC_SIZES = ['2', '4', '8', '16']
> +ELEMENTS = 8
> +
> +DIR_NAME = os.path.join("cl", "builtin", "misc")
> +
> +
> +def gen_array(size, m):
> + return [random.randint(0, m) for i in range(size)]
> +
> +
> +def permute(data1, data2, mask, ssize, dsize):
> + ret = []
> + for i in range(len(mask)):
> + idx = mask[i] % (2 * ssize)
> + src = data1 if idx < ssize else data2
> + idx = mask[i] % ssize
> + ret.append(src[idx + ((i // dsize) * ssize)])
> + return ret
please use use enumerate here, I don't know if the list comprehension would be
very readable here, so just enumerate would be good.
> +
> +
> +def ext_req(type_name):
> + if type_name[:6] == "double":
> + return "require_device_extensions: cl_khr_fp64"
> + if type_name[:4] == "half":
> + return "require_device_extensions: cl_khr_fp16"
> + return ""
> +
> +
> +def print_config(f, type_name, utype_name):
> + f.write(textwrap.dedent(("""\
> + /*!
> + [config]
> + name: shuffle2 {type_name} {utype_name}
> + dimensions: 1
> + """ + ext_req(type_name))
> + .format(type_name=type_name, utype_name=utype_name)))
> +
> +
> +def begin_test(type_name, utype_name):
> + fileName = os.path.join(DIR_NAME, 'builtin-shuffle2-{}-{}.cl'.format(type_name, utype_name))
> + print(fileName)
> + f = open(fileName, 'w')
> + print_config(f, type_name, utype_name)
> + return f
> +
> +
> +def main():
> + random.seed(0)
> + utils.safe_makedirs(DIR_NAME)
> +
> + for t, ut in TYPES.items():
consider:
for t, ut in six.iteritems(TYPES), which will ensure consitent behavior between
python2 and python3.
> + f = begin_test(t, ut)
> + for ss in VEC_SIZES:
> + for ds in VEC_SIZES:
consider itertools.product here as well
> + ssize = int(ss) * ELEMENTS
> + dsize = int(ds) * ELEMENTS
> + stype_name = t + ss
> + dtype_name = t + ds
> + utype_name = ut + ds
> + data1 = gen_array(ssize, MAX_VALUES['ushort'])
> + data2 = gen_array(ssize, MAX_VALUES['ushort'])
> + mask = gen_array(dsize, MAX_VALUES[ut])
> + perm = permute(data1, data2, mask, int(ss), int(ds))
> + f.write(textwrap.dedent("""
> + [test]
> + name: shuffle2 {stype_name} {utype_name}
> + global_size: {elements} 0 0
> + kernel_name: test_shuffle2_{stype_name}_{utype_name}
> + arg_out: 0 buffer {dtype_name}[{elements}] {perm}
> + arg_in: 1 buffer {stype_name}[{elements}] {data1}
> + arg_in: 2 buffer {stype_name}[{elements}] {data2}
> + arg_in: 3 buffer {utype_name}[{elements}] {mask}
> + """.format(stype_name=stype_name, utype_name=utype_name,
> + dtype_name=dtype_name, elements=ELEMENTS,
> + perm=' '.join([str(x) for x in perm]),
> + data1=' '.join([str(x) for x in data1]),
> + data2=' '.join([str(x) for x in data2]),
> + mask=' '.join([str(x) for x in mask]))))
> +
> + f.write(textwrap.dedent("""!*/"""))
> +
> + if t == "double":
> + f.write(textwrap.dedent("""
> + #pragma OPENCL EXTENSION cl_khr_fp64: enable
> + """))
> + if t == "half":
> + f.write(textwrap.dedent("""
> + #pragma OPENCL EXTENSION cl_khr_fp16: enable
> + """))
> +
> + for ss in VEC_SIZES:
> + for ds in VEC_SIZES:
> + type_name = t + ss
> + utype_name = ut + ds
> + f.write(textwrap.dedent("""
> + 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) {{
> + 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);
> + }}
> + """.format(type_name=t, utype_name=ut, ssize=ss, dsize=ds)))
> +
> + f.close()
> +
> +
> +if __name__ == '__main__':
> + main()
> --
> 2.13.5
>
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/piglit
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 488 bytes
Desc: signature
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20170908/4794467c/attachment.sig>
More information about the Piglit
mailing list