123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401 |
- //---------------------------------------------------------------------------//
- // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
- //
- // Distributed under the Boost Software License, Version 1.0
- // See accompanying file LICENSE_1_0.txt or copy at
- // http://www.boost.org/LICENSE_1_0.txt
- //
- // See http://boostorg.github.com/compute for more information.
- //---------------------------------------------------------------------------//
- #define BOOST_TEST_MODULE TestProgram
- #include <boost/test/unit_test.hpp>
- // disable the automatic kernel compilation debug messages. this allows the
- // test for program to check that compilation error exceptions are properly
- // thrown when invalid kernel code is passed to program::build().
- #undef BOOST_COMPUTE_DEBUG_KERNEL_COMPILATION
- #include <boost/compute/exception/program_build_failure.hpp>
- #include <boost/compute/kernel.hpp>
- #include <boost/compute/system.hpp>
- #include <boost/compute/program.hpp>
- #include <boost/compute/utility/source.hpp>
- #include "quirks.hpp"
- #include "context_setup.hpp"
- namespace compute = boost::compute;
- const char source[] =
- "__kernel void foo(__global float *x, const uint n) { }\n"
- "__kernel void bar(__global int *x, __global int *y) { }\n";
- BOOST_AUTO_TEST_CASE(get_program_info)
- {
- // create program
- boost::compute::program program =
- boost::compute::program::create_with_source(source, context);
- // build program
- program.build();
- // check program info
- #ifndef BOOST_COMPUTE_USE_OFFLINE_CACHE
- BOOST_CHECK(program.source().empty() == false);
- #endif
- BOOST_CHECK(program.get_context() == context);
- }
- BOOST_AUTO_TEST_CASE(program_source)
- {
- // create program from source
- boost::compute::program program =
- boost::compute::program::create_with_source(source, context);
- BOOST_CHECK_EQUAL(std::string(source), program.source());
- }
- BOOST_AUTO_TEST_CASE(program_multiple_sources)
- {
- std::vector<std::string> sources;
- sources.push_back("__kernel void foo(__global int* x) { }\n");
- sources.push_back("__kernel void bar(__global float* y) { }\n");
- // create program from sources
- boost::compute::program program =
- boost::compute::program::create_with_source(sources, context);
- program.build();
- boost::compute::kernel foo = program.create_kernel("foo");
- boost::compute::kernel bar = program.create_kernel("bar");
- }
- BOOST_AUTO_TEST_CASE(program_source_no_file)
- {
- // create program from a non-existant source file
- // and verifies it throws.
- BOOST_CHECK_THROW(boost::compute::program program =
- boost::compute::program::create_with_source_file
- (std::string(), context),
- std::ios_base::failure);
- }
- BOOST_AUTO_TEST_CASE(create_kernel)
- {
- boost::compute::program program =
- boost::compute::program::create_with_source(source, context);
- program.build();
- boost::compute::kernel foo = program.create_kernel("foo");
- boost::compute::kernel bar = program.create_kernel("bar");
- // try to create a kernel that doesn't exist
- BOOST_CHECK_THROW(program.create_kernel("baz"), boost::compute::opencl_error);
- }
- BOOST_AUTO_TEST_CASE(create_with_binary)
- {
- // create program from source
- boost::compute::program source_program =
- boost::compute::program::create_with_source(source, context);
- source_program.build();
- // create kernels in source program
- boost::compute::kernel source_foo_kernel = source_program.create_kernel("foo");
- boost::compute::kernel source_bar_kernel = source_program.create_kernel("bar");
- // check source kernels
- BOOST_CHECK_EQUAL(source_foo_kernel.name(), std::string("foo"));
- BOOST_CHECK_EQUAL(source_bar_kernel.name(), std::string("bar"));
- // get binary
- std::vector<unsigned char> binary = source_program.binary();
- // create program from binary
- boost::compute::program binary_program =
- boost::compute::program::create_with_binary(binary, context);
- binary_program.build();
- // create kernels in binary program
- boost::compute::kernel binary_foo_kernel = binary_program.create_kernel("foo");
- boost::compute::kernel binary_bar_kernel = binary_program.create_kernel("bar");
- // check binary kernels
- BOOST_CHECK_EQUAL(binary_foo_kernel.name(), std::string("foo"));
- BOOST_CHECK_EQUAL(binary_bar_kernel.name(), std::string("bar"));
- }
- #ifdef BOOST_COMPUTE_CL_VERSION_2_1
- BOOST_AUTO_TEST_CASE(create_with_il)
- {
- REQUIRES_OPENCL_VERSION(2, 1);
- size_t device_address_space_size = device.address_bits();
- std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH);
- if(device_address_space_size == 64)
- {
- file_path += "/program.spirv64";
- }
- else
- {
- file_path += "/program.spirv32";
- }
- // create program from il
- boost::compute::program il_program;
- BOOST_CHECK_NO_THROW(
- il_program = boost::compute::program::create_with_il_file(
- file_path, context
- )
- );
- BOOST_CHECK_NO_THROW(il_program.build());
- // create kernel (to check if program was loaded correctly)
- BOOST_CHECK_NO_THROW(il_program.create_kernel("foobar"));
- }
- BOOST_AUTO_TEST_CASE(get_program_il_binary)
- {
- REQUIRES_OPENCL_VERSION(2, 1);
- size_t device_address_space_size = device.address_bits();
- std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH);
- if(device_address_space_size == 64)
- {
- file_path += "/program.spirv64";
- }
- else
- {
- file_path += "/program.spirv32";
- }
- // create program from il
- boost::compute::program il_program;
- BOOST_CHECK_NO_THROW(
- il_program = boost::compute::program::create_with_il_file(
- file_path, context
- )
- );
- BOOST_CHECK_NO_THROW(il_program.build());
- std::vector<unsigned char> il_binary;
- BOOST_CHECK_NO_THROW(il_binary = il_program.il_binary());
- // create program from loaded il binary
- BOOST_CHECK_NO_THROW(
- il_program = boost::compute::program::create_with_il(il_binary, context)
- );
- BOOST_CHECK_NO_THROW(il_program.build());
- // create kernel (to check if program was loaded correctly)
- BOOST_CHECK_NO_THROW(il_program.create_kernel("foobar"));
- }
- BOOST_AUTO_TEST_CASE(get_program_il_binary_empty)
- {
- REQUIRES_OPENCL_VERSION(2, 1);
- boost::compute::program program;
- BOOST_CHECK_NO_THROW(
- program = boost::compute::program::create_with_source(source, context)
- );
- BOOST_CHECK_NO_THROW(program.build());
- std::vector<unsigned char> il_binary;
- il_binary = program.il_binary();
- BOOST_CHECK(il_binary.empty());
- }
- #endif // BOOST_COMPUTE_CL_VERSION_2_1
- BOOST_AUTO_TEST_CASE(create_with_source_doctest)
- {
- //! [create_with_source]
- std::string source = "__kernel void foo(__global int *data) { }";
- boost::compute::program foo_program =
- boost::compute::program::create_with_source(source, context);
- //! [create_with_source]
- foo_program.build();
- }
- #ifdef BOOST_COMPUTE_CL_VERSION_1_2
- BOOST_AUTO_TEST_CASE(compile_and_link)
- {
- REQUIRES_OPENCL_VERSION(1,2);
- if(!supports_compile_program(device) || !supports_link_program(device)) {
- return;
- }
- // create the library program
- const char library_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- // for some reason the apple opencl compilers complains if a prototype
- // for the square() function is not available, so we add it here
- T square(T);
- // generic square function definition
- T square(T x) { return x * x; }
- );
- compute::program library_program =
- compute::program::create_with_source(library_source, context);
- library_program.compile("-DT=int");
- // create the kernel program
- const char kernel_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- // forward declare square function
- extern int square(int);
- // square kernel definition
- __kernel void square_kernel(__global int *x)
- {
- x[0] = square(x[0]);
- }
- );
- compute::program square_program =
- compute::program::create_with_source(kernel_source, context);
- square_program.compile();
- // link the programs
- std::vector<compute::program> programs;
- programs.push_back(library_program);
- programs.push_back(square_program);
- compute::program linked_program =
- compute::program::link(programs, context);
- // create the square kernel
- compute::kernel square_kernel =
- linked_program.create_kernel("square_kernel");
- BOOST_CHECK_EQUAL(square_kernel.name(), "square_kernel");
- }
- BOOST_AUTO_TEST_CASE(compile_and_link_with_headers)
- {
- REQUIRES_OPENCL_VERSION(1,2);
- if(!supports_compile_program(device) || !supports_link_program(device)) {
- return;
- }
- // create the header programs
- const char square_header_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- T square(T x) { return x * x; }
- );
- const char div2_header_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
- T div2(T x) { return x / 2; }
- );
- compute::program square_header_program =
- compute::program::create_with_source(square_header_source, context);
- compute::program div2_header_program =
- compute::program::create_with_source(div2_header_source, context);
- // create the kernel program
- const char kernel_source[] =
- "#include \"square.h\"\n"
- "#include \"div2.h\"\n"
- "__kernel void squareby2_kernel(__global int *x)"
- "{"
- " x[0] = div2(square(x[0]));"
- "}";
- compute::program square_program =
- compute::program::create_with_source(kernel_source, context);
- std::vector<std::pair<std::string, compute::program> > header_programs;
- header_programs.push_back(std::make_pair("square.h", square_header_program));
- header_programs.push_back(std::make_pair("div2.h", div2_header_program));
- square_program.compile("-DT=int", header_programs);
- // link program
- std::vector<compute::program> programs;
- programs.push_back(square_program);
- compute::program linked_program =
- compute::program::link(programs, context);
- // create the square kernel
- compute::kernel square_kernel =
- linked_program.create_kernel("squareby2_kernel");
- BOOST_CHECK_EQUAL(square_kernel.name(), "squareby2_kernel");
- }
- #endif // BOOST_COMPUTE_CL_VERSION_1_2
- BOOST_AUTO_TEST_CASE(build_log)
- {
- const char invalid_source[] =
- "__kernel void foo(__global int *input) { !@#$%^&*() }";
- compute::program invalid_program =
- compute::program::create_with_source(invalid_source, context);
- try {
- invalid_program.build();
- // should not get here
- BOOST_CHECK(false);
- }
- catch(compute::opencl_error&){
- std::string log = invalid_program.build_log();
- BOOST_CHECK(!log.empty());
- }
- }
- BOOST_AUTO_TEST_CASE(program_build_exception)
- {
- const char invalid_source[] =
- "__kernel void foo(__global int *input) { !@#$%^&*() }";
- compute::program invalid_program =
- compute::program::create_with_source(invalid_source, context);
- BOOST_CHECK_THROW(invalid_program.build(),
- compute::program_build_failure);
- try {
- // POCL bug: https://github.com/pocl/pocl/issues/577
- if(pocl_bug_issue_577(device))
- {
- invalid_program =
- compute::program::create_with_source(invalid_source, context);
- }
- invalid_program.build();
- // should not get here
- BOOST_CHECK(false);
- }
- catch(compute::program_build_failure& e){
- BOOST_CHECK(e.build_log() == invalid_program.build_log());
- }
- catch(...)
- {
- // should not get here
- BOOST_CHECK(false);
- }
- }
- BOOST_AUTO_TEST_CASE(build_with_source_exception)
- {
- const char invalid_source[] =
- "__kernel void foo(__global int *input) { !@#$%^&*() }";
- BOOST_CHECK_THROW(compute::program::build_with_source(invalid_source, context),
- compute::program_build_failure);
- }
- BOOST_AUTO_TEST_CASE(build_with_source_file_exception)
- {
- std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH "/invalid_program.cl");
- BOOST_CHECK_THROW(compute::program::build_with_source_file(file_path, context),
- compute::program_build_failure);
- }
- BOOST_AUTO_TEST_SUITE_END()
|