Signed-off-by: Jan Vesely <jan.ves...@rutgers.edu> --- generated_tests/CMakeLists.txt | 4 + generated_tests/gen_cl_shuffle_builtins.py | 129 +++++++++++++++++++++++++++++ tests/cl.py | 3 + 3 files changed, 136 insertions(+) create mode 100644 generated_tests/gen_cl_shuffle_builtins.py
diff --git a/generated_tests/CMakeLists.txt b/generated_tests/CMakeLists.txt index fe82ccfa4..251c7e0b8 100644 --- a/generated_tests/CMakeLists.txt +++ b/generated_tests/CMakeLists.txt @@ -225,6 +225,9 @@ piglit_make_generated_tests( piglit_make_generated_tests( builtin_cl_common_tests.list gen_cl_common_builtins.py) +piglit_make_generated_tests( + builtin_cl_shuffle_tests.list + gen_cl_shuffle_builtins.py) # Create a custom target for generating OpenGL tests # This is not added to the default target, instead it is added @@ -272,6 +275,7 @@ add_custom_target(gen-cl-tests builtin_cl_math_tests.list builtin_cl_relational_tests.list builtin_cl_common_tests.list + builtin_cl_shuffle_tests.list cl_store_tests.list cl_vstore_tests.list cl_vload_tests.list diff --git a/generated_tests/gen_cl_shuffle_builtins.py b/generated_tests/gen_cl_shuffle_builtins.py new file mode 100644 index 000000000..98b7773ca --- /dev/null +++ b/generated_tests/gen_cl_shuffle_builtins.py @@ -0,0 +1,129 @@ +# Copyright 2013 Advanced Micro Devices, Inc. +# +# 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.stell...@amd.com> +# +# + +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', 'short', 'ushort', 'int', 'uint', 'long', 'ulong', 'half', 'float', 'double'] +UTYPES = ['uchar', 'ushort', 'uint', '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(data, mask, ssize, dsize): + ret = [] + for i in range(len(mask)): + idx = mask[i] % ssize + ret.append(data[idx + ((i // dsize) * ssize)]) + return ret + + +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: shuffle {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-shuffle-{}-{}.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 in TYPES: + for ut in UTYPES: + f = begin_test(t, ut) + for ss in VEC_SIZES: + for ds in VEC_SIZES: + ssize = int(ss) * ELEMENTS + dsize = int(ds) * ELEMENTS + stype_name = t + ss + dtype_name = t + ds + utype_name = ut + ds + data = gen_array(ssize, MAX_VALUES['ushort']) + mask = gen_array(dsize, MAX_VALUES[ut]) + perm = permute(data, mask, int(ss), int(ds)) + f.write(textwrap.dedent(""" + [test] + name: shuffle {stype_name} {utype_name} + global_size: {elements} 0 0 + kernel_name: test_shuffle_{stype_name}_{utype_name} + arg_out: 0 buffer {dtype_name}[{elements}] {perm} + arg_in: 1 buffer {stype_name}[{elements}] {data} + arg_in: 2 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]), + data=' '.join([str(x) for x in data]), + mask=' '.join([str(x) for x in mask])))) + + f.write(textwrap.dedent("""!*/""")) + 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_shuffle_{type_name}{ssize}_{utype_name}{dsize}(global {type_name}* out, global {type_name}* in, global {utype_name}* mask) {{ + vstore{dsize}(shuffle(vload{ssize}(get_global_id(0), in), 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() diff --git a/tests/cl.py b/tests/cl.py index ffaefb574..24febe532 100644 --- a/tests/cl.py +++ b/tests/cl.py @@ -139,6 +139,9 @@ add_program_test_dir(grouptools.join('program', 'execute', 'builtin'), add_program_test_dir(grouptools.join('program', 'execute', 'builtin'), os.path.join(GENERATED_TESTS_DIR, 'cl', 'builtin', 'common')) +add_program_test_dir(grouptools.join('program', 'execute', 'builtin'), + os.path.join(GENERATED_TESTS_DIR, 'cl', 'builtin', + 'misc')) add_program_test_dir(grouptools.join('program', 'execute', 'store'), os.path.join(GENERATED_TESTS_DIR, 'cl', 'store')) add_program_test_dir(grouptools.join('program', 'execute', 'vstore'), -- 2.13.5 _______________________________________________ Piglit mailing list Piglit@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/piglit