[Piglit] [PATCH 2/2] cl: Add generated tests for global and local stores

Tom Stellard tom at stellard.net
Wed Jul 31 19:16:54 PDT 2013


From: Tom Stellard <thomas.stellard at amd.com>

---
 generated_tests/CMakeLists.txt                    |  4 +
 generated_tests/cl/store/store-kernels-global.inc | 21 ++++++
 generated_tests/cl/store/store-kernels-local.inc  | 14 ++++
 generated_tests/generate-cl-store-tests.py        | 90 +++++++++++++++++++++++
 tests/all_cl.tests                                |  3 +
 5 files changed, 132 insertions(+)
 create mode 100644 generated_tests/cl/store/store-kernels-global.inc
 create mode 100644 generated_tests/cl/store/store-kernels-local.inc
 create mode 100644 generated_tests/generate-cl-store-tests.py

diff --git a/generated_tests/CMakeLists.txt b/generated_tests/CMakeLists.txt
index db3734f..6ac5a9c 100644
--- a/generated_tests/CMakeLists.txt
+++ b/generated_tests/CMakeLists.txt
@@ -54,6 +54,9 @@ piglit_make_generated_tests(
 piglit_make_generated_tests(
 	builtin_cl_int_tests.list
 	generate-cl-int-builtins.py)
+piglit_make_generated_tests(
+	cl_store_tests.list
+	generate-cl-store-tests.py)
 
 # Add a "gen-tests" target that can be used to generate all the
 # tests without doing any other compilation.
@@ -62,6 +65,7 @@ add_custom_target(gen-tests ALL
 		builtin_uniform_tests.list
 		constant_array_size_tests.list
 		builtin_cl_int_tests.list
+		cl_store_tests.list
 		interpolation_tests.list
 		non-lvalue_tests.list
 		texture_query_lod_tests.list
diff --git a/generated_tests/cl/store/store-kernels-global.inc b/generated_tests/cl/store/store-kernels-global.inc
new file mode 100644
index 0000000..b6220d0
--- /dev/null
+++ b/generated_tests/cl/store/store-kernels-global.inc
@@ -0,0 +1,21 @@
+typedef TYPE type_t;
+
+#if TYPE == double
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#endif
+
+kernel void store_global(global type_t *out, global type_t *in) {
+	out[0] = in[0];
+	out[1] = in[1];
+	out[2] = in[2];
+	out[3] = in[3];
+	out[4] = in[4];
+	out[5] = in[5];
+	out[6] = in[6];
+	out[7] = in[7];
+}
+
+kernel void store_global_wi(global type_t *out, global type_t *in) {
+	size_t id = get_global_id(0);
+	out[id] = in[id];
+}
diff --git a/generated_tests/cl/store/store-kernels-local.inc b/generated_tests/cl/store/store-kernels-local.inc
new file mode 100644
index 0000000..7d70d13
--- /dev/null
+++ b/generated_tests/cl/store/store-kernels-local.inc
@@ -0,0 +1,14 @@
+typedef TYPE type_t;
+
+#if TYPE == double
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#endif
+
+kernel void store_local(global type_t *out, global type_t *in) {
+	local type_t local_data[8];
+	size_t id = get_local_id(0);
+	size_t store_index = (id + 1) % 8;
+	local_data[store_index] = store_index;
+	barrier(CLK_LOCAL_MEM_FENCE);
+	out[id] = local_data[id];
+}
diff --git a/generated_tests/generate-cl-store-tests.py b/generated_tests/generate-cl-store-tests.py
new file mode 100644
index 0000000..488aa62
--- /dev/null
+++ b/generated_tests/generate-cl-store-tests.py
@@ -0,0 +1,90 @@
+#!/usr/bin/env python
+#
+# 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>
+#
+#
+
+import os
+
+TYPES = ['char', 'uchar', 'short', 'ushort', 'int', 'uint', 'long', 'ulong', 'float', 'double']
+VEC_SIZES = ['', '2', '4', '8', '16']
+
+dirName = os.path.join("cl", "store")
+if not os.path.exists(dirName):
+    os.makedirs(dirName)
+
+
+def gen_array(size):
+    items = size * 8
+    array = list(xrange(items))
+    return ' '.join(map(str,array))
+
+def print_config(f, type_name, addr_space):
+    f.write( \
+        '[config]\n' + \
+        'name: Store (' + type_name + ')\n' + \
+        'program_source_file: store-kernels-' + addr_space + '.inc\n' + \
+        'build_options: -D TYPE=' + type_name + '\n' + \
+        'dimensions: 1\n')
+
+def begin_test(type_name, addr_space):
+        fileName = os.path.join(dirName, 'store-' + type_name + '-' + addr_space + '.program_test')
+        print(fileName)
+        f = open(fileName, 'w')
+        print_config(f, type_name, addr_space)
+        return f
+
+
+for t in TYPES:
+    for s in VEC_SIZES:
+        if s == '':
+            size = 1
+        else:
+            size = int(s)
+        type_name = t + s
+        f = begin_test(type_name, 'global')
+        f.write( \
+            '[test]\n' + \
+            'name: global address space\n' + \
+            'global_size: 1 0 0\n' + \
+            'kernel_name: store_global\n' + \
+            'arg_out: 0 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \
+            'arg_in:  1 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \
+            '[test]\n' + \
+            'name: global address space work items\n' + \
+            'global_size: 8 0 0\n' + \
+            'kernel_name: store_global_wi\n' + \
+            'arg_out: 0 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \
+            'arg_in:  1 buffer ' + type_name + '[8] ' + gen_array(size) + '\n')
+        f.close()
+        f = begin_test(type_name, 'local')
+        f.write( \
+            '[test]\n' + \
+            'name: local address space\n' + \
+            'global_size: 8 0 0\n' + \
+            'local_size:  8 0 0\n' + \
+            'kernel_name: store_local\n' + \
+            'arg_out: 0 buffer ' + type_name + '[8] ' + gen_array(size) + '\n' + \
+            'arg_in:  1 buffer ' + type_name + '[8] ' + gen_array(size) + '\n')
+        f.close()
diff --git a/tests/all_cl.tests b/tests/all_cl.tests
index b722511..efb04e4 100644
--- a/tests/all_cl.tests
+++ b/tests/all_cl.tests
@@ -110,3 +110,6 @@ add_program_test_dir(program_execute, 'tests/cl/program/execute')
 program_execute_builtin = Group()
 program["Execute"]["Builtin"] = program_execute_builtin
 add_program_test_dir(program_execute_builtin, 'generated_tests/cl/builtin/int')
+program_execute_store = Group()
+program["Execute"]["Store"] = program_execute_store
+add_program_test_dir(program_execute_store, 'generated_tests/cl/store')
-- 
1.7.11.4



More information about the Piglit mailing list