[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.*
 

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> // for span
+#endif
 #include <string>      // for string
 #include <type_traits> // for enable_if_t, remove_refer...
 #include <utility>     // for move
@@ -114,6 +117,17 @@ class __SYCL_EXPORT device_image_plain {
 
   ur_native_handle_t getNative() const;
 
+  backend ext_oneapi_get_backend() const noexcept;
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+  std::vector<std::byte> ext_oneapi_get_backend_content() const;
+
+#ifdef __cpp_lib_span
+  std::span<std::byte> ext_oneapi_get_backend_content_view() const;
+#endif // __cpp_lib_span
+
+#endif // HAS_STD_BYTE
+
 protected:
   detail::DeviceImageImplPtr impl;
 
@@ -145,6 +159,26 @@ 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();
+  }
+
+#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();
+  }
+
+#ifdef __cpp_lib_span
+  template <sycl::bundle_state T = State,
+            typename = std::enable_if_t<T == bundle_state::executable>>
+  std::span<std::byte> ext_oneapi_get_content_backend_view() const {
+    return device_image_plain::ext_oneapi_get_backend_content_view();
+  }
+#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

@@ -112,6 +112,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_WORK_GROUP_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 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,28 @@ ur_native_handle_t device_image_plain::getNative() const {
   return impl->getNative();
 }
 
+backend device_image_plain::ext_oneapi_get_backend() const noexcept {
+  return impl->get_context().get_backend();
+}
+
+std::vector<std::byte>
+device_image_plain::ext_oneapi_get_backend_content() 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));
+}
+
+#ifdef __cpp_lib_span
+std::span<std::byte>
+device_image_plain::ext_oneapi_get_backend_content_view() const {
+  return std::span(reinterpret_cast<const std::byte *>(
+                       impl->get_bin_image_ref()->getRawData().BinaryStart),
+                   reinterpret_cast<const std::byte *>(
+                       impl->get_bin_image_ref()->getRawData().BinaryEnd));
+}
+#endif
+
 ////////////////////////////
 ///// kernel_bundle_plain
 ///////////////////////////

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/basic_test.cpp

@@ -0,0 +1,78 @@
+// RUN: %{build} -fsyntax-only -DTEST_API_VIOLATION=1 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#ifdef __cpp_lib_span
+#include <span>
+#endif
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <type_traits>
+
+class kernel;
+
+void define_kernel(sycl::queue &q) {
+  int data;
+  sycl::buffer<int> data_buf(&data, 1);
+  q.submit([&](sycl::handler &cgh) {
+    sycl::accessor data_acc(data_buf, cgh);
+    cgh.parallel_for<class kernel>(
+        sycl::nd_range{{1}, {1}},
+        [=](sycl::nd_item<> it) { data_acc[0] = 42; });
+  });
+}
+
+int main() {
+  sycl::device d;
+  sycl::queue q{d};
+  sycl::context ctxt = q.get_context();
+  sycl::kernel_id id = sycl::get_kernel_id<kernel>();
+  auto bundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctxt, {id});
+  assert(!bundle.empty());
+  sycl::backend backend;
+  std::vector<std::byte> bytes;
+#ifdef __cpp_lib_span
+  std::span<std::byte> bytes_view;
+#endif
+  for (const auto &img : bundle) {
+    // Check that all 3 functions of the api return correct types and compile.
+    // Furthermore, check that the backend corresponds to the backend of the
+    // bundle Check that the view of the content is indeed equal to the
+    // content.
+    static_assert(std::is_same_v<decltype(img.ext_oneapi_get_backend()),
+                                 decltype(backend)>);
+    static_assert(std::is_same_v<decltype(img.ext_oneapi_get_backend_content()),
+                                 decltype(bytes)>);
+    backend = img.ext_oneapi_get_backend();
+    assert(backend == bundle.get_backend());
+    bytes = img.ext_oneapi_get_backend_content();
+#ifdef __cpp_lib_span
+    static_assert(
+        std ::is_same_v<decltype(img.ext_oneapi.get_backend_content_view()),
+                        decltype(bytes_view)>);
+    bytes_view = img.ext_oneapi_get_backend_content_view();
+    assert(bytes_view.size() == bytes.size());
+    for (size_t i = 0; i < bytes.size(); ++i) {
+      assert(bytes[i] == bytes_view[i]);
+    }
+#endif
+  }
+
+#ifdef TEST_API_VIOLATION
+  // Check that the ext_oneapi_get_backend_content and the
+  // ext_oneapi_get_backend_content_view of the content functions are not
+  // available
+  // when the image is not in the executable state.
+
+  auto input_bundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::input>(ctxt, {id});
+  // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content'}}
+  bytes = (*input_bundle.begin()).ext_oneapi_get_backend_content();
+#ifdef _cpp_lib_span
+  // expected-error@+1 {{no matching member function for call to 'ext_oneapi_get_backend_content_view'}}
+  bytes_view = (*input_bundle.begin()).ext_oneapi_get_backend_content_view();
+#endif // __cpp_lib_span
+#endif // TEST_API_VIOLATION
+  return 0;
+}

sycl/test/abi/sycl_symbols_linux.dump

@@ -3524,14 +3524,14 @@ _ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail30UnsampledImageAccessorB
 _ZN4sycl3_V17handler20memcpyToDeviceGlobalEPKvS3_bmm
 _ZN4sycl3_V17handler20setKernelCacheConfigENS1_23StableKernelCacheConfigE
 _ZN4sycl3_V17handler20setStateSpecConstSetEv
+_ZN4sycl3_V17handler21setKernelWorkGroupMemEm
 _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE
 _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm
 _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm
 _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE
 _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE
 _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi
 _ZN4sycl3_V17handler22setKernelIsCooperativeEb
-_ZN4sycl3_V17handler21setKernelWorkGroupMemEm
 _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_
 _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb
 _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm
@@ -3737,6 +3737,8 @@ _ZNK4sycl3_V16detail16AccessorBaseHost6getPtrEv
 _ZNK4sycl3_V16detail16AccessorBaseHost9getOffsetEv
 _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idE
 _ZNK4sycl3_V16detail18device_image_plain10has_kernelERKNS0_9kernel_idERKNS0_6deviceE
+_ZNK4sycl3_V16detail18device_image_plain22ext_oneapi_get_backendEv
+_ZNK4sycl3_V16detail18device_image_plain30ext_oneapi_get_backend_contentEv
 _ZNK4sycl3_V16detail18device_image_plain9getNativeEv
 _ZNK4sycl3_V16detail19kernel_bundle_plain10get_kernelERKNS0_9kernel_idE
 _ZNK4sycl3_V16detail19kernel_bundle_plain10has_kernelERKNS0_9kernel_idE