Merge pull request #724 from jszuppe/pr_program_with_il

Add program::create_program_with_il()

Author: kylelutz

include/boost/compute/detail/get_object_info.hpp

@@ -183,6 +183,8 @@ struct get_object_info_impl<std::vector<T> >
             BOOST_THROW_EXCEPTION(opencl_error(ret));
         }
 
+        if(size == 0) return std::vector<T>();
+
         std::vector<T> vector(size / sizeof(T));
         ret = function(info, size, &vector[0], 0);
         if(ret != CL_SUCCESS){

include/boost/compute/program.hpp

@@ -179,6 +179,14 @@ class program
         return binary;
     }
 
+    #if defined(BOOST_COMPUTE_CL_VERSION_2_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
+    /// Returns the SPIR-V binary for the program.
+    std::vector<unsigned char> il_binary() const
+    {
+        return get_info<std::vector<unsigned char> >(CL_PROGRAM_IL);
+    }
+    #endif // BOOST_COMPUTE_CL_VERSION_2_1
+
     std::vector<device> get_devices() const
     {
         std::vector<cl_device_id> device_ids =
@@ -318,7 +326,6 @@ class program
             );
         }
 
-
         if(ret != CL_SUCCESS){
             BOOST_THROW_EXCEPTION(opencl_error(ret));
         }
@@ -574,6 +581,65 @@ class program
     }
     #endif // BOOST_COMPUTE_CL_VERSION_1_2
 
+    #if defined(BOOST_COMPUTE_CL_VERSION_2_1) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
+    /// Creates a new program with \p il_binary (SPIR-V binary)
+    /// of \p il_size size in \p context.
+    ///
+    /// \opencl_version_warning{2,1}
+    ///
+    /// \see_opencl_ref{clCreateProgramWithIL}
+    static program create_with_il(const void * il_binary,
+                                  const size_t il_size,
+                                  const context &context)
+    {
+        cl_int error = 0;
+
+        cl_program program_ = clCreateProgramWithIL(
+            context.get(), il_binary, il_size, &error
+        );
+
+        if(!program_){
+            BOOST_THROW_EXCEPTION(opencl_error(error));
+        }
+
+        return program(program_, false);
+    }
+
+    /// Creates a new program with \p il_binary (SPIR-V binary)
+    /// in \p context.
+    ///
+    /// \opencl_version_warning{2,1}
+    ///
+    /// \see_opencl_ref{clCreateProgramWithIL}
+    static program create_with_il(const std::vector<unsigned char> &il_binary,
+                                  const context &context)
+    {
+        return create_with_il(&il_binary[0], il_binary.size(), context);
+    }
+
+    /// Creates a new program in \p context using SPIR-V
+    /// binary \p file.
+    ///
+    /// \opencl_version_warning{2,1}
+    ///
+    /// \see_opencl_ref{clCreateProgramWithIL}
+    static program create_with_il_file(const std::string &file,
+                                       const context &context)
+    {
+        // open file stream
+        std::ifstream stream(file.c_str(), std::ios::in | std::ios::binary);
+
+        // read binary
+        std::vector<unsigned char> il(
+            (std::istreambuf_iterator<char>(stream)),
+            std::istreambuf_iterator<char>()
+        );
+
+        // create program
+        return create_with_il(&il[0], il.size(), context);
+    }
+    #endif // BOOST_COMPUTE_CL_VERSION_2_1
+
     /// Create a new program with \p source in \p context and builds it with \p options.
     /**
      * In case BOOST_COMPUTE_USE_OFFLINE_CACHE macro is defined,
@@ -701,6 +767,12 @@ BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS(program,
 )
 #endif // BOOST_COMPUTE_CL_VERSION_1_2
 
+#ifdef BOOST_COMPUTE_CL_VERSION_2_1
+BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS(program,
+    ((std::vector<unsigned char>, CL_PROGRAM_IL))
+)
+#endif // BOOST_COMPUTE_CL_VERSION_2_1
+
 } // end compute namespace
 } // end boost namespace
 

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]