Fix #3246: Correctly check .pyx/.pxd dependencies, and fix inline lints

.pre-commit-config.yaml

@@ -219,8 +219,18 @@ repos:
     rev: "35dbc940c2bd4499ebf39aba00ad717e05cec9b5"  # frozen: v0.22.0
     hooks:
       - id: cython-lint
-        args: [--no-pycodestyle]
-        exclude: ^cuda_bindings_(?:12|13)/
+        # Some rules (e.g. "inline" on a function declared in a .pxd) match
+        # declarations in one file against implementations in another, and the
+        # linter only remembers that within a single invocation. Pre-commit
+        # splits filenames across parallel processes and only passes changed
+        # files, so a .pyx edited without its .pxd would escape the check.
+        # Instead, whenever any Cython file under these trees changes, lint
+        # all of them in one process. Generated cuda_bindings_* and
+        # third-party submodules are not listed.
+        args: [--no-pycodestyle, cuda_core/cuda]
+        files: ^cuda_core/cuda
+        pass_filenames: false
+        require_serial: true
         additional_dependencies:
           - Cython==3.2.9
 

cuda_core/cuda/core/_memory/_legacy.py

@@ -12,7 +12,7 @@
     from cuda.core.typing import DevicePointerType
 
 from cuda.core._memory._buffer import Buffer, MemoryResource
-from cuda.core._utils.cuda_utils import (
+from cuda.core._utils.cuda_utils import (  # type: ignore[attr-defined]  # stubgen-pyx omits .pxd cpdef bodies
     _check_driver_error as raise_if_driver_error,
 )
 from cuda.core._utils.cuda_utils import (

cuda_core/cuda/core/_stream.pxd

@@ -2,8 +2,10 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+from libc.stdint cimport uintptr_t
+
 from cuda.bindings cimport cydriver
-from cuda.core._rt cimport ContextHandle, StreamHandle
+from cuda.core._rt cimport ContextHandle, StreamHandle, as_cu
 
 
 cdef class Stream:
@@ -28,6 +30,30 @@ cdef inline int Stream_check_open(Stream self) except -1:
     if not self._h_stream:
         raise RuntimeError("Stream has been closed")
     return 0
-cdef bint Stream_handle_is_default_token(cydriver.CUstream s) noexcept nogil
-cdef bint Stream_is_default_token(Stream self) noexcept nogil
-cdef bint Stream_is_legacy_default_token(Stream self) noexcept nogil
+
+
+cdef inline bint Stream_handle_is_default_token(cydriver.CUstream s) noexcept nogil:
+    """Return True for the raw CU_STREAM_LEGACY and CU_STREAM_PER_THREAD tokens.
+
+    These tokens carry no context of their own; they refer to whatever context
+    is current, so nothing resolved from one may be cached on the object.
+    """
+    cdef uintptr_t h = <uintptr_t>s
+    return h == <uintptr_t>cydriver.CU_STREAM_LEGACY or h == <uintptr_t>cydriver.CU_STREAM_PER_THREAD
+
+
+cdef inline bint Stream_is_default_token(Stream self) noexcept nogil:
+    """Return True when this stream is CU_STREAM_LEGACY or CU_STREAM_PER_THREAD."""
+    return Stream_handle_is_default_token(as_cu(self._h_stream))
+
+
+cdef inline bint Stream_is_legacy_default_token(Stream self) noexcept nogil:
+    """Return True only for CU_STREAM_LEGACY.
+
+    Unlike CU_STREAM_PER_THREAD, the legacy default stream token is rejected
+    outright (CUDA_ERROR_INVALID_VALUE) by cuMemcpyWithAttributesAsync and
+    cuMemcpyBatchAsync; CU_STREAM_PER_THREAD is a real stream to those entry
+    points and is accepted normally. Use this narrower check, not
+    Stream_is_default_token, wherever that distinction matters.
+    """
+    return <uintptr_t>as_cu(self._h_stream) == <uintptr_t>cydriver.CU_STREAM_LEGACY

cuda_core/cuda/core/_stream.pyx

@@ -517,32 +517,6 @@ cpdef Stream default_stream():
     else:
         return LEGACY_DEFAULT_STREAM
 
-cdef inline bint Stream_handle_is_default_token(cydriver.CUstream s) noexcept nogil:
-    """Return True for the raw CU_STREAM_LEGACY and CU_STREAM_PER_THREAD tokens.
-
-    These tokens carry no context of their own; they refer to whatever context
-    is current, so nothing resolved from one may be cached on the object.
-    """
-    cdef uintptr_t h = <uintptr_t>s
-    return h == <uintptr_t>cydriver.CU_STREAM_LEGACY or h == <uintptr_t>cydriver.CU_STREAM_PER_THREAD
-
-
-cdef inline bint Stream_is_default_token(Stream self) noexcept nogil:
-    """Return True when this stream is CU_STREAM_LEGACY or CU_STREAM_PER_THREAD."""
-    return Stream_handle_is_default_token(as_cu(self._h_stream))
-
-
-cdef inline bint Stream_is_legacy_default_token(Stream self) noexcept nogil:
-    """Return True only for CU_STREAM_LEGACY.
-
-    Unlike CU_STREAM_PER_THREAD, the legacy default stream token is rejected
-    outright (CUDA_ERROR_INVALID_VALUE) by cuMemcpyWithAttributesAsync and
-    cuMemcpyBatchAsync; CU_STREAM_PER_THREAD is a real stream to those entry
-    points and is accepted normally. Use this narrower check, not
-    Stream_is_default_token, wherever that distinction matters.
-    """
-    return <uintptr_t>as_cu(self._h_stream) == <uintptr_t>cydriver.CU_STREAM_LEGACY
-
 
 cdef inline int Stream_get_ctx(Stream self, ContextHandle* h_context) except?-1 nogil:
     """Resolve the stream's context handle into ``h_context``."""

cuda_core/cuda/core/_utils/cuda_utils.pxd

@@ -54,10 +54,36 @@ cdef inline int HANDLE_RETURN_NVJITLINK(
 cdef int _get_current_device_id() except? -1
 
 
+# Success sentinels compared against by the inline checks below.
+cdef object _RUNTIME_SUCCESS
+cdef object _NVRTC_SUCCESS
+
+# Out-of-line failure paths for the inline checks below; keeps the error
+# formatting (and its Python-level dependencies) out of every cimporting module.
+cdef int _raise_driver_error(cydriver.CUresult error) except -1 nogil
+cdef int _raise_runtime_error(error) except -1
+cdef int _raise_nvrtc_status_error(error, handle) except -1
+
+
 # TODO: stop exposing these within the codebase?
-cpdef int _check_driver_error(cydriver.CUresult error) except?-1 nogil
-cpdef int _check_runtime_error(error) except?-1
-cpdef int _check_nvrtc_error(error) except?-1
+cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
+    if error == cydriver.CUresult.CUDA_SUCCESS:
+        return 0
+    return _raise_driver_error(error)
+
+
+cpdef inline int _check_runtime_error(error) except?-1:
+    if error == _RUNTIME_SUCCESS:
+        return 0
+    return _raise_runtime_error(error)
+
+
+cpdef inline int _check_nvrtc_error(error, handle=None) except?-1:
+    if error == _NVRTC_SUCCESS:
+        return 0
+    return _raise_nvrtc_status_error(error, handle)
+
+
 cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1
 cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1
 cdef int _raise_nvjitlink_error(cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1

cuda_core/cuda/core/_utils/cuda_utils.pyi

@@ -2,14 +2,15 @@
 
 from typing import Any, Callable, NamedTuple
 
-from cuda.bindings import cydriver
 from cuda.bindings import driver as driver
 from cuda.bindings import nvrtc as nvrtc
 from cuda.bindings import runtime as runtime
 
 _keep_driver_in_stub: driver.CUresult
 _keep_nvrtc_in_stub: nvrtc.nvrtcResult
 _keep_runtime_in_stub: runtime.cudaError_t
+_RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
+_NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
 _fork_warning_checked = False
 
 class CUDAError(Exception):
@@ -52,9 +53,6 @@ class ComputeCapability(NamedTuple):
     minor: int
 
 def cast_to_3_tuple(label: str, cfg: int | tuple[int, ...]) -> tuple[int, int, int]: ...
-def _check_driver_error(error: cydriver.CUresult) -> int: ...
-def _check_runtime_error(error) -> int: ...
-def _check_nvrtc_error(error, handle=None) -> int: ...
 def handle_return(result: tuple[Any, ...], handle: object | None=None) -> Any: ...
 def check_or_create_options(cls: type, options: object, options_description: str='', keep_none: bool=False) -> object:
     """

cuda_core/cuda/core/_utils/cuda_utils.pyx

@@ -162,18 +162,16 @@ cdef int _raise_nvjitlink_error(
     raise exc
 
 
-cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
-cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
+_RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
+_NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
 
 
 cdef inline void _attach_detail(exc, str detail):
     # PEP 678 notes keep the detail separable from the message.
     exc.add_note(detail)
 
 
-cpdef int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
-    if error == cydriver.CUresult.CUDA_SUCCESS:
-        return 0
+cdef int _raise_driver_error(cydriver.CUresult error) except -1 nogil:
     cdef const char* name
     cdef const char* desc
     # A context-scoped helper in the handle layer may have recorded why this
@@ -197,9 +195,7 @@ cpdef int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
         raise exc
 
 
-cpdef int _check_runtime_error(error) except?-1:
-    if error == _RUNTIME_SUCCESS:
-        return 0
+cdef int _raise_runtime_error(error) except -1:
     # `_check_error()` reaches this path only for `runtime.cudaError_t` values.
     # Use the enum name directly because Windows hybrid cudart can lag that table.
     name = error.name
@@ -213,9 +209,7 @@ cpdef int _check_runtime_error(error) except?-1:
     raise CUDAError(f"{name}: {desc}")
 
 
-cpdef int _check_nvrtc_error(error, handle=None) except?-1:
-    if error == _NVRTC_SUCCESS:
-        return 0
+cdef int _raise_nvrtc_status_error(error, handle) except -1:
     err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}"
     if handle is not None:
         _, logsize = nvrtc.nvrtcGetProgramLogSize(handle)

cuda_core/cuda/core/graph/_graph_node.pxd

@@ -3,14 +3,18 @@
 # SPDX-License-Identifier: Apache-2.0
 
 from libc.stddef cimport size_t
+from libc.stdint cimport uintptr_t
 
 from cuda.bindings cimport cydriver
+from cuda.core._memory._buffer cimport Buffer, Buffer_check_open
 from cuda.core._rt cimport (
     GraphHandle,
     GraphNodeHandle,
     OpaqueHandle,
+    as_cu,
     as_intptr,
     graph_node_get_graph,
+    make_opaque_py,
 )
 
 
@@ -32,8 +36,51 @@ cdef inline int GN_check_valid(GraphNode self) except -1:
     return 0
 
 
-cdef OpaqueHandle _resolve_memcpy_operand(
-    object operand, object owner, str side, cydriver.CUdeviceptr* out_ptr) except *
+cdef inline OpaqueHandle _buffer_attachment_owner(Buffer buf, str label):
+    """Copy a Buffer's device-pointer handle into an attachment owner."""
+    Buffer_check_open(buf)
+    # The local is required: Cython permits the DevicePtrHandle -> OpaqueHandle
+    # conversion on assignment, but not directly in a return statement.
+    cdef OpaqueHandle attachment_owner = buf._h_ptr
+    return attachment_owner
+
+
+cdef inline OpaqueHandle _resolve_memcpy_operand(
+        object operand, object owner, str side,
+        cydriver.CUdeviceptr* out_ptr) except *:
+    """Resolve an operand to a pointer and optional attachment owner.
+
+    ``operand`` is a :class:`Buffer` or a raw integer address; its device
+    pointer is written to ``out_ptr``. For a :class:`Buffer` operand, returns an
+    owner that retains the underlying allocation (not the wrapper). For a raw
+    pointer, returns an owner built from ``owner`` (or an empty handle when
+    ``owner`` is ``None``).
+
+    Raises
+    ------
+    ValueError
+        If ``operand`` is a :class:`Buffer` and ``owner`` is not ``None``.
+        If a :class:`Buffer` operand or ``*_owner`` has no active allocation.
+    """
+    cdef Buffer buf
+    cdef OpaqueHandle attachment_owner
+
+    if isinstance(operand, Buffer):
+        if owner is not None:
+            raise ValueError(
+                f"{side}_owner cannot be used when {side} is a Buffer"
+            )
+        buf = operand
+        attachment_owner = _buffer_attachment_owner(buf, side)
+        out_ptr[0] = as_cu(buf._h_ptr)
+        return attachment_owner
+    out_ptr[0] = <cydriver.CUdeviceptr><uintptr_t>operand
+    if owner is None:
+        return OpaqueHandle()
+    if isinstance(owner, Buffer):
+        return _buffer_attachment_owner(owner, f"{side}_owner")
+    return make_opaque_py(owner)
+
 
 cdef cydriver.CUmemorytype _get_memcpy_memory_type(
     cydriver.CUdeviceptr ptr) except *

cuda_core/cuda/core/graph/_graph_node.pyx

@@ -20,7 +20,7 @@ from cuda.bindings cimport cydriver
 from cuda.core._event cimport Event, Event_check_open
 from cuda.core._kernel_arg_handler cimport ParamHolder
 from cuda.core._launch_config cimport LaunchConfig
-from cuda.core._memory._buffer cimport Buffer, Buffer_check_open
+from cuda.core._memory._buffer cimport Buffer
 from cuda.core._memory._location cimport cumemlocation_from_id
 from cuda.core._module cimport Kernel
 from cuda.core.graph._graph_definition cimport (
@@ -914,52 +914,6 @@ cdef inline FreeNode GN_free(GraphNode self, cydriver.CUdeviceptr c_dptr):
     return _registered(FreeNode._create_with_params(create_graph_node_handle(new_node, h_graph), c_dptr))
 
 
-cdef inline OpaqueHandle _buffer_attachment_owner(Buffer buf, str label):
-    """Copy a Buffer's device-pointer handle into an attachment owner."""
-    Buffer_check_open(buf)
-    # The local is required: Cython permits the DevicePtrHandle -> OpaqueHandle
-    # conversion on assignment, but not directly in a return statement.
-    cdef OpaqueHandle attachment_owner = buf._h_ptr
-    return attachment_owner
-
-
-cdef inline OpaqueHandle _resolve_memcpy_operand(
-        object operand, object owner, str side,
-        cydriver.CUdeviceptr* out_ptr) except *:
-    """Resolve an operand to a pointer and optional attachment owner.
-
-    ``operand`` is a :class:`Buffer` or a raw integer address; its device
-    pointer is written to ``out_ptr``. For a :class:`Buffer` operand, returns an
-    owner that retains the underlying allocation (not the wrapper). For a raw
-    pointer, returns an owner built from ``owner`` (or an empty handle when
-    ``owner`` is ``None``).
-
-    Raises
-    ------
-    ValueError
-        If ``operand`` is a :class:`Buffer` and ``owner`` is not ``None``.
-        If a :class:`Buffer` operand or ``*_owner`` has no active allocation.
-    """
-    cdef Buffer buf
-    cdef OpaqueHandle attachment_owner
-
-    if isinstance(operand, Buffer):
-        if owner is not None:
-            raise ValueError(
-                f"{side}_owner cannot be used when {side} is a Buffer"
-            )
-        buf = operand
-        attachment_owner = _buffer_attachment_owner(buf, side)
-        out_ptr[0] = as_cu(buf._h_ptr)
-        return attachment_owner
-    out_ptr[0] = <cydriver.CUdeviceptr><uintptr_t>operand
-    if owner is None:
-        return OpaqueHandle()
-    if isinstance(owner, Buffer):
-        return _buffer_attachment_owner(owner, f"{side}_owner")
-    return make_opaque_py(owner)
-
-
 cdef inline MemsetNode GN_memset(
         GraphNode self, cydriver.CUdeviceptr c_dst, OpaqueHandle dst_owner,
         unsigned int val, unsigned int elem_size,