Merge from 'sycl' to 'sycl-web' (2 commits)

Author: iclsrc

sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md

@@ -13,6 +13,7 @@ This extension provides a feature-test macro as described in the core SYCL speci
 |Value|Description|
 |---|:---|
 |1|Initial extension version.
+|2|Added support for the make_buffer() 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.
@@ -46,7 +47,6 @@ For further details see here: <https://github.com/intel/llvm/blob/sycl/sycl/doc/
         
 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
 programmatically select the Level-Zero backend.
                 
@@ -170,6 +170,24 @@ struct {
 }
 ```
 </td>
+</tr><tr>
+<td>buffer</td>
+<td>
+
+``` C++
+void *
+```
+</td>
+<td>
+
+``` C++
+struct {
+  void *NativeHandle;
+  ext::oneapi::level_zero::ownership Ownership{
+      ext::oneapi::level_zero::ownership::transfer};
+}
+```
+</td>
 </tr>
 </table>
 
@@ -183,8 +201,31 @@ auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT>
 ```
 It is currently supported for SYCL ```platform```, ```device```, ```context```, ```queue```, ```event```,
-```kernel_bundle```, and ```kernel``` classes.
+```kernel_bundle```, and ```kernel``` classes. 
+
+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
+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```
+class is the value returned from a call to <code>zeMemAllocShared()</code>, <code>zeMemAllocDevice()</code>,
+or <code>zeMemAllocHost()</code> and not necessarily directly accessible from the host.  Users may need to copy
+data to the host to access the data. Users can get type of the allocation using ```type``` data member of the 
+```ze_memory_allocation_properties_t``` struct returned by ```zeMemGetAllocProperties```.
 
+``` C++
+    Queue.submit([&](handler &CGH) {
+        auto BufferAcc = Buffer.get_access<access::mode::write>(CGH);
+        CGH.host_task([=](const interop_handle &IH) {
+            void *DevicePtr =
+                IH.get_native_mem<backend::ext_oneapi_level_zero>(BufferAcc);
+            ze_memory_allocation_properties_t MemAllocProperties{};
+            ze_result_t Res = zeMemGetAllocProperties(
+                ZeContext, DevicePtr, &MemAllocProperties, nullptr);
+            ze_memory_type_t ZeMemType = MemAllocProperties.type;
+        });
+    }).wait();
+```
 ### 4.3 Construct a SYCL object from a Level-Zero handle
         
 The following free functions defined in the ```sycl``` namespace are specialized for Level-Zero backend to allow
@@ -294,6 +335,37 @@ 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 runtime is going to destroy the input
 Level-Zero kernel</td>
+</tr><tr>
+<td>
+
+``` C++
+make_buffer(
+    const backend_input_t<backend::ext_oneapi_level_zero,
+                          buffer<T, Dimensions, AllocatorT>> &,
+    const context &Context)
+```
+</td>
+<td>This API is available starting with revision 2 of this specification.
+
+Construct a SYCL buffer instance from a pointer to a Level Zero memory allocation. The pointer must be the value returned from a previous call to <code>zeMemAllocShared()</code>, <code>zeMemAllocDevice()</code>, or <code>zeMemAllocHost()</code>. The input SYCL context <code>Context</code> must be associated with a single device, matching the device used at the prior allocation.
+The <code>Context</code> argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero memory must be allocated on the same context. Created SYCL buffer can be accessed in another contexts, not only in the provided input 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 runtime is going to free the input Level-Zero memory allocation. 
+Synchronization rules for a buffer that is created with this API are described in Section 4.5</td>
+</tr><tr>
+<td>
+
+``` C++
+make_buffer(
+    const backend_input_t<backend::ext_oneapi_level_zero,
+                          buffer<T, Dimensions, AllocatorT>> &,
+    const context &Context, event AvailableEvent)
+```
+</td>
+<td>This API is available starting with revision 2 of this specification.
+
+Construct a SYCL buffer instance from a pointer to a Level Zero memory allocation. Please refer to <code>make_buffer</code>
+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>
 </table>
 
@@ -349,6 +421,15 @@ Applications must make sure that the Level-Zero handles themselves aren't used s
 Practically speaking, and taking into account that SYCL runtime takes ownership of the Level-Zero handles,
 the application should not attempt further direct use of those handles.
 
+### 4.5 Interoperability buffer synchronization rules
+
+A SYCL buffer that is constructed with this interop API uses the Level Zero memory allocation for its full lifetime, and the contents of the Level Zero memory allocation are unspecified for the lifetime of the SYCL buffer. If the application modifies the contents of that Level Zero memory allocation during the lifetime of the SYCL buffer, the behavior is undefined. The initial contents of the SYCL buffer will be the initial contents of the Level Zero memory allocation at the time of the SYCL buffer's construction.
+
+The behavior of the SYCL buffer destructor depends on the Ownership flag. As with other SYCL buffers, this behavior is triggered only when the last reference count to the buffer is dropped, as described in the core SYCL specification section 4.7.2.3, "Buffer synchronization rules".
+
+* If the ownership is keep (i.e. the application retains ownership of the Level Zero memory allocation), then the SYCL buffer destructor blocks until all work in queues on the buffer have completed. The buffer's contents is not copied back to the Level Zero memory allocation.
+* If the ownership is transfer (i.e. the SYCL runtime has ownership of the Level Zero memory allocation), then the SYCL buffer destructor does not need to block even if work on the buffer has not completed. The SYCL runtime frees the Level Zero memory allocation asynchronously when it is no longer in use in queues.
+
 ## 5 Level-Zero additional functionality
 
 ### 5.1 Device Information Descriptors
@@ -403,3 +484,4 @@ struct free_memory {
 |5|2021-07-25|Sergey Maslov|Introduced SYCL interop for events
 |6|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions
 |7|2021-09-13|Sergey Maslov|Updated according to SYCL 2020 standard
+|8|2022-01-06|Artur Gainullin|Introduced make_buffer() API

sycl/include/CL/sycl/backend.hpp

@@ -242,7 +242,8 @@ typename std::enable_if<
 template <backend Backend, typename T, int Dimensions = 1,
           typename AllocatorT = buffer_allocator>
 typename std::enable_if<detail::InteropFeatureSupportMap<Backend>::MakeBuffer ==
-                            true,
+                                true &&
+                            Backend != backend::ext_oneapi_level_zero,
                         buffer<T, Dimensions, AllocatorT>>::type
 make_buffer(const typename backend_traits<Backend>::template input_type<
                 buffer<T, Dimensions, AllocatorT>> &BackendObject,

sycl/include/CL/sycl/buffer.hpp

@@ -26,8 +26,10 @@ template <int dimensions> class range;
 namespace detail {
 template <typename T, int Dimensions, typename AllocatorT>
 buffer<T, Dimensions, AllocatorT, void>
-make_buffer_helper(pi_native_handle Handle, const context &Ctx, event Evt) {
-  return buffer<T, Dimensions, AllocatorT, void>(Handle, Ctx, Evt);
+make_buffer_helper(pi_native_handle Handle, const context &Ctx, event Evt = {},
+                   bool OwnNativeHandle = true) {
+  return buffer<T, Dimensions, AllocatorT, void>(Handle, Ctx, OwnNativeHandle,
+                                                 Evt);
 }
 } // namespace detail
 
@@ -314,14 +316,11 @@ class buffer {
          const detail::code_location CodeLoc = detail::code_location::current())
       : Range{0} {
 
-    size_t BufSize = detail::SYCLMemObjT::getBufSizeForContext(
-        detail::getSyclObjImpl(SyclContext), MemObject);
-
-    Range[0] = BufSize / sizeof(T);
     impl = std::make_shared<detail::buffer_impl>(
-        detail::pi::cast<pi_native_handle>(MemObject), SyclContext, BufSize,
+        detail::pi::cast<pi_native_handle>(MemObject), SyclContext,
         make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT>>(),
-        AvailableEvent);
+        /* OwnNativeHandle */ true, AvailableEvent);
+    Range[0] = impl->getSize() / sizeof(T);
     impl->constructorNotification(CodeLoc, (void *)impl.get(), &MemObject,
                                   (const void *)typeid(T).name(), dimensions,
                                   sizeof(T), rangeToArray(Range).data());
@@ -531,7 +530,7 @@ class buffer {
   friend class accessor;
   template <typename HT, int HDims, typename HAllocT>
   friend buffer<HT, HDims, HAllocT, void>
-  detail::make_buffer_helper(pi_native_handle, const context &, event);
+  detail::make_buffer_helper(pi_native_handle, const context &, event, bool);
   range<dimensions> Range;
   // Offset field specifies the origin of the sub buffer inside the parent
   // buffer
@@ -541,18 +540,15 @@ class buffer {
   // Interop constructor
   template <int N = dimensions, typename = EnableIfOneDimension<N>>
   buffer(pi_native_handle MemObject, const context &SyclContext,
-         event AvailableEvent = {},
+         bool OwnNativeHandle, event AvailableEvent = {},
          const detail::code_location CodeLoc = detail::code_location::current())
       : Range{0} {
 
-    size_t BufSize = detail::SYCLMemObjT::getBufSizeForContext(
-        detail::getSyclObjImpl(SyclContext), MemObject);
-
-    Range[0] = BufSize / sizeof(T);
     impl = std::make_shared<detail::buffer_impl>(
-        MemObject, SyclContext, BufSize,
+        MemObject, SyclContext,
         make_unique_ptr<detail::SYCLMemObjAllocatorHolder<AllocatorT>>(),
-        AvailableEvent);
+        OwnNativeHandle, AvailableEvent);
+    Range[0] = impl->getSize() / sizeof(T);
     impl->constructorNotification(CodeLoc, (void *)impl.get(), &MemObject,
                                   (const void *)typeid(T).name(), dimensions,
                                   sizeof(T), rangeToArray(Range).data());

sycl/include/CL/sycl/detail/backend_traits_level_zero.hpp

@@ -138,6 +138,22 @@ template <> struct BackendInput<backend::ext_oneapi_level_zero, queue> {
   };
 };
 
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendInput<backend::ext_oneapi_level_zero,
+                    buffer<DataT, Dimensions, AllocatorT>> {
+  struct type {
+    void *NativeHandle;
+    ext::oneapi::level_zero::ownership Ownership{
+        ext::oneapi::level_zero::ownership::transfer};
+  };
+};
+
+template <typename DataT, int Dimensions, typename AllocatorT>
+struct BackendReturn<backend::ext_oneapi_level_zero,
+                     buffer<DataT, Dimensions, AllocatorT>> {
+  using type = void *;
+};
+
 template <> struct BackendReturn<backend::ext_oneapi_level_zero, queue> {
   using type = ze_command_queue_handle_t;
 };
@@ -195,7 +211,7 @@ template <> struct InteropFeatureSupportMap<backend::ext_oneapi_level_zero> {
   static constexpr bool MakeEvent = true;
   static constexpr bool MakeKernelBundle = true;
   static constexpr bool MakeKernel = true;
-  static constexpr bool MakeBuffer = false;
+  static constexpr bool MakeBuffer = true;
 };
 
 } // namespace detail

sycl/include/CL/sycl/detail/buffer_impl.hpp

@@ -138,6 +138,20 @@ class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT {
     BaseT::handleHostData(First, Last, RequiredAlign);
   }
 
+  buffer_impl(cl_mem MemObject, const context &SyclContext,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              event AvailableEvent)
+      : buffer_impl(pi::cast<pi_native_handle>(MemObject), SyclContext,
+                    std::move(Allocator), /*OwnNativeHandle*/ true,
+                    std::move(AvailableEvent)) {}
+
+  buffer_impl(pi_native_handle MemObject, const context &SyclContext,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator,
+              bool OwnNativeHandle, event AvailableEvent)
+      : BaseT(MemObject, SyclContext, OwnNativeHandle,
+              std::move(AvailableEvent), std::move(Allocator)) {}
+
+  // TODO: remove the following 2 constructors when it is allowed to break ABI.
   buffer_impl(cl_mem MemObject, const context &SyclContext,
               const size_t SizeInBytes,
               std::unique_ptr<SYCLMemObjAllocator> Allocator,

sycl/include/CL/sycl/detail/pi.h

@@ -42,10 +42,12 @@
 //   piextEventCreateWithNativeHandle
 // 6.8 Added new ownership argument to piextProgramCreateWithNativeHandle. Added
 // piQueueFlush function.
+// 7.9 Added new context and ownership arguments to
+// piextMemCreateWithNativeHandle.
 //
 #include "CL/cl.h"
-#define _PI_H_VERSION_MAJOR 6
-#define _PI_H_VERSION_MINOR 8
+#define _PI_H_VERSION_MAJOR 7
+#define _PI_H_VERSION_MINOR 9
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -941,6 +943,14 @@ typedef struct {
 
 using pi_image_format = _pi_image_format;
 using pi_image_desc = _pi_image_desc;
+
+typedef enum {
+  PI_MEM_CONTEXT = CL_MEM_CONTEXT,
+  PI_MEM_SIZE = CL_MEM_SIZE
+} _pi_mem_info;
+
+using pi_mem_info = _pi_mem_info;
+
 //
 // Following section contains SYCL RT Plugin Interface (PI) functions.
 // They are 3 distinct categories:
@@ -1163,10 +1173,9 @@ __SYCL_EXPORT pi_result piMemImageCreate(pi_context context, pi_mem_flags flags,
                                          const pi_image_desc *image_desc,
                                          void *host_ptr, pi_mem *ret_mem);
 
-__SYCL_EXPORT pi_result piMemGetInfo(
-    pi_mem mem,
-    cl_mem_info param_name, // TODO: untie from OpenCL
-    size_t param_value_size, void *param_value, size_t *param_value_size_ret);
+__SYCL_EXPORT pi_result piMemGetInfo(pi_mem mem, pi_mem_info param_name,
+                                     size_t param_value_size, void *param_value,
+                                     size_t *param_value_size_ret);
 
 __SYCL_EXPORT pi_result piMemImageGetInfo(pi_mem image,
                                           pi_image_info param_name,
@@ -1193,9 +1202,13 @@ __SYCL_EXPORT pi_result piextMemGetNativeHandle(pi_mem mem,
 /// NOTE: The created PI object takes ownership of the native handle.
 ///
 /// \param nativeHandle is the native handle to create PI mem 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 mem is the PI mem created from the native handle.
-__SYCL_EXPORT pi_result
-piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, pi_mem *mem);
+__SYCL_EXPORT pi_result piextMemCreateWithNativeHandle(
+    pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
+    pi_mem *mem);
 
 //
 // Program
@@ -1619,7 +1632,7 @@ typedef enum {
   PI_MEM_ALLOC_BASE_PTR = CL_MEM_ALLOC_BASE_PTR_INTEL,
   PI_MEM_ALLOC_SIZE = CL_MEM_ALLOC_SIZE_INTEL,
   PI_MEM_ALLOC_DEVICE = CL_MEM_ALLOC_DEVICE_INTEL,
-} _pi_mem_info;
+} _pi_mem_alloc_info;
 
 typedef enum {
   PI_MEM_TYPE_UNKNOWN = CL_MEM_TYPE_UNKNOWN_INTEL,
@@ -1637,7 +1650,7 @@ typedef enum : pi_bitfield {
 
 using pi_usm_capability_query = _pi_usm_capability_query;
 using pi_usm_capabilities = _pi_usm_capabilities;
-using pi_mem_info = _pi_mem_info;
+using pi_mem_alloc_info = _pi_mem_alloc_info;
 using pi_usm_type = _pi_usm_type;
 using pi_usm_migration_flags = _pi_usm_migration_flags;
 
@@ -1766,7 +1779,7 @@ __SYCL_EXPORT pi_result piextUSMEnqueueMemAdvise(pi_queue queue,
 /// \param param_value is the result
 /// \param param_value_size_ret is how many bytes were written
 __SYCL_EXPORT pi_result piextUSMGetMemAllocInfo(
-    pi_context context, const void *ptr, pi_mem_info param_name,
+    pi_context context, const void *ptr, pi_mem_alloc_info param_name,
     size_t param_value_size, void *param_value, size_t *param_value_size_ret);
 
 /// API to get Plugin internal data, opaque to SYCL RT. Some devices whose

sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

@@ -90,6 +90,10 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
       : SYCLMemObjT(MemObject, SyclContext, /*SizeInBytes*/ 0, AvailableEvent,
                     std::move(Allocator)) {}
 
+  SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext,
+              bool OwmNativeHandle, event AvailableEvent,
+              std::unique_ptr<SYCLMemObjAllocator> Allocator);
+
   virtual ~SYCLMemObjT() = default;
 
   const plugin &getPlugin() const;
@@ -336,10 +340,9 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
   EventImplPtr MInteropEvent;
   // Context passed by user to interoperability constructor.
   ContextImplPtr MInteropContext;
-  // OpenCL's memory object handle passed by user to interoperability
+  // Native backend memory object handle passed by user to interoperability
   // constructor.
-  // TODO update this member to support other backends.
-  cl_mem MInteropMemObject;
+  RT::PiMem MInteropMemObject;
   // Indicates whether memory object is created using interoperability
   // constructor or not.
   bool MOpenCLInterop;

sycl/include/CL/sycl/feature_test.hpp.in

@@ -69,7 +69,7 @@ namespace sycl {
 #define SYCL_EXT_INTEL_MEM_CHANNEL_PROPERTY 1
 #define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 1
 #define SYCL_EXT_INTEL_RUNTIME_BUFFER_LOCATION 1
-#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 1
+#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 2
 #define SYCL_EXT_ONEAPI_USM_DEVICE_READ_ONLY 1
 #cmakedefine01 SYCL_BUILD_PI_CUDA
 #if SYCL_BUILD_PI_CUDA