[SYCL] Implement sycl_ext_oneapi_kernel_compiler_spirv

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

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <cstddef>                         // for std::byte
 #include <sycl/backend_types.hpp>          // for backend, backend_return_t
 #include <sycl/context.hpp>                // for context
 #include <sycl/detail/export.hpp>          // for __SYCL_EXPORT
@@ -887,10 +888,14 @@ __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,
-    sycl::ext::oneapi::experimental::source_language Language,
-    const std::string &Source);
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::string &Source);
+
+__SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::vector<std::byte> &Bytes);
 
 /////////////////////////
 // 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,30 @@ 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)
-      throw sycl::exception(
-          make_error_code(errc::invalid),
-          "OpenCL C is the only supported language at this time");
+    assert((Language == syclex::source_language::opencl ||
+            Language == syclex::source_language::spirv) &&
+           "TODO: add other Languages. Must be OpenCL or SPIR-V");
 
     // 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);
+    const auto spirv = [&]() -> std::vector<uint8_t> {
+      if (Language == syclex::source_language::opencl) {
+        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 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 +704,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 {
@@ -372,7 +374,9 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) {
   // and it's 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) {
+  if ((Language == source_language::opencl ||
+       Language == source_language::spirv) &&
+      BE_Acceptable) {
     return detail::OpenCLC_Compilation_Available();
   }
 
@@ -399,6 +403,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
 /////////////////////////

sycl/test-e2e/KernelCompiler/Kernels/my_kernel.cl

@@ -0,0 +1,4 @@
+__kernel void my_kernel(__global int *in, __global int *out) {
+  size_t i = get_global_id(0);
+  out[i] = in[i]*2 + 100;
+}

sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp

@@ -0,0 +1,88 @@
+//==- kernel_compiler_spirv.cpp --------------------------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: ocloc
+
+// RUN: ocloc -spv_only -file %S/Kernels/my_kernel.cl -o %t.spv
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out %t.spv
+
+// Test case for the sycl_ext_oneapi_kernel_compiler_spirv extension. This test
+// loads a pre-compiled kernel from a SPIR-V file and runs it.
+
+#include <cassert>
+#include <fstream>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier,
+                    int added) {
+  constexpr int N = 4;
+  cl_int InputArray[N] = {0, 1, 2, 3};
+  cl_int OutputArray[N] = {};
+
+  sycl::buffer InputBuf(InputArray, sycl::range<1>(N));
+  sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N));
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.set_arg(0, InputBuf.get_access<sycl::access::mode::read>(CGH));
+    CGH.set_arg(1, OutputBuf.get_access<sycl::access::mode::write>(CGH));
+    CGH.parallel_for(sycl::range<1>{N}, Kernel);
+  });
+
+  sycl::host_accessor Out{OutputBuf};
+  for (int I = 0; I < N; I++) {
+    assert(Out[I] == ((I * multiplier) + added));
+  }
+}
+
+void testKernelFromSpvFile(std::string file_name) {
+  namespace syclex = sycl::ext::oneapi::experimental;
+
+  sycl::queue q;
+
+  // Read the SPIR-V module from disk.
+  std::ifstream spv_stream(file_name, std::ios::binary);
+  spv_stream.seekg(0, std::ios::end);
+  size_t sz = spv_stream.tellg();
+  spv_stream.seekg(0);
+  std::vector<std::byte> spv(sz);
+  spv_stream.read(reinterpret_cast<char *>(spv.data()), sz);
+
+  // Create a kernel bundle from the binary SPIR-V.
+  sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
+      syclex::create_kernel_bundle_from_source(
+          q.get_context(), syclex::source_language::spirv, spv);
+
+  // Build the SPIR-V module for our device.
+  sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe =
+      syclex::build(kb_src);
+
+  // Get a "kernel" object representing the kernel from the SPIR-V module.
+  sycl::kernel my_kernel = kb_exe.ext_oneapi_get_kernel("my_kernel");
+
+  // Test the kernel
+  auto my_num_args = my_kernel.get_info<sycl::info::kernel::num_args>();
+  assert(my_num_args == 2 && "my_kernel should take 2 args");
+  testSyclKernel(q, my_kernel, 2, 100);
+}
+
+int main(int argc, char **argv) {
+#ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV
+  static_assert(false, "KernelCompiler SPIR-V feature test macro undefined");
+#endif
+
+#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
+  assert(argc == 2 && "Usage: ./%t.out <kernel-spv-file>");
+  testKernelFromSpvFile(argv[1]);
+#else
+  static_assert(false, "Kernel Compiler feature test macro undefined");
+#endif
+  return 0;
+}