[SYCL] Implement backend content extension

sycl/doc/extensions/proposed/sycl_ext_oneapi_device_image_backend_content.asciidoc → sycl/doc/extensions/experimental/sycl_ext_oneapi_device_image_backend_content.asciidoc

@@ -44,12 +44,10 @@ the SYCL specification refer to that revision.
 
 == 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.
+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 {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
 *Shipping software products should not rely on APIs defined in this
 specification.*
 
@@ -101,7 +99,7 @@ class device_image {
   backend ext_oneapi_get_backend() const noexcept;
   std::vector<std::byte> ext_oneapi_get_backend_content() const;
 
-  std::span<std::byte> ext_oneapi_get_backend_content_view() const; // Requires C++20
+  std::span<const std::byte> ext_oneapi_get_backend_content_view() const; // Requires C++20
 
   /*...*/
 };
@@ -148,15 +146,15 @@ See below for a description of the formats used by {dpcpp}.
 a@
 [source,c++]
 ----
-std::span<std::byte> ext_oneapi_get_content_backend_view() const;
+std::span<const std::byte> ext_oneapi_get_content_backend_view() const;
 ----
 !====
 
 Minimum C++ Version: {cpp}20
 
 _Constraints:_ Available only when `State` is `bundle_state::executable`.
 
-_Returns:_ A view of the raw backend content for this device image.
+_Returns:_ An immutable view of the raw backend content for this device image.
 The data behind this view has the same lifetime as the `device_image` object.
 The format of this data is implementation-defined.
 See below for a description of the formats used by {dpcpp}.

sycl/include/sycl/kernel_bundle.hpp

@@ -26,12 +26,15 @@
 #include <sycl/ext/oneapi/properties/property.hpp>       // build_options
 #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
-#include <memory>      // for shared_ptr, operator==, hash
+#include <array>      // for array
+#include <cstddef>    // for std::byte
+#include <cstring>    // for size_t, memcpy
+#include <functional> // for function
+#include <iterator>   // for distance
+#include <memory>     // for shared_ptr, operator==, hash
+#if __has_include(<span>)
+#include <span>
+#endif
 #include <string>      // for string
 #include <type_traits> // for enable_if_t, remove_refer...
 #include <utility>     // for move
@@ -123,6 +126,15 @@ class __SYCL_EXPORT device_image_plain {
 
   template <class T>
   friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
+
+  backend ext_oneapi_get_backend_impl() const noexcept;
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+  std::vector<std::byte> ext_oneapi_get_backend_content_impl() const;
+
+  std::pair<const std::byte *, const std::byte *>
+  ext_oneapi_get_backend_content_view_impl() const;
+#endif // HAS_STD_BYTE
 };
 } // namespace detail
 
@@ -145,6 +157,28 @@ class device_image : public detail::device_image_plain,
     return device_image_plain::has_kernel(KernelID, Dev);
   }
 
+  backend ext_oneapi_get_backend() const noexcept {
+    return device_image_plain::ext_oneapi_get_backend_impl();
+  }
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+  template <sycl::bundle_state T = State,
+            typename = std::enable_if_t<T == bundle_state::executable>>
+  std::vector<std::byte> ext_oneapi_get_backend_content() const {
+    return device_image_plain::ext_oneapi_get_backend_content_impl();
+  }
+
+#ifdef __cpp_lib_span
+  template <sycl::bundle_state T = State,
+            typename = std::enable_if_t<T == bundle_state::executable>>
+  std::span<const std::byte> ext_oneapi_get_backend_content_view() const {
+    const auto view =
+        device_image_plain::ext_oneapi_get_backend_content_view_impl();
+    return std::span<const std::byte>{view.first, view.second};
+  }
+#endif // __cpp_lib_span
+#endif // _HAS_STD_BYTE
+
 private:
   device_image(detail::DeviceImageImplPtr Impl)
       : device_image_plain(std::move(Impl)) {}

sycl/source/feature_test.hpp.in

@@ -113,6 +113,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1
 #define SYCL_EXT_ONEAPI_NUM_COMPUTE_UNITS 1
+#define SYCL_EXT_ONEAPI_DEVICE_IMAGE_BACKEND_CONTENT 1
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0
 

sycl/source/kernel_bundle.cpp

@@ -45,6 +45,27 @@ ur_native_handle_t device_image_plain::getNative() const {
   return impl->getNative();
 }
 
+backend device_image_plain::ext_oneapi_get_backend_impl() const noexcept {
+  return impl->get_context().get_backend();
+}
+
+std::vector<std::byte>
+device_image_plain::ext_oneapi_get_backend_content_impl() const {
+  return std::vector(reinterpret_cast<const std::byte *>(
+                         impl->get_bin_image_ref()->getRawData().BinaryStart),
+                     reinterpret_cast<const std::byte *>(
+                         impl->get_bin_image_ref()->getRawData().BinaryEnd));
+}
+
+std::pair<const std::byte *, const std::byte *>
+device_image_plain::ext_oneapi_get_backend_content_view_impl() const {
+  return std::make_pair(
+      reinterpret_cast<const std::byte *>(
+          impl->get_bin_image_ref()->getRawData().BinaryStart),
+      reinterpret_cast<const std::byte *>(
+          impl->get_bin_image_ref()->getRawData().BinaryEnd));
+}
+
 ////////////////////////////
 ///// kernel_bundle_plain
 ///////////////////////////

sycl/test-e2e/DeviceImageBackendContent/CUDA_interop_test.cpp

@@ -0,0 +1,33 @@
+// RUN: %{build} %cuda_options -o %t.out
+// RUN: %{run} %t.out
+// REQUIRES: cuda, cuda_dev_kit
+
+#include <cuda.h>
+#include <sycl/backend.hpp>
+#include <sycl/detail/core.hpp>
+#include <vector>
+
+int main() {
+  sycl::device d([](const sycl::device &d) {
+    return d.get_backend() == sycl::backend::opencl;
+  });
+  sycl::queue q{d};
+  sycl::context ctxt = q.get_context();
+  sycl::kernel_id k_id = sycl::get_kernel_id<class mykernel>();
+  auto bundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt, {k_id});
+  assert(!bundle.empty());
+  sycl::kernel krn = bundle.get_kernel(k_id);
+  sycl::buffer<int> buf(sycl::range<1>(1));
+  q.submit([&](sycl::handler &cgh) {
+    sycl::accessor acc(buf, cgh);
+    cgh.single_task<class mykernel>(krn, [=]() { acc[0] = 42; });
+  });
+  const auto img = *(bundle.begin());
+  const auto bytes = img.ext_oneapi_get_backend_content();
+  CUmodule m;
+  CUresult res =
+      cuModuleLoadData(&m, reinterpret_cast<const void *>(bytes.data()));
+  assert(result == CUDA_SUCCESS);
+  return 0;
+}

sycl/test-e2e/DeviceImageBackendContent/L0_interop_test.cpp

@@ -0,0 +1,97 @@
+// REQUIRES: level_zero, level_zero_dev_kit, aspect-usm_shared_allocations
+// RUN: %{build} %level_zero_options -fno-sycl-dead-args-optimization -o %t.out
+// RUN: %{run} %t.out
+//
+#include <level_zero/ze_api.h>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/usm.hpp>
+#include <vector>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (syclexp::nd_range_kernel<1>)) void iota(int *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = 42;
+}
+
+int main() {
+  sycl::device d([](const sycl::device &d) {
+    return d.get_backend() == sycl::backend::ext_oneapi_level_zero;
+  });
+  sycl::queue q{d};
+  sycl::context ctxt = q.get_context();
+
+  // The following ifndef is required due to a number of limitations of free
+  // function kernels. See CMPLRLLVM-61498.
+  // TODO: Remove it once these limitations are no longer there.
+#ifndef __SYCL_DEVICE_ONLY__
+  // First, run the kernel using the SYCL API.
+  auto Bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt);
+  sycl::kernel_id iota_id = syclexp::get_kernel_id<iota>();
+  sycl::kernel k_iota = Bundle.get_kernel(iota_id);
+  int *ptr = sycl::malloc_shared<int>(1, q);
+  *ptr = 0;
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(ptr);
+     cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota);
+   }).wait();
+
+  // Now, run the kernel by first getting its image as an executable,
+  // making an L0 kernel out of it and then making a SYCL kernel out of
+  // the L0 kernel. Run this kernel on the SYCL API and verify
+  // that it has the same result as the kernel that was run directly on SYCL
+  // API. First, get a kernel bundle that contains the kernel "iota".
+  auto exe_bndl = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+      ctxt, {d},
+      [&](const sycl::device_image<sycl::bundle_state::executable> &img) {
+        return img.has_kernel(iota_id, d);
+      });
+  assert(!exe_bndl.empty());
+  std::vector<std::byte> bytes;
+  const sycl::device_image<sycl::bundle_state::executable> &img =
+      *(exe_bndl.begin());
+  bytes = img.ext_oneapi_get_backend_content();
+
+  auto ZeContext = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(ctxt);
+  auto ZeDevice = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(d);
+
+  ze_result_t status;
+  ze_module_desc_t moduleDesc = {
+      ZE_STRUCTURE_TYPE_MODULE_DESC,
+      nullptr,
+      ZE_MODULE_FORMAT_IL_SPIRV,
+      bytes.size(),
+      reinterpret_cast<unsigned char *>(bytes.data()),
+      nullptr,
+      nullptr};
+  ze_module_handle_t ZeModule;
+  status = zeModuleCreate(ZeContext, ZeDevice, &moduleDesc, &ZeModule, nullptr);
+  assert(status == ZE_RESULT_SUCCESS);
+
+  ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0,
+                                 "__sycl_kernel_iota"};
+  ze_kernel_handle_t ZeKernel;
+  status = zeKernelCreate(ZeModule, &kernelDesc, &ZeKernel);
+  assert(status == ZE_RESULT_SUCCESS);
+  sycl::kernel k_iota_twin =
+      sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>(
+          {sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero,
+                                    sycl::bundle_state::executable>({ZeModule},
+                                                                    ctxt),
+           ZeKernel},
+          ctxt);
+  int *ptr_twin = sycl::malloc_shared<int>(1, q);
+  *ptr_twin = 1;
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(ptr_twin);
+     cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota_twin);
+   }).wait();
+  assert(*ptr_twin == *ptr);
+  sycl::free(ptr, q);
+  sycl::free(ptr_twin, q);
+#endif
+}

sycl/test-e2e/DeviceImageBackendContent/OCL_interop_test.cpp

@@ -0,0 +1,83 @@
+// REQUIRES: opencl, opencl_icd, aspect-usm_shared_allocations
+// RUN: %{build} %opencl_lib -fno-sycl-dead-args-optimization -o %t.out
+// RUN: %{run} %t.out
+//
+#include <sycl/backend.hpp>
+#include <sycl/detail/cl.h>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/usm.hpp>
+#include <vector>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (syclexp::nd_range_kernel<1>)) void iota(int *ptr) {
+  size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
+  ptr[id] = 42;
+}
+
+int main() {
+  sycl::device d([](const sycl::device &d) {
+    return d.get_backend() == sycl::backend::opencl;
+  });
+  sycl::queue q{d};
+  sycl::context ctxt = q.get_context();
+
+  // The following ifndef is required due to a number of limitations of free
+  // function kernels. See CMPLRLLVM-61498.
+  // TODO: Remove it once these limitations are no longer there.
+#ifndef __SYCL_DEVICE_ONLY__
+  // First, run the kernel using the SYCL API.
+  auto Bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt);
+  sycl::kernel_id iota_id = syclexp::get_kernel_id<iota>();
+  sycl::kernel k_iota = Bundle.get_kernel(iota_id);
+
+  int *ptr = sycl::malloc_shared<int>(1, q);
+  *ptr = 0;
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(ptr);
+     cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota);
+   }).wait();
+  // Now, run the kernel by first getting its image as an executable,
+  // making an OCL kernel out of it and then making a SYCL kernel out of
+  // the OCL kernel. Run this kernel on the SYCL API and verify
+  // that it has the same result as the kernel that was run directly on SYCL
+  // API. First, get a kernel bundle that contains the kernel "iota".
+  auto exe_bndl = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+      ctxt, {d},
+      [&](const sycl::device_image<sycl::bundle_state::executable> &img) {
+        return img.has_kernel(iota_id, d);
+      });
+  assert(!exe_bndl.empty());
+  std::vector<std::byte> bytes;
+  const sycl::device_image<sycl::bundle_state::executable> &img =
+      *(exe_bndl.begin());
+  bytes = img.ext_oneapi_get_backend_content();
+
+  auto clContext = sycl::get_native<sycl::backend::opencl>(ctxt);
+  auto clDevice = sycl::get_native<sycl::backend::opencl>(d);
+
+  cl_int status;
+  auto clProgram = clCreateProgramWithIL(
+      clContext, reinterpret_cast<unsigned char *>(bytes.data()), bytes.size(),
+      &status);
+  assert(status == CL_SUCCESS);
+  status = clBuildProgram(clProgram, 1, &clDevice, "", nullptr, nullptr);
+  assert(status == CL_SUCCESS);
+  auto clKernel = clCreateKernel(clProgram, "__sycl_kernel_iota", &status);
+  assert(status == CL_SUCCESS);
+  sycl::kernel k_iota_twin =
+      sycl::make_kernel<sycl::backend::opencl>(clKernel, ctxt);
+  int *ptr_twin = sycl::malloc_shared<int>(1, q);
+  *ptr_twin = 1;
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(ptr_twin);
+     cgh.parallel_for(sycl::nd_range{{1}, {1}}, k_iota_twin);
+   }).wait();
+  assert(*ptr_twin == *ptr);
+  sycl::free(ptr, q);
+  sycl::free(ptr_twin, q);
+#endif
+}