From aeca46dab9fe3a7fcfb5a2a08b5d9d83cfcddfd1 Mon Sep 17 00:00:00 2001 From: Simon Birkholz Date: Tue, 28 Jul 2026 09:22:39 +0200 Subject: [PATCH 1/4] initial draft of cuda ipc python wrapping implementation --- cuvis/Measurement.py | 52 ++++++++- cuvis/__init__.py | 146 +++++++++++++----------- cuvis/_dlpack.py | 91 +++++++++++++++ cuvis/cube_utils.py | 138 +++++++++++++++++++++++ cuvis/cuda.py | 107 ++++++++++++++++++ cuvis/cuda_import.py | 17 +++ cuvis/ipc.py | 258 +++++++++++++++++++++++++++++++++++++++++++ 7 files changed, 746 insertions(+), 63 deletions(-) create mode 100644 cuvis/_dlpack.py create mode 100644 cuvis/cuda.py create mode 100644 cuvis/cuda_import.py create mode 100644 cuvis/ipc.py diff --git a/cuvis/Measurement.py b/cuvis/Measurement.py index c08728e..3e5eb6b 100644 --- a/cuvis/Measurement.py +++ b/cuvis/Measurement.py @@ -13,7 +13,7 @@ _utc_from_epoch_ms, ) from .cuvis_types import DataFormat, ProcessingMode, ReferenceType -from .cube_utils import ImageData +from .cube_utils import ImageData, CudaImageData import cuvis.cuvis_types as internal @@ -36,6 +36,13 @@ class Measurement(object): session_info: SessionData # read-only frame_id: int # read-only + # When False, refresh() skips fetching image data to the host via + # cuvis_measurement_get_data_image. That host fetch moves a GPU-processed cube + # into host memory and frees the device copy, which makes the same-process CUDA + # path (get_cube_cuda) unavailable. Set False before processing when you intend + # to read the cube as CUDA device memory. Default True preserves normal behaviour. + _refresh_images = True + def __init__(self, base: int | str | Path): self._handle = None self._session = None @@ -106,6 +113,10 @@ def refresh(self) -> None: ) cdtype = cuvis_il.p_cuvis_data_type_t_value(pType) if cdtype == cuvis_il.data_type_image: + if not Measurement._refresh_images: + # Skip the host fetch so a GPU-processed cube stays in device + # memory and remains reachable via get_cube_cuda. + continue data = cuvis_il.cuvis_imbuffer_t() cuvis_il.cuvis_measurement_get_data_image(self._handle, key, data) # t0 = datetime.datetime.now() @@ -274,6 +285,45 @@ def cube(self) -> ImageData: "This Measurement does not have a cube saved. Consider reprocessing with a Processing Context." ) + def get_cube_cuda(self, key: str = "cube") -> CudaImageData: + """Image data as a device-resident CUDA buffer for same-process, zero-copy use. + + Returns a CudaImageData wrapping a CUVIS_CUDA_MEM handle; wrap it with + .to_torch() (DLPack) or __cuda_array_interface__. The underlying image data + must be backed by CUDA device memory (raises SDKException otherwise). + """ + buf = cuvis_il.cuvis_cuda_imbuffer_t() + if cuvis_il.status_ok != cuvis_il.cuvis_measurement_get_data_image_cuda( + self._handle, key, buf): + raise SDKException() + return CudaImageData(buf) + + def get_cube_cuda_ipc(self, key: str = "cube", backend: int = 0) -> CudaImageData: + """Image data as a shareable CUDA buffer for cross-process use. + + Fetches the device buffer (get_cube_cuda) and then creates an IPC export on it, + filling .descriptor with the transportable bytes; send those out-of-band to + another process and open them with cuvis.ipc.open. Keep the returned object alive + until the importer is done: it is the in-process pin (legacy IPC has no cross-process + refcount). backend selects the mechanism (0=auto, 1=pool, 2=legacy, 3=VMM); + make_ipc raises SDKException if the requested backend is unavailable on this device. + """ + cimg = self.get_cube_cuda(key) + cimg.make_ipc(backend) + return cimg + + def get_cube(self, key: str = "cube"): + """Cube via the active mode. + + When CUDA mode is enabled (cuvis.cuda.enable()), returns a device-resident + CudaImageData and raises SDKException if the device path is unavailable (no + silent host fallback). Otherwise returns the host ImageData. + """ + from . import cuda as _cuda # lazy: avoids an import cycle, cheap (no torch) + if _cuda.is_enabled(): + return self.get_cube_cuda(key) + return self.cube + @property def thumbnail(self): thumb = [val for key, val in self.data.items() if "view" in key] diff --git a/cuvis/__init__.py b/cuvis/__init__.py index 065de6c..270c613 100644 --- a/cuvis/__init__.py +++ b/cuvis/__init__.py @@ -1,65 +1,87 @@ -from .cuvis_aux import ( - SessionData, - Capabilities, - MeasurementFlags, - SensorInfo, - GPSData, - CalibrationInfo, -) -from .cuvis_types import ( - OperationMode, - HardwareState, - ProcessingMode, - PanSharpeningInterpolationType, - PanSharpeningAlgorithm, - TiffCompressionMode, - TiffFormat, - ComponentType, - ReferenceType, - SessionItemType, - SessionMergeMode, -) -from .Worker import Worker, WorkerResult -from .Viewer import Viewer -from .SessionFile import SessionFile -from .ProcessingContext import ProcessingContext -from .Measurement import Measurement -from .General import init, shutdown, version, set_log_level -from .sdk_settings import SdkSettings -from .FileWriteSettings import ( - GeneralExportSettings, - SaveArgs, - ProcessingArgs, - EnviExportSettings, - TiffExportSettings, - ViewExportSettings, - WorkerSettings, - ViewerSettings, -) -from . import binding -from .binding import BindingInfo, UnavailableSDKFunction -from .Export import CubeExporter, EnviExporter, TiffExporter, ViewExporter -from .Calibration import Calibration -from .AcquisitionContext import AcquisitionContext -from .cube_utils import ImageData -import os -import platform -import sys +"""cuvis Python SDK. -lib_dir = os.getenv("CUVIS") -if lib_dir is None: - print("CUVIS environmental variable is not set!") - sys.exit(1) -if platform.system() == "Windows": - os.add_dll_directory(lib_dir) - add_il = os.path.abspath(os.path.dirname(os.path.realpath(__file__))) - os.environ["PATH"] += os.pathsep + add_il - sys.path.append(str(add_il)) -elif platform.system() == "Linux": - os.environ["PATH"] = lib_dir + os.pathsep + os.environ["PATH"] -else: - raise NotImplementedError("Invalid operating system detected!") - # sys.exit(1) +The SDK surface (Measurement, ProcessingContext, init, ...) loads lazily on first access, +so `import cuvis` has no side effects and does not load the native binding. This lets the +import-safe `cuvis.ipc` cross-process consumer utilities be used in a process that never +initialized the SDK (no CUVIS env var, no cuvis.dll). The binding and its CUVIS/DLL setup +load only when an SDK symbol is actually used, via cuvis_il's own __init__. +""" +import importlib -del os, platform, sys +# Public name -> submodule that defines it. Loaded lazily via __getattr__ so that merely +# importing `cuvis` (or `cuvis.ipc`) never pulls in the CUDA-linked binding. +_LAZY = { + # cuvis_aux + "SessionData": "cuvis_aux", + "Capabilities": "cuvis_aux", + "MeasurementFlags": "cuvis_aux", + "SensorInfo": "cuvis_aux", + "GPSData": "cuvis_aux", + "CalibrationInfo": "cuvis_aux", + # cuvis_types + "OperationMode": "cuvis_types", + "HardwareState": "cuvis_types", + "ProcessingMode": "cuvis_types", + "PanSharpeningInterpolationType": "cuvis_types", + "PanSharpeningAlgorithm": "cuvis_types", + "TiffCompressionMode": "cuvis_types", + "TiffFormat": "cuvis_types", + "ComponentType": "cuvis_types", + "ReferenceType": "cuvis_types", + "SessionItemType": "cuvis_types", + "SessionMergeMode": "cuvis_types", + # core + "Worker": "Worker", + "WorkerResult": "Worker", + "Viewer": "Viewer", + "SessionFile": "SessionFile", + "ProcessingContext": "ProcessingContext", + "Measurement": "Measurement", + "init": "General", + "shutdown": "General", + "version": "General", + "set_log_level": "General", + # FileWriteSettings + "GeneralExportSettings": "FileWriteSettings", + "SaveArgs": "FileWriteSettings", + "ProcessingArgs": "FileWriteSettings", + "EnviExportSettings": "FileWriteSettings", + "TiffExportSettings": "FileWriteSettings", + "ViewExportSettings": "FileWriteSettings", + "WorkerSettings": "FileWriteSettings", + "ViewerSettings": "FileWriteSettings", + # Export + "CubeExporter": "Export", + "EnviExporter": "Export", + "TiffExporter": "Export", + "ViewExporter": "Export", + # binding + "BindingInfo": "binding", + "UnavailableSDKFunction": "binding", + # misc + "Calibration": "Calibration", + "AcquisitionContext": "AcquisitionContext", + "SdkSettings": "sdk_settings", + "ImageData": "cube_utils", + "CudaImageData": "cube_utils", +} + +# Submodules reachable as attributes. `ipc` is import-safe without the SDK; `binding` and +# `cuda` answer what the installed SDK provides and so must be reachable before init(). +_SUBMODULES = ("ipc", "cuda", "binding") + +__all__ = list(_LAZY) + list(_SUBMODULES) + + +def __getattr__(name): + """PEP 562 lazy attribute loader - imports the owning submodule on first access.""" + if name in _LAZY: + return getattr(importlib.import_module("." + _LAZY[name], __name__), name) + if name in _SUBMODULES: + return importlib.import_module("." + name, __name__) + raise AttributeError(f"module {__name__!r} has no attribute {name!r}") + + +def __dir__(): + return sorted(list(globals()) + __all__) diff --git a/cuvis/_dlpack.py b/cuvis/_dlpack.py new file mode 100644 index 0000000..66a950c --- /dev/null +++ b/cuvis/_dlpack.py @@ -0,0 +1,91 @@ +"""Minimal DLPack producer over a raw CUDA device pointer. + +Builds a DLManagedTensor PyCapsule via ctypes so torch.from_dlpack can consume a +foreign device buffer zero-copy. The capsule deleter calls a supplied on_delete +callback when torch releases the tensor, which is how buffer lifetime is tied to +the tensor (the callback drops the SDK reference that pins the buffer). + +Ported from utils_data_cuda/examples/torch_local.py; the only change is that the +deleter calls an injected callback instead of a hard-coded ctypes SDK. +""" + +import ctypes + +_kDLCUDA = 2 +_kDLInt = 0 +_kDLUInt = 1 +_kDLFloat = 2 + + +class _DLDevice(ctypes.Structure): + _fields_ = [("device_type", ctypes.c_int), ("device_id", ctypes.c_int)] + + +class _DLDataType(ctypes.Structure): + _fields_ = [("code", ctypes.c_uint8), ("bits", ctypes.c_uint8), ("lanes", ctypes.c_uint16)] + + +class _DLTensor(ctypes.Structure): + _fields_ = [("data", ctypes.c_void_p), ("device", _DLDevice), ("ndim", ctypes.c_int), + ("dtype", _DLDataType), ("shape", ctypes.POINTER(ctypes.c_int64)), + ("strides", ctypes.POINTER(ctypes.c_int64)), ("byte_offset", ctypes.c_uint64)] + + +class _DLManagedTensor(ctypes.Structure): + pass + + +_DELETER = ctypes.CFUNCTYPE(None, ctypes.POINTER(_DLManagedTensor)) +_DLManagedTensor._fields_ = [("dl_tensor", _DLTensor), ("manager_ctx", ctypes.c_void_p), + ("deleter", _DELETER)] + +_pycapsule_new = ctypes.pythonapi.PyCapsule_New +_pycapsule_new.restype = ctypes.py_object +_pycapsule_new.argtypes = [ctypes.c_void_p, ctypes.c_char_p, ctypes.c_void_p] + +# Keep ctypes structs / deleters / callbacks alive until torch invokes the deleter. +_LIVE = {} + + +class CudaDlpack: + """DLPack producer over (ptr, nbytes) on a CUDA device. on_delete() runs when + torch releases the tensor. Exposes a flat uint8 buffer; reshape/retype in torch.""" + + def __init__(self, ptr, nbytes, device, on_delete): + self._ptr = int(ptr) + self._n = int(nbytes) + self._dev = int(device) + self._on_delete = on_delete + + def __dlpack_device__(self): + return (_kDLCUDA, self._dev) + + def __dlpack__(self, stream=None, max_version=None, dl_device=None, copy=None): + shape = (ctypes.c_int64 * 1)(self._n) + mt = _DLManagedTensor() + mt.dl_tensor.data = ctypes.c_void_p(self._ptr) + mt.dl_tensor.device = _DLDevice(_kDLCUDA, self._dev) + mt.dl_tensor.ndim = 1 + mt.dl_tensor.dtype = _DLDataType(_kDLUInt, 8, 1) + mt.dl_tensor.shape = shape + mt.dl_tensor.strides = None + mt.dl_tensor.byte_offset = 0 + + on_delete = self._on_delete + key = id(mt) + + def _del(_p): + try: + on_delete() + finally: + _LIVE.pop(key, None) + + deleter = _DELETER(_del) + mt.deleter = deleter + _LIVE[key] = (mt, shape, deleter) + return _pycapsule_new(ctypes.byref(mt), b"dltensor", None) + + +def make_cuda_dlpack(ptr, nbytes, device, on_delete): + """Return an object that torch.from_dlpack consumes into a zero-copy CUDA tensor.""" + return CudaDlpack(ptr, nbytes, device, on_delete) diff --git a/cuvis/cube_utils.py b/cuvis/cube_utils.py index cfb5fef..a9a0df4 100644 --- a/cuvis/cube_utils.py +++ b/cuvis/cube_utils.py @@ -4,6 +4,7 @@ import numpy as np import operator from .cuvis_aux import SDKException +from .cuvis_types import DataFormat _IMBUF_READERS = { 1: cuvis_il.cuvis_read_imbuf_uint8, @@ -418,3 +419,140 @@ def apply(self, other): ImageData.__abs__ = lambda self: self._wrap(abs(self.array)) del _op, _reflected, _method +class CudaImageData(object): + """Device-resident image data backed by a shareable CUDA buffer. + + Wraps a CUVIS_CUDA_MEM handle plus geometry, exposing the device memory as a + zero-copy CUDA tensor (torch via DLPack, or any consumer via + __cuda_array_interface__). Unlike ImageData, no host copy is made. + + Lifetime: this object owns the CUVIS_CUDA_MEM handle and frees it in __del__. + to_torch() takes its own SDK reference so the returned tensor can outlive this + object; the __cuda_array_interface__ path does not, so keep this object alive + until such a consumer is done with it. + + For cross-process sharing, make_ipc() (called by Measurement.get_cube_cuda_ipc) + creates a CUVIS_CUDA_IPC handle and fills .descriptor. That IPC handle is an + independent reference to the same buffer, freed in __del__; freeing the mem handle + while it is open would not release the memory. + """ + + _TYPESTR = {1: "|u1", 2: " bytes: + """A single transportable blob (IPC descriptor + geometry) for a consumer process. + + Send these bytes out-of-band; the consumer opens them with cuvis.ipc.open(payload) + and gets a correctly shaped/typed tensor. Calls make_ipc(backend) if not already done + (backend: 0=auto, 1=pool, 2=legacy, 3=VMM). Keep this CudaImageData alive until the + consumer is finished (legacy IPC has no cross-process refcount). + """ + if self.descriptor is None: + self.make_ipc(backend) + from . import ipc + return ipc.pack_payload(self.descriptor, self.width, self.height, self.channels, self._format) + + def __del__(self): + try: + if self._ipc_handle is not None: + cuvis_il.cuvis_cuda_ipc_handle_free(self._ipc_handle) + except Exception: + pass + try: + cuvis_il.cuvis_cuda_mem_free(self._handle) + except Exception: + pass diff --git a/cuvis/cuda.py b/cuvis/cuda.py new file mode 100644 index 0000000..a6703f3 --- /dev/null +++ b/cuvis/cuda.py @@ -0,0 +1,107 @@ +"""Opt-in CUDA support for cuvis. + +CUDA is off by default. Reach it explicitly: + + from cuvis import cuda + + caps = cuda.capabilities() # inspect what this binary + environment support + if caps.same_process: + cuda.enable() # BEFORE loading/processing measurements + ... + t = mesu.get_cube().to_torch() # zero-copy device tensor (DLPack, lifecycle-managed) + +`enable()` also disables the host auto-refresh (`Measurement._refresh_images = False`) so a +GPU-processed cube stays on the device instead of being copied to host and freed. + +Note: the shipped cuvis binding links the CUDA runtime, so `import cuvis` already requires a +CUDA runtime to be present. This module gates the CUDA *feature surface* and the optional +`torch` / `cuda-python` consumer dependencies, not the binding's own runtime dependency. +""" +import importlib.util +from typing import NamedTuple + +from cuvis_il import cuvis_il + +from .cube_utils import CudaImageData # re-exported; the device-image type + +# Backend codes (match cuvis.h CUVIS_CUDA_IPC_BACKEND_*). +_BACKEND_NONE = 0 +_BACKEND_POOL = 1 +_BACKEND_LEGACY = 2 +_BACKEND_VMM = 3 + +_enabled = False + + +class CudaCapabilities(NamedTuple): + """What the loaded cuvis binary and the current Python environment support. + + same_process: the CUDA boundary responds (same-process device sharing works). + ipc_pool: exportable memory pool backend (zero-copy cross-process, needs an exportable pool). + ipc_legacy: legacy cudaIpc backend (copy-on-export; available on most WDDM/Linux GPUs). + ipc_vmm: driver-API VMM backend (copy-on-export, with an exportable handle type). + torch / cuda_python: optional consumer packages importable (checked, not imported). + """ + same_process: bool + ipc_pool: bool + ipc_legacy: bool + ipc_vmm: bool + torch: bool + cuda_python: bool + + @property + def any_ipc(self) -> bool: + return self.ipc_pool or self.ipc_legacy or self.ipc_vmm + + +def _backend_available(code: int) -> bool: + try: + p = cuvis_il.new_p_int() + if cuvis_il.cuvis_cuda_ipc_backend_available(code, p) != cuvis_il.status_ok: + return False + return bool(cuvis_il.p_int_value(p)) + except Exception: + # Symbol absent or a CUDA-less binary: treat as unavailable. + return False + + +def capabilities() -> CudaCapabilities: + """Probe CUDA capabilities at runtime. Safe to call before enable().""" + return CudaCapabilities( + same_process=_backend_available(_BACKEND_NONE), + ipc_pool=_backend_available(_BACKEND_POOL), + ipc_legacy=_backend_available(_BACKEND_LEGACY), + ipc_vmm=_backend_available(_BACKEND_VMM), + torch=importlib.util.find_spec("torch") is not None, + cuda_python=importlib.util.find_spec("cuda.bindings") is not None, + ) + + +def enable() -> None: + """Turn on CUDA mode. Call this BEFORE loading or processing measurements. + + Routes `Measurement.get_cube()` through the device path and disables the host + auto-refresh so the GPU cube is kept on the device. Raises RuntimeError if this + cuvis build has no CUDA support. + """ + global _enabled + if not capabilities().same_process: + raise RuntimeError("this cuvis build has no CUDA support") + from .Measurement import Measurement + Measurement._refresh_images = False + _enabled = True + + +def disable() -> None: + """Turn off CUDA mode and restore the host auto-refresh.""" + global _enabled + from .Measurement import Measurement + Measurement._refresh_images = True + _enabled = False + + +def is_enabled() -> bool: + return _enabled + + +__all__ = ["CudaCapabilities", "capabilities", "enable", "disable", "is_enabled", "CudaImageData"] diff --git a/cuvis/cuda_import.py b/cuvis/cuda_import.py new file mode 100644 index 0000000..af91a7d --- /dev/null +++ b/cuvis/cuda_import.py @@ -0,0 +1,17 @@ +"""Deprecated: use cuvis.ipc instead. + +Kept as a thin shim for back-compat. `open_ipc(descriptor)` maps a raw descriptor (you supply +geometry via .tensor(dtype, shape)); the preferred path is a bundled payload via +cuvis.ipc.open(payload). See cuvis/ipc.py. +""" +from .ipc import ImportedCube as ImportedIpcTensor # noqa: F401 (back-compat name) +from .ipc import open_descriptor, BACKEND_NONE, BACKEND_POOL, BACKEND_LEGACY, BACKEND_VMM # noqa: F401 + + +def open_ipc(descriptor_bytes) -> ImportedIpcTensor: + """Deprecated alias of cuvis.ipc.open_descriptor().""" + return open_descriptor(descriptor_bytes) + + +__all__ = ["ImportedIpcTensor", "open_ipc", "BACKEND_NONE", "BACKEND_POOL", + "BACKEND_LEGACY", "BACKEND_VMM"] diff --git a/cuvis/ipc.py b/cuvis/ipc.py new file mode 100644 index 0000000..672fc6a --- /dev/null +++ b/cuvis/ipc.py @@ -0,0 +1,258 @@ +"""Import-safe cross-process CUDA IPC consumer utilities. + +Use this in a SEPARATE process that receives a cuvis IPC payload - it does NOT initialize +or link the cuvis SDK. It needs only cuda-python (`cuda.bindings`) and torch, imported lazily. +`import cuvis.ipc` works with no CUVIS env var and no cuvis.dll present. + +Producer (in the cuvis process): + cimg = mesu.get_cube_cuda_ipc() # keep this alive until the consumer is done + payload = cimg.export_payload() # a single transportable bytes blob (descriptor + geometry) + +Consumer (this module, any process): + import cuvis.ipc as ipc + with ipc.open(payload) as cube: + t = cube.to_torch() # correctly shaped/typed zero-copy CUDA tensor + ... # use t inside the block + # leaving the block releases this process's mapping (does not free the exporter's memory) + +The exporting process must outlive this importer: legacy IPC has no cross-process refcount. +""" +import struct + +# --- IPC descriptor wire format (locked; mirrors cuvis_cuda_ipc_descriptor_t) --- +# 48-byte header + 64-byte blob (pool OS handle) at offset 48, then ptr_blob_len(+pad) and a +# 64-byte ptr_blob (cudaMemPoolPtrExportData) at offset 120. Total 184. +_HEAD = struct.Struct(" torch dtype name / __cuda_array_interface__ typestr (1/2/3/4 = u8/u16/u32/f32) +_TORCH_DTYPE = {1: "uint8", 2: "uint16", 3: "uint32", 4: "float32"} +_TYPESTR = {1: "|u1", 2: " bytes: + """Bundle an IPC descriptor and cube geometry into one transportable blob.""" + if len(descriptor) != _DESC_LEN: + raise ValueError(f"descriptor must be {_DESC_LEN} bytes, got {len(descriptor)}") + return _PAYLOAD_HDR.pack(_MAGIC, _VERSION, int(width), int(height), int(channels), + int(format_code)) + bytes(descriptor) + + +def _unpack_payload(payload: bytes): + magic, version, width, height, channels, fmt = _PAYLOAD_HDR.unpack_from(payload, 0) + if magic != _MAGIC: + raise ValueError("not a cuvis IPC payload (bad magic)") + if version != _VERSION: + raise ValueError(f"unsupported cuvis IPC payload version {version}") + descriptor = bytes(payload[_PAYLOAD_HDR.size:]) + return (width, height, channels, fmt), descriptor + + +def _ck(ret, what): + err = ret[0] + if int(err) != 0: + raise RuntimeError(f"{what} failed: {err}") + return ret[1:] if len(ret) > 1 else None + + +class _CudaArray: + def __init__(self, ptr, nbytes): + self.__cuda_array_interface__ = { + "shape": (nbytes,), "typestr": "|u1", "data": (int(ptr), False), "version": 3, + } + + +class ImportedCube: + """A mapped IPC buffer in the consumer process. Use as a context manager. + + open()/open_descriptor() return this; leaving the `with` block releases the mapping. + to_torch() (preferred) and __cuda_array_interface__ produce zero-copy views that are + valid only while the block is open. + """ + + def __init__(self, descriptor_bytes: bytes, shape=None, format_code=None): + (self.backend, self.device, self.htype, blob_len, + self.size, self.alloc, self.offset, self.pid) = _HEAD.unpack_from(descriptor_bytes, 0) + if blob_len > _BLOB_MAX: + raise ValueError(f"blob_len {blob_len} exceeds {_BLOB_MAX}") + self._blob = bytes(descriptor_bytes[_BLOB_OFF:_BLOB_OFF + blob_len]) + (ptr_blob_len,) = struct.unpack_from(" _PTR_BLOB_MAX: + raise ValueError(f"ptr_blob_len {ptr_blob_len} exceeds {_PTR_BLOB_MAX}") + self._ptr_blob = bytes(descriptor_bytes[_PTR_BLOB_OFF:_PTR_BLOB_OFF + ptr_blob_len]) + self._shape = tuple(shape) if shape is not None else None + self._format = format_code + self._close = None + + from cuda.bindings import runtime + runtime.cudaSetDevice(self.device) + + if self.backend == BACKEND_POOL: + self._ptr, self._close = self._open_pool() + elif self.backend == BACKEND_LEGACY: + self._ptr, self._close = self._open_legacy() + elif self.backend == BACKEND_VMM: + self._ptr, self._close = self._open_vmm() + else: + raise NotImplementedError(f"backend {self.backend} is not importable cross-process") + self._ptr += self.offset # 0 for pool/legacy/vmm (import returns the exact base pointer) + + @property + def shape(self): + return self._shape + + @property + def device_ptr(self): + return self._ptr + + def _open_pool(self): + # IPC-capable memory pool: import the pool from its OS shareable handle, grant this + # device access, then import the exact pointer from the per-allocation export data. + from cuda.bindings import runtime + if self.htype == H_WIN32_KMT: + htype = runtime.cudaMemAllocationHandleType.cudaMemHandleTypeWin32Kmt + elif self.htype == H_POSIX_FD: + htype = runtime.cudaMemAllocationHandleType.cudaMemHandleTypePosixFileDescriptor + else: + raise NotImplementedError( + f"pool handle_type {self.htype} needs out-of-band duplication (DuplicateHandle / SCM_RIGHTS)") + + handle_val = int.from_bytes(self._blob, "little") + (pool,) = _ck(runtime.cudaMemPoolImportFromShareableHandle(handle_val, htype, 0), + "cudaMemPoolImportFromShareableHandle") + + acc = runtime.cudaMemAccessDesc() + acc.location.type = runtime.cudaMemLocationType.cudaMemLocationTypeDevice + acc.location.id = self.device + acc.flags = runtime.cudaMemAccessFlags.cudaMemAccessFlagsProtReadWrite + _ck(runtime.cudaMemPoolSetAccess(pool, [acc], 1), "cudaMemPoolSetAccess") + + export_data = runtime.cudaMemPoolPtrExportData() + export_data.reserved = self._ptr_blob.ljust(_PTR_BLOB_MAX, b"\x00") + (ptr,) = _ck(runtime.cudaMemPoolImportPointer(pool, export_data), "cudaMemPoolImportPointer") + + def close(): + runtime.cudaFree(ptr) # release this process's imported pointer + runtime.cudaMemPoolDestroy(pool) # release the imported pool handle + + return int(ptr), close + + def _open_legacy(self): + # Legacy cudaIpc: the descriptor blob is a self-contained 64-byte cudaIpcMemHandle_t. + from cuda.bindings import runtime + h = runtime.cudaIpcMemHandle_t() + h.reserved = self._blob.ljust(_BLOB_MAX, b"\x00") + (ptr,) = _ck(runtime.cudaIpcOpenMemHandle(h, runtime.cudaIpcMemLazyEnablePeerAccess), + "cudaIpcOpenMemHandle") + + def close(): + runtime.cudaIpcCloseMemHandle(ptr) # release this process's mapping (not the exporter's) + + return int(ptr), close + + def _open_vmm(self): + # VMM: import the generic handle from its OS shareable handle, reserve + map + grant access. + from cuda.bindings import driver + driver.cuInit(0) + if self.htype == H_WIN32_KMT: + htype = driver.CUmemAllocationHandleType.CU_MEM_HANDLE_TYPE_WIN32_KMT + elif self.htype == H_POSIX_FD: + htype = driver.CUmemAllocationHandleType.CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR + elif self.htype == H_WIN32: + htype = driver.CUmemAllocationHandleType.CU_MEM_HANDLE_TYPE_WIN32 + else: + raise NotImplementedError(f"vmm handle_type {self.htype} not supported") + + handle_val = int.from_bytes(self._blob, "little") + (gen,) = _ck(driver.cuMemImportFromShareableHandle(handle_val, htype), + "cuMemImportFromShareableHandle") + size = self.alloc + (ptr,) = _ck(driver.cuMemAddressReserve(size, 0, 0, 0), "cuMemAddressReserve") + _ck(driver.cuMemMap(ptr, size, 0, gen, 0), "cuMemMap") + + acc = driver.CUmemAccessDesc() + acc.location.type = driver.CUmemLocationType.CU_MEM_LOCATION_TYPE_DEVICE + acc.location.id = self.device + acc.flags = driver.CUmemAccess_flags.CU_MEM_ACCESS_FLAGS_PROT_READWRITE + _ck(driver.cuMemSetAccess(ptr, size, [acc], 1), "cuMemSetAccess") + + def close(): + driver.cuMemUnmap(ptr, size) + driver.cuMemAddressFree(ptr, size) + driver.cuMemRelease(gen) + + return int(ptr), close + + def _torch_dtype(self, torch): + return None if self._format is None else getattr(torch, _TORCH_DTYPE[self._format]) + + def to_torch(self, dtype=None, shape=None): + """Zero-copy CUDA torch tensor over the mapped buffer, valid inside the with-block. + + With open(payload), dtype and shape are taken from the payload geometry; overrides + may be passed explicitly (needed after open_descriptor without geometry).""" + import torch + n = self.size - self.offset + t = torch.as_tensor(_CudaArray(self._ptr, n), device=f"cuda:{self.device}") + dt = dtype if dtype is not None else self._torch_dtype(torch) + if dt is not None and dt != torch.uint8: + t = t.view(dt) + sh = shape if shape is not None else self._shape + if sh is not None: + t = t.reshape(*sh) + return t + + tensor = to_torch # back-compat alias + + @property + def __cuda_array_interface__(self): + if self._shape is None or self._format is None: + raise RuntimeError( + "__cuda_array_interface__ needs geometry; open via open(payload) or use to_torch(dtype, shape)") + return { + "shape": self._shape, + "typestr": _TYPESTR[self._format], + "data": (int(self._ptr), False), + "version": 3, + } + + def close(self): + if self._close is not None: + self._close() + self._close = None + + def __enter__(self): + return self + + def __exit__(self, *exc): + self.close() + + +def open(payload: bytes) -> ImportedCube: + """Open a payload from CudaImageData.export_payload(): maps the buffer and carries geometry.""" + (width, height, channels, fmt), descriptor = _unpack_payload(payload) + return ImportedCube(descriptor, shape=(height, width, channels), format_code=fmt) + + +def open_descriptor(descriptor_bytes: bytes, shape=None) -> ImportedCube: + """Advanced: open a raw 184-byte descriptor when you transport geometry yourself. + + Pass shape here (or dtype/shape to .to_torch()); prefer open(payload) for the easy path.""" + return ImportedCube(descriptor_bytes, shape=shape) + + +__all__ = ["ImportedCube", "open", "open_descriptor", "pack_payload", + "BACKEND_NONE", "BACKEND_POOL", "BACKEND_LEGACY", "BACKEND_VMM"] From c7b1d55873d0ef26529942b350a948787156f6af Mon Sep 17 00:00:00 2001 From: Simon Birkholz Date: Wed, 19 Aug 2026 17:28:00 +0200 Subject: [PATCH 2/4] fixup after rebase --- .gitignore | 1 + CHANGELOG.md | 26 +++++++ cuvis/Measurement.py | 10 ++- cuvis/__init__.py | 147 +++++++++++++++-------------------- cuvis/_dlpack.py | 25 ++++-- cuvis/binding.py | 47 +++++++---- cuvis/cube_utils.py | 41 +++++++--- cuvis/cuda.py | 147 +++++++++++++++++++++++++++-------- cuvis/cuda_import.py | 17 ---- cuvis/ipc.py => cuvis_ipc.py | 123 +++++++++++++++++++++-------- pyproject.toml | 3 + tests/test_binding.py | 65 ++++++++++++++++ tests/test_cuda.py | 74 ++++++++++++++++++ tests/test_cuvis_ipc.py | 63 +++++++++++++++ 14 files changed, 585 insertions(+), 204 deletions(-) delete mode 100644 cuvis/cuda_import.py rename cuvis/ipc.py => cuvis_ipc.py (73%) create mode 100644 tests/test_binding.py create mode 100644 tests/test_cuda.py create mode 100644 tests/test_cuvis_ipc.py diff --git a/.gitignore b/.gitignore index 43e869c..9793443 100644 --- a/.gitignore +++ b/.gitignore @@ -7,6 +7,7 @@ /cuvis/cuvis_il.py /cuvis/_cuvis_pyil.pyd /venv +/__pycache__ /cuvis/__pycache__ /cuvis/git-hash.txt diff --git a/CHANGELOG.md b/CHANGELOG.md index b83127d..2da645a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -20,6 +20,13 @@ Pre-releases (`b*`, `rc*`) are not listed. - `CI` - `scripts/check_changelog.py` validates this file's structure (header format, allowed section names, descending versions) and the tag/version/changelog agreement at release time. - `CONTRIBUTING.md` - documents the branch model, the version scheme, the changelog conventions and the release checklist. - `cuvis.BindingInfo` - new frozen dataclass with the fields `built_against: str`, `library_version: str`, `library_path: str` and `missing_symbols: Tuple[str, ...]`, the read-only property `is_complete: bool`, and a `__str__` rendering a report fit for a bug report. +- `cuvis.CudaImageData` - new class, device-resident image data backed by a shareable CUDA buffer, exposed zero-copy through DLPack or `__cuda_array_interface__` with no host copy. +- `cuvis.CudaImageData.export_payload` - new method, returns `bytes`. +- `cuvis.CudaImageData.make_ipc` - new method, returns `bytes`. +- `cuvis.CudaImageData.to_torch` - new method, returns `torch.Tensor`. +- `cuvis.Measurement.get_cube` - new method, returns `Union[ImageData, CudaImageData]` depending on whether `cuvis.cuda.enable` was called. +- `cuvis.Measurement.get_cube_cuda` - new method, returns `CudaImageData`. +- `cuvis.Measurement.get_cube_cuda_ipc` - new method, returns `CudaImageData`. - `cuvis.SdkSettings` - new class, a `MutableMapping` of setting id to value that writes the SDK's `cuvis.settings` file, so the SDK configuration can be built in Python instead of maintained by hand. Values are stored as strings: `bool` becomes `true`/`false`, an `Enum` becomes its value, anything else goes through `str()`, and `None` drops the entry. - `cuvis.SdkSettings.__enter__`, `cuvis.SdkSettings.__exit__` - new methods; entering the context serializes the settings into a temporary directory and returns its path as `str`, leaving the context removes the directory. @@ -32,6 +39,25 @@ Pre-releases (`b*`, `rc*`) are not listed. - `cuvis.binding.info` - new function, returns `BindingInfo`. - `cuvis.binding.missing_symbols` - new function, returns `FrozenSet[str]`. - `cuvis.binding.require` - new function, raises `UnavailableSDKFunction` naming whichever of the given functions the installed cuvis library does not provide. +- `cuvis.binding.unavailable` - new function, returns `Tuple[str, ...]`. + A function the binding never exposed is unusable as well, so availability cannot be answered from the reported-missing list alone. +- `cuvis.cuda` - new module gating the optional CUDA feature surface; CUDA stays off until `cuvis.cuda.enable` is called. +- `cuvis.cuda.BACKEND_NONE`, `cuvis.cuda.BACKEND_POOL`, `cuvis.cuda.BACKEND_LEGACY`, `cuvis.cuda.BACKEND_VMM` - new constants, the IPC backend codes from `cuvis.h`. +- `cuvis.cuda.CudaCapabilities` - new `NamedTuple` with the `bool` fields `same_process`, `ipc_pool`, `ipc_legacy`, `ipc_vmm`, `torch` and `cuda_python`, and the read-only property `any_ipc: bool`. +- `cuvis.cuda.capabilities` - new function, returns `CudaCapabilities`. +- `cuvis.cuda.disable` - new function. +- `cuvis.cuda.enable` - new function. +- `cuvis.cuda.is_enabled` - new function, returns `bool`. +- `cuvis.cuda.require_device` - new function, raises `UnavailableSDKFunction` unless the installed library provides the same-process device path. +- `cuvis.cuda.require_ipc` - new function, raises `UnavailableSDKFunction` unless the installed library provides the cross-process export path. +- `cuvis_ipc` - new top-level module, the consumer side of cross-process CUDA IPC. + It sits outside the `cuvis` package because importing `cuvis` requires the SDK and a consumer process does not have one; its only import is `struct`. +- `cuvis_ipc.ImportedCube` - new class, a mapped IPC buffer in the consumer process, usable as a context manager. +- `cuvis_ipc.open` - new function, returns `ImportedCube`. +- `cuvis_ipc.open_descriptor` - new function, returns `ImportedCube`. +- `cuvis_ipc.pack_payload` - new function, returns `bytes`. +- `pyproject.toml` - `py-modules` declaring the top-level `cuvis_ipc`. +- `tests/` - `test_binding.py`, `test_cuda.py` and `test_cuvis_ipc.py`. ### Changed diff --git a/cuvis/Measurement.py b/cuvis/Measurement.py index 3e5eb6b..b0c08a3 100644 --- a/cuvis/Measurement.py +++ b/cuvis/Measurement.py @@ -14,6 +14,7 @@ ) from .cuvis_types import DataFormat, ProcessingMode, ReferenceType from .cube_utils import ImageData, CudaImageData +from . import cuda import cuvis.cuvis_types as internal @@ -292,9 +293,11 @@ def get_cube_cuda(self, key: str = "cube") -> CudaImageData: .to_torch() (DLPack) or __cuda_array_interface__. The underlying image data must be backed by CUDA device memory (raises SDKException otherwise). """ + cuda.require_device() buf = cuvis_il.cuvis_cuda_imbuffer_t() if cuvis_il.status_ok != cuvis_il.cuvis_measurement_get_data_image_cuda( - self._handle, key, buf): + self._handle, key, buf + ): raise SDKException() return CudaImageData(buf) @@ -303,7 +306,7 @@ def get_cube_cuda_ipc(self, key: str = "cube", backend: int = 0) -> CudaImageDat Fetches the device buffer (get_cube_cuda) and then creates an IPC export on it, filling .descriptor with the transportable bytes; send those out-of-band to - another process and open them with cuvis.ipc.open. Keep the returned object alive + another process and open them with cuvis_ipc.open. Keep the returned object alive until the importer is done: it is the in-process pin (legacy IPC has no cross-process refcount). backend selects the mechanism (0=auto, 1=pool, 2=legacy, 3=VMM); make_ipc raises SDKException if the requested backend is unavailable on this device. @@ -319,8 +322,7 @@ def get_cube(self, key: str = "cube"): CudaImageData and raises SDKException if the device path is unavailable (no silent host fallback). Otherwise returns the host ImageData. """ - from . import cuda as _cuda # lazy: avoids an import cycle, cheap (no torch) - if _cuda.is_enabled(): + if cuda.is_enabled(): return self.get_cube_cuda(key) return self.cube diff --git a/cuvis/__init__.py b/cuvis/__init__.py index 270c613..23388e6 100644 --- a/cuvis/__init__.py +++ b/cuvis/__init__.py @@ -1,87 +1,66 @@ -"""cuvis Python SDK. +from .cuvis_aux import ( + SessionData, + Capabilities, + MeasurementFlags, + SensorInfo, + GPSData, + CalibrationInfo, +) +from .cuvis_types import ( + OperationMode, + HardwareState, + ProcessingMode, + PanSharpeningInterpolationType, + PanSharpeningAlgorithm, + TiffCompressionMode, + TiffFormat, + ComponentType, + ReferenceType, + SessionItemType, + SessionMergeMode, +) +from .Worker import Worker, WorkerResult +from .Viewer import Viewer +from .SessionFile import SessionFile +from .ProcessingContext import ProcessingContext +from .Measurement import Measurement +from .General import init, shutdown, version, set_log_level +from .sdk_settings import SdkSettings +from .FileWriteSettings import ( + GeneralExportSettings, + SaveArgs, + ProcessingArgs, + EnviExportSettings, + TiffExportSettings, + ViewExportSettings, + WorkerSettings, + ViewerSettings, +) +from . import binding +from .binding import BindingInfo, UnavailableSDKFunction +from .Export import CubeExporter, EnviExporter, TiffExporter, ViewExporter +from .Calibration import Calibration +from .AcquisitionContext import AcquisitionContext +from .cube_utils import ImageData, CudaImageData +from . import cuda +import os +import platform +import sys -The SDK surface (Measurement, ProcessingContext, init, ...) loads lazily on first access, -so `import cuvis` has no side effects and does not load the native binding. This lets the -import-safe `cuvis.ipc` cross-process consumer utilities be used in a process that never -initialized the SDK (no CUVIS env var, no cuvis.dll). The binding and its CUVIS/DLL setup -load only when an SDK symbol is actually used, via cuvis_il's own __init__. -""" +lib_dir = os.getenv("CUVIS") +if lib_dir is None: + print("CUVIS environmental variable is not set!") + sys.exit(1) +if platform.system() == "Windows": + os.add_dll_directory(lib_dir) + add_il = os.path.abspath(os.path.dirname(os.path.realpath(__file__))) + os.environ["PATH"] += os.pathsep + add_il + sys.path.append(str(add_il)) +elif platform.system() == "Linux": + os.environ["PATH"] = lib_dir + os.pathsep + os.environ["PATH"] +else: + raise NotImplementedError("Invalid operating system detected!") + # sys.exit(1) -import importlib -# Public name -> submodule that defines it. Loaded lazily via __getattr__ so that merely -# importing `cuvis` (or `cuvis.ipc`) never pulls in the CUDA-linked binding. -_LAZY = { - # cuvis_aux - "SessionData": "cuvis_aux", - "Capabilities": "cuvis_aux", - "MeasurementFlags": "cuvis_aux", - "SensorInfo": "cuvis_aux", - "GPSData": "cuvis_aux", - "CalibrationInfo": "cuvis_aux", - # cuvis_types - "OperationMode": "cuvis_types", - "HardwareState": "cuvis_types", - "ProcessingMode": "cuvis_types", - "PanSharpeningInterpolationType": "cuvis_types", - "PanSharpeningAlgorithm": "cuvis_types", - "TiffCompressionMode": "cuvis_types", - "TiffFormat": "cuvis_types", - "ComponentType": "cuvis_types", - "ReferenceType": "cuvis_types", - "SessionItemType": "cuvis_types", - "SessionMergeMode": "cuvis_types", - # core - "Worker": "Worker", - "WorkerResult": "Worker", - "Viewer": "Viewer", - "SessionFile": "SessionFile", - "ProcessingContext": "ProcessingContext", - "Measurement": "Measurement", - "init": "General", - "shutdown": "General", - "version": "General", - "set_log_level": "General", - # FileWriteSettings - "GeneralExportSettings": "FileWriteSettings", - "SaveArgs": "FileWriteSettings", - "ProcessingArgs": "FileWriteSettings", - "EnviExportSettings": "FileWriteSettings", - "TiffExportSettings": "FileWriteSettings", - "ViewExportSettings": "FileWriteSettings", - "WorkerSettings": "FileWriteSettings", - "ViewerSettings": "FileWriteSettings", - # Export - "CubeExporter": "Export", - "EnviExporter": "Export", - "TiffExporter": "Export", - "ViewExporter": "Export", - # binding - "BindingInfo": "binding", - "UnavailableSDKFunction": "binding", - # misc - "Calibration": "Calibration", - "AcquisitionContext": "AcquisitionContext", - "SdkSettings": "sdk_settings", - "ImageData": "cube_utils", - "CudaImageData": "cube_utils", -} - -# Submodules reachable as attributes. `ipc` is import-safe without the SDK; `binding` and -# `cuda` answer what the installed SDK provides and so must be reachable before init(). -_SUBMODULES = ("ipc", "cuda", "binding") - -__all__ = list(_LAZY) + list(_SUBMODULES) - - -def __getattr__(name): - """PEP 562 lazy attribute loader - imports the owning submodule on first access.""" - if name in _LAZY: - return getattr(importlib.import_module("." + _LAZY[name], __name__), name) - if name in _SUBMODULES: - return importlib.import_module("." + name, __name__) - raise AttributeError(f"module {__name__!r} has no attribute {name!r}") - - -def __dir__(): - return sorted(list(globals()) + __all__) +del os, platform, sys diff --git a/cuvis/_dlpack.py b/cuvis/_dlpack.py index 66a950c..1084e7b 100644 --- a/cuvis/_dlpack.py +++ b/cuvis/_dlpack.py @@ -22,13 +22,23 @@ class _DLDevice(ctypes.Structure): class _DLDataType(ctypes.Structure): - _fields_ = [("code", ctypes.c_uint8), ("bits", ctypes.c_uint8), ("lanes", ctypes.c_uint16)] + _fields_ = [ + ("code", ctypes.c_uint8), + ("bits", ctypes.c_uint8), + ("lanes", ctypes.c_uint16), + ] class _DLTensor(ctypes.Structure): - _fields_ = [("data", ctypes.c_void_p), ("device", _DLDevice), ("ndim", ctypes.c_int), - ("dtype", _DLDataType), ("shape", ctypes.POINTER(ctypes.c_int64)), - ("strides", ctypes.POINTER(ctypes.c_int64)), ("byte_offset", ctypes.c_uint64)] + _fields_ = [ + ("data", ctypes.c_void_p), + ("device", _DLDevice), + ("ndim", ctypes.c_int), + ("dtype", _DLDataType), + ("shape", ctypes.POINTER(ctypes.c_int64)), + ("strides", ctypes.POINTER(ctypes.c_int64)), + ("byte_offset", ctypes.c_uint64), + ] class _DLManagedTensor(ctypes.Structure): @@ -36,8 +46,11 @@ class _DLManagedTensor(ctypes.Structure): _DELETER = ctypes.CFUNCTYPE(None, ctypes.POINTER(_DLManagedTensor)) -_DLManagedTensor._fields_ = [("dl_tensor", _DLTensor), ("manager_ctx", ctypes.c_void_p), - ("deleter", _DELETER)] +_DLManagedTensor._fields_ = [ + ("dl_tensor", _DLTensor), + ("manager_ctx", ctypes.c_void_p), + ("deleter", _DELETER), +] _pycapsule_new = ctypes.pythonapi.PyCapsule_New _pycapsule_new.restype = ctypes.py_object diff --git a/cuvis/binding.py b/cuvis/binding.py index 9b03bf1..0cc88a0 100644 --- a/cuvis/binding.py +++ b/cuvis/binding.py @@ -29,9 +29,11 @@ derives from :class:`RuntimeError`, so a single ``except RuntimeError`` covers both, while ``except SDKException`` still catches it as an ordinary cuvis error. -Against a binding too old to report any of this (an older ``cuvis_il`` wheel), every -query answers empty: :func:`missing_symbols` is empty, :func:`available` is ``True`` -and :func:`require` never raises. Absence of evidence, not evidence of absence. +Against a binding too old to report any of this (an older ``cuvis_il`` wheel), +:func:`missing_symbols` is empty and :func:`info` reports unknown throughout: absence of +evidence, not evidence of absence. :func:`available` and :func:`require` stay meaningful +there, because a function such an old binding never exposed is unusable whether or not +anything reports it missing. """ from dataclasses import dataclass, field @@ -160,8 +162,25 @@ def missing_symbols() -> frozenset[str]: return frozenset(getattr(cuvis_il, "missing_symbols", ())) +def unavailable(*names: str) -> Tuple[str, ...]: + """Which of the named functions cannot be called, in the order given. + + A function is unusable for either of two reasons, and callers care about neither: + the binding never exposed it, which is what an older ``cuvis_il`` wheel looks like, + or the binding exposes it but the loaded library does not export it. Checking only + the second would report a function the binding does not even have as available. + + :param names: C function names as they appear in ``cuvis.h``. + :return: the subset that is unusable, empty when all of them can be called. + """ + absent = missing_symbols() + return tuple( + name for name in names if name in absent or not hasattr(cuvis_il, name) + ) + + def available(*names: str) -> bool: - """Whether every named function is provided by the installed cuvis library. + """Whether every named function can actually be called. .. code-block:: python3 @@ -169,16 +188,14 @@ def available(*names: str) -> bool: cube = mesu.get_cube_cuda() :param names: C function names as they appear in ``cuvis.h``. - :return: ``True`` if none of them is reported missing. With a binding too old to - report anything this is always ``True``, so treat it as "nothing known to be - missing" rather than a guarantee. + :return: ``True`` if the binding exposes every one of them and none is reported + missing from the loaded library. """ - absent = missing_symbols() - return not any(name in absent for name in names) + return not unavailable(*names) def require(*names: str) -> None: - """Raise unless every named function is provided by the installed cuvis library. + """Raise unless every named function can actually be called. Use it at the start of an operation to fail with a clear explanation, instead of letting a call fail deeper in with less context. @@ -188,13 +205,12 @@ def require(*names: str) -> None: binding.require("cuvis_cuda_mem_get_view", "cuvis_cuda_mem_free") :param names: C function names as they appear in ``cuvis.h``. - :raises UnavailableSDKFunction: naming whichever of them are missing; the message + :raises UnavailableSDKFunction: naming whichever of them are unusable; the message also states the loaded SDK version and the one the binding expects. """ - absent = missing_symbols() - unavailable = tuple(name for name in names if name in absent) - if unavailable: - raise UnavailableSDKFunction(*unavailable) + missing = unavailable(*names) + if missing: + raise UnavailableSDKFunction(*missing) __all__ = [ @@ -202,6 +218,7 @@ def require(*names: str) -> None: "UnavailableSDKFunction", "info", "missing_symbols", + "unavailable", "available", "require", ] diff --git a/cuvis/cube_utils.py b/cuvis/cube_utils.py index a9a0df4..1a44bd6 100644 --- a/cuvis/cube_utils.py +++ b/cuvis/cube_utils.py @@ -5,6 +5,8 @@ import operator from .cuvis_aux import SDKException from .cuvis_types import DataFormat +from . import cuda +import cuvis_ipc _IMBUF_READERS = { 1: cuvis_il.cuvis_read_imbuf_uint8, @@ -419,6 +421,8 @@ def apply(self, other): ImageData.__abs__ = lambda self: self._wrap(abs(self.array)) del _op, _reflected, _method + + class CudaImageData(object): """Device-resident image data backed by a shareable CUDA buffer. @@ -443,9 +447,10 @@ class CudaImageData(object): def __init__(self, cuda_buf): if not isinstance(cuda_buf, cuvis_il.cuvis_cuda_imbuffer_t): raise TypeError( - "Wrong data type for cuda image buffer: {}".format(type(cuda_buf))) - self._handle = cuda_buf.handle # CUVIS_CUDA_MEM (int), owned - self._ipc_handle = None # CUVIS_CUDA_IPC (int), set by make_ipc() + "Wrong data type for cuda image buffer: {}".format(type(cuda_buf)) + ) + self._handle = cuda_buf.handle # CUVIS_CUDA_MEM (int), owned + self._ipc_handle = None # CUVIS_CUDA_IPC (int), set by make_ipc() self._format = cuda_buf.format self.width = cuda_buf.width self.height = cuda_buf.height @@ -455,7 +460,8 @@ def __init__(self, cuda_buf): if cuda_buf.wavelength is not None: self.wavelength = [ cuvis_il.p_unsigned_int_getitem(cuda_buf.wavelength, z) - for z in range(self.channels)] + for z in range(self.channels) + ] # bytes of the transportable IPC descriptor, filled by make_ipc() self.descriptor = None @@ -472,6 +478,7 @@ def __cuda_array_interface__(self): # a pointer with NO lifecycle tie: the caller must keep this CudaImageData alive for # as long as the resulting array is used, or it reads freed device memory. import warnings + warnings.warn( "CudaImageData.__cuda_array_interface__ has no lifecycle management; the buffer " "is freed when this CudaImageData is dropped. Prefer to_torch() (DLPack), which " @@ -499,16 +506,20 @@ def to_torch(self): import torch except ImportError as e: raise ImportError( - "torch is required for CudaImageData.to_torch(); install 'cuvis[torch]'") from e + "torch is required for CudaImageData.to_torch(); install 'cuvis[torch]'" + ) from e from ._dlpack import make_cuda_dlpack ptr, size, dev = self._view() pref = cuvis_il.new_p_int() - if cuvis_il.status_ok != cuvis_il.cuvis_cuda_mem_copy_handle(self._handle, pref): + if cuvis_il.status_ok != cuvis_il.cuvis_cuda_mem_copy_handle( + self._handle, pref + ): raise SDKException() ref = cuvis_il.p_int_value(pref) producer = make_cuda_dlpack( - ptr, size, dev, on_delete=lambda: cuvis_il.cuvis_cuda_mem_free(ref)) + ptr, size, dev, on_delete=lambda: cuvis_il.cuvis_cuda_mem_free(ref) + ) t = torch.from_dlpack(producer) # flat uint8 t = t.view(getattr(torch, self._TORCH_DTYPE[self._format])) return t.reshape(self.height, self.width, self.channels) @@ -523,12 +534,17 @@ def make_ipc(self, backend: int = 0): The IPC handle is an independent reference kept until __del__; while it is open, freeing the mem handle does not release the device memory. Returns .descriptor. """ + cuda.require_ipc() pipc = cuvis_il.new_p_int() - if cuvis_il.status_ok != cuvis_il.cuvis_cuda_ipc_handle_create(self._handle, int(backend), pipc): + if cuvis_il.status_ok != cuvis_il.cuvis_cuda_ipc_handle_create( + self._handle, int(backend), pipc + ): raise SDKException() self._ipc_handle = cuvis_il.p_int_value(pipc) desc = cuvis_il.cuvis_cuda_ipc_descriptor_t() - if cuvis_il.status_ok != cuvis_il.cuvis_cuda_ipc_get_descriptor(self._ipc_handle, desc): + if cuvis_il.status_ok != cuvis_il.cuvis_cuda_ipc_get_descriptor( + self._ipc_handle, desc + ): raise SDKException() self.descriptor = cuvis_il.cuvis_cuda_descriptor_bytes(desc) return self.descriptor @@ -536,15 +552,16 @@ def make_ipc(self, backend: int = 0): def export_payload(self, backend: int = 0) -> bytes: """A single transportable blob (IPC descriptor + geometry) for a consumer process. - Send these bytes out-of-band; the consumer opens them with cuvis.ipc.open(payload) + Send these bytes out-of-band; the consumer opens them with cuvis_ipc.open(payload) and gets a correctly shaped/typed tensor. Calls make_ipc(backend) if not already done (backend: 0=auto, 1=pool, 2=legacy, 3=VMM). Keep this CudaImageData alive until the consumer is finished (legacy IPC has no cross-process refcount). """ if self.descriptor is None: self.make_ipc(backend) - from . import ipc - return ipc.pack_payload(self.descriptor, self.width, self.height, self.channels, self._format) + return cuvis_ipc.pack_payload( + self.descriptor, self.width, self.height, self.channels, self._format + ) def __del__(self): try: diff --git a/cuvis/cuda.py b/cuvis/cuda.py index a6703f3..60c8ce3 100644 --- a/cuvis/cuda.py +++ b/cuvis/cuda.py @@ -4,7 +4,7 @@ from cuvis import cuda - caps = cuda.capabilities() # inspect what this binary + environment support + caps = cuda.capabilities() # what this SDK, this device and this env support if caps.same_process: cuda.enable() # BEFORE loading/processing measurements ... @@ -13,35 +13,56 @@ `enable()` also disables the host auto-refresh (`Measurement._refresh_images = False`) so a GPU-processed cube stays on the device instead of being copied to host and freed. -Note: the shipped cuvis binding links the CUDA runtime, so `import cuvis` already requires a -CUDA runtime to be present. This module gates the CUDA *feature surface* and the optional -`torch` / `cuda-python` consumer dependencies, not the binding's own runtime dependency. +Three unrelated things can each deny CUDA, and they are answered separately rather than +collapsed into one boolean: the installed cuvis library may not provide the functions, which +`cuvis.binding` reports without calling anything; the device or driver may not support a +backend, which only the SDK can answer; and the optional consumer packages may be absent. +Keeping them apart is what lets `enable()` say which of the three went wrong. """ + import importlib.util from typing import NamedTuple -from cuvis_il import cuvis_il - -from .cube_utils import CudaImageData # re-exported; the device-image type - -# Backend codes (match cuvis.h CUVIS_CUDA_IPC_BACKEND_*). -_BACKEND_NONE = 0 -_BACKEND_POOL = 1 -_BACKEND_LEGACY = 2 -_BACKEND_VMM = 3 +from . import binding +from ._cuvis_il import cuvis_il + +# Backend codes, matching CUVIS_CUDA_IPC_BACKEND_* in cuvis.h. BACKEND_NONE doubles as the +# same-process probe and as "auto" where an export picks a backend itself. +BACKEND_NONE = 0 +BACKEND_POOL = 1 +BACKEND_LEGACY = 2 +BACKEND_VMM = 3 + +# Functions the same-process device path calls, as named in cuvis.h. BACKEND_PROBE answers +# for a backend and so gates every capability query, including same_process. +BACKEND_PROBE = "cuvis_cuda_ipc_backend_available" +DEVICE_FUNCTIONS = ( + "cuvis_measurement_get_data_image_cuda", + "cuvis_cuda_mem_get_view", + "cuvis_cuda_mem_copy_handle", + "cuvis_cuda_mem_free", +) + +# Needed on top of those to export a device buffer to another process. +IPC_FUNCTIONS = ( + "cuvis_cuda_ipc_handle_create", + "cuvis_cuda_ipc_get_descriptor", + "cuvis_cuda_ipc_handle_free", +) _enabled = False class CudaCapabilities(NamedTuple): - """What the loaded cuvis binary and the current Python environment support. + """What the installed cuvis library, the current device and this environment support. - same_process: the CUDA boundary responds (same-process device sharing works). + same_process: the CUDA boundary responds, so same-process device sharing works. ipc_pool: exportable memory pool backend (zero-copy cross-process, needs an exportable pool). ipc_legacy: legacy cudaIpc backend (copy-on-export; available on most WDDM/Linux GPUs). ipc_vmm: driver-API VMM backend (copy-on-export, with an exportable handle type). torch / cuda_python: optional consumer packages importable (checked, not imported). """ + same_process: bool ipc_pool: bool ipc_legacy: bool @@ -54,40 +75,85 @@ def any_ipc(self) -> bool: return self.ipc_pool or self.ipc_legacy or self.ipc_vmm -def _backend_available(code: int) -> bool: +def _installed(package: str) -> bool: + """Whether an optional consumer package is installed, without importing it. + + find_spec imports parent packages to reach a submodule, and cuvis puts its own + directory on sys.path, so probing `cuda.bindings` can resolve `cuda` to this very + module whenever cuda-python is absent. A probe that cannot resolve is an answer: + the package is not installed. + """ try: - p = cuvis_il.new_p_int() - if cuvis_il.cuvis_cuda_ipc_backend_available(code, p) != cuvis_il.status_ok: - return False - return bool(cuvis_il.p_int_value(p)) - except Exception: - # Symbol absent or a CUDA-less binary: treat as unavailable. + return importlib.util.find_spec(package) is not None + except (ImportError, ValueError): return False +def _backend_supported(code: int) -> bool: + """Ask the SDK whether this device and driver support one backend. + + Reached only once the function is known to exist, so a false answer here is the SDK's + verdict on the hardware rather than a missing symbol wearing the same disguise. + """ + out = cuvis_il.new_p_int() + if cuvis_il.status_ok != cuvis_il.cuvis_cuda_ipc_backend_available(code, out): + return False + return bool(cuvis_il.p_int_value(out)) + + def capabilities() -> CudaCapabilities: - """Probe CUDA capabilities at runtime. Safe to call before enable().""" + """Probe what is supported, here and now. Safe to call before init() or enable().""" + device = binding.available(BACKEND_PROBE, *DEVICE_FUNCTIONS) + ipc = device and binding.available(*IPC_FUNCTIONS) return CudaCapabilities( - same_process=_backend_available(_BACKEND_NONE), - ipc_pool=_backend_available(_BACKEND_POOL), - ipc_legacy=_backend_available(_BACKEND_LEGACY), - ipc_vmm=_backend_available(_BACKEND_VMM), - torch=importlib.util.find_spec("torch") is not None, - cuda_python=importlib.util.find_spec("cuda.bindings") is not None, + same_process=device and _backend_supported(BACKEND_NONE), + ipc_pool=ipc and _backend_supported(BACKEND_POOL), + ipc_legacy=ipc and _backend_supported(BACKEND_LEGACY), + ipc_vmm=ipc and _backend_supported(BACKEND_VMM), + torch=_installed("torch"), + cuda_python=_installed("cuda.bindings"), ) +def require_device() -> None: + """Raise unless the installed cuvis library provides the same-process device path. + + Guards the entry points so a library without CUDA fails by naming the functions it + lacks, rather than as an AttributeError from deep inside the binding. + + :raises cuvis.UnavailableSDKFunction: naming the functions that are unavailable. + """ + binding.require(*DEVICE_FUNCTIONS) + + +def require_ipc() -> None: + """Raise unless the installed cuvis library provides the cross-process export path. + + :raises cuvis.UnavailableSDKFunction: naming the functions that are unavailable. + """ + binding.require(*DEVICE_FUNCTIONS, *IPC_FUNCTIONS) + + def enable() -> None: """Turn on CUDA mode. Call this BEFORE loading or processing measurements. Routes `Measurement.get_cube()` through the device path and disables the host - auto-refresh so the GPU cube is kept on the device. Raises RuntimeError if this - cuvis build has no CUDA support. + auto-refresh so the GPU cube is kept on the device. + + :raises cuvis.UnavailableSDKFunction: the installed cuvis library does not provide the + CUDA functions; the message names them and both library versions. + :raises RuntimeError: the library provides them, but this device or driver reports no + CUDA support. """ global _enabled - if not capabilities().same_process: - raise RuntimeError("this cuvis build has no CUDA support") + binding.require(BACKEND_PROBE, *DEVICE_FUNCTIONS) + if not _backend_supported(BACKEND_NONE): + raise RuntimeError( + "the installed CUVIS SDK provides the CUDA functions, but this device " + "reports no CUDA support\n{}".format(binding.info()) + ) from .Measurement import Measurement + Measurement._refresh_images = False _enabled = True @@ -96,6 +162,7 @@ def disable() -> None: """Turn off CUDA mode and restore the host auto-refresh.""" global _enabled from .Measurement import Measurement + Measurement._refresh_images = True _enabled = False @@ -104,4 +171,16 @@ def is_enabled() -> bool: return _enabled -__all__ = ["CudaCapabilities", "capabilities", "enable", "disable", "is_enabled", "CudaImageData"] +__all__ = [ + "CudaCapabilities", + "capabilities", + "require_device", + "require_ipc", + "enable", + "disable", + "is_enabled", + "BACKEND_NONE", + "BACKEND_POOL", + "BACKEND_LEGACY", + "BACKEND_VMM", +] diff --git a/cuvis/cuda_import.py b/cuvis/cuda_import.py deleted file mode 100644 index af91a7d..0000000 --- a/cuvis/cuda_import.py +++ /dev/null @@ -1,17 +0,0 @@ -"""Deprecated: use cuvis.ipc instead. - -Kept as a thin shim for back-compat. `open_ipc(descriptor)` maps a raw descriptor (you supply -geometry via .tensor(dtype, shape)); the preferred path is a bundled payload via -cuvis.ipc.open(payload). See cuvis/ipc.py. -""" -from .ipc import ImportedCube as ImportedIpcTensor # noqa: F401 (back-compat name) -from .ipc import open_descriptor, BACKEND_NONE, BACKEND_POOL, BACKEND_LEGACY, BACKEND_VMM # noqa: F401 - - -def open_ipc(descriptor_bytes) -> ImportedIpcTensor: - """Deprecated alias of cuvis.ipc.open_descriptor().""" - return open_descriptor(descriptor_bytes) - - -__all__ = ["ImportedIpcTensor", "open_ipc", "BACKEND_NONE", "BACKEND_POOL", - "BACKEND_LEGACY", "BACKEND_VMM"] diff --git a/cuvis/ipc.py b/cuvis_ipc.py similarity index 73% rename from cuvis/ipc.py rename to cuvis_ipc.py index 672fc6a..0bb7ec0 100644 --- a/cuvis/ipc.py +++ b/cuvis_ipc.py @@ -2,30 +2,37 @@ Use this in a SEPARATE process that receives a cuvis IPC payload - it does NOT initialize or link the cuvis SDK. It needs only cuda-python (`cuda.bindings`) and torch, imported lazily. -`import cuvis.ipc` works with no CUVIS env var and no cuvis.dll present. + +It sits outside the `cuvis` package on purpose. Importing `cuvis` requires the SDK to be +installed and the CUVIS environment variable to be set, which a consumer process by +definition does not have; `import cuvis_ipc` needs neither, and the only import here is +`struct`. Producer (in the cuvis process): cimg = mesu.get_cube_cuda_ipc() # keep this alive until the consumer is done payload = cimg.export_payload() # a single transportable bytes blob (descriptor + geometry) Consumer (this module, any process): - import cuvis.ipc as ipc - with ipc.open(payload) as cube: + import cuvis_ipc + with cuvis_ipc.open(payload) as cube: t = cube.to_torch() # correctly shaped/typed zero-copy CUDA tensor ... # use t inside the block # leaving the block releases this process's mapping (does not free the exporter's memory) The exporting process must outlive this importer: legacy IPC has no cross-process refcount. """ + import struct # --- IPC descriptor wire format (locked; mirrors cuvis_cuda_ipc_descriptor_t) --- # 48-byte header + 64-byte blob (pool OS handle) at offset 48, then ptr_blob_len(+pad) and a # 64-byte ptr_blob (cudaMemPoolPtrExportData) at offset 120. Total 184. -_HEAD = struct.Struct(" torch dtype name / __cuda_array_interface__ typestr (1/2/3/4 = u8/u16/u32/f32) _TORCH_DTYPE = {1: "uint8", 2: "uint16", 3: "uint32", 4: "float32"} _TYPESTR = {1: "|u1", 2: " bytes: +def pack_payload( + descriptor: bytes, width: int, height: int, channels: int, format_code: int +) -> bytes: """Bundle an IPC descriptor and cube geometry into one transportable blob.""" if len(descriptor) != _DESC_LEN: raise ValueError(f"descriptor must be {_DESC_LEN} bytes, got {len(descriptor)}") - return _PAYLOAD_HDR.pack(_MAGIC, _VERSION, int(width), int(height), int(channels), - int(format_code)) + bytes(descriptor) + return _PAYLOAD_HDR.pack( + _MAGIC, _VERSION, int(width), int(height), int(channels), int(format_code) + ) + bytes(descriptor) def _unpack_payload(payload: bytes): @@ -57,7 +69,7 @@ def _unpack_payload(payload: bytes): raise ValueError("not a cuvis IPC payload (bad magic)") if version != _VERSION: raise ValueError(f"unsupported cuvis IPC payload version {version}") - descriptor = bytes(payload[_PAYLOAD_HDR.size:]) + descriptor = bytes(payload[_PAYLOAD_HDR.size :]) return (width, height, channels, fmt), descriptor @@ -71,7 +83,10 @@ def _ck(ret, what): class _CudaArray: def __init__(self, ptr, nbytes): self.__cuda_array_interface__ = { - "shape": (nbytes,), "typestr": "|u1", "data": (int(ptr), False), "version": 3, + "shape": (nbytes,), + "typestr": "|u1", + "data": (int(ptr), False), + "version": 3, } @@ -84,20 +99,31 @@ class ImportedCube: """ def __init__(self, descriptor_bytes: bytes, shape=None, format_code=None): - (self.backend, self.device, self.htype, blob_len, - self.size, self.alloc, self.offset, self.pid) = _HEAD.unpack_from(descriptor_bytes, 0) + ( + self.backend, + self.device, + self.htype, + blob_len, + self.size, + self.alloc, + self.offset, + self.pid, + ) = _HEAD.unpack_from(descriptor_bytes, 0) if blob_len > _BLOB_MAX: raise ValueError(f"blob_len {blob_len} exceeds {_BLOB_MAX}") - self._blob = bytes(descriptor_bytes[_BLOB_OFF:_BLOB_OFF + blob_len]) + self._blob = bytes(descriptor_bytes[_BLOB_OFF : _BLOB_OFF + blob_len]) (ptr_blob_len,) = struct.unpack_from(" _PTR_BLOB_MAX: raise ValueError(f"ptr_blob_len {ptr_blob_len} exceeds {_PTR_BLOB_MAX}") - self._ptr_blob = bytes(descriptor_bytes[_PTR_BLOB_OFF:_PTR_BLOB_OFF + ptr_blob_len]) + self._ptr_blob = bytes( + descriptor_bytes[_PTR_BLOB_OFF : _PTR_BLOB_OFF + ptr_blob_len] + ) self._shape = tuple(shape) if shape is not None else None self._format = format_code self._close = None from cuda.bindings import runtime + runtime.cudaSetDevice(self.device) if self.backend == BACKEND_POOL: @@ -107,8 +133,12 @@ def __init__(self, descriptor_bytes: bytes, shape=None, format_code=None): elif self.backend == BACKEND_VMM: self._ptr, self._close = self._open_vmm() else: - raise NotImplementedError(f"backend {self.backend} is not importable cross-process") - self._ptr += self.offset # 0 for pool/legacy/vmm (import returns the exact base pointer) + raise NotImplementedError( + f"backend {self.backend} is not importable cross-process" + ) + self._ptr += ( + self.offset + ) # 0 for pool/legacy/vmm (import returns the exact base pointer) @property def shape(self): @@ -122,17 +152,23 @@ def _open_pool(self): # IPC-capable memory pool: import the pool from its OS shareable handle, grant this # device access, then import the exact pointer from the per-allocation export data. from cuda.bindings import runtime + if self.htype == H_WIN32_KMT: htype = runtime.cudaMemAllocationHandleType.cudaMemHandleTypeWin32Kmt elif self.htype == H_POSIX_FD: - htype = runtime.cudaMemAllocationHandleType.cudaMemHandleTypePosixFileDescriptor + htype = ( + runtime.cudaMemAllocationHandleType.cudaMemHandleTypePosixFileDescriptor + ) else: raise NotImplementedError( - f"pool handle_type {self.htype} needs out-of-band duplication (DuplicateHandle / SCM_RIGHTS)") + f"pool handle_type {self.htype} needs out-of-band duplication (DuplicateHandle / SCM_RIGHTS)" + ) handle_val = int.from_bytes(self._blob, "little") - (pool,) = _ck(runtime.cudaMemPoolImportFromShareableHandle(handle_val, htype, 0), - "cudaMemPoolImportFromShareableHandle") + (pool,) = _ck( + runtime.cudaMemPoolImportFromShareableHandle(handle_val, htype, 0), + "cudaMemPoolImportFromShareableHandle", + ) acc = runtime.cudaMemAccessDesc() acc.location.type = runtime.cudaMemLocationType.cudaMemLocationTypeDevice @@ -142,10 +178,13 @@ def _open_pool(self): export_data = runtime.cudaMemPoolPtrExportData() export_data.reserved = self._ptr_blob.ljust(_PTR_BLOB_MAX, b"\x00") - (ptr,) = _ck(runtime.cudaMemPoolImportPointer(pool, export_data), "cudaMemPoolImportPointer") + (ptr,) = _ck( + runtime.cudaMemPoolImportPointer(pool, export_data), + "cudaMemPoolImportPointer", + ) def close(): - runtime.cudaFree(ptr) # release this process's imported pointer + runtime.cudaFree(ptr) # release this process's imported pointer runtime.cudaMemPoolDestroy(pool) # release the imported pool handle return int(ptr), close @@ -153,19 +192,25 @@ def close(): def _open_legacy(self): # Legacy cudaIpc: the descriptor blob is a self-contained 64-byte cudaIpcMemHandle_t. from cuda.bindings import runtime + h = runtime.cudaIpcMemHandle_t() h.reserved = self._blob.ljust(_BLOB_MAX, b"\x00") - (ptr,) = _ck(runtime.cudaIpcOpenMemHandle(h, runtime.cudaIpcMemLazyEnablePeerAccess), - "cudaIpcOpenMemHandle") + (ptr,) = _ck( + runtime.cudaIpcOpenMemHandle(h, runtime.cudaIpcMemLazyEnablePeerAccess), + "cudaIpcOpenMemHandle", + ) def close(): - runtime.cudaIpcCloseMemHandle(ptr) # release this process's mapping (not the exporter's) + runtime.cudaIpcCloseMemHandle( + ptr + ) # release this process's mapping (not the exporter's) return int(ptr), close def _open_vmm(self): # VMM: import the generic handle from its OS shareable handle, reserve + map + grant access. from cuda.bindings import driver + driver.cuInit(0) if self.htype == H_WIN32_KMT: htype = driver.CUmemAllocationHandleType.CU_MEM_HANDLE_TYPE_WIN32_KMT @@ -177,8 +222,10 @@ def _open_vmm(self): raise NotImplementedError(f"vmm handle_type {self.htype} not supported") handle_val = int.from_bytes(self._blob, "little") - (gen,) = _ck(driver.cuMemImportFromShareableHandle(handle_val, htype), - "cuMemImportFromShareableHandle") + (gen,) = _ck( + driver.cuMemImportFromShareableHandle(handle_val, htype), + "cuMemImportFromShareableHandle", + ) size = self.alloc (ptr,) = _ck(driver.cuMemAddressReserve(size, 0, 0, 0), "cuMemAddressReserve") _ck(driver.cuMemMap(ptr, size, 0, gen, 0), "cuMemMap") @@ -197,7 +244,9 @@ def close(): return int(ptr), close def _torch_dtype(self, torch): - return None if self._format is None else getattr(torch, _TORCH_DTYPE[self._format]) + return ( + None if self._format is None else getattr(torch, _TORCH_DTYPE[self._format]) + ) def to_torch(self, dtype=None, shape=None): """Zero-copy CUDA torch tensor over the mapped buffer, valid inside the with-block. @@ -205,6 +254,7 @@ def to_torch(self, dtype=None, shape=None): With open(payload), dtype and shape are taken from the payload geometry; overrides may be passed explicitly (needed after open_descriptor without geometry).""" import torch + n = self.size - self.offset t = torch.as_tensor(_CudaArray(self._ptr, n), device=f"cuda:{self.device}") dt = dtype if dtype is not None else self._torch_dtype(torch) @@ -221,7 +271,8 @@ def to_torch(self, dtype=None, shape=None): def __cuda_array_interface__(self): if self._shape is None or self._format is None: raise RuntimeError( - "__cuda_array_interface__ needs geometry; open via open(payload) or use to_torch(dtype, shape)") + "__cuda_array_interface__ needs geometry; open via open(payload) or use to_torch(dtype, shape)" + ) return { "shape": self._shape, "typestr": _TYPESTR[self._format], @@ -254,5 +305,13 @@ def open_descriptor(descriptor_bytes: bytes, shape=None) -> ImportedCube: return ImportedCube(descriptor_bytes, shape=shape) -__all__ = ["ImportedCube", "open", "open_descriptor", "pack_payload", - "BACKEND_NONE", "BACKEND_POOL", "BACKEND_LEGACY", "BACKEND_VMM"] +__all__ = [ + "ImportedCube", + "open", + "open_descriptor", + "pack_payload", + "BACKEND_NONE", + "BACKEND_POOL", + "BACKEND_LEGACY", + "BACKEND_VMM", +] diff --git a/pyproject.toml b/pyproject.toml index 2599ca3..841cd4d 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -38,6 +38,9 @@ Issues = "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/cubert-hyperspectral/cuvis.python/issues" [tool.setuptools] packages = ["cuvis"] +# cuvis_ipc sits outside the package because a consumer process opening a CUDA IPC +# payload has no SDK, and importing `cuvis` requires one. +py-modules = ["cuvis_ipc"] include-package-data = true [tool.setuptools.package-data] diff --git a/tests/test_binding.py b/tests/test_binding.py new file mode 100644 index 0000000..9eaeaa6 --- /dev/null +++ b/tests/test_binding.py @@ -0,0 +1,65 @@ +""" +Tests for cuvis.binding. + +Covers availability reporting. The installed library differs between CI and developer +machines, so the tests pin the two ends that are always true: a function the binding +does expose is available, and one nothing exposes is not. +""" + +import pytest + +import cuvis +from cuvis import binding + +PRESENT = "cuvis_measurement_load" +ABSENT = "cuvis_a_function_that_does_not_exist" + + +def test_info_reports_without_an_initialised_sdk(): + """Test info() is answerable before cuvis.init, which is what gates a feature check.""" + current = binding.info() + assert isinstance(current, cuvis.BindingInfo) + assert isinstance(current.is_complete, bool) + assert "cuvis binding" in str(current) + + +def test_a_function_the_binding_exposes_is_available(): + """Test a function present in the binding and not reported missing is available.""" + assert binding.available(PRESENT) + assert binding.unavailable(PRESENT) == () + binding.require(PRESENT) + + +def test_a_function_no_binding_exposes_is_unavailable(): + """Test availability accounts for absent symbols, not only reported-missing ones. + + A binding too old to report missing symbols reports none, so consulting that list + alone would call a function the binding never had and fail with an AttributeError. + """ + assert not binding.available(ABSENT) + assert binding.unavailable(ABSENT) == (ABSENT,) + + +def test_unavailable_preserves_the_requested_order(): + """Test the report names only the unusable functions, in the order asked for.""" + assert binding.unavailable(PRESENT, ABSENT, PRESENT) == (ABSENT,) + + +def test_require_names_every_unavailable_function(): + """Test the raised error carries the names, so the message says what to install.""" + with pytest.raises(cuvis.UnavailableSDKFunction) as excinfo: + binding.require(PRESENT, ABSENT) + assert excinfo.value.names == (ABSENT,) + assert ABSENT in str(excinfo.value) + + +@pytest.mark.parametrize("expected", [RuntimeError, cuvis.cuvis_aux.SDKException]) +def test_unavailable_function_is_catchable_by_either_base(expected): + """Test the exception satisfies both except clauses its docstring advertises.""" + with pytest.raises(expected): + binding.require(ABSENT) + + +def test_missing_symbols_is_a_frozenset(): + """Test the reported set is immutable, so a caller cannot corrupt it.""" + assert isinstance(binding.missing_symbols(), frozenset) diff --git a/tests/test_cuda.py b/tests/test_cuda.py new file mode 100644 index 0000000..7ef719f --- /dev/null +++ b/tests/test_cuda.py @@ -0,0 +1,74 @@ +""" +Tests for cuvis.cuda capability reporting. + +The assertions hold whether or not the installed cuvis library provides the CUDA +functions, since CI and developer machines differ on that. What is pinned is that +capabilities() answers rather than raises, and that an unavailable function is +reported by name instead of surfacing as an AttributeError from the binding. +""" + +import pytest + +import cuvis +from cuvis import binding, cuda + +_HAS_DEVICE = binding.available(cuda.BACKEND_PROBE, *cuda.DEVICE_FUNCTIONS) + + +def test_capabilities_answers_without_raising(): + """Test capabilities() reports booleans instead of raising on a CUDA-less SDK.""" + caps = cuda.capabilities() + assert isinstance(caps, cuda.CudaCapabilities) + assert all(isinstance(value, bool) for value in caps) + assert isinstance(caps.any_ipc, bool) + + +def test_capabilities_is_safe_before_init(): + """Test capabilities() needs no initialised SDK, so a caller can gate on it first.""" + assert cuda.capabilities() == cuda.capabilities() + + +def test_cuda_mode_is_off_by_default(): + """Test CUDA stays opt-in, leaving the host refresh path in place.""" + assert cuda.is_enabled() is False + assert cuvis.Measurement._refresh_images is True + + +def test_disable_restores_the_host_refresh(): + """Test disable() is safe to call unconditionally and restores the host path.""" + cuda.disable() + assert cuda.is_enabled() is False + assert cuvis.Measurement._refresh_images is True + + +@pytest.mark.skipif(_HAS_DEVICE, reason="this SDK provides the CUDA functions") +def test_capabilities_are_false_without_the_sdk_functions(): + """Test a library lacking the CUDA functions reports no CUDA support.""" + caps = cuda.capabilities() + assert caps.same_process is False + assert caps.any_ipc is False + + +@pytest.mark.skipif(_HAS_DEVICE, reason="this SDK provides the CUDA functions") +@pytest.mark.parametrize("guard", [cuda.require_device, cuda.require_ipc, cuda.enable]) +def test_missing_functions_are_reported_by_name(guard): + """Test the CUDA entry points name what the SDK lacks (see cuvis.binding).""" + with pytest.raises(cuvis.UnavailableSDKFunction) as excinfo: + guard() + assert excinfo.value.names + assert all(name in str(excinfo.value) for name in excinfo.value.names) + assert cuda.is_enabled() is False + + +@pytest.mark.skipif(_HAS_DEVICE, reason="this SDK provides the CUDA functions") +def test_get_cube_cuda_reports_the_missing_functions(test_measurement): + """Test the device path fails with the diagnostic, not an AttributeError.""" + with pytest.raises(cuvis.UnavailableSDKFunction): + test_measurement.get_cube_cuda() + + +def test_unavailable_function_is_catchable_either_way(): + """Test UnavailableSDKFunction satisfies both except clauses it advertises.""" + for expected in (RuntimeError, cuvis.cuvis_aux.SDKException): + with pytest.raises(expected): + binding.require("cuvis_a_function_that_does_not_exist") diff --git a/tests/test_cuvis_ipc.py b/tests/test_cuvis_ipc.py new file mode 100644 index 0000000..0168f9b --- /dev/null +++ b/tests/test_cuvis_ipc.py @@ -0,0 +1,63 @@ +""" +Tests for the cuvis_ipc cross-process consumer module. + +Covers the payload codec and the property the module exists for: it is importable in +a process that has no cuvis SDK. Mapping a buffer needs a CUDA device and a live +exporting process, so that is not covered here. +""" + +import os +import subprocess +import sys + +import pytest + +import cuvis_ipc + +DESCRIPTOR = bytes(range(184)) +GEOMETRY = (290, 275, 51, 2) # width, height, channels, format code + + +def test_payload_round_trip(): + """Test a packed payload decodes back to the same descriptor and geometry.""" + payload = cuvis_ipc.pack_payload(DESCRIPTOR, *GEOMETRY) + geometry, descriptor = cuvis_ipc._unpack_payload(payload) + assert geometry == GEOMETRY + assert descriptor == DESCRIPTOR + + +def test_pack_payload_rejects_a_wrong_sized_descriptor(): + """Test the descriptor length is checked, since the wire format is fixed.""" + with pytest.raises(ValueError, match="184 bytes"): + cuvis_ipc.pack_payload(DESCRIPTOR[:-1], *GEOMETRY) + + +def test_unpack_rejects_foreign_bytes(): + """Test the magic guards against anything that is not a cuvis payload.""" + with pytest.raises(ValueError, match="magic"): + cuvis_ipc._unpack_payload(b"XXXX" + bytes(200)) + + +def test_unpack_rejects_a_future_version(): + """Test a payload from a newer wire format is refused rather than misread.""" + payload = bytearray(cuvis_ipc.pack_payload(DESCRIPTOR, *GEOMETRY)) + payload[4:8] = (cuvis_ipc._VERSION + 1).to_bytes(4, "little") + with pytest.raises(ValueError, match="version"): + cuvis_ipc._unpack_payload(bytes(payload)) + + +def test_importable_without_the_sdk(): + """Test the consumer module imports with no CUVIS environment variable set. + + This is the whole reason it sits outside the cuvis package: importing cuvis + requires the SDK, and a consumer process opening a payload does not have one. + """ + env = {k: v for k, v in os.environ.items() if k != "CUVIS"} + result = subprocess.run( + [sys.executable, "-c", "import cuvis_ipc, sys; print('cuvis' in sys.modules)"], + capture_output=True, + text=True, + env=env, + ) + assert result.returncode == 0, result.stderr + assert result.stdout.strip() == "False" From 8af2feff824043fff49766d13f1db345b63f64d1 Mon Sep 17 00:00:00 2001 From: Simon Birkholz Date: Wed, 19 Aug 2026 17:50:15 +0200 Subject: [PATCH 3/4] wip --- cuvis/Measurement.py | 75 +++++++++++++++++++++++++++++++++++--------- cuvis/cube_utils.py | 6 +++- 2 files changed, 66 insertions(+), 15 deletions(-) diff --git a/cuvis/Measurement.py b/cuvis/Measurement.py index b0c08a3..56cdc1f 100644 --- a/cuvis/Measurement.py +++ b/cuvis/Measurement.py @@ -289,9 +289,28 @@ def cube(self) -> ImageData: def get_cube_cuda(self, key: str = "cube") -> CudaImageData: """Image data as a device-resident CUDA buffer for same-process, zero-copy use. - Returns a CudaImageData wrapping a CUVIS_CUDA_MEM handle; wrap it with - .to_torch() (DLPack) or __cuda_array_interface__. The underlying image data - must be backed by CUDA device memory (raises SDKException otherwise). + .. code-block:: python3 + + from cuvis import cuda + + if cuda.capabilities().same_process: + cuda.enable() # BEFORE loading or processing + tensor = mesu.get_cube_cuda().to_torch() + + Returns a :class:`cuvis.CudaImageData` wrapping a CUVIS_CUDA_MEM handle; read it + with ``.to_torch()`` (DLPack, which ties the buffer lifetime to the tensor) or + through ``__cuda_array_interface__`` (which does not, so keep the CudaImageData + alive). No host copy is made. + + The cube must still be on the device, which it is only when `cuda.enable` was + called before it was processed; the host fetch in `refresh` otherwise moves it to + host memory and frees the device copy. + + :param key: which image entry to read, `"cube"` unless the measurement carries + several. + :raises cuvis.UnavailableSDKFunction: the installed cuvis library provides no + CUDA support; call `cuda.capabilities` first to avoid this. + :raises cuvis.cuvis_aux.SDKException: the image data is not device-backed. """ cuda.require_device() buf = cuvis_il.cuvis_cuda_imbuffer_t() @@ -304,23 +323,51 @@ def get_cube_cuda(self, key: str = "cube") -> CudaImageData: def get_cube_cuda_ipc(self, key: str = "cube", backend: int = 0) -> CudaImageData: """Image data as a shareable CUDA buffer for cross-process use. - Fetches the device buffer (get_cube_cuda) and then creates an IPC export on it, - filling .descriptor with the transportable bytes; send those out-of-band to - another process and open them with cuvis_ipc.open. Keep the returned object alive - until the importer is done: it is the in-process pin (legacy IPC has no cross-process - refcount). backend selects the mechanism (0=auto, 1=pool, 2=legacy, 3=VMM); - make_ipc raises SDKException if the requested backend is unavailable on this device. + Producer side; the consumer opens the payload with :mod:`cuvis_ipc`, which needs + no SDK of its own. + + .. code-block:: python3 + + cimg = mesu.get_cube_cuda_ipc() # keep alive until the consumer is done + send(cimg.export_payload()) # descriptor + geometry, one blob + + # ... in the consumer process, no cuvis installed ... + import cuvis_ipc + with cuvis_ipc.open(payload) as cube: + tensor = cube.to_torch() + + Fetches the device buffer with `get_cube_cuda` and creates an IPC export on it, + filling `.descriptor` with the transportable bytes. The returned object is the + in-process pin, since legacy IPC carries no cross-process refcount: drop it and + the consumer is reading freed device memory. + + :param key: which image entry to read. + :param backend: which mechanism to export with, one of `cuvis.cuda.BACKEND_NONE` + (auto), `BACKEND_POOL`, `BACKEND_LEGACY` or `BACKEND_VMM`. `cuda.capabilities` + reports which of them this device supports. + :raises cuvis.UnavailableSDKFunction: the installed cuvis library provides no + CUDA IPC support. + :raises cuvis.cuvis_aux.SDKException: the requested backend is unavailable on + this device. """ cimg = self.get_cube_cuda(key) cimg.make_ipc(backend) return cimg - def get_cube(self, key: str = "cube"): - """Cube via the active mode. + def get_cube(self, key: str = "cube") -> Union[ImageData, CudaImageData]: + """Cube through whichever mode is active, so one call site serves both. + + .. code-block:: python3 + + cube = mesu.get_cube() # CudaImageData after cuda.enable(), else ImageData + + With CUDA mode on (`cuvis.cuda.enable`) this is `get_cube_cuda` and raises when + the device path is unavailable. There is deliberately no silent fallback to the + host: a zero-copy path that quietly degrades to two copies is worse than an error, + because the cost is invisible. - When CUDA mode is enabled (cuvis.cuda.enable()), returns a device-resident - CudaImageData and raises SDKException if the device path is unavailable (no - silent host fallback). Otherwise returns the host ImageData. + :param key: which image entry to read. + :return: `CudaImageData` in CUDA mode, otherwise the host `ImageData`. """ if cuda.is_enabled(): return self.get_cube_cuda(key) diff --git a/cuvis/cube_utils.py b/cuvis/cube_utils.py index 1a44bd6..e68ee8b 100644 --- a/cuvis/cube_utils.py +++ b/cuvis/cube_utils.py @@ -6,7 +6,6 @@ from .cuvis_aux import SDKException from .cuvis_types import DataFormat from . import cuda -import cuvis_ipc _IMBUF_READERS = { 1: cuvis_il.cuvis_read_imbuf_uint8, @@ -559,6 +558,11 @@ def export_payload(self, backend: int = 0) -> bytes: """ if self.descriptor is None: self.make_ipc(backend) + # Imported here, not at module scope: cuvis_ipc is the consumer half and lives + # outside the package, so a producer that never exports must not fail to import + # cuvis because it is absent. + import cuvis_ipc + return cuvis_ipc.pack_payload( self.descriptor, self.width, self.height, self.channels, self._format ) From bec088a9660d359706a258bf087c42712585068d Mon Sep 17 00:00:00 2001 From: Simon Birkholz Date: Sat, 22 Aug 2026 18:50:31 +0200 Subject: [PATCH 4/4] cuda temp files --- .gitignore | 1 + CHANGELOG.md | 14 +-- ImageData-api-notes.md | 152 ++++++++++++++++++++++++++ cuvis/Measurement.py | 2 +- cuvis/binding.py | 2 +- cuvis/cube_utils.py | 25 +++-- examples/cuda_ipc_test.py | 150 ++++++++++++++++++++++++++ examples/cuda_misuse_test.py | 188 +++++++++++++++++++++++++++++++++ examples/cuda_optin.py | 52 +++++++++ examples/cuda_smoke.py | 144 +++++++++++++++++++++++++ examples/cuda_tensor_bench.py | 183 ++++++++++++++++++++++++++++++++ examples/ipc_consumer.py | 28 +++++ tests/test_handle_lifecycle.py | 3 + 13 files changed, 924 insertions(+), 20 deletions(-) create mode 100644 ImageData-api-notes.md create mode 100644 examples/cuda_ipc_test.py create mode 100644 examples/cuda_misuse_test.py create mode 100644 examples/cuda_optin.py create mode 100644 examples/cuda_smoke.py create mode 100644 examples/cuda_tensor_bench.py create mode 100644 examples/ipc_consumer.py diff --git a/.gitignore b/.gitignore index 9793443..7a1533b 100644 --- a/.gitignore +++ b/.gitignore @@ -12,5 +12,6 @@ /cuvis/git-hash.txt /tests/__pycache__ +/examples/__pycache__ /.claude diff --git a/CHANGELOG.md b/CHANGELOG.md index 2da645a..39b467a 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -19,12 +19,12 @@ Pre-releases (`b*`, `rc*`) are not listed. - `CI` - `.github/workflows/release.yml` is driven by `v*.*.*.*` tags: it validates the tag against `pyproject.toml` and against this file, builds, publishes to TestPyPI, and publishes to PyPI plus a GitHub Release after manual approval. - `CI` - `scripts/check_changelog.py` validates this file's structure (header format, allowed section names, descending versions) and the tag/version/changelog agreement at release time. - `CONTRIBUTING.md` - documents the branch model, the version scheme, the changelog conventions and the release checklist. -- `cuvis.BindingInfo` - new frozen dataclass with the fields `built_against: str`, `library_version: str`, `library_path: str` and `missing_symbols: Tuple[str, ...]`, the read-only property `is_complete: bool`, and a `__str__` rendering a report fit for a bug report. +- `cuvis.BindingInfo` - new frozen dataclass with the fields `built_against: str`, `library_version: str`, `library_path: str` and `missing_symbols: tuple[str, ...]`, the read-only property `is_complete: bool`, and a `__str__` rendering a report fit for a bug report. - `cuvis.CudaImageData` - new class, device-resident image data backed by a shareable CUDA buffer, exposed zero-copy through DLPack or `__cuda_array_interface__` with no host copy. - `cuvis.CudaImageData.export_payload` - new method, returns `bytes`. - `cuvis.CudaImageData.make_ipc` - new method, returns `bytes`. - `cuvis.CudaImageData.to_torch` - new method, returns `torch.Tensor`. -- `cuvis.Measurement.get_cube` - new method, returns `Union[ImageData, CudaImageData]` depending on whether `cuvis.cuda.enable` was called. +- `cuvis.Measurement.get_cube` - new method, returns `ImageData | CudaImageData` depending on whether `cuvis.cuda.enable` was called. - `cuvis.Measurement.get_cube_cuda` - new method, returns `CudaImageData`. - `cuvis.Measurement.get_cube_cuda_ipc` - new method, returns `CudaImageData`. - `cuvis.SdkSettings` - new class, a `MutableMapping` of setting id to value that writes the SDK's `cuvis.settings` file, so the SDK configuration can be built in Python instead of maintained by hand. @@ -32,14 +32,14 @@ Pre-releases (`b*`, `rc*`) are not listed. - `cuvis.SdkSettings.__enter__`, `cuvis.SdkSettings.__exit__` - new methods; entering the context serializes the settings into a temporary directory and returns its path as `str`, leaving the context removes the directory. - `cuvis.SdkSettings.save` - new method, writes the settings to a file, or into a directory as `cuvis.settings`. - `cuvis.SdkSettings.xml_str` - new read-only property, returns the serialized settings document as `str`. -- `cuvis.UnavailableSDKFunction` - new exception deriving from both `cuvis.cuvis_aux.SDKException` and `RuntimeError`, with the field `names: Tuple[str, ...]`. +- `cuvis.UnavailableSDKFunction` - new exception deriving from both `cuvis.cuvis_aux.SDKException` and `RuntimeError`, with the field `names: tuple[str, ...]`. - `cuvis.binding` - new module reporting the compiled binding, the cuvis library loaded beside it, and the functions that library does not provide. Nothing in it needs the SDK to be initialised, so it can be called before `cuvis.init`. - `cuvis.binding.available` - new function, returns `bool`. - `cuvis.binding.info` - new function, returns `BindingInfo`. -- `cuvis.binding.missing_symbols` - new function, returns `FrozenSet[str]`. +- `cuvis.binding.missing_symbols` - new function, returns `frozenset[str]`. - `cuvis.binding.require` - new function, raises `UnavailableSDKFunction` naming whichever of the given functions the installed cuvis library does not provide. -- `cuvis.binding.unavailable` - new function, returns `Tuple[str, ...]`. +- `cuvis.binding.unavailable` - new function, returns `tuple[str, ...]`. A function the binding never exposed is unusable as well, so availability cannot be answered from the reported-missing list alone. - `cuvis.cuda` - new module gating the optional CUDA feature surface; CUDA stays off until `cuvis.cuda.enable` is called. - `cuvis.cuda.BACKEND_NONE`, `cuvis.cuda.BACKEND_POOL`, `cuvis.cuda.BACKEND_LEGACY`, `cuvis.cuda.BACKEND_VMM` - new constants, the IPC backend codes from `cuvis.h`. @@ -64,14 +64,14 @@ Pre-releases (`b*`, `rc*`) are not listed. - Whole tree reformatted with `ruff format`; no behaviour change. - `README.md` - documents the version scheme, and lists Python 3.14 among the supported interpreters as `pyproject.toml` already did. - `prebuild.py` - writes `cuvis/git-hash.txt` instead of `git-hash.txt` at the repository root, so the file lands inside the package that declares it as package data. -- `cuvis.General.init` - parameter `settings_path` type changed from `str` to `Union[str, Path, SdkSettings]`. +- `cuvis.General.init` - parameter `settings_path` type changed from `str` to `str | Path | SdkSettings`. An `SdkSettings` is written to a temporary directory that exists only for the duration of the call, since the SDK reads the settings once during initialisation. - `cuvis.FileWriteSettings.GeneralExportSettings.__repr__`, `cuvis.FileWriteSettings.ViewerSettings.__repr__` - the docstring that sat below the nested helper, where it was a dead expression rather than a docstring, moved to the top of the method. - `cuvis.Measurement.capture_time`, `cuvis.Measurement.factory_calibration`, `cuvis.GPSData.time`, `cuvis.SensorInfo.readout_time` - type changed from a naive `datetime.datetime` to one carrying `tzinfo=datetime.timezone.utc`. The instant is unchanged, only the `+00:00` label is added; comparing or subtracting against a naive `datetime` now raises `TypeError`, so use `datetime.datetime.now(datetime.timezone.utc)` or `.astimezone()` for local time. - `cuvis.CalibrationInfo.calibration_date` - type changed from `int` to a `datetime.datetime` carrying `tzinfo=datetime.timezone.utc`; the field was annotated as a `datetime` but returned the raw epoch milliseconds unconverted. The SDK derives this value as midnight on the calibration day in the host's standard local time, so unlike the other timestamps the instant it denotes shifts with the reading machine; treat it as a day, not as an exact moment. -- `cuvis.Measurement.factory_calibration` - type changed from `datetime.datetime` to `Optional[datetime.datetime]`, matching the existing fallback to `None` for SDK values the `datetime` range cannot represent. +- `cuvis.Measurement.factory_calibration` - type changed from `datetime.datetime` to `datetime.datetime | None`, matching the existing fallback to `None` for SDK values the `datetime` range cannot represent. Only the day carries meaning: the SDK stores it as midnight in the local time of the machine that wrote the file, so the time component is an artifact of that machine and the day can be off by one when the file is read in another timezone. - `pyproject.toml` - `requires-python` raised from `>=3.9` to `>=3.10`, Python 3.9 having reached end of life in October 2025. - Annotations throughout `cuvis` restated in the forms Python 3.10 provides: `Union[A, B]` and `Optional[A]` became `A | B` and `A | None`, and `Tuple`, `FrozenSet`, `Sequence`, `Callable` and `Awaitable` now come from `builtins` and `collections.abc` rather than `typing`. diff --git a/ImageData-api-notes.md b/ImageData-api-notes.md new file mode 100644 index 0000000..3c4e3f0 --- /dev/null +++ b/ImageData-api-notes.md @@ -0,0 +1,152 @@ +# ImageData: API notes and proposed additions + +Written 2026-08-19, against `hotfix/fix_broken_qmini_imagedata_wrapping` at version 3.5.3.2. + +This is a design note, not a plan. +It records what the class looks like after the QMini hotfix, which additions are worth making next, and which tempting ones should be refused. +Nothing here is scheduled. + +The reason to write it down now is that 3.5.3.2 introduces the whole numeric surface of `ImageData` at once. +Every member it ships becomes permanent. +The additions below are the ones that fit that surface without contradicting it, and a few of them are cheaper to do before the release than after. + +## What shipped in 3.5.3.2 + +Attributes: `array`, `width`, `height`, `channels`, `wavelength`. +Properties: `shape`, `dtype`, `is_spectrum`, `spectrum`. +Protocols: `__getitem__`, `__array__`, `__array_ufunc__`, `__repr__`, the seven arithmetic operators with their reflected forms, `__neg__`, `__abs__`. +Constructors: `__init__` from a `cuvis_imbuffer_t`, `from_array` from a NumPy array. + +Two invariants hold the design together. +`array` is always three dimensional, `(height, width, channels)`, so a point spectrometer arrives as `(1, 1, channels)` rather than as a bare vector. +`wavelength` is either `None` or exactly `channels` long, and is never guessed: a slice whose effect on the band axis cannot be determined drops the wavelengths rather than inventing them. + +## Proposed additions + +### 1. Band lookup by wavelength + +The gap a user notices first. +Hyperspectral work is expressed in nanometres, but every accessor on the class takes band indices, so callers hand-compute indices from the `wavelength` list before they can slice. + +The minimal addition is one function that converts, leaving the existing slicing to do the rest: + +```python +def band_at(self, nm: int) -> int: + """The index of the band whose centre is closest to `nm`.""" +``` + +which composes with what already exists: + +```python +red = cube[:, :, cube.band_at(650)] +window = cube[:, :, cube.band_at(600) : cube.band_at(700) + 1] +``` + +Nearest match is the only honest semantic. +The SDK reports wavelengths as `uint32_t` nanometres (`cuvis.h:939`), the grid is whatever the camera's calibration produced, and an exact-match lookup would fail for most inputs a user types. +`band_at` should raise when `wavelength is None`, because there is no defensible answer for a preview or an info layer. + +A richer alternative is an indexer object, `cube.nm[600:700]`, so that nanometres read like slicing. +It is more pleasant at the call site and considerably more machinery: a second indexing protocol to document, test and keep in step with `__getitem__`. +Not worth it for what is fundamentally a coordinate conversion. +`band_at` first; revisit only if call sites turn out to be dominated by ranges. + +### 2. Mean spectrum over a region + +After slicing, the most common hyperspectral operation is averaging a spatial region into one spectrum. +Today that loses the wavelengths, because the shape changes and the ufunc machinery correctly declines to carry metadata across a reduction: + +```python +np.mean(cube, axis=(0, 1)) # plain ndarray, shape (channels,), wavelengths gone +``` + +so the caller reassembles by hand. +A method that returns a single pixel `ImageData` closes the loop and keeps the band axis labelled: + +```python +def mean_spectrum(self) -> "ImageData": + """The spatial mean, as a (1, 1, channels) ImageData carrying this instance's wavelengths.""" +``` + +Then `cube[100:200, 50:150].mean_spectrum()` is the whole region-of-interest workflow, and the result is `is_spectrum` and plots exactly like a QMini reading. +This is the one addition that turns the existing pieces into a workflow rather than adding another spelling of something already possible. + +### 3. Comparison operators + +`np.asarray(cube) > 500` works and is documented. +`cube > 500` raises `TypeError`, and `np.greater(cube, 500)` returns a plain boolean array since the guard added in 3.5.3.2. +The asymmetry is the kind that costs a user ten minutes. + +Adding `__lt__`, `__le__`, `__gt__`, `__ge__` through the existing `_binary_op` factory, returning the plain array rather than rewrapping, removes it for about four lines. + +`__eq__` and `__ne__` must stay out. +Defining them would make `img == img` elementwise, which silently breaks every truth test on the result, and it would take `__hash__` with it unless explicitly restored. +`ImageData` is hashable today and identity comparison is the useful default for a handle-like object. +Asymmetric operator sets are unusual enough to deserve a comment at the definition site saying why. + +### 4. `spectrum` as a method taking pixel coordinates + +Discussed during the hotfix and deliberately left out of it. + +`spectrum` is currently a property restricted to single pixel images. +The general operation is "the band vector at a pixel", which `__getitem__` already performs, except that it returns `(values, wavelengths)` rather than a bare array. +So there are two ways to reach a band vector with two different return types, and the property covers only the `0, 0` case: + +```python +point.spectrum # ndarray +cube[10, 10] # (ndarray, wavelengths) +``` + +A method subsumes both with one return type and no arbitrary restriction: + +```python +def spectrum(self, y: int = 0, x: int = 0) -> np.ndarray: +``` + +`point.spectrum()` keeps reading well, `cube.spectrum(10, 10)` gains what the property could not express, and `is_spectrum` reverts to what it should have been all along: an informational shape check, not the precondition of another member. + +The catch is timing. +Turning a property into a method is a breaking change, so this is free before 3.5.3.2 ships and a deprecation cycle afterwards. +It is listed here rather than applied because it widens a hotfix, but it is the item on this list whose cost grows the fastest. + +### 5. Store the buffer format, or stop requiring it + +Not an addition so much as a wart to resolve, recorded here because it touches the constructor signature. + +`__init__` accepts `dformat`, raises `TypeError` when it is missing, and never reads it. +The format is taken from `img_buf.format` directly, two lines further down. +The three call sites do not even agree on the type they pass: `Measurement.py:108` passes a `DataFormat` enum member, `SessionFile.py:60` and `Viewer.py:46` pass the raw integer. + +Either the value is worth keeping, in which case store it as a public `format` and use it instead of re-reading the buffer, or it is not, in which case drop the parameter. +The current state is the worst of the three: a required argument, inconsistently supplied, with no effect. +Dropping it is technically a signature change, but `ImageData(img_buf, dformat)` is not something callers outside the wrapper construct. + +## Considered and refused + +**`wavelength` as an ndarray.** +It is a list of Python ints today, so callers wrap it in `np.asarray` when they want arithmetic. +Changing the type would break `wavelength == [450, 458]` comparisons, including several in `tests/test_cube_utils.py`, and turn every truth test on the result into an ambiguity error. +Adding a second `wavelength_nm` property alongside it trades that break for a permanent duplicate. +The `np.asarray` at the call site is the smaller cost. + +**Iteration and `__len__`.** +There is no defensible answer to what iterating an image yields. +Rows, pixels and bands are all plausible, and a wrong guess is worse than a `TypeError`. + +**`is_cube`, `is_image`, or other siblings of `is_spectrum`.** +`not img.is_spectrum` already says it. +`is_spectrum` earns its place because the alternative forces callers to know the `(1, 1, channels)` convention; a negation does not clear that bar. + +**More conversion spellings.** +`array`, `to_numpy()` and `np.asarray(img)` are already three ways to reach the same buffer. +The direction of travel should be fewer, not more: `to_numpy()` is the redundant one, and if anything happens here it should be a deprecation. + +## Ordering + +If these are picked up, the order that yields the most per change: + +1. `spectrum(y, x)`, if and only if it happens before 3.5.3.2 ships. Afterwards it drops to last, behind a deprecation cycle. +2. `band_at`. Largest gap, smallest implementation, no interaction with anything else. +3. `mean_spectrum`. Depends on nothing, completes the region-of-interest workflow. +4. Comparison operators. Small and self-contained. +5. The `dformat` cleanup. Internal, do it alongside whichever of the above touches the constructor. diff --git a/cuvis/Measurement.py b/cuvis/Measurement.py index 56cdc1f..8c03023 100644 --- a/cuvis/Measurement.py +++ b/cuvis/Measurement.py @@ -354,7 +354,7 @@ def get_cube_cuda_ipc(self, key: str = "cube", backend: int = 0) -> CudaImageDat cimg.make_ipc(backend) return cimg - def get_cube(self, key: str = "cube") -> Union[ImageData, CudaImageData]: + def get_cube(self, key: str = "cube") -> ImageData | CudaImageData: """Cube through whichever mode is active, so one call site serves both. .. code-block:: python3 diff --git a/cuvis/binding.py b/cuvis/binding.py index 0cc88a0..dc056a5 100644 --- a/cuvis/binding.py +++ b/cuvis/binding.py @@ -162,7 +162,7 @@ def missing_symbols() -> frozenset[str]: return frozenset(getattr(cuvis_il, "missing_symbols", ())) -def unavailable(*names: str) -> Tuple[str, ...]: +def unavailable(*names: str) -> tuple[str, ...]: """Which of the named functions cannot be called, in the order given. A function is unusable for either of two reasons, and callers care about neither: diff --git a/cuvis/cube_utils.py b/cuvis/cube_utils.py index e68ee8b..0203195 100644 --- a/cuvis/cube_utils.py +++ b/cuvis/cube_utils.py @@ -444,12 +444,15 @@ class CudaImageData(object): _TORCH_DTYPE = {1: "uint8", 2: "uint16", 3: "uint32", 4: "float32"} def __init__(self, cuda_buf): + # Set before the check that can raise, so __del__ on the half-built object + # sees None rather than an absent attribute. + self._handle = None # CUVIS_CUDA_MEM (int), owned + self._ipc_handle = None # CUVIS_CUDA_IPC (int), set by make_ipc() if not isinstance(cuda_buf, cuvis_il.cuvis_cuda_imbuffer_t): raise TypeError( "Wrong data type for cuda image buffer: {}".format(type(cuda_buf)) ) - self._handle = cuda_buf.handle # CUVIS_CUDA_MEM (int), owned - self._ipc_handle = None # CUVIS_CUDA_IPC (int), set by make_ipc() + self._handle = cuda_buf.handle self._format = cuda_buf.format self.width = cuda_buf.width self.height = cuda_buf.height @@ -568,12 +571,12 @@ def export_payload(self, backend: int = 0) -> bytes: ) def __del__(self): - try: - if self._ipc_handle is not None: - cuvis_il.cuvis_cuda_ipc_handle_free(self._ipc_handle) - except Exception: - pass - try: - cuvis_il.cuvis_cuda_mem_free(self._handle) - except Exception: - pass + if self._handle is None: + return + # Not wrapped in try/except: a device buffer that fails to free is a VRAM leak, + # and swallowing it here would hide the leak as well as the reason for it. + if self._ipc_handle is not None: + cuvis_il.cuvis_cuda_ipc_handle_free(self._ipc_handle) + self._ipc_handle = None + cuvis_il.cuvis_cuda_mem_free(self._handle) + self._handle = None diff --git a/examples/cuda_ipc_test.py b/examples/cuda_ipc_test.py new file mode 100644 index 0000000..5149c87 --- /dev/null +++ b/examples/cuda_ipc_test.py @@ -0,0 +1,150 @@ +"""Cross-process CUDA IPC end-to-end test. + +Producer: process a measurement, get the cube as a shareable IPC buffer, write the +descriptor + geometry + a host reference to a work dir, then stay alive (it is the +in-process pin; legacy IPC has no cross-process refcount). +Importer (separate process): open the descriptor, wrap the device memory zero-copy as +a torch tensor, and assert it holds the exact same data as the host reference. +Also times the open/import latency. + +Run (orchestrates both processes): + set CUVIS=C:\\Program Files\\Cuvis\\bin + set CUVIS_SETTINGS=C:\\Program Files\\Cuvis\\user\\settings + set PYTHONPATH=C:\\dev\\cuvis_sdk\\cuvis.pyil;C:\\dev\\cuvis_sdk\\cuvis.python + \\Scripts\\python.exe examples\\cuda_ipc_test.py +""" + +import os +import sys +import time +import argparse +import tempfile +from pathlib import Path + +import numpy as np + +KEY = "cube" +N = 100 + + +def _wait_for(path, timeout=60.0): + t0 = time.time() + while not path.exists(): + if time.time() - t0 > timeout: + raise TimeoutError(f"timed out waiting for {path}") + time.sleep(0.02) + + +def producer(workdir: Path): + import cuvis + + cuvis.init(settings_path=os.environ.get("CUVIS_SETTINGS", ".")) + data = os.path.join( + os.path.dirname(__file__), "..", "tests", "test_data", "test_mesu.cu3s" + ) + sess = cuvis.SessionFile(data) + mesu = sess.get_measurement(0) + pc = cuvis.ProcessingContext(sess) + pc.processing_mode = cuvis.ProcessingMode.Raw + pc.apply(mesu) + + backend = int( + os.environ.get("CUVIS_IPC_BACKEND", "0") + ) # 0=auto 1=pool 2=legacy 3=vmm + host = np.array(mesu.cube.array, copy=True) # (h, w, c) uint16 + cimg = mesu.get_cube_cuda_ipc( + KEY, backend=backend + ) # cimg is the in-process pin; keep it alive + payload = cimg.export_payload() # ONE self-contained blob: descriptor + geometry + + np.save(workdir / "ref.npy", host) + (workdir / "payload.bin").write_bytes(payload) + print( + f"[producer] cube {host.shape} {host.dtype}, payload {len(payload)} bytes, " + f"backend_req={backend}, pid {os.getpid()}" + ) + (workdir / "ready").touch() + + _wait_for(workdir / "done", timeout=120.0) + print("[producer] importer finished, releasing buffer") + + +def importer(workdir: Path): + # NOTE: no cuvis SDK init here - only the import-safe cuvis_ipc consumer utilities. + import cuvis_ipc as ipc + + payload = (workdir / "payload.bin").read_bytes() + ref = np.load(workdir / "ref.npy") + + # Correctness: geometry comes from the payload, nothing hard-coded. + with ipc.open(payload) as cube: + backend = cube.backend + got = cube.to_torch().cpu().numpy() + ok = got.shape == ref.shape and got.dtype == ref.dtype and np.array_equal(got, ref) + print( + f"[importer] imported {got.shape} {got.dtype} backend={backend} equals_host={ok}" + ) + if not ok: + (workdir / "done").touch() + return 1 + + # Latency: open + wrap + one sync, N times + import torch + + times = [] + for _ in range(N): + t0 = time.perf_counter_ns() + with ipc.open(payload) as cube: + tt = cube.to_torch() + torch.cuda.synchronize() + t1 = time.perf_counter_ns() + times.append((t1 - t0) / 1000.0) + del tt + a = np.array(times) + print( + f"[importer] open+wrap latency us: p50={np.percentile(a, 50):.1f} " + f"p99={np.percentile(a, 99):.1f} mean={a.mean():.1f} min={a.min():.1f}" + ) + + (workdir / "done").touch() + return 0 + + +def orchestrate(): + import subprocess + + workdir = Path(tempfile.mkdtemp(prefix="cuvis_ipc_")) + print(f"[main] work dir {workdir}") + env = dict(os.environ) + prod = subprocess.Popen( + [sys.executable, __file__, "--role", "producer", "--dir", str(workdir)], env=env + ) + try: + _wait_for(workdir / "ready", timeout=120.0) + rc = subprocess.call( + [sys.executable, __file__, "--role", "importer", "--dir", str(workdir)], + env=env, + ) + prod.wait(timeout=30) + finally: + if prod.poll() is None: + prod.terminate() + print(f"[main] {'PASSED' if rc == 0 else 'FAILED'} (importer rc={rc})") + return rc + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--role", choices=["producer", "importer"]) + ap.add_argument("--dir") + args = ap.parse_args() + if args.role == "producer": + producer(Path(args.dir)) + return 0 + if args.role == "importer": + return importer(Path(args.dir)) + return orchestrate() + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/examples/cuda_misuse_test.py b/examples/cuda_misuse_test.py new file mode 100644 index 0000000..b711b67 --- /dev/null +++ b/examples/cuda_misuse_test.py @@ -0,0 +1,188 @@ +"""Adversarial / misuse tests for the cuvis CUDA-mem C API. + +Exercises the boundary the way a careless caller would: multiple IPC exports of one +buffer, double free, use-after-free, wrong-vault handle mixups, invalid handles, and +whether repeated exports over-allocate device memory. Findings feed the docs. + +Run: + set CUVIS=C:\\Program Files\\Cuvis\\bin + set CUVIS_SETTINGS=C:\\Program Files\\Cuvis\\user\\settings + set PYTHONPATH=C:\\dev\\cuvis_sdk\\cuvis.pyil;C:\\dev\\cuvis_sdk\\cuvis.python + \\Scripts\\python.exe examples\\cuda_misuse_test.py +""" + +import os +import cuvis +from cuvis_il import cuvis_il as il + +OK = il.status_ok +PASS = [] + + +def check(cond, msg): + print((" ok " if cond else " FAIL") + " " + msg) + PASS.append(bool(cond)) + + +def _free_mib(): + """CUDA free memory in MiB, or None if cuda-python is unavailable.""" + try: + from cuda.bindings import runtime as rt + + rt.cudaSetDevice(0) + rt.cudaDeviceSynchronize() + err, free_b, _total = rt.cudaMemGetInfo() + return free_b / (1024 * 1024) + except Exception as e: + print(f" (cudaMemGetInfo unavailable: {e})") + return None + + +def acquire_mem(mesu, key="cube"): + buf = il.cuvis_cuda_imbuffer_t() + st = il.cuvis_measurement_get_data_image_cuda(mesu._handle, key, buf) + assert st == OK, "acquire failed" + return buf.handle + + +def make_ipc(mem): + p = il.new_p_int() + st = il.cuvis_cuda_ipc_handle_create(mem, 0, p) # backend 0 = auto + return st, (il.p_int_value(p) if st == OK else None) + + +def descriptor_bytes(ipc): + d = il.cuvis_cuda_ipc_descriptor_t() + st = il.cuvis_cuda_ipc_get_descriptor(ipc, d) + return st, (il.cuvis_cuda_descriptor_bytes(d) if st == OK else None) + + +def main(): + cuvis.init(settings_path=os.environ.get("CUVIS_SETTINGS", ".")) + data = os.path.join( + os.path.dirname(__file__), "..", "tests", "test_data", "test_mesu.cu3s" + ) + sess = cuvis.SessionFile(data) + mesu = sess.get_measurement(0) + pc = cuvis.ProcessingContext(sess) + pc.processing_mode = cuvis.ProcessingMode.Raw + pc.apply(mesu) + + # --- M1: multiple IPC exports of one buffer --- + print("[M1] export the same buffer as many IPC handles") + mem = acquire_mem(mesu) + N = 64 + before = _free_mib() + ipcs = [] + for _ in range(N): + st, h = make_ipc(mem) + if st != OK: + break + ipcs.append(h) + check( + len(ipcs) == N, f"{N} IPC exports of one buffer all succeeded (got {len(ipcs)})" + ) + check(len(set(ipcs)) == len(ipcs), "each export returned a distinct handle") + after = _free_mib() + if before is not None and after is not None: + delta = before - after + print( + f" device free before={before:.1f} MiB after={after:.1f} MiB delta={delta:.1f} MiB" + ) + check( + delta < 8.0, + f"{N} exports did not allocate N buffers (delta {delta:.1f} MiB << {N}*8 MiB)", + ) + + # all descriptors valid and identical (same underlying buffer) + sts = [descriptor_bytes(h) for h in ipcs] + check(all(st == OK for st, _ in sts), "every IPC handle yields a descriptor") + blobs = {b for _, b in sts if b is not None} + check(len(blobs) == 1, "all descriptors are byte-identical (one physical buffer)") + + # --- M2: buffer stays alive until the LAST reference is freed --- + print("[M2] refcount: free all but one, buffer still usable") + for h in ipcs[:-1]: + check( + il.cuvis_cuda_ipc_handle_free(h) == OK, None + ) if False else il.cuvis_cuda_ipc_handle_free(h) + last = ipcs[-1] + st, _ = descriptor_bytes(last) + check(st == OK, "last IPC handle still valid after the other 63 were freed") + # free the mem handle too; the buffer is still pinned by the last IPC handle + check( + il.cuvis_cuda_mem_free(mem) == OK, + "mem handle freed while one IPC handle remains", + ) + st, _ = descriptor_bytes(last) + check( + st == OK, + "descriptor still valid after mem handle freed (buffer alive via IPC handle)", + ) + check( + il.cuvis_cuda_ipc_handle_free(last) == OK, + "last IPC handle freed -> last reference gone", + ) + + # --- M3: double free is rejected, not a crash --- + print("[M3] double free") + check( + il.cuvis_cuda_ipc_handle_free(last) != OK, + "double free of an IPC handle is rejected", + ) + check(il.cuvis_cuda_mem_free(mem) != OK, "double free of a mem handle is rejected") + + # --- M4: use-after-free is rejected --- + print("[M4] use after free") + v = il.cuvis_cuda_mem_view_t() + check( + il.cuvis_cuda_mem_get_view(mem, v) != OK, + "get_view on a freed mem handle is rejected", + ) + st, _ = descriptor_bytes(last) + check(st != OK, "get_descriptor on a freed IPC handle is rejected") + + # --- M5: invalid / never-existed handles --- + print("[M5] invalid handles") + check( + il.cuvis_cuda_mem_free(999999) != OK, + "free of a never-existed mem handle is rejected", + ) + st, _ = make_ipc(999999) + check(st != OK, "handle_create on an invalid mem handle is rejected") + check( + il.cuvis_cuda_ipc_handle_free(999999) != OK, + "free of a never-existed IPC handle is rejected", + ) + + # --- M6: wrong-vault handle mixups (the two vaults have independent id spaces) --- + print("[M6] cross-vault handle mixups") + mem2 = acquire_mem(mesu) + st, ipc2 = make_ipc(mem2) + print( + f" mem2={mem2} ipc2={ipc2} (independent id spaces - may collide numerically)" + ) + # Passing an IPC handle to a mem function looks it up in the WRONG vault. + v2 = il.cuvis_cuda_mem_view_t() + st_view = il.cuvis_cuda_mem_get_view(ipc2, v2) + print(f" cuvis_cuda_mem_get_view(ipc handle) -> status {st_view}") + # Passing a mem handle to an IPC function likewise. + d2 = il.cuvis_cuda_ipc_descriptor_t() + st_desc = il.cuvis_cuda_ipc_get_descriptor(mem2, d2) + print(f" cuvis_cuda_ipc_get_descriptor(mem handle) -> status {st_desc}") + print( + " (see the hazard note in the docs: handle types are not interchangeable," + ) + print(" and because the vaults number independently a wrong-type handle can") + print(" alias an unrelated same-id handle in the other vault.)") + il.cuvis_cuda_ipc_handle_free(ipc2) + il.cuvis_cuda_mem_free(mem2) + + print() + n_fail = PASS.count(False) + print(f"MISUSE TESTS: {PASS.count(True)} ok, {n_fail} FAIL") + raise SystemExit(1 if n_fail else 0) + + +if __name__ == "__main__": + main() diff --git a/examples/cuda_optin.py b/examples/cuda_optin.py new file mode 100644 index 0000000..63def53 --- /dev/null +++ b/examples/cuda_optin.py @@ -0,0 +1,52 @@ +"""Opt-in CUDA usage: check capabilities, enable, read the cube zero-copy as a tensor. + +Run: + set CUVIS=C:\\Program Files\\Cuvis\\bin + set CUVIS_SETTINGS=C:\\Program Files\\Cuvis\\user\\settings + set PYTHONPATH=C:\\dev\\cuvis_sdk\\cuvis.pyil;C:\\dev\\cuvis_sdk\\cuvis.python + \\Scripts\\python.exe examples\\cuda_optin.py +""" + +import os +import cuvis +from cuvis import ( + cuda, +) # explicit opt-in module; importing cuvis alone does not enable CUDA + +# 1. Inspect what this binary + environment support, then decide. +caps = cuda.capabilities() +print("CUDA capabilities:", caps) +print(" same-process device sharing:", caps.same_process) +print(" cross-process IPC (pool):", caps.ipc_pool) +print(" torch installed:", caps.torch, " cuda-python installed:", caps.cuda_python) + +if not caps.same_process: + print("This build has no CUDA support; nothing to do.") + raise SystemExit(0) + +# 2. Opt in BEFORE loading/processing (also disables the host auto-refresh so the +# GPU cube stays on the device). +cuda.enable() +print("cuda mode enabled:", cuda.is_enabled()) + +cuvis.init(settings_path=os.environ.get("CUVIS_SETTINGS", ".")) +data = os.path.join( + os.path.dirname(__file__), "..", "tests", "test_data", "test_mesu.cu3s" +) +sess = cuvis.SessionFile(data) +mesu = sess.get_measurement(0) +pc = cuvis.ProcessingContext(sess) +pc.processing_mode = cuvis.ProcessingMode.Raw +pc.apply(mesu) + +# 3. get_cube() now routes through CUDA and returns a CudaImageData (raises if the +# device path is unavailable - no silent host copy). +cimg = mesu.get_cube() +print("get_cube() returned:", type(cimg).__name__) + +# 4. Preferred same-process wrap: DLPack (torch owns the lifetime). +if caps.torch: + t = cimg.to_torch() + print("to_torch (DLPack):", tuple(t.shape), t.dtype, t.device) +else: + print("install cuvis[torch] to wrap as a tensor via to_torch()") diff --git a/examples/cuda_smoke.py b/examples/cuda_smoke.py new file mode 100644 index 0000000..3fd07aa --- /dev/null +++ b/examples/cuda_smoke.py @@ -0,0 +1,144 @@ +"""Adaptive smoke test for the CUDA IPC Python wrapping. + +Run: + set CUVIS=C:\\Program Files\\Cuvis\\bin + set PYTHONPATH=C:\\dev\\cuvis_sdk\\cuvis.pyil;C:\\dev\\cuvis_sdk\\cuvis.python + \\Scripts\\python.exe examples\\cuda_smoke.py + +Tiers 1-3 need only the built binding. Tier 4 additionally needs a measurement whose +image data is CUDA-device-backed, and (for the tensor compare) torch. +""" + +import struct +import sys + +FAIL = [] + + +def check(cond, msg): + print((" ok " if cond else " FAIL") + " " + msg) + if not cond: + FAIL.append(msg) + + +# ---- Tier 1: the installed cuvis library provides the CUDA surface ---- +# Asked through cuvis.binding rather than by probing cuvis_il attributes: a function can +# be absent because the binding never wrapped it, or because the loaded cuvis.dll does +# not export it, and binding.unavailable answers both at once. binding.info() names the +# two libraries, which is what tells a version mismatch from a missing feature. +print("[tier 1] SDK provides the CUDA functions") +from cuvis import binding, cuda +from cuvis_il import cuvis_il + +print(binding.info()) +absent = binding.unavailable(cuda.BACKEND_PROBE, *cuda.DEVICE_FUNCTIONS, *cuda.IPC_FUNCTIONS) +check(not absent, "every CUDA function is provided" + (f" (missing: {', '.join(absent)})" if absent else "")) + +# SWIG helpers and structs are not cuvis.h functions, so the library cannot be asked +# about them; they either were compiled into the binding or were not. +for name in ("cuvis_cuda_view_ptr", "cuvis_cuda_descriptor_bytes", + "cuvis_cuda_imbuffer_t", "cuvis_cuda_mem_view_t", + "cuvis_cuda_ipc_descriptor_t"): + check(hasattr(cuvis_il, name), f"cuvis_il.{name} compiled in") + +if absent: + print() + print("Nothing further can run against this SDK. Stopping.") + raise SystemExit(1) + +# Removed in the Phase 16 API simplification - assert they are gone. +removed = [ + "cuvis_measurement_get_data_image_cuda_ipc", + "cuvis_cuda_mem_ref", + "cuvis_cuda_mem_export", + "cuvis_cuda_mem_get_descriptor", + "cuvis_cuda_ipc_export_free", + "cuvis_cuda_mem_get_device_ptr", + "cuvis_cuda_mem_get_device_ordinal", + "cuvis_cuda_ipc_get_last_error_msg", +] +for name in removed: + check(not hasattr(cuvis_il, name), f"cuvis_il.{name} removed") + +# ---- Tier 2: what this device and driver actually support ---- +print("[tier 2] device backends") +caps = cuda.capabilities() +check(caps.same_process, "same-process device sharing available") +for label, code in (("NONE", cuda.BACKEND_NONE), ("POOL", cuda.BACKEND_POOL), + ("LEGACY", cuda.BACKEND_LEGACY), ("VMM", cuda.BACKEND_VMM)): + p = cuvis_il.new_p_int() + st = cuvis_il.cuvis_cuda_ipc_backend_available(code, p) + check(st == cuvis_il.status_ok, f"backend_available({label}) status ok") + print(f" backend {label}: available={cuvis_il.p_int_value(p)}") +print(f" any cross-process backend: {caps.any_ipc}") +print(f" torch={caps.torch} cuda-python={caps.cuda_python}") + +# ---- Tier 3: the two %inline helpers ---- +print("[tier 3] %inline helpers") +desc = cuvis_il.cuvis_cuda_ipc_descriptor_t() +raw = cuvis_il.cuvis_cuda_descriptor_bytes(desc) +check(isinstance(raw, (bytes, bytearray)), "descriptor_bytes returns bytes") +check(len(raw) == 184, f"descriptor is 184 bytes (got {len(raw)})") +head = struct.unpack_from(" get_cube_cuda (needs a device-backed cube)") +import os +import cuvis +from cuvis.cuvis_aux import SDKException + +cuvis.init(settings_path=os.environ.get("CUVIS_SETTINGS", ".")) +data = os.path.join( + os.path.dirname(__file__), "..", "tests", "test_data", "test_mesu.cu3s" +) +if not os.path.exists(data): + print(f" skip: test data not found at {data}") +else: + sess = cuvis.SessionFile(data) + mesu = sess.get_measurement(0) + pc = cuvis.ProcessingContext(sess) + pc.processing_mode = cuvis.ProcessingMode.Raw + pc.apply(mesu) + host = mesu.cube # ImageData (host numpy copy) + print(f" host cube: shape={host.array.shape} dtype={host.array.dtype}") + try: + cimg = mesu.get_cube_cuda("cube") + print( + f" CudaImageData: {cimg.width}x{cimg.height}x{cimg.channels} dtype={cimg.dtype}" + ) + check( + (cimg.height, cimg.width, cimg.channels) == host.array.shape, + "cuda geometry matches host cube shape", + ) + cai = cimg.__cuda_array_interface__ + print( + f" __cuda_array_interface__: shape={cai['shape']} typestr={cai['typestr']}" + ) + try: + import numpy as np + import torch # noqa + + t = cimg.to_torch() + check(tuple(t.shape) == host.array.shape, "torch tensor shape matches host") + check( + np.array_equal(t.cpu().numpy(), host.array), + "torch tensor data equals host cube", + ) + except ImportError: + print(" (torch not installed - skipping zero-copy tensor compare)") + except SDKException as e: + print(f" get_cube_cuda raised SDKException: {e}") + print( + " -> plumbing reached the C function; the cube is not CUDA-device-backed," + ) + print(" so the device happy-path needs a GPU-resident measurement.") + +print() +if FAIL: + print(f"SMOKE FAILED: {len(FAIL)} check(s) failed") + sys.exit(1) +print("SMOKE PASSED (tiers that could run)") diff --git a/examples/cuda_tensor_bench.py b/examples/cuda_tensor_bench.py new file mode 100644 index 0000000..f3cf898 --- /dev/null +++ b/examples/cuda_tensor_bench.py @@ -0,0 +1,183 @@ +"""Benchmark: get a measurement's cube into a torch CUDA tensor, three ways, timed. + +Each way ends with a tensor on cuda. Before timing, every path is checked to hold +the EXACT same data as the host cube (np.array_equal), so the zero-copy paths are +proven to read the same pixels, not just the same shape. + + A - host round-trip (status quo): cuvis_measurement_get_data_image (device->host + copy in SDK) -> cuvis_read_imbuf (numpy) -> torch.from_numpy().to(cuda). 2 copies. + B - DLPack zero-copy: get_cube_cuda().to_torch(). pointer wrap, 0 copies. + C - __cuda_array_interface__ zero-copy: torch.as_tensor(cimg). pointer wrap, 0 copies. + +Run: + set CUVIS=C:\\Program Files\\Cuvis\\bin + set CUVIS_SETTINGS=C:\\Program Files\\Cuvis\\user\\settings + set PYTHONPATH=C:\\dev\\cuvis_sdk\\cuvis.pyil;C:\\dev\\cuvis_sdk\\cuvis.python + \\Scripts\\python.exe examples\\cuda_tensor_bench.py +""" + +import os +import time + +import numpy as np +import torch + +import cuvis +from cuvis.cuvis_aux import SDKException +from cuvis_il import cuvis_il + +KEY = "cube" +WARMUP = 5 +N = 100 + +_FMT_TO_READER = { + 1: cuvis_il.cuvis_read_imbuf_uint8, + 2: cuvis_il.cuvis_read_imbuf_uint16, + 3: cuvis_il.cuvis_read_imbuf_uint32, + 4: cuvis_il.cuvis_read_imbuf_float32, +} + + +def host_array(mesu): + """Fresh host fetch of the cube (device->host copy in the SDK), as a numpy copy.""" + buf = cuvis_il.cuvis_imbuffer_t() + if cuvis_il.status_ok != cuvis_il.cuvis_measurement_get_data_image( + mesu._handle, KEY, buf + ): + raise SDKException() + arr = _FMT_TO_READER[buf.format](buf) # numpy view over the SDK buffer + return np.array(arr, copy=True) # detach from the SDK buffer + + +def path_a(mesu): + """Host round-trip -> cuda tensor. Two data movements.""" + buf = cuvis_il.cuvis_imbuffer_t() + if cuvis_il.status_ok != cuvis_il.cuvis_measurement_get_data_image( + mesu._handle, KEY, buf + ): + raise SDKException() + arr = _FMT_TO_READER[buf.format](buf) + t = torch.from_numpy(arr).to("cuda") + del buf + return t + + +def path_b(mesu): + """DLPack zero-copy.""" + return mesu.get_cube_cuda(KEY).to_torch() + + +def path_c(mesu): + """__cuda_array_interface__ zero-copy. Keep cimg alive for the tensor's use.""" + cimg = mesu.get_cube_cuda(KEY) + t = torch.as_tensor(cimg, device=f"cuda:{cimg._view()[2]}") + return t, cimg # return cimg so caller keeps it alive + + +def bench(fn, n=N, warmup=WARMUP): + for _ in range(warmup): + r = fn() + del r + torch.cuda.synchronize() + times = [] + for _ in range(n): + t0 = time.perf_counter_ns() + r = fn() + torch.cuda.synchronize() + t1 = time.perf_counter_ns() + times.append((t1 - t0) / 1000.0) # us + del r + a = np.array(times) + return dict( + p50=np.percentile(a, 50), p99=np.percentile(a, 99), mean=a.mean(), min=a.min() + ) + + +def main(): + cuvis.init(settings_path=os.environ.get("CUVIS_SETTINGS", ".")) + data = os.path.join( + os.path.dirname(__file__), "..", "tests", "test_data", "test_mesu.cu3s" + ) + sess = cuvis.SessionFile(data) + mesu = sess.get_measurement(0) + pc = cuvis.ProcessingContext(sess) + pc.processing_mode = cuvis.ProcessingMode.Raw + pc.apply(mesu) + + ref = host_array(mesu) + print( + f"host reference cube: shape={ref.shape} dtype={ref.dtype} bytes={ref.nbytes}" + ) + + # Preflight: device-backed cube available? + try: + _ = mesu.get_cube_cuda(KEY) + except SDKException as e: + print(f"PREFLIGHT FAILED: get_cube_cuda raised: {e}") + print("Cube is not device-backed; running only path A (host round-trip).") + stats = bench(lambda: path_a(mesu)) + _print_table({"A host round-trip": stats}, ref) + return + + # ---- Correctness: each path holds the EXACT same data as the host cube ---- + print("\n[correctness] each path vs host cube (exact equality)") + ta = path_a(mesu) + tb = path_b(mesu) + tc, _c_keep = path_c(mesu) + + checks = { + "A host round-trip": ta, + "B DLPack zero-copy": tb, + "C CAI zero-copy": tc, + } + all_ok = True + for name, t in checks.items(): + got = t.cpu().numpy() + shape_ok = got.shape == ref.shape + dtype_ok = got.dtype == ref.dtype + data_ok = shape_ok and dtype_ok and np.array_equal(got, ref) + # cross-check the zero-copy tensors are bit-identical to A on-device too + eq_a = ( + bool(torch.equal(t.to("cuda"), ta.to("cuda"))) + if t.dtype == ta.dtype + else False + ) + status = "ok " if data_ok and eq_a else "FAIL" + all_ok = all_ok and data_ok and eq_a + print( + f" {status} {name}: shape={got.shape} dtype={got.dtype} " + f"equals_host={data_ok} equals_A={eq_a}" + ) + if not all_ok: + print( + "\nDATA MISMATCH - not all paths carry identical data. Aborting before timing." + ) + raise SystemExit(1) + print(" all paths carry identical data.") + del ta, tb, tc, _c_keep + + # ---- Timing ---- + torch.zeros(1, device="cuda") # force context + results = { + "A host round-trip": bench(lambda: path_a(mesu)), + "B DLPack zero-copy": bench(lambda: path_b(mesu)), + "C CAI zero-copy": bench(lambda: path_c(mesu)), + } + _print_table(results, ref) + + +def _print_table(results, ref): + print(f"\ncube {ref.shape} {ref.dtype} ({ref.nbytes} bytes), N={N} iters") + print(f"{'way':<22}{'p50 us':>12}{'p99 us':>12}{'mean us':>12}{'min us':>12}") + print("-" * 70) + for name, s in results.items(): + print( + f"{name:<22}{s['p50']:>12.1f}{s['p99']:>12.1f}{s['mean']:>12.1f}{s['min']:>12.1f}" + ) + print( + "\nA moves ~8 MB twice (device->host->device); B/C wrap the device pointer (no copy)." + ) + + +if __name__ == "__main__": + main() diff --git a/examples/ipc_consumer.py b/examples/ipc_consumer.py new file mode 100644 index 0000000..375fef7 --- /dev/null +++ b/examples/ipc_consumer.py @@ -0,0 +1,28 @@ +"""Standalone cross-process IPC consumer. + +Runs in a process that never initialized the cuvis SDK: it imports ONLY cuvis_ipc (plus +cuda-python + torch) and needs no CUVIS env var and no cuvis.dll. Point it at a payload file +written by a producer (cimg.export_payload()); it opens the shared buffer and reads it. + + python ipc_consumer.py +""" + +import sys +import cuvis_ipc as ipc # import-safe: no SDK init, no CUVIS env, no cuvis.dll + + +def main(): + if len(sys.argv) != 2: + print("usage: python ipc_consumer.py ") + return 2 + payload = open(sys.argv[1], "rb").read() + with ipc.open(payload) as cube: + t = cube.to_torch() # correct shape + dtype from the payload; zero-copy + print( + f"opened shape={tuple(t.shape)} dtype={t.dtype} device={t.device} sum={int(t.sum())}" + ) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tests/test_handle_lifecycle.py b/tests/test_handle_lifecycle.py index 23b4c8d..5d950d5 100644 --- a/tests/test_handle_lifecycle.py +++ b/tests/test_handle_lifecycle.py @@ -41,6 +41,9 @@ def _construct(make): "CubeExporter": lambda: cuvis.CubeExporter("not export settings"), "Worker": lambda: cuvis.Worker("not worker settings"), "Viewer": lambda: cuvis.Viewer("not viewer settings"), + # Raises TypeError on a CUDA-capable binding and AttributeError on one without the + # CUDA surface; either way __del__ must stay quiet about the half-built object. + "CudaImageData": lambda: cuvis.CudaImageData("not a cuda buffer"), }