[SYCL] Implement sycl_ext_oneapi_kernel_compiler_spirv (#12291)

Implements the extension described in #11954. This PR includes the
following changes:
- Adds a `create_kernel_bundle_from_source` overload for
`std::vector<std::byte>` kernel sources.
- Adds new `source_language::spirv`.
- Defines `SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV`.
- Adds support for SPIR-V kernels created from `std::vector<std::byte>`
sources.
- Moves `sycl_ext_oneapi_kernel_compiler_spirv.asciidoc` from `proposed`
to `experimental`.

---------

Signed-off-by: Michael Aziz <[email protected]>

Author: 0x12CC

sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc

@@ -52,14 +52,12 @@ This extension also depends on the following other SYCL extensions:
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.
-Interfaces defined in this specification may not be implemented yet or may be
-in a preliminary state.
-The specification itself may also change in incompatible ways before it is
-finalized.
-*Shipping software products should not rely on APIs defined in this
-specification.*
+This is an experimental extension specification, intended to provide early 
+access to features and gather community feedback. Interfaces defined in 
+this specification are implemented in DPC++, but they are not finalized 
+and may change incompatibly in future versions of DPC++ without prior notice. 
+*Shipping software products should not rely on APIs defined in 
+this specification.*
 
 
 == Overview

sycl/include/sycl/kernel_bundle.hpp

@@ -25,6 +25,7 @@
 #include <sycl/ext/oneapi/properties/property_value.hpp> // and log
 
 #include <array>       // for array
+#include <cstddef>     // for std::byte
 #include <cstring>     // for size_t, memcpy
 #include <functional>  // for function
 #include <iterator>    // for distance
@@ -46,7 +47,7 @@ auto get_native(const kernel_bundle<State> &Obj)
 namespace detail {
 class kernel_id_impl;
 class kernel_impl;
-}
+} // namespace detail
 
 template <typename KernelName> kernel_id get_kernel_id();
 
@@ -886,11 +887,18 @@ __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE,
 /////////////////////////
 // syclex::create_kernel_bundle_from_source
 /////////////////////////
+
+__SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::string &Source);
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
 __SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
-create_kernel_bundle_from_source(
-    const context &SyclContext,
-    sycl::ext::oneapi::experimental::source_language Language,
-    const std::string &Source);
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::vector<std::byte> &Bytes);
+#endif
 
 /////////////////////////
 // syclex::build(source_kb) => exe_kb

sycl/include/sycl/kernel_bundle_enums.hpp

@@ -20,7 +20,7 @@ enum class bundle_state : char {
 
 namespace ext::oneapi::experimental {
 
-enum class source_language : int { opencl = 0 /* sycl , spir-v, cuda */ };
+enum class source_language : int { opencl = 0, spirv = 1 /* sycl, cuda */ };
 
 } // namespace ext::oneapi::experimental
 

sycl/source/detail/kernel_bundle_impl.hpp

@@ -22,6 +22,7 @@
 
 #include <algorithm>
 #include <cassert>
+#include <cstdint>
 #include <cstring>
 #include <memory>
 #include <vector>
@@ -334,6 +335,14 @@ class kernel_bundle_impl {
       : MContext(Context), MDevices(Context.get_devices()),
         MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src) {}
 
+  // oneapi_ext_kernel_compiler
+  // construct from source bytes
+  kernel_bundle_impl(const context &Context, syclex::source_language Lang,
+                     const std::vector<std::byte> &Bytes)
+      : MContext(Context), MDevices(Context.get_devices()),
+        MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Bytes) {
+  }
+
   // oneapi_ext_kernel_compiler
   // interop constructor
   kernel_bundle_impl(context Ctx, std::vector<device> Devs,
@@ -350,17 +359,27 @@ class kernel_bundle_impl {
                     std::string *LogPtr) {
     assert(MState == bundle_state::ext_oneapi_source &&
            "bundle_state::ext_oneapi_source required");
-    assert(Language == syclex::source_language::opencl &&
-           "TODO: add other Languages. Must be OpenCL");
-    if (Language != syclex::source_language::opencl)
+
+    const auto spirv = [&]() -> std::vector<uint8_t> {
+      if (Language == syclex::source_language::opencl) {
+        // if successful, the log is empty. if failed, throws an error with the
+        // compilation log.
+        const auto &SourceStr = std::get<std::string>(this->Source);
+        return syclex::detail::OpenCLC_to_SPIRV(SourceStr, BuildOptions,
+                                                LogPtr);
+      }
+      if (Language == syclex::source_language::spirv) {
+        const auto &SourceBytes =
+            std::get<std::vector<std::byte>>(this->Source);
+        std::vector<uint8_t> Result(SourceBytes.size());
+        std::transform(SourceBytes.cbegin(), SourceBytes.cend(), Result.begin(),
+                       [](std::byte B) { return static_cast<uint8_t>(B); });
+        return Result;
+      }
       throw sycl::exception(
           make_error_code(errc::invalid),
-          "OpenCL C is the only supported language at this time");
-
-    // if successful, the log is empty. if failed, throws an error with the
-    // compilation log.
-    auto spirv =
-        syclex::detail::OpenCLC_to_SPIRV(this->Source, BuildOptions, LogPtr);
+          "OpenCL C and SPIR-V are the only supported languages at this time");
+    }();
 
     // see also program_manager.cpp::createSpirvProgram()
     using ContextImplPtr = std::shared_ptr<sycl::detail::context_impl>;
@@ -682,7 +701,7 @@ class kernel_bundle_impl {
   bundle_state MState;
   // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames
   const syclex::source_language Language = syclex::source_language::opencl;
-  const std::string Source;
+  const std::variant<std::string, std::vector<std::byte>> Source;
   // only kernel_bundles created from source have KernelNames member.
   std::vector<std::string> KernelNames;
 };

sycl/source/feature_test.hpp.in

@@ -45,6 +45,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_GROUP_SORT 1
 #define SYCL_EXT_ONEAPI_KERNEL_COMPILER 1
 #define SYCL_EXT_ONEAPI_KERNEL_COMPILER_OPENCL 1
+#define SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV 1
 #define SYCL_EXT_ONEAPI_MAX_WORK_GROUP_QUERY 1
 #define SYCL_EXT_ONEAPI_ND_RANGE_REDUCTIONS 1
 #define SYCL_EXT_ONEAPI_DEFAULT_CONTEXT 1

sycl/source/kernel_bundle.cpp

@@ -12,7 +12,9 @@
 #include <detail/kernel_id_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
 
+#include <cstddef>
 #include <set>
+#include <vector>
 
 namespace sycl {
 inline namespace _V1 {
@@ -368,12 +370,16 @@ using kernel_bundle_impl = sycl::detail::kernel_bundle_impl;
 // syclex::is_source_kernel_bundle_supported
 /////////////////////////
 bool is_source_kernel_bundle_supported(backend BE, source_language Language) {
-  // at the moment, OpenCL is the only language supported
-  // and it's support is limited to the opencl and level_zero backends.
+  // Support is limited to the opencl and level_zero backends.
   bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) ||
                        (BE == sycl::backend::opencl);
-  if ((Language == source_language::opencl) && BE_Acceptable) {
-    return detail::OpenCLC_Compilation_Available();
+  if (BE_Acceptable) {
+    // At the moment, OpenCL and SPIR-V are the only supported languages.
+    if (Language == source_language::opencl) {
+      return detail::OpenCLC_Compilation_Available();
+    } else if (Language == source_language::spirv) {
+      return true;
+    }
   }
 
   // otherwise
@@ -383,6 +389,7 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) {
 /////////////////////////
 // syclex::create_kernel_bundle_from_source
 /////////////////////////
+
 source_kb create_kernel_bundle_from_source(const context &SyclContext,
                                            source_language Language,
                                            const std::string &Source) {
@@ -399,6 +406,20 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext,
   return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
 }
 
+source_kb
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::vector<std::byte> &Bytes) {
+  backend BE = SyclContext.get_backend();
+  if (!is_source_kernel_bundle_supported(BE, Language))
+    throw sycl::exception(make_error_code(errc::invalid),
+                          "kernel_bundle creation from source not supported");
+
+  std::shared_ptr<kernel_bundle_impl> KBImpl =
+      std::make_shared<kernel_bundle_impl>(SyclContext, Language, Bytes);
+  return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
+}
+
 /////////////////////////
 // syclex::detail::build_from_source(source_kb) => exe_kb
 /////////////////////////