[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