[Piglit] [PATCH v2 1/2] cl: Add generator for shuffle builtins
Jan Vesely
jan.vesely at rutgers.edu
Fri Sep 1 19:30:51 UTC 2017
v2: "The size of each element in the mask must match the size of each element in the result."
Signed-off-by: Jan Vesely <jan.vesely at rutgers.edu>
---
generated_tests/CMakeLists.txt | 4 +
generated_tests/gen_cl_shuffle_builtins.py | 149 +++++++++++++++++++++++++++++
tests/cl.py | 3 +
3 files changed, 156 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..4e007b130
--- /dev/null
+++ b/generated_tests/gen_cl_shuffle_builtins.py
@@ -0,0 +1,149 @@
+# 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.stellard at 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, DATA_SIZES
+
+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(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, ut in TYPES.items():
+ 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("""!*/"""))
+
+ 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_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
More information about the Piglit
mailing list