[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:45:48 PDT 2013
On Wed, Jul 31, 2013 at 11:37 PM, Dylan Baker <baker.dylan.c at gmail.com>wrote:
> 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
>>
>
sorry, one more thing. The above function is double indented.
> +
>> +
>> +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/695d8718/attachment.html>
More information about the Piglit
mailing list