[Piglit] [PATCH 2/2] cl: Add generated tests for global and local stores
Tom Stellard
tom at stellard.net
Thu Aug 1 19:34:12 PDT 2013
On Wed, Jul 31, 2013 at 11:37:16PM -0700, Dylan Baker wrote:
Hi Dylan,
Thanks for your comments.
> I have some python comments for you, they're mainly style type comments,
> but some of them should help your code readability a lot.
>
>
> On Wed, Jul 31, 2013 at 7:16 PM, Tom Stellard <tom at stellard.net> wrote:
>
> > +VEC_SIZES = ['', '2', '4', '8', '16']
> > +
...
> I only see the one use of VEC_SIZES. Is there a reason element 0 is set to
> '' and then you have an if statement here to change it to 1?
>
For the first element, 1 is the vector size, but '' is the string that
needs to be appended to the type name. e.g.
Type Size
----- -----
Float 1
Float2 2
Float4 4
etc.
I have addressed all you other comments in the attached patch.
-Tom
> + else:
> > + size = int(s)
> > + type_name = t + s
> > + f = begin_test(type_name, 'global')
> > + f.write( \
> > + '[test]\n' + \
-------------- next part --------------
>From 283219710449ec1c0235ca94e018fd4bb75f4702 Mon Sep 17 00:00:00 2001
From: Tom Stellard <thomas.stellard at amd.com>
Date: Wed, 31 Jul 2013 19:10:58 -0700
Subject: [PATCH 2/2] cl: Add generated tests for global and local stores v2
v2:
- Fix coding style
---
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 | 97 +++++++++++++++++++++++
tests/all_cl.tests | 3 +
5 files changed, 139 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..58aa3a2
--- /dev/null
+++ b/generated_tests/generate-cl-store-tests.py
@@ -0,0 +1,97 @@
+#!/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
+import textwrap
+
+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):
+ return ' '.join([str(i) for i in xrange(size * 8)])
+
+
+def print_config(f, type_name, addr_space):
+ f.write(textwrap.dedent("""
+ [config]
+ name: Store {type_name}
+ program_source_file: store-kernels-{addr_space}.inc
+ build_options: -D TYPE={type_name}
+ dimensions: 1
+ """.format(type_name=type_name, addr_space=addr_space)))
+
+
+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(textwrap.dedent("""
+ [test]
+ name: global address space
+ global_size: 1 0 0
+ kernel_name: store_global
+ arg_out: 0 buffer {type_name}[8] {gen_array}
+ arg_in: 1 buffer {type_name}[8] {gen_array}
+ [test]
+ name: global address space work items
+ global_size: 8 0 0
+ kernel_name: store_global_wi
+ arg_out: 0 buffer {type_name}[8] {gen_array}
+ arg_in: 1 buffer {type_name}[8] {gen_array}
+ """.format(type_name=type_name, gen_array=gen_array(size))))
+
+ f.close()
+
+ f = begin_test(type_name, 'local')
+ f.write(textwrap.dedent("""
+ [test]
+ name: local address space
+ global_size: 8 0 0
+ local_size: 8 0 0
+ kernel_name: store_local
+ arg_out: 0 buffer {type_name}[8] {gen_array}
+ arg_in: 1 buffer {type_name}[8] {gen_array}
+ """.format(type_name=type_name, gen_array=gen_array(size))))
+
+ 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