[Piglit] [PATCH 2/2] cl: Add generated tests for global and local stores
Dylan Baker
baker.dylan.c at gmail.com
Wed Jul 31 23:37:16 PDT 2013
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:
> 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))
>
In python comma is always followed by a space.
Also, have you considered a simpler list comprehension? They are lazy, and
thus use less memory, besides not creating two useless objects. ie:
def gen_array(size):
return ' '.join([str(i) for i in xrange(size * 8)])
In python all top level function and class definitions should have exactly
two blank lines proceeding and following them: there should be two, not one
blank line here
> +
> +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')
>
and here
> +
> +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
>
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?
+ 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')
>
In python do not use \ inside of parenthesis, python will concatenate those
lines.
However, that said I would use a triple quote and format() here to make
this code more readable. I'm not familiar with CL, so I may not get the CL
syntax perfect, but something like:
f.write("""
[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)))
if you import textwrap you can use textwrap.dedent() to indent the text in
the for/if block and still have it print at the root of your text document.
and for the next one also.
+ 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
>
> _______________________________________________
> Piglit mailing list
> Piglit at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/piglit
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/piglit/attachments/20130731/6c0a8a79/attachment-0001.html>
More information about the Piglit
mailing list