[SYCL] [L0] Add user APIs to import / release host memory from USM. (#9294)

This change adds a SYCL interface to the Level Zero APIs
zexDriverImportExternalPointer and zexDriverReleaseImportedPointer.
These functions are used for importing host memory into USM for the
duration of data transfer to increase bandwidth.

Author: rdeodhar

sycl/doc/extensions/experimental/sycl_ext_oneapi_copy_optimize.asciidoc

@@ -0,0 +1,159 @@
+= sycl_ext_oneapi_copy_optimize
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2023-2023 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 7 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+
+== Status
+
+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 DPC\++ but they are not finalized and may
+change incompatibly in future versions of DPC++ without prior notice.
+Shipping software products should not rely on APIs defined in this
+specification.
+
+
+== Overview
+
+SYCL provides explicit copy APIs that copy memory between a device and the host.
+Some backends can optimize these copy operations if the backend knows in
+advance that the host memory will be used for such a copy operation.
+If a SYCL application knows that it will frequently copy to or from the same
+host memory, it can use the APIs in this extension to inform the backend
+about this, which can result in a performance benefit.
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine
+the macro `SYCL_EXT_ONEAPI_COPY_OPTIMIZE` to one of the values defined
+in the table below.  Applications can test for the existence of this macro
+to determine if the implementation supports this feature, or applications
+can test the macro's value to determine which of the extension's features
+the implementation supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+=== API of the extension
+
+This extension adds the following free functions:
+
+```c++
+namespace sycl::ext::oneapi::experimental {
+
+void prepare_for_device_copy(const void *Ptr, size_t NumBytes,
+                             const context &Context);
+void prepare_for_device_copy(const void *Ptr, size_t NumBytes,
+                             const queue &Queue);
+
+void release_from_device_copy(const void *Ptr, const context &Context);
+void release_from_device_copy(const void *Ptr, const queue &Queue);
+
+} // namespace sycl::ext::oneapi::experimental
+```
+
+Table 1. Functions added by this extension.
+|====
+| Member Function | Description
+a|
+```
+void prepare_for_device_copy(
+    const void *Ptr,
+    size_t NumBytes,
+    const context &Context);
+```
+
+| Informs the implementation that the host memory range starting at `Ptr` and
+extending for `NumBytes` bytes may be used as either the source or destination
+of an explicit copy operation to a queue using context Context. Calling this
+function may allow the implementation to accelerate these copy operations.
+Generally, it is only useful to call this function if the host memory range will
+be used for several subsequent copy operations; it is not generally useful to
+call it before each individual copy operation.
+
+The behavior is undefined when multiple calls to `prepare_for_device_copy`
+specify memory ranges that overlap, even when using different
+SYCL contexts. No error message is issued.
+
+a|
+```
+void prepare_for_device_copy(
+    const void *Ptr,
+    size_t NumBytes,
+    const queue &Queue)
+```
+| Has the same effect as
+`prepare_for_device_copy(Ptr, NumBytes, Queue.get_context())`.
+
+a|
+```
+void release_from_device_copy(
+    const void *Ptr,
+    const context &Context)
+```
+| Undoes the effect of a previous call to prepare_for_device_copy on `Ptr`.
+It is still valid to copy to or from `Ptr`, but these copies might not be
+optimized.
+
+The `Ptr` and `Context` must match the values passed to a previous call to
+`prepare_for_device_copy` for a memory range that has not yet been
+released, otherwise the behavior is undefined.
+
+a|
+```
+void release_from_device_copy(
+    const void *Ptr,
+    const queue &Queue)
+```
+| Has the same effect as
+`release_from_device_copy(Ptr, Queue.get_context())`.
+
+|====

sycl/include/sycl/detail/pi.def

@@ -159,6 +159,10 @@ _PI_API(piextEnablePeerAccess)
 _PI_API(piextDisablePeerAccess)
 _PI_API(piextPeerAccessGetInfo)
 
+// USM import/release APIs
+_PI_API(piextUSMImport)
+_PI_API(piextUSMRelease)
+
 // command-buffer Extension
 _PI_API(piextCommandBufferCreate)
 _PI_API(piextCommandBufferRetain)

sycl/include/sycl/detail/pi.h

@@ -143,9 +143,10 @@
 //         - piextDestroyExternalSemaphore
 //         - piextWaitExternalSemaphore
 //         - piextSignalExternalSemaphore
+// 14.37 Added piextUSMImportExternalPointer and piextUSMReleaseImportedPointer.
 
 #define _PI_H_VERSION_MAJOR 14
-#define _PI_H_VERSION_MINOR 36
+#define _PI_H_VERSION_MINOR 37
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -2087,6 +2088,20 @@ __SYCL_EXPORT pi_result piextUSMEnqueueMemcpy2D(
     pi_uint32 num_events_in_waitlist, const pi_event *events_waitlist,
     pi_event *event);
 
+/// Import host system memory into USM.
+///
+/// \param ptr start address of memory range to import
+/// \param size is the number of bytes to import
+/// \param context is the pi_context
+__SYCL_EXPORT pi_result piextUSMImport(const void *ptr, size_t size,
+                                       pi_context context);
+
+/// Release host system memory from USM.
+///
+/// \param ptr start address of imported memory range
+/// \param context is the pi_context
+__SYCL_EXPORT pi_result piextUSMRelease(const void *ptr, pi_context context);
+
 ///
 /// Device global variable
 ///

sycl/include/sycl/usm.hpp

@@ -296,5 +296,23 @@ T *aligned_alloc(
                           Kind, PropList, CodeLoc);
 }
 
+// Device copy enhancement APIs, prepare_for and release_from USM.
+
+namespace ext::oneapi::experimental {
+
+__SYCL_EXPORT void prepare_for_device_copy(const void *Ptr, size_t Size,
+                                           const context &Context);
+
+__SYCL_EXPORT void prepare_for_device_copy(const void *Ptr, size_t Size,
+                                           const queue &Queue);
+
+__SYCL_EXPORT void release_from_device_copy(const void *Ptr,
+                                            const context &Context);
+
+__SYCL_EXPORT void release_from_device_copy(const void *Ptr,
+                                            const queue &Queue);
+
+} // namespace ext::oneapi::experimental
+
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/plugins/cuda/pi_cuda.cpp

@@ -168,6 +168,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMEnqueueMemset2D, pi2ur::piextUSMEnqueueMemset2D)
   _PI_CL(piextUSMEnqueueMemcpy2D, pi2ur::piextUSMEnqueueMemcpy2D)
   _PI_CL(piextUSMGetMemAllocInfo, pi2ur::piextUSMGetMemAllocInfo)
+  _PI_CL(piextUSMImport, pi2ur::piextUSMImport)
+  _PI_CL(piextUSMRelease, pi2ur::piextUSMRelease)
   // Device global variable
   _PI_CL(piextEnqueueDeviceGlobalVariableWrite,
          pi2ur::piextEnqueueDeviceGlobalVariableWrite)

sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp

@@ -2045,6 +2045,14 @@ pi_result piextUSMGetMemAllocInfo(pi_context, const void *, pi_mem_alloc_info,
   DIE_NO_IMPLEMENTATION;
 }
 
+pi_result piextUSMImport(const void *ptr, size_t size, pi_context context) {
+  return PI_SUCCESS;
+}
+
+pi_result piextUSMRelease(const void *ptr, pi_context context) {
+  return PI_SUCCESS;
+}
+
 /// Host Pipes
 pi_result piextEnqueueReadHostPipe(pi_queue, pi_program, const char *, pi_bool,
                                    void *, size_t, pi_uint32, const pi_event *,

sycl/plugins/hip/pi_hip.cpp

@@ -167,6 +167,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMEnqueueFill2D, pi2ur::piextUSMEnqueueFill2D)
   _PI_CL(piextUSMEnqueueMemset2D, pi2ur::piextUSMEnqueueMemset2D)
   _PI_CL(piextUSMGetMemAllocInfo, pi2ur::piextUSMGetMemAllocInfo)
+  _PI_CL(piextUSMImport, pi2ur::piextUSMImport)
+  _PI_CL(piextUSMRelease, pi2ur::piextUSMRelease)
   // Device global variable
   _PI_CL(piextEnqueueDeviceGlobalVariableWrite,
          pi2ur::piextEnqueueDeviceGlobalVariableWrite)

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -1065,6 +1065,14 @@ pi_result piextUSMGetMemAllocInfo(pi_context Context, const void *Ptr,
                                         ParamValue, ParamValueSizeRet);
 }
 
+pi_result piextUSMImport(const void *HostPtr, size_t Size, pi_context Context) {
+  return pi2ur::piextUSMImport(HostPtr, Size, Context);
+}
+
+pi_result piextUSMRelease(const void *HostPtr, pi_context Context) {
+  return pi2ur::piextUSMRelease(HostPtr, Context);
+}
+
 /// API for writing data from host to a device global variable.
 ///
 /// \param Queue is the queue

sycl/plugins/opencl/pi_opencl.cpp

@@ -2041,6 +2041,14 @@ pi_result piextUSMGetMemAllocInfo(pi_context context, const void *ptr,
   return RetVal;
 }
 
+pi_result piextUSMImport(const void *ptr, size_t size, pi_context context) {
+  return PI_SUCCESS;
+}
+
+pi_result piextUSMRelease(const void *ptr, pi_context context) {
+  return PI_SUCCESS;
+}
+
 /// API for writing data from host to a device global variable.
 ///
 /// \param queue is the queue
@@ -2695,6 +2703,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_CL(piextUSMEnqueueMemset2D, piextUSMEnqueueMemset2D)
   _PI_CL(piextUSMEnqueueMemcpy2D, piextUSMEnqueueMemcpy2D)
   _PI_CL(piextUSMGetMemAllocInfo, piextUSMGetMemAllocInfo)
+  _PI_CL(piextUSMImport, piextUSMImport)
+  _PI_CL(piextUSMRelease, piextUSMRelease)
   // Device global variable
   _PI_CL(piextEnqueueDeviceGlobalVariableWrite,
          piextEnqueueDeviceGlobalVariableWrite)

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -3280,6 +3280,29 @@ inline pi_result piextUSMGetMemAllocInfo(pi_context Context, const void *Ptr,
   return PI_SUCCESS;
 }
 
+inline pi_result piextUSMImport(const void *HostPtr, size_t Size,
+                                pi_context Context) {
+
+  PI_ASSERT(Context, PI_ERROR_INVALID_CONTEXT);
+
+  ur_context_handle_t UrContext =
+      reinterpret_cast<ur_context_handle_t>(Context);
+
+  HANDLE_ERRORS(urUSMImportExp(UrContext, const_cast<void *>(HostPtr), Size));
+  return PI_SUCCESS;
+}
+
+inline pi_result piextUSMRelease(const void *HostPtr, pi_context Context) {
+
+  PI_ASSERT(Context, PI_ERROR_INVALID_CONTEXT);
+
+  ur_context_handle_t UrContext =
+      reinterpret_cast<ur_context_handle_t>(Context);
+
+  HANDLE_ERRORS(urUSMReleaseExp(UrContext, const_cast<void *>(HostPtr)));
+  return PI_SUCCESS;
+}
+
 inline pi_result piMemImageGetInfo(pi_mem Image, pi_image_info ParamName,
                                    size_t ParamValueSize, void *ParamValue,
                                    size_t *ParamValueSizeRet) {

sycl/plugins/unified_runtime/pi_unified_runtime.cpp

@@ -368,6 +368,16 @@ __SYCL_EXPORT pi_result piextUSMFree(pi_context Context, void *Ptr) {
   return pi2ur::piextUSMFree(Context, Ptr);
 }
 
+__SYCL_EXPORT pi_result piextUSMImport(const void *HostPtr, size_t Size,
+                                       pi_context Context) {
+  return pi2ur::piextUSMImport(HostPtr, Size, Context);
+}
+
+__SYCL_EXPORT pi_result piextUSMRelease(const void *HostPtr,
+                                        pi_context Context) {
+  return pi2ur::piextUSMRelease(HostPtr, Context);
+}
+
 __SYCL_EXPORT pi_result piContextRetain(pi_context Context) {
   return pi2ur::piContextRetain(Context);
 }
@@ -1360,6 +1370,9 @@ __SYCL_EXPORT pi_result piPluginInit(pi_plugin *PluginInit) {
   _PI_API(piextUSMSharedAlloc)
   _PI_API(piextUSMFree)
 
+  _PI_API(piextUSMImport)
+  _PI_API(piextUSMRelease)
+
   _PI_API(piEnqueueKernelLaunch)
   _PI_API(piEnqueueMemImageWrite)
   _PI_API(piEnqueueMemImageRead)

sycl/plugins/unified_runtime/ur/adapters/cuda/usm.cpp

@@ -218,3 +218,18 @@ urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem,
   }
   return Result;
 }
+
+UR_APIEXPORT ur_result_t UR_APICALL urUSMImportExp(ur_context_handle_t Context,
+                                                   void *HostPtr, size_t Size) {
+  UR_ASSERT(Context, UR_RESULT_ERROR_INVALID_CONTEXT);
+  UR_ASSERT(!HostPtr, UR_RESULT_ERROR_INVALID_VALUE);
+  UR_ASSERT(Size > 0, UR_RESULT_ERROR_INVALID_VALUE);
+  return UR_RESULT_SUCCESS;
+}
+
+UR_APIEXPORT ur_result_t UR_APICALL urUSMReleaseExp(ur_context_handle_t Context,
+                                                    void *HostPtr) {
+  UR_ASSERT(Context, UR_RESULT_ERROR_INVALID_CONTEXT);
+  UR_ASSERT(!HostPtr, UR_RESULT_ERROR_INVALID_VALUE);
+  return UR_RESULT_SUCCESS;
+}

sycl/plugins/unified_runtime/ur/adapters/hip/usm.cpp

@@ -188,3 +188,18 @@ urUSMGetMemAllocInfo(ur_context_handle_t hContext, const void *pMem,
   }
   return Result;
 }
+
+UR_APIEXPORT ur_result_t UR_APICALL urUSMImportExp(ur_context_handle_t Context,
+                                                   void *HostPtr, size_t Size) {
+  UR_ASSERT(Context, UR_RESULT_ERROR_INVALID_CONTEXT);
+  UR_ASSERT(!HostPtr, UR_RESULT_ERROR_INVALID_VALUE);
+  UR_ASSERT(Size > 0, UR_RESULT_ERROR_INVALID_VALUE);
+  return UR_RESULT_SUCCESS;
+}
+
+UR_APIEXPORT ur_result_t UR_APICALL urUSMReleaseExp(ur_context_handle_t Context,
+                                                    void *HostPtr) {
+  UR_ASSERT(Context, UR_RESULT_ERROR_INVALID_CONTEXT);
+  UR_ASSERT(!HostPtr, UR_RESULT_ERROR_INVALID_VALUE);
+  return UR_RESULT_SUCCESS;
+}