intel · againull · Apr 18, 2023 · Mar 10, 2023 · Mar 10, 2023 · Mar 10, 2023
@@ -17,6 +17,7 @@ NOTE: By necessity, this specification exposes some details about the way SYCL i
 |2|Added support for the make_buffer() API.
 |3|Added device member to backend_input_t<backend::ext_oneapi_level_zero, queue>.
 |4|Change the definition of backend_input_t and backend_return_t for the queue object, which changes the API for make_queue and get_native (when applied to queue).
+|5|Added support for make_image() API.
 
 NOTE: This extension is following SYCL 2020 backend specification. Prior API for interoperability with Level-Zero is marked
       as deprecated and will be removed in the next release.
@@ -43,15 +44,15 @@ There are multiple ways in which the Level-Zero backend can be selected by the u
 
 ### 3.1 Through an environment variable
 
-The SYCL_DEVICE_FILTER environment variable limits the SYCL runtime to use only a subset of the system's devices.
-By using ```level_zero``` for backend in SYCL_DEVICE_FILTER you can select the use of Level-Zero as a SYCL backend.
+The ONEAPI_DEVICE_SELECTOR environment variable limits the SYCL runtime to use only a subset of the system's devices.
+By using ```level_zero``` for backend in ONEAPI_DEVICE_SELECTOR you can select the use of Level-Zero as a SYCL backend.
 For further details see here: <https://github.com/intel/llvm/blob/sycl/sycl/doc/EnvironmentVariables.md>.
 
 ### 3.2 Through a programming API
 
 There is an extension that introduces a filtering device selection to SYCL described in
 [sycl\_ext\_oneapi\_filter\_selector](../supported/sycl_ext_oneapi_filter_selector.asciidoc).
-Similar to how SYCL_DEVICE_FILTER applies filtering to the entire process this device selector can be used to
+Similar to how SYCL_DEVICE_FILTER or ONEAPI_DEVICE_SELECTOR applies filtering to the entire process this device selector can be used to
 programmatically select the Level-Zero backend.
 
 When neither the environment variable nor the filtering device selector are used, the implementation chooses
@@ -247,6 +248,28 @@ struct {
 ```
 </td>
 </tr>
+<tr>
+<td>image</td>
+<td>
+
+``` C++
+ze_image_handle_t
+```
+</td>
+<td>
+
+``` C++
+struct {
+    ze_image_handle_t ZeImageHandle;
+    sycl::image_channel_order ChanOrder;
+    sycl::image_channel_type ChanType;
+    range<Dimensions> Range;
+    ext::oneapi::level_zero::ownership Ownership{
+        ext::oneapi::level_zero::ownership::transfer};
+  }
+```
+</td>
+</tr>
 </table>
 
 ### 4.2 Obtaining of native Level-Zero handles from SYCL objects
@@ -264,7 +287,7 @@ It is currently supported for SYCL ```platform```, ```device```, ```context```,
 The ```get_native(queue)``` function returns either ```ze_command_queue_handle_t``` or ```ze_command_list_handle_t``` depending on the manner in which the input argument ```queue``` had been created. Queues created with the SYCL ```queue``` constructors have a default setting for whether they use command queues or command lists. The default and how it may be changed is documented in the description for the environment variable ```SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS```. Queues created using ```make_queue()``` use either a command list or command queue depending on the input argument to ```make_queue``` and are not affected by the default for SYCL queues or the environment variable.
 
 The ```sycl::get_native<backend::ext_oneapi_level_zero>```
-free-function is not supported for SYCL ```buffer``` class. The native backend object associated with the
+free-function is not supported for SYCL ```buffer``` or ```image``` class. The native backend object associated with the
 buffer can be obtained using interop_hande class as described in the core SYCL specification section
 4.10.2, "Class interop_handle". 
 The pointer returned by ```get_native_mem<backend::ext_oneapi_level_zero>``` method of the ```interop_handle```
@@ -433,6 +456,67 @@ Construct a SYCL buffer instance from a pointer to a Level Zero memory allocatio
 description above for semantics and restrictions.
 The additional <code>AvailableEvent</code> argument must be a valid SYCL event. The instance of the SYCL buffer class template being constructed must wait for the SYCL event parameter to signal that the memory native handle is ready to be used.
 </tr>
+
+<tr>
+<td>
+
+``` C++
+make_image<backend::ext_oneapi_level_zero, Dims>(
+    const backend_input_t<backend::ext_oneapi_level_zero,
+                          image<Dimensions, AllocatorT>> &,
+    const context &Context)
+```
+</td>
+<td>This API is available starting with revision 4 of this specification.
+
+Construct a SYCL image instance from a ze_image_handle_t. 
+
+Because LevelZero has no way of getting image information from an image, it must be provided.  The  <code>backend_input_t</code> is a struct type like so:
+``` C++
+struct type {
+    ze_image_handle_t ZeImageHandle;
+    sycl::image_channel_order ChanOrder;
+    sycl::image_channel_type ChanType;
+    sycl::range<Dimensions> Range;
+    ext::oneapi::level_zero::ownership Ownership{
+        ext::oneapi::level_zero::ownership::transfer};
+  };
+```
+where the Range should be ordered (width), (width, height), or (width, height, depth) for 1D, 2D and 3D images respectively, 
+with those values matching the dimensions used in the `ze_image_desc` that was used to create the `ze_image_handle_t` initially. 
+Note that the range term ordering (width first, depth last) is true for SYCL 1.2.1 images that are supported here. But future classes like
+sampled_image and unsampled_image might have a different ordering.
+
+Example Usage
+``` C++
+sycl::backend_input_t<BE, sycl::image<2>> ImageInteropInput{ ZeHImage, ChanOrder, ChanType, ImgRange_2D, sycl::ext::oneapi::level_zero::ownership::transfer };      
+
+sycl::image<2> Image_2D  = sycl::make_image<BE, 2>(ImageInteropInput, Context);
+```
+
+ The input SYCL context <code>Context</code> must be associated with a single device, matching the device used to create the Level Zero image handle.
+The <code>Context</code> argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The created SYCL image can only be accessed from kernels that are submitted to a queue using this same context.
+The <code>Ownership</code> input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details. If the behavior is "transfer" then the SYCL runtime is going to free the input Level-Zero memory allocation, meaning the memory will be freed when the ~image destructor fires. When using "transfer" the ~image destructor may not need to block.  If the behavior is "keep", then the memory will not be freed by the ~image destructor, and the ~image destructor blocks until all work in the queues on the image have been completed. When using "keep" it is the responsibility of the caller to free the memory appropriately. 
+</td>
+</tr>
+
+<tr>
+<td>
+
+``` C++
+make_image(
+    const backend_input_t<backend::ext_oneapi_level_zero,
+                          image<Dimensions, AllocatorT>> &,
+    const context &Context, event AvailableEvent)
+```
+</td>
+<td>This API is available starting with revision 4 of this specification.
+
+Construct a SYCL image instance from a pointer to a Level Zero memory allocation. Please refer to <code>make_image</code>
+description above for semantics and restrictions.
+The additional <code>AvailableEvent</code> argument must be a valid SYCL event. The instance of the SYCL image class template being constructed must wait for the SYCL event parameter to signal that the memory native handle is ready to be used.
+</td>
+</tr>
 </table>
 
 NOTE: We shall consider adding other interoperability as needed, if possible.
@@ -509,4 +593,5 @@ The behavior of the SYCL buffer destructor depends on the Ownership flag. As wit
 |8|2022-01-06|Artur Gainullin|Introduced make_buffer() API
 |9|2022-05-12|Steffen Larsen|Added device member to queue input type
 |10|2022-08-18|Sergey Maslov|Moved free_memory device info query to be sycl_ext_intel_device_info extension
-|10|2023-03-14|Rajiv Deodhar|Added support for Level Zero immediate command lists
+|11|2023-03-14|Rajiv Deodhar|Added support for Level Zero immediate command lists
+|12|2023-04-06|Chris Perkins|Introduced make_image() API
@@ -14,6 +14,7 @@
 #include <sycl/context.hpp>
 #include <sycl/detail/backend_traits.hpp>
 #include <sycl/feature_test.hpp>
+#include <sycl/image.hpp>
 #if SYCL_BACKEND_OPENCL
 #include <sycl/detail/backend_traits_opencl.hpp>
 #endif
@@ -335,6 +336,20 @@ make_buffer(const typename backend_traits<Backend>::template input_type<
       AvailableEvent);
 }
 
+template <backend Backend, int Dimensions = 1,
+          typename AllocatorT = image_allocator>
+typename std::enable_if<detail::InteropFeatureSupportMap<Backend>::MakeImage ==
+                                true &&
+                            Backend != backend::ext_oneapi_level_zero,
+                        image<Dimensions, AllocatorT>>::type
+make_image(const typename backend_traits<Backend>::template input_type<
+               image<Dimensions, AllocatorT>> &BackendObject,
+           const context &TargetContext, event AvailableEvent = {}) {
+  return image<Dimensions, AllocatorT>(
+      detail::pi::cast<pi_native_handle>(BackendObject), TargetContext,
+      AvailableEvent);
+}
+
 template <backend Backend>
 kernel
 make_kernel(const typename backend_traits<Backend>::template input_type<kernel>

@@ -29,6 +29,8 @@ template <backend Backend> struct InteropFeatureSupportMap {
   static constexpr bool MakeEvent = false;
   static constexpr bool MakeBuffer = false;
   static constexpr bool MakeKernel = false;
+  static constexpr bool MakeKernelBundle = false;
+  static constexpr bool MakeImage = false;
 };
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)

@@ -165,6 +165,27 @@ struct BackendReturn<backend::ext_oneapi_level_zero,
   using type = void *;
 };
 
+template <int Dimensions, typename AllocatorT>
+struct BackendInput<backend::ext_oneapi_level_zero,
+                    image<Dimensions, AllocatorT>> {
+  // LevelZero has no way of getting image description FROM a ZeImageHandle so
+  // it must be provided.
+  struct type {
+    ze_image_handle_t ZeImageHandle;
+    sycl::image_channel_order ChanOrder;
+    sycl::image_channel_type ChanType;
+    range<Dimensions> Range;
+    ext::oneapi::level_zero::ownership Ownership{
+        ext::oneapi::level_zero::ownership::transfer};
+  };
+};
+
+template <int Dimensions, typename AllocatorT>
+struct BackendReturn<backend::ext_oneapi_level_zero,
+                     image<Dimensions, AllocatorT>> {
+  using type = ze_image_handle_t;
+};
+
 template <> struct BackendReturn<backend::ext_oneapi_level_zero, queue> {
   using type =
       std::variant<ze_command_queue_handle_t, ze_command_list_handle_t>;
@@ -214,6 +235,7 @@ template <> struct InteropFeatureSupportMap<backend::ext_oneapi_level_zero> {
   static constexpr bool MakeKernelBundle = true;
   static constexpr bool MakeKernel = true;
   static constexpr bool MakeBuffer = true;
+  static constexpr bool MakeImage = true;
 };
 
 } // namespace detail

@@ -155,6 +155,7 @@ template <> struct InteropFeatureSupportMap<backend::opencl> {
   static constexpr bool MakeBuffer = true;
   static constexpr bool MakeKernel = true;
   static constexpr bool MakeKernelBundle = true;
+  static constexpr bool MakeImage = false;
 };
 
 namespace pi {

@@ -61,6 +61,7 @@ _PI_API(piMemRelease)
 _PI_API(piMemBufferPartition)
 _PI_API(piextMemGetNativeHandle)
 _PI_API(piextMemCreateWithNativeHandle)
+_PI_API(piextMemImgCreateWithNativeHandle)
 // Program
 _PI_API(piProgramCreate)
 _PI_API(piclProgramCreateWithSource)

@@ -86,9 +86,11 @@
 // 12.27 Added new queue create and get APIs for immediate commandlists
 // piextQueueCreate2, piextQueueCreateWithNativeHandle2,
 // piextQueueGetNativeHandle2
+// 12.28 Added piextMemImgCreateWithNativeHandle for creating images from native
+// handles.
 
 #define _PI_H_VERSION_MAJOR 12
-#define _PI_H_VERSION_MINOR 27
+#define _PI_H_VERSION_MINOR 28
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -1308,6 +1310,19 @@ __SYCL_EXPORT pi_result piextMemCreateWithNativeHandle(
     pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
     pi_mem *mem);
 
+/// Creates PI image object from a native handle.
+/// NOTE: The created PI object takes ownership of the native handle.
+///
+/// \param nativeHandle is the native handle to create PI image from.
+/// \param context The PI context of the memory allocation.
+/// \param ownNativeHandle Indicates if we own the native memory handle or it
+/// came from interop that asked to not transfer the ownership to SYCL RT.
+/// \param img is the PI img created from the native handle.
+__SYCL_EXPORT pi_result piextMemImgCreateWithNativeHandle(
+    pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
+    const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc,
+    pi_mem *img);
+
 //
 // Program
 //

@@ -229,6 +229,24 @@ make_buffer(
       !(BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep));
 }
 
+// Specialization of sycl::make_image for Level-Zero backend.
+template <backend Backend, int Dimensions = 1,
+          typename AllocatorT = image_allocator>
+typename std::enable_if<Backend == backend::ext_oneapi_level_zero,
+                        image<Dimensions, AllocatorT>>::type
+make_image(const backend_input_t<Backend, image<Dimensions, AllocatorT>>
+               &BackendObject,
+           const context &TargetContext, event AvailableEvent) {
+
+  bool OwnNativeHandle =
+      (BackendObject.Ownership == ext::oneapi::level_zero::ownership::transfer);
+
+  return image<Dimensions, AllocatorT>(
+      detail::pi::cast<pi_native_handle>(BackendObject.ZeImageHandle),
+      TargetContext, AvailableEvent, BackendObject.ChanOrder,
+      BackendObject.ChanType, OwnNativeHandle, BackendObject.Range);
+}
+
 namespace __SYCL2020_DEPRECATED("use 'ext::oneapi::level_zero' instead")
     level_zero {
 using namespace ext::oneapi::level_zero;

@@ -120,6 +120,7 @@ template <> struct InteropFeatureSupportMap<backend::ext_oneapi_cuda> {
   static constexpr bool MakeBuffer = false;
   static constexpr bool MakeKernel = false;
   static constexpr bool MakeKernelBundle = false;
+  static constexpr bool MakeImage = false;
 };
 
 } // namespace detail

@@ -22,8 +22,18 @@
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 
+// forward declarations
 class handler;
 
+template <int D, typename A> class image;
+
+// 'friend'
+template <backend Backend, int D, typename A>
+typename std::enable_if<Backend == backend::ext_oneapi_level_zero,
+                        image<D, A>>::type
+make_image(const backend_input_t<Backend, image<D, A>> &BackendObject,
+           const context &TargetContext, event AvailableEvent = {});
+
 enum class image_channel_order : unsigned int {
   a = 0,
   r = 1,
@@ -128,6 +138,13 @@ class __SYCL_EXPORT image_plain {
               uint8_t Dimensions);
 #endif
 
+  image_plain(pi_native_handle MemObject, const context &SyclContext,
+              event AvailableEvent,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              uint8_t Dimensions, image_channel_order Order,
+              image_channel_type Type, bool OwnNativeHandle,
+              range<3> Range3WithOnes);
+
   template <typename propertyT> bool has_property() const noexcept;
 
   template <typename propertyT> propertyT get_property() const;
@@ -467,6 +484,15 @@ class image : public detail::image_plain {
   void set_write_back(bool flag = true) { image_plain::set_write_back(flag); }
 
 private:
+  image(pi_native_handle MemObject, const context &SyclContext,
+        event AvailableEvent, image_channel_order Order,
+        image_channel_type Type, bool OwnNativeHandle, range<Dimensions> Range)
+      : image_plain(MemObject, SyclContext, AvailableEvent,
+                    make_unique_ptr<
+                        detail::SYCLMemObjAllocatorHolder<AllocatorT, byte>>(),
+                    Dimensions, Order, Type, OwnNativeHandle,
+                    detail::convertToArrayOfN<3, 1>(Range)) {}
+
   // This utility api is currently used by accessor to get the element size of
   // the image. Element size is dependent on num of channels and channel type.
   // This information is not accessible from the image using any public API.
@@ -484,6 +510,23 @@ class image : public detail::image_plain {
     return image_plain::getChannelType();
   }
 
+  // Declare make_image as a friend function
+  template <backend Backend, int D, typename A>
+  friend typename std::enable_if<
+      detail::InteropFeatureSupportMap<Backend>::MakeImage == true &&
+          Backend != backend::ext_oneapi_level_zero,
+      image<D, A>>::type
+  make_image(
+      const typename backend_traits<Backend>::template input_type<image<D, A>>
+          &BackendObject,
+      const context &TargetContext, event AvailableEvent);
+
+  template <backend Backend, int D, typename A>
+  friend typename std::enable_if<Backend == backend::ext_oneapi_level_zero,
+                                 image<D, A>>::type
+  make_image(const backend_input_t<Backend, image<D, A>> &BackendObject,
+             const context &TargetContext, event AvailableEvent);
+
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);