[SYCL][Docs] Add sycl_ext_oneapi_inter_process_communication

This commit adds a new extension for inter-process communicable SYCL
object handles. As part of the initial version of this extension, only
inter-process communicable memory is exposed.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -97,6 +97,7 @@ def Aspectext_oneapi_exportable_device_mem : Aspect<"ext_oneapi_exportable_devic
 def Aspectext_oneapi_clock_sub_group : Aspect<"ext_oneapi_clock_sub_group">;
 def Aspectext_oneapi_clock_work_group : Aspect<"ext_oneapi_clock_work_group">;
 def Aspectext_oneapi_clock_device : Aspect<"ext_oneapi_clock_device">;
+def Aspectext_oneapi_ipc_memory : Aspect<"ext_oneapi_ipc_memory">;
 
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
@@ -174,7 +175,8 @@ def : TargetInfo<"__TestAspectList",
     Aspectext_oneapi_exportable_device_mem,
     Aspectext_oneapi_clock_sub_group,
     Aspectext_oneapi_clock_work_group,
-    Aspectext_oneapi_clock_device],
+    Aspectext_oneapi_clock_device,
+    Aspectext_oneapi_ipc_memory],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/doc/extensions/experimental/sycl_ext_oneapi_inter_process_communication.asciidoc

@@ -0,0 +1,189 @@
+= sycl_ext_oneapi_inter_process_communication
+
+: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++]
+:endnote: —{nbsp}end{nbsp}note
+
+// 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) 2025 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 10 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 {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.*
+
+
+== Backend support status
+
+The APIs in this extension may be used only on a device that has
+`aspect::ext_oneapi_ipc_memory`.  The application must check that the device has
+this aspect before submitting a kernel using any of the APIs in this
+extension.  If the application fails to do this, the implementation throws
+a synchronous exception with the `errc::kernel_not_supported` error code
+when the kernel is submitted to the queue.
+
+
+== Overview
+
+This extension adds the ability for SYCL programs to share device USM memory
+allocations between processes. This is done by the allocating process creating
+a new `ipc_memory` object and transferring the "handle data" to the other
+processes. The other processes can use the handle data to recreate the
+`ipc_memory` object and get a pointer to the corresponding device USM memory.
+
+
+== 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_IPC` 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.
+
+_And follow the text with a table like this *unless the extension is
+"experimental"*.  Note that your table may have more than one row if it
+has multiple versions._
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|The APIs of this experimental extension are not versioned, so the
+ feature-test macro always has this value.
+|===
+
+=== Inter-process communicable memory
+
+
+This extension adds the new `ipc_memory` class. This new class adheres to the
+common reference semantics described in
+https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:reference-semantics[Section 4.5.2.]
+in the SYCL 2020 specification.
+
+```
+namespace sycl::ext::oneapi::experimental {
+
+class ipc_memory {
+public:
+  ipc_memory(void *ptr, sycl::context &ctx);
+  ipc_memory(span<const char, sycl::dynamic_extent> ipc_memory_handle_data,
+             const sycl::context &ctx, const sycl::device &dev);
+
+  span<const char, sycl::dynamic_extent> get_handle_data() const;
+
+  void *get_ptr() const;
+};
+
+}
+```
+
+|====
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+ipc_memory(void *ptr, const sycl::context &ctx)
+----
+!====
+
+_Effects:_ Constructs an IPC memory object in `ctx` from a pointer `ptr` to
+device USM memory.
+If `ptr` is not pointing to device USM memory, the behaviors of this constructor
+and any resulting objects are undefined.
+
+!====
+a!
+[source]
+----
+ipc_memory(span<const char, sycl::dynamic_extent> ipc_memory_handle_data,
+           const sycl::context &ctx, const sycl::device &dev)
+----
+!====
+
+_Effects:_ Constructs an IPC memory object in `ctx` from the handle data
+`ipc_memory_handle_data` of returned by the `get_handle_data()` member function
+of another `ipc_memory` object.
+The `ipc_memory` object that the handle data originated from is allowed to be
+from another process on the host system.
+If the `ipc_memory` object that the handle data originated from has been
+destroyed, the behaviors of this constructor and any resulting objects are
+undefined.
+If the device USM memory the original `ipc_memory` object was created with was
+not originally allocated on `dev`, the behaviors of this constructor and any
+resulting objects are undefined.
+
+!====
+a!
+[source]
+----
+span<const char, sycl::dynamic_extent> get_handle_data() const
+----
+!====
+
+_Returns:_ The handle data of the `ipc_memory` object.
+Accessing the handle data returned by this API after the `ipc_memory` object has
+been destroyed results in undefined behavior.
+
+!====
+a!
+[source]
+----
+void *get_ptr() const
+----
+!====
+
+_Returns:_ A pointer to device USM memory corresponding to the pointer used to
+construct the original `ipc_memory` object.
+Accessing the pointer returned by this API after the `ipc_memory` object has
+been destroyed results in undefined behavior.
+
+|====
+
+

sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp

@@ -0,0 +1,59 @@
+//==------- ipc_memory.hpp --- SYCL inter-process communicable memory ------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/defines_elementary.hpp>
+#include <sycl/detail/export.hpp>
+#include <sycl/detail/owner_less_base.hpp>
+#include <sycl/sycl_span.hpp>
+
+#include <memory>
+
+namespace sycl {
+inline namespace _V1 {
+
+class context;
+class device;
+
+namespace detail {
+class ipc_memory_impl;
+}
+
+namespace ext::oneapi::experimental {
+class __SYCL_EXPORT ipc_memory
+    : public sycl::detail::OwnerLessBase<ipc_memory> {
+public:
+  ipc_memory(void *Ptr, const sycl::context &Ctx);
+  ipc_memory(const span<const char, sycl::dynamic_extent> IPCMemoryHandleData,
+             const sycl::context &Ctx, const sycl::device &Dev);
+
+  sycl::span<const char, sycl::dynamic_extent> get_handle_data() const;
+
+  void *get_ptr() const;
+
+private:
+  ipc_memory(std::shared_ptr<sycl::detail::ipc_memory_impl> IPCMemImpl)
+      : impl{IPCMemImpl} {}
+
+  std::shared_ptr<sycl::detail::ipc_memory_impl> impl;
+
+  template <class Obj>
+  friend const decltype(Obj::impl) &
+  sycl::detail::getSyclObjImpl(const Obj &SyclObject);
+
+  template <class T>
+  friend T sycl::detail::createSyclObjFromImpl(
+      std::add_rvalue_reference_t<decltype(T::impl)> ImplObj);
+  template <class T>
+  friend T sycl::detail::createSyclObjFromImpl(
+      std::add_lvalue_reference_t<const decltype(T::impl)> ImplObj);
+};
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/info/aspects.def

@@ -83,3 +83,4 @@ __SYCL_ASPECT(ext_oneapi_exportable_device_mem, 90)
 __SYCL_ASPECT(ext_oneapi_clock_sub_group, 91)
 __SYCL_ASPECT(ext_oneapi_clock_work_group, 92)
 __SYCL_ASPECT(ext_oneapi_clock_device, 93)
+__SYCL_ASPECT(ext_oneapi_ipc_memory, 94)

sycl/include/sycl/sycl.hpp

@@ -128,6 +128,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/ext/oneapi/experimental/group_helpers_sorters.hpp>
 #include <sycl/ext/oneapi/experimental/group_load_store.hpp>
 #include <sycl/ext/oneapi/experimental/group_sort.hpp>
+#include <sycl/ext/oneapi/experimental/ipc_memory.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
 #include <sycl/ext/oneapi/experimental/profiling_tag.hpp>
 #include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>

sycl/source/CMakeLists.txt

@@ -321,6 +321,7 @@ set(SYCL_COMMON_SOURCES
     "handler.cpp"
     "image.cpp"
     "interop_handle.cpp"
+    "ipc_memory.cpp"
     "kernel.cpp"
     "kernel_bundle.cpp"
     "physical_mem.cpp"

sycl/source/detail/device_impl.hpp

@@ -1591,6 +1591,10 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
       // Will be updated in a follow-up UR patch.
       return false;
     }
+    CASE(ext_oneapi_ipc_memory) {
+      return get_info_impl_nocheck<UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP>()
+          .value_or(0);
+    }
     else {
       return false; // This device aspect has not been implemented yet.
     }