diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index e141937748a..7bd81c17c52 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -58,6 +58,12 @@ repos: exclude: '(.*pixi\.lock)|(\.git_archival\.txt)|(.*\.patch$)|(^cuda_core/cuda/core/_vendored/)|(^cuda_core/third_party/)' args: ["--fix"] + - id: check-cython-not-none + name: Check cdef-class-annotated def parameters carry 'not None' + entry: python ./toolshed/check_cython_not_none.py + language: python + files: ^cuda_core/cuda/.*\.pyx$ + - id: check-generated-file-seals name: Check generated-file seals entry: python ./toolshed/check_generated_file_seals.py diff --git a/cuda_core/build_hooks.py b/cuda_core/build_hooks.py index 95deb1b44a4..5c8d8770975 100644 --- a/cuda_core/build_hooks.py +++ b/cuda_core/build_hooks.py @@ -538,7 +538,12 @@ def module_names(): "CUDA_CORE_BUILD_MAJOR": int(cuda_major), "CUDA_CORE_BINDINGS_FLOOR": tuple(_floor_for(cuda_major)), } - compiler_directives = {"embedsignature": True, "warn.deprecated.IF": False, "freethreading_compatible": True} + compiler_directives = { + "embedsignature": True, + "warn.deprecated.IF": False, + "freethreading_compatible": True, + "annotation_typing": False, + } _CythonOptions.warning_errors = True if COMPILE_FOR_COVERAGE: compiler_directives["linetrace"] = True diff --git a/cuda_core/cuda/core/_context.pyx b/cuda_core/cuda/core/_context.pyx index 7f5fba9f80f..725dd8eb7fd 100644 --- a/cuda_core/cuda/core/_context.pyx +++ b/cuda_core/cuda/core/_context.pyx @@ -8,8 +8,6 @@ from collections.abc import Sequence from dataclasses import dataclass from typing import TYPE_CHECKING -import cython - from cuda.bindings cimport cydriver from cuda.core._device_resources cimport DeviceResources, SMResource, WorkqueueResource from cuda.core._device_resources import SMResource, WorkqueueResource @@ -101,7 +99,6 @@ cdef class Context: Context_check_open(self) return DeviceResources._init_from_ctx(self._h_context, self._device_id) - @cython.annotation_typing(False) def create_stream(self, options: StreamOptions | None = None) -> Stream: """Create a new stream bound to this green context. diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index 8611cd3d819..d2119100c7b 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -30,7 +30,6 @@ from cuda.core._rt cimport ( ) from cuda.core._stream import IsStreamType, Stream, StreamOptions -from cuda.core._utils.clear_error_support import assert_type from cuda.core._utils.cuda_utils import ( ComputeCapability, CUDAError, @@ -1208,9 +1207,7 @@ class Device: return self._memory_resource @memory_resource.setter - def memory_resource(self, mr: MemoryResource) -> None: - from cuda.core._memory import MemoryResource - assert_type(mr, MemoryResource) + def memory_resource(self, MemoryResource mr not None: MemoryResource) -> None: self._memory_resource = mr @property @@ -1247,7 +1244,7 @@ class Device: def __reduce__(self) -> tuple[object, ...]: return Device, (self.device_id,) - def set_current(self, ctx: Context | None = None) -> Context | None: + def set_current(self, Context ctx: Context | None = None) -> Context | None: """Set device to be used for GPU executions. Initializes CUDA and sets the calling thread to a valid CUDA @@ -1291,8 +1288,6 @@ class Device: cdef Context prev_owned = None if ctx is not None: - # TODO: revisit once Context is cythonized - assert_type(ctx, Context) Context_check_open(ctx) if ctx._device_id != self._device_id: # The CUDA context stack is per-thread, not per-Device-object, diff --git a/cuda_core/cuda/core/_device_resources.pyx b/cuda_core/cuda/core/_device_resources.pyx index 47085a0df51..e4d5809b193 100644 --- a/cuda_core/cuda/core/_device_resources.pyx +++ b/cuda_core/cuda/core/_device_resources.pyx @@ -475,7 +475,7 @@ cdef class SMResource: def split( self, - options: SMResourceOptions, + SMResourceOptions options not None: SMResourceOptions, *, bint dry_run=False ) -> tuple[list[SMResource], SMResource]: @@ -587,7 +587,7 @@ cdef class WorkqueueResource: "WorkqueueResource requires cuda.core to be built with CUDA 13.x bindings" ) - def configure(self, options: WorkqueueResourceOptions) -> None: + def configure(self, WorkqueueResourceOptions options not None: WorkqueueResourceOptions) -> None: """Configure the workqueue resource in place. Parameters diff --git a/cuda_core/cuda/core/_event.pyx b/cuda_core/cuda/core/_event.pyx index 965c4b878eb..99dbba60721 100644 --- a/cuda_core/cuda/core/_event.pyx +++ b/cuda_core/cuda/core/_event.pyx @@ -168,7 +168,7 @@ cdef class Event: def __rsub__(self, other: object): return NotImplemented - def __sub__(self, other: Event) -> float: + def __sub__(self, Event other not None: Event) -> float: # return self - other (in milliseconds) Event_check_open(self) cdef Event other_event = Event_accept(other) @@ -228,7 +228,7 @@ cdef class Event: return self._ipc_descriptor @classmethod - def from_ipc_descriptor(cls, ipc_descriptor: IPCEventDescriptor) -> Event: + def from_ipc_descriptor(cls, IPCEventDescriptor ipc_descriptor not None: IPCEventDescriptor) -> Event: """Import an event that was exported from another process. Parameters @@ -377,7 +377,7 @@ cdef class IPCEventDescriptor: return IPCEventDescriptor._init, (self._reserved, self._is_blocking_sync) -def _reduce_event(event: Event) -> tuple[object, ...]: +def _reduce_event(Event event not None: Event) -> tuple[object, ...]: check_multiprocessing_start_method() return event.from_ipc_descriptor, (event.ipc_descriptor,) diff --git a/cuda_core/cuda/core/_graphics.pyx b/cuda_core/cuda/core/_graphics.pyx index 6e5ea061f50..cd3fe1a0995 100644 --- a/cuda_core/cuda/core/_graphics.pyx +++ b/cuda_core/cuda/core/_graphics.pyx @@ -223,7 +223,7 @@ cdef class GraphicsResource: return None return self._mapped_buffer - def map(self, *, stream: Stream) -> Buffer: + def map(self, *, Stream stream not None: Stream) -> Buffer: """Map this graphics resource for CUDA access. After mapping, a CUDA device pointer into the underlying graphics diff --git a/cuda_core/cuda/core/_launcher.pyx b/cuda_core/cuda/core/_launcher.pyx index e5947b48be2..1eef53beaf3 100644 --- a/cuda_core/cuda/core/_launcher.pyx +++ b/cuda_core/cuda/core/_launcher.pyx @@ -29,8 +29,8 @@ __all__ = ['launch'] def launch( stream: Stream | GraphBuilder | IsStreamType, - config: LaunchConfig, - kernel: Kernel, + LaunchConfig config not None: LaunchConfig, + Kernel kernel not None: Kernel, *kernel_args ) -> None: """Launches a :obj:`~_module.Kernel` @@ -59,8 +59,7 @@ def launch( cdef ParamHolder ker_args = ParamHolder(kernel_args) cdef void** args_ptr = (ker_args.ptr) - cdef Kernel ker = kernel - cdef cydriver.CUfunction func_handle = as_cu(ker._h_kernel) + cdef cydriver.CUfunction func_handle = as_cu(kernel._h_kernel) drv_cfg = conf._to_native_launch_config() drv_cfg.hStream = as_cu(s._h_stream) diff --git a/cuda_core/cuda/core/_layout.pyx b/cuda_core/cuda/core/_layout.pyx index 56f0a775a5a..1935e9969fa 100644 --- a/cuda_core/cuda/core/_layout.pyx +++ b/cuda_core/cuda/core/_layout.pyx @@ -96,7 +96,7 @@ cdef class _StridedLayout: @classmethod def dense_like( cls, - other : _StridedLayout, + _StridedLayout other not None : _StridedLayout, stride_order : str | tuple[int] = "K" ) -> _StridedLayout: """ diff --git a/cuda_core/cuda/core/_memory/_buffer.pyx b/cuda_core/cuda/core/_memory/_buffer.pyx index 4050b56c3ca..23ab10c06e9 100644 --- a/cuda_core/cuda/core/_memory/_buffer.pyx +++ b/cuda_core/cuda/core/_memory/_buffer.pyx @@ -404,8 +404,8 @@ cdef class Buffer: @classmethod def from_ipc_descriptor( cls, mr: DeviceMemoryResource | PinnedMemoryResource | VirtualMemoryResource, - ipc_descriptor: IPCBufferDescriptor, - *, stream: Stream + IPCBufferDescriptor ipc_descriptor not None: IPCBufferDescriptor, + *, Stream stream not None: Stream ) -> Buffer: """Import a buffer that was exported from another process. @@ -505,7 +505,7 @@ cdef class Buffer: self.close() return False - def copy_to(self, dst: Buffer | None = None, *, stream: Stream | GraphBuilder, + def copy_to(self, Buffer dst: Buffer | None = None, *, stream: Stream | GraphBuilder, options: CopyOptions | None = None) -> Buffer: """Copy from this buffer to the dst buffer asynchronously on the given stream. @@ -555,7 +555,7 @@ cdef class Buffer: "buffer does not have a memory_resource)") dst = self._memory_resource.allocate(src_size, stream=s) else: - Buffer_check_open(dst) + Buffer_check_open(dst) cdef size_t dst_size = dst._size if dst_size != src_size: @@ -566,7 +566,7 @@ cdef class Buffer: as_cu(dst._h_ptr), as_cu(self._h_ptr), src_size, s, options, "copy_to") return dst - def copy_from(self, src: Buffer, *, stream: Stream | GraphBuilder, + def copy_from(self, Buffer src not None: Buffer, *, stream: Stream | GraphBuilder, options: CopyOptions | None = None) -> None: """Copy from the src buffer to this buffer asynchronously on the given stream. diff --git a/cuda_core/cuda/core/_memory/_copy_ops.pyx b/cuda_core/cuda/core/_memory/_copy_ops.pyx index 73824e6c1f4..7de5a3f6abf 100644 --- a/cuda_core/cuda/core/_memory/_copy_ops.pyx +++ b/cuda_core/cuda/core/_memory/_copy_ops.pyx @@ -89,7 +89,7 @@ def _normalize_copy_options( def copy_batch( - stream: Stream, + Stream stream not None: Stream, srcs: Sequence[Buffer], dsts: Sequence[Buffer], *, diff --git a/cuda_core/cuda/core/_memory/_device_memory_resource.pyx b/cuda_core/cuda/core/_memory/_device_memory_resource.pyx index 8a0fed73eee..ed481334233 100644 --- a/cuda_core/cuda/core/_memory/_device_memory_resource.pyx +++ b/cuda_core/cuda/core/_memory/_device_memory_resource.pyx @@ -17,7 +17,6 @@ from cuda.core._utils.cuda_utils cimport ( HANDLE_RETURN, ) -import cython from dataclasses import dataclass import multiprocessing import platform # no-cython-lint @@ -144,7 +143,6 @@ cdef class DeviceMemoryResource(_MemPool): def __cinit__(self, *args, **kwargs) -> None: self._dev_id = cydriver.CU_DEVICE_INVALID - @cython.annotation_typing(False) def __init__( self, device_id: Device | int, diff --git a/cuda_core/cuda/core/_memory/_ipc.pyx b/cuda_core/cuda/core/_memory/_ipc.pyx index 5bc35fa84b0..59814314720 100644 --- a/cuda_core/cuda/core/_memory/_ipc.pyx +++ b/cuda_core/cuda/core/_memory/_ipc.pyx @@ -163,7 +163,7 @@ cdef class IPCAllocationHandle: return self._uuid -def _reduce_allocation_handle(alloc_handle: IPCAllocationHandle) -> tuple[object, ...]: +def _reduce_allocation_handle(IPCAllocationHandle alloc_handle not None: IPCAllocationHandle) -> tuple[object, ...]: IPCAllocationHandle_check_open(alloc_handle) check_multiprocessing_start_method() df = multiprocessing.reduction.DupFd(alloc_handle.handle) diff --git a/cuda_core/cuda/core/_memory/_managed_memory_resource.pyx b/cuda_core/cuda/core/_memory/_managed_memory_resource.pyx index 8c47ff487ab..f12be4c246c 100644 --- a/cuda_core/cuda/core/_memory/_managed_memory_resource.pyx +++ b/cuda_core/cuda/core/_memory/_managed_memory_resource.pyx @@ -13,7 +13,6 @@ from cuda.core._utils.cuda_utils cimport HANDLE_RETURN from cuda.core._utils.cuda_utils cimport check_or_create_options # no-cython-lint from cuda.core._utils.cuda_utils import CUDAError # no-cython-lint -import cython from dataclasses import dataclass import threading from typing import TYPE_CHECKING @@ -98,7 +97,6 @@ cdef class ManagedMemoryResource(_MemPool): memory pools. """ - @cython.annotation_typing(False) def __init__(self, options: ManagedMemoryResourceOptions | None = None) -> None: _MMR_init(self, options) diff --git a/cuda_core/cuda/core/_memory/_peer_access_utils.pyi b/cuda_core/cuda/core/_memory/_peer_access_utils.pyi index 83dde54945e..c656fd4220d 100644 --- a/cuda_core/cuda/core/_memory/_peer_access_utils.pyi +++ b/cuda_core/cuda/core/_memory/_peer_access_utils.pyi @@ -78,7 +78,7 @@ def plan_peer_access_update(owner_device_id: int, current_peer_ids: Iterable[int """Compute the peer-access target state and add/remove deltas.""" def _resolve_peer_device_id(value: Device | int | None) -> int: """Coerce ``Device | int`` into a device-ordinal int.""" -def _set_pool_access(mr: object, to_add: tuple[int, ...], to_remove: tuple[int, ...]) -> None: +def _set_pool_access(mr: DeviceMemoryResource, to_add: tuple[int, ...], to_remove: tuple[int, ...]) -> None: """Issue one ``cuMemPoolSetAccess`` for the given add/remove deltas. The thin Python-callable layer that wraps the actual driver call: building diff --git a/cuda_core/cuda/core/_memory/_peer_access_utils.pyx b/cuda_core/cuda/core/_memory/_peer_access_utils.pyx index be3426d8a69..31ed7840985 100644 --- a/cuda_core/cuda/core/_memory/_peer_access_utils.pyx +++ b/cuda_core/cuda/core/_memory/_peer_access_utils.pyx @@ -83,10 +83,9 @@ def _resolve_peer_device_id(value: Device | int | None) -> int: # ---- driver-touching helpers (cdef inline, called from .pyx code) ----------- -cdef inline DeviceMemoryResource _check_peer_access_open(object mr): - cdef DeviceMemoryResource mr_typed = mr - MP_check_open(mr_typed) - return mr_typed +cdef inline DeviceMemoryResource _check_peer_access_open(DeviceMemoryResource mr): + MP_check_open(mr) + return mr cdef inline tuple _query_peer_access_ids(DeviceMemoryResource mr): @@ -132,7 +131,7 @@ cdef inline bint _peer_access_includes(DeviceMemoryResource mr, int dev_id): return flags == cydriver.CUmemAccess_flags.CU_MEM_ACCESS_FLAGS_PROT_READWRITE -def _set_pool_access(mr: object, to_add: tuple[int, ...], to_remove: tuple[int, ...]) -> None: +def _set_pool_access(DeviceMemoryResource mr, to_add: tuple[int, ...], to_remove: tuple[int, ...]) -> None: """Issue one ``cuMemPoolSetAccess`` for the given add/remove deltas. The thin Python-callable layer that wraps the actual driver call: building @@ -143,8 +142,7 @@ def _set_pool_access(mr: object, to_add: tuple[int, ...], to_remove: tuple[int, Preconditions: ``len(to_add) + len(to_remove) > 0`` (the caller is responsible for skipping empty diffs). """ - cdef DeviceMemoryResource mr_typed = mr - MP_check_open(mr_typed) + MP_check_open(mr) cdef size_t count = len(to_add) + len(to_remove) cdef cydriver.CUmemAccessDesc* access_desc = NULL cdef size_t i = 0 @@ -166,13 +164,13 @@ def _set_pool_access(mr: object, to_add: tuple[int, ...], to_remove: tuple[int, i += 1 with nogil: - HANDLE_RETURN(cydriver.cuMemPoolSetAccess(as_cu(mr_typed._h_pool), access_desc, count)) + HANDLE_RETURN(cydriver.cuMemPoolSetAccess(as_cu(mr._h_pool), access_desc, count)) finally: if access_desc != NULL: PyMem_Free(access_desc) -def _apply_peer_access_diff(mr: DeviceMemoryResource, to_add: Iterable[int], to_remove: Iterable[int]) -> None: +def _apply_peer_access_diff(DeviceMemoryResource mr not None: DeviceMemoryResource, to_add: Iterable[int], to_remove: Iterable[int]) -> None: """Apply a peer-access diff in at most one driver call. Every write path on :class:`PeerAccessibleBySetProxy` and the @@ -233,7 +231,7 @@ class PeerAccessibleBySetProxy(MutableSet["Device"]): __slots__ = ("_mr",) - def __init__(self, mr: DeviceMemoryResource) -> None: + def __init__(self, DeviceMemoryResource mr not None: DeviceMemoryResource) -> None: self._mr = mr @classmethod diff --git a/cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx b/cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx index 18621fb08bc..8fe06a254e6 100644 --- a/cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx +++ b/cuda_core/cuda/core/_memory/_pinned_memory_resource.pyx @@ -21,7 +21,6 @@ from cuda.core._utils.cuda_utils cimport ( HANDLE_RETURN, ) -import cython from dataclasses import dataclass import multiprocessing import platform # no-cython-lint @@ -110,7 +109,6 @@ cdef class PinnedMemoryResource(_MemPool): See :class:`DeviceMemoryResource` for more details on IPC usage patterns. """ - @cython.annotation_typing(False) def __init__(self, options: PinnedMemoryResourceOptions | None = None) -> None: _PMR_init(self, options) diff --git a/cuda_core/cuda/core/_memory/_virtual_memory_resource.pyx b/cuda_core/cuda/core/_memory/_virtual_memory_resource.pyx index 9a4bb083922..827677939c1 100644 --- a/cuda_core/cuda/core/_memory/_virtual_memory_resource.pyx +++ b/cuda_core/cuda/core/_memory/_virtual_memory_resource.pyx @@ -653,7 +653,7 @@ cdef class VirtualMemoryResource(MemoryResource): return Buffer_from_deviceptr_handle(h_ptr, aligned, self, None, VirtualMemoryBuffer) def modify_allocation( - self, buf: Buffer, size_t new_size, config: VirtualMemoryResourceOptions | None = None + self, Buffer buf not None: Buffer, size_t new_size, config: VirtualMemoryResourceOptions | None = None ) -> VirtualMemoryBuffer: """ Grow a buffer of this resource to at least ``new_size`` bytes. diff --git a/cuda_core/cuda/core/_memoryview.pyi b/cuda_core/cuda/core/_memoryview.pyi index 6084734d037..ed0c2c2b3f8 100644 --- a/cuda_core/cuda/core/_memoryview.pyi +++ b/cuda_core/cuda/core/_memoryview.pyi @@ -170,7 +170,7 @@ class StridedMemoryView: individual keyword arguments directly, or provide bundled tiled options via ``options=``. """ - def copy_from(self, other: StridedMemoryView, stream: Stream, allocator: object | None=None, blocking: bool | None=None) -> None: + def copy_from(self, other: StridedMemoryView, stream: Stream | None, allocator: object | None=None, blocking: bool | None=None) -> None: """ Copies the data from the other view into this view. diff --git a/cuda_core/cuda/core/_memoryview.pyx b/cuda_core/cuda/core/_memoryview.pyx index 8a6e989044b..dee143330d2 100644 --- a/cuda_core/cuda/core/_memoryview.pyx +++ b/cuda_core/cuda/core/_memoryview.pyx @@ -291,7 +291,7 @@ cdef class StridedMemoryView: @classmethod def from_buffer( cls, - buffer : Buffer, + cyBuffer buffer not None : Buffer, shape : tuple[int, ...], strides : tuple[int, ...] | None = None, *, @@ -418,8 +418,8 @@ cdef class StridedMemoryView: def copy_from( self, - other: StridedMemoryView, - stream: Stream, + StridedMemoryView other not None: StridedMemoryView, + stream: Stream | None, allocator: object = None, blocking: bool | None = None, ) -> None: @@ -454,7 +454,7 @@ cdef class StridedMemoryView: def copy_to( self, - other: StridedMemoryView, + StridedMemoryView other not None: StridedMemoryView, stream: Stream | None = None, allocator: object = None, blocking: bool | None = None, diff --git a/cuda_core/cuda/core/_module.pyx b/cuda_core/cuda/core/_module.pyx index b8dc8ee6a03..d4de653406f 100644 --- a/cuda_core/cuda/core/_module.pyx +++ b/cuda_core/cuda/core/_module.pyx @@ -373,7 +373,7 @@ cdef class KernelOccupancy: )) return dynamic_smem_size - def max_potential_cluster_size(self, config: LaunchConfig, *, stream: Stream) -> int: + def max_potential_cluster_size(self, LaunchConfig config not None: LaunchConfig, *, Stream stream not None: Stream) -> int: """Maximum potential cluster size. The maximum potential cluster size for this kernel and given launch configuration. @@ -392,7 +392,7 @@ cdef class KernelOccupancy: int The maximum cluster size that can be launched for this kernel and launch configuration. """ - cdef cydriver.CUlaunchConfig drv_cfg = (config)._to_native_launch_config() + cdef cydriver.CUlaunchConfig drv_cfg = config._to_native_launch_config() cdef Stream s = Stream_accept(stream) drv_cfg.hStream = as_cu(s._h_stream) cdef int cluster_size @@ -401,7 +401,7 @@ cdef class KernelOccupancy: HANDLE_RETURN(cydriver.cuOccupancyMaxPotentialClusterSize(&cluster_size, func, &drv_cfg)) return cluster_size - def max_active_clusters(self, config: LaunchConfig, *, stream: Stream) -> int: + def max_active_clusters(self, LaunchConfig config not None: LaunchConfig, *, Stream stream not None: Stream) -> int: """Maximum number of active clusters on the target device. The maximum number of clusters that could concurrently execute on the target device. @@ -420,7 +420,7 @@ cdef class KernelOccupancy: int The maximum number of clusters that could co-exist on the target device. """ - cdef cydriver.CUlaunchConfig drv_cfg = (config)._to_native_launch_config() + cdef cydriver.CUlaunchConfig drv_cfg = config._to_native_launch_config() cdef Stream s = Stream_accept(stream) drv_cfg.hStream = as_cu(s._h_stream) cdef int num_clusters @@ -528,7 +528,7 @@ cdef class Kernel: return self.handle @staticmethod - def from_handle(handle, mod: ObjectCode | None = None) -> Kernel: + def from_handle(handle, ObjectCode mod: ObjectCode | None = None) -> Kernel: """Creates a new :obj:`Kernel` object from a kernel handle. Parameters @@ -554,7 +554,7 @@ cdef class Kernel: cdef LibraryHandle h_caller_lib if mod is not None: - h_caller_lib = (mod)._h_library + h_caller_lib = mod._h_library if h_existing_lib and h_caller_lib: if as_cu(h_existing_lib) != as_cu(h_caller_lib): import warnings diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index b5d3038e601..1876885e0cf 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -128,13 +128,11 @@ cdef class Stream: return Stream._from_handle(cls, get_per_thread_stream()) @classmethod - @cython.annotation_typing(False) def _init(cls, obj: IsStreamType | None = None, options: StreamOptions | None = None, - device_id: int | None = None, ctx: Context | None = None) -> Stream: + device_id: int | None = None, Context ctx: Context | None = None) -> Stream: cdef StreamHandle h_stream cdef cydriver.CUstream borrowed cdef ContextHandle h_context - cdef Context context if obj is not None and options is not None: raise ValueError("obj and options cannot be both specified") @@ -148,9 +146,8 @@ cdef class Stream: if ctx is None: raise RuntimeError("A CUDA context is required to create a stream") - context = ctx - Context_check_open(context) - h_context = context._h_context + Context_check_open(ctx) + h_context = ctx._h_context cdef StreamOptions opts = check_or_create_options(StreamOptions, options, "Stream options") nonblocking = opts.nonblocking @@ -161,7 +158,7 @@ cdef class Stream: # TODO: we might want to consider memoizing high/low per CUDA context and avoid this call cdef int high, low cdef cydriver.CUresult res_code - res_code = context_get_stream_priority_range(context._h_context, &high, &low) + res_code = context_get_stream_priority_range(ctx._h_context, &high, &low) HANDLE_RETURN(res_code) cdef int prio if priority is not None: @@ -175,7 +172,7 @@ cdef class Stream: # For green contexts, the C++ layer auto-dispatches to cuGreenCtxStreamCreate, # a 12.5 driver API. cuGreenCtxCreate itself is 12.4. The driver alone # decides availability. The gate lives here, not in C++. - if context.is_green and cy_driver_version() < (12, 5, 0): + if ctx.is_green and cy_driver_version() < (12, 5, 0): raise RuntimeError( "Green context stream creation requires CUDA driver 12.5 or newer " f"(current driver: {'.'.join(map(str, cy_driver_version()))})" diff --git a/cuda_core/cuda/core/_tensor_bridge.pyx b/cuda_core/cuda/core/_tensor_bridge.pyx index c7c3cd5c718..128200971bf 100644 --- a/cuda_core/cuda/core/_tensor_bridge.pyx +++ b/cuda_core/cuda/core/_tensor_bridge.pyx @@ -339,7 +339,7 @@ cpdef int sync_torch_stream(int32_t device_index, def view_as_torch_tensor( object obj, object stream_ptr, - view: StridedMemoryView | None = None + StridedMemoryView view: StridedMemoryView | None = None ) -> StridedMemoryView: """Create/populate a :class:`StridedMemoryView` from a ``torch.Tensor``. @@ -392,7 +392,7 @@ def view_as_torch_tensor( # -- populate StridedMemoryView -- if view is not None: - buf = view + buf = view else: buf = StridedMemoryView.__new__(StridedMemoryView) diff --git a/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyi b/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyi index 850e231a3b0..26b01d2427d 100644 --- a/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyi +++ b/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyi @@ -23,7 +23,7 @@ class AdjacencySetProxy(MutableSet[GraphNode]): def __iter__(self) -> Iterator[GraphNode]: ... def __len__(self) -> int: ... def add(self, value: GraphNode) -> None: ... - def discard(self, value: GraphNode) -> None: ... + def discard(self, value: object) -> None: ... def clear(self) -> None: """Remove all edges in a single driver call.""" def __isub__(self, it: Set[Any]) -> AdjacencySetProxy: diff --git a/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyx b/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyx index 225f75339a1..c784c1e94f5 100644 --- a/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyx +++ b/cuda_core/cuda/core/graph/_adjacency_set_proxy.pyx @@ -15,7 +15,6 @@ from cuda.core._rt cimport ( graph_node_get_graph, ) from cuda.core._utils.cuda_utils cimport HANDLE_RETURN -import cython from collections.abc import Iterable, Iterator, MutableSet, Set from typing import Any, TypeVar @@ -30,7 +29,7 @@ class AdjacencySetProxy(MutableSet[GraphNode]): __slots__ = ("_core",) - def __init__(self, node: GraphNode, bint is_fwd) -> None: + def __init__(self, GraphNode node not None: GraphNode, bint is_fwd) -> None: self._core = _AdjacencySetCore(node, is_fwd) # Used by operators such as &|^ to create non-proxy views when needed. @@ -51,7 +50,7 @@ class AdjacencySetProxy(MutableSet[GraphNode]): def __len__(self) -> int: return (<_AdjacencySetCore>self._core).count() - def add(self, value: GraphNode) -> None: + def add(self, GraphNode value not None: GraphNode) -> None: if not isinstance(value, GraphNode): raise TypeError( f"expected GraphNode, got {type(value).__name__}") @@ -60,8 +59,8 @@ class AdjacencySetProxy(MutableSet[GraphNode]): return (<_AdjacencySetCore>self._core).add_edge(value) - @cython.annotation_typing(False) - def discard(self, value: GraphNode) -> None: + def discard(self, value: object) -> None: + # Non-GraphNode values are a no-op, matching MutableSet semantics. (<_AdjacencySetCore>self._core).check_owner_mutable() if value not in self: return diff --git a/cuda_core/cuda/core/graph/_graph_builder.pyx b/cuda_core/cuda/core/graph/_graph_builder.pyx index 6d62390cb91..ce45b4d272c 100644 --- a/cuda_core/cuda/core/graph/_graph_builder.pyx +++ b/cuda_core/cuda/core/graph/_graph_builder.pyx @@ -735,7 +735,7 @@ cdef class GraphBuilder: ) return GraphCondition._from_handle(int(raw_handle)) - def if_then(self, condition: GraphCondition) -> GraphBuilder: + def if_then(self, GraphCondition condition not None: GraphCondition) -> GraphBuilder: """Adds an if condition branch and returns a new graph builder for it. The resulting if graph will only execute the branch if the @@ -770,7 +770,7 @@ cdef class GraphBuilder: node_params.conditional.ctx = self._get_conditional_context() return GB_cond_with_params(self, node_params)[0] - def if_else(self, condition: GraphCondition) -> tuple[GraphBuilder, GraphBuilder]: + def if_else(self, GraphCondition condition not None: GraphCondition) -> tuple[GraphBuilder, GraphBuilder]: """Adds an if-else condition branch and returns new graph builders for both branches. The resulting if graph will execute the branch if the condition @@ -805,7 +805,7 @@ cdef class GraphBuilder: node_params.conditional.ctx = self._get_conditional_context() return GB_cond_with_params(self, node_params) - def switch(self, condition: GraphCondition, count: int) -> tuple[GraphBuilder, ...]: + def switch(self, GraphCondition condition not None: GraphCondition, count: int) -> tuple[GraphBuilder, ...]: """Adds a switch condition branch and returns new graph builders for all cases. The resulting switch graph will execute the branch whose case index @@ -843,7 +843,7 @@ cdef class GraphBuilder: node_params.conditional.ctx = self._get_conditional_context() return GB_cond_with_params(self, node_params) - def while_loop(self, condition: GraphCondition) -> GraphBuilder: + def while_loop(self, GraphCondition condition not None: GraphCondition) -> GraphBuilder: """Adds a while loop and returns a new graph builder for it. The resulting while loop graph will execute the branch repeatedly at runtime @@ -1279,7 +1279,7 @@ cdef class Graph: @overload def __getitem__(self, node: HostCallbackNode) -> ExecutableHostCallbackNode: ... - def __getitem__(self, node: GraphNode) -> ExecutableGraphNode: + def __getitem__(self, GraphNode node not None: GraphNode) -> ExecutableGraphNode: """Return a view for updating *node* in this executable graph. *node* is a definition node from the graph used to instantiate this @@ -1329,7 +1329,7 @@ cdef class Graph: raise CUDAError(msg) HANDLE_RETURN(err) - def upload(self, stream: Stream) -> None: + def upload(self, Stream stream not None: Stream) -> None: """Uploads the graph in a stream. Parameters @@ -1345,7 +1345,7 @@ cdef class Graph: with nogil: HANDLE_RETURN(cydriver.cuGraphUpload(c_exec, c_stream)) - def launch(self, stream: Stream) -> None: + def launch(self, Stream stream not None: Stream) -> None: """Launches the graph in a stream. Parameters diff --git a/cuda_core/cuda/core/graph/_graph_definition.pyx b/cuda_core/cuda/core/graph/_graph_definition.pyx index 1c02255ba7c..f9254e8c229 100644 --- a/cuda_core/cuda/core/graph/_graph_definition.pyx +++ b/cuda_core/cuda/core/graph/_graph_definition.pyx @@ -190,7 +190,7 @@ cdef class GraphDefinition: dst, value, width, height=height, pitch=pitch, dst_owner=dst_owner ) - def launch(self, config: LaunchConfig, kernel: Kernel, *args) -> KernelNode: + def launch(self, config not None: LaunchConfig, kernel not None: Kernel, *args) -> KernelNode: """Add an entry-point kernel launch node (no dependencies). See :meth:`GraphNode.launch` for full documentation. @@ -229,21 +229,21 @@ cdef class GraphDefinition: """ return self._entry.memcpy(dst, src, size, dst_owner=dst_owner, src_owner=src_owner) - def embed(self, child: GraphDefinition) -> ChildGraphNode: + def embed(self, GraphDefinition child not None: GraphDefinition) -> ChildGraphNode: """Add an entry-point child graph node (no dependencies). See :meth:`GraphNode.embed` for full documentation. """ return self._entry.embed(child) - def record(self, event: Event) -> EventRecordNode: + def record(self, event not None: Event) -> EventRecordNode: """Add an entry-point event record node (no dependencies). See :meth:`GraphNode.record` for full documentation. """ return self._entry.record(event) - def wait(self, event: Event) -> EventWaitNode: + def wait(self, event not None: Event) -> EventWaitNode: """Add an entry-point event wait node (no dependencies). See :meth:`GraphNode.wait` for full documentation. @@ -292,28 +292,28 @@ cdef class GraphDefinition: return GraphCondition._from_handle(c_handle) - def if_then(self, condition: GraphCondition) -> IfNode: + def if_then(self, GraphCondition condition not None: GraphCondition) -> IfNode: """Add an entry-point if-conditional node (no dependencies). See :meth:`GraphNode.if_then` for full documentation. """ return self._entry.if_then(condition) - def if_else(self, condition: GraphCondition) -> IfElseNode: + def if_else(self, GraphCondition condition not None: GraphCondition) -> IfElseNode: """Add an entry-point if-else conditional node (no dependencies). See :meth:`GraphNode.if_else` for full documentation. """ return self._entry.if_else(condition) - def while_loop(self, condition: GraphCondition) -> WhileNode: + def while_loop(self, GraphCondition condition not None: GraphCondition) -> WhileNode: """Add an entry-point while-loop conditional node (no dependencies). See :meth:`GraphNode.while_loop` for full documentation. """ return self._entry.while_loop(condition) - def switch(self, condition: GraphCondition, unsigned int count) -> SwitchNode: + def switch(self, GraphCondition condition not None: GraphCondition, unsigned int count) -> SwitchNode: """Add an entry-point switch conditional node (no dependencies). See :meth:`GraphNode.switch` for full documentation. diff --git a/cuda_core/cuda/core/graph/_graph_node.pyx b/cuda_core/cuda/core/graph/_graph_node.pyx index 1b21400e693..09a52f09884 100644 --- a/cuda_core/cuda/core/graph/_graph_node.pyx +++ b/cuda_core/cuda/core/graph/_graph_node.pyx @@ -215,7 +215,7 @@ cdef class GraphNode: s.clear() s.update(value) - def launch(self, config: LaunchConfig, kernel: Kernel, *args) -> KernelNode: + def launch(self, LaunchConfig config not None: LaunchConfig, Kernel kernel not None: Kernel, *args) -> KernelNode: """Add a kernel launch node depending on this node. Clustered and cooperative launch configurations are not currently @@ -243,7 +243,7 @@ cdef class GraphNode: A new KernelNode representing the kernel launch. """ GN_check_valid(self) - return GN_launch(self, config, kernel, ParamHolder(args)) + return GN_launch(self, config, kernel, ParamHolder(args)) def join(self, *nodes: GraphNode) -> EmptyNode: """Create an empty node that depends on this node and all given nodes. @@ -446,7 +446,7 @@ cdef class GraphNode: self, c_dst, dst_attachment_owner, c_src, src_attachment_owner, size) - def embed(self, child: GraphDefinition) -> ChildGraphNode: + def embed(self, GraphDefinition child not None: GraphDefinition) -> ChildGraphNode: """Add a child graph node depending on this node. Embeds a clone of the given graph definition as a sub-graph node. @@ -463,9 +463,9 @@ cdef class GraphNode: ChildGraphNode A new ChildGraphNode representing the embedded sub-graph. """ - return GN_embed(self, child) + return GN_embed(self, child) - def record(self, event: Event) -> EventRecordNode: + def record(self, Event event not None: Event) -> EventRecordNode: """Add an event record node depending on this node. Parameters @@ -478,9 +478,9 @@ cdef class GraphNode: EventRecordNode A new EventRecordNode representing the event record operation. """ - return GN_record_event(self, event) + return GN_record_event(self, event) - def wait(self, event: Event) -> EventWaitNode: + def wait(self, Event event not None: Event) -> EventWaitNode: """Add an event wait node depending on this node. Parameters @@ -493,7 +493,7 @@ cdef class GraphNode: EventWaitNode A new EventWaitNode representing the event wait operation. """ - return GN_wait_event(self, event) + return GN_wait_event(self, event) def callback(self, fn, *, user_data=None) -> object: """Add a host callback node depending on this node. @@ -545,7 +545,7 @@ cdef class GraphNode: """ return GN_callback(self, fn, user_data) - def if_then(self, condition: GraphCondition) -> IfNode: + def if_then(self, GraphCondition condition not None: GraphCondition) -> IfNode: """Add an if-conditional node depending on this node. The body graph executes only when the condition evaluates to @@ -565,7 +565,7 @@ cdef class GraphNode: self, condition, cydriver.CU_GRAPH_COND_TYPE_IF, 1, IfNode) - def if_else(self, condition: GraphCondition) -> IfElseNode: + def if_else(self, GraphCondition condition not None: GraphCondition) -> IfElseNode: """Add an if-else conditional node depending on this node. Two body graphs: the first executes when the condition is @@ -586,7 +586,7 @@ cdef class GraphNode: self, condition, cydriver.CU_GRAPH_COND_TYPE_IF, 2, IfElseNode) - def while_loop(self, condition: GraphCondition) -> WhileNode: + def while_loop(self, GraphCondition condition not None: GraphCondition) -> WhileNode: """Add a while-loop conditional node depending on this node. The body graph executes repeatedly while the condition @@ -606,7 +606,7 @@ cdef class GraphNode: self, condition, cydriver.CU_GRAPH_COND_TYPE_WHILE, 1, WhileNode) - def switch(self, condition: GraphCondition, unsigned int count) -> SwitchNode: + def switch(self, GraphCondition condition not None: GraphCondition, unsigned int count) -> SwitchNode: """Add a switch conditional node depending on this node. The condition value selects which branch to execute. If the diff --git a/cuda_core/cuda/core/graph/_subclasses.pyx b/cuda_core/cuda/core/graph/_subclasses.pyx index c52bc6f085e..50f09ca8f6b 100644 --- a/cuda_core/cuda/core/graph/_subclasses.pyx +++ b/cuda_core/cuda/core/graph/_subclasses.pyx @@ -1025,7 +1025,7 @@ cdef class ChildGraphNode(GraphNode): return (f"") - def update(self, child: GraphDefinition) -> None: + def update(self, GraphDefinition child not None: GraphDefinition) -> None: """Replace the embedded graph with a clone of ``child``. ``child`` must belong to an independent graph hierarchy. @@ -1096,7 +1096,7 @@ cdef class EventRecordNode(GraphNode): return (f"") - def update(self, event: Event) -> None: + def update(self, Event event not None: Event) -> None: """Replace the event recorded by this node.""" GN_check_valid(self) Event_check_open(event) @@ -1149,7 +1149,7 @@ cdef class EventWaitNode(GraphNode): return (f"") - def update(self, event: Event) -> None: + def update(self, Event event not None: Event) -> None: """Replace the event waited on by this node.""" GN_check_valid(self) Event_check_open(event) @@ -1444,8 +1444,8 @@ cdef class ExecutableKernelNode(ExecutableGraphNode): def update( self, *, - config: LaunchConfig, - kernel: Kernel, + LaunchConfig config not None: LaunchConfig, + Kernel kernel not None: Kernel, args, ) -> None: """Replace all kernel launch parameters for future launches. @@ -1639,7 +1639,7 @@ cdef class ExecutableMemcpyNode(ExecutableGraphNode): cdef class ExecutableChildGraphNode(ExecutableGraphNode): """An executable child-graph-node view.""" - def update(self, child: GraphDefinition) -> None: + def update(self, GraphDefinition child not None: GraphDefinition) -> None: """Replace the embedded graph parameters for future launches.""" GD_check_valid(child) cdef cydriver.CUgraphNodeParams params @@ -1653,7 +1653,7 @@ cdef class ExecutableChildGraphNode(ExecutableGraphNode): cdef class ExecutableEventRecordNode(ExecutableGraphNode): """An executable event-record-node view.""" - def update(self, event: Event) -> None: + def update(self, Event event not None: Event) -> None: """Replace the event recorded by future launches.""" Event_check_open(event) cdef OpaqueHandle event_owner = event._h_event @@ -1668,7 +1668,7 @@ cdef class ExecutableEventRecordNode(ExecutableGraphNode): cdef class ExecutableEventWaitNode(ExecutableGraphNode): """An executable event-wait-node view.""" - def update(self, event: Event) -> None: + def update(self, Event event not None: Event) -> None: """Replace the event waited on by future launches.""" Event_check_open(event) cdef OpaqueHandle event_owner = event._h_event diff --git a/cuda_core/cuda/core/system/_device.pyx b/cuda_core/cuda/core/system/_device.pyx index bcea5ff79dd..76bfb128503 100644 --- a/cuda_core/cuda/core/system/_device.pyx +++ b/cuda_core/cuda/core/system/_device.pyx @@ -1110,7 +1110,7 @@ cdef class Device: return Utilization(nvml.device_get_utilization_rates(self._handle)) -def get_topology_common_ancestor(device1: Device, device2: Device) -> GpuTopologyLevel: +def get_topology_common_ancestor(Device device1 not None: Device, Device device2 not None: Device) -> GpuTopologyLevel: """ Retrieve the common ancestor for two devices. @@ -1136,7 +1136,7 @@ def get_topology_common_ancestor(device1: Device, device2: Device) -> GpuTopolog ] -def get_p2p_status(device1: Device, device2: Device, index: GpuP2PCapsIndex | str) -> GpuP2PStatus: +def get_p2p_status(Device device1 not None: Device, Device device2 not None: Device, index: GpuP2PCapsIndex | str) -> GpuP2PStatus: """ Retrieve the P2P status between two devices. diff --git a/cuda_core/cuda/core/system/_field_values.pxi b/cuda_core/cuda/core/system/_field_values.pxi index d109bc2bc13..65259da3768 100644 --- a/cuda_core/cuda/core/system/_field_values.pxi +++ b/cuda_core/cuda/core/system/_field_values.pxi @@ -93,7 +93,7 @@ cdef class FieldValues: def __init__(self, field_values: nvml.FieldValue): self._field_values = field_values - def __getitem__(self, idx: int) -> FieldValue: + def __getitem__(self, Py_ssize_t idx) -> FieldValue: return FieldValue(self._field_values[idx]) def __len__(self) -> int: diff --git a/cuda_core/cuda/core/system/_system_events.pyx b/cuda_core/cuda/core/system/_system_events.pyx index 865a1b74890..ad1e4bb7847 100644 --- a/cuda_core/cuda/core/system/_system_events.pyx +++ b/cuda_core/cuda/core/system/_system_events.pyx @@ -39,7 +39,7 @@ cdef class SystemEvent: """ cdef object _event_data - def __init__(self, event_data: nvml.SystemEventData_v1): + def __init__(self, event_data not None: nvml.SystemEventData_v1): assert len(event_data) == 1 self._event_data = event_data @@ -71,13 +71,13 @@ cdef class SystemEvents: """ cdef object _event_data - def __init__(self, event_data: nvml.SystemEventData_v1): + def __init__(self, event_data not None: nvml.SystemEventData_v1): self._event_data = event_data def __len__(self) -> int: return len(self._event_data) - def __getitem__(self, idx: int) -> SystemEvent: + def __getitem__(self, Py_ssize_t idx) -> SystemEvent: """ Get the :obj:`~_system_events.SystemEvent` at the specified index. """ diff --git a/cuda_core/tests/graph/test_graph_definition_errors.py b/cuda_core/tests/graph/test_graph_definition_errors.py index 40f181e5db1..5069dacf4b8 100644 --- a/cuda_core/tests/graph/test_graph_definition_errors.py +++ b/cuda_core/tests/graph/test_graph_definition_errors.py @@ -46,6 +46,24 @@ def test_conditional_rejects_non_condition(init_cuda, method, args): getattr(g, method)(*args) +@pytest.mark.agent_authored(model="claude-sonnet-5-5") +@pytest.mark.parametrize("none_arg", ["config", "kernel"]) +@pytest.mark.parametrize("target", ["GraphDefinition", "GraphNode"]) +def test_launch_rejects_none(init_cuda, target, none_arg): + """launch() on a graph or node rejects None for config and kernel.""" + kernel = compile_common_kernels().get_kernel("add_one") + g = GraphDefinition() + if target == "GraphDefinition": + owner = g + else: + owner = g.empty() + args = {"config": LaunchConfig(grid=1, block=1), "kernel": kernel} + args[none_arg] = None + + with pytest.raises(TypeError): + owner.launch(args["config"], args["kernel"]) + + def test_embed_rejects_non_graphdef(init_cuda): """embed() rejects non-GraphDefinition arguments.""" g = GraphDefinition() diff --git a/cuda_core/tests/memory/test_copy_single_options.py b/cuda_core/tests/memory/test_copy_single_options.py index 044c5a4c118..11d7aeea02f 100644 --- a/cuda_core/tests/memory/test_copy_single_options.py +++ b/cuda_core/tests/memory/test_copy_single_options.py @@ -46,6 +46,17 @@ def pinned_mr(): return LegacyPinnedMemoryResource() +@pytest.mark.agent_authored(model="claude-sonnet-5-5") +def test_copy_from_rejects_none_src(single_copy_device, single_copy_stream, pinned_mr): + """copy_from(None) raises TypeError instead of dereferencing a null source.""" + dst = pinned_mr.allocate(SIZE) + + with pytest.raises(TypeError): + dst.copy_from(None, stream=single_copy_stream) + + dst.close() + + @pytest.mark.agent_authored(model="Claude Sonnet 4.6") @thread_unsafe_on_windows def test_options_none_copy_to_data_correct(single_copy_device, single_copy_stream, pinned_mr): diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 2466d966439..a8fe53dab3d 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -493,6 +493,20 @@ def test_launch_cluster_scheduling_policy_smoke(init_cuda): stream.sync() +@pytest.mark.agent_authored(model="claude-sonnet-5-5") +@pytest.mark.parametrize("none_arg", ["config", "kernel"]) +def test_launch_rejects_none(init_cuda, none_arg): + """None for a cdef-class argument raises TypeError instead of reaching C.""" + program = Program('extern "C" __global__ void my_kernel() {}', SourceCodeType.CXX) + ker = program.compile(ObjectCodeFormatType.CUBIN).get_kernel("my_kernel") + config = LaunchConfig(grid=(1, 1, 1), block=(1, 1, 1)) + args = {"config": config, "kernel": ker} + args[none_arg] = None + + with pytest.raises(TypeError): + launch(Device().create_stream(), args["config"], args["kernel"]) + + def test_launch_invalid_values(init_cuda): code = 'extern "C" __global__ void my_kernel() {}' program = Program(code, SourceCodeType.CXX) diff --git a/toolshed/check_cython_not_none.py b/toolshed/check_cython_not_none.py new file mode 100644 index 00000000000..7f3da80b616 --- /dev/null +++ b/toolshed/check_cython_not_none.py @@ -0,0 +1,155 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Reject ``def`` parameters annotated with a bare cdef class but lacking ``not None``. + +Cython lets ``None`` through for a parameter typed as an extension type unless +the parameter says ``not None``, and an annotation Cython cannot resolve to a +type constrains nothing at all. Either way ``None`` can reach C code and crash +instead of raising ``TypeError``. Write one of:: + + def f(Buffer buf not None: Buffer) # Cython knows Buffer + def f(buf not None: Buffer) # Cython does not know Buffer + def f(buf: Buffer | None) # None is a valid value +""" + +import re +import sys +from pathlib import Path +from typing import NamedTuple + +REPO_ROOT = Path(__file__).resolve().parent.parent +CLASS_SEARCH_ROOTS = ("cuda_core/cuda", "cuda_bindings/cuda") + +_CDEF_CLASS = re.compile(r"^\s*cdef\s+(?:public\s+|final\s+)*class\s+(\w+)", re.MULTILINE) +_DEF_START = re.compile(r"^[ \t]*(?:cp)?def\b[^\n(]*\(", re.MULTILINE) +_IDENT = r"[A-Za-z_]\w*" +# [Type ]name[ not None][: annotation][ = default] +_PARAM = re.compile( + rf"""\A\s* + (?P(?:{_IDENT}(?:\.{_IDENT})*\s+)?) + (?P\*{{0,2}}{_IDENT}) + (?P\s+not\s+None)? + (?:\s*:\s*(?P.+?))? + (?:\s*=\s*(?P.+?))? + \s*\Z""", + re.VERBOSE | re.DOTALL, +) +_SINGLE_NAME = re.compile(rf"\A{_IDENT}(?:\.{_IDENT})*\Z") + + +class Violation(NamedTuple): + offset: int # start of the parameter text + name_end: int # where ``not None`` belongs + name: str + typed: bool # already has a leading C-level type, e.g. ``Buffer buf`` + annotation: str + message: str + + +def collect_cdef_classes(root=REPO_ROOT): + names = set() + for search_root in CLASS_SEARCH_ROOTS: + for path in (root / search_root).rglob("*"): + if path.suffix in (".pyx", ".pxd"): + names.update(_CDEF_CLASS.findall(path.read_text(encoding="utf-8"))) + return names + + +def _blank(match): + return re.sub(r"[^\n]", " ", match.group(0)) + + +def _strip_comments(text): + # Blank out docstrings and ``#`` comments (keeping offsets stable) so that + # example code inside them is not mistaken for a signature. + text = re.sub(r'(?s)(?:"""|\'\'\').*?(?:"""|\'\'\')', _blank, text) + return re.sub(r"#[^\n]*", _blank, text) + + +def _split_signature(text, start): + """Yield ``(offset, text)`` for each top-level parameter after ``start``.""" + depth = 0 + param_start = start + pos = start + while pos < len(text): + char = text[pos] + if char in "([{": + depth += 1 + elif char in ")]}": + if depth == 0: + yield param_start, text[param_start:pos] + return + depth -= 1 + elif char == "," and depth == 0: + yield param_start, text[param_start:pos] + param_start = pos + 1 + pos += 1 + + +def _includes_none(annotation): + return re.search(r"\bNone\b", annotation) is not None + + +def _is_overload_stub(clean, def_start): + preceding = clean[:def_start].rstrip().rsplit("\n", 1)[-1].strip() + return preceding == "@overload" + + +def find_violations(text, cdef_classes): + """Return the ``Violation`` records for one source.""" + clean = _strip_comments(text) + violations = [] + for match in _DEF_START.finditer(clean): + if _is_overload_stub(clean, match.start()): + continue + for offset, param in _split_signature(clean, match.end()): + parsed = _PARAM.match(param) + if parsed is None or parsed.group("annotation") is None: + continue + annotation = parsed.group("annotation").strip() + if not _SINGLE_NAME.match(annotation): + continue + if annotation.rsplit(".", 1)[-1] not in cdef_classes: + continue + name = parsed.group("name") + # ``self``/``cls`` are never None, and ``not None`` is not allowed + # on ``*args`` / ``**kwargs``. + if name in ("self", "cls") or name.startswith("*"): + continue + if parsed.group("default") is not None and parsed.group("default").strip() == "None": + message = f"parameter {name!r} defaults to None; annotate it as '{annotation} | None'" + elif parsed.group("notnone") is None: + message = f"parameter {name!r} is annotated with cdef class {annotation!r} but lacks 'not None'" + else: + continue + violations.append( + Violation( + offset=offset + (len(param) - len(param.lstrip())), + name_end=offset + parsed.end("name"), + name=name, + typed=bool(parsed.group("typed").strip()), + annotation=annotation, + message=message, + ) + ) + return violations + + +def main(argv): + cdef_classes = collect_cdef_classes() + failed = False + for arg in argv: + path = Path(arg) + if path.suffix != ".pyx": + continue + text = path.read_text(encoding="utf-8") + for violation in find_violations(text, cdef_classes): + lineno = text.count("\n", 0, violation.offset) + 1 + print(f"{arg}:{lineno}: {violation.message}") + failed = True + return int(failed) + + +if __name__ == "__main__": + sys.exit(main(sys.argv[1:]))