402 lines
12 KiB
C++
402 lines
12 KiB
C++
//---------------------------------------------------------------------------//
|
|
// 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()
|