1 files changed, 401 insertions, 0 deletions
diff --git a/src/boost/libs/compute/test/test_program.cpp b/src/boost/libs/compute/test/test_program.cpp
new file mode 100644
index 00000000..7172c5e2
--- /dev/null
+++ b/src/boost/libs/compute/test/test_program.cpp
@@ -0,0 +1,401 @@
+//---------------------------------------------------------------------------//
+// 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()