Commit 2a43c940 authored by Tom Stellard's avatar Tom Stellard

cl: Add generated tests for global and local stores v2

v2:
  - Fix coding style
Acked-by: Dylan Baker's avatarDylan Baker <baker.dylan.c@gmail.com>
parent 4869588d
......@@ -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
......
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];
}
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];
}
#!/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@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()
......@@ -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')
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment