diff --git a/cuda_core/AGENTS.md b/cuda_core/AGENTS.md index 9d80ab74aaa..8b3d9d9b7b9 100644 --- a/cuda_core/AGENTS.md +++ b/cuda_core/AGENTS.md @@ -101,6 +101,81 @@ and agents should flag violations. (kernel arguments, memcpy/memset operands, `dst_owner`/`src_owner`, and host-callback closures) inherit this contract. +## Failure handling + +The user-facing contract lives in `docs/source/error_handling.rst`; the rules +below are for contributors. Reviewers and agents should flag violations. + +- **Raise by default**: any failure on a path where an exception can propagate + raises. Driver statuses go through `HANDLE_RETURN` (Cython) or are returned as + `CUresult` from the C++ handle layer and then `HANDLE_RETURN`ed; never + replace a `CUresult` with a generic `RuntimeError`, and drain + `get_last_error()` immediately after a handle constructor returns empty so a + stale status cannot be misattributed later. +- **Guarantees**: a call that creates a resource must create nothing when it + raises (undo the creation if a later step fails). Every call except + `Device.set_current` must leave the calling thread's current context as it + found it. Do not hand-roll `cuCtxPush/Pop/SetCurrent` sequences in Cython; use + the handle layer's scoped-context helpers (`invoke_in_context`, + `invoke_in_context_or_undo`, `cleanup_in_context`, `context_get_device`, + `graph_node_set_params`) so the failure handling exists in one place. +- **Publish before you raise**: when a driver mutation has succeeded and a later + step can still fail, commit whatever keeps that mutation memory-safe (for + example the graph attachment that retains a node's new owners) before raising + the later error. Rolling back the retention of a live mutation creates a + dangling reference. When ownership cannot be established, retain the + resources anyway (leak) rather than release them; a leak is always preferred + to a use-after-free. +- **Non-propagating paths never raise and never discard a status**: shared_ptr + deleters, `__dealloc__` and CUDA callbacks report through one channel, `report_cuda_error()` / `report_message()` in C++ (the + `pw_*` wrappers) or `warnings.warn(..., CUDAWarning)` in Cython and Python, + which emits `cuda.core.CUDAWarning`. No `print(file=sys.stderr)` and no + `fprintf` outside that helper. `CUDA_ERROR_DEINITIALIZED` is filtered by the + helper because it means the driver is shutting down. +- **Pick the channel by where you are**: a path that can raise uses + `HANDLE_RETURN`; an `except` block whose rollback failed uses + `attach_rollback_failure()`; a deleter or cleanup path uses a `pw_*` + wrapper or `report_cuda_error()`; the same situation in Cython or Python + uses `warnings.warn(..., CUDAWarning)`; a CUDA callback thread does nothing + that needs the GIL and hands its work to the deferred-cleanup queue + (`Py_AddPendingCall` is GIL-free and allowed there). The table in + `_cpp/DESIGN.md` ("Which channel to use") spells this out. +- **`pw_*` runs user Python**: a `p_` pointer only calls the driver; its `pw_` + twin also acquires the GIL on failure and runs the warning filters, + `showwarning`, or `sys.unraisablehook`, any of which may call back into + cuda.core. Never call a `pw_*` wrapper or `report_*` while holding a C++ + lock. Take the GIL as the outermost lock, release it before taking a C++ + lock, and when a lock must stay held call `p_`, keep the status, and report + after the lock is released (`deviceptr_import_ipc` is the model). +- **Rollback failure**: the original exception propagates; the failed rollback + is attached to it with `attach_rollback_failure()` (a PEP 678 note on + Python 3.11+, reported out-of-band on 3.10), or chained with + `raise ... from` when a second exception must be raised. Catching everything + (bare `except:` or `except BaseException:`) is acceptable only for + rollback-then-`raise` blocks, where the rollback must also run for + `KeyboardInterrupt`. +- **Finalization**: once `py_is_finalizing()` is true, do no Python work from + destructors or callbacks and accept the leak (see + `_cpp/resource_handles.hpp` and `_cpp/GRAPH_ATTACHMENTS.md`). +- **Never terminate the process**: no `std::abort`, `std::terminate`, `exit`, + `Py_FatalError`, or `assert` that survives into a release build, anywhere in + `cuda.core`. A failed CUDA call, including a failed context restoration, is + raised or reported. An internal invariant violation is handled the same way: + raise a `RuntimeError` that says "internal cuda.core error, please report" + where an exception can propagate, report through the channel above where it + cannot, and leak the affected resource rather than touch state that may be + inconsistent. Users who want fail-fast behavior get it with + `warnings.filterwarnings("error", category=CUDAWarning)` and + `PYTHONFAULTHANDLER`; the library does not make that choice for them. An + *implicit* abort (an exception escaping a `noexcept` function or a deleter, + including `std::bad_alloc` from an allocation inside `noexcept` code) is a + bug (#1489, #2417), not a policy choice: `noexcept` helpers must not + allocate, or must catch what they call. +- **Testing**: inject restoration failures with + `cuda.core._resource_handles._set_context_restore_fault_for_testing`; assert + reports with `pytest.warns(CUDAWarning)` or `warnings.catch_warnings`, never + by matching stderr text. + ## API design guidelines These are some API design guidelines we try to follow when adding new APIs to diff --git a/cuda_core/cuda/core/__init__.py b/cuda_core/cuda/core/__init__.py index 7864ae794ca..5020f0f2a83 100644 --- a/cuda_core/cuda/core/__init__.py +++ b/cuda_core/cuda/core/__init__.py @@ -102,8 +102,12 @@ class _PatchedProperty(metaclass=_PatchedPropMeta): from cuda.core._stream import __all__ as _stream_all from cuda.core._tensor_map import * from cuda.core._tensor_map import __all__ as _tensor_map_all +from cuda.core._utils.cuda_utils import CUDAError, CUDAWarning, NVRTCError __all__ = [ + "CUDAError", + "CUDAWarning", + "NVRTCError", *_context_all, *_device_all, *_device_resources_all, diff --git a/cuda_core/cuda/core/_cpp/DESIGN.md b/cuda_core/cuda/core/_cpp/DESIGN.md index 6615f21c4ba..4d7f0244463 100644 --- a/cuda_core/cuda/core/_cpp/DESIGN.md +++ b/cuda_core/cuda/core/_cpp/DESIGN.md @@ -226,6 +226,20 @@ Handle destructors may run from any thread. The implementation includes RAII gua The handle API functions are safe to call with or without the GIL held. They will release the GIL (if necessary) before calling CUDA driver API functions. +**The GIL is the outermost lock.** Code that holds a C++ lock (a registry's +mutex, `ipc_import_mutex`, any `std::mutex`) must not acquire or reacquire the +GIL while the lock is held: no `report_*` or `pw_*` calls, no +`GILAcquireGuard`, and no `GILReleaseGuard` whose destructor runs inside the +locked region. Code that needs a C++ lock and may run with the GIL held +releases the GIL first (`GILReleaseGuard` before `lock_guard`). Otherwise a +thread blocked on the lock while holding the GIL deadlocks with the lock holder +waiting for the GIL (#2840). Collect statuses under the lock and report after it +is released, as `deviceptr_import_ipc` does: `cleanup_in_context` takes an +`after_cleanup` hook that runs once the cleanup is done and before anything that +may run user code, and the deleter passes one that unlocks its +`std::unique_lock`. The registries store `weak_ptr`s, +so erasing an entry under a registry lock never runs a deleter. + ### Static Initialization and Deadlock Hazards When writing C++ code that interacts with Python, a subtle deadlock can occur @@ -275,6 +289,95 @@ Related functions: - `peek_last_error()`: Returns the error without clearing it - `clear_last_error()`: Clears the error state +The C++ layer never raises Python exceptions: it runs `nogil` and `noexcept`, +and is called from deleters, CUDA callbacks and GIL-released code where raising +is impossible. Status is turned into `CUDAError` in one place, `HANDLE_RETURN` +in the Cython layer. Which status convention a function uses is decided by its +return value. Factories return the handle, so their status goes to thread-local +`err` and is read with `get_last_error()`. Functions that do not produce a +handle (`context_synchronize`, `context_get_device`, `graph_node_set_params`, +the `graph_*_attachment` family, `deviceptr_alloc_raw`) return the `CUresult` +directly and deliver results through out-parameters, mirroring the driver API; +their callers `HANDLE_RETURN` the value. The two conventions never mix. + +### Context-scoped operations + +Operations that must run in a specific context use `invoke_in_context` / +`invoke_in_context_or_undo` (propagating paths) and `cleanup_in_context` +(deleters). They switch the current context, run the operation, and restore the +caller's context. `cleanup_in_context` emits its reports only after that +restoration, so the user code a `CUDAWarning` runs (filters, `showwarning`) +observes the caller's context. When restoration fails after the operation +succeeded, the creation is undone and the restoration status is returned. When both fail, the +operation status is returned. Either way the helper records a thread-local +detail keyed to the returned status (`take_last_error_detail(status)`) that +`_check_driver_error` attaches to the raised `CUDAError` as a PEP 678 note +(appended to the message on Python 3.10), so the user learns that the caller's +context was not restored, which context is current and, for a double failure, +why restoration failed. Keying the detail to its status narrows, but does not +remove, misattribution: a caller that drops the status (an empty handle raised +as a generic error) leaves the detail behind, and a later error on the same +thread with the same status code picks it up. `enter_context` clears stale +detail at the next context-scoped operation. Issue #2760 removes this +thread-local state in favor of explicit status returns. Tests inject restoration failures with +`set_context_restore_fault_for_testing()`. + +### Reporting from non-propagating paths + +Deleters and CUDA callbacks cannot raise. They report through +`report_cuda_error()` / `report_message()` (the `pw_*` wrappers decorate +destroy calls with it and name the resource handle in the message, so Python's +warning registry does not collapse independent failures of one call), which emit a `cuda.core.CUDAWarning` through +the Python warnings machinery when the interpreter is usable, deliver an +escalated warning as an unraisable exception, and fall back to stderr when the +GIL cannot be taken (for example during finalization). `CUDA_ERROR_DEINITIALIZED` +is never reported because it means the driver is shutting down. No status is +discarded silently anywhere in this layer, and nothing in this layer may +terminate the process; see `docs/source/error_handling.rst` and the "Failure handling" +section of `AGENTS.md` for the policy. + +A rollback that fails inside a Cython `except` block is not a non-propagating +path: `attach_rollback_failure()` attaches it as a note to the exception being +handled (`PyErr_GetHandledException`, Python 3.11+) and falls back to a report +only when there is no such exception or notes are unavailable. + +### Which channel to use + +Pick the channel by where the failure happens. Every failure goes through +exactly one of these; none is ever dropped. + +| Where you are | Use | Result | +|---|---|---| +| Cython, on a path that can raise | `HANDLE_RETURN(status)` | Raises `CUDAError`. A restoration detail recorded by the C++ helper becomes a note on the exception. | +| Cython, after a handle constructor returned an empty handle | `HANDLE_RETURN(get_last_error())`, immediately | Same. Transitional: #2760 makes constructors return the status instead. | +| C++, a helper that runs an operation in another context | Return the `CUresult`; `exit_context` records the restoration detail | Cython raises it. Transitional: #2760 returns the restoration status as a second out-parameter. | +| Cython, inside an `except` block whose rollback failed | `attach_rollback_failure(op, status, detail)` | Adds a note to the exception being handled. Reports instead if nothing is being handled or notes do not exist (Python 3.10). | +| C++, a deleter or deferred cleanup | A `pw_*` wrapper, or `report_cuda_error()` / `report_message()` | Emits `CUDAWarning`. Never raises. | +| Cython or Python, a `__dealloc__` or destructor-path callback | `warnings.warn(msg, CUDAWarning, stacklevel=2)` | Same. | +| A CUDA callback thread | Nothing that needs the GIL. Hand the work to the deferred-cleanup queue with `Py_AddPendingCall` | CUDA forbids driver calls there, and acquiring the GIL there can deadlock with a GIL holder blocked in a driver call. GIL-free C API that only schedules work is fine. | + +### `p_` versus `pw_` + +A `p_` function pointer calls the driver and nothing else. Its `pw_` twin calls +the driver and, if the call fails, acquires the GIL and runs Python: the warning +filters, `showwarning`, or `sys.unraisablehook`. Any of those can be user code, +and user code can call back into cuda.core. This is the one place where the +handle layer runs code it does not control, and it is the entry point through +which a thread holding a C++ lock can deadlock (see "GIL Management"). + +Python exceptions raised by that code never become C++ exceptions: the C API +reports them as return codes, and `report_message` hands them to +`sys.unraisablehook`. Nothing on the report path may allocate or throw, since a +deleter is `noexcept`. + +So: use `pw_` only in deleters and cleanup paths that hold no C++ lock and have +finished updating the layer's own state. Where a lock must stay held, call +`p_`, keep the status, and report after the lock is released, as +`deviceptr_import_ipc` does. CUDA callback threads need no extra rule for +`pw_`: the driver call is forbidden there, so the wrapper is too. The general +rule for those threads is no GIL and no Python objects; GIL-free scheduling +calls such as `Py_AddPendingCall` are how work leaves them. + ## Usage from Cython ```cython diff --git a/cuda_core/cuda/core/_cpp/resource_handles.cpp b/cuda_core/cuda/core/_cpp/resource_handles.cpp index 132822d65a7..55e184dfa84 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.cpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.cpp @@ -45,6 +45,8 @@ decltype(&cuCtxGetCurrent) p_cuCtxGetCurrent = nullptr; decltype(&cuCtxSetCurrent) p_cuCtxSetCurrent = nullptr; decltype(&cuCtxSynchronize) p_cuCtxSynchronize = nullptr; decltype(&cuCtxGetStreamPriorityRange) p_cuCtxGetStreamPriorityRange = nullptr; +decltype(&cuCtxGetDevice) p_cuCtxGetDevice = nullptr; +decltype(&cuGraphNodeSetParams) p_cuGraphNodeSetParams = nullptr; decltype(&cuGreenCtxCreate) p_cuGreenCtxCreate = nullptr; decltype(&cuGreenCtxDestroy) p_cuGreenCtxDestroy = nullptr; decltype(&cuCtxFromGreenCtx) p_cuCtxFromGreenCtx = nullptr; @@ -197,14 +199,208 @@ class GILAcquireGuard { bool acquired_; }; -void warn_on_cuda_error(const char* operation, CUresult status, const char* detail = nullptr) noexcept; +// ---------------------------------------------------------------------------- +// Non-propagating error reporting +// +// Deleters, CUDA callbacks and other non-propagating paths cannot raise. They +// report through report_cuda_error()/report_message(), which emit a +// cuda.core.CUDAWarning when the interpreter is usable and fall back to stderr +// otherwise. See docs/source/error_handling.rst for the policy. +// ---------------------------------------------------------------------------- + +// Warning category registered by _resource_handles.pyx (cuda.core.CUDAWarning). +std::atomic warning_category{nullptr}; + +// Thread-local detail attached to the next raised CUDAError with a matching +// status (see take_last_error_detail()). Written only by propagating helpers. +// The taken copy stays valid until the next take on the same thread. +thread_local char last_error_detail[512] = {0}; +thread_local char taken_error_detail[512] = {0}; +thread_local CUresult last_error_detail_status = CUDA_SUCCESS; + +// Thread-local fault injected into the next context restoration (tests only). +thread_local CUresult context_restore_fault = CUDA_SUCCESS; + +// Format " : : " for a failed CUDA call. +void format_cuda_error(char* buffer, size_t size, const char* operation, CUresult status, + const char* detail) noexcept { + const char* error_name = nullptr; + const char* error_description = nullptr; + bool decoded = p_cuGetErrorName && p_cuGetErrorString + && p_cuGetErrorName(status, &error_name) == CUDA_SUCCESS + && p_cuGetErrorString(status, &error_description) == CUDA_SUCCESS; + const char* outcome = detail ? detail : "failed"; + if (decoded) { + std::snprintf(buffer, size, "%s %s: %s: %s", operation, outcome, error_name, error_description); + } else { + std::snprintf(buffer, size, "%s %s (CUDA error %d)", operation, outcome, static_cast(status)); + } +} + +// Bits of a resource handle for diagnostics: a pointer as its address, an +// integer handle (CUdeviceptr, CUtexObject, ...) as its value, and a pointer +// to a handle (nvrtcDestroyProgram(&prog) and friends) as the handle it points +// to. +template +unsigned long long handle_bits(const T& value) noexcept { + using U = std::remove_cv_t>; + if constexpr (std::is_pointer_v) { + using P = std::remove_cv_t>; + if constexpr (std::is_pointer_v

) { + return value ? handle_bits(*value) : 0ull; + } else { + return static_cast(reinterpret_cast(value)); + } + } else if constexpr (std::is_integral_v || std::is_enum_v) { + return static_cast(value); + } else { + return 0ull; + } +} + +// "()": naming the resource keeps independent failures of +// the same call distinct. Python's warning registry collapses repeated warnings +// with identical text from one call site, so without the handle only the first +// of several leaked resources would be reported. +void format_operation(char* buffer, size_t size, const char* operation, + unsigned long long handle) noexcept { + std::snprintf(buffer, size, "%s(%#llx)", operation, handle); +} + +} // namespace + +// Report a message that could not be raised. Emits cuda.core.CUDAWarning via +// the Python warnings machinery; if that itself fails (for example because the +// warning was promoted to an error), the failure is written as an unraisable +// exception, the CPython convention for exceptions in destructors. Falls back +// to stderr when the interpreter cannot be used. +void report_message(const char* message) noexcept { + PyObject* category = warning_category.load(std::memory_order_acquire); + if (category && Py_IsInitialized() && !py_is_finalizing()) { + GILAcquireGuard gil; + if (gil.acquired()) { + // Deleters can run while a Python exception is propagating; keep it. +#if PY_VERSION_HEX >= 0x030C0000 + PyObject* pending = PyErr_GetRaisedException(); +#else + PyObject *pending_type, *pending_value, *pending_tb; + PyErr_Fetch(&pending_type, &pending_value, &pending_tb); +#endif + if (PyErr_WarnEx(category, message, 1) != 0) { + PyObject* subject = PyUnicode_FromString(message); + PyErr_WriteUnraisable(subject); + Py_XDECREF(subject); + } +#if PY_VERSION_HEX >= 0x030C0000 + PyErr_SetRaisedException(pending); +#else + PyErr_Restore(pending_type, pending_value, pending_tb); +#endif + return; + } + } + std::fprintf(stderr, "%s\n", message); + std::fflush(stderr); +} + +// Report a failed non-CUDA call (NVRTC, NVVM, nvJitLink) from a path that +// cannot raise. +void report_status_code(const char* operation, long code) noexcept { + char message[256]; + std::snprintf(message, sizeof(message), "%s failed (status %ld)", operation, code); + report_message(message); +} + +void register_warning_category(PyObject* category) noexcept { + warning_category.store(category, std::memory_order_release); +} + +// Report a failed CUDA call from a path that cannot raise. CUDA_ERROR_DEINITIALIZED +// is not reported: it means the driver is shutting down, which makes cleanup +// failures expected and uninteresting. +void report_cuda_error(const char* operation, CUresult status, const char* detail) noexcept { + if (status == CUDA_SUCCESS || status == CUDA_ERROR_DEINITIALIZED) { + return; + } + char message[512]; + format_cuda_error(message, sizeof(message), operation, status, detail); + report_message(message); +} + +namespace { + +// Attach `message` as a PEP 678 note to the exception currently being handled. +// Returns false when there is none or the interpreter cannot be used. +bool add_note_to_handled_exception(const char* message) noexcept { +#if PY_VERSION_HEX >= 0x030B0000 + if (!Py_IsInitialized() || py_is_finalizing()) { + return false; + } + GILAcquireGuard gil; + if (!gil.acquired()) { + return false; + } + PyObject* exc = PyErr_GetHandledException(); + if (!exc) { + return false; + } + PyObject* result = PyObject_CallMethod(exc, "add_note", "s", message); + Py_DECREF(exc); + if (!result) { + PyErr_Clear(); + return false; + } + Py_DECREF(result); + return true; +#else + (void)message; + return false; +#endif +} + +} // namespace + +void attach_rollback_failure(const char* operation, CUresult status, const char* detail) noexcept { + if (status == CUDA_SUCCESS || status == CUDA_ERROR_DEINITIALIZED) { + return; + } + char message[512]; + format_cuda_error(message, sizeof(message), operation, status, detail); + if (!add_note_to_handled_exception(message)) { + report_message(message); + } +} + +const char* take_last_error_detail(CUresult status) noexcept { + if (!last_error_detail[0] || status != last_error_detail_status) { + return nullptr; + } + std::memcpy(taken_error_detail, last_error_detail, sizeof(taken_error_detail)); + clear_last_error_detail(); + return taken_error_detail; +} + +void clear_last_error_detail() noexcept { + last_error_detail[0] = 0; + last_error_detail_status = CUDA_SUCCESS; +} + +void set_context_restore_fault_for_testing(CUresult status) noexcept { + context_restore_fault = status; +} + +namespace { // Make a context current and record the state needed to restore it. // An empty handle is a no-op: the operation runs in the caller's current -// context, and nothing is restored on exit. +// context, and nothing is restored on exit. invoke_in_context and +// invoke_in_context_or_undo reject empty handles before getting here; only +// graph_node_set_params relies on the no-op (pre-13.2 node updates run in the +// caller's context). CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* changed) noexcept { *previous = nullptr; *changed = 0; + clear_last_error_detail(); CUcontext target = as_cu(h_context); if (!target) { return CUDA_SUCCESS; @@ -220,16 +416,59 @@ CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* return status; } -// Restore the previous context and preserve an earlier operation error. -CUresult exit_context(CUcontext previous, int changed, CUresult operation_status) noexcept { - CUresult restore_status = CUDA_SUCCESS; - if (changed) { - GILReleaseGuard gil; - restore_status = p_cuCtxSetCurrent(previous); +// Restore the caller's context. Returns the restoration status. +CUresult restore_context(CUcontext previous) noexcept { + if (context_restore_fault != CUDA_SUCCESS) { + // Test hook: behave as if cuCtxSetCurrent(previous) failed, leaving the + // target context current exactly as a real failure would. + CUresult fault = context_restore_fault; + context_restore_fault = CUDA_SUCCESS; + return fault; } - if (operation_status != CUDA_SUCCESS && restore_status != CUDA_SUCCESS) { - warn_on_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restore_status); + GILReleaseGuard gil; + return p_cuCtxSetCurrent(previous); +} + +// Record that the caller's context was not restored as the detail of the +// CUresult about to be returned and raised: the operation status if the +// operation failed too, else the restoration status. For a double failure the +// detail also names the restoration error, which the raised error does not. +void note_context_not_restored(CUcontext previous, CUresult operation_status, + CUresult restore_status) noexcept { + CUcontext current = nullptr; + if (p_cuCtxGetCurrent(¤t) != CUDA_SUCCESS) { + current = nullptr; + } + char cause[128] = {0}; + if (operation_status != CUDA_SUCCESS) { + const char* error_name = nullptr; + if (p_cuGetErrorName && p_cuGetErrorName(restore_status, &error_name) == CUDA_SUCCESS) { + std::snprintf(cause, sizeof(cause), " after this failure (cuCtxSetCurrent: %s)", error_name); + } else { + std::snprintf(cause, sizeof(cause), " after this failure (cuCtxSetCurrent: CUDA error %d)", + static_cast(restore_status)); + } } + std::snprintf(last_error_detail, sizeof(last_error_detail), + "the calling thread's CUDA context (%#llx) could not be restored%s; " + "context %#llx is now current. Call Device.set_current() before issuing " + "further CUDA work on this thread", + static_cast(reinterpret_cast(previous)), + cause, + static_cast(reinterpret_cast(current))); + last_error_detail_status = operation_status != CUDA_SUCCESS ? operation_status : restore_status; +} + +// Restore the previous context and preserve an earlier operation error. The +// operation error, if any, is returned; otherwise the restoration status is. +// Either way a restoration failure is recorded as the detail of the returned +// status, so the eventual CUDAError explains it (see take_last_error_detail()). +CUresult exit_context(CUcontext previous, int changed, CUresult operation_status) noexcept { + CUresult restore_status = changed ? restore_context(previous) : CUDA_SUCCESS; + if (restore_status == CUDA_SUCCESS) { + return operation_status; + } + note_context_not_restored(previous, operation_status, restore_status); return operation_status != CUDA_SUCCESS ? operation_status : restore_status; } @@ -257,12 +496,11 @@ ContextHandle deallocation_context(const DeallocationStream& stream) noexcept { } if (stream.ptds_tid != std::thread::id{} && stream.ptds_tid != std::this_thread::get_id()) { - std::fprintf( - stderr, - "Warning: Buffer deallocation for a per-thread default stream " + report_message( + "Buffer deallocation for a per-thread default stream " "is running on a different host thread than the one that recorded " "the deallocation stream; ordering relative to the allocating " - "thread's PTDS is not preserved\n"); + "thread's PTDS is not preserved"); } return get_stream_context(stream.h_stream); } @@ -313,7 +551,7 @@ CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operatio if (undo_ok) { std::invoke(std::forward(undo)); } else { - warn_on_cuda_error( + report_cuda_error( "cuCtxSetCurrent (restoring the caller's context)", composite, "failed; cleanup of the new resource skipped because its context " "is no longer current (resource leaked)"); @@ -322,81 +560,110 @@ CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operatio return composite; } -// Write a warning that includes the CUDA error name and description. -void warn_on_cuda_error(const char* operation, CUresult status, const char* detail) noexcept { - const char* error_name = nullptr; - const char* error_description = nullptr; - CUresult name_status = p_cuGetErrorName(status, &error_name); - CUresult description_status = p_cuGetErrorString(status, &error_description); - - if (name_status == CUDA_SUCCESS && description_status == CUDA_SUCCESS) { - if (detail) { - std::fprintf(stderr, "Warning: %s %s: %s: %s\n", - operation, detail, error_name, error_description); - } else { - std::fprintf(stderr, "Warning: %s failed: %s: %s\n", - operation, error_name, error_description); - } - } else { - if (detail) { - std::fprintf(stderr, "Warning: %s %s (CUDA error %d)\n", - operation, detail, static_cast(status)); - } else { - std::fprintf(stderr, "Warning: %s failed (CUDA error %d)\n", - operation, static_cast(status)); - } - } -} - -// Run cleanup with the requested context current. Warn and skip the operation -// if activation fails, and independently warn on operation or restoration -// failure. Return the operation or activation status; restoration never -// changes the return value. -template +// Run cleanup with the requested context current, restore the caller's +// context, call `after_cleanup`, then report any failure (activation or +// operation failure first, then restoration failure). `after_cleanup` marks +// the point from which code we do not control may run: the report emits a +// CUDAWarning, which acquires the GIL and runs user code, and whatever the +// caller does next may do the same. It is called unconditionally, so a lock +// the cleanup had to run under is released at one fixed point regardless of +// outcome; pass a hook that unlocks it. Returns the operation or activation +// status; restoration never changes it. +template CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, - Fn&& operation, Args&&... args) noexcept { - ASSERT_NOTHROW_INVOCABLE(Fn&&, Args&&...); + unsigned long long handle, Fn&& operation, + AfterCleanup&& after_cleanup) noexcept { + ASSERT_NOTHROW_INVOCABLE(Fn&&); + ASSERT_NOTHROW_INVOCABLE(AfterCleanup&&); CUcontext previous = nullptr; int changed = 0; + const char* detail = nullptr; CUresult status = enter_context(h_context, &previous, &changed); if (status != CUDA_SUCCESS) { - warn_on_cuda_error(name, status, - "skipped (context activation failed; resource leaked)"); + detail = "skipped (context activation failed; resource leaked)"; } else { - status = std::invoke(std::forward(operation), std::forward(args)...); - if (status != CUDA_SUCCESS) { - warn_on_cuda_error(name, status); - } + status = std::invoke(std::forward(operation)); } CUresult restore = exit_context(previous, changed, CUDA_SUCCESS); if (restore != CUDA_SUCCESS) { - warn_on_cuda_error(name, restore, "failed while restoring the caller's context"); + // Nothing is raised here, so the detail exit_context recorded has no + // exception to attach to; drop it. + clear_last_error_detail(); + } + std::invoke(std::forward(after_cleanup)); + if (status != CUDA_SUCCESS || restore != CUDA_SUCCESS) { + char operation_name[160]; + format_operation(operation_name, sizeof(operation_name), name, handle); + if (status != CUDA_SUCCESS) { + report_cuda_error(operation_name, status, detail); + } + if (restore != CUDA_SUCCESS) { + report_cuda_error(operation_name, restore, "failed while restoring the caller's context"); + } } return status; } +// Same, with nothing to do after the cleanup. +template +CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, + unsigned long long handle, Fn&& operation) noexcept { + return cleanup_in_context(h_context, name, handle, std::forward(operation), + []() noexcept {}); +} + #undef ASSERT_NOTHROW_INVOCABLE -// Decorate a CUDA operation to warn whenever it returns an error. +// Decorate a status-returning cleanup call to report whenever it fails. CUDA +// calls (CUresult) are reported with the error name and description; NVRTC, +// NVVM and nvJitLink calls (integer status codes) with the raw code. +// +// A pw_ wrapper is not a p_ pointer with logging. On failure it acquires the +// GIL and runs Python: the warning filters, showwarning, or sys.unraisablehook, +// any of which may be user code that calls back into cuda.core. Never invoke +// one while holding a C++ lock; the GIL must be the outermost lock. Where a +// lock must stay held, call the p_ pointer, keep the status, and report after +// the lock is released (see deviceptr_import_ipc and DESIGN.md). template class WarnOnFailure { public: explicit WarnOnFailure(const char* operation) noexcept : operation_(operation) {} - template - CUresult operator()(Args&&... args) const noexcept { - CUresult status = Function(std::forward(args)...); - if (status != CUDA_SUCCESS) { - warn_on_cuda_error(operation_, status); - } + // The first argument is the resource being released; it is named in the + // report so that independent failures are not collapsed by the warning + // registry (see format_operation). + template + auto operator()(First&& first, Rest&&... rest) const noexcept { + const unsigned long long handle = handle_bits(first); + auto status = Function(std::forward(first), std::forward(rest)...); + report(status, handle); return status; } private: + void report(CUresult status, unsigned long long handle) const noexcept { + if (status == CUDA_SUCCESS || status == CUDA_ERROR_DEINITIALIZED) { + return; + } + char operation[160]; + format_operation(operation, sizeof(operation), operation_, handle); + report_cuda_error(operation, status); + } + + template + void report(Status status, unsigned long long handle) const noexcept { + if (static_cast(status) != 0) { + char operation[160]; + format_operation(operation, sizeof(operation), operation_, handle); + report_status_code(operation, static_cast(status)); + } + } + const char* operation_; }; -// Warning-decorated CUDA operations used by non-throwing cleanup paths. +// Warning-decorated CUDA operations for deleters and cleanup paths. Each one +// may run user Python on failure (see WarnOnFailure above): no C++ lock held. const WarnOnFailure pw_cuStreamDestroy{"cuStreamDestroy"}; const WarnOnFailure pw_cuEventDestroy{"cuEventDestroy"}; const WarnOnFailure pw_cuMemFree{"cuMemFree"}; @@ -405,6 +672,18 @@ const WarnOnFailure pw_cuArrayDestroy{"cuArrayDestroy"}; const WarnOnFailure pw_cuMipmappedArrayDestroy{"cuMipmappedArrayDestroy"}; const WarnOnFailure pw_cuTexObjectDestroy{"cuTexObjectDestroy"}; const WarnOnFailure pw_cuSurfObjectDestroy{"cuSurfObjectDestroy"}; +const WarnOnFailure pw_cuGreenCtxDestroy{"cuGreenCtxDestroy"}; +const WarnOnFailure pw_cuMemPoolDestroy{"cuMemPoolDestroy"}; +const WarnOnFailure pw_cuMemFreeHost{"cuMemFreeHost"}; +const WarnOnFailure pw_cuGraphDestroy{"cuGraphDestroy"}; +const WarnOnFailure pw_cuGraphExecDestroy{"cuGraphExecDestroy"}; +const WarnOnFailure pw_cuGraphicsUnregisterResource{"cuGraphicsUnregisterResource"}; +const WarnOnFailure pw_cuLinkDestroy{"cuLinkDestroy"}; +const WarnOnFailure pw_cuUserObjectRelease{"cuUserObjectRelease"}; +const WarnOnFailure pw_cuGraphReleaseUserObject{"cuGraphReleaseUserObject"}; +const WarnOnFailure pw_nvrtcDestroyProgram{"nvrtcDestroyProgram"}; +const WarnOnFailure pw_nvvmDestroyProgram{"nvvmDestroyProgram"}; +const WarnOnFailure pw_nvJitLinkDestroy{"nvJitLinkDestroy"}; } // namespace @@ -426,6 +705,50 @@ CUresult context_get_stream_priority_range(const ContextHandle& h_context, }); } +// Query the device of the provided context. +CUresult context_get_device(const ContextHandle& h_context, CUdevice* device) noexcept { + return invoke_in_context(h_context, [&]() noexcept { + return p_cuCtxGetDevice(device); + }); +} + +// Set a graph node's parameters with h_context current (an empty handle runs in +// the caller's context). Returns the cuGraphNodeSetParams status. A failure to +// restore the caller's context is returned separately in *restore_status so the +// caller can publish the metadata that depends on the successful update before +// raising it; if the update itself failed, its status is returned with the +// restoration failure recorded as its detail and *restore_status is CUDA_SUCCESS. +CUresult graph_node_set_params(CUgraphNode node, CUgraphNodeParams* params, + const ContextHandle& h_context, + CUresult* restore_status) noexcept { + *restore_status = CUDA_SUCCESS; + if (!p_cuGraphNodeSetParams) { + return CUDA_ERROR_NOT_SUPPORTED; + } + CUcontext previous = nullptr; + int changed = 0; + CUresult status = enter_context(h_context, &previous, &changed); + if (status != CUDA_SUCCESS) { + return status; + } + { + GILReleaseGuard gil; + status = p_cuGraphNodeSetParams(node, params); + } + if (!changed) { + return status; + } + CUresult restored = restore_context(previous); + if (restored == CUDA_SUCCESS) { + return status; + } + note_context_not_restored(previous, status, restored); + if (status == CUDA_SUCCESS) { + *restore_status = restored; + } + return status; +} + // ============================================================================ // CUDA user-object deferred cleanup // @@ -767,7 +1090,7 @@ GreenCtxHandle create_green_ctx_handle(CUdevResource* resources, unsigned int nb new GreenCtxBox{green_ctx}, [](const GreenCtxBox* b) { GILReleaseGuard gil; - p_cuGreenCtxDestroy(b->resource); + pw_cuGreenCtxDestroy(b->resource); delete b; } ); @@ -1187,7 +1510,7 @@ static MemoryPoolHandle wrap_mempool_owned(CUmemoryPool pool) { [](const MemoryPoolBox* b) { GILReleaseGuard gil; clear_mempool_peer_access(b->resource); - p_cuMemPoolDestroy(b->resource); + pw_cuMemPoolDestroy(b->resource); delete b; } ); @@ -1290,7 +1613,7 @@ DevicePtrHandle deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle& h GILReleaseGuard gil; const DeallocationStream& stream = b->deallocation; cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", + deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); @@ -1320,7 +1643,7 @@ DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) GILReleaseGuard gil; const DeallocationStream& stream = b->deallocation; cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", + deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); @@ -1353,7 +1676,7 @@ DevicePtrHandle deviceptr_alloc_host(size_t size) { new DevicePtrBox{reinterpret_cast(ptr), DeallocationStream{}}, [](DevicePtrBox* b) { GILReleaseGuard gil; - p_cuMemFreeHost(reinterpret_cast(b->resource)); + pw_cuMemFreeHost(reinterpret_cast(b->resource)); delete b; } ); @@ -1405,7 +1728,7 @@ DevicePtrHandle deviceptr_create_mapped_graphics( CUgraphicsResource resource = as_cu(h_resource); const DeallocationStream& stream = b->deallocation; cleanup_in_context( - deallocation_context(stream), "cuGraphicsUnmapResources", + deallocation_context(stream), "cuGraphicsUnmapResources", handle_bits(resource), [&]() noexcept { return p_cuGraphicsUnmapResources( 1, &resource, as_cu(stream.h_stream)); @@ -1444,7 +1767,7 @@ DevicePtrHandle deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject* if (mr_dealloc_cb) { const DeallocationStream& stream = b->deallocation; cleanup_in_context( - deallocation_context(stream), "MemoryResource.deallocate", + deallocation_context(stream), "MemoryResource.deallocate", handle_bits(b->resource), [&]() noexcept { mr_dealloc_cb(mr, b->resource, size, stream.h_stream); return CUDA_SUCCESS; @@ -1524,48 +1847,71 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* ExportDataKey key; std::memcpy(&key.data, data, sizeof(key.data)); + // The mutex makes lookup, import and registration one step, so two + // threads cannot import the same descriptor twice. Release the GIL + // before taking it: a thread blocked on the mutex while holding the + // GIL would deadlock with a holder that needs the GIL back (#2840). + // Nothing under the mutex may acquire the GIL, so a failed discard is + // reported only after the lock is released (see DESIGN.md). GILReleaseGuard gil; - std::lock_guard lock(ipc_import_mutex); - - if (auto h = ipc_ptr_cache.lookup(key)) { - return h; - } + CUresult discard_status = CUDA_SUCCESS; + CUdeviceptr discarded = 0; + { + std::lock_guard lock(ipc_import_mutex); - CUdeviceptr ptr; - if (CUDA_SUCCESS != (err = p_cuMemPoolImportPointer(&ptr, *h_pool, data))) { - return {}; - } + if (auto h = ipc_ptr_cache.lookup(key)) { + return h; + } - DeallocationStream ds; - if (!make_deallocation_stream(h_stream, ds)) { - pw_cuMemFreeAsync(ptr, as_cu(h_stream)); - return {}; - } + CUdeviceptr ptr; + if (CUDA_SUCCESS != (err = p_cuMemPoolImportPointer(&ptr, *h_pool, data))) { + return {}; + } - auto box = std::shared_ptr( - new DevicePtrBox{ptr, std::move(ds)}, - [h_pool, key](DevicePtrBox* b) { - // Release the GIL first (the GIL is the outermost lock), then hold - // the mutex across unregister + free. A concurrent import that finds - // this entry expired must wait until the mapping is gone; otherwise - // it re-imports the same allocation and the first cuMemFreeAsync - // unmaps it for both (nvbug 5570902). - GILReleaseGuard gil; - std::lock_guard lock(ipc_import_mutex); - ipc_ptr_cache.unregister_handle(key); - const DeallocationStream& stream = b->deallocation; - cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", - [&]() noexcept { - return p_cuMemFreeAsync( - b->resource, as_cu(stream.h_stream)); - }); - delete b; + DeallocationStream ds; + if (make_deallocation_stream(h_stream, ds)) { + auto box = std::shared_ptr( + new DevicePtrBox{ptr, std::move(ds)}, + [h_pool, key](DevicePtrBox* b) { + // Release the GIL first (the GIL is the outermost lock), then hold the + // mutex across unregister + free: a concurrent import that finds this + // entry expired must wait until the mapping is gone, or it re-imports + // the same allocation and the first cuMemFreeAsync unmaps it for both + // (nvbug 5570902). Nothing under the mutex may acquire the GIL, so the + // deallocation context is resolved before the lock (it may report) and + // the lock is released as soon as the cleanup is done. + GILReleaseGuard gil; + const DeallocationStream& stream = b->deallocation; + ContextHandle h_dealloc = deallocation_context(stream); + std::unique_lock lock(ipc_import_mutex); + ipc_ptr_cache.unregister_handle(key); + cleanup_in_context( + h_dealloc, "cuMemFreeAsync", handle_bits(b->resource), + [&]() noexcept { + return p_cuMemFreeAsync(b->resource, as_cu(stream.h_stream)); + }, + [&]() noexcept { lock.unlock(); }); + delete b; + } + ); + DevicePtrHandle h(box, &box->resource); + ipc_ptr_cache.register_handle(key, h); + return h; } - ); - DevicePtrHandle h(box, &box->resource); - ipc_ptr_cache.register_handle(key, h); - return h; + + // No deallocation stream could be recorded: discard the import with + // the raw call (a pw_ report would acquire the GIL under the mutex). + discard_status = p_cuMemFreeAsync(ptr, as_cu(h_stream)); + discarded = ptr; + } + if (discard_status != CUDA_SUCCESS) { + char operation[160]; + format_operation(operation, sizeof(operation), "cuMemFreeAsync", discarded); + report_cuda_error(operation, discard_status, + "failed while discarding an IPC import that could not record a " + "deallocation stream; the mapping leaked"); + } + return {}; } else { GILReleaseGuard gil; @@ -1586,7 +1932,7 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* GILReleaseGuard gil; const DeallocationStream& stream = b->deallocation; cleanup_in_context( - deallocation_context(stream), "cuMemFreeAsync", + deallocation_context(stream), "cuMemFreeAsync", handle_bits(b->resource), [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); @@ -1915,7 +2261,7 @@ void rollback_prepared_attachment( GraphBox* box = get_box(state->h_graph); if (box->resource) { GILReleaseGuard gil; - p_cuGraphReleaseUserObject( + pw_cuGraphReleaseUserObject( box->resource, state->replacement->object, 1); } } @@ -1961,7 +2307,7 @@ GraphHandle create_graph_handle(CUgraph graph) { GraphBox* root = hierarchy->root(); if (root && root->resource) { GILReleaseGuard gil; - p_cuGraphDestroy(root->resource); + pw_cuGraphDestroy(root->resource); } retry_deferred_cleanup(); delete hierarchy; @@ -2193,7 +2539,7 @@ CUresult graph_prepare_attachment( if (status != CUDA_SUCCESS) { prepared->replacement_entry.mapped() = nullptr; prepared->replacement = nullptr; - p_cuUserObjectRelease(object, 1); + pw_cuUserObjectRelease(object, 1); return status; } } @@ -2324,7 +2670,7 @@ struct GraphExecBox { ~GraphExecBox() noexcept { if (resource) { GILReleaseGuard gil; - p_cuGraphExecDestroy(resource); + pw_cuGraphExecDestroy(resource); } // The accumulator fields may be dangling after exec destruction. retry_deferred_cleanup(); @@ -2344,7 +2690,7 @@ GraphExecHandle make_graph_exec_handle( ~RawGraphExecGuard() noexcept { if (resource) { GILReleaseGuard gil; - p_cuGraphExecDestroy(resource); + pw_cuGraphExecDestroy(resource); } retry_deferred_cleanup(); } @@ -2366,7 +2712,8 @@ struct ExecAttachmentStaging { ExecAttachments* accumulator = nullptr; ~ExecAttachmentStaging() noexcept { - release(); + report_cuda_error("cuGraphReleaseUserObject", release(), + "failed while dropping a staged graph attachment"); } CUresult release() noexcept { @@ -2412,7 +2759,7 @@ CUresult stage_exec_attachments( *h_source, object, 1, CU_GRAPH_USER_OBJECT_MOVE); if (status != CUDA_SUCCESS) { // Dropping the last reference retires the accumulator. - p_cuUserObjectRelease(object, 1); + pw_cuUserObjectRelease(object, 1); return status; } } @@ -2682,7 +3029,7 @@ GraphicsResourceHandle create_graphics_resource_handle(CUgraphicsResource resour new GraphicsResourceBox{resource}, [](const GraphicsResourceBox* b) { GILReleaseGuard gil; - p_cuGraphicsUnregisterResource(b->resource); + pw_cuGraphicsUnregisterResource(b->resource); delete b; } ); @@ -2705,8 +3052,10 @@ NvrtcProgramHandle create_nvrtc_program_handle(nvrtcProgram prog) { [](NvrtcProgramBox* b) { // Note: nvrtcDestroyProgram takes nvrtcProgram* and nulls it, // but we're deleting the box anyway so nulling is harmless. - // Errors are ignored (standard destructor practice). - p_nvrtcDestroyProgram(&b->resource); + if (p_nvrtcDestroyProgram) { + GILReleaseGuard gil; + pw_nvrtcDestroyProgram(&b->resource); + } delete b; } ); @@ -2736,7 +3085,8 @@ NvvmProgramHandle create_nvvm_program_handle(nvvmProgram prog) { // but we're deleting the box anyway so nulling is harmless. // If NVVM is not available, the function pointer is null. if (p_nvvmDestroyProgram) { - p_nvvmDestroyProgram(&b->resource.raw); + GILReleaseGuard gil; + pw_nvvmDestroyProgram(&b->resource.raw); } delete b; } @@ -2767,7 +3117,8 @@ NvJitLinkHandle create_nvjitlink_handle(nvJitLink_t handle) { // but we're deleting the box anyway so nulling is harmless. // If nvJitLink is not available, the function pointer is null. if (p_nvJitLinkDestroy) { - p_nvJitLinkDestroy(&b->resource.raw); + GILReleaseGuard gil; + pw_nvJitLinkDestroy(&b->resource.raw); } delete b; } @@ -2795,9 +3146,9 @@ CuLinkHandle create_culink_handle(CUlinkState state) { new CuLinkBox{state}, [](CuLinkBox* b) { // cuLinkDestroy takes CUlinkState by value (not pointer). - // Errors are ignored (standard destructor practice). if (p_cuLinkDestroy) { - p_cuLinkDestroy(b->resource); + GILReleaseGuard gil; + pw_cuLinkDestroy(b->resource); } delete b; } @@ -2820,7 +3171,12 @@ FileDescriptorHandle create_fd_handle(int fd) { #else return FileDescriptorHandle( new int(fd), - [](const int* p) { ::close(*p); delete p; } + [](const int* p) { + if (::close(*p) != 0) { + report_message("close() failed for an IPC file descriptor; the descriptor may have leaked"); + } + delete p; + } ); #endif } @@ -3000,7 +3356,7 @@ TexObjectHandle make_tex_object_handle(const CUDA_RESOURCE_DESC& res, new TexObjectBox{TexObjectValue{obj}, std::move(h_backing), h_context}, [](const TexObjectBox* b) { GILReleaseGuard gil; - cleanup_in_context(b->h_context, "cuTexObjectDestroy", [&]() noexcept { + cleanup_in_context(b->h_context, "cuTexObjectDestroy", handle_bits(b->resource.raw), [&]() noexcept { return p_cuTexObjectDestroy(b->resource.raw); }); delete b; @@ -3048,7 +3404,7 @@ SurfObjectHandle create_surf_object_handle(const ContextHandle& h_context, new SurfObjectBox{SurfObjectValue{obj}, h_backing, h_context}, [](const SurfObjectBox* b) { GILReleaseGuard gil; - cleanup_in_context(b->h_context, "cuSurfObjectDestroy", [&]() noexcept { + cleanup_in_context(b->h_context, "cuSurfObjectDestroy", handle_bits(b->resource.raw), [&]() noexcept { return p_cuSurfObjectDestroy(b->resource.raw); }); delete b; diff --git a/cuda_core/cuda/core/_cpp/resource_handles.hpp b/cuda_core/cuda/core/_cpp/resource_handles.hpp index 419710ea0d9..b6136673889 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.hpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.hpp @@ -57,6 +57,52 @@ CUresult peek_last_error() noexcept; // Explicitly clear the last error void clear_last_error() noexcept; +// ============================================================================ +// Non-propagating error reporting +// +// Paths that cannot raise (shared_ptr deleters, __dealloc__) report failures +// through these functions instead of discarding them. They emit a +// cuda.core.CUDAWarning when the interpreter can be used and write to stderr +// otherwise; they never raise. Emitting the warning acquires the GIL and runs +// user Python (warning filters, showwarning, sys.unraisablehook), so never call +// them while holding a C++ lock. See docs/source/error_handling.rst and the +// "Which channel to use" table in DESIGN.md. +// ============================================================================ + +// Register the Python warning category used by report_* (cuda.core.CUDAWarning). +void register_warning_category(PyObject* category) noexcept; + +// Report a failed CUDA call. `detail` replaces the default "failed" wording, +// e.g. "skipped (context activation failed; resource leaked)". +// CUDA_ERROR_DEINITIALIZED (driver shutting down) is never reported. +void report_cuda_error(const char* operation, CUresult status, const char* detail = nullptr) noexcept; + +// Report a message that is not tied to a CUresult. +void report_message(const char* message) noexcept; + +// Report a failed NVRTC/NVVM/nvJitLink call by raw status code. +void report_status_code(const char* operation, long code) noexcept; + +// Attach a failed CUDA call to the Python exception currently being handled +// (PEP 678 note, Python 3.11+): for rollback failures inside `except` blocks +// whose original exception is about to be re-raised. When no exception is +// being handled or notes are unavailable, falls back to report_cuda_error(). +void attach_rollback_failure(const char* operation, CUresult status, const char* detail = nullptr) noexcept; + +// Detail recorded by a context-scoped helper for the CUresult it is about to +// return, e.g. that the caller's context could not be restored. The Cython +// error path attaches it to the raised CUDAError as a note. Thread-local and +// keyed by status: take_ returns the detail (valid until the next take on this +// thread) and clears it when `status` is the CUresult it was recorded for, and +// returns nullptr otherwise, so a detail whose status was never raised cannot +// attach to an unrelated error. +const char* take_last_error_detail(CUresult status) noexcept; +void clear_last_error_detail() noexcept; + +// Tests only: make the next context restoration on this thread fail with +// `status`, leaving the target context current as a real failure would. +void set_context_restore_fault_for_testing(CUresult status) noexcept; + // ============================================================================ // CUDA driver function pointers // @@ -73,6 +119,8 @@ extern decltype(&cuCtxGetCurrent) p_cuCtxGetCurrent; extern decltype(&cuCtxSetCurrent) p_cuCtxSetCurrent; extern decltype(&cuCtxSynchronize) p_cuCtxSynchronize; extern decltype(&cuCtxGetStreamPriorityRange) p_cuCtxGetStreamPriorityRange; +extern decltype(&cuCtxGetDevice) p_cuCtxGetDevice; +extern decltype(&cuGraphNodeSetParams) p_cuGraphNodeSetParams; extern decltype(&cuGreenCtxCreate) p_cuGreenCtxCreate; extern decltype(&cuGreenCtxDestroy) p_cuGreenCtxDestroy; extern decltype(&cuCtxFromGreenCtx) p_cuCtxFromGreenCtx; @@ -263,6 +311,21 @@ CUresult context_get_stream_priority_range( int* least_priority, int* greatest_priority) noexcept; +// Query the device of the provided context. +// Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. +CUresult context_get_device(const ContextHandle& h_context, CUdevice* device) noexcept; + +// Call cuGraphNodeSetParams with h_context current (empty handle: the caller's +// context). Returns the update status; *restore_status receives a failure to +// restore the caller's context after a successful update, which the caller +// raises only after publishing the metadata that depends on the update. +// Returns CUDA_ERROR_NOT_SUPPORTED when the driver lacks cuGraphNodeSetParams. +CUresult graph_node_set_params( + CUgraphNode node, + CUgraphNodeParams* params, + const ContextHandle& h_context, + CUresult* restore_status) noexcept; + // ============================================================================ // Stream handle functions // ============================================================================ diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index a7e7d59e04a..170bedc0034 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -1315,8 +1315,12 @@ class Device: HANDLE_RETURN(cydriver.cuCtxGetCurrent(&prev_ctx)) if prev_ctx != NULL: HANDLE_RETURN(cydriver.cuCtxGetDevice(&prev_dev)) - HANDLE_RETURN(cydriver.cuCtxPopCurrent(&prev_ctx)) - HANDLE_RETURN(cydriver.cuCtxPushCurrent(curr_ctx)) + # cuCtxSetCurrent replaces the top of the thread's context stack + # in one driver call (or binds ctx when nothing is current), so + # a failure leaves the previous context current instead of + # leaving the thread with no context, as a failed pop-then-push + # would. + HANDLE_RETURN(cydriver.cuCtxSetCurrent(curr_ctx)) self._has_inited = True self._context = ctx # Store owning context reference if prev_ctx != NULL: diff --git a/cuda_core/cuda/core/_memory/_buffer.pyx b/cuda_core/cuda/core/_memory/_buffer.pyx index 52539b8c314..740f40b7de5 100644 --- a/cuda_core/cuda/core/_memory/_buffer.pyx +++ b/cuda_core/cuda/core/_memory/_buffer.pyx @@ -36,11 +36,12 @@ IF CUDA_CORE_BUILD_MAJOR >= 13: from cuda.core._stream cimport Stream, Stream_accept, Stream_is_legacy_default_token, default_stream from cuda.core._utils.cuda_utils cimport HANDLE_RETURN, _parse_fill_value -import sys +import warnings from collections.abc import Sequence from typing import TYPE_CHECKING from cuda.core._memory._copy_enums import CopyOptions, _reject_unsupported_during_api_call +from cuda.core._utils.cuda_utils import CUDAWarning from cuda.core._utils.pycompat import BufferProtocol from cuda.core._dlpack import classify_dl_device, make_py_capsule from cuda.core._device import Device @@ -59,7 +60,11 @@ cdef void _mr_dealloc_callback( size_t size, const StreamHandle& h_stream, ) noexcept: - """Called by the C++ deleter to deallocate via MemoryResource.deallocate.""" + """Called by the C++ deleter to deallocate via MemoryResource.deallocate. + + Runs from a destructor, so nothing can be raised here; failures are reported + as :class:`~cuda.core.CUDAWarning` (see the error handling policy). + """ cdef Stream stream try: if not h_stream: @@ -71,8 +76,12 @@ cdef void _mr_dealloc_callback( stream = Stream._from_handle(Stream, h_stream) mr.deallocate(int(ptr), size, stream=stream) except Exception as exc: - print(f"Warning: mr.deallocate() failed during Buffer destruction: {exc}", - file=sys.stderr) + warnings.warn( + f"mr.deallocate({int(ptr):#x}) failed during Buffer destruction; " + f"the allocation may have leaked: {exc}", + CUDAWarning, + stacklevel=2, + ) register_mr_dealloc_callback(_mr_dealloc_callback) diff --git a/cuda_core/cuda/core/_resource_handles.pxd b/cuda_core/cuda/core/_resource_handles.pxd index fe075b6414b..5c69ac3a2a8 100644 --- a/cuda_core/cuda/core/_resource_handles.pxd +++ b/cuda_core/cuda/core/_resource_handles.pxd @@ -2,6 +2,7 @@ # # SPDX-License-Identifier: Apache-2.0 +from cpython.object cimport PyObject from libc.stddef cimport size_t from libc.stdint cimport intptr_t @@ -168,6 +169,18 @@ cdef cydriver.CUresult get_last_error() noexcept nogil cdef cydriver.CUresult peek_last_error() noexcept nogil cdef void clear_last_error() noexcept nogil +# Non-propagating error reporting (never raises; emits cuda.core.CUDAWarning +# when possible, else writes to stderr) +cdef void register_warning_category(PyObject* category) noexcept +cdef void report_cuda_error( + const char* operation, cydriver.CUresult status, const char* detail) noexcept nogil +cdef void report_message(const char* message) noexcept nogil +cdef void report_status_code(const char* operation, long code) noexcept nogil +cdef void attach_rollback_failure( + const char* operation, cydriver.CUresult status, const char* detail) noexcept nogil +cdef const char* take_last_error_detail(cydriver.CUresult status) noexcept nogil +cdef void clear_last_error_detail() noexcept nogil + # Context handles cdef ContextHandle create_context_handle_ref(cydriver.CUcontext ctx) except+ nogil cdef ContextHandle create_context_handle_from_green_ctx(const GreenCtxHandle& h_green_ctx) except+ nogil @@ -184,6 +197,11 @@ cdef cydriver.CUresult context_get_stream_priority_range( const ContextHandle& h_context, int* least_priority, int* greatest_priority) noexcept nogil +cdef cydriver.CUresult context_get_device( + const ContextHandle& h_context, cydriver.CUdevice* device) noexcept nogil +cdef cydriver.CUresult graph_node_set_params( + cydriver.CUgraphNode node, cydriver.CUgraphNodeParams* params, + const ContextHandle& h_context, cydriver.CUresult* restore_status) noexcept nogil # Stream handles cdef StreamHandle create_stream_handle( diff --git a/cuda_core/cuda/core/_resource_handles.pyi b/cuda_core/cuda/core/_resource_handles.pyi index 66cbf80761a..3e3c1977fe5 100644 --- a/cuda_core/cuda/core/_resource_handles.pyi +++ b/cuda_core/cuda/core/_resource_handles.pyi @@ -41,3 +41,18 @@ PreparedAttachmentDeleter: TypeAlias = Incomplete PreparedChildGraphUpdateState: TypeAlias = Incomplete PreparedExecAttachmentState: TypeAlias = Incomplete PreparedExecAttachmentDeleter: TypeAlias = Incomplete + +def _set_context_restore_fault_for_testing(status: int): + """Make the next context restoration on this thread fail with ``status``. + + Test hook for the context save/restore paths in the handle layer. The + injected failure leaves the target context current, exactly as a failing + ``cuCtxSetCurrent`` would, so callers must restore the context themselves. + """ +def _attach_rollback_failure_for_testing(status: int): + """Attach a failed CUDA call to the exception being handled, or report it. + + Test hook for ``attach_rollback_failure()``. Called inside an ``except`` + block it adds a note to the exception being handled (Python 3.11+); anywhere + else it emits a ``CUDAWarning``. + """ diff --git a/cuda_core/cuda/core/_resource_handles.pyx b/cuda_core/cuda/core/_resource_handles.pyx index 0f8d6e15cde..4f003fdbdcb 100644 --- a/cuda_core/cuda/core/_resource_handles.pyx +++ b/cuda_core/cuda/core/_resource_handles.pyx @@ -10,6 +10,7 @@ # The cdef extern from declarations below satisfy the .pxd declarations directly, # without needing separate wrapper functions. +from cpython.object cimport PyObject from cpython.pycapsule cimport PyCapsule_GetName, PyCapsule_GetPointer from libc.stddef cimport size_t @@ -36,6 +37,26 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": cydriver.CUresult peek_last_error "cuda_core::peek_last_error" () noexcept nogil void clear_last_error "cuda_core::clear_last_error" () noexcept nogil + # Non-propagating error reporting + void register_warning_category "cuda_core::register_warning_category" ( + PyObject* category) noexcept + void report_cuda_error "cuda_core::report_cuda_error" ( + const char* operation, cydriver.CUresult status, const char* detail) noexcept nogil + void report_message "cuda_core::report_message" (const char* message) noexcept nogil + void report_status_code "cuda_core::report_status_code" ( + const char* operation, long code) noexcept nogil + void attach_rollback_failure "cuda_core::attach_rollback_failure" ( + const char* operation, cydriver.CUresult status, const char* detail) noexcept nogil + # Alias for calls made from this module: calling the pxd-declared name here + # would make Cython emit a conflicting static prototype for it. + void _attach_rollback_failure_local "cuda_core::attach_rollback_failure" ( + const char* operation, cydriver.CUresult status, const char* detail) noexcept nogil + const char* take_last_error_detail "cuda_core::take_last_error_detail" ( + cydriver.CUresult status) noexcept nogil + void clear_last_error_detail "cuda_core::clear_last_error_detail" () noexcept nogil + void set_context_restore_fault_for_testing "cuda_core::set_context_restore_fault_for_testing" ( + cydriver.CUresult status) noexcept nogil + # Context handles ContextHandle create_context_handle_ref "cuda_core::create_context_handle_ref" ( cydriver.CUcontext ctx) except+ nogil @@ -57,6 +78,11 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": const ContextHandle& h_context, int* least_priority, int* greatest_priority) noexcept nogil + cydriver.CUresult context_get_device "cuda_core::context_get_device" ( + const ContextHandle& h_context, cydriver.CUdevice* device) noexcept nogil + cydriver.CUresult graph_node_set_params "cuda_core::graph_node_set_params" ( + cydriver.CUgraphNode node, cydriver.CUgraphNodeParams* params, + const ContextHandle& h_context, cydriver.CUresult* restore_status) noexcept nogil # Stream handles StreamHandle create_stream_handle "cuda_core::create_stream_handle" ( @@ -323,6 +349,8 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": void* p_cuCtxSetCurrent "reinterpret_cast(cuda_core::p_cuCtxSetCurrent)" void* p_cuCtxSynchronize "reinterpret_cast(cuda_core::p_cuCtxSynchronize)" void* p_cuCtxGetStreamPriorityRange "reinterpret_cast(cuda_core::p_cuCtxGetStreamPriorityRange)" + void* p_cuCtxGetDevice "reinterpret_cast(cuda_core::p_cuCtxGetDevice)" + void* p_cuGraphNodeSetParams "reinterpret_cast(cuda_core::p_cuGraphNodeSetParams)" void* p_cuGreenCtxCreate "reinterpret_cast(cuda_core::p_cuGreenCtxCreate)" void* p_cuGreenCtxDestroy "reinterpret_cast(cuda_core::p_cuGreenCtxDestroy)" void* p_cuCtxFromGreenCtx "reinterpret_cast(cuda_core::p_cuCtxFromGreenCtx)" @@ -433,6 +461,7 @@ cdef void _init_driver_fn_pointers() noexcept: global p_cuGetErrorName, p_cuGetErrorString global p_cuDevicePrimaryCtxRetain, p_cuDevicePrimaryCtxRelease, p_cuCtxGetCurrent global p_cuCtxSetCurrent, p_cuCtxSynchronize, p_cuCtxGetStreamPriorityRange + global p_cuCtxGetDevice, p_cuGraphNodeSetParams global p_cuGreenCtxCreate, p_cuGreenCtxDestroy, p_cuCtxFromGreenCtx global p_cuDevResourceGenerateDesc, p_cuGreenCtxStreamCreate global p_cuStreamCreateWithPriority, p_cuStreamDestroy, p_cuStreamGetCtx @@ -469,6 +498,9 @@ cdef void _init_driver_fn_pointers() noexcept: p_cuCtxSetCurrent = _get_driver_fn("cuCtxSetCurrent") p_cuCtxSynchronize = _get_driver_fn("cuCtxSynchronize") p_cuCtxGetStreamPriorityRange = _get_driver_fn("cuCtxGetStreamPriorityRange") + p_cuCtxGetDevice = _get_driver_fn("cuCtxGetDevice") + # Graph node parameter updates need CUDA 12.2+ (checked again at the call site). + p_cuGraphNodeSetParams = _get_optional_driver_fn("cuGraphNodeSetParams") p_cuGreenCtxCreate = _get_optional_driver_fn("cuGreenCtxCreate") p_cuGreenCtxDestroy = _get_optional_driver_fn("cuGreenCtxDestroy") p_cuCtxFromGreenCtx = _get_optional_driver_fn("cuCtxFromGreenCtx") @@ -554,6 +586,27 @@ cdef void _init_driver_fn_pointers() noexcept: _init_driver_fn_pointers() initialize_deferred_cleanup() + +def _set_context_restore_fault_for_testing(int status): + """Make the next context restoration on this thread fail with ``status``. + + Test hook for the context save/restore paths in the handle layer. The + injected failure leaves the target context current, exactly as a failing + ``cuCtxSetCurrent`` would, so callers must restore the context themselves. + """ + set_context_restore_fault_for_testing(status) + + +def _attach_rollback_failure_for_testing(int status): + """Attach a failed CUDA call to the exception being handled, or report it. + + Test hook for ``attach_rollback_failure()``. Called inside an ``except`` + block it adds a note to the exception being handled (Python 3.11+); anywhere + else it emits a ``CUDAWarning``. + """ + _attach_rollback_failure_local( + b"cuTestOperation", status, b"failed while testing") + # ============================================================================= # NVRTC function pointer initialization # ============================================================================= diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index e662d67c87f..916a6eb01fe 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -27,6 +27,7 @@ from cuda.core._context cimport ( from cuda.core._device_resources cimport DeviceResources from cuda.core._event import Event, EventOptions +from cuda.core._resource_handles cimport context_get_device from cuda.core._resource_handles cimport ( ContextHandle, EventHandle, @@ -36,7 +37,6 @@ from cuda.core._resource_handles cimport ( create_stream_handle, create_stream_handle_with_owner, context_get_stream_priority_range, - get_current_context, get_last_error, get_legacy_stream, get_per_thread_stream, @@ -564,10 +564,7 @@ cdef inline int Stream_get_ctx(Stream self, ContextHandle* h_context) except?-1 cdef inline int Stream_get_ctx_device(Stream self, ContextHandle* h_context, int* device_id) except?-1: """Resolve the stream's context handle and device ID.""" - cdef cydriver.CUcontext ctx cdef cydriver.CUdevice target_dev - cdef ContextHandle current_context - cdef bint switch_context cdef bint is_default = Stream_is_default_token(self) with nogil: @@ -575,14 +572,9 @@ cdef inline int Stream_get_ctx_device(Stream self, ContextHandle* h_context, int if self._device_id >= 0 and not is_default: device_id[0] = self._device_id else: - # Get device ID from context, switching context temporarily if needed - current_context = get_current_context() - switch_context = (as_cu(current_context) != as_cu(h_context[0])) - if switch_context: - HANDLE_RETURN(cydriver.cuCtxPushCurrent(as_cu(h_context[0]))) - HANDLE_RETURN(cydriver.cuCtxGetDevice(&target_dev)) - if switch_context: - HANDLE_RETURN(cydriver.cuCtxPopCurrent(&ctx)) + # Query the device with the stream's context current. The handle + # layer restores the caller's context, including on failure. + HANDLE_RETURN(context_get_device(h_context[0], &target_dev)) device_id[0] = target_dev if not is_default: self._device_id = device_id[0] diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyi b/cuda_core/cuda/core/_utils/cuda_utils.pyi index 51f992fa238..a883328e605 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyi +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyi @@ -12,9 +12,40 @@ _keep_nvrtc_in_stub: nvrtc.nvrtcResult _keep_runtime_in_stub: runtime.cudaError_t _fork_warning_checked = False -class CUDAError(Exception): ... +class CUDAError(Exception): + """Raised when a CUDA driver or runtime call fails. + + The message names the CUDA error and, when one is known, explains it. A + secondary failure observed while the error was being raised, such as a + failed restoration of the caller's CUDA context, is attached as a note + (``__notes__``) on Python 3.11 and newer and appended to the message on + older interpreters. See the error handling page of the ``cuda.core`` + documentation for the guarantees an exception provides. + """ + +class CUDAWarning(RuntimeWarning): + """Warning issued when ``cuda.core`` hits a CUDA error it cannot raise. + + ``cuda.core`` raises exceptions for failures in ordinary calls. Some failures + happen where no exception can propagate: while a resource is released by the + garbage collector or by a CUDA callback, including the driver calls that + switch and restore the CUDA context around such a release. Those failures + are reported as this warning instead, and the affected resource may have + leaked. + + Filter on this category to make such failures fatal in tests:: + + warnings.filterwarnings("error", category=cuda.core.CUDAWarning) + + Because the report comes from a destructor, an escalated warning cannot be + raised into user code; it is delivered through :func:`sys.unraisablehook` + (which pytest surfaces as ``PytestUnraisableExceptionWarning``). + + .. versionadded:: 1.3.0 + """ -class NVRTCError(CUDAError): ... +class NVRTCError(CUDAError): + """Raised when an NVRTC call fails; the compiler log is appended when available.""" class ComputeCapability(NamedTuple): """A named tuple of (major, minor) CUDA compute capability version numbers.""" diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyx b/cuda_core/cuda/core/_utils/cuda_utils.pyx index ce75746de56..7e96a06d026 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyx +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyx @@ -27,17 +27,55 @@ from cuda.bindings.nvjitlink import nvJitLinkError from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink +from cuda.core._resource_handles cimport ( + register_warning_category, + take_last_error_detail, +) from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS class CUDAError(Exception): - pass + """Raised when a CUDA driver or runtime call fails. + + The message names the CUDA error and, when one is known, explains it. A + secondary failure observed while the error was being raised, such as a + failed restoration of the caller's CUDA context, is attached as a note + (``__notes__``) on Python 3.11 and newer and appended to the message on + older interpreters. See the error handling page of the ``cuda.core`` + documentation for the guarantees an exception provides. + """ + + +class CUDAWarning(RuntimeWarning): + """Warning issued when ``cuda.core`` hits a CUDA error it cannot raise. + + ``cuda.core`` raises exceptions for failures in ordinary calls. Some failures + happen where no exception can propagate: while a resource is released by the + garbage collector or by a CUDA callback, including the driver calls that + switch and restore the CUDA context around such a release. Those failures + are reported as this warning instead, and the affected resource may have + leaked. + + Filter on this category to make such failures fatal in tests:: + + warnings.filterwarnings("error", category=cuda.core.CUDAWarning) + + Because the report comes from a destructor, an escalated warning cannot be + raised into user code; it is delivered through :func:`sys.unraisablehook` + (which pytest surfaces as ``PytestUnraisableExceptionWarning``). + + .. versionadded:: 1.3.0 + """ + + +# Route the C++ handle layer's non-propagating reports through this category. +register_warning_category(CUDAWarning) class NVRTCError(CUDAError): - pass + """Raised when an NVRTC call fails; the compiler log is appended when available.""" @@ -131,23 +169,40 @@ cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS +cdef inline void _attach_detail(exc, str detail): + # PEP 678 notes (Python 3.11+) keep the detail separable from the message; + # older interpreters get it appended to the message instead. + add_note = getattr(exc, "add_note", None) + if add_note is not None: + add_note(detail) + else: + exc.args = (f"{exc.args[0]} ({detail})", *exc.args[1:]) + + cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil: if error == cydriver.CUresult.CUDA_SUCCESS: return 0 cdef const char* name + cdef const char* desc + # A context-scoped helper in the handle layer may have recorded why this + # status needs more explanation (e.g. the caller's context was not restored). + cdef const char* detail = take_last_error_detail(error) name_err = cydriver.cuGetErrorName(error, &name) if name_err != cydriver.CUresult.CUDA_SUCCESS: raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") + desc_err = cydriver.cuGetErrorString(error, &desc) with gil: # TODO: consider lower this to Cython expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) if expl is not None: - raise CUDAError(f"{name.decode()}: {expl}") - cdef const char* desc - desc_err = cydriver.cuGetErrorString(error, &desc) - if desc_err != cydriver.CUresult.CUDA_SUCCESS: - raise CUDAError(f"{name.decode()}") - raise CUDAError(f"{name.decode()}: {desc.decode()}") + exc = CUDAError(f"{name.decode()}: {expl}") + elif desc_err != cydriver.CUresult.CUDA_SUCCESS: + exc = CUDAError(name.decode()) + else: + exc = CUDAError(f"{name.decode()}: {desc.decode()}") + if detail != NULL: + _attach_detail(exc, detail.decode()) + raise exc cpdef inline int _check_runtime_error(error) except?-1: diff --git a/cuda_core/cuda/core/graph/_graph_builder.pyx b/cuda_core/cuda/core/graph/_graph_builder.pyx index 071fff38386..40d74bbbeba 100644 --- a/cuda_core/cuda/core/graph/_graph_builder.pyx +++ b/cuda_core/cuda/core/graph/_graph_builder.pyx @@ -20,6 +20,7 @@ from cuda.core.graph._subclasses cimport ( ExecutableGraphNode, create_executable_node_view, ) +from cuda.core._resource_handles cimport attach_rollback_failure, report_cuda_error from cuda.core._resource_handles cimport ( GraphExecHandle, GraphHandle, @@ -860,6 +861,12 @@ cdef class GraphBuilder: if rollback_status == cydriver.CUDA_SUCCESS: invalidate_child_graph_state( self._h_graph, c_new_node) + else: + # The original exception propagates with the failed rollback + # attached as a note (error handling policy). + attach_rollback_failure( + b"cuGraphDestroyNode", rollback_status, + b"failed while rolling back a child graph node; the node remains in the graph") raise deps_info_update = [[new_node]] + [None] * (len(deps_info_out) - 1) @@ -990,8 +997,8 @@ cdef inline int GB_end_capture_if_needed(GraphBuilder gb, bint check_status) exc capture. A FORKED builder must not call cuStreamEndCapture: the driver requires forked streams to be joined first. - check_status=True checks the driver return (close()); False ignores it - (__dealloc__). + check_status=True raises on a driver error (close()); False reports it as + a CUDAWarning instead, because nothing can be raised from __dealloc__. """ cdef cydriver.CUgraph c_graph cdef cydriver.CUresult err @@ -1002,6 +1009,10 @@ cdef inline int GB_end_capture_if_needed(GraphBuilder gb, bint check_status) exc err = cydriver.cuStreamEndCapture(c_stream, &c_graph) if check_status: HANDLE_RETURN(err) + else: + report_cuda_error( + b"cuStreamEndCapture", err, + b"failed while releasing a GraphBuilder that was still building") return 0 diff --git a/cuda_core/cuda/core/graph/_graph_node.pyx b/cuda_core/cuda/core/graph/_graph_node.pyx index 7295d786089..7ce1588fd8b 100644 --- a/cuda_core/cuda/core/graph/_graph_node.pyx +++ b/cuda_core/cuda/core/graph/_graph_node.pyx @@ -45,6 +45,7 @@ from cuda.core.graph._subclasses cimport ( SwitchNode, WhileNode, ) +from cuda.core._resource_handles cimport attach_rollback_failure from cuda.core._resource_handles cimport ( GraphHandle, GraphNodeHandle, @@ -1114,6 +1115,12 @@ cdef inline ChildGraphNode GN_embed(GraphNode self, GraphDefinition child_def): rollback_status = cydriver.cuGraphDestroyNode(new_node) if rollback_status == cydriver.CUDA_SUCCESS: invalidate_child_graph_state(h_graph, new_node) + else: + # The original exception propagates with the failed rollback + # attached as a note (error handling policy). + attach_rollback_failure( + b"cuGraphDestroyNode", rollback_status, + b"failed while rolling back a child graph node; the node remains in the graph") raise return _registered(ChildGraphNode._create_with_params( diff --git a/cuda_core/cuda/core/graph/_subclasses.pyx b/cuda_core/cuda/core/graph/_subclasses.pyx index 2201f7babf0..4967983d947 100644 --- a/cuda_core/cuda/core/graph/_subclasses.pyx +++ b/cuda_core/cuda/core/graph/_subclasses.pyx @@ -29,6 +29,11 @@ from cuda.core.graph._graph_node cimport ( _init_memcpy_params, _resolve_memcpy_operand, ) +from cuda.core._resource_handles cimport ( + ContextHandle, + create_context_handle_ref, + graph_node_set_params, +) from cuda.core._resource_handles cimport ( EventHandle, GraphExecHandle, @@ -124,26 +129,25 @@ cdef void _set_definition_node_params( if node == NULL: raise RuntimeError("GraphNode has been destroyed") _require_graph_node_update_support() - cdef cydriver.CUcontext previous_ctx = NULL - cdef bint restore_ctx = False cdef PreparedAttachment prepared + cdef ContextHandle h_update_ctx + cdef cydriver.CUresult status + cdef cydriver.CUresult restore_status = cydriver.CUresult.CUDA_SUCCESS HANDLE_RETURN(graph_prepare_attachment( h_graph, owner0, owner1, &prepared)) if update_ctx != NULL: - with nogil: - HANDLE_RETURN(cydriver.cuCtxGetCurrent(&previous_ctx)) - if previous_ctx != update_ctx: - HANDLE_RETURN(cydriver.cuCtxSetCurrent(update_ctx)) - restore_ctx = True + h_update_ctx = create_context_handle_ref(update_ctx) + with nogil: + status = graph_node_set_params(node, params, h_update_ctx, &restore_status) + HANDLE_RETURN(status) + # The driver node now references the new owners. Publish their attachment + # before raising anything else: an exception here would roll back the + # prepared retention and leave the node pointing at released resources. try: - with nogil: - HANDLE_RETURN(cydriver.cuGraphNodeSetParams(node, params)) + HANDLE_RETURN(graph_commit_attachment(prepared, node)) finally: - if restore_ctx: - with nogil: - HANDLE_RETURN(cydriver.cuCtxSetCurrent(previous_ctx)) - HANDLE_RETURN(graph_commit_attachment(prepared, node)) + HANDLE_RETURN(restore_status) cdef void _set_executable_node_params( diff --git a/cuda_core/docs/source/api.rst b/cuda_core/docs/source/api.rst index 5ee34d34f54..76a228625e8 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -312,6 +312,23 @@ The associated enumerations — alongside the other ``cuda.core`` enumerations. +Errors and warnings +------------------- + +Failed CUDA calls raise exceptions; see :doc:`error_handling` for the +guarantees an exception provides and for the situations in which a failure is +reported as a warning instead. + +.. currentmodule:: cuda.core + +.. autosummary:: + :toctree: generated/ + + CUDAError + NVRTCError + CUDAWarning + + CUDA process checkpointing -------------------------- diff --git a/cuda_core/docs/source/error_handling.rst b/cuda_core/docs/source/error_handling.rst new file mode 100644 index 00000000000..519b0cacc47 --- /dev/null +++ b/cuda_core/docs/source/error_handling.rst @@ -0,0 +1,147 @@ +.. SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +.. SPDX-License-Identifier: Apache-2.0 + +.. currentmodule:: cuda.core + +Error Handling +============== + +``cuda.core`` reports failures with Python exceptions. This page describes what +an exception from ``cuda.core`` guarantees about the state it leaves behind, +what happens when a failure occurs where no exception can be raised, and the +few situations in which ``cuda.core`` cannot fully undo a failed operation. + +Exceptions +---------- + +A CUDA driver, runtime, NVRTC, NVVM or nvJitLink call that fails raises an +exception whose message contains the error name and its description: +:class:`CUDAError` for driver and runtime failures, its subclass +:class:`NVRTCError` for NVRTC failures, and the ``cuda.bindings`` error types +for NVVM and nvJitLink failures. Both ``cuda.core`` classes are importable from +the top-level ``cuda.core`` namespace. Invalid arguments and misuse raise the +usual Python exception types (``TypeError``, ``ValueError``, ``RuntimeError``). + +When a ``cuda.core`` call raises, the following hold: + +- A call that creates a resource creates nothing. If a later step of the call + fails after the resource was created, the resource is destroyed before the + exception propagates. +- The calling thread's current CUDA context is the one that was current when + the call began. The only method that changes the current context on purpose + is :meth:`Device.set_current`; every other method that must run in a + different context restores the caller's context before returning, whether it + succeeds or fails. See `Context restoration failures`_ for the one case in + which the driver refuses to restore it. +- Objects that were modified by a call that failed midway remain usable and + consistent, but some operations do not have an all-or-nothing outcome. Their + documentation says so where it applies (for example the graph mutation + methods that add several driver edges). + +``cuda.core`` does not swallow driver errors. When a second failure occurs +while an exception is being raised, for example the caller's context cannot be +restored after a failed call, or the rollback of a partially built graph node +fails, the second failure is attached to the exception as a note +(:meth:`BaseException.add_note`), which appears in the traceback and in +``__notes__``. Python 3.10 has no exception notes; there the information is +appended to the message when ``cuda.core`` constructs the exception, and +reported as described in the next section otherwise. + +Failures that cannot be raised +------------------------------ + +Some ``cuda.core`` code runs where no Python exception can propagate: + +- resources released by the garbage collector or by the deferred cleanup of + CUDA graphs, and the CUDA driver calls those releases make, including the + context switch and restoration around such a release; +- callbacks invoked by CUDA. + +A CUDA error in one of these places is reported as a :class:`CUDAWarning`. The +message names the failed driver call, the handle of the resource involved, and +the CUDA error. Python shows a warning with a given text once per call site by +default; because the handle differs per resource, independent failures are not +collapsed into one report. The warning means the +affected resource may have leaked; ``cuda.core`` never leaves a resource in use +by CUDA with its memory released (it prefers a leak to a dangling pointer). + +:class:`CUDAWarning` derives from :class:`RuntimeWarning`, so it is shown by +default and can be filtered like any other warning. To make these failures +loud in a test suite:: + + import warnings + import cuda.core + + warnings.filterwarnings("error", category=cuda.core.CUDAWarning) + +Because the report comes from a destructor or callback, an escalated warning is +delivered through :func:`sys.unraisablehook` rather than raised into user code. +pytest reports it as ``PytestUnraisableExceptionWarning``, which its +``-W error`` option turns into a test failure. + +``CUDA_ERROR_DEINITIALIZED`` is not reported. It means the CUDA driver is +shutting down, which happens during process exit; cleanup failures at that +point are expected and there is nothing left to clean up. + +Context restoration failures +---------------------------- + +Methods that run in a context other than the current one, such as +:meth:`Device.create_stream` when another device is current, switch the current +context, perform the driver call, and switch back. Restoring the caller's +context can fail only when the driver is shutting down +(``CUDA_ERROR_DEINITIALIZED``), when the caller's context was destroyed in the +meantime (``CUDA_ERROR_INVALID_CONTEXT``), or when the driver is reporting an +earlier, unrecoverable kernel fault (see `Sticky errors`_). None of these can +be fixed by retrying, so ``cuda.core`` does not retry. + +When restoration fails in an ordinary call, the resource created by the call is +destroyed and a ``CUDAError`` is raised for the failed ``cuCtxSetCurrent``, +with a note stating that the caller's context could not be restored and which +context is now current. If the call itself failed as well, its own error is +raised and the restoration failure is the note. Call +:meth:`Device.set_current` before issuing further CUDA work on that thread. + +When restoration fails inside a destructor or callback, a :class:`CUDAWarning` +is issued and the thread keeps the context that the cleanup used. + +Sticky errors +------------- + +Some CUDA errors mark the process as unusable for further CUDA work, for +example ``CUDA_ERROR_ILLEGAL_ADDRESS`` or ``CUDA_ERROR_LAUNCH_FAILED`` after a +kernel fault. The CUDA documentation calls for the process to be terminated and +relaunched after such an error, and every later CUDA call returns the same +error. Because these faults are detected asynchronously, the call that first +raises the error is often unrelated to the kernel that caused it. + +``cuda.core`` raises these errors like any other and does not attempt to +recover from them. It does not terminate the process for you: the exception +carries the Python traceback of the call that observed the fault, and your +application decides how to shut down. + +Interpreter shutdown +-------------------- + +Once the interpreter starts finalizing, ``cuda.core`` no longer touches Python +objects from CUDA callbacks or destructors. Resources whose release would +require Python at that point are intentionally leaked; the operating system and +the driver reclaim them when the process exits. Release all ``cuda.core`` +objects explicitly (with ``close()`` or a ``with`` block) if their deterministic +release matters. + +Process termination +------------------- + +``cuda.core`` never terminates the process. A CUDA error is raised, or reported +as a :class:`CUDAWarning` where nothing can be raised; this includes errors in +destructors and callbacks and failures to restore the caller's context. An +internal error in ``cuda.core`` itself is handled the same way: it raises a +``RuntimeError`` that asks you to report it, or is reported as a warning, and +the affected resource is leaked rather than released in an inconsistent state. +If you want such failures to stop your program, escalate the warning category:: + + warnings.filterwarnings("error", category=cuda.core.CUDAWarning) + +A process exit caused by ``cuda.core``, for example through an exception that +escapes a destructor, is a bug; please report it. diff --git a/cuda_core/docs/source/index.rst b/cuda_core/docs/source/index.rst index 34c0933ffb8..373e6b61665 100644 --- a/cuda_core/docs/source/index.rst +++ b/cuda_core/docs/source/index.rst @@ -16,6 +16,7 @@ Welcome to the documentation for ``cuda.core``. examples interoperability concurrency + error_handling api api_nvml environment_variables diff --git a/cuda_core/docs/source/release/1.3.0-notes.rst b/cuda_core/docs/source/release/1.3.0-notes.rst index a1432459e58..6280f127285 100644 --- a/cuda_core/docs/source/release/1.3.0-notes.rst +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -18,6 +18,18 @@ New features completion of the asynchronous prefetch operation. (`#2109 `__) +- Added :class:`CUDAWarning`, the warning category ``cuda.core`` uses for CUDA + errors that cannot be raised, such as a failed driver call while a resource + is released by the garbage collector. Filter on it with + ``warnings.filterwarnings("error", category=cuda.core.CUDAWarning)`` to make + such failures loud. The new :doc:`error handling <../error_handling>` page + documents what an exception from ``cuda.core`` guarantees, how failures that + cannot be raised are reported, and how context restoration failures and + sticky CUDA errors are handled. + +- :class:`CUDAError` and :class:`NVRTCError` are now importable from the + top-level ``cuda.core`` namespace, alongside :class:`CUDAWarning`. Previously + they were only available from a private module. Fixes and enhancements ---------------------- @@ -48,3 +60,47 @@ Fixes and enhancements ``RuntimeError``. :attr:`Buffer.device_id` on such a buffer returns ``-1`` as well, which also lets a pinned buffer back a linear or pitched texture resource. + +- Cleanup failures are now reported as :class:`CUDAWarning` instead of being + written to ``stderr`` with ``print`` or ``fprintf``, so they can be filtered, + captured with :func:`warnings.catch_warnings`, and escalated. Failures of + ``cuStreamDestroy``, ``cuEventDestroy``, ``cuMemFree``, ``cuMemFreeAsync``, + ``cuMemFreeHost``, ``cuMemPoolDestroy``, ``cuGreenCtxDestroy``, + ``cuGraphDestroy``, ``cuGraphExecDestroy``, ``cuGraphicsUnregisterResource``, + ``cuLinkDestroy``, ``cuArrayDestroy``, ``cuMipmappedArrayDestroy``, + ``cuTexObjectDestroy``, ``cuSurfObjectDestroy``, user-object releases and the + NVRTC, NVVM and nvJitLink destroy calls made from destructors were previously + discarded; they are now reported, and each message names the failed call, + the handle of the affected resource and the CUDA error, so Python's + once-per-call-site warning filter does not collapse independent failures. + ``CUDA_ERROR_DEINITIALIZED`` (the driver is shutting down) is not reported. Test code that matched the old ``stderr`` + text uses ``pytest.warns(CUDAWarning)`` instead. + +- When a :class:`Device` method has to run in the device's context and the + caller's context cannot be restored afterwards, the created resource is + destroyed and the raised ``CUDAError`` now carries a note (Python 3.11+; + appended to the message on 3.10) stating that the caller's context could not + be restored and which context is current. Previously the error named only + the driver status of the failed ``cuCtxSetCurrent`` call. If the call itself + failed as well, its error is raised and the restoration failure is the note. + A restoration failure during resource cleanup is reported as + :class:`CUDAWarning`. + +- Updating a memcpy or memset graph node whose context differs from the current + one no longer risks a dangling node parameter when the caller's context + cannot be restored after the update: the resources referenced by the new + parameters are now retained before the restoration failure is raised. + +- :meth:`Device.set_current` with an explicit :class:`Context` now switches + contexts with a single driver call, so a failure leaves the previous context + current instead of leaving the thread with no context. It also works when no + context is current, returning ``None``. + +- A failed rollback of a partially embedded child graph node is now attached + as a note to the exception that triggered the rollback (reported as + :class:`CUDAWarning` on Python 3.10), and a failed ``cuStreamEndCapture`` + made when a still-building :class:`~graph.GraphBuilder` is garbage collected + is now reported as :class:`CUDAWarning`; both were silent. + +- :attr:`Stream.device` and related queries on a stream whose context is not + current now restore the caller's context even when the device query fails. diff --git a/cuda_core/tests/helpers/contexts.py b/cuda_core/tests/helpers/contexts.py index 7ee01bb255f..f613f0ac336 100644 --- a/cuda_core/tests/helpers/contexts.py +++ b/cuda_core/tests/helpers/contexts.py @@ -1,18 +1,36 @@ # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 +import warnings from contextlib import contextmanager +from cuda.core import CUDAWarning from cuda.core._utils.cuda_utils import driver, handle_return __all__ = [ "assert_device_operations_use_bound_context", + "assert_no_cuda_warning", "current_context_handle", "no_current_context", "use_context", ] +@contextmanager +def assert_no_cuda_warning(): + """Fail if a :class:`CUDAWarning` is issued inside the block. + + Cleanup paths cannot raise, so a driver failure there surfaces only as a + warning; this makes such a failure a test failure. Tests using it must be + marked ``thread_unsafe``: warning capture is process-global. + """ + with warnings.catch_warnings(record=True) as records: + warnings.simplefilter("always", CUDAWarning) + yield + cuda_warnings = [str(record.message) for record in records if issubclass(record.category, CUDAWarning)] + assert not cuda_warnings, f"unexpected CUDAWarning(s): {cuda_warnings}" + + def current_context_handle(): """Return the current CUDA context handle, or zero if none is current.""" return int(handle_return(driver.cuCtxGetCurrent())) diff --git a/cuda_core/tests/test_error_handling.py b/cuda_core/tests/test_error_handling.py new file mode 100644 index 00000000000..dc5d7d56b7b --- /dev/null +++ b/cuda_core/tests/test_error_handling.py @@ -0,0 +1,324 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Tests for the error handling policy (docs/source/error_handling.rst). + +Ordinary calls raise and leave the caller's context untouched; failures that +cannot be raised are reported as CUDAWarning; a failure to restore the caller's +context is raised (or reported) with an explanation rather than swallowed or +turned into a process abort; and a secondary failure that occurs while an +exception is being raised is attached to that exception as a note. Restoration +failures are injected with the handle layer's test hook, which leaves the target +context current exactly as a real ``cuCtxSetCurrent`` failure would, so every +test here restores the context stack itself. +""" + +import ctypes +import sys +import warnings +from contextlib import contextmanager + +import pytest +from helpers.constants import POOL_SIZE +from helpers.contexts import assert_no_cuda_warning, current_context_handle + +import cuda.core +from cuda.core import ( + CUDAWarning, + DeviceMemoryResource, + DeviceMemoryResourceOptions, + LegacyPinnedMemoryResource, +) +from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource +from cuda.core._resource_handles import ( + _attach_rollback_failure_for_testing, + _set_context_restore_fault_for_testing, +) +from cuda.core._stream import default_stream +from cuda.core._utils.cuda_utils import CUDAError, driver, handle_return +from cuda.core._utils.version import binding_version, driver_version +from cuda.core.graph import GraphDefinition + +INVALID_CONTEXT = int(driver.CUresult.CUDA_ERROR_INVALID_CONTEXT) +INVALID_VALUE = int(driver.CUresult.CUDA_ERROR_INVALID_VALUE) +DEINITIALIZED = int(driver.CUresult.CUDA_ERROR_DEINITIALIZED) + +# PEP 678 exception notes; on 3.10 the same information lands in the message. +HAS_NOTES = sys.version_info >= (3, 11) + +thread_unsafe_context_fault = pytest.mark.thread_unsafe( + reason="injects a thread-local restoration fault and mutates the CUDA context stack" +) +thread_unsafe_warning_capture = pytest.mark.thread_unsafe(reason="warning capture is process-global") + + +def error_text(exc): + """The message plus any notes, wherever the detail lives on this interpreter.""" + return "\n".join([str(exc), *getattr(exc, "__notes__", [])]) + + +@contextmanager +def no_context_with_restore_fault(status=INVALID_CONTEXT): + """Pop the current context and make the next restoration fail with ``status``. + + On exit, drop whatever the failed restoration left current, clear an unused + fault, and push the popped context back. + """ + previous = handle_return(driver.cuCtxPopCurrent()) + assert current_context_handle() == 0 + _set_context_restore_fault_for_testing(status) + try: + yield + finally: + _set_context_restore_fault_for_testing(0) + handle_return(driver.cuCtxSetCurrent(driver.CUcontext(0))) + handle_return(driver.cuCtxPushCurrent(previous)) + + +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_cudawarning_is_public_and_shown_by_default(): + assert "CUDAWarning" in cuda.core.__all__ + assert issubclass(CUDAWarning, RuntimeWarning) + + +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_cuda_errors_are_public(): + assert "CUDAError" in cuda.core.__all__ + assert "NVRTCError" in cuda.core.__all__ + assert cuda.core.CUDAError is CUDAError + assert issubclass(cuda.core.NVRTCError, CUDAError) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_create_stream_raises_when_context_cannot_be_restored(init_cuda): + """Creation is undone and the error explains the context state; no abort, no warning.""" + dev = init_cuda + with no_context_with_restore_fault(): + with assert_no_cuda_warning(), pytest.raises(CUDAError) as excinfo: + dev.create_stream() + text = error_text(excinfo.value) + assert "could not be restored" in text + assert "CUDA_ERROR_INVALID_CONTEXT" in text + assert "Device.set_current()" in text + if HAS_NOTES: + # The explanation is a note, separable from the driver error message. + assert "could not be restored" not in str(excinfo.value) + assert any("could not be restored" in note for note in excinfo.value.__notes__) + # As documented, a failed restoration leaves the device's context current. + assert current_context_handle() == int(dev.context.handle) + assert current_context_handle() == int(dev.context.handle) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_sync_raises_when_context_cannot_be_restored(init_cuda): + """A context-scoped call without a created resource raises the same explanation.""" + dev = init_cuda + with no_context_with_restore_fault(): + with pytest.raises(CUDAError) as excinfo: + dev.sync() + assert "could not be restored" in error_text(excinfo.value) + assert current_context_handle() == int(dev.context.handle) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_failed_call_raises_its_own_error_with_the_restore_failure_attached(init_cuda): + """When the call and the restoration both fail, the call's error is raised and the + restoration failure is attached to it; nothing is reported out of band.""" + dev = init_cuda + mr = _SynchronousMemoryResource(dev.device_id) + with no_context_with_restore_fault(): + with assert_no_cuda_warning(), pytest.raises(CUDAError) as excinfo: + mr.allocate(1 << 62) + message = str(excinfo.value) + text = error_text(excinfo.value) + # The allocation failure is the primary error, not the restoration failure. + assert not message.startswith("CUDA_ERROR_INVALID_CONTEXT") + assert "could not be restored after this failure" in text + assert "cuCtxSetCurrent: CUDA_ERROR_INVALID_CONTEXT" in text + assert "Device.set_current()" in text + assert current_context_handle() == int(dev.context.handle) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_unused_restore_fault_does_not_fire_without_a_context_switch(init_cuda): + """The hook only affects restorations; a call in the current context never restores.""" + dev = init_cuda + _set_context_restore_fault_for_testing(INVALID_CONTEXT) + try: + stream = dev.create_stream() + stream.close() + finally: + _set_context_restore_fault_for_testing(0) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_cleanup_reports_restore_failure_as_warning(mempool_device): + """A restoration failure inside a destructor cannot raise, so it is reported.""" + dev = mempool_device + mr = DeviceMemoryResource(dev, DeviceMemoryResourceOptions(max_size=POOL_SIZE)) + # A default-stream deallocation records the allocating context, so freeing + # with no context current switches to it and must switch back. + buf = mr.allocate(256, stream=default_stream()) + ptr = int(buf.handle) + with no_context_with_restore_fault(): + with pytest.warns(CUDAWarning, match="restoring the caller's context") as records: + buf.close() + messages = [str(record.message) for record in records] + assert any("CUDA_ERROR_INVALID_CONTEXT" in message for message in messages) + # The report names the resource so that independent failures stay distinct. + assert any(f"{ptr:#x}" in message for message in messages), messages + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_independent_cleanup_failures_are_reported_separately(mempool_device): + """Python's default filter shows a given warning text once per call site. Two + resources failing the same call from the same line must still produce two + reports, which the handle in the message guarantees.""" + dev = mempool_device + mr = DeviceMemoryResource(dev, DeviceMemoryResourceOptions(max_size=POOL_SIZE)) + buffers = [mr.allocate(256, stream=default_stream()) for _ in range(2)] + handles = [int(buf.handle) for buf in buffers] + previous = handle_return(driver.cuCtxPopCurrent()) + try: + with warnings.catch_warnings(record=True) as records: + warnings.simplefilter("default", CUDAWarning) + for buf in buffers: + _set_context_restore_fault_for_testing(INVALID_CONTEXT) + buf.close() + # The failed restoration left the device's context current; clear it + # so the next release also has to switch and fail the same way. + handle_return(driver.cuCtxSetCurrent(driver.CUcontext(0))) + finally: + _set_context_restore_fault_for_testing(0) + handle_return(driver.cuCtxSetCurrent(driver.CUcontext(0))) + handle_return(driver.cuCtxPushCurrent(previous)) + messages = [str(record.message) for record in records if issubclass(record.category, CUDAWarning)] + assert len(messages) == 2, messages + assert all(f"{handle:#x}" in message for handle, message in zip(handles, messages)), messages + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_escalated_cudawarning_from_cleanup_is_not_a_crash(mempool_device): + """With CUDAWarning promoted to an error, a destructor-path report cannot be raised. + + It is delivered through sys.unraisablehook instead; the process continues and + the resource release still runs. pytest surfaces the hook as a warning, so the + hook is replaced here to keep the test's own outcome deterministic. + """ + import sys + import warnings + + dev = mempool_device + mr = DeviceMemoryResource(dev, DeviceMemoryResourceOptions(max_size=POOL_SIZE)) + buf = mr.allocate(256, stream=default_stream()) + unraisable = [] + previous_hook = sys.unraisablehook + sys.unraisablehook = unraisable.append + try: + with no_context_with_restore_fault(), warnings.catch_warnings(): + warnings.simplefilter("error", CUDAWarning) + buf.close() + finally: + sys.unraisablehook = previous_hook + assert len(unraisable) == 1 + assert issubclass(unraisable[0].exc_type, CUDAWarning) + assert "restoring the caller's context" in str(unraisable[0].exc_value) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_set_current_with_context_works_without_a_current_context(init_cuda): + """set_current(ctx) binds in one driver call; no previous context means None.""" + dev = init_cuda + ctx = dev.context + previous = handle_return(driver.cuCtxPopCurrent()) + try: + assert current_context_handle() == 0 + assert dev.set_current(ctx) is None + assert current_context_handle() == int(ctx.handle) + finally: + # Leave exactly one context on the stack, as the fixture expects. + handle_return(driver.cuCtxSetCurrent(driver.CUcontext(0))) + handle_return(driver.cuCtxPushCurrent(previous)) + + +@thread_unsafe_context_fault +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_memset_update_keeps_new_owners_alive_when_context_cannot_be_restored(device_x2): + """The node's new parameters stay valid: the attachment is published before the + restoration failure is raised, so the updated graph instantiates and runs.""" + if driver_version() < (13, 2, 0) or binding_version() < (13, 2, 0): + pytest.skip("node contexts are only recorded by cuGraphNodeGetParams on CUDA 13.2+") + node_dev, other_dev = device_x2 + node_dev.set_current() + memory_resource = LegacyPinnedMemoryResource() + dst = memory_resource.allocate(4) + replacement = memory_resource.allocate(4) + graph_def = GraphDefinition() + node = graph_def.memset(dst, 0x11, 4) + + # Updating from another device's context switches to the node's context and + # must switch back; make that restoration fail. + other_dev.set_current() + _set_context_restore_fault_for_testing(INVALID_CONTEXT) + try: + with pytest.raises(CUDAError) as excinfo: + node.update(dst=replacement, value=0x22) + assert "could not be restored" in error_text(excinfo.value) + finally: + _set_context_restore_fault_for_testing(0) + node_dev.set_current() + + def as_bytes(buffer): + return (ctypes.c_uint8 * 4).from_address(int(buffer.handle)) + + as_bytes(dst)[:] = [0] * 4 + as_bytes(replacement)[:] = [0] * 4 + graph = graph_def.instantiate() + stream = node_dev.create_stream() + graph.launch(stream) + stream.sync() + # The driver applied the update, and the replacement buffer it references + # is still retained by the graph rather than dangling. + assert list(as_bytes(replacement)) == [0x22] * 4 + assert list(as_bytes(dst)) == [0] * 4 + graph.close() + stream.close() + + +@thread_unsafe_warning_capture +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_rollback_failure_is_attached_to_the_propagating_exception(): + """A failed rollback inside an except block becomes a note on the exception being + handled (Python 3.11+); on 3.10, or with no exception being handled, it is reported + as a CUDAWarning. CUDA_ERROR_DEINITIALIZED is neither attached nor reported.""" + with pytest.raises(RuntimeError) as excinfo: + try: + raise RuntimeError("primary failure") + except RuntimeError: + if HAS_NOTES: + with assert_no_cuda_warning(): + _attach_rollback_failure_for_testing(INVALID_VALUE) + else: + with pytest.warns(CUDAWarning, match="cuTestOperation failed while testing"): + _attach_rollback_failure_for_testing(INVALID_VALUE) + with assert_no_cuda_warning(): + _attach_rollback_failure_for_testing(DEINITIALIZED) + raise + exc = excinfo.value + assert str(exc) == "primary failure" + if HAS_NOTES: + assert len(exc.__notes__) == 1 + assert "cuTestOperation failed while testing: CUDA_ERROR_INVALID_VALUE" in exc.__notes__[0] + else: + assert not hasattr(exc, "__notes__") + # With no exception being handled there is nothing to attach to. + with pytest.warns(CUDAWarning, match="cuTestOperation failed while testing"): + _attach_rollback_failure_for_testing(INVALID_VALUE) diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index befc3031e90..b61a2d5b2f3 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -27,7 +27,7 @@ ) from helpers.child_processes import child_timeout_sec, kill_subprocesses from helpers.constants import POOL_SIZE -from helpers.contexts import current_context_handle, no_current_context +from helpers.contexts import assert_no_cuda_warning, current_context_handle, no_current_context from helpers.memory import ( create_managed_memory_resource_or_skip, create_pinned_memory_resource_or_xfail, @@ -37,6 +37,7 @@ from cuda.core import ( Buffer, + CUDAWarning, Device, DeviceMemoryResource, DeviceMemoryResourceOptions, @@ -824,9 +825,10 @@ def test_from_handle_mr_explicit_stream_without_current_context(buffer_type): ] +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="claude-fable-5-1") @pytest.mark.parametrize("mr_cls", _HOST_ONLY_MRS) -def test_from_handle_host_only_mr_without_current_context(mr_cls, capfd): +def test_from_handle_host_only_mr_without_current_context(mr_cls): """Host-only memory needs no current context to create or free a Buffer.""" device = Device() device.set_current() @@ -836,28 +838,30 @@ def test_from_handle_host_only_mr_without_current_context(mr_cls, capfd): assert int(previous) != 0 try: assert int(handle_return(driver.cuCtxGetCurrent())) == 0 - buf = mr.allocate(64) - assert buf.is_host_accessible - buf.close() + with assert_no_cuda_warning(): + buf = mr.allocate(64) + assert buf.is_host_accessible + buf.close() assert int(handle_return(driver.cuCtxGetCurrent())) == 0 finally: handle_return(driver.cuCtxSetCurrent(previous)) - assert "Warning" not in capfd.readouterr().err - def _host_only_child_main(mr_cls): """Allocate and free host-only memory in a process that never initialized CUDA.""" - buf = mr_cls().allocate(64) - assert buf.is_host_accessible - buf.close() + # Warnings do not cross processes: check for a CUDAWarning here, where a + # failed assertion becomes a non-zero exit code for the parent to see. + with assert_no_cuda_warning(): + buf = mr_cls().allocate(64) + assert buf.is_host_accessible + buf.close() err, _ = driver.cuCtxGetCurrent() assert err == driver.CUresult.CUDA_ERROR_NOT_INITIALIZED, err @pytest.mark.agent_authored(model="claude-fable-5-1") @pytest.mark.parametrize("mr_cls", _HOST_ONLY_MRS) -def test_from_handle_host_only_mr_without_cuda_init(mr_cls, capfd): +def test_from_handle_host_only_mr_without_cuda_init(mr_cls): """Host-only buffers work in a spawned process that never initializes CUDA.""" process = mp.Process(target=_host_only_child_main, args=(mr_cls,)) process.start() @@ -865,26 +869,28 @@ def test_from_handle_host_only_mr_without_cuda_init(mr_cls, capfd): survivors = kill_subprocesses(process) assert not survivors, "child did not exit within timeout" assert process.exitcode == 0, f"child exited with {process.exitcode}" - assert "Warning" not in capfd.readouterr().err +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="gpt-5.6") -def test_mr_deallocation_failure_warns(capfd): - """Destructor-path MR failures are contained and reported.""" +def test_mr_deallocation_failure_warns(): + """Destructor-path MR failures are contained and reported as CUDAWarning.""" device = Device() device.set_current() FailingMR, _ = make_instrumented_memory_resource(deallocate_error=RuntimeError("expected deallocation failure")) buf = Buffer.from_handle(1, 1024, mr=FailingMR(device)) - buf.close() - assert ( - "Warning: mr.deallocate() failed during Buffer destruction: expected deallocation failure" - ) in capfd.readouterr().err + with pytest.warns( + CUDAWarning, + match=r"mr\.deallocate\(0x[0-9a-f]+\) failed during Buffer destruction.*expected deallocation failure", + ): + buf.close() +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="cursor-grok-4.5") @pytest.mark.parametrize("replace_stream", [False, True]) -def test_mr_deallocation_without_current_context(init_cuda, capsys, replace_stream): +def test_mr_deallocation_without_current_context(init_cuda, replace_stream): """MR-backed Buffer teardown activates the recorded context when none is current.""" TrackingMR, telemetry = make_instrumented_memory_resource(DummyDeviceMemoryResource, track_active=True) mr = TrackingMR(init_cuda) @@ -895,16 +901,17 @@ def test_mr_deallocation_without_current_context(init_cuda, capsys, replace_stre with no_current_context(): assert current_context_handle() == 0 - buf.close(stream) + with assert_no_cuda_warning(): + buf.close(stream) assert len(telemetry["active"]) == 0 assert current_context_handle() == 0 - assert "mr.deallocate() failed" not in capsys.readouterr().err +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="cursor-grok-4.5") @pytest.mark.parametrize("replace_stream", [False, True]) -def test_mr_deallocation_with_foreign_context(device_x2, capsys, replace_stream): +def test_mr_deallocation_with_foreign_context(device_x2, replace_stream): """MR-backed Buffer teardown switches away from an unrelated current context.""" alloc_dev, foreign_dev = device_x2 alloc_dev.set_current() @@ -921,11 +928,11 @@ def test_mr_deallocation_with_foreign_context(device_x2, capsys, replace_stream) assert foreign_ctx != alloc_ctx try: - buf.close(stream) + with assert_no_cuda_warning(): + buf.close(stream) assert len(telemetry["active"]) == 0 assert current_context_handle() == foreign_ctx - assert "mr.deallocate() failed" not in capsys.readouterr().err finally: alloc_dev.set_current() @@ -945,8 +952,9 @@ def test_mr_deallocate_raises_on_driver_error(mempool_device): mr.deallocate(0xDEADBEEF, 256, stream=stream) +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="cursor-grok-4.5") -def test_pool_buffer_deallocates_without_current_context(mempool_device, capfd): +def test_pool_buffer_deallocates_without_current_context(mempool_device): """Pool Buffer.close frees on the recorded stream with no current context.""" dev = mempool_device stream = dev.create_stream() @@ -959,18 +967,17 @@ def test_pool_buffer_deallocates_without_current_context(mempool_device, capfd): with no_current_context(): assert current_context_handle() == 0 - buf.close() + with assert_no_cuda_warning(): + buf.close() stream.sync() assert mr.attributes.used_mem_current < used_after_alloc assert current_context_handle() == 0 - err = capfd.readouterr().err - assert "cuMemFreeAsync failed" not in err - assert "mr.deallocate() failed" not in err +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="cursor-grok-4.5") -def test_pool_buffer_deallocates_with_foreign_context(mempool_device_x2, capfd): +def test_pool_buffer_deallocates_with_foreign_context(mempool_device_x2): """Pool Buffer.close frees under the recorded context while another is current.""" alloc_dev, foreign_dev = mempool_device_x2 alloc_dev.set_current() @@ -988,7 +995,8 @@ def test_pool_buffer_deallocates_with_foreign_context(mempool_device_x2, capfd): assert foreign_ctx != alloc_ctx try: - buf.close() + with assert_no_cuda_warning(): + buf.close() assert current_context_handle() == foreign_ctx # Observe the free on the allocation device, then restore the foreign context. @@ -996,9 +1004,6 @@ def test_pool_buffer_deallocates_with_foreign_context(mempool_device_x2, capfd): stream.sync() assert mr.attributes.used_mem_current < used_after_alloc foreign_dev.set_current() - - err = capfd.readouterr().err - assert "cuMemFreeAsync failed" not in err finally: alloc_dev.set_current() @@ -2354,8 +2359,9 @@ def test_synchronous_memory_resource_restores_context_after_failure(device_x2): assert current_context_handle() == current_context +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="claude-sonnet-5") -def test_synchronous_memory_resource_default_stream_deallocates_in_own_context(device_x2, capsys): +def test_synchronous_memory_resource_default_stream_deallocates_in_own_context(device_x2): """Buffer teardown with no explicit stream frees in the resource's own context, not whatever context happens to be current at close() time.""" from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource @@ -2370,13 +2376,14 @@ def test_synchronous_memory_resource_default_stream_deallocates_in_own_context(d buf = mr.allocate(64) # no explicit stream: records a context-bound default token assert current_context_handle() == current_context - buf.close() # no explicit stream: reuses the recorded token + with assert_no_cuda_warning(): + buf.close() # no explicit stream: reuses the recorded token assert current_context_handle() == current_context - assert capsys.readouterr().err == "" +@pytest.mark.thread_unsafe(reason="records process-global warnings and mutates the context stack") @pytest.mark.agent_authored(model="claude-sonnet-5") -def test_synchronous_memory_resource_allocate_without_current_context(device_x2, capsys): +def test_synchronous_memory_resource_allocate_without_current_context(device_x2): """allocate()/close() with no explicit stream succeed with no context current, instead of raising or leaking the allocation (#2311).""" from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource @@ -2386,14 +2393,12 @@ def test_synchronous_memory_resource_allocate_without_current_context(device_x2, mr = _SynchronousMemoryResource(alloc_dev.device_id, alloc_dev.context) current_dev.set_current() - with no_current_context(): + with no_current_context(), assert_no_cuda_warning(): buf = mr.allocate(64) assert current_context_handle() == 0 buf.close() assert current_context_handle() == 0 - assert capsys.readouterr().err == "" - @pytest.mark.parametrize( ("method", "spec", "match"),