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 <[email protected]> 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' + \
>From 283219710449ec1c0235ca94e018fd4bb75f4702 Mon Sep 17 00:00:00 2001 From: Tom Stellard <[email protected]> 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 <[email protected]> +# +# + +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
_______________________________________________ Piglit mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/piglit
