Skip to content

Commit 1e53cc6

Browse files
rparolinclaude
andauthored
cuda.core: add CUDAArray, TextureObject, SurfaceObject, MipmappedArray (texture/surface API, NVIDIA#467) (NVIDIA#2095)
* Add Array and ArrayFormat to cuda.core (refs NVIDIA#467) Introduce a Pythonic wrapper around CUarray as a prerequisite for TextureObject / SurfaceObject support. This initial slice covers plain 1D/2D/3D allocations via cuArrayCreate / cuArray3DCreate, with an opt-in surface_load_store flag for binding as a SurfaceObject. Layered, cubemap, sparse, and texture-gather variants are intentionally deferred. _from_handle is provided for graphics-interop borrowing and queries shape, format, and channel count from the driver. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add copy_from / copy_to to cuda.core.Array (refs NVIDIA#467) Full-array async copies between an Array and either a Buffer or any buffer-protocol host object (numpy, bytes, bytearray, array.array). Implemented as a single cuMemcpy3DAsync path so 1D/2D/3D arrays share one code path. Also exposes a size_bytes property used to size matching host or device buffers. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add TextureObject and ResourceDescriptor/TextureDescriptor (refs NVIDIA#467) Wraps cuTexObjectCreate with a Pythonic descriptor pair: - ResourceDescriptor.from_array(array) is the only resource kind supported in this initial slice; from_linear and from_pitch2d will follow once Buffer carries format/channel metadata. - TextureDescriptor mirrors CUDA_TEXTURE_DESC: per-axis AddressMode, FilterMode, ReadMode, normalized coords, sRGB, border color, mipmap params, anisotropy. - TextureObject holds a strong ref to the ResourceDescriptor (and transitively the backing Array) for the lifetime of the handle to prevent dangling-pointer kernel launches. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add SurfaceObject for kernel-side typed load/store (refs NVIDIA#467) Completes the second half of NVIDIA#467 alongside the existing TextureObject: - SurfaceObject wraps cuSurfObjectCreate / cuSurfObjectDestroy. Unlike a texture it has no sampling state (no filter mode, no addressing, no normalization); kernels read and write through it with integer pixel coordinates. - Track CUDA_ARRAY3D_SURFACE_LDST on Array as a new surface_load_store property, populated in both Array.from_descriptor and Array._from_handle. SurfaceObject.from_array validates this upfront rather than letting the driver surface CUDA_ERROR_INVALID_VALUE late. - Add a convenience SurfaceObject.from_array shortcut next to from_descriptor so the common case skips building a ResourceDescriptor by hand. Covered by tests/test_texture_surface.py (14 tests: array shape/format/ flag plumbing, texture + surface creation, surface_load_store validation, unsupported-resource-kind guard). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add ResourceDescriptor.from_linear and .from_pitch2d (refs NVIDIA#467) Widens the texture-resource surface to cover the two Buffer-backed variants from CUDA_RESOURCE_DESC: - ResourceDescriptor.from_linear(buffer, *, format, num_channels, size_bytes=None) wraps a Buffer as a typed 1D fetch. Defaults size_bytes to buffer.size; validates against it. - ResourceDescriptor.from_pitch2d(buffer, *, format, num_channels, width, height, pitch_bytes) wraps a Buffer as a row-pitched 2D image. Validates pitch_bytes >= width * element_size and pitch_bytes * height <= buffer.size; the driver enforces its own CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT on top. - TextureObject.from_descriptor handles the three resType branches (ARRAY, LINEAR, PITCH2D); SurfaceObject continues to require an array-backed resource. - ResourceDescriptor gains format/num_channels read-only properties (None for array-backed) and a kind-aware __repr__. Tests: 9 new (linear/pitch2D creation, validation paths, surface rejection of non-array resources) on top of the existing 14. Full test-core suite green (3287 passed). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add docs for texture/surface APIs (refs NVIDIA#467) Wire the newly public Array, ArrayFormat, TextureObject, SurfaceObject, ResourceDescriptor, TextureDescriptor, AddressMode, FilterMode, and ReadMode symbols into the cuda.core Sphinx reference under a new "Textures and surfaces" section in api.rst. No source docstring changes; documentation is rendered via the existing autosummary templates and the enum_documenter extension. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add MipmappedArray and ResourceDescriptor.from_mipmapped_array (refs NVIDIA#467) Introduces a MipmappedArray cdef class wrapping CUmipmappedArray with the same lifetime model as Array (close/__dealloc__/context-manager). Levels are obtained via get_level(L), which returns a non-owning Array that holds a strong ref back to the parent MipmappedArray via a new Array._parent_ref slot, ensuring level views cannot outlive the underlying storage. Surfaces continue to require a single-Array backing; the existing kind != "array" check in SurfaceObject.from_descriptor naturally rejects mipmapped resources (covered by a new test). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add texture sampling example (refs NVIDIA#467) End-to-end example that builds a 2D Array with a known pattern, binds it as a bindless TextureObject with LINEAR/CLAMP/non-normalized sampling, and launches a kernel that samples both texel-center and half-integer coordinates. Verifies POINT-exact returns at texel centers and analytical bilinear blends at half-pixel positions. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Address code review feedback on texture/surface stack (refs NVIDIA#467) Safety and correctness: - Validate buffer sizes against array extent in Array.copy_from/copy_to; undersized host or device Buffer inputs were previously silent stomps via cuMemcpy3DAsync. Both branches now raise ValueError before issuing the copy. - Zero the underlying handle BEFORE calling cuXxxDestroy in close() for Array, MipmappedArray, TextureObject, SurfaceObject. Prevents a double-destroy via __dealloc__ if the driver call raises. - ResourceDescriptor.from_linear: require size_bytes >= element_size and size_bytes % element_size == 0; previously accepted zero and arbitrary non-multiples. - Reject bool in num_channels across Array, MipmappedArray, and the two Buffer-backed ResourceDescriptor factories (True was silently treated as 1 channel). API polish: - Rename TextureObject.from_descriptor params resource_desc/texture_desc to resource/texture_descriptor so they match the .resource and .texture_descriptor properties; same rename in SurfaceObject. Both factories are now keyword-only, consistent with Array.from_descriptor and MipmappedArray.from_descriptor. - Add four ResourceDescriptor properties (size_bytes, width, height, pitch_bytes) so values shown in __repr__ are reachable programmatically. - Add MipmappedArray to docs/source/api.rst (was exported but unlinked). - Align error message style across new files: type(x).__name__ instead of type(x); include got <type> in three previously-bare TypeErrors in TextureObject.from_descriptor. Refactor: - Extract _get_current_context_ptr and _get_current_device_id to cuda_utils.{pxd,pyx} and share across all four new files (was duplicated four times). Generic error message keeps the helper reusable for the 9+ remaining duplicate sites in cuda.core. - Hoist the buffer-protocol path in _fill_linear_endpoint into a new _fill_host_endpoint helper. Original function becomes a thin Buffer-vs-host router. - Type Array._format and MipmappedArray._format as cydriver.CUarray_format instead of int (was a comment-typed int; now C-level type-checked). - Drop unused `field` import from _texture.pyx. Tests (+28, total 62 in this file): - Undersized host/device buffer rejection in Array.copy_from/copy_to. - ResourceDescriptor.from_linear rejects size_bytes=0 and non-multiples. - _normalize_address_modes unit tests now make explicit assertions instead of only smoke-testing TextureObject creation. - Negative-path coverage for Array.from_descriptor (bad format, non- iterable shape, zero dim), MipmappedArray.from_descriptor, all TextureObject.from_descriptor validation branches (filter_mode, read_mode, mipmap_filter_mode, max_anisotropy, border_color length), address-mode normalization (scalar non-AddressMode, empty/4-entry tuples, mixed-type entries), ResourceDescriptor.from_pitch2d, and copy_from/copy_to non-Stream rejection. - TextureObject and SurfaceObject keepalive lifetime tests verifying the _source_ref chain holds after gc.collect() (mirrors the existing MipmappedArray level keepalive test). - copy_from must not mutate the source buffer (round-trip test now also asserts list(src) is unchanged). Example: - texture_sample.py uses `with` blocks for Array and TextureObject so the user-facing demo shows the idiomatic context-manager pattern rather than manual try/finally. Full cuda_core suite: 3326 passed, 199 skipped, 2 xfailed. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add 9 cuda.core texture/surface examples (refs NVIDIA#467) These graphical examples demonstrate the new Array, TextureObject, SurfaceObject, MipmappedArray, and ResourceDescriptor APIs in increasing order of complexity. All use the existing GraphicsResource + GL PBO pattern for display (matching gl_interop_plasma.py); CI is gated on has_display so headless runners skip them. Minimum-API examples: - gl_interop_image_show.py Hello-world for the stack: 64x64 Array, TextureObject, key F toggles POINT/LINEAR. Read this file first. - gl_interop_texture_filter.py POINT vs LINEAR side-by-side on one Array with two TextureObjects; mouse pan/zoom, key M cycles AddressMode. Simulation examples (Array + SurfaceObject + TextureObject ping-pong): - gl_interop_reaction_diffusion.py Gray-Scott with FLOAT32 x 2 channels; LINEAR + WRAP for toroidal diffusion. - gl_interop_lenia.py Continuous-state CA with bell-curve convolution; FLOAT32 x 1 channel. - gl_interop_fire.py Canonical Doom fire (37-color indexed palette, UINT8 intensity 0..36, gather equivalent of the original scatter algorithm); exercises ArrayFormat.UINT8. - gl_interop_ocean.py Animated Gerstner-wave ocean with normal mapping via finite-difference texture reads and Phong + Fresnel shading. Visualization examples: - gl_interop_mandelbrot.py Real-time deep-zoom using a 1D Array as a color LUT (TextureObject for palette lookup, not simulation). - gl_interop_mipmap_lod.py Procedural mipmap pyramid built with a SurfaceObject per level; trilinear sampling via tex2DLod and TextureDescriptor mipmap fields. - gl_interop_sdf_volume.py 3D ray-marched gyroid via a 128^3 Array, surf3Dwrite for bake, tex3D for trilinear SDF sampling. Only example exercising the 3D side of the API. Every public symbol added in this PR is exercised by at least one of these examples. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * cuda.core: rename Array->CUDAArray, surface_load_store->is_surface_load_store (NVIDIA#467) Applies design decisions resolved in NVIDIA#2188: - #1: rename public `Array` class to `CUDAArray` (PEP 8 CapWords; aligns with CuPy's `cupy.cuda.texture.CUDAarray`). `ArrayFormat` left unchanged (open detail). - #6: rename the bool property `surface_load_store` -> `is_surface_load_store` to follow the repo's `is_<x>` convention. Constructor keyword `surface_load_store=` kept as-is (open detail). Private field `_surface_load_store` unchanged. GL interop examples retained (decision #7 reversed) and updated to the new names. Verified: cuda.core builds in the cu12 env and the renamed public API imports (`CUDAArray`, `is_surface_load_store` present; old `Array`/`surface_load_store` gone). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: fix lint/mypy after merging main; add generated .pyi stubs Post-merge cleanup so pre-commit.ci passes on the texture/surface stack: - Switch ReadMode/AddressMode/FilterMode/ArrayFormat to `from enum import IntEnum` so stubgen-pyx preserves the IntEnum base in the generated stubs (qualified `enum.IntEnum` was dropped, making members infer as `int` and failing mypy assignment checks). - Annotate TextureDescriptor.border_color as `tuple[float, ...] | None` (disallow_any_generics flagged the bare `tuple`). - Prefix unused pyglet event-handler args with `_` (ARG001) and lowercase in-function locals (N806) across the GL interop examples; drop dead pre-try buffer inits in texture_sample.main (F841). - Commit the auto-generated .pyi stubs for the new _array/_texture/_surface/ _mipmapped_array modules. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: rename surface_load_store ctor keyword to is_surface_load_store Closes the open detail on design issue NVIDIA#2188 item #6: the read-back property is already `is_surface_load_store`, so rename the `from_descriptor` keyword on both `CUDAArray` and `MipmappedArray` to match, giving one symmetric name for set and read-back (following the existing `StridedMemoryView(is_readonly=...)` precedent). Updates call sites in tests and GL examples, the SurfaceObject error message + docstrings, and regenerates the .pyi stubs. The unrelated GraphicsResource `"surface_load_store"` register-flag string is left untouched. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: add 7 texture/surface graphics examples Add seven new cuda.core GL-interop examples exercising the new CUDAArray / MipmappedArray / SurfaceObject / TextureObject / GraphicsResource APIs, each centered on a distinct feature and verified on-GPU: - gl_interop_fluid.py Stable-Fluids ink: LINEAR advection, float4 dye, frame-rate-independent stepping - gl_interop_physarum.py slime-mold: Buffer agents, surface deposit / texture sense, direction-hued veins - gl_interop_clouds.py 3D CUDAArray + tex3D trilinear volumetric raymarch with HG forward-scattering - gl_interop_particles.py VBO interop (from_gl_buffer) + baked curl-noise TextureObject, additive points - gl_interop_bloom.py MipmappedArray get_level + per-level surface downsample + tex2DLod composite, live LOD - gl_interop_jfa_voronoi.py POINT-filtered JFA, AddressMode.BORDER + border_color sentinel - gl_interop_caustics.py UINT8 background sampled LINEAR + MIRROR + sRGB + max_anisotropy, chromatic dispersion Each example documents which cuda.core APIs it uses via a code->API comment map and a live config string in the window caption. All seven are registered (display-gated) in test_basic_examples.py. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: dedup texture/array validation; fix docstring + address_mode type - Extract shared _validate_format_channels / _validate_array_shape helpers in _array.pyx; adopt them in CUDAArray, MipmappedArray, and the texture from_linear/from_pitch2d factories (removes 4x num_channels, 4x format, and 2x shape duplicate validators). - ResourceDescriptor docstring now lists from_mipmapped_array (was 3 of 4 factories). - TextureDescriptor.address_mode annotated AddressMode | tuple[AddressMode, ...] instead of object (CLAUDE.md: avoid Any). - Regenerate .pyi stubs. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: remove dead _context field and orphaned helper The _context field (raw intptr_t) was captured in CUDAArray, MipmappedArray, SurfaceObject, and TextureObject constructors but never read. It could not be safely used to gate destruction (it is not a refcounted context-handle ref like Stream._h_context), so it was pure dead state plus an extra cuCtxGetCurrent per construction. - Drop the _context slot from all four .pxd files (and the now-unused intptr_t cimports). - Drop the self._context assignments and the _get_current_context_ptr cimport from all four .pyx files. - Remove the now-orphaned _get_current_context_ptr helper from cuda_utils (.pyx + .pxd); _tensor_map.pyx keeps its own local copy and is unaffected. _get_current_device_id stays (still used for the .device property). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * caustics improvements * cuda.core: add numba-cuda-mlir port of the Stable Fluids example Add examples/gl_interop_fluid_numba_cuda_mlir.py, a numba-cuda port of gl_interop_fluid.py. numba has no texture/surface support, so the same Stable Fluids solver runs on plain linear device arrays with a hand-written bilinear sampler; the physics and pipeline are unchanged from the CUDA C++ version, only the memory model and read path differ. It targets the MLIR numba-cuda backend (numba-cuda-mlir), which tracks the current cuda.bindings/cuda.core API. Classic numba-cuda lags the editable dev tip (cuda.core.graph relocation per NVIDIA#1858, property-style kernel attributes), so it cannot launch kernels against it. Declare numba-cuda-mlir as a linux-gated pypi dependency of the test feature so `pixi run python examples/...` works in the default env. Upstream ships only manylinux cp311-cp314 wheels (no win-64), hence the linux target gating. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: satisfy ruff on the numba-cuda-mlir fluid example Apply ruff-format and suppress a false-positive RUF046 on the int(math.floor(...)) index conversion: stock numba.cuda's math.floor returns a float (unlike CPython), so the int() cast is required for array indexing. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: isolate numba-cuda-mlir in its own env (fix texture examples in default) numba-cuda-mlir's pypi dependency closure pulls *published* cuda-core / cuda-bindings wheels, which overwrite the editable dev path builds in site-packages and strip the texture/surface modules. That broke every gl_interop_* texture example (e.g. gl_interop_caustics.py) in the default environment. Remove numba-cuda-mlir from the test/default feature and give it a dedicated, isolated `numba-mlir` environment that intentionally uses published cuda-core (not the local-deps path builds). Its [cu13] extra makes the env self-contained (nvvm/cudart/nvrtc from nvidia pypi wheels), so no conda CUDA toolkit or CUDA_HOME wiring is needed. default: dev path cuda.core -> texture/surface examples work numba-mlir: published cuda.core -> gl_interop_fluid_numba_cuda_mlir.py works via `pixi run -e numba-mlir python examples/...` Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: route texture/surface/array lifetime through resource-handle layer The four texture/surface types (CUDAArray, MipmappedArray, TextureObject, SurfaceObject) managed their CUDA resources with raw driver handles, ad-hoc _owning flags, Python-reference keepalives (_parent_ref / _source_ref), and __dealloc__ methods that re-called cu*Destroy and swallowed errors. That diverged from the rest of cuda.core, where every resource-owning type routes lifetime through the C++ std::shared_ptr handle layer (see _cpp/DESIGN.md), re-introducing the GC-ordering, interpreter-shutdown, and double-free hazards that layer exists to prevent. Bring all four into the handle architecture: - _cpp/resource_handles.{hpp,cpp}: add ArrayHandle, MipmappedArrayHandle, TexObjectHandle, SurfObjectHandle. CUtexObject/CUsurfObject are both `unsigned long long` (as is CUdeviceptr), so the latter two are TaggedHandle- wrapped to keep each handle a distinct C++ type (mirrors NVVM/nvJitLink). Boxes embed structural dependencies (mipmap-level -> parent mipmap, texture -> type-erased backing, surface -> backing array); deleters call the matching cu*Destroy with the GIL released; create_* return an empty handle + thread-local error on failure. Add as_cu/as_intptr/as_py overloads and 9 driver function pointers. - _resource_handles.{pxd,pyx}: declare the new aliases, create_* functions, and accessors; populate the new driver pointers from cydriver.__pyx_capi__. - _array/_mipmapped_array/_texture/_surface (.pyx+.pxd): store a *Handle; close() is now self._handle.reset(); all four __dealloc__ removed; _owning and _parent_ref dropped (ownership and the parent dependency live in the box); _source_ref/_texture_desc retained for introspection only. Creation goes through create_*_handle + HANDLE_RETURN(get_last_error()). - tests: rewrite the mip-level keepalive test to assert behavior (a level CUDAArray survives a round-trip copy after its parent MipmappedArray is dropped and GC'd, proving the structural keepalive) and add an idempotent- close guard for all four types. Public API is unchanged. Verified: texture/surface suite 63 passed; full cuda_core suite 3449 passed / 228 skipped / 3 xfailed; cython ABI test passed. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: address NVIDIA#2188 review feedback — copy parity + cuda.core.textures namespace Three items from the NVIDIA#2188 design review (Andy-Jost): A+B. CUDAArray.copy_from / copy_to now match the Buffer copy API: - Accept Stream | GraphBuilder (via Stream_accept), so array copies can be captured into a CUDA graph — previously only a concrete Stream was accepted. - copy_to returns the destination object, for parity with Buffer.copy_to. The internal _copy3d helper now takes an already-coerced Stream. C. Group the texture/surface object model under a new cuda.core.textures namespace instead of the flat cuda.core namespace, before v1.1.0 ships (it is hard to move post-release). Adds cuda/core/textures.py re-exporting CUDAArray, ArrayFormat, MipmappedArray, ResourceDescriptor, TextureDescriptor, TextureObject, SurfaceObject, AddressMode, FilterMode, ReadMode; removes the flat exports from cuda.core.__init__ (grouped-only — no deprecation needed pre-release); migrates the texture/surface examples, texture_sample, and the docs api.rst section to the new namespace. Tests: update the stream-rejection match; add copy_to-returns-dst and a GraphBuilder-capture test for CUDAArray copies. Verified: texture/surface suite 65 passed; full cuda_core suite 3449 passed / 230 skipped / 3 xfailed; all examples byte-compile; cuda.core.textures exports the 10 symbols and cuda.core.CUDAArray is gone. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: document CUDAArray copy-only interop contract (NVIDIA#2188 decision 2) Decision 2 of NVIDIA#2188 resolved to "ship copy_from/copy_to only, and document the copy-only contract." The code shipped; this adds the documentation. State plainly in the CUDAArray class docstring (and the .pyi stub) and the api.rst "Textures and surfaces" section that a CUDAArray has an opaque, hardware-defined layout with no linear device pointer, so it cannot expose __cuda_array_interface__ / DLPack or share memory zero-copy; data is moved in and out only by copying via copy_from / copy_to against a linear Buffer or a host buffer-protocol object, and there is no allocation helper. Docs only; no behavior change. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: scope texture/surface PR to the core API + fluid examples (NVIDIA#2188 decision 7) Per NVIDIA#2188 decision 7, the texture/surface example showcase is orthogonal to the core API and lands in a follow-up. Move the 15 standalone gl_interop_*.py demos (bloom, caustics, clouds, fire, image_show, jfa_voronoi, lenia, mandelbrot, mipmap_lod, ocean, particles, physarum, reaction_diffusion, sdf_volume, texture_filter) onto the feature/cuda-core-texture-examples-467 branch for continued iteration and a future examples PR. Kept here as the representative end-to-end coverage of the new API: - examples/texture_sample.py (canonical sampling smoke test), - examples/gl_interop_fluid.py (cuda.core CUDAArray/TextureObject/SurfaceObject), - examples/gl_interop_fluid_numba_cuda_mlir.py (numba-cuda-mlir port) + its isolated pixi env. Drop the moved examples from the example-test requirements map; gl_interop_fluid and the pre-existing gl_interop_plasma stay wired. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: keep the mipmap/LOD example in the texture/surface PR MipmappedArray is a headline type of this PR but was only exercised by unit tests after the example split. Pull gl_interop_mipmap_lod.py back in so the mipmap pyramid + get_level + trilinear (LINEAR mipmap) sampling with LOD bias path has an end-to-end demonstration; re-wire it in the example-test requirements map (has_display-gated). The other gl_interop_*.py demos stay on feature/cuda-core-texture-examples-467 for the follow-up examples PR. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: fix test_utils parametrize collection error on pytest 8.4+ `@pytest.mark.parametrize("in_arr,", _cpu_array_samples())` has a stray trailing comma in the names string. Older pytest tolerated it, but pytest 8.4+ (used in CI) treats `"in_arr,"` as a request to unpack each argvalue across multiple names — so a 3-element sample array becomes "3 values for 1 name" and the whole test session aborts at collection ("1 error during collection"), failing every GPU test job. Drop the trailing comma so each array sample is one parameter. Pre-existing (introduced in NVIDIA#1894, also on main); fixing here to unblock CI. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: satisfy pre-commit (cython-lint, stubs, formatting) - Drop unused `intptr_t` cimports in _mipmapped_array.pyx / _surface.pyx (flagged by cython-lint; the handle ints now go through as_intptr). - Regenerate the .pyi stubs (stubgen-pyx) to match the texture/surface .pyx after the resource-handle + copy-API changes. - isort/ruff-format fixups (cuda.core.textures import ordering, etc.). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: speed up numba-cuda-mlir fluid kernels (~3.6x FPS) Replace the manual bilinear sampler with direct clamped integer-coordinate loads on every stencil read (divergence, pressure_jacobi, subtract_gradient, curl) and the self-velocity / colorize reads. At integer coordinates the sampler's fractional weights are zero, so it returned one exact texel only after a floor, four clamps, two redundant corner loads, and the blend FMAs; the new load_scalar/load_vec/load_color do just the one clamped load. Output is bit-identical (verified vs a numpy reference); pressure_jacobi, the dominant kernel, drops from 46.4us to 7.07us per launch (6.6x). Split vorticity_confinement into compute_curl (writes curl once to a scratch field) + apply_vorticity (reads it), eliminating the 5x-redundant curl recompute and the register pressure that capped occupancy at 31%. The true bilinear sampler is retained on the genuinely fractional back-traced coordinates in advect_*/splat. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: add numba-classic pixi env for a same-source compiler comparison Mirrors the numba-mlir env but runs the fluid example through the classic numba-cuda compiler (stock numba + llvmlite -> NVVM) on CUDA 13, so the same kernels can be profiled against the MLIR backend to isolate codegen differences. Pinned to Python 3.12 because llvmlite has no 3.14 wheels yet. Regenerates pixi.lock for the new environment. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: rename ArrayHandle -> CUDAArrayHandle; fix CUDAArray article in error msgs Addresses two review nits on NVIDIA#2095: - Rename the C++/Cython resource-handle type alias ArrayHandle to CUDAArrayHandle so it matches the public CUDAArray name (NVIDIA#2188 decision 1). Pure type-alias rename across the .hpp/.cpp and the Cython .pyx/.pxd/.pyi; the lowercase create_array_handle* functions and MipmappedArrayHandle are unchanged. - Fix grammar in the SurfaceObject/TextureObject from_array error messages: "must be an CUDAArray" -> "must be a CUDAArray". Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: dedupe owning CUDAArray-handle creation (NVIDIA#2095 review) create_array_handle(desc) and create_array_handle_owning(arr) duplicated the owning shared_ptr box plus its destroy-on-last-ref deleter, differing only in who calls cuArray3DCreate. Have the descriptor variant allocate and then delegate to create_array_handle_owning so the owning box/deleter is defined in exactly one place. Both entry points remain: allocate-from-descriptor (normal path) and adopt-an-existing-CUarray (graphics interop). cuArray3DCreate stays in the C++ lifetime layer, consistent with _cpp/DESIGN.md. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: move texture/surface types into a cuda.core.texture subpackage (NVIDIA#2095 review) The texture/surface implementation modules (_array, _texture, _surface, _mipmapped_array) previously sat flat at cuda/core/ with textures.py as a re-export aggregator. Per review, mirror the cuda.core.graph layout: relocate the four {.pyx,.pxd,.pyi} into a cuda/core/texture/ subpackage and turn textures.py into texture/__init__.py. The public namespace is now cuda.core.texture (singular), matching CuPy's cupy.cuda.texture which groups the same texture+surface+array set. Intra-subpackage cimports/imports rewired to cuda.core.texture._*; the shared _resource_handles/_memory/_stream/_utils imports stay at cuda/core/ level. Build and packaging need no changes: build_hooks globs cuda/core/**/*.pyx and names modules by path, and packaging uses cuda.core* discovery with recursive package-data. Tests and docs (api.rst) updated to the new namespace. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: rename CUDAArray -> OpaqueArray (and CUDAArrayHandle -> OpaqueArrayHandle) Rename the public texture/surface backing-storage class CUDAArray to OpaqueArray, and the internal C++/Cython handle type CUDAArrayHandle to OpaqueArrayHandle so the handle keeps mirroring the class name (per the comment #1 rationale on NVIDIA#2095). Updates the class, error messages, docstrings, .pyi stubs, tests, and docs (api.rst). Driver types (CUarray, cuArray3DCreate) are unaffected. This supersedes NVIDIA#2188 decision 1 (which had settled on CUDAArray); the public type is now cuda.core.texture.OpaqueArray. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: re-sort __all__ and imports after OpaqueArray rename (pre-commit) ruff isort (I001) and RUF022 require alphabetical ordering; renaming CUDAArray -> OpaqueArray moved it past FilterMode/MipmappedArray, so the __all__ list in texture/__init__.py and the import block in tests/test_texture_surface.py needed re-sorting. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: update texture examples for OpaqueArray / cuda.core.texture rename The examples still imported `from cuda.core.textures import (CUDAArray, ...)`, which broke after the namespace/class rename. test_basic_examples runs texture_sample.py as a subprocess (no system requirement gating it), so the dead import failed "Run cuda.core tests" across the GPU matrix. Renamed cuda.core.textures -> cuda.core.texture and CUDAArray -> OpaqueArray in all four texture examples. Verified: texture_sample.py runs to completion on a GPU (texel-center + half-integer samples verified, exit 0). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * cuda.core: re-sort example imports after OpaqueArray rename (pre-commit) ruff I001 (isort): OpaqueArray sorts after FilterMode where CUDAArray used to sit, so the `from cuda.core.texture import (...)` blocks in gl_interop_fluid.py and texture_sample.py needed reordering. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
1 parent 875f192 commit 1e53cc6

30 files changed

Lines changed: 11276 additions & 31 deletions

cuda_core/cuda/core/__init__.py

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -110,8 +110,12 @@ class _PatchedProperty(metaclass=_PatchedPropMeta):
110110
from cuda.core._tensor_map import TensorMapDescriptor, TensorMapDescriptorOptions
111111

112112
# isort: split
113+
# Texture/surface types live under the cuda.core.texture namespace (not the
114+
# flat cuda.core namespace); import the subpackage so it is available as
115+
# `cuda.core.texture` after `import cuda.core`.
113116
# Must come after the cuda.core._* extension imports above: loading graph
114117
# earlier interacts badly with the merged-wheel __path__ rewrite and leaves
115118
# Graph/GraphBuilder/GraphCompleteOptions/GraphDebugPrintOptions missing from
116119
# cuda.core.graph.
117120
import cuda.core.graph
121+
import cuda.core.texture

cuda_core/cuda/core/_cpp/resource_handles.cpp

Lines changed: 167 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -77,6 +77,16 @@ decltype(&cuLinkDestroy) p_cuLinkDestroy = nullptr;
7777
decltype(&cuGraphicsUnmapResources) p_cuGraphicsUnmapResources = nullptr;
7878
decltype(&cuGraphicsUnregisterResource) p_cuGraphicsUnregisterResource = nullptr;
7979

80+
decltype(&cuArray3DCreate) p_cuArray3DCreate = nullptr;
81+
decltype(&cuArrayDestroy) p_cuArrayDestroy = nullptr;
82+
decltype(&cuMipmappedArrayCreate) p_cuMipmappedArrayCreate = nullptr;
83+
decltype(&cuMipmappedArrayDestroy) p_cuMipmappedArrayDestroy = nullptr;
84+
decltype(&cuMipmappedArrayGetLevel) p_cuMipmappedArrayGetLevel = nullptr;
85+
decltype(&cuTexObjectCreate) p_cuTexObjectCreate = nullptr;
86+
decltype(&cuTexObjectDestroy) p_cuTexObjectDestroy = nullptr;
87+
decltype(&cuSurfObjectCreate) p_cuSurfObjectCreate = nullptr;
88+
decltype(&cuSurfObjectDestroy) p_cuSurfObjectDestroy = nullptr;
89+
8090
// SM resource split (13.1+ — may be null on older drivers/bindings)
8191
#if CUDA_VERSION >= 13010
8292
decltype(&cuDevSmResourceSplit) p_cuDevSmResourceSplit = nullptr;
@@ -1337,6 +1347,163 @@ FileDescriptorHandle create_fd_handle_ref(int fd) {
13371347
#endif
13381348
}
13391349

1350+
// ============================================================================
1351+
// Array / mipmapped-array / texture / surface handles (PR #467)
1352+
// ============================================================================
1353+
1354+
namespace {
1355+
struct ArrayBox {
1356+
CUarray resource;
1357+
// Non-null only for a mipmap-level view: keeps the parent mipmap (the real
1358+
// owner of the level's storage) alive for as long as the level is held.
1359+
MipmappedArrayHandle h_parent;
1360+
};
1361+
1362+
struct MipmappedArrayBox {
1363+
CUmipmappedArray resource;
1364+
};
1365+
1366+
struct TexObjectBox {
1367+
// Tagged so TexObjectHandle is a distinct C++ type from DevicePtrHandle /
1368+
// SurfObjectHandle (all wrap `unsigned long long`).
1369+
TexObjectValue resource;
1370+
// Type-erased backing dependency (OpaqueArrayHandle / MipmappedArrayHandle /
1371+
// DevicePtrHandle). The texture's resource is a union; we only need to keep
1372+
// whichever backing it was built from alive, never to dereference it.
1373+
std::shared_ptr<const void> h_backing;
1374+
};
1375+
1376+
struct SurfObjectBox {
1377+
SurfObjectValue resource;
1378+
OpaqueArrayHandle h_array; // surfaces are always array-backed
1379+
};
1380+
} // namespace
1381+
1382+
OpaqueArrayHandle create_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc) {
1383+
GILReleaseGuard gil;
1384+
CUarray arr;
1385+
if (CUDA_SUCCESS != (err = p_cuArray3DCreate(&arr, &desc))) {
1386+
return {};
1387+
}
1388+
// Allocation and adoption share the same owning lifetime; the only
1389+
// difference is who calls cuArray3DCreate. Delegate so the owning box and
1390+
// its destroy-on-last-ref deleter are defined in exactly one place.
1391+
return create_array_handle_owning(arr);
1392+
}
1393+
1394+
OpaqueArrayHandle create_array_handle_ref(CUarray arr) {
1395+
if (!arr) {
1396+
return {};
1397+
}
1398+
auto box = std::make_shared<const ArrayBox>(ArrayBox{arr, {}});
1399+
return OpaqueArrayHandle(box, &box->resource);
1400+
}
1401+
1402+
OpaqueArrayHandle create_array_handle_owning(CUarray arr) {
1403+
if (!arr) {
1404+
return {};
1405+
}
1406+
auto box = std::shared_ptr<const ArrayBox>(
1407+
new ArrayBox{arr, {}},
1408+
[](const ArrayBox* b) {
1409+
GILReleaseGuard gil;
1410+
p_cuArrayDestroy(b->resource);
1411+
delete b;
1412+
}
1413+
);
1414+
return OpaqueArrayHandle(box, &box->resource);
1415+
}
1416+
1417+
OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level) {
1418+
GILReleaseGuard gil;
1419+
CUarray arr;
1420+
if (CUDA_SUCCESS != (err = p_cuMipmappedArrayGetLevel(&arr, as_cu(h_mip), level))) {
1421+
return {};
1422+
}
1423+
// Non-owning level view: storage belongs to the mipmap. Embed the mipmap
1424+
// handle so the parent outlives this level; the deleter does not destroy.
1425+
auto box = std::shared_ptr<const ArrayBox>(
1426+
new ArrayBox{arr, h_mip},
1427+
[](const ArrayBox* b) { delete b; }
1428+
);
1429+
return OpaqueArrayHandle(box, &box->resource);
1430+
}
1431+
1432+
MipmappedArrayHandle create_mipmapped_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc,
1433+
unsigned int num_levels) {
1434+
GILReleaseGuard gil;
1435+
CUmipmappedArray mip;
1436+
if (CUDA_SUCCESS != (err = p_cuMipmappedArrayCreate(&mip, &desc, num_levels))) {
1437+
return {};
1438+
}
1439+
auto box = std::shared_ptr<const MipmappedArrayBox>(
1440+
new MipmappedArrayBox{mip},
1441+
[](const MipmappedArrayBox* b) {
1442+
GILReleaseGuard gil;
1443+
p_cuMipmappedArrayDestroy(b->resource);
1444+
delete b;
1445+
}
1446+
);
1447+
return MipmappedArrayHandle(box, &box->resource);
1448+
}
1449+
1450+
namespace {
1451+
TexObjectHandle make_tex_object_handle(const CUDA_RESOURCE_DESC& res,
1452+
const CUDA_TEXTURE_DESC& tex,
1453+
std::shared_ptr<const void> h_backing) {
1454+
GILReleaseGuard gil;
1455+
CUtexObject obj;
1456+
if (CUDA_SUCCESS != (err = p_cuTexObjectCreate(&obj, &res, &tex, nullptr))) {
1457+
return {};
1458+
}
1459+
auto box = std::shared_ptr<const TexObjectBox>(
1460+
new TexObjectBox{TexObjectValue{obj}, std::move(h_backing)},
1461+
[](const TexObjectBox* b) {
1462+
GILReleaseGuard gil;
1463+
p_cuTexObjectDestroy(b->resource.raw);
1464+
delete b;
1465+
}
1466+
);
1467+
return TexObjectHandle(box, &box->resource);
1468+
}
1469+
} // namespace
1470+
1471+
TexObjectHandle create_tex_object_handle_array(const CUDA_RESOURCE_DESC& res,
1472+
const CUDA_TEXTURE_DESC& tex,
1473+
const OpaqueArrayHandle& h_backing) {
1474+
return make_tex_object_handle(res, tex, h_backing);
1475+
}
1476+
1477+
TexObjectHandle create_tex_object_handle_mipmap(const CUDA_RESOURCE_DESC& res,
1478+
const CUDA_TEXTURE_DESC& tex,
1479+
const MipmappedArrayHandle& h_backing) {
1480+
return make_tex_object_handle(res, tex, h_backing);
1481+
}
1482+
1483+
TexObjectHandle create_tex_object_handle_linear(const CUDA_RESOURCE_DESC& res,
1484+
const CUDA_TEXTURE_DESC& tex,
1485+
const DevicePtrHandle& h_backing) {
1486+
return make_tex_object_handle(res, tex, h_backing);
1487+
}
1488+
1489+
SurfObjectHandle create_surf_object_handle(const CUDA_RESOURCE_DESC& res,
1490+
const OpaqueArrayHandle& h_backing) {
1491+
GILReleaseGuard gil;
1492+
CUsurfObject obj;
1493+
if (CUDA_SUCCESS != (err = p_cuSurfObjectCreate(&obj, &res))) {
1494+
return {};
1495+
}
1496+
auto box = std::shared_ptr<const SurfObjectBox>(
1497+
new SurfObjectBox{SurfObjectValue{obj}, h_backing},
1498+
[](const SurfObjectBox* b) {
1499+
GILReleaseGuard gil;
1500+
p_cuSurfObjectDestroy(b->resource.raw);
1501+
delete b;
1502+
}
1503+
);
1504+
return SurfObjectHandle(box, &box->resource);
1505+
}
1506+
13401507
// ============================================================================
13411508
// SM resource split wrapper
13421509
// ============================================================================

cuda_core/cuda/core/_cpp/resource_handles.hpp

Lines changed: 128 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,14 @@ struct TaggedHandle {
3636
using NvvmProgramValue = TaggedHandle<nvvmProgram, 0>;
3737
using NvJitLinkValue = TaggedHandle<nvJitLink_t, 1>;
3838

39+
// CUtexObject, CUsurfObject and CUdeviceptr are all `unsigned long long`, so
40+
// shared_ptr<const CUtexObject> et al. would be the *same* C++ type as
41+
// DevicePtrHandle (and each other), collapsing the as_cu/as_intptr/as_py
42+
// overload sets. Tag them to keep each handle type distinct, exactly as the
43+
// NVVM / nvJitLink handles above do.
44+
using TexObjectValue = TaggedHandle<CUtexObject, 2>;
45+
using SurfObjectValue = TaggedHandle<CUsurfObject, 3>;
46+
3947
// ============================================================================
4048
// Thread-local error handling
4149
// ============================================================================
@@ -108,6 +116,17 @@ extern decltype(&cuLinkDestroy) p_cuLinkDestroy;
108116
extern decltype(&cuGraphicsUnmapResources) p_cuGraphicsUnmapResources;
109117
extern decltype(&cuGraphicsUnregisterResource) p_cuGraphicsUnregisterResource;
110118

119+
// Texture / surface / array (PR #467)
120+
extern decltype(&cuArray3DCreate) p_cuArray3DCreate;
121+
extern decltype(&cuArrayDestroy) p_cuArrayDestroy;
122+
extern decltype(&cuMipmappedArrayCreate) p_cuMipmappedArrayCreate;
123+
extern decltype(&cuMipmappedArrayDestroy) p_cuMipmappedArrayDestroy;
124+
extern decltype(&cuMipmappedArrayGetLevel) p_cuMipmappedArrayGetLevel;
125+
extern decltype(&cuTexObjectCreate) p_cuTexObjectCreate;
126+
extern decltype(&cuTexObjectDestroy) p_cuTexObjectDestroy;
127+
extern decltype(&cuSurfObjectCreate) p_cuSurfObjectCreate;
128+
extern decltype(&cuSurfObjectDestroy) p_cuSurfObjectDestroy;
129+
111130
// SM resource split (13.1+ — may be null on older drivers/bindings)
112131
#if CUDA_VERSION >= 13010
113132
extern decltype(&cuDevSmResourceSplit) p_cuDevSmResourceSplit;
@@ -171,6 +190,10 @@ using NvvmProgramHandle = std::shared_ptr<const NvvmProgramValue>;
171190
using NvJitLinkHandle = std::shared_ptr<const NvJitLinkValue>;
172191
using CuLinkHandle = std::shared_ptr<const CUlinkState>;
173192
using FileDescriptorHandle = std::shared_ptr<const int>;
193+
using OpaqueArrayHandle = std::shared_ptr<const CUarray>;
194+
using MipmappedArrayHandle = std::shared_ptr<const CUmipmappedArray>;
195+
using TexObjectHandle = std::shared_ptr<const TexObjectValue>;
196+
using SurfObjectHandle = std::shared_ptr<const SurfObjectValue>;
174197

175198

176199
// ============================================================================
@@ -530,6 +553,61 @@ FileDescriptorHandle create_fd_handle(int fd);
530553
// Create a non-owning file descriptor handle (caller manages the fd).
531554
FileDescriptorHandle create_fd_handle_ref(int fd);
532555

556+
// ============================================================================
557+
// Array / mipmapped-array / texture / surface handle functions (PR #467)
558+
//
559+
// These resources are managed exactly like every other cuda.core resource:
560+
// the owning handle's deleter calls the matching cu*Destroy with the GIL
561+
// released, structural dependencies are embedded in the box (so a backing
562+
// resource always outlives a texture/surface/level built on it), and
563+
// creation returns an empty handle + thread-local error on failure.
564+
// ============================================================================
565+
566+
// Create an owning CUDA array via cuArray3DCreate.
567+
// When the last reference is released, cuArrayDestroy is called automatically.
568+
// Returns empty handle on error (caller must check).
569+
OpaqueArrayHandle create_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc);
570+
571+
// Create a non-owning array handle (references an existing CUarray).
572+
// Use for arrays owned elsewhere (e.g. graphics interop). Never destroyed here.
573+
OpaqueArrayHandle create_array_handle_ref(CUarray arr);
574+
575+
// Create an owning array handle adopting an existing CUarray.
576+
// When the last reference is released, cuArrayDestroy is called automatically.
577+
OpaqueArrayHandle create_array_handle_owning(CUarray arr);
578+
579+
// Create a non-owning handle to a mipmap level via cuMipmappedArrayGetLevel.
580+
// The level CUarray is owned by the mipmap; the parent MipmappedArrayHandle is
581+
// embedded in the box so it outlives the level view. No destroy in the deleter.
582+
// Returns empty handle on error (caller must check).
583+
OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level);
584+
585+
// Create an owning mipmapped array via cuMipmappedArrayCreate.
586+
// When the last reference is released, cuMipmappedArrayDestroy is called.
587+
// Returns empty handle on error (caller must check).
588+
MipmappedArrayHandle create_mipmapped_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc,
589+
unsigned int num_levels);
590+
591+
// Create an owning texture object via cuTexObjectCreate, embedding the backing
592+
// resource handle (array / mipmapped array / linear-or-pitch2d device pointer)
593+
// so the backing always outlives the texture. cuTexObjectDestroy runs in the
594+
// deleter. Returns empty handle on error (caller must check).
595+
TexObjectHandle create_tex_object_handle_array(const CUDA_RESOURCE_DESC& res,
596+
const CUDA_TEXTURE_DESC& tex,
597+
const OpaqueArrayHandle& h_backing);
598+
TexObjectHandle create_tex_object_handle_mipmap(const CUDA_RESOURCE_DESC& res,
599+
const CUDA_TEXTURE_DESC& tex,
600+
const MipmappedArrayHandle& h_backing);
601+
TexObjectHandle create_tex_object_handle_linear(const CUDA_RESOURCE_DESC& res,
602+
const CUDA_TEXTURE_DESC& tex,
603+
const DevicePtrHandle& h_backing);
604+
605+
// Create an owning surface object via cuSurfObjectCreate, embedding the backing
606+
// array handle so it outlives the surface. cuSurfObjectDestroy runs in the
607+
// deleter. Returns empty handle on error (caller must check).
608+
SurfObjectHandle create_surf_object_handle(const CUDA_RESOURCE_DESC& res,
609+
const OpaqueArrayHandle& h_backing);
610+
533611
// ============================================================================
534612
// Overloaded helper functions to extract raw resources from handles
535613
// ============================================================================
@@ -595,6 +673,24 @@ inline CUlinkState as_cu(const CuLinkHandle& h) noexcept {
595673
return h ? *h : nullptr;
596674
}
597675

676+
inline CUarray as_cu(const OpaqueArrayHandle& h) noexcept {
677+
return h ? *h : nullptr;
678+
}
679+
680+
inline CUmipmappedArray as_cu(const MipmappedArrayHandle& h) noexcept {
681+
return h ? *h : nullptr;
682+
}
683+
684+
// CUtexObject / CUsurfObject are integer-valued (like CUdeviceptr); null is 0.
685+
// The raw value lives in the tagged wrapper's `raw` field.
686+
inline CUtexObject as_cu(const TexObjectHandle& h) noexcept {
687+
return h ? h->raw : 0;
688+
}
689+
690+
inline CUsurfObject as_cu(const SurfObjectHandle& h) noexcept {
691+
return h ? h->raw : 0;
692+
}
693+
598694
// as_intptr() - extract handle as intptr_t for Python interop
599695
// Using signed intptr_t per C standard convention and issue #1342
600696
inline std::intptr_t as_intptr(const ContextHandle& h) noexcept {
@@ -661,6 +757,22 @@ inline std::intptr_t as_intptr(const FileDescriptorHandle& h) noexcept {
661757
return h ? static_cast<std::intptr_t>(*h) : -1;
662758
}
663759

760+
inline std::intptr_t as_intptr(const OpaqueArrayHandle& h) noexcept {
761+
return reinterpret_cast<std::intptr_t>(as_cu(h));
762+
}
763+
764+
inline std::intptr_t as_intptr(const MipmappedArrayHandle& h) noexcept {
765+
return reinterpret_cast<std::intptr_t>(as_cu(h));
766+
}
767+
768+
inline std::intptr_t as_intptr(const TexObjectHandle& h) noexcept {
769+
return static_cast<std::intptr_t>(as_cu(h));
770+
}
771+
772+
inline std::intptr_t as_intptr(const SurfObjectHandle& h) noexcept {
773+
return static_cast<std::intptr_t>(as_cu(h));
774+
}
775+
664776
// as_py() - convert handle to Python wrapper object (returns new reference)
665777
#if PY_VERSION_HEX < 0x030D0000
666778
extern "C" int _Py_IsFinalizing(void);
@@ -776,6 +888,22 @@ inline PyObject* as_py(const FileDescriptorHandle& h) noexcept {
776888
return PyLong_FromSsize_t(as_intptr(h));
777889
}
778890

891+
inline PyObject* as_py(const OpaqueArrayHandle& h) noexcept {
892+
return detail::make_py("cuda.bindings.driver", "CUarray", as_intptr(h));
893+
}
894+
895+
inline PyObject* as_py(const MipmappedArrayHandle& h) noexcept {
896+
return detail::make_py("cuda.bindings.driver", "CUmipmappedArray", as_intptr(h));
897+
}
898+
899+
inline PyObject* as_py(const TexObjectHandle& h) noexcept {
900+
return detail::make_py("cuda.bindings.driver", "CUtexObject", as_intptr(h));
901+
}
902+
903+
inline PyObject* as_py(const SurfObjectHandle& h) noexcept {
904+
return detail::make_py("cuda.bindings.driver", "CUsurfObject", as_intptr(h));
905+
}
906+
779907
// ============================================================================
780908
// SM resource split wrapper (13.1+)
781909
//

0 commit comments

Comments
 (0)