diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 07d494f0463..f049061a54d 100644 --- a/.pre-commit-config.yaml +++ b/.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 diff --git a/cuda_core/cuda/core/_memory/_legacy.py b/cuda_core/cuda/core/_memory/_legacy.py index f3dff33a133..05ebcb557b5 100644 --- a/cuda_core/cuda/core/_memory/_legacy.py +++ b/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 ( diff --git a/cuda_core/cuda/core/_stream.pxd b/cuda_core/cuda/core/_stream.pxd index 1b1c9837eb4..4a7d4510d1a 100644 --- a/cuda_core/cuda/core/_stream.pxd +++ b/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 = s + return h == cydriver.CU_STREAM_LEGACY or h == 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 as_cu(self._h_stream) == cydriver.CU_STREAM_LEGACY diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index 1876885e0cf..4f40cdb9eda 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/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 = s - return h == cydriver.CU_STREAM_LEGACY or h == 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 as_cu(self._h_stream) == 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``.""" diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pxd b/cuda_core/cuda/core/_utils/cuda_utils.pxd index 9b485597912..26495b23477 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pxd +++ b/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 diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyi b/cuda_core/cuda/core/_utils/cuda_utils.pyi index 040f718428b..d8fa035ea93 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyi +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyi @@ -2,7 +2,6 @@ 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 @@ -10,6 +9,8 @@ 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: """ diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyx b/cuda_core/cuda/core/_utils/cuda_utils.pyx index bf0b32c8db7..6d1109c5154 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyx +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyx @@ -162,8 +162,8 @@ 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): @@ -171,9 +171,7 @@ cdef inline void _attach_detail(exc, str detail): 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) diff --git a/cuda_core/cuda/core/graph/_graph_node.pxd b/cuda_core/cuda/core/graph/_graph_node.pxd index b34885fde68..a13854e9a2e 100644 --- a/cuda_core/cuda/core/graph/_graph_node.pxd +++ b/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] = 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 * diff --git a/cuda_core/cuda/core/graph/_graph_node.pyx b/cuda_core/cuda/core/graph/_graph_node.pyx index 09a52f09884..4c81860d39a 100644 --- a/cuda_core/cuda/core/graph/_graph_node.pyx +++ b/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] = 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,