[SYCL] native image handle support for LevelZero. (#8603)

New interop support for images for LevelZero. Includes make_image and
interop_handle::get_native_mem.

Tests are present in this PR.

---------

Signed-off-by: Chris Perkins <[email protected]>
Co-authored-by: Steffen Larsen <[email protected]>

Author: cperkinsintel

sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md

@@ -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;
+    sycl::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```
@@ -299,39 +322,39 @@ an application to create a SYCL object that encapsulates a corresponding Level-Z
 <td>
 
 ``` C++
-make_platform<backend::ext_oneapi_level_zero>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero, platform> &)
+template <backend Backend>
+platform make_platform(
+    const backend_input_t<Backend, platform> &)
 ```
 </td>
 <td>Constructs a SYCL platform instance from a Level-Zero <code>ze_driver_handle_t</code>. The SYCL execution environment contains a fixed number of platforms that are enumerated via <code>sycl::platform::get_platforms()</code>. Calling this function does not create a new platform. Rather it merely creates a <code>sycl::platform</code> object that is a copy of one of the platforms from that enumeration.</td>
 </tr><tr>
 <td>
 
 ``` C++
-make_device<backend::ext_oneapi_level_zero>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero, device> &)
+template <backend Backend>
+device make_device(
+    const backend_input_t<Backend, device> &)
 ```
 </td>
 <td>Constructs a SYCL device instance from a Level-Zero <code>ze_device_handle_t</code>. The SYCL execution environment for the Level Zero backend contains a fixed number of devices that are enumerated via <code>sycl::device::get_devices()</code> and a fixed number of sub-devices that are enumerated via <code>sycl::device::create_sub_devices(...)</code>. Calling this function does not create a new device. Rather it merely creates a <code>sycl::device</code> object that is a copy of one of the devices from those enumerations.</td>
 </tr><tr>
 <td>
 
 ``` C++
-make_context<backend::ext_oneapi_level_zero>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero, context> &)
+template <backend Backend>
+context make_context(
+    const backend_input_t<Backend, context> &)
 ```
 </td>
 <td>Constructs a SYCL context instance from a Level-Zero <code>ze_context_handle_t</code>. The context is created against the devices passed in <code>DeviceList</code> structure member. There must be at least one device given and all the devices must be from the same SYCL platform and thus from the same Level-Zero driver. 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.</td>
 </tr><tr>
 <td>
 
 ``` C++
-make_queue<backend::ext_oneapi_level_zero>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero, queue> &,
+template <backend Backend>
+queue make_queue(
+    const backend_input_t<Backend, queue> &,
     const context &Context)
 ```
 </td>
@@ -346,9 +369,9 @@ the ```compute_index``` property which is built into the command queue or comman
 <td>
 
 ``` C++
-make_event<backend::ext_oneapi_level_zero>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero, event> &,
+template <backend Backend>
+event make_event(
+    const backend_input_t<Backend, event> &,
     const context &Context)
 ```
 </td>
@@ -357,11 +380,11 @@ make_event<backend::ext_oneapi_level_zero>(
 <td>
 
 ``` C++
-make_kernel_bundle<backend::ext_oneapi_level_zero,
-                   bundle_state::executable>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero,
-        kernel_bundle<bundle_state::executable>> &,
+// State must be bundle_state::executable
+template <backend Backend, bundle_state State>
+kernel_bundle<State> make_kernel_bundle(
+    const backend_input_t<Backend,
+        kernel_bundle<State>> &,
     const context &Context)
 ```
 </td>
@@ -383,9 +406,9 @@ interoperability <code>kernel_bundle</code> destructor is called.</td>
 <td>
 
 ``` C++
-make_kernel<backend::ext_oneapi_level_zero>(
-    const backend_input_t<
-        backend::ext_oneapi_level_zero, kernel> &,
+template <backend Backend>
+kernel make_kernel(
+    const backend_input_t<Backend, kernel> &,
     const context &Context)
 ```
 </td>
@@ -405,9 +428,15 @@ Level-Zero kernel</td>
 <td>
 
 ``` C++
-make_buffer(
-    const backend_input_t<backend::ext_oneapi_level_zero,
-                          buffer<T, Dimensions, AllocatorT>> &,
+template <backend Backend, 
+          typename T, int Dimensions = 1,
+          typename AllocatorT = 
+            buffer_allocator<std::remove_const_t<T>>>
+buffer<T, Dimensions, AllocatorT> make_buffer(
+    const backend_input_t<Backend,
+                          buffer<T, 
+                                 Dimensions, 
+                                 AllocatorT>> &,
     const context &Context)
 ```
 </td>
@@ -421,9 +450,15 @@ Synchronization rules for a buffer that is created with this API are described i
 <td>
 
 ``` C++
-make_buffer(
-    const backend_input_t<backend::ext_oneapi_level_zero,
-                          buffer<T, Dimensions, AllocatorT>> &,
+template <backend Backend, 
+          typename T, int Dimensions = 1,
+          typename AllocatorT = 
+            buffer_allocator<std::remove_const_t<T>>>
+buffer<T, Dimensions, AllocatorT> make_buffer(
+    const backend_input_t<Backend,
+                          buffer<T, 
+                                 Dimensions, 
+                                 AllocatorT>> &,
     const context &Context, event AvailableEvent)
 ```
 </td>
@@ -433,6 +468,94 @@ 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++
+template<backend Backend, int Dimensions = 1, 
+         typename AllocrT = sycl::image_allocator>
+image<Dimensions, AllocrT> make_image(
+    const backend_input_t<Backend, 
+                          image<Dimensions, 
+                            AllocrT>> &backendObject,
+    const context &targetContext);
+```
+</td>
+<td>This API is available starting with revision 5 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++
+ze_image_handle_t ZeHImage; 
+// ... user provided LevelZero ZeHImage image 
+// handle gotten somehow (possibly zeImageCreate)
+
+// the informational data that matches ZeHImage
+sycl::image_channel_order ChanOrder 
+     = sycl::image_channel_order::rgba;
+sycl::image_channel_type ChanType 
+     = sycl::image_channel_type::unsigned_int8;
+size_t width  = 4;
+size_t height = 2;
+sycl::range<2> ImgRange_2D(width, height);
+
+constexpr sycl::backend BE 
+       = sycl::backend::ext_oneapi_level_zero;
+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 image can only be used on the single device where it was created. This limitation may be relaxed in the future.
+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++
+template<backend Backend, int Dimensions = 1, 
+         typename AllocrT = sycl::image_allocator>
+image<Dimensions, AllocrT> make_image(
+    const backend_input_t<Backend, 
+                          image<Dimensions, 
+                            AllocrT>> &backendObject,
+    const context &targetContext, event availableEvent);
+```
+</td>
+<td>This API is available starting with revision 5 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 +632,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

sycl/include/sycl/backend.hpp

@@ -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>

sycl/include/sycl/detail/backend_traits.hpp

@@ -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)

sycl/include/sycl/detail/backend_traits_level_zero.hpp

@@ -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

sycl/include/sycl/detail/backend_traits_opencl.hpp

@@ -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 {

sycl/include/sycl/detail/pi.def

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