diff options
| -rw-r--r-- | src/CMakeLists.txt | 14 | ||||
| -rw-r--r-- | src/core/cpu/builtins.cpp | 11 | ||||
| -rw-r--r-- | src/runtime/CMakeLists.txt | 70 | ||||
| -rw-r--r-- | src/runtime/builtins.def | 18 | ||||
| -rwxr-xr-x | src/runtime/builtins.py | 365 | ||||
| -rwxr-xr-x | src/runtime/embed.py | 40 | ||||
| -rw-r--r-- | src/runtime/stdlib.c | 5 | ||||
| -rw-r--r-- | tests/test_builtins.cpp | 36 |
8 files changed, 523 insertions, 36 deletions
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 348d4bc..14927bf 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -45,6 +45,10 @@ set(COAL_SRC_FILES ${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib.h.embed.h ${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib.c.bc.embed.h + ${CMAKE_CURRENT_BINARY_DIR}/runtime/builtins_impl.h + ${CMAKE_CURRENT_BINARY_DIR}/runtime/builtins_def.h + ${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib_impl.h + ${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib_def.h ) add_subdirectory(runtime) @@ -55,8 +59,18 @@ set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib.h.embed.h PROPERTIES GENERATED 1) set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib.c.bc.embed.h PROPERTIES GENERATED 1) +set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/runtime/builtins_impl.h + PROPERTIES GENERATED 1) +set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/runtime/builtins_def.h + PROPERTIES GENERATED 1) +set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib_impl.h + PROPERTIES GENERATED 1) +set_source_files_properties(${CMAKE_CURRENT_BINARY_DIR}/runtime/stdlib_def.h + PROPERTIES GENERATED 1) + add_dependencies(OpenCL generate_stdlib_h) add_dependencies(OpenCL generate_stdlib_c) +add_dependencies(OpenCL generate_builtins) SET(LIBRARY_OUTPUT_PATH ${Coal_BINARY_DIR}/lib) diff --git a/src/core/cpu/builtins.cpp b/src/core/cpu/builtins.cpp index 98a6e65..1993e10 100644 --- a/src/core/cpu/builtins.cpp +++ b/src/core/cpu/builtins.cpp @@ -47,8 +47,10 @@ #include <signal.h> #include <llvm/Function.h> + #include <iostream> #include <cstring> +#include <cmath> #include <stdio.h> @@ -396,6 +398,12 @@ static void read_imageuif(uint32_t *result, Image2D *image, float x, float y, } /* + * Built-in functions generated by src/runtime/builtins.py + */ + +#include <runtime/builtins_impl.h> + +/* * Bridge between LLVM and us */ static void unimplemented_stub() @@ -456,6 +464,9 @@ void *getBuiltin(const std::string &name) else if (name == "__cpu_read_imageuif") return (void *)&read_imageuif; + // Built-in functions generated by src/runtime/builtins.py +#include <runtime/builtins_def.h> + else if (name == "debug") return (void *)&printf; diff --git a/src/runtime/CMakeLists.txt b/src/runtime/CMakeLists.txt index e381ebf..f825825 100644 --- a/src/runtime/CMakeLists.txt +++ b/src/runtime/CMakeLists.txt @@ -1,35 +1,45 @@ -macro(EMBED_FILE INFILE INNAME) - add_custom_command( - OUTPUT ${INFILE}.embed.h - COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/embed.py - ${CMAKE_CURRENT_SOURCE_DIR}/${INFILE} - ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.embed.h - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/embed.py - ${CMAKE_CURRENT_SOURCE_DIR}/${INFILE}) +add_custom_command( + OUTPUT stdlib.h.embed.h + COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/embed.py + ${CMAKE_CURRENT_BINARY_DIR}/stdlib.h.embed.h + ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.h + ${CMAKE_CURRENT_BINARY_DIR}/stdlib_def.h + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/embed.py + ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.h + ${CMAKE_CURRENT_BINARY_DIR}/stdlib_def.h) - add_custom_target(generate_${INNAME} DEPENDS - ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.embed.h) -endmacro(EMBED_FILE) +add_custom_target(generate_stdlib_h DEPENDS + ${CMAKE_CURRENT_BINARY_DIR}/stdlib.h.embed.h) -macro(COMPILE_BITCODE INFILE INNAME) - add_custom_command( - OUTPUT ${INFILE}.bc - COMMAND clang -c -emit-llvm -x cl -O2 - ${CMAKE_CURRENT_SOURCE_DIR}/${INFILE} - -o ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.bc - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/${INFILE}) +add_custom_command( + OUTPUT stdlib.c.bc + COMMAND clang -c -emit-llvm -x cl -O2 -nobuiltininc -nostdinc -fno-builtin + ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.c + -I${CMAKE_CURRENT_BINARY_DIR} + -o ${CMAKE_CURRENT_BINARY_DIR}/stdlib.c.bc + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/stdlib.c) - add_custom_command( - OUTPUT ${INFILE}.bc.embed.h - COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/embed.py - ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.bc - ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.bc.embed.h - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/embed.py - ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.bc) +add_custom_command( + OUTPUT stdlib.c.bc.embed.h + COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/embed.py + ${CMAKE_CURRENT_BINARY_DIR}/stdlib.c.bc.embed.h + ${CMAKE_CURRENT_BINARY_DIR}/stdlib.c.bc + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/embed.py + ${CMAKE_CURRENT_BINARY_DIR}/stdlib.c.bc) - add_custom_target(generate_${INNAME} DEPENDS - ${CMAKE_CURRENT_BINARY_DIR}/${INFILE}.bc.embed.h) -endmacro(COMPILE_BITCODE) +add_custom_target(generate_stdlib_c DEPENDS + ${CMAKE_CURRENT_BINARY_DIR}/stdlib.c.bc.embed.h) -EMBED_FILE(stdlib.h stdlib_h) -COMPILE_BITCODE(stdlib.c stdlib_c) +add_custom_command( + OUTPUT builtins_def.h stdlib_def.h builtins_impl.h stdlib_impl.h + COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/builtins.py + ${CMAKE_CURRENT_SOURCE_DIR}/builtins.def + ${CMAKE_CURRENT_BINARY_DIR} + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/builtins.py + ${CMAKE_CURRENT_SOURCE_DIR}/builtins.def) + +add_custom_target(generate_builtins DEPENDS + ${CMAKE_CURRENT_BINARY_DIR}/builtins_def.h + ${CMAKE_CURRENT_BINARY_DIR}/builtins_impl.h + ${CMAKE_CURRENT_BINARY_DIR}/stdlib_def.h + ${CMAKE_CURRENT_BINARY_DIR}/stdlib_impl.h) diff --git a/src/runtime/builtins.def b/src/runtime/builtins.def new file mode 100644 index 0000000..39e705b --- /dev/null +++ b/src/runtime/builtins.def @@ -0,0 +1,18 @@ +def vecf : float2 float3 float4 float8 float16 +def veci : int2 int3 int4 int8 int16 + +def vec : $vecf $veci +def gentype : float $vecf + +func $type fmin $gentype : x:$type y:$type + return (x < y ? x : y); +end + +native float cos float : x:float + return std::cos(x); +end + +native $type cos $vecf : x:$type + for (unsigned int i=0; i<$vecdim; ++i) + result[i] = std::cos(x[i]); +end diff --git a/src/runtime/builtins.py b/src/runtime/builtins.py new file mode 100755 index 0000000..b38f7d1 --- /dev/null +++ b/src/runtime/builtins.py @@ -0,0 +1,365 @@ +#!/usr/bin/python +# -*- coding: utf-8 -*- +# +# Copyright (c) 2011, Denis Steckelmacher <steckdenis@yahoo.fr> +# All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions are met: +# * Redistributions of source code must retain the above copyright +# notice, this list of conditions and the following disclaimer. +# * Redistributions in binary form must reproduce the above copyright +# notice, this list of conditions and the following disclaimer in the +# documentation and/or other materials provided with the distribution. +# * Neither the name of the copyright holder nor the +# names of its contributors may be used to endorse or promote products +# derived from this software without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +# ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +# WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +# DISCLAIMED. IN NO EVENT SHALL <COPYRIGHT HOLDER> BE LIABLE FOR ANY +# DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES +# (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; +# LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND +# ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +# SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + + +# builtins.py <def> <outdir> + +import sys + +class Function: + class Arg: + def __init__(self, name, t): + self.name = name + self.t = t + + KIND_BUILTINS_IMPL = 0 # static function in builtins.cpp + KIND_BUILTINS_DEF = 1 # if (name == '__cpu_$name') return (void *)&name; + KIND_STDLIB_IMPL = 2 # OpenCL C function in stdlib.c + KIND_STDLIB_DEF = 3 # Header in stdlib.h + KIND_STDLIB_STUB = 4 # OpenCL C stub in stdlib.c: calls __cpu_$name + KIND_STDLIB_STUB_DEF = 5 # __cpu_$name declared in stdlib.c + + def __init__(self, name, native): + self.name = name + self.native = native + + self.args = [] # Array <Arg> + self.types = [] # Array <str> + self.return_type = '' + self.body = '' + + def set_return_type(self, ty): + self.return_type = ty + + def append_body(self, body): + self.body += body + + def add_arg(self, name, ty): + self.args.append(self.Arg(name, ty)) + + def add_type(self, ty): + self.types.append(ty) + + def mangled_name(self, current_type): + return_type = self.return_type + if return_type == '$type': + return_type = current_type + + rs = return_type + '_' + self.name + first = True + + for a in self.args: + if first: + rs += '_' + first = False + + arg_type = a.t + if arg_type == '$type': + arg_type = current_type + rs += arg_type + + return rs + + def arg_list(self, current_type, handle_first_arg): + rs = '' + first = True + append_arg = None + + # We may need a first "result" arg + if handle_first_arg: + return_type = self.return_type + if return_type == '$type': + return_type = current_type + + if return_type[-1].isdigit(): + # Return is a vector + append_arg = self.Arg('result', return_type) + + if append_arg: + args = [append_arg] + args.extend(self.args) + else: + args = self.args + + for arg in args: + # Resolve type + arg_type = arg.t + if arg_type == '$type': + arg_type = current_type + + # We need to pass vector arguments as pointers + arg_vector = False + if handle_first_arg: + arg_vector = arg_type[-1].isdigit() + arg_type = arg_type.rstrip('0123456789') + + # Build the string + if not first: + rs += ', ' + first = False + + rs += arg_type + ' ' + + if arg_vector: + rs += '*' + + rs += arg.name + + return rs + + def write(self, current_type, kind): + # Template: + # (static) $ret_type $name($args) { + # $body + # } + rs = '' + + if kind == self.KIND_BUILTINS_IMPL: + rs = 'static ' + elif kind == self.KIND_BUILTINS_DEF: + rs += ' else if (name == "__cpu_' + self.mangled_name(current_type) + '")\n' + rs += ' return (void *)&' + self.mangled_name(current_type) + ';\n' + return rs + + # Calculate return type + return_type = self.return_type + if return_type == '$type': + return_type = current_type + + if (kind == self.KIND_BUILTINS_IMPL or kind == self.KIND_STDLIB_STUB_DEF) \ + and return_type[-1].isdigit(): + return_type = 'void' # We'll use a 'result' argument + + rs += return_type + ' ' + + # Append mangled name if needed + if kind == self.KIND_BUILTINS_IMPL: + rs += self.mangled_name(current_type) + elif kind == self.KIND_STDLIB_STUB_DEF: + rs += '__cpu_' + self.mangled_name(current_type) + else: + # No need to mangle the name, but add OVERLOAD + rs += 'OVERLOAD ' + self.name + + # Print function args + rs += '(' + rs += self.arg_list(current_type, kind == self.KIND_BUILTINS_IMPL or \ + kind == self.KIND_STDLIB_STUB_DEF) + rs += ')' + + # If only a declaration, end it + if kind == self.KIND_STDLIB_DEF or kind == self.KIND_STDLIB_STUB_DEF: + rs += ';\n' + return rs + + # Add the body + rs += '\n{\n' + + if kind == self.KIND_STDLIB_STUB: + # Special body : call __cpu_$name + return_is_vector = return_type[-1].isdigit() + if return_is_vector: + # Need to create a temporary + rs += ' ' + return_type + ' result;\n' + rs += '\n' + + # Call the cpu stub + rs += ' ' + if not return_is_vector: + rs += 'return ' + + rs += '__cpu_' + self.mangled_name(current_type) + '(' + + # Pass the result if needed + first = True + if return_is_vector: + rs += '(' + return_type.rstrip('0123456789') + ' *)&result' + first = False + + # Append the args + for arg in self.args: + # Resolve type + arg_type = arg.t + if arg_type == '$type': + arg_type = current_type + arg_vector = arg_type[-1].isdigit() + + if not first: + rs += ', ' + first = False + + # We need to pass vector arguments as pointers + if arg_vector: + rs += '(' + arg_type.rstrip('0123456789') + ' *)&' + arg.name + else: + rs += arg.name + + # End the call + rs += ');\n' + + if return_is_vector: + rs += '\n return result;\n' + + rs += '}\n\n' + else: + # Simply copy the body + vecdim = '1' + + if current_type[-1].isdigit(): + if current_type[-2].isdigit(): + vecdim = current_type[-2:] + else: + vecdim = current_type[-1] + + rs += self.body.replace('$type', current_type) \ + .replace('$vecdim', vecdim) + rs += '\n}\n\n' + + return rs + +class Generator: + builtins_impl_file = 'builtins_impl.h' # static functions + builtins_def_file = 'builtins_def.h' # if () in getBuiltin + stdlib_impl_file = 'stdlib_impl.h' # stdlib.c functions + stdlib_def_file = 'stdlib_def.h' # stdlib.h definitions + + def __init__(self, out_path): + self.out_path = out_path + + # Buffers + self.builtins_impl_buffer = '' + self.builtins_def_buffer = '' + self.stdlib_impl_buffer = '' + self.stdlib_def_buffer = '' + + def add_function(self, function): + for t in function.types: + if function.native: + self.stdlib_impl_buffer += function.write(t, function.KIND_STDLIB_STUB_DEF) + self.stdlib_impl_buffer += function.write(t, function.KIND_STDLIB_STUB) + self.stdlib_def_buffer += function.write(t, function.KIND_STDLIB_DEF) + self.builtins_impl_buffer += function.write(t, function.KIND_BUILTINS_IMPL) + self.builtins_def_buffer += function.write(t, function.KIND_BUILTINS_DEF) + else: + self.stdlib_def_buffer += function.write(t, function.KIND_STDLIB_DEF) + self.stdlib_impl_buffer += function.write(t, function.KIND_STDLIB_IMPL) + + def write(self): + of = open(self.out_path + '/' + self.stdlib_def_file, 'w') + of.write(self.stdlib_def_buffer) + of.close() + + of = open(self.out_path + '/' + self.stdlib_impl_file, 'w') + of.write(self.stdlib_impl_buffer) + of.close() + + of = open(self.out_path + '/' + self.builtins_def_file, 'w') + of.write(self.builtins_def_buffer) + of.close() + + of = open(self.out_path + '/' + self.builtins_impl_file, 'w') + of.write(self.builtins_impl_buffer) + of.close() + +class Parser: + def __init__(self, generator, def_file_name): + self.generator = generator + self.def_file_name = def_file_name + + self.defs = {} + + def replace_variable(self, token): + result = [] + + if token[0] == '$': + for tok in self.defs[token[1:]]: + result.extend(self.replace_variable(tok)) + else: + result.append(token) + + return result + + def parse(self): + def_file = open(self.def_file_name, 'rb') + current_function = None + + for line in def_file: + if current_function: + # End if we encounter an end + if line.startswith('end'): + self.generator.add_function(current_function) + current_function = None + else: + # Add a line to the body + current_function.append_body(line) + else: + line = line.strip() + tokens = line.split(' ') + tok = tokens[0] + + if tok == 'def': + # A definition : def <variable> : [values] + name = tokens[1] + values = [] + + for token in tokens[3:]: + values.extend(self.replace_variable(token)) + + self.defs[name] = values + elif tok == 'func' or tok == 'native': + # Function : func|native <ret_type> <name> [types] : [args] + current_function = Function(tokens[2], \ + tokens[0] == 'native') + + current_function.set_return_type(tokens[1]) + + # Explore the types and args + in_types = True + + for token in tokens[3:]: + if token == ':': + in_types = False + elif in_types: + for ty in self.replace_variable(token): + current_function.add_type(ty) + else: + # Parameters + parts = token.split(':') + current_function.add_arg(parts[0], parts[1]) + + def_file.close() + +if __name__ == '__main__': + def_file = sys.argv[1] + out_dir = sys.argv[2] + + gen = Generator(out_dir) + parser = Parser(gen, def_file) + + parser.parse() + gen.write() diff --git a/src/runtime/embed.py b/src/runtime/embed.py index 81a4462..7d40cbb 100755 --- a/src/runtime/embed.py +++ b/src/runtime/embed.py @@ -1,16 +1,44 @@ #!/usr/bin/python # -*- coding: utf-8 -*- +# +# Copyright (c) 2011, Denis Steckelmacher <steckdenis@yahoo.fr> +# All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions are met: +# * Redistributions of source code must retain the above copyright +# notice, this list of conditions and the following disclaimer. +# * Redistributions in binary form must reproduce the above copyright +# notice, this list of conditions and the following disclaimer in the +# documentation and/or other materials provided with the distribution. +# * Neither the name of the copyright holder nor the +# names of its contributors may be used to endorse or promote products +# derived from this software without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND +# ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED +# WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +# DISCLAIMED. IN NO EVENT SHALL <COPYRIGHT HOLDER> BE LIABLE FOR ANY +# DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES +# (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; +# LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND +# ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +# SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. -# embed.py <filename> <outfile> -# <filename> => <outfile> +# embed.py <outfile> <filenames..> +# <filenames> => <outfile> import sys -infile = open(sys.argv[1], 'rb') -name = sys.argv[1].split('/')[-1].replace('.', '_') -outfile = open(sys.argv[2], 'w') +outfile = open(sys.argv[1], 'w') +name = sys.argv[1].split('/')[-1].replace('.embed.h', '').replace('.', '_') -data = infile.read() +data = '' + +for i in xrange(len(sys.argv) - 1): + infile = open(sys.argv[i + 1], 'rb') + data += infile.read() # Header outfile.write('#ifndef __%s__\n' % name.upper()) diff --git a/src/runtime/stdlib.c b/src/runtime/stdlib.c index 0f10f0d..c2bb2eb 100644 --- a/src/runtime/stdlib.c +++ b/src/runtime/stdlib.c @@ -267,3 +267,8 @@ int OVERLOAD get_image_channel_order(image3d_t image) return __cpu_get_image_channel_order(image); } +/* + * Built-in functions generated by src/runtime/builtins.py + */ + +#include <stdlib_impl.h> diff --git a/tests/test_builtins.cpp b/tests/test_builtins.cpp index 94e836e..d8a5826 100644 --- a/tests/test_builtins.cpp +++ b/tests/test_builtins.cpp @@ -95,6 +95,17 @@ const char image_source[] = " fcolor = read_imagef(image3, sampler, fcoords);\n" "}\n"; +const char builtins_source[] = + "__kernel void test_case(__global uint *rs) {\n" + " float2 f2;\n" + "\n" + " f2.x = 1.0f;\n" + " f2.y = 0.0f;\n" + "\n" + " if (cos(f2).y != 1.0f) { *rs = 1; return; }\n" + " if (cos(0.0f) != 1.0f) { *rs = 2; return; }\n" + "}\n"; + enum TestCaseKind { NormalKind, @@ -357,6 +368,30 @@ START_TEST (test_image) } END_TEST +START_TEST (test_builtins) +{ + uint32_t rs = run_kernel(builtins_source, NormalKind); + const char *errstr = 0; + + switch (rs) + { + case 1: + errstr = "float2 cos(float2) doesn't behave correctly"; + break; + case 2: + errstr = "float cos(float) doesn't behave correctly"; + break; + default: + errstr = default_error(rs); + } + + fail_if( + errstr != 0, + errstr + ); +} +END_TEST + TCase *cl_builtins_tcase_create(void) { TCase *tc = NULL; @@ -364,5 +399,6 @@ TCase *cl_builtins_tcase_create(void) tcase_add_test(tc, test_sampler); tcase_add_test(tc, test_barrier); tcase_add_test(tc, test_image); + tcase_add_test(tc, test_builtins); return tc; } |