diff --git a/CMakeLists.txt b/CMakeLists.txt index 7ef8113778..0f84348251 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -115,7 +115,7 @@ install(DIRECTORY ) # find Python before enabling pybind11 -find_package(Python REQUIRED COMPONENTS Development.Module) +find_package(Python REQUIRED COMPONENTS Interpreter Development.Module) # Define CMAKE_INSTALL_xxx: LIBDIR, INCLUDEDIR include(GNUInstallDirs) diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 2f8911fdc5..91629e2947 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -101,6 +101,7 @@ cdef extern from "syclinterface/dpctl_sycl_enum_types.h": _emulated "emulated", _is_component "is_component", _is_composite "is_composite", + _ext_oneapi_ipc_memory "ext_oneapi_ipc_memory", ctypedef enum _partition_affinity_domain_type \ "DPCTLPartitionAffinityDomainType": @@ -595,6 +596,7 @@ cdef extern from "syclinterface/dpctl_sycl_extension_interface.h": DPCTLSyclWorkGroupMemoryRef Ref) cdef bint DPCTLWorkGroupMemory_Available() + cdef bint DPCTLIPCMem_Available() cdef struct DPCTLOpaqueRawKernelArg ctypedef DPCTLOpaqueRawKernelArg *DPCTLSyclRawKernelArgRef @@ -606,3 +608,19 @@ cdef extern from "syclinterface/dpctl_sycl_extension_interface.h": DPCTLSyclRawKernelArgRef Ref) cdef bint DPCTLRawKernelArg_Available() + +cdef extern from "syclinterface/dpctl_sycl_ipc_memory_interface.h": + cdef int DPCTLIPCMem_GetHandle( + DPCTLSyclUSMRef Ptr, + DPCTLSyclContextRef CRef, + char **DataOut, + size_t *SizeOut) + cdef DPCTLSyclUSMRef DPCTLIPCMem_OpenHandle( + const char *HandleData, + size_t HandleDataSize, + DPCTLSyclContextRef CRef, + DPCTLSyclDeviceRef DRef) + cdef void DPCTLIPCMem_CloseHandle( + DPCTLSyclUSMRef MappedPtr, + DPCTLSyclContextRef CRef) + cdef void DPCTLIPCMem_FreeHandleData(char *Data) diff --git a/dpctl/_sycl_device.pyx b/dpctl/_sycl_device.pyx index deda5a6ff8..506664ab18 100644 --- a/dpctl/_sycl_device.pyx +++ b/dpctl/_sycl_device.pyx @@ -885,6 +885,18 @@ cdef class SyclDevice(_SyclDevice): cdef _aspect_type AT = _aspect_type._is_composite return DPCTLDevice_HasAspect(self._device_ref, AT) + @property + def has_aspect_ext_oneapi_ipc_memory(self): + """ Returns ``True`` if this device supports inter-process + communication (IPC) memory handles, ``False`` otherwise. + + Returns: + bool: + Indicates if device supports IPC memory. + """ + cdef _aspect_type AT = _aspect_type._ext_oneapi_ipc_memory + return DPCTLDevice_HasAspect(self._device_ref, AT) + @property def image_2d_max_width(self): """ Returns the maximum width of a 2D image or 1D image in pixels. diff --git a/dpctl/memory/__init__.py b/dpctl/memory/__init__.py index fc5fc9170d..69fa3caab4 100644 --- a/dpctl/memory/__init__.py +++ b/dpctl/memory/__init__.py @@ -30,17 +30,25 @@ """ from ._memory import ( + MemoryIPCDevice, MemoryUSMDevice, MemoryUSMHost, MemoryUSMShared, + SyclIPCCloseMemHandle, + SyclIPCGetMemHandle, + SyclIPCOpenMemHandle, USMAllocationError, as_usm_memory, ) __all__ = [ + "MemoryIPCDevice", "MemoryUSMDevice", "MemoryUSMHost", "MemoryUSMShared", "USMAllocationError", "as_usm_memory", + "SyclIPCGetMemHandle", + "SyclIPCOpenMemHandle", + "SyclIPCCloseMemHandle", ] diff --git a/dpctl/memory/_memory.pxd b/dpctl/memory/_memory.pxd index d8842a565b..ce1156f5c0 100644 --- a/dpctl/memory/_memory.pxd +++ b/dpctl/memory/_memory.pxd @@ -90,3 +90,12 @@ cdef public api class MemoryUSMDevice(_Memory) [ object PyMemoryUSMDeviceObject, type PyMemoryUSMDeviceType ]: pass + + +cdef class MemoryIPCDevice(MemoryUSMDevice): + @staticmethod + cdef object create_ipc_from_usm_pointer_size_qref( + DPCTLSyclUSMRef USMRef, + Py_ssize_t nbytes, + DPCTLSyclQueueRef QRef, + ) diff --git a/dpctl/memory/_memory.pyx b/dpctl/memory/_memory.pyx index 99223039ba..488115ebb5 100644 --- a/dpctl/memory/_memory.pyx +++ b/dpctl/memory/_memory.pyx @@ -39,6 +39,11 @@ from dpctl._backend cimport ( # noqa: E211 DPCTLDevice_Copy, DPCTLEvent_Delete, DPCTLEvent_Wait, + DPCTLIPCMem_Available, + DPCTLIPCMem_CloseHandle, + DPCTLIPCMem_FreeHandleData, + DPCTLIPCMem_GetHandle, + DPCTLIPCMem_OpenHandle, DPCTLmalloc_device, DPCTLmalloc_host, DPCTLmalloc_shared, @@ -73,7 +78,11 @@ __all__ = [ "MemoryUSMShared", "MemoryUSMHost", "MemoryUSMDevice", + "MemoryIPCDevice", "USMAllocationError", + "SyclIPCGetMemHandle", + "SyclIPCOpenMemHandle", + "SyclIPCCloseMemHandle", ] include "_sycl_usm_array_interface_utils.pxi" @@ -892,6 +901,60 @@ cdef class MemoryUSMDevice(_Memory): ) +cdef class MemoryIPCDevice(MemoryUSMDevice): + """ + Class representing emory object backed by an IPC-mapped USM device pointer. + + This class is not intended to be instantiated directly. + """ + + @property + def is_closed(self): + """Whether the IPC mapping has been closed.""" + return self._memory_ptr is NULL + + def __dealloc__(self): + cdef DPCTLSyclContextRef ctx_ref + if self._memory_ptr is not NULL: + ctx_ref = DPCTLQueue_GetContext(self.queue.get_queue_ref()) + DPCTLIPCMem_CloseHandle(self._memory_ptr, ctx_ref) + DPCTLContext_Delete(ctx_ref) + self._memory_ptr = NULL + + @staticmethod + cdef object create_ipc_from_usm_pointer_size_qref( + DPCTLSyclUSMRef USMRef, + Py_ssize_t nbytes, + DPCTLSyclQueueRef QRef, + ): + """Create a MemoryIPCDevice wrapping an IPC-mapped pointer.""" + cdef DPCTLSyclQueueRef QRef_copy = NULL + cdef _Memory base + + if nbytes <= 0: + raise ValueError("Number of bytes must be positive") + if QRef is NULL: + raise TypeError("Argument DPCTLSyclQueueRef is NULL") + + base = _Memory.__new__(_Memory) + base._cinit_empty() + base._memory_ptr = USMRef + base.nbytes = nbytes + + QRef_copy = DPCTLQueue_Copy(QRef) + if QRef_copy is NULL: + raise ValueError("Referenced queue could not be copied.") + try: + base.queue = SyclQueue._create(QRef_copy) + except dpctl.SyclQueueCreationError as sqce: + raise ValueError( + "SyclQueue object could not be created from " + "copy of referenced queue" + ) from sqce + + return MemoryIPCDevice(base) + + def as_usm_memory(obj): """ as_usm_memory(obj) @@ -961,3 +1024,306 @@ cdef api object Memory_Make( return _Memory.create_from_usm_pointer_size_qref( ptr, nbytes, QRef, memory_owner=owner ) + + +# ─── IPC Memory Handle ──────────────────────────────────────────────── + +cdef class IPCMemoryHandle: + """Wrapper around a SYCL IPC memory handle. + + Instances are created by passing a :class:`MemoryUSMDevice` to the + constructor. The resulting object exposes :meth:`to_bytes` which + returns an opaque ``bytes`` payload suitable for inter-process + transport (e.g. via pickle, ZMQ, shared-memory). + + On the receiving side, call :meth:`IPCMemoryHandle.open` with the + payload and a target device to obtain a :class:`MemoryIPCDevice` + backed by the IPC-mapped memory. + + Parameters + ---------- + usm_memory : MemoryUSMDevice + USM device memory object whose pointer to export. + context : dpctl.SyclContext, optional + SYCL context to use. Defaults to the context of *usm_memory*'s + queue. + + Raises + ------ + TypeError + If *usm_memory* is not a MemoryUSMDevice instance. + RuntimeError + If IPC memory is not supported or the handle export fails. + + Examples + -------- + Sender process:: + + from dpctl.memory import MemoryUSMDevice, IPCMemoryHandle + mem = MemoryUSMDevice(4096) + handle = IPCMemoryHandle(mem) + raw = handle.to_bytes() # send to another process + + Receiver process:: + + from dpctl.memory import IPCMemoryHandle + mapped_mem = IPCMemoryHandle.open(raw, device, nbytes=4096) + # ... use mapped_mem ... + IPCMemoryHandle.close_mapping(mapped_mem) + """ + + cdef bytes _handle_bytes + cdef SyclContext _ctx + cdef void *_opaque_ptr + cdef bint _closed + + def __cinit__(self): + self._handle_bytes = None + self._ctx = None + self._opaque_ptr = NULL + self._closed = False + + def __init__(self, _Memory usm_memory not None, SyclContext context=None): + if not DPCTLIPCMem_Available(): + raise RuntimeError("IPC memory not supported in this build") + + if not isinstance(usm_memory, MemoryUSMDevice): + raise TypeError( + "IPC handles are only supported for USM device " + "allocations, not " + type(usm_memory).__name__ + ) + + cdef DPCTLSyclUSMRef ptr = usm_memory._memory_ptr + if ptr is NULL: + raise ValueError("USM memory object has a null pointer") + + cdef void *src_opaque = usm_memory._opaque_ptr + if src_opaque is NULL: + raise ValueError( + "USM memory object does not own its allocation" + ) + + cdef SyclDevice dev = usm_memory.sycl_device + if not dev.has_aspect_ext_oneapi_ipc_memory: + raise RuntimeError( + "Device does not support IPC memory " + "(aspect::ext_oneapi_ipc_memory)" + ) + + cdef SyclQueue q = usm_memory.queue + if context is None: + context = q.sycl_context + + cdef DPCTLSyclContextRef ctx_ref = context.get_context_ref() + cdef char *data_out = NULL + cdef size_t size_out = 0 + + cdef int rc = DPCTLIPCMem_GetHandle(ptr, ctx_ref, &data_out, &size_out) + if rc != 0: + raise RuntimeError( + "DPCTLIPCMem_GetHandle failed \u2014 IPC handle export failed" + ) + + try: + self._handle_bytes = PyBytes_FromStringAndSize( + data_out, size_out) + finally: + DPCTLIPCMem_FreeHandleData(data_out) + + self._opaque_ptr = OpaqueSmartPtr_Copy(src_opaque) + self._ctx = context + self._closed = False + + def to_bytes(self): + """Return the raw IPC handle data as ``bytes``. + + The returned object can be pickled, sent over a socket, or + written to shared memory for another process to consume via + :meth:`open`. + """ + if self._closed: + raise RuntimeError("IPC handle has already been closed") + return self._handle_bytes + + @staticmethod + def open(bytes handle_bytes not None, + SyclDevice device not None, + Py_ssize_t nbytes, + SyclContext context=None): + """Open an IPC handle in this process. + + Parameters + ---------- + handle_bytes : bytes + Opaque payload from :meth:`to_bytes` (possibly from another + process). + device : dpctl.SyclDevice + Device to map the memory on. + nbytes : int + Byte size of the original allocation. Must be > 0. + context : dpctl.SyclContext, optional + SYCL context to use. Defaults to the default context for + *device*'s platform. + + Returns + ------- + MemoryIPCDevice + A USM device memory object backed by the IPC-mapped pointer. + Call :meth:`close_mapping` when done. + + Raises + ------ + RuntimeError + If the device does not support IPC memory or the handle + cannot be opened. + ValueError + If *nbytes* <= 0. + """ + if not DPCTLIPCMem_Available(): + raise RuntimeError("IPC memory not supported in this build") + + if not device.has_aspect_ext_oneapi_ipc_memory: + raise RuntimeError( + "Device does not support IPC memory " + "(aspect::ext_oneapi_ipc_memory)" + ) + + if nbytes <= 0: + raise ValueError("nbytes must be > 0 for IPC open") + + cdef const char *raw = PyBytes_AS_STRING(handle_bytes) + cdef size_t raw_size = len(handle_bytes) + + if context is None: + context = device.sycl_platform.default_context + + cdef DPCTLSyclContextRef ctx_ref = context.get_context_ref() + cdef DPCTLSyclDeviceRef dev_ref = device.get_device_ref() + + cdef DPCTLSyclUSMRef mapped_ptr = DPCTLIPCMem_OpenHandle( + raw, raw_size, ctx_ref, dev_ref) + if mapped_ptr is NULL: + raise RuntimeError("DPCTLIPCMem_OpenHandle failed") + + cdef SyclQueue q + try: + q = dpctl.SyclQueue(context, device) + mem = MemoryIPCDevice.create_ipc_from_usm_pointer_size_qref( + mapped_ptr, nbytes, q.get_queue_ref()) + except Exception: + DPCTLIPCMem_CloseHandle(mapped_ptr, ctx_ref) + raise + + return mem + + @staticmethod + def close_mapping(_Memory usm_memory not None): + """Explicitly close an IPC mapping. + + After calling this, *usm_memory* is invalidated and must not be + used again. + + Parameters + ---------- + usm_memory : _Memory + The memory object returned by :meth:`open`. + """ + cdef DPCTLSyclUSMRef ptr + cdef DPCTLSyclContextRef ctx_ref + + if not DPCTLIPCMem_Available(): + raise RuntimeError("IPC memory not supported in this build") + + if not isinstance(usm_memory, MemoryIPCDevice): + raise RuntimeError( + "close_mapping called on a non-IPC memory object; " + "this method is only valid for IPC-mapped memory" + ) + + ptr = usm_memory._memory_ptr + if ptr is NULL: + return + + ctx_ref = DPCTLQueue_GetContext(usm_memory.queue.get_queue_ref()) + DPCTLIPCMem_CloseHandle(ptr, ctx_ref) + DPCTLContext_Delete(ctx_ref) + + usm_memory._memory_ptr = NULL + usm_memory.nbytes = 0 + + def close(self): + """Mark this handle object as closed.""" + self._closed = True + + def __dealloc__(self): + if self._opaque_ptr is not NULL: + OpaqueSmartPtr_Delete(self._opaque_ptr) + self._opaque_ptr = NULL + self._closed = True + + def __repr__(self): + cdef Py_ssize_t sz = 0 + if self._handle_bytes is not None: + sz = len(self._handle_bytes) + return ( + f"IPCMemoryHandle(size={sz}, closed={self._closed})" + ) + + +# ─── SYCL IPC free functions ────────────────────────────────────────── + +def SyclIPCGetMemHandle(_Memory usm_memory not None, SyclContext context=None): + """Export a USM device allocation as IPC handle bytes. + + Parameters + ---------- + usm_memory : MemoryUSMDevice + USM device memory to export. + context : dpctl.SyclContext, optional + SYCL context. Defaults to the memory's queue context. + + Returns + ------- + bytes + Opaque IPC handle data for cross-process transport. + """ + cdef IPCMemoryHandle h = IPCMemoryHandle(usm_memory, context) + return h.to_bytes() + + +def SyclIPCOpenMemHandle(bytes handle_bytes not None, + SyclDevice device not None, + Py_ssize_t nbytes, + SyclContext context=None): + """Open an IPC handle and return a mapped MemoryIPCDevice. + + Parameters + ---------- + handle_bytes : bytes + Opaque payload from :func:`SyclIPCGetMemHandle`. + device : dpctl.SyclDevice + Device to map the memory on. + nbytes : int + Byte size of the original allocation. + context : dpctl.SyclContext, optional + SYCL context. Defaults to the platform default context. + + Returns + ------- + MemoryIPCDevice + Mapped IPC memory. Must be closed with :func:`SyclIPCCloseMemHandle`. + """ + return IPCMemoryHandle.open(handle_bytes, device, nbytes, context) + + +def SyclIPCCloseMemHandle(_Memory usm_memory not None): + """Close an IPC memory mapping. + + After this call, *usm_memory* is invalidated. + + Parameters + ---------- + usm_memory : _Memory + The memory object returned by :func:`SyclIPCOpenMemHandle`. + """ + IPCMemoryHandle.close_mapping(usm_memory) diff --git a/dpctl/tests/test_sycl_ipc_memory.py b/dpctl/tests/test_sycl_ipc_memory.py new file mode 100644 index 0000000000..fae13c8287 --- /dev/null +++ b/dpctl/tests/test_sycl_ipc_memory.py @@ -0,0 +1,257 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2026 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Tests for IPC memory handle functionality.""" + +import gc + +import pytest + +import dpctl +from dpctl.memory import ( + MemoryIPCDevice, + MemoryUSMDevice, + MemoryUSMHost, + MemoryUSMShared, + SyclIPCCloseMemHandle, + SyclIPCGetMemHandle, + SyclIPCOpenMemHandle, +) + + +def _get_ipc_device(): + """Return a device with IPC memory support, or skip the test.""" + try: + dev = dpctl.SyclDevice() + except dpctl.SyclDeviceCreationError: + pytest.skip("Default device not available") + if not dev.has_aspect_ext_oneapi_ipc_memory: + pytest.skip("Device does not support IPC memory") + return dev + + +def _get_ipc_queue(): + """Return a queue on a device with IPC memory support.""" + dev = _get_ipc_device() + return dpctl.SyclQueue(dev) + + +# ─── Handle creation tests ───────────────────────────────────────────── + + +class TestIPCHandleCreation: + def test_get_handle_from_device_memory(self): + q = _get_ipc_queue() + mem = MemoryUSMDevice(1024, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + assert isinstance(handle_bytes, bytes) + assert len(handle_bytes) > 0 + + def test_get_handle_rejects_shared_memory(self): + q = _get_ipc_queue() + mem = MemoryUSMShared(1024, queue=q) + with pytest.raises(TypeError): + SyclIPCGetMemHandle(mem) + + def test_get_handle_rejects_host_memory(self): + q = _get_ipc_queue() + mem = MemoryUSMHost(1024, queue=q) + with pytest.raises(TypeError): + SyclIPCGetMemHandle(mem) + + +# ─── Open / close tests ─────────────────────────────────────────────── + + +class TestIPCOpenClose: + def test_open_returns_ipc_device_memory(self): + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + assert isinstance(mapped, MemoryIPCDevice) + assert isinstance(mapped, MemoryUSMDevice) + assert mapped.nbytes == 4096 + + def test_close_mapping_nulls_pointer(self): + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + assert mapped._pointer != 0 + + SyclIPCCloseMemHandle(mapped) + assert mapped._pointer == 0 + assert mapped.nbytes == 0 + + def test_close_mapping_idempotent(self): + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + SyclIPCCloseMemHandle(mapped) + # Second close should be a no-op (ptr is NULL) + SyclIPCCloseMemHandle(mapped) + + def test_close_mapping_rejects_non_ipc_memory(self): + q = _get_ipc_queue() + mem = MemoryUSMDevice(1024, queue=q) + with pytest.raises(RuntimeError, match="non-IPC"): + SyclIPCCloseMemHandle(mem) + + +# ─── Type identity tests ────────────────────────────────────────────── + + +class TestIPCTypeIdentity: + def test_normal_memory_is_not_ipc(self): + q = _get_ipc_queue() + mem = MemoryUSMDevice(1024, queue=q) + assert not isinstance(mem, MemoryIPCDevice) + + def test_shared_memory_is_not_ipc(self): + q = _get_ipc_queue() + mem = MemoryUSMShared(1024, queue=q) + assert not isinstance(mem, MemoryIPCDevice) + + def test_host_memory_is_not_ipc(self): + q = _get_ipc_queue() + mem = MemoryUSMHost(1024, queue=q) + assert not isinstance(mem, MemoryIPCDevice) + + def test_ipc_mapped_memory_is_ipc_type(self): + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + assert isinstance(mapped, MemoryIPCDevice) + + def test_type_persists_after_close_mapping(self): + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + SyclIPCCloseMemHandle(mapped) + assert isinstance(mapped, MemoryIPCDevice) + + def test_is_closed_property(self): + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + assert mapped.is_closed is False + + SyclIPCCloseMemHandle(mapped) + assert mapped.is_closed is True + + +# ─── Lifetime / deletion tests ──────────────────────────────────────── + + +class TestIPCLifetime: + def test_source_memory_outlives_mapped(self): + """Source memory must stay alive while receiver uses the mapping.""" + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + assert mapped.nbytes == 4096 + + # Source is still alive — mapping is valid + SyclIPCCloseMemHandle(mapped) + del mem + gc.collect() + + def test_del_mapped_memory_does_not_crash(self): + """Deleting IPC-mapped memory should call CloseHandle, not + sycl::free, and not crash.""" + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + assert isinstance(mapped, MemoryIPCDevice) + + # del triggers __dealloc__ → CloseHandle path + del mapped + gc.collect() + + def test_del_mapped_memory_after_explicit_close(self): + """If close_mapping was already called, __dealloc__ should + be a no-op (ptr is NULL).""" + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + SyclIPCCloseMemHandle(mapped) + + # __dealloc__ sees NULL pointer, does nothing + del mapped + gc.collect() + + def test_multiple_opens_from_same_handle(self): + """Each open() produces an independent mapping that must be + closed independently.""" + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(4096, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped1 = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + mapped2 = SyclIPCOpenMemHandle(handle_bytes, dev, 4096) + + assert isinstance(mapped1, MemoryIPCDevice) + assert isinstance(mapped2, MemoryIPCDevice) + + SyclIPCCloseMemHandle(mapped1) + # mapped2 is still valid + assert mapped2._pointer != 0 + SyclIPCCloseMemHandle(mapped2) + + def test_del_mapped_via_gc(self): + """Full lifecycle: get handle, open, use, let GC collect mapped.""" + q = _get_ipc_queue() + dev = q.sycl_device + mem = MemoryUSMDevice(2048, queue=q) + handle_bytes = SyclIPCGetMemHandle(mem) + + mapped = SyclIPCOpenMemHandle(handle_bytes, dev, 2048) + assert mapped.nbytes == 2048 + assert isinstance(mapped, MemoryIPCDevice) + + # GC collects mapped → MemoryIPCDevice.__dealloc__ → CloseHandle + del mapped + gc.collect() + + # Source is still valid after receiver is gone + assert mem.nbytes == 2048 diff --git a/libsyclinterface/CMakeLists.txt b/libsyclinterface/CMakeLists.txt index b4331bffbb..3ad4ae5418 100644 --- a/libsyclinterface/CMakeLists.txt +++ b/libsyclinterface/CMakeLists.txt @@ -111,6 +111,35 @@ endif() message(STATUS "LIB_ZE: ${LIBZE_LOADER_FILENAME}") message(STATUS "LIB_CL: ${LIBCL_LOADER_FILENAME}") + +# --- IPC Memory support detection --- +option(DPCTL_ENABLE_IPC_MEMORY + "Enable IPC memory support (requires SYCL IPC runtime)" + ON +) +if(DPCTL_ENABLE_IPC_MEMORY) + include(CheckCXXSourceCompiles) + set(CMAKE_REQUIRED_FLAGS "${SYCL_FLAGS}") + set(CMAKE_REQUIRED_INCLUDES "${SYCL_INCLUDE_DIR}") + check_cxx_source_compiles(" + #include + int main() { return 0; } + " DPCTL_IPC_MEMORY_HEADER_FOUND) + if(DPCTL_IPC_MEMORY_HEADER_FOUND) + # Header found; enable IPC memory support + set(DPCTL_HAS_IPC_MEMORY 1) + set(DPCTL_HAS_IPC_MEMORY 1 PARENT_SCOPE) + message(STATUS "SYCL IPC memory support: ENABLED (ipc_memory.hpp found)") + else() + message(STATUS "SYCL IPC memory support: DISABLED (ipc_memory.hpp not found)") + endif() + unset(CMAKE_REQUIRED_FLAGS) + unset(CMAKE_REQUIRED_INCLUDES) + unset(CMAKE_REQUIRED_LIBRARIES) +else() + message(STATUS "SYCL IPC memory support: DISABLED (DPCTL_ENABLE_IPC_MEMORY=OFF)") +endif() + configure_file( ${CMAKE_CURRENT_SOURCE_DIR}/include/syclinterface/Config/dpctl_config.h.in ${CMAKE_CURRENT_SOURCE_DIR}/include/syclinterface/Config/dpctl_config.h @@ -222,8 +251,22 @@ file(GLOB_RECURSE sources list(REMOVE_ITEM sources "${CMAKE_CURRENT_SOURCE_DIR}/source/dpctl_vector_templ.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/source/dpctl_sycl_ipc_memory_interface.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/source/dpctl_sycl_ipc_memory_stubs.cpp" ) +# IPC memory: use real implementation or stubs depending on header availability. +if(DPCTL_HAS_IPC_MEMORY) + list(APPEND sources + "${CMAKE_CURRENT_SOURCE_DIR}/source/dpctl_sycl_ipc_memory_interface.cpp" + ) +else() + list(APPEND sources + "${CMAKE_CURRENT_SOURCE_DIR}/source/dpctl_sycl_ipc_memory_stubs.cpp" + ) + message(STATUS "Using IPC memory stubs (header not found)") +endif() + file(GLOB_RECURSE helper_sources ${CMAKE_CURRENT_SOURCE_DIR}/helper/source/*.cpp ) diff --git a/libsyclinterface/helper/source/dpctl_utils_helper.cpp b/libsyclinterface/helper/source/dpctl_utils_helper.cpp index 3b7a1eb3c8..347d036458 100644 --- a/libsyclinterface/helper/source/dpctl_utils_helper.cpp +++ b/libsyclinterface/helper/source/dpctl_utils_helper.cpp @@ -224,6 +224,11 @@ std::string DPCTL_AspectToStr(aspect aspectTy) case aspect::ext_oneapi_is_composite: ss << "is_composite"; break; +#ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION + case aspect::ext_oneapi_ipc_memory: + ss << "ext_oneapi_ipc_memory"; + break; +#endif default: throw std::runtime_error("Unsupported aspect type"); } @@ -299,6 +304,11 @@ aspect DPCTL_StrToAspectType(const std::string &aspectTyStr) else if (aspectTyStr == "is_composite") { aspectTy = aspect::ext_oneapi_is_composite; } +#ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION + else if (aspectTyStr == "ext_oneapi_ipc_memory") { + aspectTy = aspect::ext_oneapi_ipc_memory; + } +#endif else { // \todo handle the error throw std::runtime_error("Unsupported aspect type"); @@ -351,6 +361,10 @@ aspect DPCTL_DPCTLAspectTypeToSyclAspect(DPCTLSyclAspectType AspectTy) return aspect::ext_oneapi_is_component; case DPCTLSyclAspectType::is_composite: return aspect::ext_oneapi_is_composite; +#ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION + case DPCTLSyclAspectType::ext_oneapi_ipc_memory: + return aspect::ext_oneapi_ipc_memory; +#endif default: throw std::runtime_error("Unsupported aspect type"); } @@ -401,6 +415,10 @@ DPCTLSyclAspectType DPCTL_SyclAspectToDPCTLAspectType(aspect Aspect) return DPCTLSyclAspectType::is_composite; case aspect::ext_oneapi_is_component: return DPCTLSyclAspectType::is_component; +#ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION + case aspect::ext_oneapi_ipc_memory: + return DPCTLSyclAspectType::ext_oneapi_ipc_memory; +#endif default: throw std::runtime_error("Unsupported aspect type"); } diff --git a/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in b/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in index 0568f38507..97e28d1219 100644 --- a/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in +++ b/libsyclinterface/include/syclinterface/Config/dpctl_config.h.in @@ -35,3 +35,5 @@ #define DPCTL_LIBZE_LOADER_FILENAME "@LIBZE_LOADER_FILENAME@" #define DPCTL_LIBCL_LOADER_FILENAME "@LIBCL_LOADER_FILENAME@" +/* Defined when the SYCL runtime supports IPC memory (ipc::memory API). */ +#cmakedefine DPCTL_HAS_IPC_MEMORY 1 diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h index e8c4cba7e9..e6ee311980 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_enum_types.h @@ -133,7 +133,8 @@ typedef enum host_debuggable, emulated, is_component, - is_composite + is_composite, + ext_oneapi_ipc_memory } DPCTLSyclAspectType; /*! diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_extension_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_extension_interface.h index c7ff7463ea..fa0b1c15df 100644 --- a/libsyclinterface/include/syclinterface/dpctl_sycl_extension_interface.h +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_extension_interface.h @@ -53,6 +53,9 @@ void DPCTLWorkGroupMemory_Delete(__dpctl_take DPCTLSyclWorkGroupMemoryRef Ref); DPCTL_API bool DPCTLWorkGroupMemory_Available(); +DPCTL_API +bool DPCTLIPCMem_Available(); + typedef struct DPCTLOpaqueSyclRawKernelArg *DPCTLSyclRawKernelArgRef; DPCTL_API diff --git a/libsyclinterface/include/syclinterface/dpctl_sycl_ipc_memory_interface.h b/libsyclinterface/include/syclinterface/dpctl_sycl_ipc_memory_interface.h new file mode 100644 index 0000000000..a62265af1c --- /dev/null +++ b/libsyclinterface/include/syclinterface/dpctl_sycl_ipc_memory_interface.h @@ -0,0 +1,105 @@ +//===- dpctl_sycl_ipc_memory_interface.h - C API for SYCL IPC mem -*-C++-*-===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2026 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This header declares a C interface to +/// sycl::ext::oneapi::experimental::ipc::memory functions. +/// +//===----------------------------------------------------------------------===// + +#pragma once + +#include "Support/DllExport.h" +#include "Support/ExternC.h" +#include "Support/MemOwnershipAttrs.h" +#include "dpctl_data_types.h" +#include "dpctl_sycl_types.h" + +DPCTL_C_EXTERN_C_BEGIN + +/** + * @defgroup IPCMemoryInterface IPC Memory Interface + */ + +/*! + * @brief Get an IPC memory handle for a USM device pointer. + * + * Wraps ``sycl::ext::oneapi::experimental::ipc::memory::get()``. + * The returned handle bytes are copied out and the driver-side handle + * resource is released (via ``put``) before returning. + * + * @param Ptr USM device pointer to export. + * @param CRef Sycl context associated with the pointer. + * @param DataOut [out] Pointer to receive a malloc'd byte buffer + * containing the IPC handle data. Caller must free + * with DPCTLIPCMem_FreeHandleData(). + * @param SizeOut [out] Pointer to receive the byte count of DataOut. + * @return 0 on success, non-zero on failure. + * @ingroup IPCMemoryInterface + */ +DPCTL_API +int DPCTLIPCMem_GetHandle(__dpctl_keep DPCTLSyclUSMRef Ptr, + __dpctl_keep const DPCTLSyclContextRef CRef, + char **DataOut, + size_t *SizeOut); + +/*! + * @brief Open an IPC memory handle in the receiving process. + * + * Wraps ``sycl::ext::oneapi::experimental::ipc::memory::open()``. + * + * @param HandleData Byte buffer from DPCTLIPCMem_GetHandle. + * @param HandleDataSize Size of HandleData in bytes. + * @param CRef Sycl context for the receiving side. + * @param DRef Sycl device to map the memory on. + * @return A USM pointer to the IPC-mapped memory, or nullptr on failure. + * @ingroup IPCMemoryInterface + */ +DPCTL_API +__dpctl_give DPCTLSyclUSMRef +DPCTLIPCMem_OpenHandle(const char *HandleData, + size_t HandleDataSize, + __dpctl_keep const DPCTLSyclContextRef CRef, + __dpctl_keep const DPCTLSyclDeviceRef DRef); + +/*! + * @brief Close an IPC memory mapping opened by DPCTLIPCMem_OpenHandle. + * + * Wraps ``sycl::ext::oneapi::experimental::ipc::memory::close()``. + * + * @param MappedPtr The USM pointer returned by DPCTLIPCMem_OpenHandle. + * @param CRef Sycl context used when opening the handle. + * @ingroup IPCMemoryInterface + */ +DPCTL_API +void DPCTLIPCMem_CloseHandle(__dpctl_keep DPCTLSyclUSMRef MappedPtr, + __dpctl_keep const DPCTLSyclContextRef CRef); + +/*! + * @brief Free a handle data buffer returned by DPCTLIPCMem_GetHandle. + * + * @param Data Pointer previously returned via the DataOut parameter + * of DPCTLIPCMem_GetHandle. + * @ingroup IPCMemoryInterface + */ +DPCTL_API +void DPCTLIPCMem_FreeHandleData(char *Data); + +DPCTL_C_EXTERN_C_END diff --git a/libsyclinterface/source/dpctl_sycl_extension_interface.cpp b/libsyclinterface/source/dpctl_sycl_extension_interface.cpp index fa33fd8997..20638f4d11 100644 --- a/libsyclinterface/source/dpctl_sycl_extension_interface.cpp +++ b/libsyclinterface/source/dpctl_sycl_extension_interface.cpp @@ -63,6 +63,16 @@ bool DPCTLWorkGroupMemory_Available() #endif } +DPCTL_API +bool DPCTLIPCMem_Available() +{ +#ifdef SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION + return true; +#else + return false; +#endif +} + using raw_kernel_arg_t = std::vector; DPCTL_API diff --git a/libsyclinterface/source/dpctl_sycl_ipc_memory_interface.cpp b/libsyclinterface/source/dpctl_sycl_ipc_memory_interface.cpp new file mode 100644 index 0000000000..bb70c57594 --- /dev/null +++ b/libsyclinterface/source/dpctl_sycl_ipc_memory_interface.cpp @@ -0,0 +1,165 @@ +//===- dpctl_sycl_ipc_memory_interface.cpp - C API for SYCL IPC memory ----===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2026 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// This file implements the functions declared in +/// dpctl_sycl_ipc_memory_interface.h. +/// +//===----------------------------------------------------------------------===// + +#include "dpctl_sycl_ipc_memory_interface.h" +#include "Config/dpctl_config.h" +#include "dpctl_error_handlers.h" +#include "dpctl_sycl_type_casters.hpp" +#include +#include +#include +#include +#include + +using namespace sycl; + +// Support both the new namespace (ipc::memory, oneAPI >= 2026.1) and the +// deprecated namespace (ipc_memory, oneAPI 2026.0). +#if __has_include() +// New layout: ipc::memory namespace with separate ipc::handle/handle_data_t +namespace ipc = sycl::ext::oneapi::experimental::ipc::memory; +using ipc_handle_data_t = sycl::ext::oneapi::experimental::ipc::handle_data_t; +#else +// Old layout: everything in ipc_memory namespace +namespace ipc = sycl::ext::oneapi::experimental::ipc_memory; +using ipc_handle_data_t = ipc::handle_data_t; +#endif + +namespace +{ +static_assert(__SYCL_COMPILER_VERSION >= __SYCL_COMPILER_VERSION_REQUIRED, + "The compiler does not meet minimum version requirement"); + +using namespace dpctl::syclinterface; +} // end of anonymous namespace + +int DPCTLIPCMem_GetHandle(__dpctl_keep DPCTLSyclUSMRef Ptr, + __dpctl_keep const DPCTLSyclContextRef CRef, + char **DataOut, + size_t *SizeOut) +{ + if (!Ptr) { + error_handler("Input Ptr is nullptr.", __FILE__, __func__, __LINE__); + return 1; + } + if (!CRef) { + error_handler("Input CRef is nullptr.", __FILE__, __func__, __LINE__); + return 1; + } + if (!DataOut || !SizeOut) { + error_handler("Output pointers are nullptr.", __FILE__, __func__, + __LINE__); + return 1; + } + + try { + auto *RawPtr = unwrap(Ptr); + auto *Ctx = unwrap(CRef); + + // Obtain the IPC handle from the SYCL runtime. + auto Handle = ipc::get(RawPtr, *Ctx); + + // Copy handle data into a malloc'd buffer for the caller. + auto HandleData = Handle.data(); // std::vector + + size_t Size = HandleData.size(); + char *Buf = static_cast(std::malloc(Size)); + if (!Buf) { + error_handler("Failed to allocate handle data buffer.", __FILE__, + __func__, __LINE__); + return 1; + } + std::memcpy(Buf, HandleData.data(), Size); + + *DataOut = Buf; + *SizeOut = Size; + return 0; + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return 1; + } +} + +__dpctl_give DPCTLSyclUSMRef +DPCTLIPCMem_OpenHandle(const char *HandleData, + size_t HandleDataSize, + __dpctl_keep const DPCTLSyclContextRef CRef, + __dpctl_keep const DPCTLSyclDeviceRef DRef) +{ + if (!HandleData) { + error_handler("Input HandleData is nullptr.", __FILE__, __func__, + __LINE__); + return nullptr; + } + if (!CRef) { + error_handler("Input CRef is nullptr.", __FILE__, __func__, __LINE__); + return nullptr; + } + if (!DRef) { + error_handler("Input DRef is nullptr.", __FILE__, __func__, __LINE__); + return nullptr; + } + + try { + auto *Ctx = unwrap(CRef); + auto *Dev = unwrap(DRef); + + // Rebuild handle_data_t (vector) from the raw byte buffer. + ipc_handle_data_t HData( + reinterpret_cast(HandleData), + reinterpret_cast(HandleData) + HandleDataSize); + + void *MappedPtr = ipc::open(HData, *Ctx, *Dev); + return wrap(MappedPtr); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + return nullptr; + } +} + +void DPCTLIPCMem_CloseHandle(__dpctl_keep DPCTLSyclUSMRef MappedPtr, + __dpctl_keep const DPCTLSyclContextRef CRef) +{ + if (!MappedPtr) { + error_handler("Input MappedPtr is nullptr.", __FILE__, __func__, + __LINE__); + return; + } + if (!CRef) { + error_handler("Input CRef is nullptr.", __FILE__, __func__, __LINE__); + return; + } + + try { + auto *RawPtr = unwrap(MappedPtr); + auto *Ctx = unwrap(CRef); + ipc::close(RawPtr, *Ctx); + } catch (std::exception const &e) { + error_handler(e, __FILE__, __func__, __LINE__); + } +} + +void DPCTLIPCMem_FreeHandleData(char *Data) { std::free(Data); } diff --git a/libsyclinterface/source/dpctl_sycl_ipc_memory_stubs.cpp b/libsyclinterface/source/dpctl_sycl_ipc_memory_stubs.cpp new file mode 100644 index 0000000000..417fa80674 --- /dev/null +++ b/libsyclinterface/source/dpctl_sycl_ipc_memory_stubs.cpp @@ -0,0 +1,60 @@ +//===- dpctl_sycl_ipc_memory_stubs.cpp - Stub IPC functions ---------------===// +// +// Data Parallel Control (dpctl) +// +// Copyright 2020-2026 Intel Corporation +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +//===----------------------------------------------------------------------===// +/// +/// \file +/// Stub implementations of DPCTLIPCMem_* functions for builds where the +/// SYCL IPC memory extension is not available. These allow _memory.pyx to +/// always link; the functions return error codes at runtime. +/// +//===----------------------------------------------------------------------===// + +#include "dpctl_error_handlers.h" +#include "dpctl_sycl_ipc_memory_interface.h" +#include + +int DPCTLIPCMem_GetHandle(__dpctl_keep DPCTLSyclUSMRef, + __dpctl_keep const DPCTLSyclContextRef, + char **, + size_t *) +{ + error_handler("IPC memory not supported in this build.", __FILE__, __func__, + __LINE__); + return 1; +} + +__dpctl_give DPCTLSyclUSMRef +DPCTLIPCMem_OpenHandle(const char *, + size_t, + __dpctl_keep const DPCTLSyclContextRef, + __dpctl_keep const DPCTLSyclDeviceRef) +{ + error_handler("IPC memory not supported in this build.", __FILE__, __func__, + __LINE__); + return nullptr; +} + +void DPCTLIPCMem_CloseHandle(__dpctl_keep DPCTLSyclUSMRef, + __dpctl_keep const DPCTLSyclContextRef) +{ + error_handler("IPC memory not supported in this build.", __FILE__, __func__, + __LINE__); +} + +void DPCTLIPCMem_FreeHandleData(char *Data) { (void)Data; }