diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index d5fdc5d7551..14893fbd3c0 100644 --- a/cuda_core/cuda/core/_device.pyi +++ b/cuda_core/cuda/core/_device.pyi @@ -12,6 +12,10 @@ from cuda.core._memory._buffer import Buffer, MemoryResource from cuda.core._stream import IsStreamType, Stream from cuda.core._utils.cuda_utils import ComputeCapability from cuda.core.graph import GraphBuilder +from cuda.core.texture import (MipmappedArray, MipmappedArrayOptions, + OpaqueArray, OpaqueArrayOptions, + ResourceDescriptor, SurfaceObject, + TextureObject, TextureObjectOptions) class DeviceProperties: @@ -909,5 +913,114 @@ class Device: Newly created graph builder object. """ + + def create_opaque_array(self, options: OpaqueArrayOptions) -> OpaqueArray: + """Create an :obj:`~cuda.core.texture.OpaqueArray` on the current device. + + Allocates an opaque, hardware-laid-out CUDA array for texture/surface + access. The array is created in the current CUDA context, so make this + device current with :meth:`set_current` before calling (mirroring + :meth:`create_stream` / :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + options : :obj:`~cuda.core.texture.OpaqueArrayOptions` + Allocation options (shape, format, channels, surface flag). + + Returns + ------- + :obj:`~cuda.core.texture.OpaqueArray` + Newly created opaque array. + + .. versionadded:: 1.1.0 + """ + + def create_mipmapped_array(self, options: MipmappedArrayOptions) -> MipmappedArray: + """Create a :obj:`~cuda.core.texture.MipmappedArray` on the current device. + + Allocates a mipmapped CUDA array for texture/surface access across + levels. The array is created in the current CUDA context, so make this + device current with :meth:`set_current` before calling (mirroring + :meth:`create_stream` / :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + options : :obj:`~cuda.core.texture.MipmappedArrayOptions` + Allocation options (shape, format, channels, levels, surface flag). + + Returns + ------- + :obj:`~cuda.core.texture.MipmappedArray` + Newly created mipmapped array. + + .. versionadded:: 1.1.0 + """ + + def create_texture_object(self, *, resource: ResourceDescriptor, options: TextureObjectOptions | None=None) -> TextureObject: + """Create a :obj:`~cuda.core.texture.TextureObject` on the current device. + + Binds a resource (an :obj:`~cuda.core.texture.OpaqueArray` / + :obj:`~cuda.core.texture.MipmappedArray` / linear or pitch2d + :obj:`~cuda.core.Buffer`, wrapped in a + :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless texture for + kernel-side sampled reads. The object is created in the current CUDA + context, so make this device current with :meth:`set_current` before + calling (mirroring :meth:`create_stream` / :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + resource : :obj:`~cuda.core.texture.ResourceDescriptor` + The memory backing the texture. + options : :obj:`~cuda.core.texture.TextureObjectOptions` + Sampling state (address/filter/read modes, normalization, etc.). + + Returns + ------- + :obj:`~cuda.core.texture.TextureObject` + Newly created texture object. + + .. versionadded:: 1.1.0 + """ + + def create_surface_object(self, *, resource: ResourceDescriptor) -> SurfaceObject: + """Create a :obj:`~cuda.core.texture.SurfaceObject` on the current device. + + Binds an :obj:`~cuda.core.texture.OpaqueArray` (via a + :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless surface for + kernel-side typed load/store. The backing array must have been created + with ``is_surface_load_store=True``. The object is created in the + current CUDA context, so make this device current with + :meth:`set_current` before calling (mirroring :meth:`create_stream` / + :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + resource : :obj:`~cuda.core.texture.ResourceDescriptor` + Must wrap an :obj:`~cuda.core.texture.OpaqueArray` allocated with + ``is_surface_load_store=True``. + + Returns + ------- + :obj:`~cuda.core.texture.SurfaceObject` + Newly created surface object. + + .. versionadded:: 1.1.0 + """ _tls = threading.local() _lock = threading.Lock() \ No newline at end of file diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index 7b27884b312..ea5bb62cc1b 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -44,6 +44,16 @@ from typing import TYPE_CHECKING if TYPE_CHECKING: import cuda.core.system # no-cython-lint from cuda.core.graph import GraphBuilder + from cuda.core.texture import ( + MipmappedArray, + MipmappedArrayOptions, + OpaqueArray, + OpaqueArrayOptions, + ResourceDescriptor, + SurfaceObject, + TextureObject, + TextureObjectOptions, + ) # TODO: I prefer to type these as "cdef object" and avoid accessing them from within Python, # but it seems it is very convenient to expose them for testing purposes... @@ -1456,6 +1466,133 @@ class Device: self._check_context_initialized() return GraphBuilder._init(self.create_stream()) + def create_opaque_array(self, options: OpaqueArrayOptions) -> OpaqueArray: + """Create an :obj:`~cuda.core.texture.OpaqueArray` on the current device. + + Allocates an opaque, hardware-laid-out CUDA array for texture/surface + access. The array is created in the current CUDA context, so make this + device current with :meth:`set_current` before calling (mirroring + :meth:`create_stream` / :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + options : :obj:`~cuda.core.texture.OpaqueArrayOptions` + Allocation options (shape, format, channels, surface flag). + + Returns + ------- + :obj:`~cuda.core.texture.OpaqueArray` + Newly created opaque array. + + .. versionadded:: 1.1.0 + """ + from cuda.core.texture._array import _create_opaque_array + + self._check_context_initialized() + return _create_opaque_array(options) + + def create_mipmapped_array(self, options: MipmappedArrayOptions) -> MipmappedArray: + """Create a :obj:`~cuda.core.texture.MipmappedArray` on the current device. + + Allocates a mipmapped CUDA array for texture/surface access across + levels. The array is created in the current CUDA context, so make this + device current with :meth:`set_current` before calling (mirroring + :meth:`create_stream` / :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + options : :obj:`~cuda.core.texture.MipmappedArrayOptions` + Allocation options (shape, format, channels, levels, surface flag). + + Returns + ------- + :obj:`~cuda.core.texture.MipmappedArray` + Newly created mipmapped array. + + .. versionadded:: 1.1.0 + """ + from cuda.core.texture._mipmapped_array import _create_mipmapped_array + + self._check_context_initialized() + return _create_mipmapped_array(options) + + def create_texture_object( + self, *, resource: ResourceDescriptor, options: TextureObjectOptions | None = None + ) -> TextureObject: + """Create a :obj:`~cuda.core.texture.TextureObject` on the current device. + + Binds a resource (an :obj:`~cuda.core.texture.OpaqueArray` / + :obj:`~cuda.core.texture.MipmappedArray` / linear or pitch2d + :obj:`~cuda.core.Buffer`, wrapped in a + :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless texture for + kernel-side sampled reads. The object is created in the current CUDA + context, so make this device current with :meth:`set_current` before + calling (mirroring :meth:`create_stream` / :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + resource : :obj:`~cuda.core.texture.ResourceDescriptor` + The memory backing the texture. + options : :obj:`~cuda.core.texture.TextureObjectOptions` + Sampling state (address/filter/read modes, normalization, etc.). + + Returns + ------- + :obj:`~cuda.core.texture.TextureObject` + Newly created texture object. + + .. versionadded:: 1.1.0 + """ + from cuda.core.texture._texture import _create_texture_object + + self._check_context_initialized() + return _create_texture_object(resource, options) + + def create_surface_object(self, *, resource: ResourceDescriptor) -> SurfaceObject: + """Create a :obj:`~cuda.core.texture.SurfaceObject` on the current device. + + Binds an :obj:`~cuda.core.texture.OpaqueArray` (via a + :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless surface for + kernel-side typed load/store. The backing array must have been created + with ``is_surface_load_store=True``. The object is created in the + current CUDA context, so make this device current with + :meth:`set_current` before calling (mirroring :meth:`create_stream` / + :meth:`create_event`). + + Note + ---- + Device must be initialized. + + Parameters + ---------- + resource : :obj:`~cuda.core.texture.ResourceDescriptor` + Must wrap an :obj:`~cuda.core.texture.OpaqueArray` allocated with + ``is_surface_load_store=True``. + + Returns + ------- + :obj:`~cuda.core.texture.SurfaceObject` + Newly created surface object. + + .. versionadded:: 1.1.0 + """ + from cuda.core.texture._surface import _create_surface_object + + self._check_context_initialized() + return _create_surface_object(resource) + cdef inline int Device_ensure_cuda_initialized() except? -1: """Initialize CUDA driver and check version compatibility (once per process).""" diff --git a/cuda_core/cuda/core/texture/__init__.py b/cuda_core/cuda/core/texture/__init__.py index e4414ed28ad..cd3e355f324 100644 --- a/cuda_core/cuda/core/texture/__init__.py +++ b/cuda_core/cuda/core/texture/__init__.py @@ -10,30 +10,31 @@ Import these types from here, e.g.:: - from cuda.core.texture import OpaqueArray, TextureObject, TextureDescriptor + from cuda.core.texture import OpaqueArray, TextureObject, TextureObjectOptions + +The associated enumerations (:class:`~cuda.core.typing.ArrayFormatType`, +:class:`~cuda.core.typing.AddressModeType`, +:class:`~cuda.core.typing.FilterModeType`, +:class:`~cuda.core.typing.ReadModeType`) live in :mod:`cuda.core.typing` +alongside the other ``cuda.core`` enumerations. """ -from cuda.core.texture._array import ArrayFormat, OpaqueArray -from cuda.core.texture._mipmapped_array import MipmappedArray +from cuda.core.texture._array import OpaqueArray, OpaqueArrayOptions +from cuda.core.texture._mipmapped_array import MipmappedArray, MipmappedArrayOptions from cuda.core.texture._surface import SurfaceObject from cuda.core.texture._texture import ( - AddressMode, - FilterMode, - ReadMode, ResourceDescriptor, - TextureDescriptor, TextureObject, + TextureObjectOptions, ) __all__ = [ - "AddressMode", - "ArrayFormat", - "FilterMode", "MipmappedArray", + "MipmappedArrayOptions", "OpaqueArray", - "ReadMode", + "OpaqueArrayOptions", "ResourceDescriptor", "SurfaceObject", - "TextureDescriptor", "TextureObject", + "TextureObjectOptions", ] diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 75dadb46d98..380c2fe1c10 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -2,24 +2,41 @@ from __future__ import annotations -from enum import IntEnum +from dataclasses import dataclass +import numpy from cuda.bindings import cydriver - - -class ArrayFormat(IntEnum): - """Element format for a :class:`OpaqueArray` allocation. - - Mirrors ``CUarray_format`` from the CUDA driver API. +from cuda.core.typing import ArrayFormatType + + +@dataclass +class OpaqueArrayOptions: + """Options for :meth:`cuda.core.Device.create_opaque_array`. + + Attributes + ---------- + shape : tuple of int + ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in + elements. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. + num_channels : int + Channels per element. Must be 1, 2, or 4. + is_surface_load_store : bool + If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so the array can be + bound as a :class:`~cuda.core.texture.SurfaceObject` for kernel-side + writes. Default False. + + .. versionadded:: 1.1.0 """ - UINT8 = cydriver.CU_AD_FORMAT_UNSIGNED_INT8 - UINT16 = cydriver.CU_AD_FORMAT_UNSIGNED_INT16 - UINT32 = cydriver.CU_AD_FORMAT_UNSIGNED_INT32 - INT8 = cydriver.CU_AD_FORMAT_SIGNED_INT8 - INT16 = cydriver.CU_AD_FORMAT_SIGNED_INT16 - INT32 = cydriver.CU_AD_FORMAT_SIGNED_INT32 - FLOAT16 = cydriver.CU_AD_FORMAT_HALF - FLOAT32 = cydriver.CU_AD_FORMAT_FLOAT + shape: tuple[int, ...] + format: object + num_channels: int + is_surface_load_store: bool = False + + def __post_init__(self): + ... class OpaqueArray: """An opaque, hardware-laid-out GPU allocation for texture/surface access. @@ -38,9 +55,11 @@ class OpaqueArray: linear :class:`Buffer` yourself (e.g. ``mr.allocate(arr.size_bytes, stream=s)``) and copy. - Construct via :meth:`from_descriptor`. Only plain 1D/2D/3D allocations are - supported in this initial version; layered/cubemap/sparse variants will - follow once their shape semantics are settled. + Construct via :meth:`cuda.core.Device.create_opaque_array`. Only plain + 1D/2D/3D allocations are supported in this initial version; layered/cubemap/ + sparse variants will follow once their shape semantics are settled. + + .. versionadded:: 1.1.0 """ def close(self): @@ -55,29 +74,6 @@ class OpaqueArray: def __init__(self, *args, **kwargs): ... - @classmethod - def from_descriptor(cls, *, shape, format, num_channels, is_surface_load_store=False): - """Allocate a new CUDA array. - - Parameters - ---------- - shape : tuple of int - ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` - in elements. - format : ArrayFormat - Element format. - num_channels : int - Channels per element. Must be 1, 2, or 4. - is_surface_load_store : bool - If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so the array - can be bound as a :class:`SurfaceObject` for kernel-side writes. - Default False. - - Returns - ------- - OpaqueArray - """ - @classmethod def _from_handle(cls, handle: int, owning: bool, *, device_id=None): """Wrap an externally-allocated ``CUarray``. @@ -98,14 +94,14 @@ class OpaqueArray: @property def format(self): - """The element :class:`ArrayFormat`.""" + """The element :class:`~cuda.core.typing.ArrayFormatType`.""" @property def num_channels(self): """Channels per element (1, 2, or 4).""" @property - def element_size(self): + def element_bytes(self): """Bytes per element (format size * channels).""" @property @@ -153,7 +149,7 @@ class OpaqueArray: @property def size_bytes(self): - """Total bytes of array storage (``prod(shape) * element_size``).""" + """Total bytes of array storage (``prod(shape) * element_bytes``).""" def __enter__(self): ... @@ -163,12 +159,38 @@ class OpaqueArray: def __repr__(self): ... -_FORMAT_ELEM_SIZE = {int(ArrayFormat.UINT8): 1, int(ArrayFormat.INT8): 1, int(ArrayFormat.UINT16): 2, int(ArrayFormat.INT16): 2, int(ArrayFormat.FLOAT16): 2, int(ArrayFormat.UINT32): 4, int(ArrayFormat.INT32): 4, int(ArrayFormat.FLOAT32): 4} +_ARRAYFORMAT_TO_CU = {ArrayFormatType.UINT8: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT8), ArrayFormatType.UINT16: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT16), ArrayFormatType.UINT32: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT32), ArrayFormatType.INT8: int(cydriver.CU_AD_FORMAT_SIGNED_INT8), ArrayFormatType.INT16: int(cydriver.CU_AD_FORMAT_SIGNED_INT16), ArrayFormatType.INT32: int(cydriver.CU_AD_FORMAT_SIGNED_INT32), ArrayFormatType.FLOAT16: int(cydriver.CU_AD_FORMAT_HALF), ArrayFormatType.FLOAT32: int(cydriver.CU_AD_FORMAT_FLOAT)} +_CU_TO_ARRAYFORMAT = {cu: fmt for fmt, cu in _ARRAYFORMAT_TO_CU.items()} +_NUMPY_DTYPE_TO_ARRAYFORMAT = {numpy.dtype(fmt.value): fmt for fmt in ArrayFormatType} +_FORMAT_ELEM_SIZE = {_ARRAYFORMAT_TO_CU[ArrayFormatType.UINT8]: 1, _ARRAYFORMAT_TO_CU[ArrayFormatType.INT8]: 1, _ARRAYFORMAT_TO_CU[ArrayFormatType.UINT16]: 2, _ARRAYFORMAT_TO_CU[ArrayFormatType.INT16]: 2, _ARRAYFORMAT_TO_CU[ArrayFormatType.FLOAT16]: 2, _ARRAYFORMAT_TO_CU[ArrayFormatType.UINT32]: 4, _ARRAYFORMAT_TO_CU[ArrayFormatType.INT32]: 4, _ARRAYFORMAT_TO_CU[ArrayFormatType.FLOAT32]: 4} + +def _normalize_array_format(format): + """Coerce ``format`` to an :class:`ArrayFormatType`. + + Accepts, in order of preference: + + * an :class:`ArrayFormatType`; + * a plain ``str`` naming one of its values (e.g. ``"float32"``); + * a NumPy dtype object (or anything ``numpy.dtype()`` accepts, such as + ``numpy.float32``) whose canonical dtype maps 1:1 to one of the eight + supported formats. + + Raises :class:`ValueError` on anything else.""" def _validate_format_channels(format, num_channels): """Validate the ``(format, num_channels)`` pair shared by the array, - mipmap, and texture factories. Raises on an invalid combination.""" + mipmap, and texture factories. Returns the normalized + :class:`ArrayFormatType`. Raises on an invalid combination.""" def _validate_array_shape(shape): """Coerce ``shape`` to a tuple of ints and validate rank (1-3) and that - every extent is >= 1. Returns the normalized tuple.""" \ No newline at end of file + every extent is >= 1. Returns the normalized tuple.""" + +def _create_opaque_array(options): + """Allocate a new :class:`OpaqueArray` on the current device. + + Backs :meth:`cuda.core.Device.create_opaque_array`. ``options`` is an + :class:`OpaqueArrayOptions` (or a mapping accepted by it); it is validated + at construction, so ``shape`` is already a normalized tuple and ``format`` + an :class:`~cuda.core.typing.ArrayFormatType`. + """ \ No newline at end of file diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index 82a08b045dc..0a1cb671daf 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -25,44 +25,100 @@ from cuda.core._utils.cuda_utils cimport ( _get_current_device_id, ) -from enum import IntEnum - +import numpy + +from dataclasses import dataclass + +from cuda.core._utils.cuda_utils import check_or_create_options +from cuda.core.typing import ArrayFormatType + + +# Bridge between the public ArrayFormatType StrEnum and the driver +# CUarray_format integer values. OpaqueArray stores the driver int internally +# (see ._format), so all conversions funnel through these two maps. +_ARRAYFORMAT_TO_CU = { + ArrayFormatType.UINT8: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT8), + ArrayFormatType.UINT16: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT16), + ArrayFormatType.UINT32: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT32), + ArrayFormatType.INT8: int(cydriver.CU_AD_FORMAT_SIGNED_INT8), + ArrayFormatType.INT16: int(cydriver.CU_AD_FORMAT_SIGNED_INT16), + ArrayFormatType.INT32: int(cydriver.CU_AD_FORMAT_SIGNED_INT32), + ArrayFormatType.FLOAT16: int(cydriver.CU_AD_FORMAT_HALF), + ArrayFormatType.FLOAT32: int(cydriver.CU_AD_FORMAT_FLOAT), +} +_CU_TO_ARRAYFORMAT = {cu: fmt for fmt, cu in _ARRAYFORMAT_TO_CU.items()} -class ArrayFormat(IntEnum): - """Element format for a :class:`OpaqueArray` allocation. - Mirrors ``CUarray_format`` from the CUDA driver API. - """ - UINT8 = cydriver.CU_AD_FORMAT_UNSIGNED_INT8 - UINT16 = cydriver.CU_AD_FORMAT_UNSIGNED_INT16 - UINT32 = cydriver.CU_AD_FORMAT_UNSIGNED_INT32 - INT8 = cydriver.CU_AD_FORMAT_SIGNED_INT8 - INT16 = cydriver.CU_AD_FORMAT_SIGNED_INT16 - INT32 = cydriver.CU_AD_FORMAT_SIGNED_INT32 - FLOAT16 = cydriver.CU_AD_FORMAT_HALF - FLOAT32 = cydriver.CU_AD_FORMAT_FLOAT +# Every ArrayFormatType value is spelled as a NumPy dtype name, so the eight +# formats map 1:1 to NumPy dtypes. This lets callers pass a dtype object (or +# anything numpy.dtype() accepts) instead of the enum, matching the precedent +# set by TensorMapDescriptorOptions.data_type. +_NUMPY_DTYPE_TO_ARRAYFORMAT = { + numpy.dtype(fmt.value): fmt for fmt in ArrayFormatType +} -# Bytes per element (single channel) for each format. +# Bytes per element (single channel), keyed by the driver CUarray_format int. _FORMAT_ELEM_SIZE = { - int(ArrayFormat.UINT8): 1, - int(ArrayFormat.INT8): 1, - int(ArrayFormat.UINT16): 2, - int(ArrayFormat.INT16): 2, - int(ArrayFormat.FLOAT16): 2, - int(ArrayFormat.UINT32): 4, - int(ArrayFormat.INT32): 4, - int(ArrayFormat.FLOAT32): 4, + _ARRAYFORMAT_TO_CU[ArrayFormatType.UINT8]: 1, + _ARRAYFORMAT_TO_CU[ArrayFormatType.INT8]: 1, + _ARRAYFORMAT_TO_CU[ArrayFormatType.UINT16]: 2, + _ARRAYFORMAT_TO_CU[ArrayFormatType.INT16]: 2, + _ARRAYFORMAT_TO_CU[ArrayFormatType.FLOAT16]: 2, + _ARRAYFORMAT_TO_CU[ArrayFormatType.UINT32]: 4, + _ARRAYFORMAT_TO_CU[ArrayFormatType.INT32]: 4, + _ARRAYFORMAT_TO_CU[ArrayFormatType.FLOAT32]: 4, } +def _normalize_array_format(format): + """Coerce ``format`` to an :class:`ArrayFormatType`. + + Accepts, in order of preference: + + * an :class:`ArrayFormatType`; + * a plain ``str`` naming one of its values (e.g. ``"float32"``); + * a NumPy dtype object (or anything ``numpy.dtype()`` accepts, such as + ``numpy.float32``) whose canonical dtype maps 1:1 to one of the eight + supported formats. + + Raises :class:`ValueError` on anything else.""" + if isinstance(format, ArrayFormatType): + return format + if isinstance(format, str): + try: + return ArrayFormatType(format) + except ValueError as e: + valid = ", ".join(repr(f.value) for f in ArrayFormatType) + raise ValueError( + f"format must be an ArrayFormatType or one of {{{valid}}}, got {format!r}" + ) from e + # Fall back to interpreting ``format`` as a NumPy dtype (dtype object, + # scalar type, etc.). Unknown dtypes are reported against the supported set. + try: + dt = numpy.dtype(format) + except TypeError as e: + raise ValueError( + f"format must be an ArrayFormatType, str, or NumPy dtype, got {format!r}" + ) from e + try: + return _NUMPY_DTYPE_TO_ARRAYFORMAT[dt] + except KeyError as e: + valid = ", ".join(repr(f.value) for f in ArrayFormatType) + raise ValueError( + f"NumPy dtype {dt!r} has no ArrayFormatType equivalent; " + f"supported formats: {{{valid}}}" + ) from e + + def _validate_format_channels(format, num_channels): """Validate the ``(format, num_channels)`` pair shared by the array, - mipmap, and texture factories. Raises on an invalid combination.""" - if not isinstance(format, ArrayFormat): - raise TypeError(f"format must be an ArrayFormat, got {type(format).__name__}") + mipmap, and texture factories. Returns the normalized + :class:`ArrayFormatType`. Raises on an invalid combination.""" + fmt = _normalize_array_format(format) if isinstance(num_channels, bool) or num_channels not in (1, 2, 4): raise ValueError(f"num_channels must be 1, 2, or 4, got {num_channels!r}") + return fmt def _validate_array_shape(shape): @@ -80,6 +136,38 @@ def _validate_array_shape(shape): return shape_t +@dataclass +class OpaqueArrayOptions: + """Options for :meth:`cuda.core.Device.create_opaque_array`. + + Attributes + ---------- + shape : tuple of int + ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in + elements. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. + num_channels : int + Channels per element. Must be 1, 2, or 4. + is_surface_load_store : bool + If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so the array can be + bound as a :class:`~cuda.core.texture.SurfaceObject` for kernel-side + writes. Default False. + + .. versionadded:: 1.1.0 + """ + + shape: tuple[int, ...] + format: object + num_channels: int + is_surface_load_store: bool = False + + def __post_init__(self): + self.format = _validate_format_channels(self.format, self.num_channels) + self.shape = _validate_array_shape(self.shape) + + cdef void _fill_array_endpoint( cydriver.CUDA_MEMCPY3D* p, OpaqueArray arr, bint is_src ) noexcept: @@ -245,71 +333,19 @@ cdef class OpaqueArray: linear :class:`Buffer` yourself (e.g. ``mr.allocate(arr.size_bytes, stream=s)``) and copy. - Construct via :meth:`from_descriptor`. Only plain 1D/2D/3D allocations are - supported in this initial version; layered/cubemap/sparse variants will - follow once their shape semantics are settled. + Construct via :meth:`cuda.core.Device.create_opaque_array`. Only plain + 1D/2D/3D allocations are supported in this initial version; layered/cubemap/ + sparse variants will follow once their shape semantics are settled. + + .. versionadded:: 1.1.0 """ def __init__(self, *args, **kwargs): raise RuntimeError( - "OpaqueArray cannot be instantiated directly. Use OpaqueArray.from_descriptor()." + "OpaqueArray cannot be instantiated directly. " + "Use Device.create_opaque_array()." ) - @classmethod - def from_descriptor(cls, *, shape, format, num_channels, is_surface_load_store=False): - """Allocate a new CUDA array. - - Parameters - ---------- - shape : tuple of int - ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` - in elements. - format : ArrayFormat - Element format. - num_channels : int - Channels per element. Must be 1, 2, or 4. - is_surface_load_store : bool - If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so the array - can be bound as a :class:`SurfaceObject` for kernel-side writes. - Default False. - - Returns - ------- - OpaqueArray - """ - _validate_format_channels(format, num_channels) - shape_t = _validate_array_shape(shape) - - cdef cydriver.CUarray_format c_format = format - cdef cydriver.CUDA_ARRAY3D_DESCRIPTOR desc3d - cdef int rank = len(shape_t) - cdef unsigned int flags = ( - cydriver.CUDA_ARRAY3D_SURFACE_LDST if is_surface_load_store else 0 - ) - - # cuArray3DCreate handles 1D/2D/3D uniformly (Height/Depth 0 sentinels), - # so a single descriptor + create_array_handle covers every shape. - memset(&desc3d, 0, sizeof(desc3d)) - desc3d.Width = shape_t[0] - desc3d.Height = (shape_t[1] if rank >= 2 else 0) - desc3d.Depth = (shape_t[2] if rank >= 3 else 0) - desc3d.Format = c_format - desc3d.NumChannels = num_channels - desc3d.Flags = flags - - cdef OpaqueArrayHandle h = create_array_handle(desc3d) - if not h: - HANDLE_RETURN(get_last_error()) - - cdef OpaqueArray self = cls.__new__(cls) - self._handle = h - self._shape = shape_t - self._format = c_format - self._num_channels = num_channels - self._surface_load_store = bool(is_surface_load_store) - self._device_id = _get_current_device_id() - return self - @classmethod def _from_handle(cls, intptr_t handle, bint owning, *, device_id=None): """Wrap an externally-allocated ``CUarray``. @@ -340,8 +376,8 @@ cdef class OpaqueArray: @property def format(self): - """The element :class:`ArrayFormat`.""" - return ArrayFormat(self._format) + """The element :class:`~cuda.core.typing.ArrayFormatType`.""" + return _CU_TO_ARRAYFORMAT[self._format] @property def num_channels(self): @@ -349,7 +385,7 @@ cdef class OpaqueArray: return self._num_channels @property - def element_size(self): + def element_bytes(self): """Bytes per element (format size * channels).""" return _FORMAT_ELEM_SIZE[self._format] * self._num_channels @@ -411,7 +447,7 @@ cdef class OpaqueArray: @property def size_bytes(self): - """Total bytes of array storage (``prod(shape) * element_size``).""" + """Total bytes of array storage (``prod(shape) * element_bytes``).""" cdef size_t n = 1 for s in self._shape: n *= s @@ -436,7 +472,7 @@ cdef class OpaqueArray: def __repr__(self): return ( f"OpaqueArray(shape={self._shape}, " - f"format={ArrayFormat(self._format).name}, " + f"format={_CU_TO_ARRAYFORMAT[self._format].name}, " f"num_channels={self._num_channels})" ) @@ -470,3 +506,47 @@ cdef OpaqueArray _array_from_handle(OpaqueArrayHandle h, int device_id): self._num_channels = desc.NumChannels self._surface_load_store = bool(desc.Flags & cydriver.CUDA_ARRAY3D_SURFACE_LDST) return self + + +def _create_opaque_array(options): + """Allocate a new :class:`OpaqueArray` on the current device. + + Backs :meth:`cuda.core.Device.create_opaque_array`. ``options`` is an + :class:`OpaqueArrayOptions` (or a mapping accepted by it); it is validated + at construction, so ``shape`` is already a normalized tuple and ``format`` + an :class:`~cuda.core.typing.ArrayFormatType`. + """ + cdef object opts = check_or_create_options( + OpaqueArrayOptions, options, "Opaque array options" + ) + shape_t = opts.shape + + cdef cydriver.CUarray_format c_format = _ARRAYFORMAT_TO_CU[opts.format] + cdef cydriver.CUDA_ARRAY3D_DESCRIPTOR desc3d + cdef int rank = len(shape_t) + cdef unsigned int flags = ( + cydriver.CUDA_ARRAY3D_SURFACE_LDST if opts.is_surface_load_store else 0 + ) + + # cuArray3DCreate handles 1D/2D/3D uniformly (Height/Depth 0 sentinels), + # so a single descriptor + create_array_handle covers every shape. + memset(&desc3d, 0, sizeof(desc3d)) + desc3d.Width = shape_t[0] + desc3d.Height = (shape_t[1] if rank >= 2 else 0) + desc3d.Depth = (shape_t[2] if rank >= 3 else 0) + desc3d.Format = c_format + desc3d.NumChannels = opts.num_channels + desc3d.Flags = flags + + cdef OpaqueArrayHandle h = create_array_handle(desc3d) + if not h: + HANDLE_RETURN(get_last_error()) + + cdef OpaqueArray self = OpaqueArray.__new__(OpaqueArray) + self._handle = h + self._shape = shape_t + self._format = c_format + self._num_channels = opts.num_channels + self._surface_load_store = bool(opts.is_surface_load_store) + self._device_id = _get_current_device_id() + return self diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index 0741df1781e..db4413dbaf4 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -2,6 +2,43 @@ from __future__ import annotations +from dataclasses import dataclass + + +@dataclass +class MipmappedArrayOptions: + """Options for :meth:`cuda.core.Device.create_mipmapped_array`. + + Attributes + ---------- + shape : tuple of int + ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in + elements, for the base (level 0) mip. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. + num_channels : int + Channels per element. Must be 1, 2, or 4. + num_levels : int + Number of mip levels to allocate; must be >= 1. The driver caps this at + the log2 of the largest dimension; passing a larger value yields a + driver error. + is_surface_load_store : bool + If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so individual + levels (obtained via :meth:`MipmappedArray.get_level`) can be bound as + a :class:`~cuda.core.texture.SurfaceObject` for kernel-side writes. + Default False. + + .. versionadded:: 1.1.0 + """ + shape: tuple[int, ...] + format: object + num_channels: int + num_levels: int + is_surface_load_store: bool = False + + def __post_init__(self): + ... class MipmappedArray: """A mipmapped CUDA array for texture/surface access across levels. @@ -13,7 +50,9 @@ class MipmappedArray: implicitly, so the :class:`OpaqueArray` instances returned by :meth:`get_level` are non-owning and hold a strong reference back to their parent. - Construct via :meth:`from_descriptor`. + Construct via :meth:`cuda.core.Device.create_mipmapped_array`. + + .. versionadded:: 1.1.0 """ def close(self): @@ -28,33 +67,6 @@ class MipmappedArray: def __init__(self, *args, **kwargs): ... - @classmethod - def from_descriptor(cls, *, shape, format, num_channels, num_levels, is_surface_load_store=False): - """Allocate a new mipmapped CUDA array. - - Parameters - ---------- - shape : tuple of int - ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` - in elements, for the base (level 0) mip. - format : ArrayFormat - Element format. - num_channels : int - Channels per element. Must be 1, 2, or 4. - num_levels : int - Number of mip levels to allocate; must be >= 1. The driver caps - this at the log2 of the largest dimension; passing a larger value - yields a driver error. - is_surface_load_store : bool - If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so individual - levels (obtained via :meth:`get_level`) can be bound as - :class:`SurfaceObject` for kernel-side writes. Default False. - - Returns - ------- - MipmappedArray - """ - def get_level(self, level): """Return a non-owning :class:`OpaqueArray` view of the given mip level. @@ -82,7 +94,7 @@ class MipmappedArray: @property def format(self): - """The element :class:`ArrayFormat`.""" + """The element :class:`~cuda.core.typing.ArrayFormatType`.""" @property def num_channels(self): @@ -108,4 +120,12 @@ class MipmappedArray: ... def __repr__(self): - ... \ No newline at end of file + ... + +def _create_mipmapped_array(options): + """Allocate a new :class:`MipmappedArray` on the current device. + + Backs :meth:`cuda.core.Device.create_mipmapped_array`. ``options`` is a + :class:`MipmappedArrayOptions` (or a mapping accepted by it); its fields are + validated at construction. + """ \ No newline at end of file diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyx b/cuda_core/cuda/core/texture/_mipmapped_array.pyx index 0847203fe99..3f151f7bb9f 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyx +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyx @@ -8,7 +8,12 @@ from libc.string cimport memset from cuda.bindings cimport cydriver from cuda.core.texture._array cimport _array_from_handle -from cuda.core.texture._array import ArrayFormat, _validate_array_shape, _validate_format_channels +from cuda.core.texture._array import ( + _ARRAYFORMAT_TO_CU, + _CU_TO_ARRAYFORMAT, + _validate_array_shape, + _validate_format_channels, +) from cuda.core._resource_handles cimport ( OpaqueArrayHandle, MipmappedArrayHandle, @@ -22,6 +27,51 @@ from cuda.core._utils.cuda_utils cimport ( _get_current_device_id, ) +from dataclasses import dataclass + +from cuda.core._utils.cuda_utils import check_or_create_options + + +@dataclass +class MipmappedArrayOptions: + """Options for :meth:`cuda.core.Device.create_mipmapped_array`. + + Attributes + ---------- + shape : tuple of int + ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in + elements, for the base (level 0) mip. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. + num_channels : int + Channels per element. Must be 1, 2, or 4. + num_levels : int + Number of mip levels to allocate; must be >= 1. The driver caps this at + the log2 of the largest dimension; passing a larger value yields a + driver error. + is_surface_load_store : bool + If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so individual + levels (obtained via :meth:`MipmappedArray.get_level`) can be bound as + a :class:`~cuda.core.texture.SurfaceObject` for kernel-side writes. + Default False. + + .. versionadded:: 1.1.0 + """ + + shape: tuple[int, ...] + format: object + num_channels: int + num_levels: int + is_surface_load_store: bool = False + + def __post_init__(self): + self.format = _validate_format_channels(self.format, self.num_channels) + self.shape = _validate_array_shape(self.shape) + self.num_levels = int(self.num_levels) + if self.num_levels < 1: + raise ValueError(f"num_levels must be >= 1, got {self.num_levels}") + cdef class MipmappedArray: """A mipmapped CUDA array for texture/surface access across levels. @@ -33,81 +83,16 @@ cdef class MipmappedArray: implicitly, so the :class:`OpaqueArray` instances returned by :meth:`get_level` are non-owning and hold a strong reference back to their parent. - Construct via :meth:`from_descriptor`. + Construct via :meth:`cuda.core.Device.create_mipmapped_array`. + + .. versionadded:: 1.1.0 """ def __init__(self, *args, **kwargs): raise RuntimeError( "MipmappedArray cannot be instantiated directly. " - "Use MipmappedArray.from_descriptor()." - ) - - @classmethod - def from_descriptor( - cls, *, shape, format, num_channels, num_levels, is_surface_load_store=False - ): - """Allocate a new mipmapped CUDA array. - - Parameters - ---------- - shape : tuple of int - ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` - in elements, for the base (level 0) mip. - format : ArrayFormat - Element format. - num_channels : int - Channels per element. Must be 1, 2, or 4. - num_levels : int - Number of mip levels to allocate; must be >= 1. The driver caps - this at the log2 of the largest dimension; passing a larger value - yields a driver error. - is_surface_load_store : bool - If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST`` so individual - levels (obtained via :meth:`get_level`) can be bound as - :class:`SurfaceObject` for kernel-side writes. Default False. - - Returns - ------- - MipmappedArray - """ - _validate_format_channels(format, num_channels) - shape_t = _validate_array_shape(shape) - - levels = int(num_levels) - if levels < 1: - raise ValueError(f"num_levels must be >= 1, got {levels}") - - cdef cydriver.CUarray_format c_format = format - cdef cydriver.CUDA_ARRAY3D_DESCRIPTOR desc3d - cdef int rank = len(shape_t) - cdef unsigned int flags = ( - cydriver.CUDA_ARRAY3D_SURFACE_LDST if is_surface_load_store else 0 + "Use Device.create_mipmapped_array()." ) - cdef unsigned int c_levels = levels - - # Mipmap creation uses the 3D descriptor regardless of rank; lower-rank - # shapes use Height=0/Depth=0 sentinels, matching cuArray3DCreate. - memset(&desc3d, 0, sizeof(desc3d)) - desc3d.Width = shape_t[0] - desc3d.Height = (shape_t[1] if rank >= 2 else 0) - desc3d.Depth = (shape_t[2] if rank >= 3 else 0) - desc3d.Format = c_format - desc3d.NumChannels = num_channels - desc3d.Flags = flags - - cdef MipmappedArrayHandle h = create_mipmapped_array_handle(desc3d, c_levels) - if not h: - HANDLE_RETURN(get_last_error()) - - cdef MipmappedArray self = cls.__new__(cls) - self._handle = h - self._shape = shape_t - self._format = c_format - self._num_channels = num_channels - self._num_levels = levels - self._surface_load_store = bool(is_surface_load_store) - self._device_id = _get_current_device_id() - return self def get_level(self, level): """Return a non-owning :class:`OpaqueArray` view of the given mip level. @@ -153,8 +138,8 @@ cdef class MipmappedArray: @property def format(self): - """The element :class:`ArrayFormat`.""" - return ArrayFormat(self._format) + """The element :class:`~cuda.core.typing.ArrayFormatType`.""" + return _CU_TO_ARRAYFORMAT[self._format] @property def num_channels(self): @@ -197,7 +182,52 @@ cdef class MipmappedArray: def __repr__(self): return ( f"MipmappedArray(shape={self._shape}, " - f"format={ArrayFormat(self._format).name}, " + f"format={_CU_TO_ARRAYFORMAT[self._format].name}, " f"num_channels={self._num_channels}, " f"num_levels={self._num_levels})" ) + + +def _create_mipmapped_array(options): + """Allocate a new :class:`MipmappedArray` on the current device. + + Backs :meth:`cuda.core.Device.create_mipmapped_array`. ``options`` is a + :class:`MipmappedArrayOptions` (or a mapping accepted by it); its fields are + validated at construction. + """ + cdef object opts = check_or_create_options( + MipmappedArrayOptions, options, "Mipmapped array options" + ) + shape_t = opts.shape + + cdef cydriver.CUarray_format c_format = _ARRAYFORMAT_TO_CU[opts.format] + cdef cydriver.CUDA_ARRAY3D_DESCRIPTOR desc3d + cdef int rank = len(shape_t) + cdef unsigned int flags = ( + cydriver.CUDA_ARRAY3D_SURFACE_LDST if opts.is_surface_load_store else 0 + ) + cdef unsigned int c_levels = opts.num_levels + + # Mipmap creation uses the 3D descriptor regardless of rank; lower-rank + # shapes use Height=0/Depth=0 sentinels, matching cuArray3DCreate. + memset(&desc3d, 0, sizeof(desc3d)) + desc3d.Width = shape_t[0] + desc3d.Height = (shape_t[1] if rank >= 2 else 0) + desc3d.Depth = (shape_t[2] if rank >= 3 else 0) + desc3d.Format = c_format + desc3d.NumChannels = opts.num_channels + desc3d.Flags = flags + + cdef MipmappedArrayHandle h = create_mipmapped_array_handle(desc3d, c_levels) + if not h: + HANDLE_RETURN(get_last_error()) + + cdef MipmappedArray self = MipmappedArray.__new__(MipmappedArray) + self._handle = h + self._shape = shape_t + self._format = c_format + self._num_channels = opts.num_channels + self._num_levels = opts.num_levels + self._surface_load_store = bool(opts.is_surface_load_store) + self._device_id = _get_current_device_id() + return self diff --git a/cuda_core/cuda/core/texture/_surface.pyi b/cuda_core/cuda/core/texture/_surface.pyi index be61e1099ba..977268abd5f 100644 --- a/cuda_core/cuda/core/texture/_surface.pyi +++ b/cuda_core/cuda/core/texture/_surface.pyi @@ -14,8 +14,10 @@ class SurfaceObject: ``is_surface_load_store=True`` and is kept alive for the lifetime of this object to prevent dangling handles. - Construct via :meth:`from_array` or :meth:`from_descriptor`. Passes to + Construct via :meth:`cuda.core.Device.create_surface_object`. Passes to kernels as a 64-bit handle (via the ``handle`` property). + + .. versionadded:: 1.1.0 """ def close(self): @@ -29,25 +31,6 @@ class SurfaceObject: def __init__(self, *args, **kwargs): ... - @classmethod - def from_array(cls, array): - """Create a surface object directly from an :class:`OpaqueArray`. - - The array must have been created with ``is_surface_load_store=True``. - """ - - @classmethod - def from_descriptor(cls, *, resource): - """Create a surface object from a :class:`ResourceDescriptor`. - - Parameters - ---------- - resource : ResourceDescriptor - Must wrap an :class:`OpaqueArray` allocated with - ``is_surface_load_store=True``. Linear/pitch2d resources are not - valid surface backings. - """ - @property def handle(self): """The underlying ``CUsurfObject`` as an integer (64-bit kernel arg).""" @@ -67,4 +50,13 @@ class SurfaceObject: ... def __repr__(self): - ... \ No newline at end of file + ... + +def _create_surface_object(resource): + """Create a :class:`SurfaceObject` on the current device. + + Backs :meth:`cuda.core.Device.create_surface_object`. ``resource`` must be a + :class:`ResourceDescriptor` wrapping an :class:`OpaqueArray` allocated with + ``is_surface_load_store=True``; linear/pitch2d resources are not valid + surface backings. + """ \ No newline at end of file diff --git a/cuda_core/cuda/core/texture/_surface.pyx b/cuda_core/cuda/core/texture/_surface.pyx index 93836dfeab8..ac61fddd357 100644 --- a/cuda_core/cuda/core/texture/_surface.pyx +++ b/cuda_core/cuda/core/texture/_surface.pyx @@ -33,70 +33,18 @@ cdef class SurfaceObject: ``is_surface_load_store=True`` and is kept alive for the lifetime of this object to prevent dangling handles. - Construct via :meth:`from_array` or :meth:`from_descriptor`. Passes to + Construct via :meth:`cuda.core.Device.create_surface_object`. Passes to kernels as a 64-bit handle (via the ``handle`` property). + + .. versionadded:: 1.1.0 """ def __init__(self, *args, **kwargs): raise RuntimeError( "SurfaceObject cannot be instantiated directly. " - "Use SurfaceObject.from_array() or SurfaceObject.from_descriptor()." + "Use Device.create_surface_object()." ) - @classmethod - def from_array(cls, array): - """Create a surface object directly from an :class:`OpaqueArray`. - - The array must have been created with ``is_surface_load_store=True``. - """ - if not isinstance(array, OpaqueArray): - raise TypeError(f"array must be a OpaqueArray, got {type(array).__name__}") - return cls.from_descriptor(resource=ResourceDescriptor.from_array(array)) - - @classmethod - def from_descriptor(cls, *, resource): - """Create a surface object from a :class:`ResourceDescriptor`. - - Parameters - ---------- - resource : ResourceDescriptor - Must wrap an :class:`OpaqueArray` allocated with - ``is_surface_load_store=True``. Linear/pitch2d resources are not - valid surface backings. - """ - if not isinstance(resource, ResourceDescriptor): - raise TypeError( - f"resource must be a ResourceDescriptor, got " - f"{type(resource).__name__}" - ) - if resource.kind != "array": - raise ValueError( - f"SurfaceObject requires an array-backed ResourceDescriptor, " - f"got kind={resource.kind!r}" - ) - - cdef OpaqueArray arr = resource.source - if not arr.is_surface_load_store: - raise ValueError( - "OpaqueArray must be created with is_surface_load_store=True to be " - "bound as a SurfaceObject" - ) - - cdef cydriver.CUDA_RESOURCE_DESC res_desc - memset(&res_desc, 0, sizeof(res_desc)) - res_desc.resType = cydriver.CU_RESOURCE_TYPE_ARRAY - res_desc.res.array.hArray = as_cu(arr._handle) - - cdef SurfObjectHandle h = create_surf_object_handle(res_desc, arr._handle) - if not h: - HANDLE_RETURN(get_last_error()) - - cdef SurfaceObject self = cls.__new__(cls) - self._handle = h - self._source_ref = resource - self._device_id = _get_current_device_id() - return self - @property def handle(self): """The underlying ``CUsurfObject`` as an integer (64-bit kernel arg).""" @@ -130,3 +78,45 @@ cdef class SurfaceObject: def __repr__(self): return f"SurfaceObject(handle=0x{as_intptr(self._handle):x})" + + +def _create_surface_object(resource): + """Create a :class:`SurfaceObject` on the current device. + + Backs :meth:`cuda.core.Device.create_surface_object`. ``resource`` must be a + :class:`ResourceDescriptor` wrapping an :class:`OpaqueArray` allocated with + ``is_surface_load_store=True``; linear/pitch2d resources are not valid + surface backings. + """ + if not isinstance(resource, ResourceDescriptor): + raise TypeError( + f"resource must be a ResourceDescriptor, got " + f"{type(resource).__name__}" + ) + if resource.kind != "array": + raise ValueError( + f"SurfaceObject requires an array-backed ResourceDescriptor, " + f"got kind={resource.kind!r}" + ) + + cdef OpaqueArray arr = resource.source + if not arr.is_surface_load_store: + raise ValueError( + "OpaqueArray must be created with is_surface_load_store=True to be " + "bound as a SurfaceObject" + ) + + cdef cydriver.CUDA_RESOURCE_DESC res_desc + memset(&res_desc, 0, sizeof(res_desc)) + res_desc.resType = cydriver.CU_RESOURCE_TYPE_ARRAY + res_desc.res.array.hArray = as_cu(arr._handle) + + cdef SurfObjectHandle h = create_surf_object_handle(res_desc, arr._handle) + if not h: + HANDLE_RETURN(get_last_error()) + + cdef SurfaceObject self = SurfaceObject.__new__(SurfaceObject) + self._handle = h + self._source_ref = resource + self._device_id = _get_current_device_id() + return self diff --git a/cuda_core/cuda/core/texture/_texture.pxd b/cuda_core/cuda/core/texture/_texture.pxd index f8532a4dad5..6d0871fe848 100644 --- a/cuda_core/cuda/core/texture/_texture.pxd +++ b/cuda_core/cuda/core/texture/_texture.pxd @@ -13,7 +13,7 @@ cdef class TextureObject: # structurally by the C++ box behind this handle, not by _source_ref. TexObjectHandle _handle object _source_ref # ResourceDescriptor, retained for introspection - object _texture_desc # original TextureDescriptor for introspection + object _options # original TextureObjectOptions for introspection int _device_id cpdef close(self) diff --git a/cuda_core/cuda/core/texture/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index b04ca384989..7840585bb4d 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -3,41 +3,17 @@ from __future__ import annotations from dataclasses import dataclass -from enum import IntEnum from cuda.bindings import cydriver +from cuda.core.typing import AddressModeType, FilterModeType, ReadModeType -class AddressMode(IntEnum): - """Boundary behavior for out-of-range texture coordinates.""" - WRAP = cydriver.CU_TR_ADDRESS_MODE_WRAP - CLAMP = cydriver.CU_TR_ADDRESS_MODE_CLAMP - MIRROR = cydriver.CU_TR_ADDRESS_MODE_MIRROR - BORDER = cydriver.CU_TR_ADDRESS_MODE_BORDER - -class FilterMode(IntEnum): - """Texel sampling mode.""" - POINT = cydriver.CU_TR_FILTER_MODE_POINT - LINEAR = cydriver.CU_TR_FILTER_MODE_LINEAR - -class ReadMode(IntEnum): - """How sampled values are returned to the kernel. - - - ``ELEMENT_TYPE``: return the raw element value (integer formats stay - integer, float stays float). - - ``NORMALIZED_FLOAT``: integer formats are promoted to a normalized - ``float`` in ``[0, 1]`` (unsigned) or ``[-1, 1]`` (signed). - Float formats are unaffected. - """ - ELEMENT_TYPE = 0 - NORMALIZED_FLOAT = 1 - class ResourceDescriptor: """Describes the memory backing a :class:`TextureObject`. Construct via the ``from_*`` classmethods: - - :meth:`from_array` wraps a :class:`OpaqueArray` (works for both + - :meth:`from_opaque_array` wraps a :class:`OpaqueArray` (works for both :class:`TextureObject` and :class:`SurfaceObject`). - :meth:`from_mipmapped_array` wraps a :class:`MipmappedArray` for mipmapped sampling (texture only, not surface). @@ -49,6 +25,8 @@ class ResourceDescriptor: Linear and pitch2D resources cannot back a :class:`SurfaceObject` — those require an :class:`OpaqueArray` allocated with ``is_surface_load_store=True``. + + .. versionadded:: 1.1.0 """ __slots__ = ('_kind', '_source', '_format', '_num_channels', '_size_bytes', '_width', '_height', '_pitch_bytes') @@ -56,7 +34,7 @@ class ResourceDescriptor: ... @classmethod - def from_array(cls, array): + def from_opaque_array(cls, array): """Build a resource descriptor backed by a :class:`OpaqueArray`.""" @classmethod @@ -78,8 +56,9 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat - Element format. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. num_channels : int Channels per element. Must be 1, 2, or 4. size_bytes : int, optional @@ -89,7 +68,7 @@ class ResourceDescriptor: Notes ----- Texture objects built from a linear resource ignore the - :class:`TextureDescriptor` addressing/filtering fields — kernels read + :class:`TextureObjectOptions` addressing/filtering fields — kernels read through a typed 1D fetch with bounds checking only. """ @@ -102,8 +81,9 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat - Element format. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. num_channels : int Channels per element. Must be 1, 2, or 4. width : int @@ -126,7 +106,7 @@ class ResourceDescriptor: @property def format(self): - """The element :class:`ArrayFormat` (``None`` for array-backed).""" + """The element :class:`~cuda.core.typing.ArrayFormatType` (``None`` for array-backed).""" @property def num_channels(self): @@ -152,18 +132,20 @@ class ResourceDescriptor: ... @dataclass -class TextureDescriptor: +class TextureObjectOptions: """Sampling state for a :class:`TextureObject` (mirrors ``CUDA_TEXTURE_DESC``). Attributes ---------- - address_mode : tuple of AddressMode - Boundary behavior per axis. May be a single :class:`AddressMode` (applied - to all axes) or a tuple of 1-3 entries (one per dimension). - filter_mode : FilterMode - Texel sampling mode. Default ``POINT``. - read_mode : ReadMode + address_mode : AddressModeType or tuple of AddressModeType + Boundary behavior per axis. May be a single + :class:`~cuda.core.typing.AddressModeType` (applied to all axes) or a + tuple of 1-3 entries (one per dimension). Plain strings are accepted. + filter_mode : FilterModeType + Texel sampling mode. Default ``POINT``. Plain strings are accepted. + read_mode : ReadModeType How sampled integer values are returned. Default ``ELEMENT_TYPE``. + Plain strings are accepted. normalized_coords : bool If True, coordinates are in ``[0, 1]`` instead of pixel indices. srgb : bool @@ -174,29 +156,35 @@ class TextureDescriptor: If True, enable seamless cubemap edge filtering. max_anisotropy : int Maximum anisotropy; 0 disables anisotropic filtering. - mipmap_filter_mode : FilterMode - Filtering between mipmap levels. Default ``POINT``. + mipmap_filter_mode : FilterModeType + Filtering between mipmap levels. Default ``POINT``. Plain strings are + accepted. mipmap_level_bias : float min_mipmap_level_clamp : float max_mipmap_level_clamp : float border_color : tuple of float or None 4-tuple used when ``address_mode`` includes ``BORDER``; ``None`` means zero. + + .. versionadded:: 1.1.0 """ - address_mode: AddressMode | tuple[AddressMode, ...] = AddressMode.CLAMP - filter_mode: FilterMode = FilterMode.POINT - read_mode: ReadMode = ReadMode.ELEMENT_TYPE + address_mode: AddressModeType | str | tuple[AddressModeType | str, ...] = AddressModeType.CLAMP + filter_mode: FilterModeType | str = FilterModeType.POINT + read_mode: ReadModeType | str = ReadModeType.ELEMENT_TYPE normalized_coords: bool = False srgb: bool = False disable_trilinear_optimization: bool = False seamless_cubemap: bool = False max_anisotropy: int = 0 - mipmap_filter_mode: FilterMode = FilterMode.POINT + mipmap_filter_mode: FilterModeType | str = FilterModeType.POINT mipmap_level_bias: float = 0.0 min_mipmap_level_clamp: float = 0.0 max_mipmap_level_clamp: float = 0.0 border_color: tuple[float, ...] | None = None + def __post_init__(self): + ... + class TextureObject: """A bindless texture handle for kernel-side sampled reads. @@ -204,8 +192,10 @@ class TextureObject: :class:`OpaqueArray` referenced by the descriptor) is kept alive for the lifetime of this object to prevent dangling handles. - Construct via :meth:`from_descriptor`. Passes to kernels as a 64-bit - handle (via the ``handle`` property). + Construct via :meth:`cuda.core.Device.create_texture_object`. Passes to + kernels as a 64-bit handle (via the ``handle`` property). + + .. versionadded:: 1.1.0 """ def close(self): @@ -219,16 +209,6 @@ class TextureObject: def __init__(self, *args, **kwargs): ... - @classmethod - def from_descriptor(cls, *, resource, texture_descriptor): - """Create a texture object from a resource + sampling descriptor. - - Parameters - ---------- - resource : ResourceDescriptor - texture_descriptor : TextureDescriptor - """ - @property def handle(self): """The underlying ``CUtexObject`` as an integer (64-bit kernel arg).""" @@ -238,8 +218,8 @@ class TextureObject: """The :class:`ResourceDescriptor` this texture was built from.""" @property - def texture_descriptor(self): - """The :class:`TextureDescriptor` this texture was built from.""" + def options(self): + """The :class:`TextureObjectOptions` this texture was built from.""" @property def device(self): @@ -258,6 +238,20 @@ _TRSF_NORMALIZED_COORDINATES = 2 _TRSF_SRGB = 16 _TRSF_DISABLE_TRILINEAR_OPTIMIZATION = 32 _TRSF_SEAMLESS_CUBEMAP = 64 +_ADDRESSMODE_TO_CU = {AddressModeType.WRAP: int(cydriver.CU_TR_ADDRESS_MODE_WRAP), AddressModeType.CLAMP: int(cydriver.CU_TR_ADDRESS_MODE_CLAMP), AddressModeType.MIRROR: int(cydriver.CU_TR_ADDRESS_MODE_MIRROR), AddressModeType.BORDER: int(cydriver.CU_TR_ADDRESS_MODE_BORDER)} +_FILTERMODE_TO_CU = {FilterModeType.POINT: int(cydriver.CU_TR_FILTER_MODE_POINT), FilterModeType.LINEAR: int(cydriver.CU_TR_FILTER_MODE_LINEAR)} + +def _normalize_enum(name, value, enum_type): + """Coerce ``value`` to ``enum_type`` (a StrEnum), accepting a plain str.""" def _normalize_address_modes(address_mode): - """Return a 3-tuple of AddressMode values from a scalar or 1-3 tuple.""" \ No newline at end of file + """Return a 3-tuple of :class:`AddressModeType` values from a scalar or + 1-3 tuple. Individual entries may be plain strings.""" + +def _create_texture_object(resource, options): + """Create a :class:`TextureObject` on the current device. + + Backs :meth:`cuda.core.Device.create_texture_object`. ``resource`` is a + :class:`ResourceDescriptor`; ``options`` is a :class:`TextureObjectOptions` + (or a mapping accepted by it). + """ \ No newline at end of file diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index e88ec2d2a36..8d09c727c31 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -9,7 +9,12 @@ from libc.string cimport memset from cuda.bindings cimport cydriver from cuda.core.texture._array cimport OpaqueArray -from cuda.core.texture._array import ArrayFormat, _FORMAT_ELEM_SIZE, _validate_format_channels +from cuda.core.texture._array import ( + _ARRAYFORMAT_TO_CU, + _CU_TO_ARRAYFORMAT, + _FORMAT_ELEM_SIZE, + _validate_format_channels, +) from cuda.core._memory._buffer cimport Buffer from cuda.core.texture._mipmapped_array cimport MipmappedArray from cuda.core.texture._mipmapped_array import MipmappedArray as _PyMipmappedArray @@ -27,8 +32,11 @@ from cuda.core._utils.cuda_utils cimport ( _get_current_device_id, ) +from cuda.core.typing import AddressModeType, FilterModeType, ReadModeType + from dataclasses import dataclass -from enum import IntEnum + +from cuda.core._utils.cuda_utils import check_or_create_options # Driver texture-descriptor flag bits (CU_TRSF_*). @@ -39,31 +47,30 @@ _TRSF_DISABLE_TRILINEAR_OPTIMIZATION = 0x20 _TRSF_SEAMLESS_CUBEMAP = 0x40 -class AddressMode(IntEnum): - """Boundary behavior for out-of-range texture coordinates.""" - WRAP = cydriver.CU_TR_ADDRESS_MODE_WRAP - CLAMP = cydriver.CU_TR_ADDRESS_MODE_CLAMP - MIRROR = cydriver.CU_TR_ADDRESS_MODE_MIRROR - BORDER = cydriver.CU_TR_ADDRESS_MODE_BORDER - +# Bridge between the public sampling StrEnums and the driver integer values. +_ADDRESSMODE_TO_CU = { + AddressModeType.WRAP: int(cydriver.CU_TR_ADDRESS_MODE_WRAP), + AddressModeType.CLAMP: int(cydriver.CU_TR_ADDRESS_MODE_CLAMP), + AddressModeType.MIRROR: int(cydriver.CU_TR_ADDRESS_MODE_MIRROR), + AddressModeType.BORDER: int(cydriver.CU_TR_ADDRESS_MODE_BORDER), +} +_FILTERMODE_TO_CU = { + FilterModeType.POINT: int(cydriver.CU_TR_FILTER_MODE_POINT), + FilterModeType.LINEAR: int(cydriver.CU_TR_FILTER_MODE_LINEAR), +} -class FilterMode(IntEnum): - """Texel sampling mode.""" - POINT = cydriver.CU_TR_FILTER_MODE_POINT - LINEAR = cydriver.CU_TR_FILTER_MODE_LINEAR - -class ReadMode(IntEnum): - """How sampled values are returned to the kernel. - - - ``ELEMENT_TYPE``: return the raw element value (integer formats stay - integer, float stays float). - - ``NORMALIZED_FLOAT``: integer formats are promoted to a normalized - ``float`` in ``[0, 1]`` (unsigned) or ``[-1, 1]`` (signed). - Float formats are unaffected. - """ - ELEMENT_TYPE = 0 - NORMALIZED_FLOAT = 1 +def _normalize_enum(name, value, enum_type): + """Coerce ``value`` to ``enum_type`` (a StrEnum), accepting a plain str.""" + if isinstance(value, enum_type): + return value + try: + return enum_type(value) + except ValueError as e: + valid = ", ".join(repr(m.value) for m in enum_type) + raise ValueError( + f"{name} must be a {enum_type.__name__} or one of {{{valid}}}, got {value!r}" + ) from e class ResourceDescriptor: @@ -71,7 +78,7 @@ class ResourceDescriptor: Construct via the ``from_*`` classmethods: - - :meth:`from_array` wraps a :class:`OpaqueArray` (works for both + - :meth:`from_opaque_array` wraps a :class:`OpaqueArray` (works for both :class:`TextureObject` and :class:`SurfaceObject`). - :meth:`from_mipmapped_array` wraps a :class:`MipmappedArray` for mipmapped sampling (texture only, not surface). @@ -83,6 +90,8 @@ class ResourceDescriptor: Linear and pitch2D resources cannot back a :class:`SurfaceObject` — those require an :class:`OpaqueArray` allocated with ``is_surface_load_store=True``. + + .. versionadded:: 1.1.0 """ __slots__ = ( @@ -99,7 +108,7 @@ class ResourceDescriptor: ) @classmethod - def from_array(cls, array): + def from_opaque_array(cls, array): """Build a resource descriptor backed by a :class:`OpaqueArray`.""" if not isinstance(array, OpaqueArray): raise TypeError(f"array must be a OpaqueArray, got {type(array).__name__}") @@ -148,8 +157,9 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat - Element format. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. num_channels : int Channels per element. Must be 1, 2, or 4. size_bytes : int, optional @@ -159,15 +169,16 @@ class ResourceDescriptor: Notes ----- Texture objects built from a linear resource ignore the - :class:`TextureDescriptor` addressing/filtering fields — kernels read + :class:`TextureObjectOptions` addressing/filtering fields — kernels read through a typed 1D fetch with bounds checking only. """ if not isinstance(buffer, Buffer): raise TypeError(f"buffer must be a Buffer, got {type(buffer).__name__}") - _validate_format_channels(format, num_channels) + fmt = _validate_format_channels(format, num_channels) + cu_format = _ARRAYFORMAT_TO_CU[fmt] buf_size = int(buffer.size) - elem = _FORMAT_ELEM_SIZE[int(format)] * int(num_channels) + elem = _FORMAT_ELEM_SIZE[cu_format] * int(num_channels) if size_bytes is None: size = buf_size else: @@ -183,13 +194,13 @@ class ResourceDescriptor: if size % elem != 0: raise ValueError( f"size_bytes ({size}) must be a multiple of element size " - f"({elem} bytes for {format.name} x {num_channels})" + f"({elem} bytes for {fmt.name} x {num_channels})" ) self = cls.__new__(cls) self._kind = "linear" self._source = buffer - self._format = int(format) + self._format = cu_format self._num_channels = int(num_channels) self._size_bytes = size self._width = None @@ -208,8 +219,9 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat - Element format. + format : ArrayFormatType, str, or numpy.dtype + Element format. Accepts an :class:`~cuda.core.typing.ArrayFormatType`, + a plain string (e.g. ``"float32"``), or a NumPy dtype object. num_channels : int Channels per element. Must be 1, 2, or 4. width : int @@ -223,7 +235,8 @@ class ResourceDescriptor: """ if not isinstance(buffer, Buffer): raise TypeError(f"buffer must be a Buffer, got {type(buffer).__name__}") - _validate_format_channels(format, num_channels) + fmt = _validate_format_channels(format, num_channels) + cu_format = _ARRAYFORMAT_TO_CU[fmt] w = int(width) h = int(height) @@ -232,11 +245,11 @@ class ResourceDescriptor: raise ValueError(f"width must be >= 1, got {w}") if h < 1: raise ValueError(f"height must be >= 1, got {h}") - elem = _FORMAT_ELEM_SIZE[int(format)] * int(num_channels) + elem = _FORMAT_ELEM_SIZE[cu_format] * int(num_channels) min_pitch = w * elem if p < min_pitch: raise ValueError( - f"pitch_bytes ({p}) must be >= width * element_size ({min_pitch})" + f"pitch_bytes ({p}) must be >= width * element_bytes ({min_pitch})" ) if p * h > int(buffer.size): raise ValueError( @@ -246,7 +259,7 @@ class ResourceDescriptor: self = cls.__new__(cls) self._kind = "pitch2d" self._source = buffer - self._format = int(format) + self._format = cu_format self._num_channels = int(num_channels) self._size_bytes = None self._width = w @@ -264,8 +277,8 @@ class ResourceDescriptor: @property def format(self): - """The element :class:`ArrayFormat` (``None`` for array-backed).""" - return None if self._format is None else ArrayFormat(self._format) + """The element :class:`~cuda.core.typing.ArrayFormatType` (``None`` for array-backed).""" + return None if self._format is None else _CU_TO_ARRAYFORMAT[self._format] @property def num_channels(self): @@ -309,18 +322,20 @@ class ResourceDescriptor: @dataclass -class TextureDescriptor: +class TextureObjectOptions: """Sampling state for a :class:`TextureObject` (mirrors ``CUDA_TEXTURE_DESC``). Attributes ---------- - address_mode : tuple of AddressMode - Boundary behavior per axis. May be a single :class:`AddressMode` (applied - to all axes) or a tuple of 1-3 entries (one per dimension). - filter_mode : FilterMode - Texel sampling mode. Default ``POINT``. - read_mode : ReadMode + address_mode : AddressModeType or tuple of AddressModeType + Boundary behavior per axis. May be a single + :class:`~cuda.core.typing.AddressModeType` (applied to all axes) or a + tuple of 1-3 entries (one per dimension). Plain strings are accepted. + filter_mode : FilterModeType + Texel sampling mode. Default ``POINT``. Plain strings are accepted. + read_mode : ReadModeType How sampled integer values are returned. Default ``ELEMENT_TYPE``. + Plain strings are accepted. normalized_coords : bool If True, coordinates are in ``[0, 1]`` instead of pixel indices. srgb : bool @@ -331,50 +346,61 @@ class TextureDescriptor: If True, enable seamless cubemap edge filtering. max_anisotropy : int Maximum anisotropy; 0 disables anisotropic filtering. - mipmap_filter_mode : FilterMode - Filtering between mipmap levels. Default ``POINT``. + mipmap_filter_mode : FilterModeType + Filtering between mipmap levels. Default ``POINT``. Plain strings are + accepted. mipmap_level_bias : float min_mipmap_level_clamp : float max_mipmap_level_clamp : float border_color : tuple of float or None 4-tuple used when ``address_mode`` includes ``BORDER``; ``None`` means zero. + + .. versionadded:: 1.1.0 """ - address_mode: AddressMode | tuple[AddressMode, ...] = AddressMode.CLAMP - filter_mode: FilterMode = FilterMode.POINT - read_mode: ReadMode = ReadMode.ELEMENT_TYPE + address_mode: AddressModeType | str | tuple[AddressModeType | str, ...] = AddressModeType.CLAMP + filter_mode: FilterModeType | str = FilterModeType.POINT + read_mode: ReadModeType | str = ReadModeType.ELEMENT_TYPE normalized_coords: bool = False srgb: bool = False disable_trilinear_optimization: bool = False seamless_cubemap: bool = False max_anisotropy: int = 0 - mipmap_filter_mode: FilterMode = FilterMode.POINT + mipmap_filter_mode: FilterModeType | str = FilterModeType.POINT mipmap_level_bias: float = 0.0 min_mipmap_level_clamp: float = 0.0 max_mipmap_level_clamp: float = 0.0 border_color: tuple[float, ...] | None = None + def __post_init__(self): + self.filter_mode = _normalize_enum("filter_mode", self.filter_mode, FilterModeType) + self.read_mode = _normalize_enum("read_mode", self.read_mode, ReadModeType) + self.mipmap_filter_mode = _normalize_enum( + "mipmap_filter_mode", self.mipmap_filter_mode, FilterModeType + ) + def _normalize_address_modes(address_mode): - """Return a 3-tuple of AddressMode values from a scalar or 1-3 tuple.""" - if isinstance(address_mode, AddressMode): - return (address_mode, address_mode, address_mode) + """Return a 3-tuple of :class:`AddressModeType` values from a scalar or + 1-3 tuple. Individual entries may be plain strings.""" + if isinstance(address_mode, (AddressModeType, str)): + m = _normalize_enum("address_mode", address_mode, AddressModeType) + return (m, m, m) try: modes = tuple(address_mode) except TypeError as e: raise TypeError( - "address_mode must be an AddressMode or a tuple of AddressMode" + "address_mode must be an AddressModeType or a tuple of AddressModeType" ) from e if not 1 <= len(modes) <= 3: raise ValueError( f"address_mode tuple must have 1-3 entries, got {len(modes)}" ) - for i, m in enumerate(modes): - if not isinstance(m, AddressMode): - raise TypeError( - f"address_mode[{i}] must be an AddressMode, got {type(m).__name__}" - ) + modes = tuple( + _normalize_enum(f"address_mode[{i}]", m, AddressModeType) + for i, m in enumerate(modes) + ) # Pad to 3 entries by repeating the last one. padded = list(modes) + [modes[-1]] * (3 - len(modes)) return tuple(padded) @@ -387,155 +413,18 @@ cdef class TextureObject: :class:`OpaqueArray` referenced by the descriptor) is kept alive for the lifetime of this object to prevent dangling handles. - Construct via :meth:`from_descriptor`. Passes to kernels as a 64-bit - handle (via the ``handle`` property). + Construct via :meth:`cuda.core.Device.create_texture_object`. Passes to + kernels as a 64-bit handle (via the ``handle`` property). + + .. versionadded:: 1.1.0 """ def __init__(self, *args, **kwargs): raise RuntimeError( "TextureObject cannot be instantiated directly. " - "Use TextureObject.from_descriptor()." + "Use Device.create_texture_object()." ) - @classmethod - def from_descriptor(cls, *, resource, texture_descriptor): - """Create a texture object from a resource + sampling descriptor. - - Parameters - ---------- - resource : ResourceDescriptor - texture_descriptor : TextureDescriptor - """ - if not isinstance(resource, ResourceDescriptor): - raise TypeError( - f"resource must be a ResourceDescriptor, got " - f"{type(resource).__name__}" - ) - if not isinstance(texture_descriptor, TextureDescriptor): - raise TypeError( - f"texture_descriptor must be a TextureDescriptor, got " - f"{type(texture_descriptor).__name__}" - ) - - cdef cydriver.CUDA_RESOURCE_DESC res_desc - cdef cydriver.CUDA_TEXTURE_DESC tex_desc - memset(&res_desc, 0, sizeof(res_desc)) - memset(&tex_desc, 0, sizeof(tex_desc)) - - # --- Resource descriptor --- - cdef OpaqueArray arr - cdef MipmappedArray mip - cdef Buffer buf - cdef intptr_t devptr - if resource.kind == "array": - arr = resource.source - res_desc.resType = cydriver.CU_RESOURCE_TYPE_ARRAY - res_desc.res.array.hArray = as_cu(arr._handle) - elif resource.kind == "mipmapped_array": - mip = resource.source - res_desc.resType = cydriver.CU_RESOURCE_TYPE_MIPMAPPED_ARRAY - res_desc.res.mipmap.hMipmappedArray = as_cu(mip._handle) - elif resource.kind == "linear": - buf = resource.source - devptr = int(buf.handle) - res_desc.resType = cydriver.CU_RESOURCE_TYPE_LINEAR - res_desc.res.linear.devPtr = devptr - res_desc.res.linear.format = resource._format - res_desc.res.linear.numChannels = resource._num_channels - res_desc.res.linear.sizeInBytes = resource._size_bytes - elif resource.kind == "pitch2d": - buf = resource.source - devptr = int(buf.handle) - res_desc.resType = cydriver.CU_RESOURCE_TYPE_PITCH2D - res_desc.res.pitch2D.devPtr = devptr - res_desc.res.pitch2D.format = resource._format - res_desc.res.pitch2D.numChannels = resource._num_channels - res_desc.res.pitch2D.width = resource._width - res_desc.res.pitch2D.height = resource._height - res_desc.res.pitch2D.pitchInBytes = resource._pitch_bytes - else: - raise NotImplementedError( - f"ResourceDescriptor kind {resource.kind!r} is not yet supported" - ) - - # --- Texture descriptor --- - modes = _normalize_address_modes(texture_descriptor.address_mode) - tex_desc.addressMode[0] = modes[0] - tex_desc.addressMode[1] = modes[1] - tex_desc.addressMode[2] = modes[2] - - if not isinstance(texture_descriptor.filter_mode, FilterMode): - raise TypeError( - f"filter_mode must be a FilterMode, got " - f"{type(texture_descriptor.filter_mode).__name__}" - ) - tex_desc.filterMode = texture_descriptor.filter_mode - - if not isinstance(texture_descriptor.read_mode, ReadMode): - raise TypeError( - f"read_mode must be a ReadMode, got " - f"{type(texture_descriptor.read_mode).__name__}" - ) - - cdef unsigned int flags = 0 - # CU_TRSF_READ_AS_INTEGER suppresses normalization, so it maps to - # ReadMode.ELEMENT_TYPE. - if texture_descriptor.read_mode == ReadMode.ELEMENT_TYPE: - flags |= _TRSF_READ_AS_INTEGER - if texture_descriptor.normalized_coords: - flags |= _TRSF_NORMALIZED_COORDINATES - if texture_descriptor.srgb: - flags |= _TRSF_SRGB - if texture_descriptor.disable_trilinear_optimization: - flags |= _TRSF_DISABLE_TRILINEAR_OPTIMIZATION - if texture_descriptor.seamless_cubemap: - flags |= _TRSF_SEAMLESS_CUBEMAP - tex_desc.flags = flags - - if texture_descriptor.max_anisotropy < 0: - raise ValueError("max_anisotropy must be >= 0") - tex_desc.maxAnisotropy = texture_descriptor.max_anisotropy - - if not isinstance(texture_descriptor.mipmap_filter_mode, FilterMode): - raise TypeError( - f"mipmap_filter_mode must be a FilterMode, got " - f"{type(texture_descriptor.mipmap_filter_mode).__name__}" - ) - tex_desc.mipmapFilterMode = texture_descriptor.mipmap_filter_mode - tex_desc.mipmapLevelBias = texture_descriptor.mipmap_level_bias - tex_desc.minMipmapLevelClamp = texture_descriptor.min_mipmap_level_clamp - tex_desc.maxMipmapLevelClamp = texture_descriptor.max_mipmap_level_clamp - - cdef int i - if texture_descriptor.border_color is None: - for i in range(4): - tex_desc.borderColor[i] = 0.0 - else: - bc = tuple(texture_descriptor.border_color) - if len(bc) != 4: - raise ValueError( - f"border_color must have 4 elements, got {len(bc)}" - ) - for i in range(4): - tex_desc.borderColor[i] = bc[i] - - cdef TexObjectHandle h - if resource.kind == "array": - h = create_tex_object_handle_array(res_desc, tex_desc, arr._handle) - elif resource.kind == "mipmapped_array": - h = create_tex_object_handle_mipmap(res_desc, tex_desc, mip._handle) - else: # linear or pitch2d — both backed by a device Buffer - h = create_tex_object_handle_linear(res_desc, tex_desc, buf._h_ptr) - if not h: - HANDLE_RETURN(get_last_error()) - - cdef TextureObject self = cls.__new__(cls) - self._handle = h - self._source_ref = resource - self._texture_desc = texture_descriptor - self._device_id = _get_current_device_id() - return self - @property def handle(self): """The underlying ``CUtexObject`` as an integer (64-bit kernel arg).""" @@ -547,9 +436,9 @@ cdef class TextureObject: return self._source_ref @property - def texture_descriptor(self): - """The :class:`TextureDescriptor` this texture was built from.""" - return self._texture_desc + def options(self): + """The :class:`TextureObjectOptions` this texture was built from.""" + return self._options @property def device(self): @@ -574,3 +463,126 @@ cdef class TextureObject: def __repr__(self): return f"TextureObject(handle=0x{as_intptr(self._handle):x})" + + +def _create_texture_object(resource, options): + """Create a :class:`TextureObject` on the current device. + + Backs :meth:`cuda.core.Device.create_texture_object`. ``resource`` is a + :class:`ResourceDescriptor`; ``options`` is a :class:`TextureObjectOptions` + (or a mapping accepted by it). + """ + if not isinstance(resource, ResourceDescriptor): + raise TypeError( + f"resource must be a ResourceDescriptor, got " + f"{type(resource).__name__}" + ) + cdef object opts = check_or_create_options( + TextureObjectOptions, options, "Texture object options" + ) + + cdef cydriver.CUDA_RESOURCE_DESC res_desc + cdef cydriver.CUDA_TEXTURE_DESC tex_desc + memset(&res_desc, 0, sizeof(res_desc)) + memset(&tex_desc, 0, sizeof(tex_desc)) + + # --- Resource descriptor --- + cdef OpaqueArray arr + cdef MipmappedArray mip + cdef Buffer buf + cdef intptr_t devptr + if resource.kind == "array": + arr = resource.source + res_desc.resType = cydriver.CU_RESOURCE_TYPE_ARRAY + res_desc.res.array.hArray = as_cu(arr._handle) + elif resource.kind == "mipmapped_array": + mip = resource.source + res_desc.resType = cydriver.CU_RESOURCE_TYPE_MIPMAPPED_ARRAY + res_desc.res.mipmap.hMipmappedArray = as_cu(mip._handle) + elif resource.kind == "linear": + buf = resource.source + devptr = int(buf.handle) + res_desc.resType = cydriver.CU_RESOURCE_TYPE_LINEAR + res_desc.res.linear.devPtr = devptr + res_desc.res.linear.format = resource._format + res_desc.res.linear.numChannels = resource._num_channels + res_desc.res.linear.sizeInBytes = resource._size_bytes + elif resource.kind == "pitch2d": + buf = resource.source + devptr = int(buf.handle) + res_desc.resType = cydriver.CU_RESOURCE_TYPE_PITCH2D + res_desc.res.pitch2D.devPtr = devptr + res_desc.res.pitch2D.format = resource._format + res_desc.res.pitch2D.numChannels = resource._num_channels + res_desc.res.pitch2D.width = resource._width + res_desc.res.pitch2D.height = resource._height + res_desc.res.pitch2D.pitchInBytes = resource._pitch_bytes + else: + raise NotImplementedError( + f"ResourceDescriptor kind {resource.kind!r} is not yet supported" + ) + + # --- Texture descriptor --- + # filter_mode/read_mode/mipmap_filter_mode are normalized to their + # StrEnum types by TextureObjectOptions.__post_init__; address_mode is + # normalized (and str-coerced) here. + modes = _normalize_address_modes(opts.address_mode) + tex_desc.addressMode[0] = _ADDRESSMODE_TO_CU[modes[0]] + tex_desc.addressMode[1] = _ADDRESSMODE_TO_CU[modes[1]] + tex_desc.addressMode[2] = _ADDRESSMODE_TO_CU[modes[2]] + + tex_desc.filterMode = _FILTERMODE_TO_CU[opts.filter_mode] + + cdef unsigned int flags = 0 + # CU_TRSF_READ_AS_INTEGER suppresses normalization, so it maps to + # ReadModeType.ELEMENT_TYPE. + if opts.read_mode == ReadModeType.ELEMENT_TYPE: + flags |= _TRSF_READ_AS_INTEGER + if opts.normalized_coords: + flags |= _TRSF_NORMALIZED_COORDINATES + if opts.srgb: + flags |= _TRSF_SRGB + if opts.disable_trilinear_optimization: + flags |= _TRSF_DISABLE_TRILINEAR_OPTIMIZATION + if opts.seamless_cubemap: + flags |= _TRSF_SEAMLESS_CUBEMAP + tex_desc.flags = flags + + if opts.max_anisotropy < 0: + raise ValueError("max_anisotropy must be >= 0") + tex_desc.maxAnisotropy = opts.max_anisotropy + + tex_desc.mipmapFilterMode = _FILTERMODE_TO_CU[opts.mipmap_filter_mode] + tex_desc.mipmapLevelBias = opts.mipmap_level_bias + tex_desc.minMipmapLevelClamp = opts.min_mipmap_level_clamp + tex_desc.maxMipmapLevelClamp = opts.max_mipmap_level_clamp + + cdef int i + if opts.border_color is None: + for i in range(4): + tex_desc.borderColor[i] = 0.0 + else: + bc = tuple(opts.border_color) + if len(bc) != 4: + raise ValueError( + f"border_color must have 4 elements, got {len(bc)}" + ) + for i in range(4): + tex_desc.borderColor[i] = bc[i] + + cdef TexObjectHandle h + if resource.kind == "array": + h = create_tex_object_handle_array(res_desc, tex_desc, arr._handle) + elif resource.kind == "mipmapped_array": + h = create_tex_object_handle_mipmap(res_desc, tex_desc, mip._handle) + else: # linear or pitch2d — both backed by a device Buffer + h = create_tex_object_handle_linear(res_desc, tex_desc, buf._h_ptr) + if not h: + HANDLE_RETURN(get_last_error()) + + cdef TextureObject self = TextureObject.__new__(TextureObject) + self._handle = h + self._source_ref = resource + self._options = opts + self._device_id = _get_current_device_id() + return self diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 278b516217e..2d253d8ca5e 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -32,9 +32,12 @@ class StrEnum(str, Enum): from cuda.core._utils.cuda_utils import driver __all__ = [ + "AddressModeType", + "ArrayFormatType", "CompilerBackendType", "DevicePointerType", "DeviceResourcesType", + "FilterModeType", "GraphConditionalType", "GraphMemoryType", "IsStreamType", @@ -42,6 +45,7 @@ class StrEnum(str, Enum): "ObjectCodeFormatType", "PCHStatusType", "ProcessStateType", + "ReadModeType", "SourceCodeType", "VirtualMemoryAccessType", "VirtualMemoryAllocationType", @@ -219,4 +223,77 @@ class VirtualMemoryAllocationType(StrEnum): MANAGED = "managed" +class ArrayFormatType(StrEnum): + """Element format for an :class:`~cuda.core.texture.OpaqueArray` allocation. + + Corresponds to ``CUarray_format`` from the CUDA driver API. Each value maps + 1:1 to a NumPy dtype; the enum is retained as an explicit escape hatch. + + * ``UINT8`` / ``UINT16`` / ``UINT32`` — unsigned integer elements. + * ``INT8`` / ``INT16`` / ``INT32`` — signed integer elements. + * ``FLOAT16`` / ``FLOAT32`` — half- and single-precision float elements. + + .. versionadded:: 1.1.0 + """ + + UINT8 = "uint8" + UINT16 = "uint16" + UINT32 = "uint32" + INT8 = "int8" + INT16 = "int16" + INT32 = "int32" + FLOAT16 = "float16" + FLOAT32 = "float32" + + +class AddressModeType(StrEnum): + """Boundary behavior for out-of-range texture coordinates. + + Corresponds to ``CUaddress_mode`` from the CUDA driver API. + + * ``WRAP`` — wrap coordinates around (tiling). + * ``CLAMP`` — clamp to the edge texel. + * ``MIRROR`` — reflect coordinates at the boundary. + * ``BORDER`` — return the configured border color. + + .. versionadded:: 1.1.0 + """ + + WRAP = "wrap" + CLAMP = "clamp" + MIRROR = "mirror" + BORDER = "border" + + +class FilterModeType(StrEnum): + """Texel sampling mode for a :class:`~cuda.core.texture.TextureObject`. + + Corresponds to ``CUfilter_mode`` from the CUDA driver API. + + * ``POINT`` — nearest-texel sampling. + * ``LINEAR`` — (bi/tri)linear interpolation. + + .. versionadded:: 1.1.0 + """ + + POINT = "point" + LINEAR = "linear" + + +class ReadModeType(StrEnum): + """How sampled values are returned to the kernel. + + * ``ELEMENT_TYPE`` — return the raw element value (integer formats stay + integer, float stays float). + * ``NORMALIZED_FLOAT`` — integer formats are promoted to a normalized + ``float`` in ``[0, 1]`` (unsigned) or ``[-1, 1]`` (signed). Float + formats are unaffected. + + .. versionadded:: 1.1.0 + """ + + ELEMENT_TYPE = "element_type" + NORMALIZED_FLOAT = "normalized_float" + + del StrEnum diff --git a/cuda_core/docs/source/api.rst b/cuda_core/docs/source/api.rst index 5008e66ab74..d0078b073d4 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -167,10 +167,13 @@ Textures and surfaces CUDA arrays back bindless texture and surface objects for kernel-side sampled reads and typed load/store. These types live in the :mod:`cuda.core.texture` namespace. :class:`OpaqueArray` is allocated through -:meth:`OpaqueArray.from_descriptor` and bound through a :class:`ResourceDescriptor` -factory; linear (1D) and row-pitched 2D :class:`Buffer` views as well as -mipmapped allocations (:class:`MipmappedArray`) are also supported as texture -backings. +:meth:`cuda.core.Device.create_opaque_array` and bound through a +:class:`ResourceDescriptor` factory; linear (1D) and row-pitched 2D +:class:`Buffer` views as well as mipmapped allocations (:class:`MipmappedArray`, +via :meth:`cuda.core.Device.create_mipmapped_array`) are also supported as +texture backings. Bindless handles are created with +:meth:`cuda.core.Device.create_texture_object` and +:meth:`cuda.core.Device.create_surface_object`. A :class:`OpaqueArray` has an opaque, hardware-defined layout with no linear device pointer, so it cannot participate in ``__cuda_array_interface__`` / @@ -193,15 +196,16 @@ DLPack zero-copy interop. Data is moved in and out only by copying — use :template: dataclass.rst - TextureDescriptor + OpaqueArrayOptions + MipmappedArrayOptions + TextureObjectOptions -.. autosummary:: - :toctree: generated/ - - ArrayFormat - AddressMode - FilterMode - ReadMode +The associated enumerations — +:class:`~cuda.core.typing.ArrayFormatType`, +:class:`~cuda.core.typing.AddressModeType`, +:class:`~cuda.core.typing.FilterModeType`, and +:class:`~cuda.core.typing.ReadModeType` — live in :mod:`cuda.core.typing` +alongside the other ``cuda.core`` enumerations. CUDA compilation toolchain diff --git a/cuda_core/docs/source/api_private.rst b/cuda_core/docs/source/api_private.rst index f8ce51d6a3a..907fc2f5bcf 100644 --- a/cuda_core/docs/source/api_private.rst +++ b/cuda_core/docs/source/api_private.rst @@ -21,15 +21,19 @@ CUDA runtime _module.KernelOccupancy _module.MaxPotentialBlockSizeOccupancyResult _module.ParamInfo + typing.AddressModeType + typing.ArrayFormatType typing.CompilerBackendType typing.DevicePointerType typing.DeviceResourcesType + typing.FilterModeType typing.GraphConditionalType typing.GraphMemoryType typing.ManagedMemoryLocationType typing.ObjectCodeFormatType typing.PCHStatusType typing.ProcessStateType + typing.ReadModeType typing.SourceCodeType typing.VirtualMemoryAccessType typing.VirtualMemoryAllocationType diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index e2cdb1037e2..68259d32e3b 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -108,15 +108,15 @@ launch, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, - OpaqueArray, - ReadMode, + OpaqueArrayOptions, ResourceDescriptor, - SurfaceObject, - TextureDescriptor, - TextureObject, + TextureObjectOptions, +) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, ) # --------------------------------------------------------------------------- @@ -398,13 +398,13 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # ============================ API MAP (cuda.core) =========================== # # The three helpers below are where every OpaqueArray / ResourceDescriptor / -# TextureDescriptor / TextureObject / SurfaceObject knob in this example is set. +# TextureObjectOptions / TextureObject / SurfaceObject knob in this example is set. # Each visible setting maps to a concrete piece of cuda.core / CUDA behavior: # -# OpaqueArray.from_descriptor(...) -> allocates a CUDA *array* (opaque, tiled +# Device.create_opaque_array(...) -> allocates a CUDA *array* (opaque, tiled # layout optimized for 2D texture fetches), # not linear device memory. -# ArrayFormat.FLOAT32 -> each channel is a 32-bit float texel. +# ArrayFormatType.FLOAT32 -> each channel is a 32-bit float texel. # num_channels=2 / num_channels=1 -> float2 (vx, vy) vs scalar (pressure / # divergence / dye); also fixes the # surf2Dwrite byte offset per element. @@ -414,22 +414,22 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # is what lets each field be sampled and # then written back in the ping-pong. # -# ResourceDescriptor.from_array(arr) -> wraps the OpaqueArray as the resource a -# TextureObject reads from. -# FilterMode.LINEAR -> free HARDWARE bilinear interpolation; +# ResourceDescriptor.from_opaque_array -> wraps the OpaqueArray as the resource a +# TextureObject reads from. +# FilterModeType.LINEAR -> free HARDWARE bilinear interpolation; # this is what makes semi-Lagrangian # advection a single tex2D fetch at a # fractional back-traced position (no # manual lerp, no neighbor gather). -# AddressMode.CLAMP -> bounded box boundary: out-of-range traces +# AddressModeType.CLAMP -> bounded box boundary: out-of-range traces # read the edge texel (ink piles up at the # walls instead of wrapping like a torus). -# ReadMode.ELEMENT_TYPE -> return the stored float value as-is (no +# ReadModeType.ELEMENT_TYPE -> return the stored float value as-is (no # integer->[0,1] normalization of texels). # normalized_coords=True -> sample in [0, 1) so CLAMP is well-defined # and texel centers are (i + 0.5) / N. # -# SurfaceObject.from_array(arr) -> binds the array for surf2Dread/surf2Dwrite. +# make_surface(arr) -> binds the array for surf2Dread/surf2Dwrite. # The x coordinate is in BYTES, so it is # x * sizeof(elem): sizeof(float2)=8 for # velocity, sizeof(float)=4 for the scalars. @@ -438,21 +438,25 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): def make_velocity_array(): """Allocate a `float2` velocity CUDA array (channel 0 = vx, channel 1 = vy).""" - return OpaqueArray.from_descriptor( - shape=(WIDTH, HEIGHT), - format=ArrayFormat.FLOAT32, - num_channels=2, - is_surface_load_store=True, + return Device().create_opaque_array( + OpaqueArrayOptions( + shape=(WIDTH, HEIGHT), + format=ArrayFormatType.FLOAT32, + num_channels=2, + is_surface_load_store=True, + ) ) def make_scalar_array(): """Allocate a single-channel `float` CUDA array (pressure / divergence / dye).""" - return OpaqueArray.from_descriptor( - shape=(WIDTH, HEIGHT), - format=ArrayFormat.FLOAT32, - num_channels=1, - is_surface_load_store=True, + return Device().create_opaque_array( + OpaqueArrayOptions( + shape=(WIDTH, HEIGHT), + format=ArrayFormatType.FLOAT32, + num_channels=1, + is_surface_load_store=True, + ) ) @@ -464,11 +468,13 @@ def make_color_array(): surface-write machinery as the scalar fields -- only the channel count (and the surf2Dwrite byte stride, sizeof(float4) = 16) differ. """ - return OpaqueArray.from_descriptor( - shape=(WIDTH, HEIGHT), - format=ArrayFormat.FLOAT32, - num_channels=4, - is_surface_load_store=True, + return Device().create_opaque_array( + OpaqueArrayOptions( + shape=(WIDTH, HEIGHT), + format=ArrayFormatType.FLOAT32, + num_channels=4, + is_surface_load_store=True, + ) ) @@ -479,16 +485,22 @@ def make_texture(arr): needs the bilinear interpolation, and the stencil reads (divergence, Jacobi, gradient) sample exactly at texel centers so LINEAR returns the exact value. """ - res_desc = ResourceDescriptor.from_array(arr) - tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + res_desc = ResourceDescriptor.from_opaque_array(arr) + tex_desc = TextureObjectOptions( + address_mode=AddressModeType.CLAMP, + filter_mode=FilterModeType.LINEAR, + read_mode=ReadModeType.ELEMENT_TYPE, # Normalized coordinates keep CLAMP addressing well-defined and let us # sample at texel centers as (i + 0.5) / N. normalized_coords=True, ) - return TextureObject.from_descriptor(resource=res_desc, texture_descriptor=tex_desc) + return Device().create_texture_object(resource=res_desc, options=tex_desc) + + +def make_surface(arr): + """Bind `arr` as a SurfaceObject for raw surf2Dwrite writes.""" + res_desc = ResourceDescriptor.from_opaque_array(arr) + return Device().create_surface_object(resource=res_desc) def seed_field(stream, kernels, config, vel_surf, dye_surf, prs_surf, seed_value): @@ -552,25 +564,25 @@ def main(): # up front and keep them alive for the whole run. # API MAP: make_texture binds an array as a read-only TextureObject # (LINEAR + CLAMP + normalized; see the API MAP block above), while - # SurfaceObject.from_array binds the SAME array for raw surf2Dwrite + # make_surface binds the SAME array for raw surf2Dwrite # writes -- the read/write halves of one ping-pong buffer. vel_tex_a = make_texture(vel_a) vel_tex_b = make_texture(vel_b) - vel_surf_a = SurfaceObject.from_array(vel_a) - vel_surf_b = SurfaceObject.from_array(vel_b) + vel_surf_a = make_surface(vel_a) + vel_surf_b = make_surface(vel_b) prs_tex_a = make_texture(prs_a) prs_tex_b = make_texture(prs_b) - prs_surf_a = SurfaceObject.from_array(prs_a) - prs_surf_b = SurfaceObject.from_array(prs_b) + prs_surf_a = make_surface(prs_a) + prs_surf_b = make_surface(prs_b) div_tex = make_texture(div) - div_surf = SurfaceObject.from_array(div) + div_surf = make_surface(div) dye_tex_a = make_texture(dye_a) dye_tex_b = make_texture(dye_b) - dye_surf_a = SurfaceObject.from_array(dye_a) - dye_surf_b = SurfaceObject.from_array(dye_b) + dye_surf_a = make_surface(dye_a) + dye_surf_b = make_surface(dye_b) # --- Step 8: Seed the initial field (curl into vel_a, zero pressure/dye) --- seed_field(stream, kernels, config, vel_surf_a, dye_surf_a, prs_surf_a, seed_value=0) diff --git a/cuda_core/examples/gl_interop_fluid_numba_cuda_mlir.py b/cuda_core/examples/gl_interop_fluid_numba_cuda_mlir.py index 8500c3d54c9..791b3bd67df 100644 --- a/cuda_core/examples/gl_interop_fluid_numba_cuda_mlir.py +++ b/cuda_core/examples/gl_interop_fluid_numba_cuda_mlir.py @@ -23,7 +23,7 @@ # -------------------------------------- ----------------------------------- # OpaqueArray(num_channels=2) as texture -> cuda.device_array((H, W, 2), f32) # tex2D(tex, u, v) [HW LINEAR] -> sample_vec(arr, px, py) [manual lerp] -# AddressMode.CLAMP -> index clamp inside sample_*() +# AddressModeType.CLAMP -> index clamp inside sample_*() # surf2Dwrite(v, surf, x*8, y) -> arr[y, x, 0] = v.x; arr[y, x, 1] = v.y # TextureObject + SurfaceObject pair -> one device array; ping-pong by swap # GraphicsResource PBO (zero-copy) -> copy_to_host + glTexSubImage2D @@ -91,7 +91,7 @@ @cuda.jit(device=True, inline=True) def _clampi(i, n): - # AddressMode.CLAMP: out-of-range coordinates read the border texel. + # AddressModeType.CLAMP: out-of-range coordinates read the border texel. if i < 0: return 0 if i > n - 1: diff --git a/cuda_core/examples/gl_interop_mipmap_lod.py b/cuda_core/examples/gl_interop_mipmap_lod.py index ea82ddbd640..8c13b17c8ed 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -101,15 +101,15 @@ launch, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, - MipmappedArray, - ReadMode, + MipmappedArrayOptions, ResourceDescriptor, - SurfaceObject, - TextureDescriptor, - TextureObject, + TextureObjectOptions, +) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, ) # --------------------------------------------------------------------------- @@ -183,7 +183,7 @@ def build_mipmap_pyramid(mip, num_levels, stream, kernels): """ # ---- Level 0: seed the base image ------------------------------------- base_arr = mip.get_level(0) # non-owning view; do NOT use a `with` block - with SurfaceObject.from_array(base_arr) as base_surf: + with Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(base_arr)) as base_surf: block = (16, 16, 1) grid = ( (BASE_SIZE + block[0] - 1) // block[0], @@ -205,10 +205,10 @@ def build_mipmap_pyramid(mip, num_levels, stream, kernels): # Each iteration reads level (L-1) through a temporary TextureObject and # writes level L through a temporary SurfaceObject. Both close cleanly # at the end of their `with` blocks. - src_tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.POINT, # explicit per-texel reads - read_mode=ReadMode.ELEMENT_TYPE, + src_tex_desc = TextureObjectOptions( + address_mode=AddressModeType.CLAMP, + filter_mode=FilterModeType.POINT, # explicit per-texel reads + read_mode=ReadModeType.ELEMENT_TYPE, normalized_coords=False, # integer pixel coordinates ) for level in range(1, num_levels): @@ -219,10 +219,10 @@ def build_mipmap_pyramid(mip, num_levels, stream, kernels): src_arr = mip.get_level(level - 1) dst_arr = mip.get_level(level) - src_res = ResourceDescriptor.from_array(src_arr) + src_res = ResourceDescriptor.from_opaque_array(src_arr) with ( - TextureObject.from_descriptor(resource=src_res, texture_descriptor=src_tex_desc) as src_tex, - SurfaceObject.from_array(dst_arr) as dst_surf, + Device().create_texture_object(resource=src_res, options=src_tex_desc) as src_tex, + Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(dst_arr)) as dst_surf, ): block = (16, 16, 1) grid = ( @@ -408,12 +408,14 @@ def main(): # --- Step 2: Allocate the mipmap pyramid and build every level --- # is_surface_load_store=True is required for kernel-side writes. num_levels = int(math.log2(BASE_SIZE)) + 1 - mip = MipmappedArray.from_descriptor( - shape=(BASE_SIZE, BASE_SIZE), - format=ArrayFormat.FLOAT32, - num_channels=4, - num_levels=num_levels, - is_surface_load_store=True, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(BASE_SIZE, BASE_SIZE), + format=ArrayFormatType.FLOAT32, + num_channels=4, + num_levels=num_levels, + is_surface_load_store=True, + ) ) build_mipmap_pyramid(mip, num_levels, stream, kernels) @@ -423,19 +425,19 @@ def main(): # receives the user-controlled bias as a kernel argument and folds # it into the tex2DLod call (avoids rebuilding the TextureObject # whenever the user changes the bias). - display_tex_desc = TextureDescriptor( - address_mode=AddressMode.WRAP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + display_tex_desc = TextureObjectOptions( + address_mode=AddressModeType.WRAP, + filter_mode=FilterModeType.LINEAR, + read_mode=ReadModeType.ELEMENT_TYPE, normalized_coords=True, - mipmap_filter_mode=FilterMode.LINEAR, # trilinear + mipmap_filter_mode=FilterModeType.LINEAR, # trilinear mipmap_level_bias=0.0, min_mipmap_level_clamp=0.0, max_mipmap_level_clamp=float(num_levels - 1), ) - display_tex = TextureObject.from_descriptor( + display_tex = Device().create_texture_object( resource=ResourceDescriptor.from_mipmapped_array(mip), - texture_descriptor=display_tex_desc, + options=display_tex_desc, ) # --- Step 4: Open a window and set up the GL/CUDA bridge --- diff --git a/cuda_core/examples/texture_sample.py b/cuda_core/examples/texture_sample.py index e69655eb479..952c77c48ca 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -30,14 +30,15 @@ launch, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, - OpaqueArray, - ReadMode, + OpaqueArrayOptions, ResourceDescriptor, - TextureDescriptor, - TextureObject, + TextureObjectOptions, +) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, ) # Kernel reads N (x, y) coordinates from `coords` (interleaved float pairs) and @@ -66,14 +67,16 @@ def main(): pinned_mr = LegacyPinnedMemoryResource() try: # Allocate a 2D OpaqueArray: shape=(W, H), single-channel float32. - # Note: OpaqueArray.from_descriptor takes shape=(width, height), so the host + # Note: create_opaque_array takes shape=(width, height), so the host # buffer fed into copy_from must be laid out as H rows of W elements # (row-major), i.e. host_pattern.shape == (H, W). width, height = 16, 16 - with OpaqueArray.from_descriptor( - shape=(width, height), - format=ArrayFormat.FLOAT32, - num_channels=1, + with Device().create_opaque_array( + OpaqueArrayOptions( + shape=(width, height), + format=ArrayFormatType.FLOAT32, + num_channels=1, + ) ) as arr: # Plant a known pattern: pattern[y, x] = x + 100*y. # Cast to float32 so the byte count matches the array's storage. @@ -87,14 +90,14 @@ def main(): arr.copy_from(pattern, stream=stream) # Build a linear-filtering, clamped, non-normalized texture. - res_desc = ResourceDescriptor.from_array(arr) - tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + res_desc = ResourceDescriptor.from_opaque_array(arr) + tex_desc = TextureObjectOptions( + address_mode=AddressModeType.CLAMP, + filter_mode=FilterModeType.LINEAR, + read_mode=ReadModeType.ELEMENT_TYPE, normalized_coords=False, ) - with TextureObject.from_descriptor(resource=res_desc, texture_descriptor=tex_desc) as tex: + with Device().create_texture_object(resource=res_desc, options=tex_desc) as tex: _run_kernel_and_verify(dev, stream, tex, pattern, width, height, pinned_mr) finally: stream.close() diff --git a/cuda_core/tests/test_enum_coverage.py b/cuda_core/tests/test_enum_coverage.py index 308614679c8..09c12bbaf52 100644 --- a/cuda_core/tests/test_enum_coverage.py +++ b/cuda_core/tests/test_enum_coverage.py @@ -98,6 +98,25 @@ }, set(), ), + # AddressModeType / FilterModeType are 1:1 wrappers of small driver enums. + # The texture mapping dicts (_ADDRESSMODE_TO_CU / _FILTERMODE_TO_CU) store + # plain ints keyed on cydriver values, not driver. members, so they + # are incompatible with the isinstance-based mapping check. Use the + # mapping=None form, which count-checks the StrEnum against the binding enum. + ( + driver.CUaddress_mode, + cuda.core.typing.AddressModeType, + None, + set(), + set(), + ), + ( + driver.CUfilter_mode, + cuda.core.typing.FilterModeType, + None, + set(), + set(), + ), ] if system.CUDA_BINDINGS_NVML_IS_COMPATIBLE: @@ -278,6 +297,16 @@ cuda.core.typing.SourceCodeType, # This enum is dynamic depending on the version of CTK installed. cuda.core.typing.VirtualMemoryAllocationType, + # ArrayFormatType wraps CUarray_format, but intentionally exposes only the + # 8 NumPy-representable formats out of the driver's ~67 members (BC1..BC7, + # UNORM_INTx, packed/planar YUV, etc. are deliberately unsupported for now). + # It is a curated subset, not a 1:1 wrapper, so a count/coverage check + # against the binding enum does not apply. + cuda.core.typing.ArrayFormatType, + # ReadModeType (ELEMENT_TYPE / NORMALIZED_FLOAT) has no backing cuda_binding + # enum: it maps to the presence/absence of the CU_TRSF_READ_AS_INTEGER + # texture-descriptor flag bit, not to a CUenum. + cuda.core.typing.ReadModeType, } diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 1a6d05e65a2..bee60104be7 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -3,6 +3,7 @@ import gc +import numpy as np import pytest import cuda.core @@ -10,16 +11,17 @@ Device, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, - MipmappedArray, + MipmappedArrayOptions, OpaqueArray, - ReadMode, + OpaqueArrayOptions, ResourceDescriptor, - SurfaceObject, - TextureDescriptor, - TextureObject, + TextureObjectOptions, +) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, ) @@ -44,12 +46,14 @@ def test_resource_descriptor_init_disabled(): def test_array_2d_create_and_properties(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(32, 16), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1) + ) try: assert arr.shape == (32, 16) - assert arr.format == ArrayFormat.FLOAT32 + assert arr.format == ArrayFormatType.FLOAT32 assert arr.num_channels == 1 - assert arr.element_size == 4 + assert arr.element_bytes == 4 assert arr.size_bytes == 32 * 16 * 4 assert arr.is_surface_load_store is False assert arr.handle != 0 @@ -59,28 +63,66 @@ def test_array_2d_create_and_properties(init_cuda): def test_array_3d_with_surface_flag(init_cuda): - arr = OpaqueArray.from_descriptor( - shape=(8, 8, 4), - format=ArrayFormat.UINT8, - num_channels=4, - is_surface_load_store=True, + arr = Device().create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8, 4), + format=ArrayFormatType.UINT8, + num_channels=4, + is_surface_load_store=True, + ) ) try: assert arr.shape == (8, 8, 4) assert arr.is_surface_load_store is True - assert arr.element_size == 4 + assert arr.element_bytes == 4 + finally: + arr.close() + + +@pytest.mark.agent_authored(model="claude-opus-4.8") +@pytest.mark.parametrize( + "dtype, expected", + [ + (np.float32, ArrayFormatType.FLOAT32), + (np.dtype("float16"), ArrayFormatType.FLOAT16), + (np.uint8, ArrayFormatType.UINT8), + (np.dtype("i4"), ArrayFormatType.INT32), + ], +) +def test_array_accepts_numpy_dtype_format(init_cuda, dtype, expected): + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=dtype, num_channels=1)) + try: + assert arr.format == expected + finally: + arr.close() + + +@pytest.mark.agent_authored(model="claude-opus-4.8") +def test_array_accepts_str_format(init_cuda): + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format="float32", num_channels=1)) + try: + assert arr.format == ArrayFormatType.FLOAT32 finally: arr.close() +@pytest.mark.agent_authored(model="claude-opus-4.8") +def test_array_rejects_unsupported_dtype_format(init_cuda): + # float64 has no ArrayFormatType equivalent. + with pytest.raises(ValueError, match="no ArrayFormatType equivalent"): + Device().create_opaque_array(OpaqueArrayOptions(shape=(8,), format=np.float64, num_channels=1)) + + def test_array_rejects_bad_channels(init_cuda): with pytest.raises(ValueError, match="num_channels"): - OpaqueArray.from_descriptor(shape=(8,), format=ArrayFormat.UINT8, num_channels=3) + Device().create_opaque_array(OpaqueArrayOptions(shape=(8,), format=ArrayFormatType.UINT8, num_channels=3)) def test_array_rejects_bad_rank(init_cuda): with pytest.raises(ValueError, match="shape rank"): - OpaqueArray.from_descriptor(shape=(2, 2, 2, 2), format=ArrayFormat.UINT8, num_channels=1) + Device().create_opaque_array( + OpaqueArrayOptions(shape=(2, 2, 2, 2), format=ArrayFormatType.UINT8, num_channels=1) + ) def test_array_roundtrip_copy(init_cuda): @@ -88,7 +130,7 @@ def test_array_roundtrip_copy(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormat.UINT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1)) try: src = _array.array("I", list(range(16))) dst = _array.array("I", [0] * 16) @@ -108,7 +150,7 @@ def test_array_copy_rejects_undersized_host_buffer(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormat.UINT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1)) try: # arr is 16 * 4 = 64 bytes; pass an 8-element (32-byte) host buffer. too_small = _array.array("I", [0] * 8) @@ -124,7 +166,7 @@ def test_array_copy_rejects_undersized_host_buffer(init_cuda): def test_array_copy_rejects_undersized_device_buffer(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormat.UINT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1)) # arr is 64 bytes; allocate a 32-byte device buffer. small_buf = device.memory_resource.allocate(32, stream=device.default_stream) try: @@ -139,20 +181,22 @@ def test_array_copy_rejects_undersized_device_buffer(init_cuda): def test_texture_object_create(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(32, 16), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1) + ) try: - res = ResourceDescriptor.from_array(arr) - tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + res = ResourceDescriptor.from_opaque_array(arr) + tex_desc = TextureObjectOptions( + address_mode=AddressModeType.CLAMP, + filter_mode=FilterModeType.LINEAR, + read_mode=ReadModeType.ELEMENT_TYPE, normalized_coords=True, ) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=tex_desc) + tex = Device().create_texture_object(resource=res, options=tex_desc) try: assert tex.handle != 0 assert tex.resource is res - assert tex.texture_descriptor is tex_desc + assert tex.options is tex_desc finally: tex.close() finally: @@ -160,14 +204,16 @@ def test_texture_object_create(init_cuda): def test_surface_object_create(init_cuda): - arr = OpaqueArray.from_descriptor( - shape=(8, 8), - format=ArrayFormat.UINT8, - num_channels=4, - is_surface_load_store=True, + arr = Device().create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + is_surface_load_store=True, + ) ) try: - surf = SurfaceObject.from_array(arr) + surf = Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(arr)) try: assert surf.handle != 0 assert isinstance(surf.resource, ResourceDescriptor) @@ -178,10 +224,10 @@ def test_surface_object_create(init_cuda): def test_surface_requires_ldst_flag(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.UINT8, num_channels=4) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4)) try: with pytest.raises(ValueError, match="is_surface_load_store=True"): - SurfaceObject.from_array(arr) + Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(arr)) finally: arr.close() @@ -191,28 +237,30 @@ def test_address_mode_normalization(init_cuda): # 3-tuple; a shorter tuple should be padded by repeating the last entry. from cuda.core.texture._texture import _normalize_address_modes - assert _normalize_address_modes(AddressMode.WRAP) == ( - AddressMode.WRAP, - AddressMode.WRAP, - AddressMode.WRAP, + assert _normalize_address_modes(AddressModeType.WRAP) == ( + AddressModeType.WRAP, + AddressModeType.WRAP, + AddressModeType.WRAP, ) - assert _normalize_address_modes((AddressMode.WRAP, AddressMode.CLAMP)) == ( - AddressMode.WRAP, - AddressMode.CLAMP, - AddressMode.CLAMP, + assert _normalize_address_modes((AddressModeType.WRAP, AddressModeType.CLAMP)) == ( + AddressModeType.WRAP, + AddressModeType.CLAMP, + AddressModeType.CLAMP, ) - assert _normalize_address_modes((AddressMode.WRAP, AddressMode.CLAMP, AddressMode.MIRROR)) == ( - AddressMode.WRAP, - AddressMode.CLAMP, - AddressMode.MIRROR, + assert _normalize_address_modes((AddressModeType.WRAP, AddressModeType.CLAMP, AddressModeType.MIRROR)) == ( + AddressModeType.WRAP, + AddressModeType.CLAMP, + AddressModeType.MIRROR, ) # Smoke test: a 2-entry tuple is also accepted end-to-end. - arr = OpaqueArray.from_descriptor(shape=(8, 8, 4), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1) + ) try: - res = ResourceDescriptor.from_array(arr) - tex_desc = TextureDescriptor(address_mode=(AddressMode.WRAP, AddressMode.CLAMP)) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=tex_desc) + res = ResourceDescriptor.from_opaque_array(arr) + tex_desc = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) + tex = Device().create_texture_object(resource=res, options=tex_desc) try: assert tex.handle != 0 finally: @@ -233,9 +281,9 @@ def test_resource_descriptor_from_linear_defaults_size(init_cuda): device = Device() buf = _alloc_device_buffer(device, 4096) try: - res = ResourceDescriptor.from_linear(buf, format=ArrayFormat.FLOAT32, num_channels=1) + res = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.FLOAT32, num_channels=1) assert res.kind == "linear" - assert res.format == ArrayFormat.FLOAT32 + assert res.format == ArrayFormatType.FLOAT32 assert res.num_channels == 1 assert res.source is buf # repr should include the kind/format hint @@ -248,7 +296,7 @@ def test_resource_descriptor_from_linear_size_override(init_cuda): device = Device() buf = _alloc_device_buffer(device, 4096) try: - res = ResourceDescriptor.from_linear(buf, format=ArrayFormat.UINT32, num_channels=1, size_bytes=2048) + res = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT32, num_channels=1, size_bytes=2048) assert res._size_bytes == 2048 finally: buf.close() @@ -259,7 +307,7 @@ def test_resource_descriptor_from_linear_rejects_oversize(init_cuda): buf = _alloc_device_buffer(device, 1024) try: with pytest.raises(ValueError, match="exceeds buffer.size"): - ResourceDescriptor.from_linear(buf, format=ArrayFormat.UINT8, num_channels=1, size_bytes=2048) + ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT8, num_channels=1, size_bytes=2048) finally: buf.close() @@ -269,14 +317,14 @@ def test_resource_descriptor_from_linear_rejects_bad_channels(init_cuda): buf = _alloc_device_buffer(device, 1024) try: with pytest.raises(ValueError, match="num_channels"): - ResourceDescriptor.from_linear(buf, format=ArrayFormat.UINT8, num_channels=3) + ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT8, num_channels=3) finally: buf.close() def test_resource_descriptor_from_linear_rejects_non_buffer(): with pytest.raises(TypeError, match="Buffer"): - ResourceDescriptor.from_linear(object(), format=ArrayFormat.UINT8, num_channels=1) + ResourceDescriptor.from_linear(object(), format=ArrayFormatType.UINT8, num_channels=1) def test_resource_descriptor_from_linear_rejects_zero_size(init_cuda): @@ -284,7 +332,7 @@ def test_resource_descriptor_from_linear_rejects_zero_size(init_cuda): buf = _alloc_device_buffer(device, 1024) try: with pytest.raises(ValueError, match="at least one element"): - ResourceDescriptor.from_linear(buf, format=ArrayFormat.UINT32, num_channels=1, size_bytes=0) + ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT32, num_channels=1, size_bytes=0) finally: buf.close() @@ -295,7 +343,7 @@ def test_resource_descriptor_from_linear_rejects_non_multiple(init_cuda): try: # UINT32 x 1 channel = 4 bytes/element; 10 bytes is not a multiple. with pytest.raises(ValueError, match="multiple of element size"): - ResourceDescriptor.from_linear(buf, format=ArrayFormat.UINT32, num_channels=1, size_bytes=10) + ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT32, num_channels=1, size_bytes=10) finally: buf.close() @@ -307,8 +355,8 @@ def test_texture_object_from_linear(init_cuda): # 1024 float elements buf = _alloc_device_buffer(device, 1024 * 4) try: - res = ResourceDescriptor.from_linear(buf, format=ArrayFormat.FLOAT32, num_channels=1) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) + res = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.FLOAT32, num_channels=1) + tex = Device().create_texture_object(resource=res, options=TextureObjectOptions()) try: assert tex.handle != 0 assert tex.resource is res @@ -322,15 +370,15 @@ def test_resource_descriptor_from_pitch2d_validates_pitch(init_cuda): device = Device() buf = _alloc_device_buffer(device, 64 * 1024) try: - # element_size = 4 (UINT32 * 1 channel); width=16 -> min_pitch=64 + # element_bytes = 4 (UINT32 * 1 channel); width=16 -> min_pitch=64 with pytest.raises(ValueError, match="pitch_bytes"): ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT32, + format=ArrayFormatType.UINT32, num_channels=1, width=16, height=8, - pitch_bytes=32, # < 64 = width*element_size + pitch_bytes=32, # < 64 = width*element_bytes ) finally: buf.close() @@ -343,7 +391,7 @@ def test_resource_descriptor_from_pitch2d_validates_buffer_size(init_cuda): with pytest.raises(ValueError, match="exceeds buffer.size"): ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=4, width=64, height=128, @@ -370,7 +418,7 @@ def test_texture_object_from_pitch2d(init_cuda): try: res = ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=4, width=32, height=height, @@ -378,7 +426,7 @@ def test_texture_object_from_pitch2d(init_cuda): ) assert res.kind == "pitch2d" assert "pitch2d" in repr(res) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) + tex = Device().create_texture_object(resource=res, options=TextureObjectOptions()) try: assert tex.handle != 0 finally: @@ -391,20 +439,20 @@ def test_surface_rejects_linear_and_pitch2d(init_cuda): device = Device() buf = _alloc_device_buffer(device, 4096) try: - res_lin = ResourceDescriptor.from_linear(buf, format=ArrayFormat.UINT32, num_channels=1) + res_lin = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT32, num_channels=1) with pytest.raises(ValueError, match="array-backed"): - SurfaceObject.from_descriptor(resource=res_lin) + Device().create_surface_object(resource=res_lin) res_p2 = ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=4, width=8, height=8, pitch_bytes=64, ) with pytest.raises(ValueError, match="array-backed"): - SurfaceObject.from_descriptor(resource=res_p2) + Device().create_surface_object(resource=res_p2) finally: buf.close() @@ -417,16 +465,18 @@ def test_mipmapped_array_init_disabled(): cuda.core.texture._mipmapped_array.MipmappedArray() -def test_mipmapped_array_from_descriptor_2d(init_cuda): - mip = MipmappedArray.from_descriptor( - shape=(64, 32), - format=ArrayFormat.FLOAT32, - num_channels=1, - num_levels=4, +def test_mipmapped_array_create_2d(init_cuda): + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(64, 32), + format=ArrayFormatType.FLOAT32, + num_channels=1, + num_levels=4, + ) ) try: assert mip.shape == (64, 32) - assert mip.format == ArrayFormat.FLOAT32 + assert mip.format == ArrayFormatType.FLOAT32 assert mip.num_channels == 1 assert mip.num_levels == 4 assert mip.is_surface_load_store is False @@ -438,11 +488,13 @@ def test_mipmapped_array_from_descriptor_2d(init_cuda): def test_mipmapped_array_get_level_zero_matches_shape(init_cuda): shape = (64, 32) - mip = MipmappedArray.from_descriptor( - shape=shape, - format=ArrayFormat.UINT8, - num_channels=4, - num_levels=4, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=shape, + format=ArrayFormatType.UINT8, + num_channels=4, + num_levels=4, + ) ) try: lvl0 = mip.get_level(0) @@ -450,7 +502,7 @@ def test_mipmapped_array_get_level_zero_matches_shape(init_cuda): assert isinstance(lvl0, OpaqueArray) # Level 0 must match the base shape and rank. assert lvl0.shape == shape - assert lvl0.format == ArrayFormat.UINT8 + assert lvl0.format == ArrayFormatType.UINT8 assert lvl0.num_channels == 4 assert lvl0.handle != 0 finally: @@ -462,11 +514,13 @@ def test_mipmapped_array_get_level_zero_matches_shape(init_cuda): def test_mipmapped_array_get_level_halves_dims(init_cuda): shape = (64, 32) num_levels = 4 - mip = MipmappedArray.from_descriptor( - shape=shape, - format=ArrayFormat.UINT8, - num_channels=1, - num_levels=num_levels, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=shape, + format=ArrayFormatType.UINT8, + num_channels=1, + num_levels=num_levels, + ) ) try: for level in range(num_levels): @@ -482,11 +536,13 @@ def test_mipmapped_array_get_level_halves_dims(init_cuda): def test_mipmapped_array_get_level_out_of_range(init_cuda): - mip = MipmappedArray.from_descriptor( - shape=(16, 16), - format=ArrayFormat.UINT8, - num_channels=1, - num_levels=2, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(16, 16), + format=ArrayFormatType.UINT8, + num_channels=1, + num_levels=2, + ) ) try: with pytest.raises(ValueError, match="num_levels"): @@ -499,20 +555,24 @@ def test_mipmapped_array_get_level_out_of_range(init_cuda): def test_mipmapped_array_rejects_zero_levels(init_cuda): with pytest.raises(ValueError, match="num_levels"): - MipmappedArray.from_descriptor( - shape=(8, 8), - format=ArrayFormat.UINT8, - num_channels=1, - num_levels=0, + Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=1, + num_levels=0, + ) ) def test_resource_descriptor_from_mipmapped_array(init_cuda): - mip = MipmappedArray.from_descriptor( - shape=(32, 16), - format=ArrayFormat.FLOAT32, - num_channels=1, - num_levels=3, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(32, 16), + format=ArrayFormatType.FLOAT32, + num_channels=1, + num_levels=3, + ) ) try: res = ResourceDescriptor.from_mipmapped_array(mip) @@ -528,25 +588,27 @@ def test_resource_descriptor_from_mipmapped_array_rejects_non_mipmap(): def test_texture_object_from_mipmapped_array(init_cuda): - mip = MipmappedArray.from_descriptor( - shape=(32, 32), - format=ArrayFormat.FLOAT32, - num_channels=1, - num_levels=3, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(32, 32), + format=ArrayFormatType.FLOAT32, + num_channels=1, + num_levels=3, + ) ) try: res = ResourceDescriptor.from_mipmapped_array(mip) # Use non-default mipmap params so the driver exercises that path. - tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, + tex_desc = TextureObjectOptions( + address_mode=AddressModeType.CLAMP, + filter_mode=FilterModeType.LINEAR, normalized_coords=True, - mipmap_filter_mode=FilterMode.LINEAR, + mipmap_filter_mode=FilterModeType.LINEAR, mipmap_level_bias=0.0, min_mipmap_level_clamp=0.0, max_mipmap_level_clamp=float(mip.num_levels - 1), ) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=tex_desc) + tex = Device().create_texture_object(resource=res, options=tex_desc) try: assert tex.handle != 0 assert tex.resource is res @@ -557,17 +619,19 @@ def test_texture_object_from_mipmapped_array(init_cuda): def test_surface_rejects_mipmapped_array(init_cuda): - mip = MipmappedArray.from_descriptor( - shape=(16, 16), - format=ArrayFormat.UINT8, - num_channels=4, - num_levels=2, - is_surface_load_store=True, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(16, 16), + format=ArrayFormatType.UINT8, + num_channels=4, + num_levels=2, + is_surface_load_store=True, + ) ) try: res = ResourceDescriptor.from_mipmapped_array(mip) with pytest.raises(ValueError, match="array-backed"): - SurfaceObject.from_descriptor(resource=res) + Device().create_surface_object(resource=res) finally: mip.close() @@ -587,11 +651,13 @@ def test_mipmapped_array_level_outlives_dropped_parent(init_cuda): device = Device() stream = device.create_stream() - mip = MipmappedArray.from_descriptor( - shape=(16, 16), - format=ArrayFormat.UINT32, - num_channels=1, - num_levels=3, + mip = Device().create_mipmapped_array( + MipmappedArrayOptions( + shape=(16, 16), + format=ArrayFormatType.UINT32, + num_channels=1, + num_levels=3, + ) ) lvl = mip.get_level(1) # (8, 8) # Drop the only Python reference to the parent and force GC. @@ -615,29 +681,33 @@ def test_texture_surface_close_is_idempotent(init_cuda): """close() drops this object's handle reference; destruction happens once via the handle deleter. A second close() (and the later __dealloc__) must be a safe no-op, and handle must read back as 0.""" - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.UINT8, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=1)) arr.close() assert arr.handle == 0 arr.close() # second close must not raise - mip = MipmappedArray.from_descriptor(shape=(8, 8), format=ArrayFormat.UINT8, num_channels=1, num_levels=2) + mip = Device().create_mipmapped_array( + MipmappedArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=1, num_levels=2) + ) mip.close() assert mip.handle == 0 mip.close() - surf_arr = OpaqueArray.from_descriptor( - shape=(8, 8), format=ArrayFormat.UINT8, num_channels=4, is_surface_load_store=True + surf_arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True) ) - surf = SurfaceObject.from_array(surf_arr) + surf = Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(surf_arr)) surf.close() assert surf.handle == 0 surf.close() surf_arr.close() - tex_arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) - tex = TextureObject.from_descriptor( - resource=ResourceDescriptor.from_array(tex_arr), - texture_descriptor=TextureDescriptor(), + tex_arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + ) + tex = Device().create_texture_object( + resource=ResourceDescriptor.from_opaque_array(tex_arr), + options=TextureObjectOptions(), ) tex.close() assert tex.handle == 0 @@ -648,23 +718,23 @@ def test_texture_surface_close_is_idempotent(init_cuda): # --- Negative-path validation tests ------------------------------------------ -def test_array_from_descriptor_rejects_bad_format(init_cuda): - with pytest.raises(TypeError, match="format must be an ArrayFormat"): - OpaqueArray.from_descriptor(shape=(8,), format=0, num_channels=1) +def test_array_create_rejects_bad_format(init_cuda): + with pytest.raises(ValueError, match="format must be an ArrayFormatType"): + Device().create_opaque_array(OpaqueArrayOptions(shape=(8,), format=0, num_channels=1)) -def test_array_from_descriptor_rejects_non_iterable_shape(init_cuda): +def test_array_create_rejects_non_iterable_shape(init_cuda): with pytest.raises(TypeError, match="shape must be a tuple"): - OpaqueArray.from_descriptor(shape=8, format=ArrayFormat.UINT8, num_channels=1) + Device().create_opaque_array(OpaqueArrayOptions(shape=8, format=ArrayFormatType.UINT8, num_channels=1)) -def test_array_from_descriptor_rejects_zero_dim(init_cuda): +def test_array_create_rejects_zero_dim(init_cuda): with pytest.raises(ValueError, match=r"shape\[1\] must be >= 1"): - OpaqueArray.from_descriptor(shape=(8, 0), format=ArrayFormat.UINT8, num_channels=1) + Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 0), format=ArrayFormatType.UINT8, num_channels=1)) def test_array_copy_rejects_non_stream(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8,), format=ArrayFormat.UINT8, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8,), format=ArrayFormatType.UINT8, num_channels=1)) try: import array as _array @@ -683,7 +753,7 @@ def test_array_copy_to_returns_dst(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormat.UINT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1)) try: dst = _array.array("I", [0] * 16) returned = arr.copy_to(dst, stream=stream) @@ -699,7 +769,7 @@ def test_array_copy_accepts_graph_builder(init_cuda): into a CUDA graph (parity with Buffer, which accepts Stream | GraphBuilder).""" device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormat.UINT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1)) buf_in = device.memory_resource.allocate(arr.size_bytes, stream=stream) buf_out = device.memory_resource.allocate(arr.size_bytes, stream=stream) stream.sync() @@ -722,7 +792,7 @@ def test_resource_descriptor_from_pitch2d_rejects_non_buffer(): with pytest.raises(TypeError, match="buffer must be a Buffer"): ResourceDescriptor.from_pitch2d( object(), - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=1, width=8, height=8, @@ -734,7 +804,7 @@ def test_resource_descriptor_from_pitch2d_rejects_bad_format(init_cuda): device = Device() buf = _alloc_device_buffer(device, 4096) try: - with pytest.raises(TypeError, match="format must be an ArrayFormat"): + with pytest.raises(ValueError, match="format must be an ArrayFormatType"): ResourceDescriptor.from_pitch2d( buf, format=0, @@ -754,7 +824,7 @@ def test_resource_descriptor_from_pitch2d_rejects_bad_channels(init_cuda): with pytest.raises(ValueError, match="num_channels"): ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=3, width=8, height=8, @@ -771,7 +841,7 @@ def test_resource_descriptor_from_pitch2d_rejects_zero_dims(init_cuda): with pytest.raises(ValueError, match="width"): ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=1, width=0, height=8, @@ -780,7 +850,7 @@ def test_resource_descriptor_from_pitch2d_rejects_zero_dims(init_cuda): with pytest.raises(ValueError, match="height"): ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=1, width=8, height=0, @@ -791,130 +861,122 @@ def test_resource_descriptor_from_pitch2d_rejects_zero_dims(init_cuda): def test_mipmapped_array_rejects_bad_format(init_cuda): - with pytest.raises(TypeError, match="format must be an ArrayFormat"): - MipmappedArray.from_descriptor(shape=(8, 8), format=0, num_channels=1, num_levels=2) + with pytest.raises(ValueError, match="format must be an ArrayFormatType"): + Device().create_mipmapped_array(MipmappedArrayOptions(shape=(8, 8), format=0, num_channels=1, num_levels=2)) def test_mipmapped_array_rejects_bad_channels(init_cuda): with pytest.raises(ValueError, match="num_channels"): - MipmappedArray.from_descriptor(shape=(8, 8), format=ArrayFormat.UINT8, num_channels=3, num_levels=2) + Device().create_mipmapped_array( + MipmappedArrayOptions(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=3, num_levels=2) + ) def test_mipmapped_array_rejects_zero_dim(init_cuda): with pytest.raises(ValueError, match=r"shape\[0\] must be >= 1"): - MipmappedArray.from_descriptor(shape=(0, 8), format=ArrayFormat.UINT8, num_channels=1, num_levels=1) + Device().create_mipmapped_array( + MipmappedArrayOptions(shape=(0, 8), format=ArrayFormatType.UINT8, num_channels=1, num_levels=1) + ) def test_texture_object_rejects_non_resource_descriptor(init_cuda): with pytest.raises(TypeError, match="resource must be a ResourceDescriptor"): - TextureObject.from_descriptor(resource=object(), texture_descriptor=TextureDescriptor()) + Device().create_texture_object(resource=object(), options=TextureObjectOptions()) -def test_texture_object_rejects_non_texture_descriptor(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) +def test_texture_object_rejects_bad_options_type(init_cuda): + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - with pytest.raises(TypeError, match="texture_descriptor must be a TextureDescriptor"): - TextureObject.from_descriptor(resource=res, texture_descriptor="nope") + res = ResourceDescriptor.from_opaque_array(arr) + with pytest.raises(TypeError, match="must be provided as an object of type TextureObjectOptions"): + Device().create_texture_object(resource=res, options="nope") finally: arr.close() def test_texture_object_rejects_bad_filter_mode(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) - try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(filter_mode=0) # int, not FilterMode - with pytest.raises(TypeError, match="filter_mode must be a FilterMode"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) - finally: - arr.close() + # Invalid enum values are rejected at TextureObjectOptions construction. + with pytest.raises(ValueError, match="filter_mode must be a FilterModeType"): + TextureObjectOptions(filter_mode=0) # int, not FilterModeType def test_texture_object_rejects_bad_read_mode(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) - try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(read_mode=0) # int, not ReadMode - with pytest.raises(TypeError, match="read_mode must be a ReadMode"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) - finally: - arr.close() + with pytest.raises(ValueError, match="read_mode must be a ReadModeType"): + TextureObjectOptions(read_mode=0) # int, not ReadModeType def test_texture_object_rejects_bad_mipmap_filter_mode(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) - try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(mipmap_filter_mode=0) # int, not FilterMode - with pytest.raises(TypeError, match="mipmap_filter_mode must be a FilterMode"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) - finally: - arr.close() + with pytest.raises(ValueError, match="mipmap_filter_mode must be a FilterModeType"): + TextureObjectOptions(mipmap_filter_mode=0) # int, not FilterModeType def test_texture_object_rejects_negative_anisotropy(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(max_anisotropy=-1) + res = ResourceDescriptor.from_opaque_array(arr) + td = TextureObjectOptions(max_anisotropy=-1) with pytest.raises(ValueError, match="max_anisotropy"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + Device().create_texture_object(resource=res, options=td) finally: arr.close() def test_texture_object_rejects_bad_border_color_length(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(border_color=(0.0, 0.0)) # length 2, not 4 + res = ResourceDescriptor.from_opaque_array(arr) + td = TextureObjectOptions(border_color=(0.0, 0.0)) # length 2, not 4 with pytest.raises(ValueError, match="border_color must have 4"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + Device().create_texture_object(resource=res, options=td) finally: arr.close() def test_address_mode_rejects_non_addressmode_scalar(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(address_mode=42) # int, not AddressMode / iterable + res = ResourceDescriptor.from_opaque_array(arr) + td = TextureObjectOptions(address_mode=42) # int, not AddressModeType / iterable with pytest.raises(TypeError, match="address_mode"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + Device().create_texture_object(resource=res, options=td) finally: arr.close() def test_address_mode_rejects_empty_tuple(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(address_mode=()) + res = ResourceDescriptor.from_opaque_array(arr) + td = TextureObjectOptions(address_mode=()) with pytest.raises(ValueError, match="address_mode tuple must have 1-3"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + Device().create_texture_object(resource=res, options=td) finally: arr.close() def test_address_mode_rejects_too_long_tuple(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(address_mode=(AddressMode.WRAP, AddressMode.WRAP, AddressMode.WRAP, AddressMode.WRAP)) + res = ResourceDescriptor.from_opaque_array(arr) + td = TextureObjectOptions( + address_mode=(AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP) + ) with pytest.raises(ValueError, match="address_mode tuple must have 1-3"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + Device().create_texture_object(resource=res, options=td) finally: arr.close() -def test_address_mode_rejects_non_addressmode_entry(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) +def test_address_mode_rejects_invalid_entry(init_cuda): + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: - res = ResourceDescriptor.from_array(arr) - td = TextureDescriptor(address_mode=(AddressMode.WRAP, "bad", AddressMode.CLAMP)) - with pytest.raises(TypeError, match=r"address_mode\[1\]"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + res = ResourceDescriptor.from_opaque_array(arr) + # "bad" is a valid *type* (str is accepted anywhere an AddressModeType + # is) but an invalid *value*, so it is rejected with a ValueError naming + # the offending tuple position. + td = TextureObjectOptions(address_mode=(AddressModeType.WRAP, "bad", AddressModeType.CLAMP)) + with pytest.raises(ValueError, match=r"address_mode\[1\]"): + Device().create_texture_object(resource=res, options=td) finally: arr.close() @@ -923,9 +985,9 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): """Dropping the local references to the backing OpaqueArray and the ResourceDescriptor must NOT invalidate an existing TextureObject. The TextureObject holds a strong ref through its _source_ref slot.""" - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) - res = ResourceDescriptor.from_array(arr) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) + res = ResourceDescriptor.from_opaque_array(arr) + tex = Device().create_texture_object(resource=res, options=TextureObjectOptions()) # Verify the keepalive chain via gc referents: TextureObject -> _source_ref # -> ResourceDescriptor -> _source -> OpaqueArray. We can only walk one level # at a time, so check tex's referents include the ResourceDescriptor. @@ -949,13 +1011,15 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): def test_surface_object_keeps_backing_array_alive(init_cuda): - arr = OpaqueArray.from_descriptor( - shape=(8, 8), - format=ArrayFormat.UINT8, - num_channels=4, - is_surface_load_store=True, + arr = Device().create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + is_surface_load_store=True, + ) ) - surf = SurfaceObject.from_array(arr) + surf = Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(arr)) arr_id = id(arr) del arr gc.collect()