[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