From 01df88efee5eb8fb0dc24378f9200c4d53ffc7b5 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 16:42:39 -0700 Subject: [PATCH 01/11] cuda.core.texture: rename from_array -> from_opaque_array, element_size -> element_bytes (#2292) Addresses items 5 and 7 of the pre-v1.1.0 texture API-style review. - ResourceDescriptor.from_array -> ResourceDescriptor.from_opaque_array: the Array -> OpaqueArray rename removed "array" ambiguity, so the factory name should follow suit. - OpaqueArray.element_size -> OpaqueArray.element_bytes: the accessor returns a byte count, matching the _bytes convention already used by size_bytes and pitch_bytes. Both are renames of APIs not yet released (first ship in v1.1.0), so no deprecation shim is needed; the new names carry a versionadded:: 1.1.0 note. Call sites in the texture tests and the four texture examples are updated. --- cuda_core/cuda/core/texture/_array.pyi | 9 ++++-- cuda_core/cuda/core/texture/_array.pyx | 9 ++++-- cuda_core/cuda/core/texture/_surface.pyx | 2 +- cuda_core/cuda/core/texture/_texture.pyi | 9 ++++-- cuda_core/cuda/core/texture/_texture.pyx | 9 ++++-- cuda_core/examples/gl_interop_fluid.py | 4 +-- cuda_core/examples/gl_interop_mipmap_lod.py | 2 +- cuda_core/examples/texture_sample.py | 2 +- cuda_core/tests/test_texture_surface.py | 36 ++++++++++----------- 9 files changed, 47 insertions(+), 35 deletions(-) diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 75dadb46d98..90422a6383d 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -105,8 +105,11 @@ class OpaqueArray: """Channels per element (1, 2, or 4).""" @property - def element_size(self): - """Bytes per element (format size * channels).""" + def element_bytes(self): + """Bytes per element (format size * channels). + + .. versionadded:: 1.1.0 + """ @property def device(self): @@ -153,7 +156,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): ... diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index 82a08b045dc..d0ef5213ecf 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -349,8 +349,11 @@ cdef class OpaqueArray: return self._num_channels @property - def element_size(self): - """Bytes per element (format size * channels).""" + def element_bytes(self): + """Bytes per element (format size * channels). + + .. versionadded:: 1.1.0 + """ return _FORMAT_ELEM_SIZE[self._format] * self._num_channels @property @@ -411,7 +414,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 diff --git a/cuda_core/cuda/core/texture/_surface.pyx b/cuda_core/cuda/core/texture/_surface.pyx index 93836dfeab8..67fdfaf560c 100644 --- a/cuda_core/cuda/core/texture/_surface.pyx +++ b/cuda_core/cuda/core/texture/_surface.pyx @@ -51,7 +51,7 @@ cdef class SurfaceObject: """ 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)) + return cls.from_descriptor(resource=ResourceDescriptor.from_opaque_array(array)) @classmethod def from_descriptor(cls, *, resource): diff --git a/cuda_core/cuda/core/texture/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index b04ca384989..c386762da68 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -37,7 +37,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). @@ -56,8 +56,11 @@ class ResourceDescriptor: ... @classmethod - def from_array(cls, array): - """Build a resource descriptor backed by a :class:`OpaqueArray`.""" + def from_opaque_array(cls, array): + """Build a resource descriptor backed by a :class:`OpaqueArray`. + + .. versionadded:: 1.1.0 + """ @classmethod def from_mipmapped_array(cls, mipmapped_array): diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index e88ec2d2a36..70434b0aa4f 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -71,7 +71,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). @@ -99,8 +99,11 @@ class ResourceDescriptor: ) @classmethod - def from_array(cls, array): - """Build a resource descriptor backed by a :class:`OpaqueArray`.""" + def from_opaque_array(cls, array): + """Build a resource descriptor backed by a :class:`OpaqueArray`. + + .. versionadded:: 1.1.0 + """ if not isinstance(array, OpaqueArray): raise TypeError(f"array must be a OpaqueArray, got {type(array).__name__}") self = cls.__new__(cls) diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index e2cdb1037e2..30e94a88358 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -414,7 +414,7 @@ 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 +# ResourceDescriptor.from_opaque_array(arr) -> wraps the OpaqueArray as the resource a # TextureObject reads from. # FilterMode.LINEAR -> free HARDWARE bilinear interpolation; # this is what makes semi-Lagrangian @@ -479,7 +479,7 @@ 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) + res_desc = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor( address_mode=AddressMode.CLAMP, filter_mode=FilterMode.LINEAR, diff --git a/cuda_core/examples/gl_interop_mipmap_lod.py b/cuda_core/examples/gl_interop_mipmap_lod.py index ea82ddbd640..1925e30dcf9 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -219,7 +219,7 @@ 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, diff --git a/cuda_core/examples/texture_sample.py b/cuda_core/examples/texture_sample.py index e69655eb479..2687420ec3e 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -87,7 +87,7 @@ def main(): arr.copy_from(pattern, stream=stream) # Build a linear-filtering, clamped, non-normalized texture. - res_desc = ResourceDescriptor.from_array(arr) + res_desc = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor( address_mode=AddressMode.CLAMP, filter_mode=FilterMode.LINEAR, diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 1a6d05e65a2..a98948e28ae 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -49,7 +49,7 @@ def test_array_2d_create_and_properties(init_cuda): assert arr.shape == (32, 16) assert arr.format == ArrayFormat.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 @@ -68,7 +68,7 @@ def test_array_3d_with_surface_flag(init_cuda): 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() @@ -141,7 +141,7 @@ 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) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor( address_mode=AddressMode.CLAMP, filter_mode=FilterMode.LINEAR, @@ -210,7 +210,7 @@ def test_address_mode_normalization(init_cuda): # 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) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor(address_mode=(AddressMode.WRAP, AddressMode.CLAMP)) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=tex_desc) try: @@ -322,7 +322,7 @@ 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, @@ -330,7 +330,7 @@ def test_resource_descriptor_from_pitch2d_validates_pitch(init_cuda): num_channels=1, width=16, height=8, - pitch_bytes=32, # < 64 = width*element_size + pitch_bytes=32, # < 64 = width*element_bytes ) finally: buf.close() @@ -636,7 +636,7 @@ def test_texture_surface_close_is_idempotent(init_cuda): tex_arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) tex = TextureObject.from_descriptor( - resource=ResourceDescriptor.from_array(tex_arr), + resource=ResourceDescriptor.from_opaque_array(tex_arr), texture_descriptor=TextureDescriptor(), ) tex.close() @@ -813,7 +813,7 @@ def test_texture_object_rejects_non_resource_descriptor(init_cuda): def test_texture_object_rejects_non_texture_descriptor(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) with pytest.raises(TypeError, match="texture_descriptor must be a TextureDescriptor"): TextureObject.from_descriptor(resource=res, texture_descriptor="nope") finally: @@ -823,7 +823,7 @@ def test_texture_object_rejects_non_texture_descriptor(init_cuda): 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) + res = ResourceDescriptor.from_opaque_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) @@ -834,7 +834,7 @@ def test_texture_object_rejects_bad_filter_mode(init_cuda): 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) + res = ResourceDescriptor.from_opaque_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) @@ -845,7 +845,7 @@ def test_texture_object_rejects_bad_read_mode(init_cuda): 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) + res = ResourceDescriptor.from_opaque_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) @@ -856,7 +856,7 @@ def test_texture_object_rejects_bad_mipmap_filter_mode(init_cuda): def test_texture_object_rejects_negative_anisotropy(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(max_anisotropy=-1) with pytest.raises(ValueError, match="max_anisotropy"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) @@ -867,7 +867,7 @@ def test_texture_object_rejects_negative_anisotropy(init_cuda): def test_texture_object_rejects_bad_border_color_length(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(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) @@ -878,7 +878,7 @@ def test_texture_object_rejects_bad_border_color_length(init_cuda): def test_address_mode_rejects_non_addressmode_scalar(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(address_mode=42) # int, not AddressMode / iterable with pytest.raises(TypeError, match="address_mode"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) @@ -889,7 +889,7 @@ def test_address_mode_rejects_non_addressmode_scalar(init_cuda): def test_address_mode_rejects_empty_tuple(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(address_mode=()) with pytest.raises(ValueError, match="address_mode tuple must have 1-3"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) @@ -900,7 +900,7 @@ def test_address_mode_rejects_empty_tuple(init_cuda): def test_address_mode_rejects_too_long_tuple(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(address_mode=(AddressMode.WRAP, AddressMode.WRAP, AddressMode.WRAP, AddressMode.WRAP)) with pytest.raises(ValueError, match="address_mode tuple must have 1-3"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) @@ -911,7 +911,7 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): def test_address_mode_rejects_non_addressmode_entry(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) try: - res = ResourceDescriptor.from_array(arr) + res = ResourceDescriptor.from_opaque_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) @@ -924,7 +924,7 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): 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) + res = ResourceDescriptor.from_opaque_array(arr) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) # Verify the keepalive chain via gc referents: TextureObject -> _source_ref # -> ResourceDescriptor -> _source -> OpaqueArray. We can only walk one level From 8b52b02a228de22cbf8686528c915e82a09da112 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 16:54:13 -0700 Subject: [PATCH 02/11] cuda.core.texture: move texture enums to cuda.core.typing as StrEnum (#2292) Addresses item 4 of the pre-v1.1.0 texture API-style review. The texture/surface enums were IntEnum subclasses backed by cydriver.CU_* values, defined in the cuda.core.texture root namespace and lacking the Type suffix. Every other cuda.core enum unified pre-1.0 lives in cuda.core.typing as a StrEnum with semantic string values and a Type suffix. Bring the texture enums in line: - ArrayFormat -> cuda.core.typing.ArrayFormatType - AddressMode -> cuda.core.typing.AddressModeType - FilterMode -> cuda.core.typing.FilterModeType - ReadMode -> cuda.core.typing.ReadModeType Internally the driver integer values are still needed, so _array and _texture carry private CU<->StrEnum bridge maps; OpaqueArray/MipmappedArray keep storing the driver int and convert on the .format boundary. Per the enum convention, plain strings are accepted anywhere an enum is expected (via _normalize_enum / _normalize_array_format), and invalid values raise ValueError. TextureDescriptor gains a __post_init__ that normalizes its filter_mode/read_mode/ mipmap_filter_mode fields at construction. Enums are dropped from the cuda.core.texture __all__ and documented under cuda.core.typing. Tests, examples, and docs are updated to the new names; the negative-path tests that previously expected a TypeError at texture creation now expect a ValueError at descriptor construction. --- cuda_core/cuda/core/texture/__init__.py | 15 +- cuda_core/cuda/core/texture/_array.pyi | 35 ++- cuda_core/cuda/core/texture/_array.pyx | 84 ++++--- .../cuda/core/texture/_mipmapped_array.pyi | 4 +- .../cuda/core/texture/_mipmapped_array.pyx | 19 +- cuda_core/cuda/core/texture/_texture.pyi | 65 ++---- cuda_core/cuda/core/texture/_texture.pyx | 162 ++++++------- cuda_core/cuda/core/typing.py | 69 ++++++ cuda_core/docs/source/api.rst | 13 +- cuda_core/docs/source/api_private.rst | 4 + cuda_core/examples/gl_interop_fluid.py | 30 +-- .../gl_interop_fluid_numba_cuda_mlir.py | 4 +- cuda_core/examples/gl_interop_mipmap_lod.py | 26 ++- cuda_core/examples/texture_sample.py | 18 +- cuda_core/tests/test_texture_surface.py | 213 ++++++++---------- 15 files changed, 414 insertions(+), 347 deletions(-) diff --git a/cuda_core/cuda/core/texture/__init__.py b/cuda_core/cuda/core/texture/__init__.py index e4414ed28ad..6cab3068ae6 100644 --- a/cuda_core/cuda/core/texture/__init__.py +++ b/cuda_core/cuda/core/texture/__init__.py @@ -11,27 +11,26 @@ Import these types from here, e.g.:: from cuda.core.texture import OpaqueArray, TextureObject, TextureDescriptor + +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._array import OpaqueArray from cuda.core.texture._mipmapped_array import MipmappedArray from cuda.core.texture._surface import SurfaceObject from cuda.core.texture._texture import ( - AddressMode, - FilterMode, - ReadMode, ResourceDescriptor, TextureDescriptor, TextureObject, ) __all__ = [ - "AddressMode", - "ArrayFormat", - "FilterMode", "MipmappedArray", "OpaqueArray", - "ReadMode", "ResourceDescriptor", "SurfaceObject", "TextureDescriptor", diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 90422a6383d..b1c6ba1b28b 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -2,24 +2,8 @@ from __future__ import annotations -from enum import IntEnum +from cuda.core.typing import ArrayFormatType -from cuda.bindings import cydriver - - -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 class OpaqueArray: """An opaque, hardware-laid-out GPU allocation for texture/surface access. @@ -64,7 +48,7 @@ class OpaqueArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -98,7 +82,7 @@ class OpaqueArray: @property def format(self): - """The element :class:`ArrayFormat`.""" + """The element :class:`~cuda.core.typing.ArrayFormatType`.""" @property def num_channels(self): @@ -166,11 +150,20 @@ 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: dict +_CU_TO_ARRAYFORMAT: dict +_FORMAT_ELEM_SIZE: dict + +def _normalize_array_format(format): + """Coerce ``format`` to an :class:`ArrayFormatType`. + + Accepts an :class:`ArrayFormatType` or a plain ``str`` naming one of its + values (e.g. ``"float32"``). Raises 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 diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index d0ef5213ecf..36d7821c75d 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -25,44 +25,62 @@ from cuda.core._utils.cuda_utils cimport ( _get_current_device_id, ) -from enum import IntEnum +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. +# Bytes per element (single channel), keyed by the driver CUarray_format int. +_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, +} - 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 +def _normalize_array_format(format): + """Coerce ``format`` to an :class:`ArrayFormatType`. -# Bytes per element (single channel) for each format. -_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, -} + Accepts an :class:`ArrayFormatType` or a plain ``str`` naming one of its + values (e.g. ``"float32"``). Raises on anything else.""" + if isinstance(format, ArrayFormatType): + return format + 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 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): @@ -264,7 +282,7 @@ cdef class OpaqueArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -277,10 +295,10 @@ cdef class OpaqueArray: ------- OpaqueArray """ - _validate_format_channels(format, num_channels) + fmt = _validate_format_channels(format, num_channels) shape_t = _validate_array_shape(shape) - cdef cydriver.CUarray_format c_format = format + cdef cydriver.CUarray_format c_format = _ARRAYFORMAT_TO_CU[fmt] cdef cydriver.CUDA_ARRAY3D_DESCRIPTOR desc3d cdef int rank = len(shape_t) cdef unsigned int flags = ( @@ -340,8 +358,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): @@ -439,7 +457,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})" ) diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index 0741df1781e..9f00b365bdc 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -37,7 +37,7 @@ class MipmappedArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements, for the base (level 0) mip. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -82,7 +82,7 @@ class MipmappedArray: @property def format(self): - """The element :class:`ArrayFormat`.""" + """The element :class:`~cuda.core.typing.ArrayFormatType`.""" @property def num_channels(self): diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyx b/cuda_core/cuda/core/texture/_mipmapped_array.pyx index 0847203fe99..d42d313327e 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, @@ -53,7 +58,7 @@ cdef class MipmappedArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements, for the base (level 0) mip. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -70,14 +75,14 @@ cdef class MipmappedArray: ------- MipmappedArray """ - _validate_format_channels(format, num_channels) + fmt = _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.CUarray_format c_format = _ARRAYFORMAT_TO_CU[fmt] cdef cydriver.CUDA_ARRAY3D_DESCRIPTOR desc3d cdef int rank = len(shape_t) cdef unsigned int flags = ( @@ -153,8 +158,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 +202,7 @@ 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})" ) diff --git a/cuda_core/cuda/core/texture/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index c386762da68..9bdec4a246c 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -3,35 +3,10 @@ from __future__ import annotations from dataclasses import dataclass -from enum import IntEnum -from cuda.bindings import cydriver +from cuda.core.typing import AddressModeType, ArrayFormatType, 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`. @@ -81,7 +56,7 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -105,7 +80,7 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -129,7 +104,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): @@ -160,12 +135,13 @@ class TextureDescriptor: 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 + 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``. - read_mode : ReadMode + read_mode : ReadModeType How sampled integer values are returned. Default ``ELEMENT_TYPE``. normalized_coords : bool If True, coordinates are in ``[0, 1]`` instead of pixel indices. @@ -177,7 +153,7 @@ class TextureDescriptor: If True, enable seamless cubemap edge filtering. max_anisotropy : int Maximum anisotropy; 0 disables anisotropic filtering. - mipmap_filter_mode : FilterMode + mipmap_filter_mode : FilterModeType Filtering between mipmap levels. Default ``POINT``. mipmap_level_bias : float min_mipmap_level_clamp : float @@ -186,20 +162,23 @@ class TextureDescriptor: 4-tuple used when ``address_mode`` includes ``BORDER``; ``None`` means zero. """ - address_mode: AddressMode | tuple[AddressMode, ...] = AddressMode.CLAMP - filter_mode: FilterMode = FilterMode.POINT - read_mode: ReadMode = ReadMode.ELEMENT_TYPE + address_mode: AddressModeType | tuple[AddressModeType, ...] = AddressModeType.CLAMP + filter_mode: FilterModeType = FilterModeType.POINT + read_mode: ReadModeType = 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 = 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) -> None: + ... + class TextureObject: """A bindless texture handle for kernel-side sampled reads. @@ -261,6 +240,12 @@ _TRSF_NORMALIZED_COORDINATES = 2 _TRSF_SRGB = 16 _TRSF_DISABLE_TRILINEAR_OPTIMIZATION = 32 _TRSF_SEAMLESS_CUBEMAP = 64 +_ADDRESSMODE_TO_CU: dict +_FILTERMODE_TO_CU: dict + +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.""" \ 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 70434b0aa4f..dada9e47d89 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -9,7 +9,13 @@ 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, + _normalize_array_format, + _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 +33,9 @@ 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 # Driver texture-descriptor flag bits (CU_TRSF_*). @@ -39,31 +46,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: @@ -151,7 +157,7 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -167,10 +173,11 @@ 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] 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: @@ -186,13 +193,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 @@ -211,7 +218,7 @@ class ResourceDescriptor: buffer : Buffer Device-memory backing. Must remain alive for the lifetime of any :class:`TextureObject` built from this descriptor. - format : ArrayFormat + format : ArrayFormatType or str Element format. num_channels : int Channels per element. Must be 1, 2, or 4. @@ -226,7 +233,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) @@ -235,11 +243,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( @@ -249,7 +257,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 @@ -267,8 +275,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): @@ -317,12 +325,13 @@ class TextureDescriptor: 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 + 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``. - read_mode : ReadMode + read_mode : ReadModeType How sampled integer values are returned. Default ``ELEMENT_TYPE``. normalized_coords : bool If True, coordinates are in ``[0, 1]`` instead of pixel indices. @@ -334,7 +343,7 @@ class TextureDescriptor: If True, enable seamless cubemap edge filtering. max_anisotropy : int Maximum anisotropy; 0 disables anisotropic filtering. - mipmap_filter_mode : FilterMode + mipmap_filter_mode : FilterModeType Filtering between mipmap levels. Default ``POINT``. mipmap_level_bias : float min_mipmap_level_clamp : float @@ -344,40 +353,48 @@ class TextureDescriptor: zero. """ - address_mode: AddressMode | tuple[AddressMode, ...] = AddressMode.CLAMP - filter_mode: FilterMode = FilterMode.POINT - read_mode: ReadMode = ReadMode.ELEMENT_TYPE + address_mode: AddressModeType | tuple[AddressModeType, ...] = AddressModeType.CLAMP + filter_mode: FilterModeType = FilterModeType.POINT + read_mode: ReadModeType = 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 = 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) @@ -462,28 +479,20 @@ cdef class TextureObject: ) # --- Texture descriptor --- + # filter_mode/read_mode/mipmap_filter_mode are normalized to their + # StrEnum types by TextureDescriptor.__post_init__; address_mode is + # normalized (and str-coerced) here. 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 + 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]] - 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__}" - ) + tex_desc.filterMode = _FILTERMODE_TO_CU[texture_descriptor.filter_mode] 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: + # ReadModeType.ELEMENT_TYPE. + if texture_descriptor.read_mode == ReadModeType.ELEMENT_TYPE: flags |= _TRSF_READ_AS_INTEGER if texture_descriptor.normalized_coords: flags |= _TRSF_NORMALIZED_COORDINATES @@ -499,12 +508,7 @@ cdef class TextureObject: 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.mipmapFilterMode = _FILTERMODE_TO_CU[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 diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 278b516217e..0bf5f93d391 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,69 @@ 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. + """ + + 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. + """ + + 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. + """ + + 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. + """ + + 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..aafd965f85a 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -195,13 +195,12 @@ DLPack zero-copy interop. Data is moved in and out only by copying — use TextureDescriptor -.. 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 30e94a88358..95217ca884b 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -108,16 +108,18 @@ launch, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, OpaqueArray, - ReadMode, ResourceDescriptor, SurfaceObject, TextureDescriptor, TextureObject, ) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, +) # --------------------------------------------------------------------------- # Simulation parameters (feel free to change these) @@ -404,7 +406,7 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # OpaqueArray.from_descriptor(...) -> 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. @@ -416,15 +418,15 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # # ResourceDescriptor.from_opaque_array(arr) -> wraps the OpaqueArray as the resource a # TextureObject reads from. -# FilterMode.LINEAR -> free HARDWARE bilinear interpolation; +# 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. @@ -440,7 +442,7 @@ 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, + format=ArrayFormatType.FLOAT32, num_channels=2, is_surface_load_store=True, ) @@ -450,7 +452,7 @@ def make_scalar_array(): """Allocate a single-channel `float` CUDA array (pressure / divergence / dye).""" return OpaqueArray.from_descriptor( shape=(WIDTH, HEIGHT), - format=ArrayFormat.FLOAT32, + format=ArrayFormatType.FLOAT32, num_channels=1, is_surface_load_store=True, ) @@ -466,7 +468,7 @@ def make_color_array(): """ return OpaqueArray.from_descriptor( shape=(WIDTH, HEIGHT), - format=ArrayFormat.FLOAT32, + format=ArrayFormatType.FLOAT32, num_channels=4, is_surface_load_store=True, ) @@ -481,9 +483,9 @@ def make_texture(arr): """ res_desc = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + 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, 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 1925e30dcf9..8e83b0ac5fc 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -101,16 +101,18 @@ launch, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, MipmappedArray, - ReadMode, ResourceDescriptor, SurfaceObject, TextureDescriptor, TextureObject, ) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, +) # --------------------------------------------------------------------------- # Configuration (feel free to change these) @@ -206,9 +208,9 @@ def build_mipmap_pyramid(mip, num_levels, stream, kernels): # 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, + 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): @@ -410,7 +412,7 @@ def main(): num_levels = int(math.log2(BASE_SIZE)) + 1 mip = MipmappedArray.from_descriptor( shape=(BASE_SIZE, BASE_SIZE), - format=ArrayFormat.FLOAT32, + format=ArrayFormatType.FLOAT32, num_channels=4, num_levels=num_levels, is_surface_load_store=True, @@ -424,11 +426,11 @@ def main(): # 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, + 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), diff --git a/cuda_core/examples/texture_sample.py b/cuda_core/examples/texture_sample.py index 2687420ec3e..9cc99918d0d 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -30,15 +30,17 @@ launch, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, OpaqueArray, - ReadMode, ResourceDescriptor, TextureDescriptor, TextureObject, ) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, +) # Kernel reads N (x, y) coordinates from `coords` (interleaved float pairs) and # writes tex2D(tex, x, y) to out[i]. Compiled as C++ so the templated @@ -72,7 +74,7 @@ def main(): width, height = 16, 16 with OpaqueArray.from_descriptor( shape=(width, height), - format=ArrayFormat.FLOAT32, + format=ArrayFormatType.FLOAT32, num_channels=1, ) as arr: # Plant a known pattern: pattern[y, x] = x + 100*y. @@ -89,9 +91,9 @@ def main(): # Build a linear-filtering, clamped, non-normalized texture. res_desc = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + 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: diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index a98948e28ae..0ee691c7dcd 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -10,17 +10,19 @@ Device, ) from cuda.core.texture import ( - AddressMode, - ArrayFormat, - FilterMode, MipmappedArray, OpaqueArray, - ReadMode, ResourceDescriptor, SurfaceObject, TextureDescriptor, TextureObject, ) +from cuda.core.typing import ( + AddressModeType, + ArrayFormatType, + FilterModeType, + ReadModeType, +) def test_array_init_disabled(): @@ -44,10 +46,10 @@ 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 = OpaqueArray.from_descriptor(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_bytes == 4 assert arr.size_bytes == 32 * 16 * 4 @@ -61,7 +63,7 @@ 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, + format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True, ) @@ -75,12 +77,12 @@ def test_array_3d_with_surface_flag(init_cuda): 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) + OpaqueArray.from_descriptor(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) + OpaqueArray.from_descriptor(shape=(2, 2, 2, 2), format=ArrayFormatType.UINT8, num_channels=1) def test_array_roundtrip_copy(init_cuda): @@ -88,7 +90,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 = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1) try: src = _array.array("I", list(range(16))) dst = _array.array("I", [0] * 16) @@ -108,7 +110,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 = OpaqueArray.from_descriptor(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 +126,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 = OpaqueArray.from_descriptor(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,13 +141,13 @@ 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 = OpaqueArray.from_descriptor(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureDescriptor( - address_mode=AddressMode.CLAMP, - filter_mode=FilterMode.LINEAR, - read_mode=ReadMode.ELEMENT_TYPE, + 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) @@ -162,7 +164,7 @@ def test_texture_object_create(init_cuda): def test_surface_object_create(init_cuda): arr = OpaqueArray.from_descriptor( shape=(8, 8), - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True, ) @@ -178,7 +180,7 @@ 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 = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4) try: with pytest.raises(ValueError, match="is_surface_load_store=True"): SurfaceObject.from_array(arr) @@ -191,27 +193,27 @@ 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 = OpaqueArray.from_descriptor(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - tex_desc = TextureDescriptor(address_mode=(AddressMode.WRAP, AddressMode.CLAMP)) + tex_desc = TextureDescriptor(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=tex_desc) try: assert tex.handle != 0 @@ -233,9 +235,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 +250,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 +261,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 +271,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 +286,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 +297,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,7 +309,7 @@ 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) + res = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.FLOAT32, num_channels=1) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) try: assert tex.handle != 0 @@ -326,7 +328,7 @@ def test_resource_descriptor_from_pitch2d_validates_pitch(init_cuda): with pytest.raises(ValueError, match="pitch_bytes"): ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT32, + format=ArrayFormatType.UINT32, num_channels=1, width=16, height=8, @@ -343,7 +345,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 +372,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, @@ -391,13 +393,13 @@ 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) res_p2 = ResourceDescriptor.from_pitch2d( buf, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=4, width=8, height=8, @@ -420,13 +422,13 @@ def test_mipmapped_array_init_disabled(): def test_mipmapped_array_from_descriptor_2d(init_cuda): mip = MipmappedArray.from_descriptor( shape=(64, 32), - format=ArrayFormat.FLOAT32, + 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 @@ -440,7 +442,7 @@ def test_mipmapped_array_get_level_zero_matches_shape(init_cuda): shape = (64, 32) mip = MipmappedArray.from_descriptor( shape=shape, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=4, num_levels=4, ) @@ -450,7 +452,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: @@ -464,7 +466,7 @@ def test_mipmapped_array_get_level_halves_dims(init_cuda): num_levels = 4 mip = MipmappedArray.from_descriptor( shape=shape, - format=ArrayFormat.UINT8, + format=ArrayFormatType.UINT8, num_channels=1, num_levels=num_levels, ) @@ -484,7 +486,7 @@ 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, + format=ArrayFormatType.UINT8, num_channels=1, num_levels=2, ) @@ -501,7 +503,7 @@ 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, + format=ArrayFormatType.UINT8, num_channels=1, num_levels=0, ) @@ -510,7 +512,7 @@ def test_mipmapped_array_rejects_zero_levels(init_cuda): def test_resource_descriptor_from_mipmapped_array(init_cuda): mip = MipmappedArray.from_descriptor( shape=(32, 16), - format=ArrayFormat.FLOAT32, + format=ArrayFormatType.FLOAT32, num_channels=1, num_levels=3, ) @@ -530,7 +532,7 @@ 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, + format=ArrayFormatType.FLOAT32, num_channels=1, num_levels=3, ) @@ -538,10 +540,10 @@ def test_texture_object_from_mipmapped_array(init_cuda): 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, + 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), @@ -559,7 +561,7 @@ 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, + format=ArrayFormatType.UINT8, num_channels=4, num_levels=2, is_surface_load_store=True, @@ -589,7 +591,7 @@ def test_mipmapped_array_level_outlives_dropped_parent(init_cuda): stream = device.create_stream() mip = MipmappedArray.from_descriptor( shape=(16, 16), - format=ArrayFormat.UINT32, + format=ArrayFormatType.UINT32, num_channels=1, num_levels=3, ) @@ -615,18 +617,18 @@ 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 = OpaqueArray.from_descriptor(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 = MipmappedArray.from_descriptor(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 + shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True ) surf = SurfaceObject.from_array(surf_arr) surf.close() @@ -634,7 +636,7 @@ def test_texture_surface_close_is_idempotent(init_cuda): surf.close() surf_arr.close() - tex_arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + tex_arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) tex = TextureObject.from_descriptor( resource=ResourceDescriptor.from_opaque_array(tex_arr), texture_descriptor=TextureDescriptor(), @@ -649,22 +651,22 @@ def test_texture_surface_close_is_idempotent(init_cuda): def test_array_from_descriptor_rejects_bad_format(init_cuda): - with pytest.raises(TypeError, match="format must be an ArrayFormat"): + with pytest.raises(ValueError, match="format must be an ArrayFormatType"): OpaqueArray.from_descriptor(shape=(8,), format=0, num_channels=1) def test_array_from_descriptor_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) + OpaqueArray.from_descriptor(shape=8, format=ArrayFormatType.UINT8, num_channels=1) def test_array_from_descriptor_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) + OpaqueArray.from_descriptor(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 = OpaqueArray.from_descriptor(shape=(8,), format=ArrayFormatType.UINT8, num_channels=1) try: import array as _array @@ -683,7 +685,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 = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormatType.UINT32, num_channels=1) try: dst = _array.array("I", [0] * 16) returned = arr.copy_to(dst, stream=stream) @@ -699,7 +701,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 = OpaqueArray.from_descriptor(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 +724,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 +736,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 +756,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 +773,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 +782,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,18 +793,18 @@ 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"): + with pytest.raises(ValueError, match="format must be an ArrayFormatType"): MipmappedArray.from_descriptor(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) + MipmappedArray.from_descriptor(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) + MipmappedArray.from_descriptor(shape=(0, 8), format=ArrayFormatType.UINT8, num_channels=1, num_levels=1) def test_texture_object_rejects_non_resource_descriptor(init_cuda): @@ -811,7 +813,7 @@ def test_texture_object_rejects_non_resource_descriptor(init_cuda): def test_texture_object_rejects_non_texture_descriptor(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) with pytest.raises(TypeError, match="texture_descriptor must be a TextureDescriptor"): @@ -821,40 +823,23 @@ def test_texture_object_rejects_non_texture_descriptor(init_cuda): 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_opaque_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 TextureDescriptor construction. + with pytest.raises(ValueError, match="filter_mode must be a FilterModeType"): + TextureDescriptor(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_opaque_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"): + TextureDescriptor(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_opaque_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"): + TextureDescriptor(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 = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(max_anisotropy=-1) @@ -865,7 +850,7 @@ def test_texture_object_rejects_negative_anisotropy(init_cuda): 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 = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(border_color=(0.0, 0.0)) # length 2, not 4 @@ -876,10 +861,10 @@ def test_texture_object_rejects_bad_border_color_length(init_cuda): def test_address_mode_rejects_non_addressmode_scalar(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=42) # int, not AddressMode / iterable + td = TextureDescriptor(address_mode=42) # int, not AddressModeType / iterable with pytest.raises(TypeError, match="address_mode"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) finally: @@ -887,7 +872,7 @@ def test_address_mode_rejects_non_addressmode_scalar(init_cuda): def test_address_mode_rejects_empty_tuple(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureDescriptor(address_mode=()) @@ -898,10 +883,10 @@ def test_address_mode_rejects_empty_tuple(init_cuda): def test_address_mode_rejects_too_long_tuple(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=(AddressMode.WRAP, AddressMode.WRAP, AddressMode.WRAP, AddressMode.WRAP)) + td = TextureDescriptor(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) finally: @@ -909,10 +894,10 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): def test_address_mode_rejects_non_addressmode_entry(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormat.FLOAT32, num_channels=1) + arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=(AddressMode.WRAP, "bad", AddressMode.CLAMP)) + td = TextureDescriptor(address_mode=(AddressModeType.WRAP, "bad", AddressModeType.CLAMP)) with pytest.raises(TypeError, match=r"address_mode\[1\]"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) finally: @@ -923,7 +908,7 @@ 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) + arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) res = ResourceDescriptor.from_opaque_array(arr) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) # Verify the keepalive chain via gc referents: TextureObject -> _source_ref @@ -951,7 +936,7 @@ 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, + format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True, ) From b15a9f80fbd344f4fded7aaa1d6e8e7d0ae2b110 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 17:13:39 -0700 Subject: [PATCH 03/11] cuda.core.texture: accept NumPy dtypes for array format (#2292) Addresses item 3 of the pre-v1.1.0 texture API-style review. TensorMapDescriptorOptions.data_type already accepts NumPy dtype objects and keeps its enum as a fallback. Apply the same precedent to the texture format argument: OpaqueArray / MipmappedArray creation and ResourceDescriptor.from_linear / from_pitch2d now accept an ArrayFormatType, a plain str, or a NumPy dtype object (anything numpy.dtype() accepts). Each ArrayFormatType value is already spelled as a NumPy dtype name, so the eight formats map 1:1 via a new _NUMPY_DTYPE_TO_ARRAYFORMAT table. All four factories funnel through _validate_format_channels -> _normalize_array_format, so dtype support is added in one place. ml_dtypes is intentionally not imported: every current format is a standard NumPy dtype, and the ml_dtypes path only becomes relevant if/when a bfloat16 format is added. --- cuda_core/cuda/core/texture/_array.pyi | 6 ++- cuda_core/cuda/core/texture/_array.pyx | 50 ++++++++++++++++--- .../cuda/core/texture/_mipmapped_array.pyi | 5 +- .../cuda/core/texture/_mipmapped_array.pyx | 5 +- cuda_core/cuda/core/texture/_texture.pyi | 10 ++-- cuda_core/cuda/core/texture/_texture.pyx | 10 ++-- cuda_core/tests/test_texture_surface.py | 35 +++++++++++++ 7 files changed, 100 insertions(+), 21 deletions(-) diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index b1c6ba1b28b..6033b41ea0c 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -48,8 +48,9 @@ class OpaqueArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements. - format : ArrayFormatType or str - 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. is_surface_load_store : bool @@ -152,6 +153,7 @@ class OpaqueArray: ... _ARRAYFORMAT_TO_CU: dict _CU_TO_ARRAYFORMAT: dict +_NUMPY_DTYPE_TO_ARRAYFORMAT: dict _FORMAT_ELEM_SIZE: dict def _normalize_array_format(format): diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index 36d7821c75d..122244752e4 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -25,6 +25,8 @@ from cuda.core._utils.cuda_utils cimport ( _get_current_device_id, ) +import numpy + from cuda.core.typing import ArrayFormatType @@ -44,6 +46,15 @@ _ARRAYFORMAT_TO_CU = { _CU_TO_ARRAYFORMAT = {cu: fmt for fmt, cu in _ARRAYFORMAT_TO_CU.items()} +# 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), keyed by the driver CUarray_format int. _FORMAT_ELEM_SIZE = { _ARRAYFORMAT_TO_CU[ArrayFormatType.UINT8]: 1, @@ -60,16 +71,40 @@ _FORMAT_ELEM_SIZE = { def _normalize_array_format(format): """Coerce ``format`` to an :class:`ArrayFormatType`. - Accepts an :class:`ArrayFormatType` or a plain ``str`` naming one of its - values (e.g. ``"float32"``). Raises on anything else.""" + 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 ArrayFormatType(format) - except ValueError as e: + return _NUMPY_DTYPE_TO_ARRAYFORMAT[dt] + except KeyError 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}" + f"NumPy dtype {dt!r} has no ArrayFormatType equivalent; " + f"supported formats: {{{valid}}}" ) from e @@ -282,8 +317,9 @@ cdef class OpaqueArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements. - format : ArrayFormatType or str - 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. is_surface_load_store : bool diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index 9f00b365bdc..957357751f2 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -37,8 +37,9 @@ class MipmappedArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements, for the base (level 0) mip. - format : ArrayFormatType or str - 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. num_levels : int diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyx b/cuda_core/cuda/core/texture/_mipmapped_array.pyx index d42d313327e..fbe17661244 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyx +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyx @@ -58,8 +58,9 @@ cdef class MipmappedArray: shape : tuple of int ``(width,)``, ``(width, height)``, or ``(width, height, depth)`` in elements, for the base (level 0) mip. - format : ArrayFormatType or str - 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. num_levels : int diff --git a/cuda_core/cuda/core/texture/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index 9bdec4a246c..c972175d140 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -56,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 : ArrayFormatType or str - 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 @@ -80,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 : ArrayFormatType or str - 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 diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index dada9e47d89..fa2f0551bd1 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -157,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 : ArrayFormatType or str - 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 @@ -218,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 : ArrayFormatType or str - 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 diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 0ee691c7dcd..63b4597cb2d 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 @@ -75,6 +76,40 @@ def test_array_3d_with_surface_flag(init_cuda): 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 = OpaqueArray.from_descriptor(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 = OpaqueArray.from_descriptor(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"): + OpaqueArray.from_descriptor(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=ArrayFormatType.UINT8, num_channels=3) From 63336a009bd6d70cd2041e0ae47370de80e84139 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 17:16:03 -0700 Subject: [PATCH 04/11] cuda.core.texture: rename TextureDescriptor -> TextureObjectOptions (#2292) Addresses the rename portion of item 2 of the pre-v1.1.0 texture API-style review. CUDA_TEXTURE_DESC is pure creation input (it cannot be queried back from a texture object), so "descriptor" borrowed the driver naming inappropriately. TextureObjectOptions matches the XxxOptions convention used throughout cuda.core (StreamOptions, EventOptions, TensorMapDescriptorOptions, ...). Pure rename of the dataclass and all references across the implementation, stubs, __init__ __all__, docs, tests, and examples. The from_descriptor texture_descriptor= keyword is left untouched here; the whole creation entry point is replaced by Device.create_texture_object(*, resource, options) in a later commit. --- cuda_core/cuda/core/texture/__init__.py | 6 ++-- cuda_core/cuda/core/texture/_texture.pxd | 2 +- cuda_core/cuda/core/texture/_texture.pyi | 8 ++--- cuda_core/cuda/core/texture/_texture.pyx | 14 ++++---- cuda_core/docs/source/api.rst | 2 +- cuda_core/examples/gl_interop_fluid.py | 6 ++-- cuda_core/examples/gl_interop_mipmap_lod.py | 6 ++-- cuda_core/examples/texture_sample.py | 4 +-- cuda_core/tests/test_texture_surface.py | 40 ++++++++++----------- 9 files changed, 44 insertions(+), 44 deletions(-) diff --git a/cuda_core/cuda/core/texture/__init__.py b/cuda_core/cuda/core/texture/__init__.py index 6cab3068ae6..fd31abd0aea 100644 --- a/cuda_core/cuda/core/texture/__init__.py +++ b/cuda_core/cuda/core/texture/__init__.py @@ -10,7 +10,7 @@ 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`, @@ -24,8 +24,8 @@ from cuda.core.texture._surface import SurfaceObject from cuda.core.texture._texture import ( ResourceDescriptor, - TextureDescriptor, TextureObject, + TextureObjectOptions, ) __all__ = [ @@ -33,6 +33,6 @@ "OpaqueArray", "ResourceDescriptor", "SurfaceObject", - "TextureDescriptor", "TextureObject", + "TextureObjectOptions", ] diff --git a/cuda_core/cuda/core/texture/_texture.pxd b/cuda_core/cuda/core/texture/_texture.pxd index f8532a4dad5..fcac323af91 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 _texture_desc # 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 c972175d140..db654b11c87 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -68,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. """ @@ -132,7 +132,7 @@ class ResourceDescriptor: ... @dataclass -class TextureDescriptor: +class TextureObjectOptions: """Sampling state for a :class:`TextureObject` (mirrors ``CUDA_TEXTURE_DESC``). Attributes @@ -210,7 +210,7 @@ class TextureObject: Parameters ---------- resource : ResourceDescriptor - texture_descriptor : TextureDescriptor + texture_descriptor : TextureObjectOptions """ @property @@ -223,7 +223,7 @@ class TextureObject: @property def texture_descriptor(self): - """The :class:`TextureDescriptor` this texture was built from.""" + """The :class:`TextureObjectOptions` this texture was built from.""" @property def device(self): diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index fa2f0551bd1..83a28e387d6 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -169,7 +169,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. """ if not isinstance(buffer, Buffer): @@ -322,7 +322,7 @@ class ResourceDescriptor: @dataclass -class TextureDescriptor: +class TextureObjectOptions: """Sampling state for a :class:`TextureObject` (mirrors ``CUDA_TEXTURE_DESC``). Attributes @@ -426,16 +426,16 @@ cdef class TextureObject: Parameters ---------- resource : ResourceDescriptor - texture_descriptor : TextureDescriptor + texture_descriptor : TextureObjectOptions """ if not isinstance(resource, ResourceDescriptor): raise TypeError( f"resource must be a ResourceDescriptor, got " f"{type(resource).__name__}" ) - if not isinstance(texture_descriptor, TextureDescriptor): + if not isinstance(texture_descriptor, TextureObjectOptions): raise TypeError( - f"texture_descriptor must be a TextureDescriptor, got " + f"texture_descriptor must be a TextureObjectOptions, got " f"{type(texture_descriptor).__name__}" ) @@ -482,7 +482,7 @@ cdef class TextureObject: # --- Texture descriptor --- # filter_mode/read_mode/mipmap_filter_mode are normalized to their - # StrEnum types by TextureDescriptor.__post_init__; address_mode is + # StrEnum types by TextureObjectOptions.__post_init__; address_mode is # normalized (and str-coerced) here. modes = _normalize_address_modes(texture_descriptor.address_mode) tex_desc.addressMode[0] = _ADDRESSMODE_TO_CU[modes[0]] @@ -557,7 +557,7 @@ cdef class TextureObject: @property def texture_descriptor(self): - """The :class:`TextureDescriptor` this texture was built from.""" + """The :class:`TextureObjectOptions` this texture was built from.""" return self._texture_desc @property diff --git a/cuda_core/docs/source/api.rst b/cuda_core/docs/source/api.rst index aafd965f85a..dc191a58186 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -193,7 +193,7 @@ DLPack zero-copy interop. Data is moved in and out only by copying — use :template: dataclass.rst - TextureDescriptor + TextureObjectOptions The associated enumerations — :class:`~cuda.core.typing.ArrayFormatType`, diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index 95217ca884b..e3f3cb7a7a5 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -111,8 +111,8 @@ OpaqueArray, ResourceDescriptor, SurfaceObject, - TextureDescriptor, TextureObject, + TextureObjectOptions, ) from cuda.core.typing import ( AddressModeType, @@ -400,7 +400,7 @@ 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 @@ -482,7 +482,7 @@ def make_texture(arr): gradient) sample exactly at texel centers so LINEAR returns the exact value. """ res_desc = ResourceDescriptor.from_opaque_array(arr) - tex_desc = TextureDescriptor( + tex_desc = TextureObjectOptions( address_mode=AddressModeType.CLAMP, filter_mode=FilterModeType.LINEAR, read_mode=ReadModeType.ELEMENT_TYPE, diff --git a/cuda_core/examples/gl_interop_mipmap_lod.py b/cuda_core/examples/gl_interop_mipmap_lod.py index 8e83b0ac5fc..b7477dad1f9 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -104,8 +104,8 @@ MipmappedArray, ResourceDescriptor, SurfaceObject, - TextureDescriptor, TextureObject, + TextureObjectOptions, ) from cuda.core.typing import ( AddressModeType, @@ -207,7 +207,7 @@ 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( + src_tex_desc = TextureObjectOptions( address_mode=AddressModeType.CLAMP, filter_mode=FilterModeType.POINT, # explicit per-texel reads read_mode=ReadModeType.ELEMENT_TYPE, @@ -425,7 +425,7 @@ 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( + display_tex_desc = TextureObjectOptions( address_mode=AddressModeType.WRAP, filter_mode=FilterModeType.LINEAR, read_mode=ReadModeType.ELEMENT_TYPE, diff --git a/cuda_core/examples/texture_sample.py b/cuda_core/examples/texture_sample.py index 9cc99918d0d..cbfbd3a0585 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -32,8 +32,8 @@ from cuda.core.texture import ( OpaqueArray, ResourceDescriptor, - TextureDescriptor, TextureObject, + TextureObjectOptions, ) from cuda.core.typing import ( AddressModeType, @@ -90,7 +90,7 @@ def main(): # Build a linear-filtering, clamped, non-normalized texture. res_desc = ResourceDescriptor.from_opaque_array(arr) - tex_desc = TextureDescriptor( + tex_desc = TextureObjectOptions( address_mode=AddressModeType.CLAMP, filter_mode=FilterModeType.LINEAR, read_mode=ReadModeType.ELEMENT_TYPE, diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 63b4597cb2d..1bf51ac660a 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -15,8 +15,8 @@ OpaqueArray, ResourceDescriptor, SurfaceObject, - TextureDescriptor, TextureObject, + TextureObjectOptions, ) from cuda.core.typing import ( AddressModeType, @@ -179,7 +179,7 @@ def test_texture_object_create(init_cuda): arr = OpaqueArray.from_descriptor(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - tex_desc = TextureDescriptor( + tex_desc = TextureObjectOptions( address_mode=AddressModeType.CLAMP, filter_mode=FilterModeType.LINEAR, read_mode=ReadModeType.ELEMENT_TYPE, @@ -248,7 +248,7 @@ def test_address_mode_normalization(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - tex_desc = TextureDescriptor(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) + tex_desc = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=tex_desc) try: assert tex.handle != 0 @@ -345,7 +345,7 @@ def test_texture_object_from_linear(init_cuda): buf = _alloc_device_buffer(device, 1024 * 4) try: res = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.FLOAT32, num_channels=1) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) + tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureObjectOptions()) try: assert tex.handle != 0 assert tex.resource is res @@ -415,7 +415,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 = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureObjectOptions()) try: assert tex.handle != 0 finally: @@ -574,7 +574,7 @@ def test_texture_object_from_mipmapped_array(init_cuda): try: res = ResourceDescriptor.from_mipmapped_array(mip) # Use non-default mipmap params so the driver exercises that path. - tex_desc = TextureDescriptor( + tex_desc = TextureObjectOptions( address_mode=AddressModeType.CLAMP, filter_mode=FilterModeType.LINEAR, normalized_coords=True, @@ -674,7 +674,7 @@ def test_texture_surface_close_is_idempotent(init_cuda): tex_arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) tex = TextureObject.from_descriptor( resource=ResourceDescriptor.from_opaque_array(tex_arr), - texture_descriptor=TextureDescriptor(), + texture_descriptor=TextureObjectOptions(), ) tex.close() assert tex.handle == 0 @@ -844,40 +844,40 @@ def test_mipmapped_array_rejects_zero_dim(init_cuda): 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()) + TextureObject.from_descriptor(resource=object(), texture_descriptor=TextureObjectOptions()) def test_texture_object_rejects_non_texture_descriptor(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - with pytest.raises(TypeError, match="texture_descriptor must be a TextureDescriptor"): + with pytest.raises(TypeError, match="texture_descriptor must be a TextureObjectOptions"): TextureObject.from_descriptor(resource=res, texture_descriptor="nope") finally: arr.close() def test_texture_object_rejects_bad_filter_mode(init_cuda): - # Invalid enum values are rejected at TextureDescriptor construction. + # Invalid enum values are rejected at TextureObjectOptions construction. with pytest.raises(ValueError, match="filter_mode must be a FilterModeType"): - TextureDescriptor(filter_mode=0) # int, not FilterModeType + TextureObjectOptions(filter_mode=0) # int, not FilterModeType def test_texture_object_rejects_bad_read_mode(init_cuda): with pytest.raises(ValueError, match="read_mode must be a ReadModeType"): - TextureDescriptor(read_mode=0) # int, not ReadModeType + TextureObjectOptions(read_mode=0) # int, not ReadModeType def test_texture_object_rejects_bad_mipmap_filter_mode(init_cuda): with pytest.raises(ValueError, match="mipmap_filter_mode must be a FilterModeType"): - TextureDescriptor(mipmap_filter_mode=0) # int, not 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=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(max_anisotropy=-1) + td = TextureObjectOptions(max_anisotropy=-1) with pytest.raises(ValueError, match="max_anisotropy"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) finally: @@ -888,7 +888,7 @@ def test_texture_object_rejects_bad_border_color_length(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(border_color=(0.0, 0.0)) # length 2, not 4 + 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) finally: @@ -899,7 +899,7 @@ def test_address_mode_rejects_non_addressmode_scalar(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=42) # int, not AddressModeType / iterable + td = TextureObjectOptions(address_mode=42) # int, not AddressModeType / iterable with pytest.raises(TypeError, match="address_mode"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) finally: @@ -910,7 +910,7 @@ def test_address_mode_rejects_empty_tuple(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=()) + td = TextureObjectOptions(address_mode=()) with pytest.raises(ValueError, match="address_mode tuple must have 1-3"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) finally: @@ -921,7 +921,7 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=(AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP)) + 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) finally: @@ -932,7 +932,7 @@ def test_address_mode_rejects_non_addressmode_entry(init_cuda): arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureDescriptor(address_mode=(AddressModeType.WRAP, "bad", AddressModeType.CLAMP)) + td = TextureObjectOptions(address_mode=(AddressModeType.WRAP, "bad", AddressModeType.CLAMP)) with pytest.raises(TypeError, match=r"address_mode\[1\]"): TextureObject.from_descriptor(resource=res, texture_descriptor=td) finally: @@ -945,7 +945,7 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): TextureObject holds a strong ref through its _source_ref slot.""" arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) res = ResourceDescriptor.from_opaque_array(arr) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureDescriptor()) + tex = TextureObject.from_descriptor(resource=res, texture_descriptor=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. From c373a6d7d9dc20f3f1011be9814fb941db078390 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 17:28:42 -0700 Subject: [PATCH 05/11] cuda.core.texture: add Device.create_opaque_array / create_mipmapped_array (#2292) Addresses item 1 (array allocation) and item 2 (Options dataclasses) of the pre-v1.1.0 texture API-style review. Every other user-allocatable resource in cuda.core is created via Device.create_* (create_stream, create_event, ...) with an XxxOptions dataclass. The texture arrays used per-type classmethod factories (OpaqueArray.from_descriptor, MipmappedArray.from_descriptor) with loose kwargs instead. Move them in line: - add OpaqueArrayOptions and MipmappedArrayOptions dataclasses (validated in __post_init__, mirroring TensorMapDescriptorOptions); - add Device.create_opaque_array(options) and Device.create_mipmapped_array( options), which gate on _check_context_initialized() and delegate to private module factories that bind to the current device (same pattern as create_event / create_graph_builder); - remove the OpaqueArray.from_descriptor / MipmappedArray.from_descriptor classmethods; the private _from_handle graphics-interop path is retained. New entry points carry versionadded:: 1.1.0. Options are exported from cuda.core.texture and documented in api.rst. Tests and the texture examples are updated to Device().create_*_array((...)). --- cuda_core/cuda/core/_device.pyi | 18 ++ cuda_core/cuda/core/_device.pyx | 56 ++++++ cuda_core/cuda/core/texture/__init__.py | 6 +- cuda_core/cuda/core/texture/_array.pyi | 67 +++++--- cuda_core/cuda/core/texture/_array.pyx | 142 +++++++++------- .../cuda/core/texture/_mipmapped_array.pyi | 68 ++++---- .../cuda/core/texture/_mipmapped_array.pyx | 160 ++++++++++-------- cuda_core/docs/source/api.rst | 11 +- cuda_core/examples/gl_interop_fluid.py | 16 +- cuda_core/examples/gl_interop_mipmap_lod.py | 6 +- cuda_core/examples/texture_sample.py | 8 +- cuda_core/tests/test_texture_surface.py | 118 ++++++------- 12 files changed, 409 insertions(+), 267 deletions(-) diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index 70e28139ba5..71ebb2c5c72 100644 --- a/cuda_core/cuda/core/_device.pyi +++ b/cuda_core/cuda/core/_device.pyi @@ -12,6 +12,12 @@ 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, +) class DeviceProperties: @@ -909,5 +915,17 @@ class Device: Newly created graph builder object. """ + + def create_opaque_array(self, options: OpaqueArrayOptions | None = None) -> OpaqueArray: + """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. + + .. versionadded:: 1.1.0 + """ + + def create_mipmapped_array(self, options: MipmappedArrayOptions | None = None) -> MipmappedArray: + """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. + + .. 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 a90f246d381..6b8e05fe651 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -1456,6 +1456,62 @@ class Device: self._check_context_initialized() return GraphBuilder._init(self.create_stream()) + def create_opaque_array(self, options: object = None): + """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. + + Allocates an opaque, hardware-laid-out CUDA array for texture/surface + access. The array is bound to the device/context that is current at + construction. + + 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: object = None): + """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. + + Allocates a mipmapped CUDA array for texture/surface access across + levels. The array is bound to the device/context that is current at + construction. + + 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) + 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 fd31abd0aea..cd3e355f324 100644 --- a/cuda_core/cuda/core/texture/__init__.py +++ b/cuda_core/cuda/core/texture/__init__.py @@ -19,8 +19,8 @@ alongside the other ``cuda.core`` enumerations. """ -from cuda.core.texture._array import 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 ( ResourceDescriptor, @@ -30,7 +30,9 @@ __all__ = [ "MipmappedArray", + "MipmappedArrayOptions", "OpaqueArray", + "OpaqueArrayOptions", "ResourceDescriptor", "SurfaceObject", "TextureObject", diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 6033b41ea0c..02ecafe4b12 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -2,9 +2,39 @@ from __future__ import annotations +from dataclasses import dataclass + 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. + """ + shape: tuple[int, ...] + format: object + num_channels: int + is_surface_load_store: bool = False + + def __post_init__(self) -> None: + ... + + class OpaqueArray: """An opaque, hardware-laid-out GPU allocation for texture/surface access. @@ -22,9 +52,9 @@ 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. """ def close(self): @@ -39,30 +69,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 : 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:`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``. @@ -169,4 +175,9 @@ def _validate_format_channels(format, num_channels): 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`.""" \ 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 122244752e4..b2f0499c684 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -27,6 +27,9 @@ from cuda.core._utils.cuda_utils cimport ( import numpy +from dataclasses import dataclass + +from cuda.core._utils.cuda_utils import check_or_create_options from cuda.core.typing import ArrayFormatType @@ -133,6 +136,36 @@ 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. + """ + + 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: @@ -298,72 +331,17 @@ 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. """ def __init__(self, *args, **kwargs): raise RuntimeError( - "OpaqueArray cannot be instantiated directly. Use OpaqueArray.from_descriptor()." - ) - - @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 : 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:`SurfaceObject` for kernel-side writes. - Default False. - - Returns - ------- - OpaqueArray - """ - fmt = _validate_format_channels(format, num_channels) - shape_t = _validate_array_shape(shape) - - cdef cydriver.CUarray_format c_format = _ARRAYFORMAT_TO_CU[fmt] - 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 + "OpaqueArray cannot be instantiated directly. " + "Use Device.create_opaque_array()." ) - # 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``. @@ -527,3 +505,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 957357751f2..bb9cf6db9ff 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -2,6 +2,37 @@ 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. + is_surface_load_store : bool + If True, allocate with ``CUDA_ARRAY3D_SURFACE_LDST``. Default False. + """ + shape: tuple[int, ...] + format: object + num_channels: int + num_levels: int + is_surface_load_store: bool = False + + def __post_init__(self) -> None: + ... + class MipmappedArray: """A mipmapped CUDA array for texture/surface access across levels. @@ -13,7 +44,7 @@ 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`. """ def close(self): @@ -28,34 +59,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 : 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:`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. @@ -109,4 +112,9 @@ 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`.""" \ 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 fbe17661244..4338b1fa2a7 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyx +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyx @@ -27,6 +27,49 @@ 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. + """ + + 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. @@ -38,82 +81,14 @@ 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`. """ 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 : 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:`get_level`) can be bound as - :class:`SurfaceObject` for kernel-side writes. Default False. - - Returns - ------- - MipmappedArray - """ - fmt = _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 = _ARRAYFORMAT_TO_CU[fmt] - 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. @@ -207,3 +182,48 @@ cdef class MipmappedArray: 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/docs/source/api.rst b/cuda_core/docs/source/api.rst index dc191a58186..c96f6b77c90 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -167,10 +167,11 @@ 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. A :class:`OpaqueArray` has an opaque, hardware-defined layout with no linear device pointer, so it cannot participate in ``__cuda_array_interface__`` / @@ -193,6 +194,8 @@ DLPack zero-copy interop. Data is moved in and out only by copying — use :template: dataclass.rst + OpaqueArrayOptions + MipmappedArrayOptions TextureObjectOptions The associated enumerations — diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index e3f3cb7a7a5..f2cb80d346e 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -108,7 +108,7 @@ launch, ) from cuda.core.texture import ( - OpaqueArray, + OpaqueArrayOptions, ResourceDescriptor, SurfaceObject, TextureObject, @@ -403,7 +403,7 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # 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(OpaqueArrayOptions(...)) -> allocates a CUDA *array* (opaque, tiled # layout optimized for 2D texture fetches), # not linear device memory. # ArrayFormatType.FLOAT32 -> each channel is a 32-bit float texel. @@ -440,22 +440,22 @@ 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( + 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( + return Device().create_opaque_array(OpaqueArrayOptions( shape=(WIDTH, HEIGHT), format=ArrayFormatType.FLOAT32, num_channels=1, is_surface_load_store=True, - ) + )) def make_color_array(): @@ -466,12 +466,12 @@ 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( + return Device().create_opaque_array(OpaqueArrayOptions( shape=(WIDTH, HEIGHT), format=ArrayFormatType.FLOAT32, num_channels=4, is_surface_load_store=True, - ) + )) def make_texture(arr): diff --git a/cuda_core/examples/gl_interop_mipmap_lod.py b/cuda_core/examples/gl_interop_mipmap_lod.py index b7477dad1f9..26877b093f6 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -101,7 +101,7 @@ launch, ) from cuda.core.texture import ( - MipmappedArray, + MipmappedArrayOptions, ResourceDescriptor, SurfaceObject, TextureObject, @@ -410,13 +410,13 @@ 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( + 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) # --- Step 3: Bind the WHOLE pyramid as a trilinear-filtered texture --- diff --git a/cuda_core/examples/texture_sample.py b/cuda_core/examples/texture_sample.py index cbfbd3a0585..0df09c23a3e 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -30,7 +30,7 @@ launch, ) from cuda.core.texture import ( - OpaqueArray, + OpaqueArrayOptions, ResourceDescriptor, TextureObject, TextureObjectOptions, @@ -68,15 +68,15 @@ 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( + with Device().create_opaque_array(OpaqueArrayOptions( shape=(width, height), format=ArrayFormatType.FLOAT32, num_channels=1, - ) as arr: + )) as arr: # Plant a known pattern: pattern[y, x] = x + 100*y. # Cast to float32 so the byte count matches the array's storage. ys, xs = np.meshgrid( diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 1bf51ac660a..10e12edc059 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -12,7 +12,9 @@ ) from cuda.core.texture import ( MipmappedArray, + MipmappedArrayOptions, OpaqueArray, + OpaqueArrayOptions, ResourceDescriptor, SurfaceObject, TextureObject, @@ -47,7 +49,7 @@ def test_resource_descriptor_init_disabled(): def test_array_2d_create_and_properties(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(32, 16), format=ArrayFormatType.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 == ArrayFormatType.FLOAT32 @@ -62,12 +64,12 @@ def test_array_2d_create_and_properties(init_cuda): def test_array_3d_with_surface_flag(init_cuda): - arr = OpaqueArray.from_descriptor( + 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 @@ -87,7 +89,7 @@ def test_array_3d_with_surface_flag(init_cuda): ], ) def test_array_accepts_numpy_dtype_format(init_cuda, dtype, expected): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=dtype, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=dtype, num_channels=1)) try: assert arr.format == expected finally: @@ -96,7 +98,7 @@ def test_array_accepts_numpy_dtype_format(init_cuda, dtype, expected): @pytest.mark.agent_authored(model="claude-opus-4.8") def test_array_accepts_str_format(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format="float32", num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format="float32", num_channels=1)) try: assert arr.format == ArrayFormatType.FLOAT32 finally: @@ -107,17 +109,17 @@ def test_array_accepts_str_format(init_cuda): def test_array_rejects_unsupported_dtype_format(init_cuda): # float64 has no ArrayFormatType equivalent. with pytest.raises(ValueError, match="no ArrayFormatType equivalent"): - OpaqueArray.from_descriptor(shape=(8,), format=np.float64, num_channels=1) + 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=ArrayFormatType.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=ArrayFormatType.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): @@ -125,7 +127,7 @@ def test_array_roundtrip_copy(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormatType.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) @@ -145,7 +147,7 @@ def test_array_copy_rejects_undersized_host_buffer(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormatType.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) @@ -161,7 +163,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=ArrayFormatType.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: @@ -176,7 +178,7 @@ 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=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureObjectOptions( @@ -197,12 +199,12 @@ def test_texture_object_create(init_cuda): def test_surface_object_create(init_cuda): - arr = OpaqueArray.from_descriptor( + 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) try: @@ -215,7 +217,7 @@ def test_surface_object_create(init_cuda): def test_surface_requires_ldst_flag(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.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) @@ -245,7 +247,7 @@ def test_address_mode_normalization(init_cuda): ) # Smoke test: a 2-entry tuple is also accepted end-to-end. - arr = OpaqueArray.from_descriptor(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) @@ -455,12 +457,12 @@ def test_mipmapped_array_init_disabled(): def test_mipmapped_array_from_descriptor_2d(init_cuda): - mip = MipmappedArray.from_descriptor( + 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 == ArrayFormatType.FLOAT32 @@ -475,12 +477,12 @@ 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( + mip = Device().create_mipmapped_array(MipmappedArrayOptions( shape=shape, format=ArrayFormatType.UINT8, num_channels=4, num_levels=4, - ) + )) try: lvl0 = mip.get_level(0) try: @@ -499,12 +501,12 @@ 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( + 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): lvl = mip.get_level(level) @@ -519,12 +521,12 @@ 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( + 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"): mip.get_level(mip.num_levels) @@ -536,21 +538,21 @@ 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( + 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( + 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) assert res.kind == "mipmapped_array" @@ -565,12 +567,12 @@ def test_resource_descriptor_from_mipmapped_array_rejects_non_mipmap(): def test_texture_object_from_mipmapped_array(init_cuda): - mip = MipmappedArray.from_descriptor( + 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. @@ -594,13 +596,13 @@ def test_texture_object_from_mipmapped_array(init_cuda): def test_surface_rejects_mipmapped_array(init_cuda): - mip = MipmappedArray.from_descriptor( + 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"): @@ -624,12 +626,12 @@ def test_mipmapped_array_level_outlives_dropped_parent(init_cuda): device = Device() stream = device.create_stream() - mip = MipmappedArray.from_descriptor( + 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. del mip @@ -652,26 +654,26 @@ 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=ArrayFormatType.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=ArrayFormatType.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( + 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.close() assert surf.handle == 0 surf.close() surf_arr.close() - tex_arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + tex_arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) tex = TextureObject.from_descriptor( resource=ResourceDescriptor.from_opaque_array(tex_arr), texture_descriptor=TextureObjectOptions(), @@ -687,21 +689,21 @@ def test_texture_surface_close_is_idempotent(init_cuda): def test_array_from_descriptor_rejects_bad_format(init_cuda): with pytest.raises(ValueError, match="format must be an ArrayFormatType"): - OpaqueArray.from_descriptor(shape=(8,), format=0, num_channels=1) + Device().create_opaque_array(OpaqueArrayOptions(shape=(8,), format=0, num_channels=1)) def test_array_from_descriptor_rejects_non_iterable_shape(init_cuda): with pytest.raises(TypeError, match="shape must be a tuple"): - OpaqueArray.from_descriptor(shape=8, format=ArrayFormatType.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): with pytest.raises(ValueError, match=r"shape\[1\] must be >= 1"): - OpaqueArray.from_descriptor(shape=(8, 0), format=ArrayFormatType.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=ArrayFormatType.UINT8, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8,), format=ArrayFormatType.UINT8, num_channels=1)) try: import array as _array @@ -720,7 +722,7 @@ def test_array_copy_to_returns_dst(init_cuda): device = Device() stream = device.create_stream() - arr = OpaqueArray.from_descriptor(shape=(16,), format=ArrayFormatType.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) @@ -736,7 +738,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=ArrayFormatType.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() @@ -829,17 +831,17 @@ def test_resource_descriptor_from_pitch2d_rejects_zero_dims(init_cuda): def test_mipmapped_array_rejects_bad_format(init_cuda): with pytest.raises(ValueError, match="format must be an ArrayFormatType"): - MipmappedArray.from_descriptor(shape=(8, 8), format=0, num_channels=1, num_levels=2) + 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=ArrayFormatType.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=ArrayFormatType.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): @@ -848,7 +850,7 @@ def test_texture_object_rejects_non_resource_descriptor(init_cuda): def test_texture_object_rejects_non_texture_descriptor(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) with pytest.raises(TypeError, match="texture_descriptor must be a TextureObjectOptions"): @@ -874,7 +876,7 @@ def test_texture_object_rejects_bad_mipmap_filter_mode(init_cuda): def test_texture_object_rejects_negative_anisotropy(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(max_anisotropy=-1) @@ -885,7 +887,7 @@ def test_texture_object_rejects_negative_anisotropy(init_cuda): def test_texture_object_rejects_bad_border_color_length(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(border_color=(0.0, 0.0)) # length 2, not 4 @@ -896,7 +898,7 @@ def test_texture_object_rejects_bad_border_color_length(init_cuda): def test_address_mode_rejects_non_addressmode_scalar(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(address_mode=42) # int, not AddressModeType / iterable @@ -907,7 +909,7 @@ def test_address_mode_rejects_non_addressmode_scalar(init_cuda): def test_address_mode_rejects_empty_tuple(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(address_mode=()) @@ -918,7 +920,7 @@ def test_address_mode_rejects_empty_tuple(init_cuda): def test_address_mode_rejects_too_long_tuple(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP)) @@ -929,7 +931,7 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): def test_address_mode_rejects_non_addressmode_entry(init_cuda): - arr = OpaqueArray.from_descriptor(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(address_mode=(AddressModeType.WRAP, "bad", AddressModeType.CLAMP)) @@ -943,7 +945,7 @@ 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=ArrayFormatType.FLOAT32, num_channels=1) + arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) res = ResourceDescriptor.from_opaque_array(arr) tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureObjectOptions()) # Verify the keepalive chain via gc referents: TextureObject -> _source_ref @@ -969,12 +971,12 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): def test_surface_object_keeps_backing_array_alive(init_cuda): - arr = OpaqueArray.from_descriptor( + 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) arr_id = id(arr) del arr From b28a1efb0158ad756073d87394a7b3d31e5606af Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 17:36:11 -0700 Subject: [PATCH 06/11] cuda.core.texture: add Device.create_texture_object / create_surface_object (#2292) Addresses item 1 (object creation) and item 6 (drop SurfaceObject.from_array sugar) of the pre-v1.1.0 texture API-style review, completing the move to the Device.create_* allocation shape. - add Device.create_texture_object(*, resource, options) and Device.create_surface_object(*, resource), gating on _check_context_initialized() and delegating to private module factories that bind to the current device (same pattern as create_opaque_array); - remove TextureObject.from_descriptor, SurfaceObject.from_descriptor, and the speculative SurfaceObject.from_array sugar (a surface is now built from an explicit ResourceDescriptor.from_opaque_array(...) like every other backing); - rename the sampling-options argument from texture_descriptor to options on the new entry point, matching the create_* + XxxOptions convention (the queryable TextureObject.texture_descriptor property is unchanged). New entry points carry versionadded:: 1.1.0. Tests, examples, and api.rst are updated; a make_surface() helper mirrors make_texture() in the fluid example. --- cuda_core/cuda/core/_device.pyi | 16 ++ cuda_core/cuda/core/_device.pyx | 64 +++++ cuda_core/cuda/core/texture/_surface.pyi | 28 +-- cuda_core/cuda/core/texture/_surface.pyx | 100 ++++---- cuda_core/cuda/core/texture/_texture.pyi | 21 +- cuda_core/cuda/core/texture/_texture.pyx | 257 ++++++++++---------- cuda_core/docs/source/api.rst | 4 +- cuda_core/examples/gl_interop_fluid.py | 28 ++- cuda_core/examples/gl_interop_mipmap_lod.py | 16 +- cuda_core/examples/texture_sample.py | 3 +- cuda_core/tests/test_texture_surface.py | 46 ++-- 11 files changed, 319 insertions(+), 264 deletions(-) diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index 71ebb2c5c72..6d202ee5923 100644 --- a/cuda_core/cuda/core/_device.pyi +++ b/cuda_core/cuda/core/_device.pyi @@ -17,6 +17,10 @@ from cuda.core.texture import ( MipmappedArrayOptions, OpaqueArray, OpaqueArrayOptions, + ResourceDescriptor, + SurfaceObject, + TextureObject, + TextureObjectOptions, ) @@ -925,6 +929,18 @@ class Device: def create_mipmapped_array(self, options: MipmappedArrayOptions | None = None) -> MipmappedArray: """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. + .. versionadded:: 1.1.0 + """ + + def create_texture_object(self, *, resource: ResourceDescriptor, options: TextureObjectOptions | None = None) -> TextureObject: + """Create a :obj:`~cuda.core.texture.TextureObject` on this device. + + .. versionadded:: 1.1.0 + """ + + def create_surface_object(self, *, resource: ResourceDescriptor) -> SurfaceObject: + """Create a :obj:`~cuda.core.texture.SurfaceObject` on this device. + .. versionadded:: 1.1.0 """ _tls = threading.local() diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index 6b8e05fe651..5cecf47624b 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -1512,6 +1512,70 @@ class Device: self._check_context_initialized() return _create_mipmapped_array(options) + def create_texture_object(self, *, resource, options: object = None): + """Create a :obj:`~cuda.core.texture.TextureObject` on this 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 bound to the device/context + that is current at construction. + + 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): + """Create a :obj:`~cuda.core.texture.SurfaceObject` on this 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 bound to the + device/context that is current at construction. + + 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/_surface.pyi b/cuda_core/cuda/core/texture/_surface.pyi index be61e1099ba..c7b910a6670 100644 --- a/cuda_core/cuda/core/texture/_surface.pyi +++ b/cuda_core/cuda/core/texture/_surface.pyi @@ -14,7 +14,7 @@ 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). """ @@ -29,25 +29,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 +48,9 @@ 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`.""" \ 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 67fdfaf560c..63b0b125db3 100644 --- a/cuda_core/cuda/core/texture/_surface.pyx +++ b/cuda_core/cuda/core/texture/_surface.pyx @@ -33,70 +33,16 @@ 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). """ 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_opaque_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 +76,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.pyi b/cuda_core/cuda/core/texture/_texture.pyi index db654b11c87..78cccb1d00f 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -188,8 +188,8 @@ 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). """ def close(self): @@ -203,16 +203,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 : TextureObjectOptions - """ - @property def handle(self): """The underlying ``CUtexObject`` as an integer (64-bit kernel arg).""" @@ -250,4 +240,9 @@ def _normalize_enum(name, value, enum_type): def _normalize_address_modes(address_mode): """Return a 3-tuple of :class:`AddressModeType` values from a scalar or - 1-3 tuple. Individual entries may be plain strings.""" \ No newline at end of file + 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`.""" \ 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 83a28e387d6..53d85c35de5 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -37,6 +37,8 @@ from cuda.core.typing import AddressModeType, FilterModeType, ReadModeType from dataclasses import dataclass +from cuda.core._utils.cuda_utils import check_or_create_options + # Driver texture-descriptor flag bits (CU_TRSF_*). _TRSF_READ_AS_INTEGER = 0x01 @@ -409,142 +411,16 @@ 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). """ 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 : TextureObjectOptions - """ - if not isinstance(resource, ResourceDescriptor): - raise TypeError( - f"resource must be a ResourceDescriptor, got " - f"{type(resource).__name__}" - ) - if not isinstance(texture_descriptor, TextureObjectOptions): - raise TypeError( - f"texture_descriptor must be a TextureObjectOptions, 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 --- - # 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(texture_descriptor.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[texture_descriptor.filter_mode] - - cdef unsigned int flags = 0 - # CU_TRSF_READ_AS_INTEGER suppresses normalization, so it maps to - # ReadModeType.ELEMENT_TYPE. - if texture_descriptor.read_mode == ReadModeType.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 - - tex_desc.mipmapFilterMode = _FILTERMODE_TO_CU[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).""" @@ -583,3 +459,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._texture_desc = opts + self._device_id = _get_current_device_id() + return self diff --git a/cuda_core/docs/source/api.rst b/cuda_core/docs/source/api.rst index c96f6b77c90..d0078b073d4 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -171,7 +171,9 @@ namespace. :class:`OpaqueArray` is allocated through :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. +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__`` / diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index f2cb80d346e..41a9257c372 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -110,8 +110,6 @@ from cuda.core.texture import ( OpaqueArrayOptions, ResourceDescriptor, - SurfaceObject, - TextureObject, TextureObjectOptions, ) from cuda.core.typing import ( @@ -431,7 +429,7 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # 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. @@ -490,7 +488,13 @@ def make_texture(arr): # 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): @@ -554,25 +558,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_mipmap_lod.py b/cuda_core/examples/gl_interop_mipmap_lod.py index 26877b093f6..dd75199c29e 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -103,8 +103,6 @@ from cuda.core.texture import ( MipmappedArrayOptions, ResourceDescriptor, - SurfaceObject, - TextureObject, TextureObjectOptions, ) from cuda.core.typing import ( @@ -185,7 +183,9 @@ 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], @@ -223,8 +223,10 @@ def build_mipmap_pyramid(mip, num_levels, stream, kernels): dst_arr = mip.get_level(level) 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 = ( @@ -435,9 +437,9 @@ def main(): 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 0df09c23a3e..f0291db6747 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -32,7 +32,6 @@ from cuda.core.texture import ( OpaqueArrayOptions, ResourceDescriptor, - TextureObject, TextureObjectOptions, ) from cuda.core.typing import ( @@ -96,7 +95,7 @@ def main(): 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_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 10e12edc059..78ee4b1ceeb 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -187,7 +187,7 @@ def test_texture_object_create(init_cuda): 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 @@ -206,7 +206,7 @@ def test_surface_object_create(init_cuda): 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) @@ -220,7 +220,7 @@ def test_surface_requires_ldst_flag(init_cuda): 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() @@ -251,7 +251,7 @@ def test_address_mode_normalization(init_cuda): try: res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) - 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 finally: @@ -347,7 +347,7 @@ def test_texture_object_from_linear(init_cuda): buf = _alloc_device_buffer(device, 1024 * 4) try: res = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.FLOAT32, num_channels=1) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureObjectOptions()) + tex = Device().create_texture_object(resource=res, options=TextureObjectOptions()) try: assert tex.handle != 0 assert tex.resource is res @@ -417,7 +417,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=TextureObjectOptions()) + tex = Device().create_texture_object(resource=res, options=TextureObjectOptions()) try: assert tex.handle != 0 finally: @@ -432,7 +432,7 @@ def test_surface_rejects_linear_and_pitch2d(init_cuda): try: 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, @@ -443,7 +443,7 @@ def test_surface_rejects_linear_and_pitch2d(init_cuda): 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() @@ -585,7 +585,7 @@ def test_texture_object_from_mipmapped_array(init_cuda): 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 @@ -606,7 +606,7 @@ def test_surface_rejects_mipmapped_array(init_cuda): 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() @@ -667,16 +667,16 @@ def test_texture_surface_close_is_idempotent(init_cuda): 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 = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) - tex = TextureObject.from_descriptor( + tex = Device().create_texture_object( resource=ResourceDescriptor.from_opaque_array(tex_arr), - texture_descriptor=TextureObjectOptions(), + options=TextureObjectOptions(), ) tex.close() assert tex.handle == 0 @@ -846,7 +846,7 @@ def test_mipmapped_array_rejects_zero_dim(init_cuda): 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=TextureObjectOptions()) + Device().create_texture_object(resource=object(), options=TextureObjectOptions()) def test_texture_object_rejects_non_texture_descriptor(init_cuda): @@ -854,7 +854,7 @@ def test_texture_object_rejects_non_texture_descriptor(init_cuda): try: res = ResourceDescriptor.from_opaque_array(arr) with pytest.raises(TypeError, match="texture_descriptor must be a TextureObjectOptions"): - TextureObject.from_descriptor(resource=res, texture_descriptor="nope") + Device().create_texture_object(resource=res, options="nope") finally: arr.close() @@ -881,7 +881,7 @@ def test_texture_object_rejects_negative_anisotropy(init_cuda): 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() @@ -892,7 +892,7 @@ def test_texture_object_rejects_bad_border_color_length(init_cuda): 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() @@ -903,7 +903,7 @@ def test_address_mode_rejects_non_addressmode_scalar(init_cuda): 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() @@ -914,7 +914,7 @@ def test_address_mode_rejects_empty_tuple(init_cuda): 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() @@ -925,7 +925,7 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): 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() @@ -936,7 +936,7 @@ def test_address_mode_rejects_non_addressmode_entry(init_cuda): res = ResourceDescriptor.from_opaque_array(arr) td = TextureObjectOptions(address_mode=(AddressModeType.WRAP, "bad", AddressModeType.CLAMP)) with pytest.raises(TypeError, match=r"address_mode\[1\]"): - TextureObject.from_descriptor(resource=res, texture_descriptor=td) + Device().create_texture_object(resource=res, options=td) finally: arr.close() @@ -947,7 +947,7 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): TextureObject holds a strong ref through its _source_ref slot.""" arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) res = ResourceDescriptor.from_opaque_array(arr) - tex = TextureObject.from_descriptor(resource=res, texture_descriptor=TextureObjectOptions()) + 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. @@ -977,7 +977,7 @@ def test_surface_object_keeps_backing_array_alive(init_cuda): 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() From 0767000da96c67cd76536ab34a629962877f0379 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Mon, 6 Jul 2026 17:38:35 -0700 Subject: [PATCH 07/11] cuda.core.texture: apply pre-commit (ruff, ruff-format, stubgen) (#2292) Mechanical pre-commit pass over the #2292 texture API series: - ruff: drop now-unused imports in the texture example/test files and remove an unused _normalize_array_format import from _texture.pyx (cython-lint); - ruff-format: reflow the rewritten call sites in the tests and examples; - stubgen-pyx: regenerate the .pyi stubs from the updated .pyx sources. --- cuda_core/cuda/core/_device.pyi | 98 ++++++-- cuda_core/cuda/core/texture/_array.pyi | 30 ++- .../cuda/core/texture/_mipmapped_array.pyi | 17 +- cuda_core/cuda/core/texture/_surface.pyi | 6 +- cuda_core/cuda/core/texture/_texture.pyi | 14 +- cuda_core/cuda/core/texture/_texture.pyx | 1 - cuda_core/examples/gl_interop_fluid.py | 42 ++-- cuda_core/examples/gl_interop_mipmap_lod.py | 24 +- cuda_core/examples/texture_sample.py | 12 +- cuda_core/tests/test_texture_surface.py | 215 +++++++++++------- 10 files changed, 300 insertions(+), 159 deletions(-) diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index 6d202ee5923..76f59c4823c 100644 --- a/cuda_core/cuda/core/_device.pyi +++ b/cuda_core/cuda/core/_device.pyi @@ -12,16 +12,6 @@ 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: @@ -920,27 +910,107 @@ class Device: """ - def create_opaque_array(self, options: OpaqueArrayOptions | None = None) -> OpaqueArray: + def create_opaque_array(self, options: object=None): """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. + Allocates an opaque, hardware-laid-out CUDA array for texture/surface + access. The array is bound to the device/context that is current at + construction. + + 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 | None = None) -> MipmappedArray: + def create_mipmapped_array(self, options: object=None): """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. + Allocates a mipmapped CUDA array for texture/surface access across + levels. The array is bound to the device/context that is current at + construction. + + 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: + def create_texture_object(self, *, resource, options: object=None): """Create a :obj:`~cuda.core.texture.TextureObject` on this 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 bound to the device/context + that is current at construction. + + 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: + def create_surface_object(self, *, resource): """Create a :obj:`~cuda.core.texture.SurfaceObject` on this 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 bound to the + device/context that is current at construction. + + 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() diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 02ecafe4b12..a4ab361f914 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -4,6 +4,8 @@ from __future__ import annotations from dataclasses import dataclass +import numpy +from cuda.bindings import cydriver from cuda.core.typing import ArrayFormatType @@ -31,10 +33,9 @@ class OpaqueArrayOptions: num_channels: int is_surface_load_store: bool = False - def __post_init__(self) -> None: + def __post_init__(self): ... - class OpaqueArray: """An opaque, hardware-laid-out GPU allocation for texture/surface access. @@ -157,16 +158,23 @@ class OpaqueArray: def __repr__(self): ... -_ARRAYFORMAT_TO_CU: dict -_CU_TO_ARRAYFORMAT: dict -_NUMPY_DTYPE_TO_ARRAYFORMAT: dict -_FORMAT_ELEM_SIZE: dict +_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 an :class:`ArrayFormatType` or a plain ``str`` naming one of its - values (e.g. ``"float32"``). Raises on anything else.""" + 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, @@ -180,4 +188,8 @@ def _validate_array_shape(shape): def _create_opaque_array(options): """Allocate a new :class:`OpaqueArray` on the current device. - Backs :meth:`cuda.core.Device.create_opaque_array`.""" \ No newline at end of file + 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/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index bb9cf6db9ff..704578bbb34 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -20,9 +20,14 @@ class MipmappedArrayOptions: num_channels : int Channels per element. Must be 1, 2, or 4. num_levels : int - Number of mip levels to allocate; must be >= 1. + 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``. Default False. + 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. """ shape: tuple[int, ...] format: object @@ -30,10 +35,9 @@ class MipmappedArrayOptions: num_levels: int is_surface_load_store: bool = False - def __post_init__(self) -> None: + def __post_init__(self): ... - class MipmappedArray: """A mipmapped CUDA array for texture/surface access across levels. @@ -117,4 +121,7 @@ class MipmappedArray: def _create_mipmapped_array(options): """Allocate a new :class:`MipmappedArray` on the current device. - Backs :meth:`cuda.core.Device.create_mipmapped_array`.""" \ No newline at end of file + 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/_surface.pyi b/cuda_core/cuda/core/texture/_surface.pyi index c7b910a6670..c4c06387868 100644 --- a/cuda_core/cuda/core/texture/_surface.pyi +++ b/cuda_core/cuda/core/texture/_surface.pyi @@ -53,4 +53,8 @@ class SurfaceObject: def _create_surface_object(resource): """Create a :class:`SurfaceObject` on the current device. - Backs :meth:`cuda.core.Device.create_surface_object`.""" \ No newline at end of file + 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/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index 78cccb1d00f..99cfca67f20 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -4,7 +4,8 @@ from __future__ import annotations from dataclasses import dataclass -from cuda.core.typing import AddressModeType, ArrayFormatType, FilterModeType, ReadModeType +from cuda.bindings import cydriver +from cuda.core.typing import AddressModeType, FilterModeType, ReadModeType class ResourceDescriptor: @@ -178,7 +179,7 @@ class TextureObjectOptions: max_mipmap_level_clamp: float = 0.0 border_color: tuple[float, ...] | None = None - def __post_init__(self) -> None: + def __post_init__(self): ... class TextureObject: @@ -232,8 +233,8 @@ _TRSF_NORMALIZED_COORDINATES = 2 _TRSF_SRGB = 16 _TRSF_DISABLE_TRILINEAR_OPTIMIZATION = 32 _TRSF_SEAMLESS_CUBEMAP = 64 -_ADDRESSMODE_TO_CU: dict -_FILTERMODE_TO_CU: dict +_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.""" @@ -245,4 +246,7 @@ def _normalize_address_modes(address_mode): def _create_texture_object(resource, options): """Create a :class:`TextureObject` on the current device. - Backs :meth:`cuda.core.Device.create_texture_object`.""" \ No newline at end of file + 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 53d85c35de5..04d44ac4525 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -13,7 +13,6 @@ from cuda.core.texture._array import ( _ARRAYFORMAT_TO_CU, _CU_TO_ARRAYFORMAT, _FORMAT_ELEM_SIZE, - _normalize_array_format, _validate_format_channels, ) from cuda.core._memory._buffer cimport Buffer diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index 41a9257c372..eb6c2fc3fe3 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -438,22 +438,26 @@ 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 Device().create_opaque_array(OpaqueArrayOptions( - shape=(WIDTH, HEIGHT), - format=ArrayFormatType.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 Device().create_opaque_array(OpaqueArrayOptions( - shape=(WIDTH, HEIGHT), - format=ArrayFormatType.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, + ) + ) def make_color_array(): @@ -464,12 +468,14 @@ 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 Device().create_opaque_array(OpaqueArrayOptions( - shape=(WIDTH, HEIGHT), - format=ArrayFormatType.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, + ) + ) def make_texture(arr): diff --git a/cuda_core/examples/gl_interop_mipmap_lod.py b/cuda_core/examples/gl_interop_mipmap_lod.py index dd75199c29e..8c13b17c8ed 100644 --- a/cuda_core/examples/gl_interop_mipmap_lod.py +++ b/cuda_core/examples/gl_interop_mipmap_lod.py @@ -183,9 +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 Device().create_surface_object( - resource=ResourceDescriptor.from_opaque_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], @@ -224,9 +222,7 @@ def build_mipmap_pyramid(mip, num_levels, stream, kernels): src_res = ResourceDescriptor.from_opaque_array(src_arr) with ( 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, + Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(dst_arr)) as dst_surf, ): block = (16, 16, 1) grid = ( @@ -412,13 +408,15 @@ 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 = 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, - )) + 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) # --- Step 3: Bind the WHOLE pyramid as a trilinear-filtered texture --- diff --git a/cuda_core/examples/texture_sample.py b/cuda_core/examples/texture_sample.py index f0291db6747..952c77c48ca 100644 --- a/cuda_core/examples/texture_sample.py +++ b/cuda_core/examples/texture_sample.py @@ -71,11 +71,13 @@ def main(): # 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 Device().create_opaque_array(OpaqueArrayOptions( - shape=(width, height), - format=ArrayFormatType.FLOAT32, - num_channels=1, - )) as arr: + 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. ys, xs = np.meshgrid( diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 78ee4b1ceeb..fb797dedba7 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -11,13 +11,10 @@ Device, ) from cuda.core.texture import ( - MipmappedArray, MipmappedArrayOptions, OpaqueArray, OpaqueArrayOptions, ResourceDescriptor, - SurfaceObject, - TextureObject, TextureObjectOptions, ) from cuda.core.typing import ( @@ -49,7 +46,9 @@ def test_resource_descriptor_init_disabled(): def test_array_2d_create_and_properties(init_cuda): - arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.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 == ArrayFormatType.FLOAT32 @@ -64,12 +63,14 @@ def test_array_2d_create_and_properties(init_cuda): def test_array_3d_with_surface_flag(init_cuda): - arr = Device().create_opaque_array(OpaqueArrayOptions( - shape=(8, 8, 4), - format=ArrayFormatType.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 @@ -119,7 +120,9 @@ def test_array_rejects_bad_channels(init_cuda): def test_array_rejects_bad_rank(init_cuda): with pytest.raises(ValueError, match="shape rank"): - Device().create_opaque_array(OpaqueArrayOptions(shape=(2, 2, 2, 2), format=ArrayFormatType.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): @@ -178,7 +181,9 @@ def test_array_copy_rejects_undersized_device_buffer(init_cuda): def test_texture_object_create(init_cuda): - arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1)) + arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1) + ) try: res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureObjectOptions( @@ -199,12 +204,14 @@ def test_texture_object_create(init_cuda): def test_surface_object_create(init_cuda): - arr = Device().create_opaque_array(OpaqueArrayOptions( - shape=(8, 8), - format=ArrayFormatType.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 = Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(arr)) try: @@ -247,7 +254,9 @@ def test_address_mode_normalization(init_cuda): ) # Smoke test: a 2-entry tuple is also accepted end-to-end. - arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1)) + arr = Device().create_opaque_array( + OpaqueArrayOptions(shape=(8, 8, 4), format=ArrayFormatType.FLOAT32, num_channels=1) + ) try: res = ResourceDescriptor.from_opaque_array(arr) tex_desc = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.CLAMP)) @@ -457,12 +466,14 @@ def test_mipmapped_array_init_disabled(): def test_mipmapped_array_from_descriptor_2d(init_cuda): - mip = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(64, 32), - format=ArrayFormatType.FLOAT32, - num_channels=1, - num_levels=4, - )) + 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 == ArrayFormatType.FLOAT32 @@ -477,12 +488,14 @@ def test_mipmapped_array_from_descriptor_2d(init_cuda): def test_mipmapped_array_get_level_zero_matches_shape(init_cuda): shape = (64, 32) - mip = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=shape, - format=ArrayFormatType.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) try: @@ -501,12 +514,14 @@ 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 = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=shape, - format=ArrayFormatType.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): lvl = mip.get_level(level) @@ -521,12 +536,14 @@ def test_mipmapped_array_get_level_halves_dims(init_cuda): def test_mipmapped_array_get_level_out_of_range(init_cuda): - mip = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(16, 16), - format=ArrayFormatType.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"): mip.get_level(mip.num_levels) @@ -538,21 +555,25 @@ 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"): - Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(8, 8), - format=ArrayFormatType.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 = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(32, 16), - format=ArrayFormatType.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) assert res.kind == "mipmapped_array" @@ -567,12 +588,14 @@ def test_resource_descriptor_from_mipmapped_array_rejects_non_mipmap(): def test_texture_object_from_mipmapped_array(init_cuda): - mip = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(32, 32), - format=ArrayFormatType.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. @@ -596,13 +619,15 @@ def test_texture_object_from_mipmapped_array(init_cuda): def test_surface_rejects_mipmapped_array(init_cuda): - mip = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(16, 16), - format=ArrayFormatType.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"): @@ -626,12 +651,14 @@ def test_mipmapped_array_level_outlives_dropped_parent(init_cuda): device = Device() stream = device.create_stream() - mip = Device().create_mipmapped_array(MipmappedArrayOptions( - shape=(16, 16), - format=ArrayFormatType.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. del mip @@ -659,21 +686,25 @@ def test_texture_surface_close_is_idempotent(init_cuda): assert arr.handle == 0 arr.close() # second close must not raise - mip = Device().create_mipmapped_array(MipmappedArrayOptions(shape=(8, 8), format=ArrayFormatType.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 = Device().create_opaque_array(OpaqueArrayOptions( - shape=(8, 8), format=ArrayFormatType.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 = Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(surf_arr)) surf.close() assert surf.handle == 0 surf.close() surf_arr.close() - tex_arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) + 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(), @@ -836,12 +867,16 @@ def test_mipmapped_array_rejects_bad_format(init_cuda): def test_mipmapped_array_rejects_bad_channels(init_cuda): with pytest.raises(ValueError, match="num_channels"): - Device().create_mipmapped_array(MipmappedArrayOptions(shape=(8, 8), format=ArrayFormatType.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"): - Device().create_mipmapped_array(MipmappedArrayOptions(shape=(0, 8), format=ArrayFormatType.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): @@ -923,7 +958,9 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): arr = Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 8), format=ArrayFormatType.FLOAT32, num_channels=1)) try: res = ResourceDescriptor.from_opaque_array(arr) - td = TextureObjectOptions(address_mode=(AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP)) + td = TextureObjectOptions( + address_mode=(AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP, AddressModeType.WRAP) + ) with pytest.raises(ValueError, match="address_mode tuple must have 1-3"): Device().create_texture_object(resource=res, options=td) finally: @@ -971,12 +1008,14 @@ def test_texture_object_keeps_backing_array_alive(init_cuda): def test_surface_object_keeps_backing_array_alive(init_cuda): - arr = Device().create_opaque_array(OpaqueArrayOptions( - shape=(8, 8), - format=ArrayFormatType.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 = Device().create_surface_object(resource=ResourceDescriptor.from_opaque_array(arr)) arr_id = id(arr) del arr From 16940d76e78a28c40abfffd90e53443219d1087e Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Tue, 7 Jul 2026 10:22:32 -0700 Subject: [PATCH 08/11] cuda.core.texture: address API-review follow-ups (#2292) Apply the review findings from the texture API-style follow-up: - Rename TextureObject.texture_descriptor -> options (and the internal _texture_desc slot -> _options) to drop the borrowed "descriptor" naming, matching the TextureObjectOptions rename. - Make options a required argument on create_opaque_array / create_mipmapped_array (shape/format/num_channels have no defaults), so omitting it raises a clear error instead of a raw dataclass TypeError. create_texture_object keeps options=None (all-default sampling state). - Add return and resource/options type annotations to the four Device.create_* texture factories via a TYPE_CHECKING texture import (mirroring the existing GraphBuilder pattern). - Clarify in the create_* docstrings that the resource is created in the current CUDA context (call set_current first), mirroring create_stream / create_event. - Add versionadded:: 1.1.0 to OpaqueArrayOptions, MipmappedArrayOptions, TextureObjectOptions, and the new options property. - Fix a stale negative-path test whose regex no longer matched the shared check_or_create_options message; rename stale from_descriptor test names to create_*; realign the gl_interop_fluid API-MAP comment table after the enum renames. - Regenerate .pyi stubs. --- cuda_core/cuda/core/_device.pyi | 41 ++++++++++------ cuda_core/cuda/core/_device.pyx | 49 +++++++++++++------ cuda_core/cuda/core/texture/_array.pyi | 2 + cuda_core/cuda/core/texture/_array.pyx | 2 + .../cuda/core/texture/_mipmapped_array.pyi | 2 + .../cuda/core/texture/_mipmapped_array.pyx | 2 + cuda_core/cuda/core/texture/_texture.pxd | 2 +- cuda_core/cuda/core/texture/_texture.pyi | 9 +++- cuda_core/cuda/core/texture/_texture.pyx | 13 +++-- cuda_core/examples/gl_interop_fluid.py | 14 +++--- cuda_core/tests/test_texture_surface.py | 14 +++--- 11 files changed, 97 insertions(+), 53 deletions(-) diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index 76f59c4823c..9e76f1c8961 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: @@ -910,12 +914,13 @@ class Device: """ - def create_opaque_array(self, options: object=None): - """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. + 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 bound to the device/context that is current at - construction. + 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 ---- @@ -934,12 +939,13 @@ class Device: .. versionadded:: 1.1.0 """ - def create_mipmapped_array(self, options: object=None): - """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. + 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 bound to the device/context that is current at - construction. + 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 ---- @@ -958,15 +964,16 @@ class Device: .. versionadded:: 1.1.0 """ - def create_texture_object(self, *, resource, options: object=None): - """Create a :obj:`~cuda.core.texture.TextureObject` on this device. + 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 bound to the device/context - that is current at construction. + 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 ---- @@ -987,14 +994,16 @@ class Device: .. versionadded:: 1.1.0 """ - def create_surface_object(self, *, resource): - """Create a :obj:`~cuda.core.texture.SurfaceObject` on this device. + 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 bound to the - device/context that is current at construction. + 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 ---- diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index 5cecf47624b..412223bcc05 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,12 +1466,13 @@ class Device: self._check_context_initialized() return GraphBuilder._init(self.create_stream()) - def create_opaque_array(self, options: object = None): - """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. + 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 bound to the device/context that is current at - construction. + 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 ---- @@ -1484,12 +1495,13 @@ class Device: self._check_context_initialized() return _create_opaque_array(options) - def create_mipmapped_array(self, options: object = None): - """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. + 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 bound to the device/context that is current at - construction. + 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 ---- @@ -1512,15 +1524,18 @@ class Device: self._check_context_initialized() return _create_mipmapped_array(options) - def create_texture_object(self, *, resource, options: object = None): - """Create a :obj:`~cuda.core.texture.TextureObject` on this device. + 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 bound to the device/context - that is current at construction. + 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 ---- @@ -1545,14 +1560,16 @@ class Device: self._check_context_initialized() return _create_texture_object(resource, options) - def create_surface_object(self, *, resource): - """Create a :obj:`~cuda.core.texture.SurfaceObject` on this device. + 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 bound to the - device/context that is current at construction. + 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 ---- diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index a4ab361f914..6467a69ba8b 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -27,6 +27,8 @@ class OpaqueArrayOptions: 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 diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index b2f0499c684..794082618d8 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -154,6 +154,8 @@ class OpaqueArrayOptions: 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, ...] diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index 704578bbb34..8cf53f02e78 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -28,6 +28,8 @@ class MipmappedArrayOptions: 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 diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyx b/cuda_core/cuda/core/texture/_mipmapped_array.pyx index 4338b1fa2a7..43edf9ef179 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyx +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyx @@ -55,6 +55,8 @@ class MipmappedArrayOptions: 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, ...] diff --git a/cuda_core/cuda/core/texture/_texture.pxd b/cuda_core/cuda/core/texture/_texture.pxd index fcac323af91..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 TextureObjectOptions 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 99cfca67f20..7e1362bf25b 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -164,6 +164,8 @@ class TextureObjectOptions: border_color : tuple of float or None 4-tuple used when ``address_mode`` includes ``BORDER``; ``None`` means zero. + + .. versionadded:: 1.1.0 """ address_mode: AddressModeType | tuple[AddressModeType, ...] = AddressModeType.CLAMP filter_mode: FilterModeType = FilterModeType.POINT @@ -213,8 +215,11 @@ class TextureObject: """The :class:`ResourceDescriptor` this texture was built from.""" @property - def texture_descriptor(self): - """The :class:`TextureObjectOptions` this texture was built from.""" + def options(self): + """The :class:`TextureObjectOptions` this texture was built from. + + .. versionadded:: 1.1.0 + """ @property def device(self): diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index 04d44ac4525..ee0df71a138 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -354,6 +354,8 @@ class TextureObjectOptions: border_color : tuple of float or None 4-tuple used when ``address_mode`` includes ``BORDER``; ``None`` means zero. + + .. versionadded:: 1.1.0 """ address_mode: AddressModeType | tuple[AddressModeType, ...] = AddressModeType.CLAMP @@ -431,9 +433,12 @@ cdef class TextureObject: return self._source_ref @property - def texture_descriptor(self): - """The :class:`TextureObjectOptions` this texture was built from.""" - return self._texture_desc + def options(self): + """The :class:`TextureObjectOptions` this texture was built from. + + .. versionadded:: 1.1.0 + """ + return self._options @property def device(self): @@ -578,6 +583,6 @@ def _create_texture_object(resource, options): cdef TextureObject self = TextureObject.__new__(TextureObject) self._handle = h self._source_ref = resource - self._texture_desc = opts + self._options = opts self._device_id = _get_current_device_id() return self diff --git a/cuda_core/examples/gl_interop_fluid.py b/cuda_core/examples/gl_interop_fluid.py index eb6c2fc3fe3..68259d32e3b 100644 --- a/cuda_core/examples/gl_interop_fluid.py +++ b/cuda_core/examples/gl_interop_fluid.py @@ -401,10 +401,10 @@ def draw_fullscreen_quad(gl, shader_prog, vao_id, tex_id): # TextureObjectOptions / TextureObject / SurfaceObject knob in this example is set. # Each visible setting maps to a concrete piece of cuda.core / CUDA behavior: # -# Device().create_opaque_array(OpaqueArrayOptions(...)) -> 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. -# ArrayFormatType.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,17 +414,17 @@ 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_opaque_array(arr) -> wraps the OpaqueArray as the resource a -# TextureObject reads from. -# FilterModeType.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). -# AddressModeType.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). -# ReadModeType.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. diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index fb797dedba7..d786eac0c65 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -196,7 +196,7 @@ def test_texture_object_create(init_cuda): 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: @@ -465,7 +465,7 @@ def test_mipmapped_array_init_disabled(): cuda.core.texture._mipmapped_array.MipmappedArray() -def test_mipmapped_array_from_descriptor_2d(init_cuda): +def test_mipmapped_array_create_2d(init_cuda): mip = Device().create_mipmapped_array( MipmappedArrayOptions( shape=(64, 32), @@ -718,17 +718,17 @@ def test_texture_surface_close_is_idempotent(init_cuda): # --- Negative-path validation tests ------------------------------------------ -def test_array_from_descriptor_rejects_bad_format(init_cuda): +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"): 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"): Device().create_opaque_array(OpaqueArrayOptions(shape=(8, 0), format=ArrayFormatType.UINT8, num_channels=1)) @@ -884,11 +884,11 @@ def test_texture_object_rejects_non_resource_descriptor(init_cuda): Device().create_texture_object(resource=object(), options=TextureObjectOptions()) -def test_texture_object_rejects_non_texture_descriptor(init_cuda): +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_opaque_array(arr) - with pytest.raises(TypeError, match="texture_descriptor must be a TextureObjectOptions"): + 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() From 57bcf4cd62dffd827bf41c58591e9640a073eba6 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Tue, 7 Jul 2026 10:24:36 -0700 Subject: [PATCH 09/11] cuda.core.texture: fix stale address_mode entry-rejection test (#2292) The StrEnum migration (item 4 of #2292) makes a plain str acceptable anywhere an AddressModeType is expected. A bad tuple entry like "bad" is therefore a valid *type* but an invalid *value*, so _normalize_enum rejects it with a ValueError (naming the offending position address_mode[1]), not a TypeError. Update the test to assert ValueError and rename it to test_address_mode_rejects_invalid_entry. This is the second of the two CI failures on PR #2307; the first (the stale options-type message) was fixed in the previous commit. --- cuda_core/tests/test_texture_surface.py | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index d786eac0c65..bee60104be7 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -967,12 +967,15 @@ def test_address_mode_rejects_too_long_tuple(init_cuda): arr.close() -def test_address_mode_rejects_non_addressmode_entry(init_cuda): +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_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(TypeError, match=r"address_mode\[1\]"): + with pytest.raises(ValueError, match=r"address_mode\[1\]"): Device().create_texture_object(resource=res, options=td) finally: arr.close() From 60d36b797fc76d642838ea18982a4192da2882f1 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Tue, 7 Jul 2026 10:35:52 -0700 Subject: [PATCH 10/11] cuda.core.texture: register new StrEnums in enum-coverage test (#2292) test_enum_coverage.py::test_all_str_enums_in_cases requires every StrEnum in cuda.core to be declared either in _CASES (bound to a cuda_binding enum) or in _UNBOUND_STR_ENUMS. The four texture enums moved into cuda.core.typing (item 4 of #2292) were not registered, failing that guard. - AddressModeType / FilterModeType are 1:1 wrappers of CUaddress_mode / CUfilter_mode, so add mapping=None _CASES entries (count-check only; the texture mapping dicts store cydriver-derived ints, not driver. members, so the isinstance-based mapping check does not apply). - ArrayFormatType exposes only the 8 NumPy-representable formats out of CUarray_format's ~67 members, so it is a curated subset rather than a 1:1 wrapper -> _UNBOUND_STR_ENUMS. - ReadModeType maps to the CU_TRSF_READ_AS_INTEGER flag bit, not a CUenum -> _UNBOUND_STR_ENUMS. This is the third local test failure; the other two (stale texture negative-path tests) were fixed in the two preceding commits. --- cuda_core/tests/test_enum_coverage.py | 29 +++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) 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, } From fa7cb36b9803dde5369e1c21509e933d04e30768 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Tue, 7 Jul 2026 16:39:03 -0700 Subject: [PATCH 11/11] cuda.core.texture: add versionadded markers and widen enum str annotations Address review feedback (Ralf): - Add ``.. versionadded:: 1.1.0`` to the four new texture enums in typing.py (ArrayFormatType, AddressModeType, FilterModeType, ReadModeType) and to the texture classes that lacked a class-level marker (ResourceDescriptor, TextureObject, OpaqueArray, MipmappedArray, SurfaceObject). Remove the redundant method-level markers now covered at class level. - Widen the enum-bearing TextureObjectOptions field annotations to include ``str`` (address_mode, filter_mode, read_mode, mipmap_filter_mode), matching the runtime string coercion this PR adds and existing convention elsewhere in cuda.core. Note "Plain strings are accepted" on the affected fields. - Regenerate .pyi stubs. --- cuda_core/cuda/core/texture/_array.pyi | 7 ++--- cuda_core/cuda/core/texture/_array.pyx | 7 ++--- .../cuda/core/texture/_mipmapped_array.pyi | 2 ++ .../cuda/core/texture/_mipmapped_array.pyx | 2 ++ cuda_core/cuda/core/texture/_surface.pyi | 2 ++ cuda_core/cuda/core/texture/_surface.pyx | 2 ++ cuda_core/cuda/core/texture/_texture.pyi | 28 +++++++++---------- cuda_core/cuda/core/texture/_texture.pyx | 28 +++++++++---------- cuda_core/cuda/core/typing.py | 8 ++++++ 9 files changed, 50 insertions(+), 36 deletions(-) diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 6467a69ba8b..380c2fe1c10 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -58,6 +58,8 @@ class OpaqueArray: 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): @@ -100,10 +102,7 @@ class OpaqueArray: @property def element_bytes(self): - """Bytes per element (format size * channels). - - .. versionadded:: 1.1.0 - """ + """Bytes per element (format size * channels).""" @property def device(self): diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index 794082618d8..0a1cb671daf 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -336,6 +336,8 @@ cdef class OpaqueArray: 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): @@ -384,10 +386,7 @@ cdef class OpaqueArray: @property def element_bytes(self): - """Bytes per element (format size * channels). - - .. versionadded:: 1.1.0 - """ + """Bytes per element (format size * channels).""" return _FORMAT_ELEM_SIZE[self._format] * self._num_channels @property diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index 8cf53f02e78..db4413dbaf4 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -51,6 +51,8 @@ class MipmappedArray: are non-owning and hold a strong reference back to their parent. Construct via :meth:`cuda.core.Device.create_mipmapped_array`. + + .. versionadded:: 1.1.0 """ def close(self): diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyx b/cuda_core/cuda/core/texture/_mipmapped_array.pyx index 43edf9ef179..3f151f7bb9f 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyx +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyx @@ -84,6 +84,8 @@ cdef class MipmappedArray: are non-owning and hold a strong reference back to their parent. Construct via :meth:`cuda.core.Device.create_mipmapped_array`. + + .. versionadded:: 1.1.0 """ def __init__(self, *args, **kwargs): diff --git a/cuda_core/cuda/core/texture/_surface.pyi b/cuda_core/cuda/core/texture/_surface.pyi index c4c06387868..977268abd5f 100644 --- a/cuda_core/cuda/core/texture/_surface.pyi +++ b/cuda_core/cuda/core/texture/_surface.pyi @@ -16,6 +16,8 @@ class SurfaceObject: 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): diff --git a/cuda_core/cuda/core/texture/_surface.pyx b/cuda_core/cuda/core/texture/_surface.pyx index 63b0b125db3..ac61fddd357 100644 --- a/cuda_core/cuda/core/texture/_surface.pyx +++ b/cuda_core/cuda/core/texture/_surface.pyx @@ -35,6 +35,8 @@ cdef class SurfaceObject: 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): diff --git a/cuda_core/cuda/core/texture/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index 7e1362bf25b..7840585bb4d 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -25,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') @@ -33,10 +35,7 @@ class ResourceDescriptor: @classmethod def from_opaque_array(cls, array): - """Build a resource descriptor backed by a :class:`OpaqueArray`. - - .. versionadded:: 1.1.0 - """ + """Build a resource descriptor backed by a :class:`OpaqueArray`.""" @classmethod def from_mipmapped_array(cls, mipmapped_array): @@ -143,9 +142,10 @@ class TextureObjectOptions: :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``. + 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 @@ -157,7 +157,8 @@ class TextureObjectOptions: max_anisotropy : int Maximum anisotropy; 0 disables anisotropic filtering. mipmap_filter_mode : FilterModeType - Filtering between mipmap levels. Default ``POINT``. + Filtering between mipmap levels. Default ``POINT``. Plain strings are + accepted. mipmap_level_bias : float min_mipmap_level_clamp : float max_mipmap_level_clamp : float @@ -167,15 +168,15 @@ class TextureObjectOptions: .. versionadded:: 1.1.0 """ - address_mode: AddressModeType | tuple[AddressModeType, ...] = AddressModeType.CLAMP - filter_mode: FilterModeType = FilterModeType.POINT - read_mode: ReadModeType = ReadModeType.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: FilterModeType = FilterModeType.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 @@ -193,6 +194,8 @@ class TextureObject: 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): @@ -216,10 +219,7 @@ class TextureObject: @property def options(self): - """The :class:`TextureObjectOptions` this texture was built from. - - .. versionadded:: 1.1.0 - """ + """The :class:`TextureObjectOptions` this texture was built from.""" @property def device(self): diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index ee0df71a138..8d09c727c31 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -90,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__ = ( @@ -107,10 +109,7 @@ class ResourceDescriptor: @classmethod def from_opaque_array(cls, array): - """Build a resource descriptor backed by a :class:`OpaqueArray`. - - .. versionadded:: 1.1.0 - """ + """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__}") self = cls.__new__(cls) @@ -333,9 +332,10 @@ class TextureObjectOptions: :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``. + 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 @@ -347,7 +347,8 @@ class TextureObjectOptions: max_anisotropy : int Maximum anisotropy; 0 disables anisotropic filtering. mipmap_filter_mode : FilterModeType - Filtering between mipmap levels. Default ``POINT``. + Filtering between mipmap levels. Default ``POINT``. Plain strings are + accepted. mipmap_level_bias : float min_mipmap_level_clamp : float max_mipmap_level_clamp : float @@ -358,15 +359,15 @@ class TextureObjectOptions: .. versionadded:: 1.1.0 """ - address_mode: AddressModeType | tuple[AddressModeType, ...] = AddressModeType.CLAMP - filter_mode: FilterModeType = FilterModeType.POINT - read_mode: ReadModeType = ReadModeType.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: FilterModeType = FilterModeType.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 @@ -414,6 +415,8 @@ cdef class TextureObject: 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): @@ -434,10 +437,7 @@ cdef class TextureObject: @property def options(self): - """The :class:`TextureObjectOptions` this texture was built from. - - .. versionadded:: 1.1.0 - """ + """The :class:`TextureObjectOptions` this texture was built from.""" return self._options @property diff --git a/cuda_core/cuda/core/typing.py b/cuda_core/cuda/core/typing.py index 0bf5f93d391..2d253d8ca5e 100644 --- a/cuda_core/cuda/core/typing.py +++ b/cuda_core/cuda/core/typing.py @@ -232,6 +232,8 @@ class ArrayFormatType(StrEnum): * ``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" @@ -253,6 +255,8 @@ class AddressModeType(StrEnum): * ``CLAMP`` — clamp to the edge texel. * ``MIRROR`` — reflect coordinates at the boundary. * ``BORDER`` — return the configured border color. + + .. versionadded:: 1.1.0 """ WRAP = "wrap" @@ -268,6 +272,8 @@ class FilterModeType(StrEnum): * ``POINT`` — nearest-texel sampling. * ``LINEAR`` — (bi/tri)linear interpolation. + + .. versionadded:: 1.1.0 """ POINT = "point" @@ -282,6 +288,8 @@ class ReadModeType(StrEnum): * ``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"