[Piglit] [PATCH 2/2] cl: Add generator for shuffle2 builtins
Jan Vesely
jan.vesely at rutgers.edu
Fri Sep 8 19:50:32 UTC 2017
On Fri, 2017-09-08 at 11:32 -0700, Dylan Baker wrote:
> 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
I never know what the rules are if I take an existing file and modify
it. In this case there isn't much of the original left so I dropped
these (both author and license)
>
> > +#
> > +#
> > +
> > +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
thanks. I learn new things about python every time.
I applied the recommendation in few more locations.
Does the RB extend to shuffle2 as well?
I can repost if you want to see the modified versions.
It's also available here:
https://github.com/jvesely/piglit
thanks,
Jan
>
> > + 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: 833 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/piglit/attachments/20170908/58a99781/attachment.sig>
More information about the Piglit
mailing list