Add tests for creating program with IL

For testing boost::compute::program::create_with_il() 32- and 64-bit
SPIR-V binaries were added to test/data dir. File test/data/program.cl
contains OpenCL C source code of those binaries.

Author: jszuppe

test/CMakeLists.txt

@@ -50,6 +50,9 @@ if(${BOOST_COMPUTE_ENABLE_COVERAGE} AND "${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU
   add_definitions(-fprofile-arcs -ftest-coverage)
 endif()
 
+# add path to test data dir
+add_definitions(-DBOOST_COMPUTE_TEST_DATA_PATH="${CMAKE_CURRENT_SOURCE_DIR}/data")
+
 function(add_compute_test TEST_NAME TEST_SOURCE)
   get_filename_component(TEST_TARGET ${TEST_SOURCE} NAME_WE)
   add_executable(${TEST_TARGET} ${TEST_SOURCE})

test/data/program.cl

@@ -0,0 +1,15 @@
+//---------------------------------------------------------------------------//
+// Copyright (c) 2017 Jakub Szuppe <j.szuppe@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.
+//---------------------------------------------------------------------------//
+
+__kernel void foobar(__global int* x)
+{
+    const int gid = get_global_id(0);
+    x[gid] = gid;
+}

test/data/program.spirv32

@@ -126,6 +126,82 @@ BOOST_AUTO_TEST_CASE(create_with_binary)
     BOOST_CHECK_EQUAL(binary_bar_kernel.name(), std::string("bar"));
 }
 
+#ifdef BOOST_COMPUTE_CL_VERSION_2_1
+BOOST_AUTO_TEST_CASE(create_with_il)
+{
+    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)
+{
+    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)
+{
+    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]