From 83a3d466908a6fa0a01a74671285b899e5d27567 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Tue, 1 Sep 2026 13:38:43 -0700 Subject: [PATCH 1/9] cuda.core: make Device methods use their bound context Run context-sensitive Device operations against the Device's bound context while preserving caller state. Centralize context-aware cleanup and synchronous allocation handling so resource lifetimes remain correct. --- cuda_core/cuda/core/_cpp/resource_handles.cpp | 591 ++++++++++++------ cuda_core/cuda/core/_cpp/resource_handles.hpp | 48 +- cuda_core/cuda/core/_device.pyi | 44 +- cuda_core/cuda/core/_device.pyx | 81 ++- .../core/_memory/_device_memory_resource.pyi | 17 + .../core/_memory/_device_memory_resource.pyx | 72 +++ cuda_core/cuda/core/_memory/_legacy.py | 44 -- cuda_core/cuda/core/_resource_handles.pxd | 33 +- cuda_core/cuda/core/_resource_handles.pyx | 48 +- cuda_core/cuda/core/_stream.pyx | 27 +- cuda_core/cuda/core/texture/_array.pyi | 5 +- cuda_core/cuda/core/texture/_array.pyx | 9 +- .../cuda/core/texture/_mipmapped_array.pyi | 6 +- .../cuda/core/texture/_mipmapped_array.pyx | 15 +- cuda_core/cuda/core/texture/_surface.pyi | 7 +- cuda_core/cuda/core/texture/_surface.pyx | 25 +- cuda_core/cuda/core/texture/_texture.pyi | 5 +- cuda_core/cuda/core/texture/_texture.pyx | 40 +- cuda_core/docs/source/interoperability.rst | 4 + cuda_core/docs/source/release/1.2.0-notes.rst | 7 + cuda_core/tests/conftest.py | 9 + cuda_core/tests/helpers/contexts.py | 90 +++ .../tests/memory_ipc/test_peer_access.py | 4 + cuda_core/tests/test_device.py | 82 ++- cuda_core/tests/test_green_context.py | 52 +- cuda_core/tests/test_launcher.py | 2 +- cuda_core/tests/test_memory.py | 129 ++-- cuda_core/tests/test_stream.py | 2 +- cuda_core/tests/test_texture_surface.py | 105 ++++ 29 files changed, 1158 insertions(+), 445 deletions(-) create mode 100644 cuda_core/tests/helpers/contexts.py diff --git a/cuda_core/cuda/core/_cpp/resource_handles.cpp b/cuda_core/cuda/core/_cpp/resource_handles.cpp index ee116a9f353..ee4aef8e2e9 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.cpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.cpp @@ -12,12 +12,15 @@ #include #include #include +#include #include #include #include #include #include +#include #include +#include #include #ifndef _WIN32 @@ -33,10 +36,15 @@ namespace cuda_core { // function pointers extracted from cuda.bindings.cydriver.__pyx_capi__. // ============================================================================ +decltype(&cuGetErrorName) p_cuGetErrorName = nullptr; +decltype(&cuGetErrorString) p_cuGetErrorString = nullptr; + decltype(&cuDevicePrimaryCtxRetain) p_cuDevicePrimaryCtxRetain = nullptr; decltype(&cuDevicePrimaryCtxRelease) p_cuDevicePrimaryCtxRelease = nullptr; decltype(&cuCtxGetCurrent) p_cuCtxGetCurrent = nullptr; decltype(&cuCtxSetCurrent) p_cuCtxSetCurrent = nullptr; +decltype(&cuCtxSynchronize) p_cuCtxSynchronize = nullptr; +decltype(&cuCtxGetStreamPriorityRange) p_cuCtxGetStreamPriorityRange = nullptr; decltype(&cuGreenCtxCreate) p_cuGreenCtxCreate = nullptr; decltype(&cuGreenCtxDestroy) p_cuGreenCtxDestroy = nullptr; decltype(&cuCtxFromGreenCtx) p_cuCtxFromGreenCtx = nullptr; @@ -128,46 +136,34 @@ NvvmDestroyProgramFn p_nvvmDestroyProgram = nullptr; NvJitLinkDestroyFn p_nvJitLinkDestroy = nullptr; // ============================================================================ -// GIL management helpers +// GIL and scoped-context management helpers // ============================================================================ namespace { -// Helper to release the GIL while calling into the CUDA driver. -// This guard is *conditional*: if the caller already dropped the GIL, -// we avoid calling PyEval_SaveThread (which requires holding the GIL). -// It also handles the case where Python is finalizing and GIL operations -// are no longer safe. +// Conditionally release the GIL while calling into the CUDA driver. class GILReleaseGuard { public: - GILReleaseGuard() : tstate_(nullptr), released_(false) { - // Don't try to manipulate GIL if Python is finalizing + GILReleaseGuard() noexcept { if (!Py_IsInitialized() || py_is_finalizing()) { return; } - // PyGILState_Check() returns 1 if the GIL is held by this thread. if (PyGILState_Check()) { tstate_ = PyEval_SaveThread(); - released_ = true; } - // Note: If the GIL is not released (finalizing, or not held): - // - Reduces parallelism (other Python threads remain blocked) - // - No deadlock risk as long as the guarded code doesn't call back into Python } ~GILReleaseGuard() { - if (released_) { + if (tstate_) { PyEval_RestoreThread(tstate_); } } - // Non-copyable, non-movable GILReleaseGuard(const GILReleaseGuard&) = delete; GILReleaseGuard& operator=(const GILReleaseGuard&) = delete; private: - PyThreadState* tstate_; - bool released_; + PyThreadState* tstate_ = nullptr; }; // Helper to acquire the GIL when we might not hold it. @@ -200,55 +196,228 @@ class GILAcquireGuard { bool acquired_; }; -// Temporarily make a context current, restoring the caller's prior binding -// (including having no context current) on scope exit. The handle is held for -// the duration so the context cannot be destroyed mid-scope. -class ScopedCurrentContext { -public: - explicit ScopedCurrentContext(ContextHandle h_context) noexcept - : h_context_(std::move(h_context)) { - CUcontext target = as_cu(h_context_); - if (!target) { - return; - } +void warn_on_cuda_error(const char* operation, CUresult status, const char* detail = nullptr) noexcept; + +// 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. +CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* changed) noexcept { + *previous = nullptr; + *changed = 0; + CUcontext target = as_cu(h_context); + if (!target) { + return CUDA_SUCCESS; + } + + GILReleaseGuard gil; + CUresult status = p_cuCtxGetCurrent(previous); + if (status != CUDA_SUCCESS || *previous == target) { + return status; + } + status = p_cuCtxSetCurrent(target); + *changed = status == CUDA_SUCCESS; + 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; - status_ = p_cuCtxGetCurrent(&previous_); - if (status_ != CUDA_SUCCESS || previous_ == target) { - return; + restore_status = p_cuCtxSetCurrent(previous); + } + if (operation_status != CUDA_SUCCESS && restore_status != CUDA_SUCCESS) { + warn_on_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restore_status); + } + return operation_status != CUDA_SUCCESS ? operation_status : restore_status; +} + +// Require a callable to be invocable without throwing. +#define ASSERT_NOTHROW_INVOCABLE(...) \ + static_assert(std::is_nothrow_invocable_v<__VA_ARGS__>, "operation must be noexcept") + +// Store a stream and any state needed to preserve deallocation ordering. +struct DeallocationStream { + StreamHandle h_stream; + std::thread::id ptds_tid{}; +}; + +// Return whether a stream handle needs a current context to resolve it. +bool is_default_stream(CUstream stream) noexcept { + return stream == nullptr || stream == CU_STREAM_LEGACY || stream == CU_STREAM_PER_THREAD; +} + +// Return the context a deallocation-stream token must run under. Real streams +// resolve their own context; default-stream tokens use the context bound at +// allocation time. Warn when PTDS deallocation crosses host threads. +ContextHandle deallocation_context(const DeallocationStream& stream) noexcept { + if (!is_default_stream(as_cu(stream.h_stream))) { + return {}; + } + 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 " + "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"); + } + return get_stream_context(stream.h_stream); +} + +// Run an operation with the requested context current. +template +CUresult invoke_in_context(const ContextHandle& h_context, Fn&& operation, Args&&... args) noexcept { + ASSERT_NOTHROW_INVOCABLE(Fn&&, Args&&...); + CUcontext previous = nullptr; + int changed = 0; + CUresult status = enter_context(h_context, &previous, &changed); + if (status == CUDA_SUCCESS) { + status = std::invoke(std::forward(operation), std::forward(args)...); + } + return exit_context(previous, changed, status); +} + +// Run a creation operation and undo it if context restoration fails. +// Context-independent undo always runs. Context-sensitive undo runs only +// after verifying that the target context remains current; otherwise the +// resource leaks rather than risking cleanup in the wrong context. +template +CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operation, + Undo&& undo, bool undo_requires_target_context) noexcept { + ASSERT_NOTHROW_INVOCABLE(Fn&&); + ASSERT_NOTHROW_INVOCABLE(Undo&&); + CUcontext previous = nullptr; + int changed = 0; + CUresult status = enter_context(h_context, &previous, &changed); + if (status != CUDA_SUCCESS) { + return status; + } + status = std::invoke(std::forward(operation)); + CUresult composite = exit_context(previous, changed, status); + if (status == CUDA_SUCCESS && composite != CUDA_SUCCESS) { + bool undo_ok = true; + if (undo_requires_target_context) { + CUcontext current = nullptr; + undo_ok = p_cuCtxGetCurrent(¤t) == CUDA_SUCCESS + && current == as_cu(h_context); + } + if (undo_ok) { + std::invoke(std::forward(undo)); } - status_ = p_cuCtxSetCurrent(target); - changed_ = status_ == CUDA_SUCCESS; } + return composite; +} - ~ScopedCurrentContext() { - if (changed_) { - GILReleaseGuard gil; - CUresult status = p_cuCtxSetCurrent(previous_); - if (status != CUDA_SUCCESS) { - std::fprintf( - stderr, - "Warning: cuCtxSetCurrent (restoring the caller's context) " - "failed (CUDA error %d)\n", - static_cast(status)); - } +// 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 +CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, + Fn&& operation, Args&&... args) noexcept { + ASSERT_NOTHROW_INVOCABLE(Fn&&, Args&&...); + CUcontext previous = nullptr; + int changed = 0; + CUresult status = enter_context(h_context, &previous, &changed); + if (status != CUDA_SUCCESS) { + warn_on_cuda_error(name, status, + "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); } } + 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"); + } + return status; +} + +#undef ASSERT_NOTHROW_INVOCABLE - CUresult status() const noexcept { return status_; } +// Decorate a CUDA operation to warn whenever it returns an error. +template +class WarnOnFailure { +public: + explicit WarnOnFailure(const char* operation) noexcept : operation_(operation) {} - ScopedCurrentContext(const ScopedCurrentContext&) = delete; - ScopedCurrentContext& operator=(const ScopedCurrentContext&) = delete; + template + CUresult operator()(Args&&... args) const noexcept { + CUresult status = Function(std::forward(args)...); + if (status != CUDA_SUCCESS) { + warn_on_cuda_error(operation_, status); + } + return status; + } private: - ContextHandle h_context_; - CUcontext previous_ = nullptr; - bool changed_ = false; - CUresult status_ = CUDA_SUCCESS; + const char* operation_; }; +// Warning-decorated CUDA operations used by non-throwing cleanup paths. +const WarnOnFailure pw_cuStreamDestroy{"cuStreamDestroy"}; +const WarnOnFailure pw_cuEventDestroy{"cuEventDestroy"}; +const WarnOnFailure pw_cuMemFree{"cuMemFree"}; +const WarnOnFailure pw_cuMemFreeAsync{"cuMemFreeAsync"}; +const WarnOnFailure pw_cuArrayDestroy{"cuArrayDestroy"}; +const WarnOnFailure pw_cuMipmappedArrayDestroy{"cuMipmappedArrayDestroy"}; +const WarnOnFailure pw_cuTexObjectDestroy{"cuTexObjectDestroy"}; +const WarnOnFailure pw_cuSurfObjectDestroy{"cuSurfObjectDestroy"}; + } // namespace +// Synchronize the provided context. +CUresult context_synchronize(const ContextHandle& h_context) noexcept { + if (!h_context) { + return CUDA_ERROR_INVALID_CONTEXT; + } + return invoke_in_context(h_context, []() noexcept { + return p_cuCtxSynchronize(); + }); +} + +// Query the stream priority range for the provided context. +CUresult context_get_stream_priority_range(const ContextHandle& h_context, + int* least_priority, + int* greatest_priority) noexcept { + if (!h_context) { + return CUDA_ERROR_INVALID_CONTEXT; + } + return invoke_in_context(h_context, [&]() noexcept { + return p_cuCtxGetStreamPriorityRange(least_priority, greatest_priority); + }); +} + // ============================================================================ // CUDA user-object deferred cleanup // @@ -478,12 +647,14 @@ class HandleRegistry { // Thread-local status of the most recent CUDA API call in this module. static thread_local CUresult err = CUDA_SUCCESS; +// Return and clear the calling thread's most recent CUDA error. CUresult get_last_error() noexcept { CUresult e = err; err = CUDA_SUCCESS; return e; } +// Return the calling thread's most recent CUDA error without clearing it. CUresult peek_last_error() noexcept { return err; } @@ -676,22 +847,21 @@ static HandleRegistry stream_registry; StreamHandle create_stream_handle(const ContextHandle& h_ctx, unsigned int flags, int priority) { GILReleaseGuard gil; - CUstream stream; - - // Dispatch: green context uses cuGreenCtxStreamCreate, primary uses cuStreamCreateWithPriority + CUstream stream = nullptr; GreenCtxHandle h_green = get_context_green_ctx(h_ctx); if (h_green) { - if (!p_cuGreenCtxStreamCreate) { - err = CUDA_ERROR_NOT_SUPPORTED; - return {}; - } - if (CUDA_SUCCESS != (err = p_cuGreenCtxStreamCreate(&stream, as_cu(h_green), flags, priority))) { - return {}; - } + err = p_cuGreenCtxStreamCreate + ? p_cuGreenCtxStreamCreate(&stream, as_cu(h_green), flags, priority) + : CUDA_ERROR_NOT_SUPPORTED; } else { - if (CUDA_SUCCESS != (err = p_cuStreamCreateWithPriority(&stream, flags, priority))) { - return {}; - } + err = invoke_in_context_or_undo( + h_ctx, + [&]() noexcept { return p_cuStreamCreateWithPriority(&stream, flags, priority); }, + [&]() noexcept { pw_cuStreamDestroy(stream); }, + /*undo_requires_target_context=*/false); + } + if (err != CUDA_SUCCESS) { + return {}; } auto box = std::shared_ptr( @@ -699,7 +869,7 @@ StreamHandle create_stream_handle(const ContextHandle& h_ctx, unsigned int flags [](const StreamBox* b) { stream_registry.unregister_handle(b->resource); GILReleaseGuard gil; - p_cuStreamDestroy(b->resource); + pw_cuStreamDestroy(b->resource); delete b; } ); @@ -769,6 +939,7 @@ void py_object_user_object_destroy(void* py_object) noexcept { Py_DECREF(reinterpret_cast(py_object)); } +// Return the context retained by a stream handle. ContextHandle get_stream_context(const StreamHandle& h) noexcept { return h ? get_box(h)->h_context : ContextHandle{}; } @@ -797,12 +968,6 @@ StreamHandle get_per_thread_stream() { // detected and warnings can be issued. // ============================================================================ -// ptds_tid is std::thread::id{} except for CU_STREAM_PER_THREAD. -struct DeallocationStream { - StreamHandle h_stream; - std::thread::id ptds_tid{}; -}; - // Real streams are copied unchanged. Default-stream tokens without an embedded // context are bound to the current context. Returns false (and sets err) when a // default-stream token cannot be bound because no context is current. @@ -814,9 +979,7 @@ static bool make_deallocation_stream( } const CUstream stream = as_cu(h); - if (stream != nullptr - && stream != CU_STREAM_LEGACY - && stream != CU_STREAM_PER_THREAD) { + if (!is_default_stream(stream)) { out = DeallocationStream{h, {}}; return true; } @@ -845,35 +1008,6 @@ static bool make_deallocation_stream( return true; } -template -CUresult with_deallocation_context( - const DeallocationStream& stream, - const char* operation, - Fn&& fn) 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 " - "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"); - } - ScopedCurrentContext context(get_stream_context(stream.h_stream)); - CUresult status = context.status(); - if (status == CUDA_SUCCESS) { - status = fn(stream); - } - if (status != CUDA_SUCCESS) { - std::fprintf( - stderr, - "Warning: %s failed during resource destruction (CUDA error %d)\n", - operation, - static_cast(status)); - } - return status; -} - // ============================================================================ // Event Handles // ============================================================================ @@ -912,6 +1046,7 @@ int get_event_device_id(const EventHandle& h) noexcept { return h ? get_box(h)->device_id : -1; } +// Return the context retained by an event handle. ContextHandle get_event_context(const EventHandle& h) noexcept { return h ? get_box(h)->h_context : ContextHandle{}; } @@ -923,17 +1058,22 @@ EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, bool timing_enabled, bool is_blocking_sync, bool ipc_enabled, int device_id) { GILReleaseGuard gil; - CUevent event; - if (CUDA_SUCCESS != (err = p_cuEventCreate(&event, flags))) { + CUevent event = nullptr; + err = invoke_in_context_or_undo( + h_ctx, + [&]() noexcept { return p_cuEventCreate(&event, flags); }, + [&]() noexcept { pw_cuEventDestroy(event); }, + /*undo_requires_target_context=*/false); + if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( new EventBox{event, timing_enabled, is_blocking_sync, ipc_enabled, device_id, h_ctx}, - [h_ctx](const EventBox* b) { + [](const EventBox* b) { event_registry.unregister_handle(b->resource); GILReleaseGuard gil; - p_cuEventDestroy(b->resource); + pw_cuEventDestroy(b->resource); delete b; } ); @@ -967,7 +1107,7 @@ EventHandle create_event_handle_ipc(const CUipcEventHandle& ipc_handle, [](const EventBox* b) { event_registry.unregister_handle(b->resource); GILReleaseGuard gil; - p_cuEventDestroy(b->resource); + pw_cuEventDestroy(b->resource); delete b; } ); @@ -1080,12 +1220,13 @@ static DevicePtrBox* get_box(const DevicePtrHandle& h) { ); } +// Return the stream that orders a device pointer's deallocation. StreamHandle deallocation_stream(const DevicePtrHandle& h) noexcept { return get_box(h)->deallocation.h_stream; } -CUresult set_deallocation_stream( - const DevicePtrHandle& h, const StreamHandle& h_stream) noexcept { +// Replace the stream that orders a device pointer's deallocation. +CUresult set_deallocation_stream(const DevicePtrHandle& h, const StreamHandle& h_stream) noexcept { if (!h) { return CUDA_ERROR_INVALID_VALUE; } @@ -1106,7 +1247,7 @@ DevicePtrHandle deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle& h DeallocationStream ds; if (!make_deallocation_stream(h_stream, ds)) { - p_cuMemFreeAsync(ptr, as_cu(h_stream)); + pw_cuMemFreeAsync(ptr, as_cu(h_stream)); return {}; } @@ -1114,10 +1255,10 @@ DevicePtrHandle deviceptr_alloc_from_pool(size_t size, const MemoryPoolHandle& h new DevicePtrBox{ptr, std::move(ds)}, [h_pool](DevicePtrBox* b) { GILReleaseGuard gil; - with_deallocation_context( - b->deallocation, - "cuMemFreeAsync", - [b](const DeallocationStream& stream) { + const DeallocationStream& stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), "cuMemFreeAsync", + [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); }); @@ -1136,7 +1277,7 @@ DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) DeallocationStream ds; if (!make_deallocation_stream(h_stream, ds)) { - p_cuMemFreeAsync(ptr, as_cu(h_stream)); + pw_cuMemFreeAsync(ptr, as_cu(h_stream)); return {}; } @@ -1144,10 +1285,10 @@ DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) new DevicePtrBox{ptr, std::move(ds)}, [](DevicePtrBox* b) { GILReleaseGuard gil; - with_deallocation_context( - b->deallocation, - "cuMemFreeAsync", - [b](const DeallocationStream& stream) { + const DeallocationStream& stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), "cuMemFreeAsync", + [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); }); @@ -1157,22 +1298,18 @@ DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) return DevicePtrHandle(box, &box->resource); } -DevicePtrHandle deviceptr_alloc(size_t size) { - GILReleaseGuard gil; - CUdeviceptr ptr; - if (CUDA_SUCCESS != (err = p_cuMemAlloc(&ptr, size))) { - return {}; +// Allocate device memory synchronously with the provided context current. +CUresult deviceptr_alloc_raw(CUdeviceptr* ptr, size_t size, + const ContextHandle& h_context) noexcept { + if (!h_context) { + return CUDA_ERROR_INVALID_CONTEXT; } - - auto box = std::shared_ptr( - new DevicePtrBox{ptr, DeallocationStream{}}, - [](DevicePtrBox* b) { - GILReleaseGuard gil; - p_cuMemFree(b->resource); - delete b; - } - ); - return DevicePtrHandle(box, &box->resource); + GILReleaseGuard gil; + return invoke_in_context_or_undo( + h_context, + [&]() noexcept { return p_cuMemAlloc(ptr, size); }, + [&]() noexcept { pw_cuMemFree(*ptr); }, + /*undo_requires_target_context=*/false); } DevicePtrHandle deviceptr_alloc_host(size_t size) { @@ -1236,10 +1373,10 @@ DevicePtrHandle deviceptr_create_mapped_graphics( [h_resource](DevicePtrBox* b) { GILReleaseGuard gil; CUgraphicsResource resource = as_cu(h_resource); - with_deallocation_context( - b->deallocation, - "cuGraphicsUnmapResources", - [b, &resource](const DeallocationStream& stream) { + const DeallocationStream& stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), "cuGraphicsUnmapResources", + [&]() noexcept { return p_cuGraphicsUnmapResources( 1, &resource, as_cu(stream.h_stream)); }); @@ -1275,12 +1412,11 @@ DevicePtrHandle deviceptr_create_with_mr(CUdeviceptr ptr, size_t size, PyObject* GILAcquireGuard gil; if (gil.acquired()) { if (mr_dealloc_cb) { - with_deallocation_context( - b->deallocation, - "MemoryResource deallocate", - [mr, size, b](const DeallocationStream& stream) { - mr_dealloc_cb( - mr, b->resource, size, stream.h_stream); + const DeallocationStream& stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), "MemoryResource.deallocate", + [&]() noexcept { + mr_dealloc_cb(mr, b->resource, size, stream.h_stream); return CUDA_SUCCESS; }); } @@ -1372,7 +1508,7 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* DeallocationStream ds; if (!make_deallocation_stream(h_stream, ds)) { - p_cuMemFreeAsync(ptr, as_cu(h_stream)); + pw_cuMemFreeAsync(ptr, as_cu(h_stream)); return {}; } @@ -1381,10 +1517,10 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* [h_pool, key](DevicePtrBox* b) { ipc_ptr_cache.unregister_handle(key); GILReleaseGuard gil; - with_deallocation_context( - b->deallocation, - "cuMemFreeAsync", - [b](const DeallocationStream& stream) { + const DeallocationStream& stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), "cuMemFreeAsync", + [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); }); @@ -1404,7 +1540,7 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* DeallocationStream ds; if (!make_deallocation_stream(h_stream, ds)) { - p_cuMemFreeAsync(ptr, as_cu(h_stream)); + pw_cuMemFreeAsync(ptr, as_cu(h_stream)); return {}; } @@ -1412,10 +1548,10 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void* new DevicePtrBox{ptr, std::move(ds)}, [h_pool](DevicePtrBox* b) { GILReleaseGuard gil; - with_deallocation_context( - b->deallocation, - "cuMemFreeAsync", - [b](const DeallocationStream& stream) { + const DeallocationStream& stream = b->deallocation; + cleanup_in_context( + deallocation_context(stream), "cuMemFreeAsync", + [&]() noexcept { return p_cuMemFreeAsync( b->resource, as_cu(stream.h_stream)); }); @@ -2671,12 +2807,18 @@ struct ArrayBox { // Non-null only for a mipmap-level view: keeps the parent mipmap (the real // owner of the level's storage) alive for as long as the level is held. MipmappedArrayHandle h_parent; + ContextHandle h_context; }; struct MipmappedArrayBox { CUmipmappedArray resource; + ContextHandle h_context; }; +// Texture and surface objects are per-context pool indices. Destroying one +// with the wrong context current can silently succeed without freeing it or +// can free an unrelated object, so destruction must enter the creating +// context. Handle-based resources resolve their own context and must not. struct TexObjectBox { // Tagged so TexObjectHandle is a distinct C++ type from DevicePtrHandle / // SurfObjectHandle (all wrap `unsigned long long`). @@ -2685,31 +2827,64 @@ struct TexObjectBox { // DevicePtrHandle). The texture's resource is a union; we only need to keep // whichever backing it was built from alive, never to dereference it. std::shared_ptr h_backing; + ContextHandle h_context; }; struct SurfObjectBox { SurfObjectValue resource; OpaqueArrayHandle h_array; // surfaces are always array-backed + ContextHandle h_context; }; + +// Recover an array's owning box from its aliased resource pointer. +const ArrayBox* get_array_box(const OpaqueArrayHandle& h) noexcept { + const CUarray* p = h.get(); + return reinterpret_cast( + reinterpret_cast(p) - offsetof(ArrayBox, resource)); +} + +// Recover a mipmapped array's owning box from its aliased resource pointer. +const MipmappedArrayBox* get_mipmapped_array_box(const MipmappedArrayHandle& h) noexcept { + const CUmipmappedArray* p = h.get(); + return reinterpret_cast( + reinterpret_cast(p) + - offsetof(MipmappedArrayBox, resource)); +} + +// Wrap an array with shared owning-destruction behavior. +static OpaqueArrayHandle wrap_array_owned(CUarray arr, ContextHandle h_context) { + auto box = std::shared_ptr( + new ArrayBox{arr, {}, std::move(h_context)}, + [](const ArrayBox* b) { + GILReleaseGuard gil; + pw_cuArrayDestroy(b->resource); + delete b; + } + ); + return OpaqueArrayHandle(box, &box->resource); +} + } // namespace -OpaqueArrayHandle create_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc) { +OpaqueArrayHandle create_array_handle(const ContextHandle& h_context, const CUDA_ARRAY3D_DESCRIPTOR& desc) { GILReleaseGuard gil; - CUarray arr; - if (CUDA_SUCCESS != (err = p_cuArray3DCreate(&arr, &desc))) { + CUarray arr = nullptr; + err = invoke_in_context_or_undo( + h_context, + [&]() noexcept { return p_cuArray3DCreate(&arr, &desc); }, + [&]() noexcept { pw_cuArrayDestroy(arr); }, + /*undo_requires_target_context=*/false); + if (err != CUDA_SUCCESS) { return {}; } - // Allocation and adoption share the same owning lifetime; the only - // difference is who calls cuArray3DCreate. Delegate so the owning box and - // its destroy-on-last-ref deleter are defined in exactly one place. - return create_array_handle_owning(arr); + return wrap_array_owned(arr, h_context); } OpaqueArrayHandle create_array_handle_ref(CUarray arr) { if (!arr) { return {}; } - auto box = std::make_shared(ArrayBox{arr, {}}); + auto box = std::make_shared(ArrayBox{arr, {}, {}}); return OpaqueArrayHandle(box, &box->resource); } @@ -2717,64 +2892,81 @@ OpaqueArrayHandle create_array_handle_owning(CUarray arr) { if (!arr) { return {}; } - auto box = std::shared_ptr( - new ArrayBox{arr, {}}, - [](const ArrayBox* b) { - GILReleaseGuard gil; - p_cuArrayDestroy(b->resource); - delete b; - } - ); - return OpaqueArrayHandle(box, &box->resource); + return wrap_array_owned(arr, {}); +} + +// Return the context retained by an array handle. +ContextHandle get_array_context(const OpaqueArrayHandle& h) noexcept { + return h ? get_array_box(h)->h_context : ContextHandle{}; } OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level) { GILReleaseGuard gil; CUarray arr; + ContextHandle h_context = h_mip ? get_mipmapped_array_box(h_mip)->h_context : ContextHandle{}; if (CUDA_SUCCESS != (err = p_cuMipmappedArrayGetLevel(&arr, as_cu(h_mip), level))) { return {}; } // Non-owning level view: storage belongs to the mipmap. Embed the mipmap // handle so the parent outlives this level; the deleter does not destroy. auto box = std::shared_ptr( - new ArrayBox{arr, h_mip}, + new ArrayBox{arr, h_mip, h_context}, [](const ArrayBox* b) { delete b; } ); return OpaqueArrayHandle(box, &box->resource); } -MipmappedArrayHandle create_mipmapped_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc, +MipmappedArrayHandle create_mipmapped_array_handle(const ContextHandle& h_context, + const CUDA_ARRAY3D_DESCRIPTOR& desc, unsigned int num_levels) { GILReleaseGuard gil; - CUmipmappedArray mip; - if (CUDA_SUCCESS != (err = p_cuMipmappedArrayCreate(&mip, &desc, num_levels))) { + CUmipmappedArray mip = nullptr; + err = invoke_in_context_or_undo( + h_context, + [&]() noexcept { return p_cuMipmappedArrayCreate(&mip, &desc, num_levels); }, + [&]() noexcept { pw_cuMipmappedArrayDestroy(mip); }, + /*undo_requires_target_context=*/false); + if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( - new MipmappedArrayBox{mip}, + new MipmappedArrayBox{mip, h_context}, [](const MipmappedArrayBox* b) { GILReleaseGuard gil; - p_cuMipmappedArrayDestroy(b->resource); + pw_cuMipmappedArrayDestroy(b->resource); delete b; } ); return MipmappedArrayHandle(box, &box->resource); } +// Return the context retained by a mipmapped array handle. +ContextHandle get_mipmapped_array_context(const MipmappedArrayHandle& h) noexcept { + return h ? get_mipmapped_array_box(h)->h_context : ContextHandle{}; +} + namespace { TexObjectHandle make_tex_object_handle(const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, - std::shared_ptr h_backing) { + std::shared_ptr h_backing, + const ContextHandle& h_context) { GILReleaseGuard gil; - CUtexObject obj; - if (CUDA_SUCCESS != (err = p_cuTexObjectCreate(&obj, &res, &tex, nullptr))) { + CUtexObject obj = 0; + err = invoke_in_context_or_undo( + h_context, + [&]() noexcept { return p_cuTexObjectCreate(&obj, &res, &tex, nullptr); }, + [&]() noexcept { pw_cuTexObjectDestroy(obj); }, + /*undo_requires_target_context=*/true); + if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( - new TexObjectBox{TexObjectValue{obj}, std::move(h_backing)}, + new TexObjectBox{TexObjectValue{obj}, std::move(h_backing), h_context}, [](const TexObjectBox* b) { GILReleaseGuard gil; - p_cuTexObjectDestroy(b->resource.raw); + cleanup_in_context(b->h_context, "cuTexObjectDestroy", [&]() noexcept { + return p_cuTexObjectDestroy(b->resource.raw); + }); delete b; } ); @@ -2782,36 +2974,47 @@ TexObjectHandle make_tex_object_handle(const CUDA_RESOURCE_DESC& res, } } // namespace -TexObjectHandle create_tex_object_handle_array(const CUDA_RESOURCE_DESC& res, +TexObjectHandle create_tex_object_handle_array(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, const OpaqueArrayHandle& h_backing) { - return make_tex_object_handle(res, tex, h_backing); + return make_tex_object_handle(res, tex, h_backing, h_context); } -TexObjectHandle create_tex_object_handle_mipmap(const CUDA_RESOURCE_DESC& res, +TexObjectHandle create_tex_object_handle_mipmap(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, const MipmappedArrayHandle& h_backing) { - return make_tex_object_handle(res, tex, h_backing); + return make_tex_object_handle(res, tex, h_backing, h_context); } -TexObjectHandle create_tex_object_handle_linear(const CUDA_RESOURCE_DESC& res, +TexObjectHandle create_tex_object_handle_linear(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, const DevicePtrHandle& h_backing) { - return make_tex_object_handle(res, tex, h_backing); + return make_tex_object_handle(res, tex, h_backing, h_context); } -SurfObjectHandle create_surf_object_handle(const CUDA_RESOURCE_DESC& res, +SurfObjectHandle create_surf_object_handle(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const OpaqueArrayHandle& h_backing) { GILReleaseGuard gil; - CUsurfObject obj; - if (CUDA_SUCCESS != (err = p_cuSurfObjectCreate(&obj, &res))) { + CUsurfObject obj = 0; + err = invoke_in_context_or_undo( + h_context, + [&]() noexcept { return p_cuSurfObjectCreate(&obj, &res); }, + [&]() noexcept { pw_cuSurfObjectDestroy(obj); }, + /*undo_requires_target_context=*/true); + if (err != CUDA_SUCCESS) { return {}; } auto box = std::shared_ptr( - new SurfObjectBox{SurfObjectValue{obj}, h_backing}, + new SurfObjectBox{SurfObjectValue{obj}, h_backing, h_context}, [](const SurfObjectBox* b) { GILReleaseGuard gil; - p_cuSurfObjectDestroy(b->resource.raw); + cleanup_in_context(b->h_context, "cuSurfObjectDestroy", [&]() 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 ff1a12a4618..57ac9a244d6 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.hpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.hpp @@ -64,10 +64,15 @@ void clear_last_error() noexcept; // function pointers extracted from cuda.bindings.cydriver.__pyx_capi__. // ============================================================================ +extern decltype(&cuGetErrorName) p_cuGetErrorName; +extern decltype(&cuGetErrorString) p_cuGetErrorString; + extern decltype(&cuDevicePrimaryCtxRetain) p_cuDevicePrimaryCtxRetain; extern decltype(&cuDevicePrimaryCtxRelease) p_cuDevicePrimaryCtxRelease; extern decltype(&cuCtxGetCurrent) p_cuCtxGetCurrent; extern decltype(&cuCtxSetCurrent) p_cuCtxSetCurrent; +extern decltype(&cuCtxSynchronize) p_cuCtxSynchronize; +extern decltype(&cuCtxGetStreamPriorityRange) p_cuCtxGetStreamPriorityRange; extern decltype(&cuGreenCtxCreate) p_cuGreenCtxCreate; extern decltype(&cuGreenCtxDestroy) p_cuGreenCtxDestroy; extern decltype(&cuCtxFromGreenCtx) p_cuCtxFromGreenCtx; @@ -246,6 +251,17 @@ ContextHandle get_primary_context(int device_id); // Returns empty handle if no context is current (caller must check) ContextHandle get_current_context(); +// Synchronize the provided context. +// Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. +CUresult context_synchronize(const ContextHandle& h_context) noexcept; + +// Query the stream priority range for the provided context. +// Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. +CUresult context_get_stream_priority_range( + const ContextHandle& h_context, + int* least_priority, + int* greatest_priority) noexcept; + // ============================================================================ // Stream handle functions // ============================================================================ @@ -371,10 +387,11 @@ DevicePtrHandle deviceptr_alloc_from_pool( // Returns empty handle on error (caller must check). DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream); -// Allocate device memory synchronously via cuMemAlloc. -// When the last reference is released, cuMemFree is called. -// Returns empty handle on error (caller must check). -DevicePtrHandle deviceptr_alloc(size_t size); +// Allocate device memory synchronously via cuMemAlloc with the provided +// context current. The caller owns the pointer and releases it with cuMemFree. +// Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. +CUresult deviceptr_alloc_raw(CUdeviceptr* ptr, size_t size, + const ContextHandle& h_context) noexcept; // Allocate pinned host memory via cuMemAllocHost. // When the last reference is released, cuMemFreeHost is called. @@ -739,7 +756,7 @@ FileDescriptorHandle create_fd_handle_ref(int fd); // Create an owning CUDA array via cuArray3DCreate. // When the last reference is released, cuArrayDestroy is called automatically. // Returns empty handle on error (caller must check). -OpaqueArrayHandle create_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc); +OpaqueArrayHandle create_array_handle(const ContextHandle& h_context, const CUDA_ARRAY3D_DESCRIPTOR& desc); // Create a non-owning array handle (references an existing CUarray). // Use for arrays owned elsewhere (e.g. graphics interop). Never destroyed here. @@ -749,6 +766,9 @@ OpaqueArrayHandle create_array_handle_ref(CUarray arr); // When the last reference is released, cuArrayDestroy is called automatically. OpaqueArrayHandle create_array_handle_owning(CUarray arr); +// Return the context dependency associated with an array, if known. +ContextHandle get_array_context(const OpaqueArrayHandle& h) noexcept; + // Create a non-owning handle to a mipmap level via cuMipmappedArrayGetLevel. // The level CUarray is owned by the mipmap; the parent MipmappedArrayHandle is // embedded in the box so it outlives the level view. No destroy in the deleter. @@ -758,27 +778,35 @@ OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, u // Create an owning mipmapped array via cuMipmappedArrayCreate. // When the last reference is released, cuMipmappedArrayDestroy is called. // Returns empty handle on error (caller must check). -MipmappedArrayHandle create_mipmapped_array_handle(const CUDA_ARRAY3D_DESCRIPTOR& desc, +MipmappedArrayHandle create_mipmapped_array_handle(const ContextHandle& h_context, + const CUDA_ARRAY3D_DESCRIPTOR& desc, unsigned int num_levels); +// Return the context dependency associated with a mipmapped array, if known. +ContextHandle get_mipmapped_array_context(const MipmappedArrayHandle& h) noexcept; + // Create an owning texture object via cuTexObjectCreate, embedding the backing // resource handle (array / mipmapped array / linear-or-pitch2d device pointer) // so the backing always outlives the texture. cuTexObjectDestroy runs in the // deleter. Returns empty handle on error (caller must check). -TexObjectHandle create_tex_object_handle_array(const CUDA_RESOURCE_DESC& res, +TexObjectHandle create_tex_object_handle_array(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, const OpaqueArrayHandle& h_backing); -TexObjectHandle create_tex_object_handle_mipmap(const CUDA_RESOURCE_DESC& res, +TexObjectHandle create_tex_object_handle_mipmap(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, const MipmappedArrayHandle& h_backing); -TexObjectHandle create_tex_object_handle_linear(const CUDA_RESOURCE_DESC& res, +TexObjectHandle create_tex_object_handle_linear(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const CUDA_TEXTURE_DESC& tex, const DevicePtrHandle& h_backing); // Create an owning surface object via cuSurfObjectCreate, embedding the backing // array handle so it outlives the surface. cuSurfObjectDestroy runs in the // deleter. Returns empty handle on error (caller must check). -SurfObjectHandle create_surf_object_handle(const CUDA_RESOURCE_DESC& res, +SurfObjectHandle create_surf_object_handle(const ContextHandle& h_context, + const CUDA_RESOURCE_DESC& res, const OpaqueArrayHandle& h_backing); // ============================================================================ diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index 369f2b198d8..39ff7d3a26e 100644 --- a/cuda_core/cuda/core/_device.pyi +++ b/cuda_core/cuda/core/_device.pyi @@ -579,7 +579,7 @@ class Device: def memory_resource(self, mr: MemoryResource) -> None: ... @property def default_stream(self) -> Stream: - """Return default CUDA :obj:`~_stream.Stream` associated with this device. + """Return a default CUDA :obj:`~_stream.Stream` token. The type of default stream returned depends on if the environment variable CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM is set. @@ -587,6 +587,9 @@ class Device: If set, returns a per-thread default stream. Otherwise returns the legacy stream. + A default-stream token uses the device that is current when the token + is used. + """ def __int__(self) -> int: """Return device_id.""" @@ -611,7 +614,9 @@ class Device: Returns ------- :obj:`~_context.Context`, optional - Popped context. + The previous context, or ``None`` if no context was current. When + returned, its ``device_id`` identifies the device that was + previously current. Examples -------- @@ -643,7 +648,7 @@ class Device: """ def create_stream(self, obj: IsStreamType | None=None, options: StreamOptions | None=None) -> Stream: - """Create a :obj:`~_stream.Stream` object. + """Create or wrap a :obj:`~_stream.Stream` object. New stream objects can be created in two different ways: @@ -655,7 +660,7 @@ class Device: Note ---- - Device must be initialized. + Device must be initialized. New streams are created on this device. Parameters ---------- @@ -671,7 +676,7 @@ class Device: """ def create_event(self, options: EventOptions | None=None) -> Event: - """Create an :obj:`~_event.Event` object without recording it to a :obj:`~_stream.Stream`. + """Create an :obj:`~_event.Event` on this device without recording it to a :obj:`~_stream.Stream`. Note ---- @@ -714,7 +719,7 @@ class Device: """ def sync(self) -> None: - """Synchronize the device. + """Synchronize this device. Note ---- @@ -722,7 +727,7 @@ class Device: """ def create_graph_builder(self) -> GraphBuilder: - """Create a new :obj:`~graph.GraphBuilder` object. + """Create a new :obj:`~graph.GraphBuilder` on this device. Returns ------- @@ -731,12 +736,10 @@ class Device: """ def create_opaque_array(self, options: OpaqueArrayOptions) -> OpaqueArray: - """Create an :obj:`~cuda.core.texture.OpaqueArray` on the current device. + """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. Allocates an opaque, hardware-laid-out CUDA array for texture/surface - access. The array is created in the current CUDA context, so make this - device current with :meth:`set_current` before calling (mirroring - :meth:`create_stream` / :meth:`create_event`). + access. Note ---- @@ -755,12 +758,10 @@ class Device: .. versionadded:: 1.1.0 """ def create_mipmapped_array(self, options: MipmappedArrayOptions) -> MipmappedArray: - """Create a :obj:`~cuda.core.texture.MipmappedArray` on the current device. + """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. Allocates a mipmapped CUDA array for texture/surface access across - levels. The array is created in the current CUDA context, so make this - device current with :meth:`set_current` before calling (mirroring - :meth:`create_stream` / :meth:`create_event`). + levels. Note ---- @@ -779,15 +780,13 @@ class Device: .. versionadded:: 1.1.0 """ def create_texture_object(self, *, resource: ResourceDescriptor, options: TextureObjectOptions | None=None) -> TextureObject: - """Create a :obj:`~cuda.core.texture.TextureObject` on the current device. + """Create a :obj:`~cuda.core.texture.TextureObject` on this device. Binds a resource (an :obj:`~cuda.core.texture.OpaqueArray` / :obj:`~cuda.core.texture.MipmappedArray` / linear or pitch2d :obj:`~cuda.core.Buffer`, wrapped in a :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless texture for - kernel-side sampled reads. The object is created in the current CUDA - context, so make this device current with :meth:`set_current` before - calling (mirroring :meth:`create_stream` / :meth:`create_event`). + kernel-side sampled reads. The resource must belong to this device. Note ---- @@ -808,15 +807,12 @@ class Device: .. versionadded:: 1.1.0 """ def create_surface_object(self, *, resource: ResourceDescriptor) -> SurfaceObject: - """Create a :obj:`~cuda.core.texture.SurfaceObject` on the current device. + """Create a :obj:`~cuda.core.texture.SurfaceObject` on this device. Binds an :obj:`~cuda.core.texture.OpaqueArray` (via a :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless surface for kernel-side typed load/store. The backing array must have been created - with ``is_surface_load_store=True``. The object is created in the - current CUDA context, so make this device current with - :meth:`set_current` before calling (mirroring :meth:`create_stream` / - :meth:`create_event`). + with ``is_surface_load_store=True`` and must belong to this device. Note ---- diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index a52287a2aed..78245c6107d 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -23,6 +23,7 @@ from cuda.core._resource_handles cimport ( GreenCtxHandle, create_context_handle_ref, create_green_ctx_handle, + context_synchronize, get_primary_context, get_last_error, as_cu, @@ -37,7 +38,9 @@ from cuda.core._utils.cuda_utils import ( handle_return, runtime, ) -from cuda.core._stream cimport default_stream +from cuda.core._stream cimport ( + default_stream, +) from typing import TYPE_CHECKING @@ -1021,6 +1024,7 @@ class Device: raise CUDAError( f"Device {self._device_id} is not yet initialized, perhaps you forgot to call .set_current() first?" ) + Context_check_open(self._context) @classmethod @@ -1202,8 +1206,11 @@ class Device: from cuda.core._memory import DeviceMemoryResource self._memory_resource = DeviceMemoryResource(self._device_id) else: - from cuda.core._memory._legacy import _SynchronousMemoryResource - self._memory_resource = _SynchronousMemoryResource(self._device_id) + from cuda.core._memory._device_memory_resource import ( + _SynchronousMemoryResource, + ) + self._memory_resource = _SynchronousMemoryResource( + self._device_id, self._context) return self._memory_resource @@ -1215,7 +1222,7 @@ class Device: @property def default_stream(self) -> Stream: - """Return default CUDA :obj:`~_stream.Stream` associated with this device. + """Return a default CUDA :obj:`~_stream.Stream` token. The type of default stream returned depends on if the environment variable CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM is set. @@ -1223,6 +1230,9 @@ class Device: If set, returns a per-thread default stream. Otherwise returns the legacy stream. + A default-stream token uses the device that is current when the token + is used. + """ return default_stream() @@ -1261,7 +1271,9 @@ class Device: Returns ------- :obj:`~_context.Context`, optional - Popped context. + The previous context, or ``None`` if no context was current. When + returned, its ``device_id`` identifies the device that was + previously current. Examples -------- @@ -1276,6 +1288,7 @@ class Device: """ cdef ContextHandle h_context cdef cydriver.CUcontext prev_ctx, curr_ctx + cdef cydriver.CUdevice prev_dev cdef Context prev_owned = None if ctx is not None: @@ -1289,10 +1302,12 @@ class Device: ) if self._has_inited and self._context is not None: prev_owned = self._context - # prev_ctx is the previous context curr_ctx = as_cu(ctx._h_context) prev_ctx = NULL with nogil: + 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)) self._has_inited = True @@ -1300,7 +1315,8 @@ class Device: if prev_ctx != NULL: if prev_owned is not None and as_cu(prev_owned._h_context) == prev_ctx: return prev_owned - return Context._from_handle(Context, create_context_handle_ref(prev_ctx), self._device_id) + return Context._from_handle( + Context, create_context_handle_ref(prev_ctx), prev_dev) else: # use primary ctx h_context = get_primary_context(self._device_id) @@ -1381,7 +1397,7 @@ class Device: return Context._from_green_ctx(Context, h_green, self._device_id) def create_stream(self, obj: IsStreamType | None = None, options: StreamOptions | None = None) -> Stream: - """Create a :obj:`~_stream.Stream` object. + """Create or wrap a :obj:`~_stream.Stream` object. New stream objects can be created in two different ways: @@ -1393,7 +1409,7 @@ class Device: Note ---- - Device must be initialized. + Device must be initialized. New streams are created on this device. Parameters ---------- @@ -1412,7 +1428,7 @@ class Device: return Stream._init(obj=obj, options=options, device_id=self._device_id, ctx=self._context) def create_event(self, options: EventOptions | None = None) -> Event: - """Create an :obj:`~_event.Event` object without recording it to a :obj:`~_stream.Stream`. + """Create an :obj:`~_event.Event` on this device without recording it to a :obj:`~_stream.Stream`. Note ---- @@ -1462,7 +1478,7 @@ class Device: return self.memory_resource.allocate(size, stream=stream) def sync(self) -> None: - """Synchronize the device. + """Synchronize this device. Note ---- @@ -1470,10 +1486,14 @@ class Device: """ self._check_context_initialized() - handle_return(runtime.cudaDeviceSynchronize()) + cdef Context ctx = self._context + cdef cydriver.CUresult status + with nogil: + status = context_synchronize(ctx._h_context) + HANDLE_RETURN(status) def create_graph_builder(self) -> GraphBuilder: - """Create a new :obj:`~graph.GraphBuilder` object. + """Create a new :obj:`~graph.GraphBuilder` on this device. Returns ------- @@ -1487,12 +1507,10 @@ class Device: return GraphBuilder._init(self.create_stream()) def create_opaque_array(self, options: OpaqueArrayOptions) -> OpaqueArray: - """Create an :obj:`~cuda.core.texture.OpaqueArray` on the current device. + """Create an :obj:`~cuda.core.texture.OpaqueArray` on this device. Allocates an opaque, hardware-laid-out CUDA array for texture/surface - access. The array is created in the current CUDA context, so make this - device current with :meth:`set_current` before calling (mirroring - :meth:`create_stream` / :meth:`create_event`). + access. Note ---- @@ -1513,15 +1531,13 @@ class Device: from cuda.core.texture._array import _create_opaque_array self._check_context_initialized() - return _create_opaque_array(options) + return _create_opaque_array(options, self._context, self._device_id) def create_mipmapped_array(self, options: MipmappedArrayOptions) -> MipmappedArray: - """Create a :obj:`~cuda.core.texture.MipmappedArray` on the current device. + """Create a :obj:`~cuda.core.texture.MipmappedArray` on this device. Allocates a mipmapped CUDA array for texture/surface access across - levels. The array is created in the current CUDA context, so make this - device current with :meth:`set_current` before calling (mirroring - :meth:`create_stream` / :meth:`create_event`). + levels. Note ---- @@ -1542,20 +1558,18 @@ class Device: from cuda.core.texture._mipmapped_array import _create_mipmapped_array self._check_context_initialized() - return _create_mipmapped_array(options) + return _create_mipmapped_array(options, self._context, self._device_id) def create_texture_object( self, *, resource: ResourceDescriptor, options: TextureObjectOptions | None = None ) -> TextureObject: - """Create a :obj:`~cuda.core.texture.TextureObject` on the current device. + """Create a :obj:`~cuda.core.texture.TextureObject` on this device. Binds a resource (an :obj:`~cuda.core.texture.OpaqueArray` / :obj:`~cuda.core.texture.MipmappedArray` / linear or pitch2d :obj:`~cuda.core.Buffer`, wrapped in a :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless texture for - kernel-side sampled reads. The object is created in the current CUDA - context, so make this device current with :meth:`set_current` before - calling (mirroring :meth:`create_stream` / :meth:`create_event`). + kernel-side sampled reads. The resource must belong to this device. Note ---- @@ -1578,18 +1592,16 @@ class Device: from cuda.core.texture._texture import _create_texture_object self._check_context_initialized() - return _create_texture_object(resource, options) + return _create_texture_object( + resource, options, self._context, self._device_id) def create_surface_object(self, *, resource: ResourceDescriptor) -> SurfaceObject: - """Create a :obj:`~cuda.core.texture.SurfaceObject` on the current device. + """Create a :obj:`~cuda.core.texture.SurfaceObject` on this device. Binds an :obj:`~cuda.core.texture.OpaqueArray` (via a :obj:`~cuda.core.texture.ResourceDescriptor`) as a bindless surface for kernel-side typed load/store. The backing array must have been created - with ``is_surface_load_store=True``. The object is created in the - current CUDA context, so make this device current with - :meth:`set_current` before calling (mirroring :meth:`create_stream` / - :meth:`create_event`). + with ``is_surface_load_store=True`` and must belong to this device. Note ---- @@ -1611,7 +1623,8 @@ class Device: from cuda.core.texture._surface import _create_surface_object self._check_context_initialized() - return _create_surface_object(resource) + return _create_surface_object( + resource, self._context, self._device_id) cdef inline int Device_ensure_cuda_initialized() except? -1: diff --git a/cuda_core/cuda/core/_memory/_device_memory_resource.pyi b/cuda_core/cuda/core/_memory/_device_memory_resource.pyi index 897e1d03302..b74c344d2cd 100644 --- a/cuda_core/cuda/core/_memory/_device_memory_resource.pyi +++ b/cuda_core/cuda/core/_memory/_device_memory_resource.pyi @@ -4,9 +4,13 @@ import uuid from dataclasses import dataclass from cuda.core._device import Device +from cuda.core._memory._buffer import Buffer, MemoryResource from cuda.core._memory._ipc import IPCAllocationHandle from cuda.core._memory._memory_pool import _MemPool from cuda.core._memory._peer_access_utils import PeerAccessibleBySetProxy +from cuda.core._stream import Stream +from cuda.core.graph import GraphBuilder +from cuda.core.typing import DevicePointerType __all__ = ['DeviceMemoryResource', 'DeviceMemoryResourceOptions'] @@ -28,6 +32,19 @@ class DeviceMemoryResourceOptions: ipc_enabled: bool = False max_size: int = 0 +class _SynchronousMemoryResource(MemoryResource): + __slots__ = ('_context', '_device_id') + + def __init__(self, device_id: int, context=None) -> None: ... + def allocate(self, size: int, *, stream: Stream | GraphBuilder | None=None) -> Buffer: ... + def deallocate(self, ptr: DevicePointerType, size: int, *, stream: Stream | GraphBuilder | None=None) -> None: ... + @property + def is_device_accessible(self) -> bool: ... + @property + def is_host_accessible(self) -> bool: ... + @property + def device_id(self) -> int: ... + class DeviceMemoryResource(_MemPool): """ A device memory resource managing a stream-ordered memory pool. diff --git a/cuda_core/cuda/core/_memory/_device_memory_resource.pyx b/cuda_core/cuda/core/_memory/_device_memory_resource.pyx index d72b0e45ebc..426ddb6d8dc 100644 --- a/cuda_core/cuda/core/_memory/_device_memory_resource.pyx +++ b/cuda_core/cuda/core/_memory/_device_memory_resource.pyx @@ -4,7 +4,11 @@ from __future__ import annotations +from libc.stdint cimport uintptr_t + from cuda.bindings cimport cydriver +from cuda.core._context cimport Context +from cuda.core._memory._buffer cimport Buffer, MemoryResource from cuda.core._memory._location cimport cumemlocation_from_id from cuda.core._memory._memory_pool cimport ( _MemPool, MP_check_open, MP_init_create_pool, MP_raise_release_threshold, @@ -12,10 +16,14 @@ from cuda.core._memory._memory_pool cimport ( from cuda.core._memory cimport _ipc from cuda.core._memory._ipc cimport IPCAllocationHandle from cuda.core._resource_handles cimport ( + ContextHandle, as_cu, + deviceptr_alloc_raw, get_device_mempool, get_last_error, + get_primary_context, ) +from cuda.core._stream cimport Stream, Stream_accept from cuda.core._utils.cuda_utils cimport ( check_or_create_options, HANDLE_RETURN, @@ -34,6 +42,8 @@ from typing import TYPE_CHECKING if TYPE_CHECKING: from cuda.core._device import Device + from cuda.core.graph import GraphBuilder + from cuda.core.typing import DevicePointerType __all__ = ['DeviceMemoryResource', 'DeviceMemoryResourceOptions'] @@ -57,6 +67,68 @@ cdef class DeviceMemoryResourceOptions: max_size : int = 0 +class _SynchronousMemoryResource(MemoryResource): + __slots__ = ("_context", "_device_id") + + def __init__(self, device_id: int, context=None) -> None: + cdef ContextHandle h_context + from .._device import Device + + self._device_id = Device(device_id).device_id + if context is None: + h_context = get_primary_context(self._device_id) + if not h_context: + HANDLE_RETURN(get_last_error()) + context = Context._from_handle( + Context, h_context, self._device_id) + self._context = context + + def allocate( + self, + size_t size, + *, + stream: Stream | GraphBuilder | None = None, + ) -> Buffer: + # cuMemAlloc is synchronous; stream is accepted (and validated) + # for interface conformance but not used. + if stream is not None: + Stream_accept(stream) + + cdef Context context = self._context + cdef cydriver.CUdeviceptr ptr = 0 + if size: + with nogil: + HANDLE_RETURN(deviceptr_alloc_raw(&ptr, size, context._h_context)) + return Buffer._init(ptr, size, self) + + def deallocate( + self, + ptr: DevicePointerType, + size_t size, + *, + stream: Stream | GraphBuilder | None = None, + ) -> None: + if stream is not None: + Stream_accept(stream).sync() + cdef cydriver.CUdeviceptr devptr + if size: + devptr = int(ptr) + with nogil: + HANDLE_RETURN(cydriver.cuMemFree(devptr)) + + @property + def is_device_accessible(self) -> bool: + return True + + @property + def is_host_accessible(self) -> bool: + return False + + @property + def device_id(self) -> int: + return self._device_id + + cdef class DeviceMemoryResource(_MemPool): """ A device memory resource managing a stream-ordered memory pool. diff --git a/cuda_core/cuda/core/_memory/_legacy.py b/cuda_core/cuda/core/_memory/_legacy.py index 4acbcb54e3a..a2a8843a448 100644 --- a/cuda_core/cuda/core/_memory/_legacy.py +++ b/cuda_core/cuda/core/_memory/_legacy.py @@ -96,47 +96,3 @@ def is_host_accessible(self) -> bool: def device_id(self) -> int: """This memory resource is not bound to any GPU.""" raise RuntimeError("a pinned memory resource is not bound to any GPU") - - -class _SynchronousMemoryResource(MemoryResource): - __slots__ = ("_device_id",) - - def __init__(self, device_id: int) -> None: - from .._device import Device - - self._device_id = Device(device_id).device_id - - def allocate(self, size: int, *, stream: Stream | GraphBuilder | None = None) -> Buffer: - # cuMemAlloc is synchronous; stream is accepted (and validated) - # for interface conformance but not used. - from cuda.core._stream import Stream_accept - - if stream is not None: - Stream_accept(stream) - if size: - err, ptr = driver.cuMemAlloc(size) - raise_if_driver_error(err) - else: - ptr = 0 - return Buffer._init(ptr, size, self) - - def deallocate(self, ptr: DevicePointerType, size: int, *, stream: Stream | GraphBuilder | None = None) -> None: - from cuda.core._stream import Stream_accept - - if stream is not None: - Stream_accept(stream).sync() - if size: - (err,) = driver.cuMemFree(ptr) - raise_if_driver_error(err) - - @property - def is_device_accessible(self) -> bool: - return True - - @property - def is_host_accessible(self) -> bool: - return False - - @property - def device_id(self) -> int: - return self._device_id diff --git a/cuda_core/cuda/core/_resource_handles.pxd b/cuda_core/cuda/core/_resource_handles.pxd index 568af27ac2e..acf10e0fa2c 100644 --- a/cuda_core/cuda/core/_resource_handles.pxd +++ b/cuda_core/cuda/core/_resource_handles.pxd @@ -178,6 +178,12 @@ cdef GreenCtxHandle create_green_ctx_handle( cdef GreenCtxHandle create_green_ctx_handle_ref(cydriver.CUgreenCtx ctx) except+ nogil cdef ContextHandle get_primary_context(int device_id) except+ nogil cdef ContextHandle get_current_context() except+ nogil +cdef cydriver.CUresult context_synchronize( + const ContextHandle& h_context) noexcept nogil +cdef cydriver.CUresult context_get_stream_priority_range( + const ContextHandle& h_context, + int* least_priority, + int* greatest_priority) noexcept nogil # Stream handles cdef StreamHandle create_stream_handle( @@ -219,7 +225,8 @@ cdef MemoryPoolHandle create_mempool_handle_ipc( cdef DevicePtrHandle deviceptr_alloc_from_pool( size_t size, const MemoryPoolHandle& h_pool, const StreamHandle& h_stream) except+ nogil cdef DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) except+ nogil -cdef DevicePtrHandle deviceptr_alloc(size_t size) except+ nogil +cdef cydriver.CUresult deviceptr_alloc_raw( + cydriver.CUdeviceptr* ptr, size_t size, const ContextHandle& h_context) noexcept nogil cdef DevicePtrHandle deviceptr_alloc_host(size_t size) except+ nogil cdef DevicePtrHandle deviceptr_create_ref(cydriver.CUdeviceptr ptr) except+ nogil cdef DevicePtrHandle deviceptr_create_with_owner(cydriver.CUdeviceptr ptr, object owner) except+ nogil @@ -325,23 +332,29 @@ cdef FileDescriptorHandle create_fd_handle(int fd) except+ nogil cdef FileDescriptorHandle create_fd_handle_ref(int fd) except+ nogil # Array / mipmapped-array / texture / surface handles (PR #467) -cdef OpaqueArrayHandle create_array_handle(const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc) except+ nogil +cdef OpaqueArrayHandle create_array_handle( + const ContextHandle& h_context, const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc) except+ nogil cdef OpaqueArrayHandle create_array_handle_ref(cydriver.CUarray arr) except+ nogil cdef OpaqueArrayHandle create_array_handle_owning(cydriver.CUarray arr) except+ nogil +cdef ContextHandle get_array_context(const OpaqueArrayHandle& h) noexcept nogil cdef OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level) except+ nogil cdef MipmappedArrayHandle create_mipmapped_array_handle( - const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc, unsigned int num_levels) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc, + unsigned int num_levels) except+ nogil +cdef ContextHandle get_mipmapped_array_context( + const MipmappedArrayHandle& h) noexcept nogil cdef TexObjectHandle create_tex_object_handle_array( - const cydriver.CUDA_RESOURCE_DESC& res, const cydriver.CUDA_TEXTURE_DESC& tex, - const OpaqueArrayHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const cydriver.CUDA_TEXTURE_DESC& tex, const OpaqueArrayHandle& h_backing) except+ nogil cdef TexObjectHandle create_tex_object_handle_mipmap( - const cydriver.CUDA_RESOURCE_DESC& res, const cydriver.CUDA_TEXTURE_DESC& tex, - const MipmappedArrayHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const cydriver.CUDA_TEXTURE_DESC& tex, const MipmappedArrayHandle& h_backing) except+ nogil cdef TexObjectHandle create_tex_object_handle_linear( - const cydriver.CUDA_RESOURCE_DESC& res, const cydriver.CUDA_TEXTURE_DESC& tex, - const DevicePtrHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const cydriver.CUDA_TEXTURE_DESC& tex, const DevicePtrHandle& h_backing) except+ nogil cdef SurfObjectHandle create_surf_object_handle( - const cydriver.CUDA_RESOURCE_DESC& res, const OpaqueArrayHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const OpaqueArrayHandle& h_backing) except+ nogil # SM resource split (13.1+ — calls through function pointer, safe on older bindings) # groupParams is void* here to avoid referencing CU_DEV_SM_RESOURCE_GROUP_PARAMS diff --git a/cuda_core/cuda/core/_resource_handles.pyx b/cuda_core/cuda/core/_resource_handles.pyx index c7de24666f8..beecb4b745a 100644 --- a/cuda_core/cuda/core/_resource_handles.pyx +++ b/cuda_core/cuda/core/_resource_handles.pyx @@ -51,6 +51,12 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": ContextHandle get_primary_context "cuda_core::get_primary_context" ( int device_id) except+ nogil ContextHandle get_current_context "cuda_core::get_current_context" () except+ nogil + cydriver.CUresult context_synchronize "cuda_core::context_synchronize" ( + const ContextHandle& h_context) noexcept nogil + cydriver.CUresult context_get_stream_priority_range "cuda_core::context_get_stream_priority_range" ( + const ContextHandle& h_context, + int* least_priority, + int* greatest_priority) noexcept nogil # Stream handles StreamHandle create_stream_handle "cuda_core::create_stream_handle" ( @@ -107,7 +113,8 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": size_t size, const MemoryPoolHandle& h_pool, const StreamHandle& h_stream) except+ nogil DevicePtrHandle deviceptr_alloc_async "cuda_core::deviceptr_alloc_async" ( size_t size, const StreamHandle& h_stream) except+ nogil - DevicePtrHandle deviceptr_alloc "cuda_core::deviceptr_alloc" (size_t size) except+ nogil + cydriver.CUresult deviceptr_alloc_raw "cuda_core::deviceptr_alloc_raw" ( + cydriver.CUdeviceptr* ptr, size_t size, const ContextHandle& h_context) noexcept nogil DevicePtrHandle deviceptr_alloc_host "cuda_core::deviceptr_alloc_host" (size_t size) except+ nogil DevicePtrHandle deviceptr_create_ref "cuda_core::deviceptr_create_ref" ( cydriver.CUdeviceptr ptr) except+ nogil @@ -253,26 +260,32 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": # Array / mipmapped-array / texture / surface handles (PR #467) OpaqueArrayHandle create_array_handle "cuda_core::create_array_handle" ( - const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc) except+ nogil OpaqueArrayHandle create_array_handle_ref "cuda_core::create_array_handle_ref" ( cydriver.CUarray arr) except+ nogil OpaqueArrayHandle create_array_handle_owning "cuda_core::create_array_handle_owning" ( cydriver.CUarray arr) except+ nogil + ContextHandle get_array_context "cuda_core::get_array_context" ( + const OpaqueArrayHandle& h) noexcept nogil OpaqueArrayHandle create_array_level_handle "cuda_core::create_array_level_handle" ( const MipmappedArrayHandle& h_mip, unsigned int level) except+ nogil MipmappedArrayHandle create_mipmapped_array_handle "cuda_core::create_mipmapped_array_handle" ( - const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc, unsigned int num_levels) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_ARRAY3D_DESCRIPTOR& desc, + unsigned int num_levels) except+ nogil + ContextHandle get_mipmapped_array_context "cuda_core::get_mipmapped_array_context" ( + const MipmappedArrayHandle& h) noexcept nogil TexObjectHandle create_tex_object_handle_array "cuda_core::create_tex_object_handle_array" ( - const cydriver.CUDA_RESOURCE_DESC& res, const cydriver.CUDA_TEXTURE_DESC& tex, - const OpaqueArrayHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const cydriver.CUDA_TEXTURE_DESC& tex, const OpaqueArrayHandle& h_backing) except+ nogil TexObjectHandle create_tex_object_handle_mipmap "cuda_core::create_tex_object_handle_mipmap" ( - const cydriver.CUDA_RESOURCE_DESC& res, const cydriver.CUDA_TEXTURE_DESC& tex, - const MipmappedArrayHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const cydriver.CUDA_TEXTURE_DESC& tex, const MipmappedArrayHandle& h_backing) except+ nogil TexObjectHandle create_tex_object_handle_linear "cuda_core::create_tex_object_handle_linear" ( - const cydriver.CUDA_RESOURCE_DESC& res, const cydriver.CUDA_TEXTURE_DESC& tex, - const DevicePtrHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const cydriver.CUDA_TEXTURE_DESC& tex, const DevicePtrHandle& h_backing) except+ nogil SurfObjectHandle create_surf_object_handle "cuda_core::create_surf_object_handle" ( - const cydriver.CUDA_RESOURCE_DESC& res, const OpaqueArrayHandle& h_backing) except+ nogil + const ContextHandle& h_context, const cydriver.CUDA_RESOURCE_DESC& res, + const OpaqueArrayHandle& h_backing) except+ nogil # ============================================================================= @@ -297,11 +310,17 @@ cdef const char* _CUDA_DRIVER_API_V1_NAME = b"cuda.core._resource_handles._CUDA_ # Declare extern variables with reinterpret_cast to allow void* assignment cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": + # Error formatting + void* p_cuGetErrorName "reinterpret_cast(cuda_core::p_cuGetErrorName)" + void* p_cuGetErrorString "reinterpret_cast(cuda_core::p_cuGetErrorString)" + # Context void* p_cuDevicePrimaryCtxRetain "reinterpret_cast(cuda_core::p_cuDevicePrimaryCtxRetain)" void* p_cuDevicePrimaryCtxRelease "reinterpret_cast(cuda_core::p_cuDevicePrimaryCtxRelease)" void* p_cuCtxGetCurrent "reinterpret_cast(cuda_core::p_cuCtxGetCurrent)" 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_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)" @@ -408,8 +427,9 @@ cdef void* _get_optional_driver_fn(str name): cdef void _init_driver_fn_pointers() noexcept: + global p_cuGetErrorName, p_cuGetErrorString global p_cuDevicePrimaryCtxRetain, p_cuDevicePrimaryCtxRelease, p_cuCtxGetCurrent - global p_cuCtxSetCurrent + global p_cuCtxSetCurrent, p_cuCtxSynchronize, p_cuCtxGetStreamPriorityRange global p_cuGreenCtxCreate, p_cuGreenCtxDestroy, p_cuCtxFromGreenCtx global p_cuDevResourceGenerateDesc, p_cuGreenCtxStreamCreate global p_cuStreamCreateWithPriority, p_cuStreamDestroy @@ -435,11 +455,17 @@ cdef void _init_driver_fn_pointers() noexcept: global p_cuTexObjectCreate, p_cuTexObjectDestroy global p_cuSurfObjectCreate, p_cuSurfObjectDestroy + # Error formatting + p_cuGetErrorName = _get_driver_fn("cuGetErrorName") + p_cuGetErrorString = _get_driver_fn("cuGetErrorString") + # Context p_cuDevicePrimaryCtxRetain = _get_driver_fn("cuDevicePrimaryCtxRetain") p_cuDevicePrimaryCtxRelease = _get_driver_fn("cuDevicePrimaryCtxRelease") p_cuCtxGetCurrent = _get_driver_fn("cuCtxGetCurrent") p_cuCtxSetCurrent = _get_driver_fn("cuCtxSetCurrent") + p_cuCtxSynchronize = _get_driver_fn("cuCtxSynchronize") + p_cuCtxGetStreamPriorityRange = _get_driver_fn("cuCtxGetStreamPriorityRange") p_cuGreenCtxCreate = _get_optional_driver_fn("cuGreenCtxCreate") p_cuGreenCtxDestroy = _get_optional_driver_fn("cuGreenCtxDestroy") p_cuCtxFromGreenCtx = _get_optional_driver_fn("cuCtxFromGreenCtx") diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index 76aee36cdd3..9c51c488a6c 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -20,7 +20,10 @@ import warnings from dataclasses import dataclass from typing import Protocol, TYPE_CHECKING -from cuda.core._context cimport Context +from cuda.core._context cimport ( + Context, + Context_check_open, +) from cuda.core._device_resources cimport DeviceResources from cuda.core._event import Event, EventOptions @@ -32,6 +35,7 @@ from cuda.core._resource_handles cimport ( create_event_handle_noctx, create_stream_handle, create_stream_handle_with_owner, + context_get_stream_priority_range, get_current_context, get_last_error, get_legacy_stream, @@ -129,10 +133,7 @@ cdef class Stream: cdef StreamHandle h_stream cdef cydriver.CUstream borrowed cdef ContextHandle h_context - - # Extract context handle if provided - if ctx is not None: - h_context = (ctx)._h_context + cdef Context context if obj is not None and options is not None: raise ValueError("obj and options cannot be both specified") @@ -144,6 +145,12 @@ cdef class Stream: h_stream = create_stream_handle_with_owner(borrowed, obj) return Stream._from_handle(cls, h_stream) + if ctx is None: + raise RuntimeError("A CUDA context is required to create a stream") + context = ctx + Context_check_open(context) + h_context = context._h_context + cdef StreamOptions opts = check_or_create_options(StreamOptions, options, "Stream options") nonblocking = opts.nonblocking priority = opts.priority @@ -154,13 +161,9 @@ cdef class Stream: cdef int high, low cdef cydriver.CUresult res_code with nogil: - res_code = cydriver.cuCtxGetStreamPriorityRange(&high, &low) - if res_code != cydriver.CUresult.CUDA_SUCCESS: - if res_code == cydriver.CUresult.CUDA_ERROR_INVALID_CONTEXT: - raise RuntimeError( - "No current CUDA context. Call dev.set_current() before creating streams." - ) - HANDLE_RETURN(res_code) + res_code = context_get_stream_priority_range( + context._h_context, &high, &low) + HANDLE_RETURN(res_code) cdef int prio if priority is not None: prio = priority diff --git a/cuda_core/cuda/core/texture/_array.pyi b/cuda_core/cuda/core/texture/_array.pyi index 1b48b8afc82..f4258d5cbd5 100644 --- a/cuda_core/cuda/core/texture/_array.pyi +++ b/cuda_core/cuda/core/texture/_array.pyi @@ -3,6 +3,7 @@ from dataclasses import dataclass from cuda.bindings import cydriver +from cuda.core._context import Context from cuda.core.typing import ArrayFormatType _ARRAYFORMAT_TO_CU = {ArrayFormatType.UINT8: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT8), ArrayFormatType.UINT16: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT16), ArrayFormatType.UINT32: int(cydriver.CU_AD_FORMAT_UNSIGNED_INT32), ArrayFormatType.INT8: int(cydriver.CU_AD_FORMAT_SIGNED_INT8), ArrayFormatType.INT16: int(cydriver.CU_AD_FORMAT_SIGNED_INT16), ArrayFormatType.INT32: int(cydriver.CU_AD_FORMAT_SIGNED_INT32), ArrayFormatType.FLOAT16: int(cydriver.CU_AD_FORMAT_HALF), ArrayFormatType.FLOAT32: int(cydriver.CU_AD_FORMAT_FLOAT)} @@ -161,8 +162,8 @@ def _validate_format_channels(format, num_channels): def _validate_array_shape(shape): """Coerce ``shape`` to a tuple of ints and validate rank (1-3) and that every extent is >= 1. Returns the normalized tuple.""" -def _create_opaque_array(options): - """Allocate a new :class:`OpaqueArray` on the current device. +def _create_opaque_array(options, ctx: Context, device_id: int): + """Allocate a new :class:`OpaqueArray` on the specified device. Backs :meth:`cuda.core.Device.create_opaque_array`. ``options`` is an :class:`OpaqueArrayOptions` (or a mapping accepted by it); it is validated diff --git a/cuda_core/cuda/core/texture/_array.pyx b/cuda_core/cuda/core/texture/_array.pyx index fbfa908c526..e5fc3f6c9e2 100644 --- a/cuda_core/cuda/core/texture/_array.pyx +++ b/cuda_core/cuda/core/texture/_array.pyx @@ -9,6 +9,7 @@ from libc.stdint cimport intptr_t from libc.string cimport memset from cuda.bindings cimport cydriver +from cuda.core._context cimport Context from cuda.core._memory._buffer cimport Buffer, Buffer_check_open from cuda.core._resource_handles cimport ( OpaqueArrayHandle, @@ -515,8 +516,8 @@ cdef OpaqueArray _array_from_handle(OpaqueArrayHandle h, int device_id): return self -def _create_opaque_array(options): - """Allocate a new :class:`OpaqueArray` on the current device. +def _create_opaque_array(options, Context ctx, int device_id): + """Allocate a new :class:`OpaqueArray` on the specified device. Backs :meth:`cuda.core.Device.create_opaque_array`. ``options`` is an :class:`OpaqueArrayOptions` (or a mapping accepted by it); it is validated @@ -545,7 +546,7 @@ def _create_opaque_array(options): Flags=flags, ) - cdef OpaqueArrayHandle h = create_array_handle(desc3d) + cdef OpaqueArrayHandle h = create_array_handle(ctx._h_context, desc3d) if not h: HANDLE_RETURN(get_last_error()) @@ -555,5 +556,5 @@ def _create_opaque_array(options): self._format = c_format self._num_channels = opts.num_channels self._surface_load_store = bool(opts.is_surface_load_store) - self._device_id = _get_current_device_id() + self._device_id = device_id return self diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyi b/cuda_core/cuda/core/texture/_mipmapped_array.pyi index e4ec707b458..3d843abe70a 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyi +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyi @@ -2,6 +2,8 @@ from dataclasses import dataclass +from cuda.core._context import Context + @dataclass class MipmappedArrayOptions: @@ -105,8 +107,8 @@ class MipmappedArray: def __exit__(self, exc_type, exc, tb): ... def __repr__(self): ... -def _create_mipmapped_array(options): - """Allocate a new :class:`MipmappedArray` on the current device. +def _create_mipmapped_array(options, ctx: Context, device_id: int): + """Allocate a new :class:`MipmappedArray` on the specified device. Backs :meth:`cuda.core.Device.create_mipmapped_array`. ``options`` is a :class:`MipmappedArrayOptions` (or a mapping accepted by it); its fields are diff --git a/cuda_core/cuda/core/texture/_mipmapped_array.pyx b/cuda_core/cuda/core/texture/_mipmapped_array.pyx index abc30c6b25c..8d6bf5a2589 100644 --- a/cuda_core/cuda/core/texture/_mipmapped_array.pyx +++ b/cuda_core/cuda/core/texture/_mipmapped_array.pyx @@ -5,6 +5,7 @@ from __future__ import annotations from cuda.bindings cimport cydriver +from cuda.core._context cimport Context from cuda.core.texture._array cimport _array_from_handle from cuda.core.texture._array import ( _ARRAYFORMAT_TO_CU, @@ -20,10 +21,7 @@ from cuda.core._resource_handles cimport ( create_mipmapped_array_handle, get_last_error, ) -from cuda.core._utils.cuda_utils cimport ( - HANDLE_RETURN, - _get_current_device_id, -) +from cuda.core._utils.cuda_utils cimport HANDLE_RETURN from dataclasses import dataclass @@ -191,8 +189,8 @@ cdef class MipmappedArray: f"num_levels={self._num_levels})" ) -def _create_mipmapped_array(options): - """Allocate a new :class:`MipmappedArray` on the current device. +def _create_mipmapped_array(options, Context ctx, int device_id): + """Allocate a new :class:`MipmappedArray` on the specified device. Backs :meth:`cuda.core.Device.create_mipmapped_array`. ``options`` is a :class:`MipmappedArrayOptions` (or a mapping accepted by it); its fields are @@ -221,7 +219,8 @@ def _create_mipmapped_array(options): Flags=flags, ) - cdef MipmappedArrayHandle h = create_mipmapped_array_handle(desc3d, c_levels) + cdef MipmappedArrayHandle h = create_mipmapped_array_handle( + ctx._h_context, desc3d, c_levels) if not h: HANDLE_RETURN(get_last_error()) @@ -232,5 +231,5 @@ def _create_mipmapped_array(options): self._num_channels = opts.num_channels self._num_levels = opts.num_levels self._surface_load_store = bool(opts.is_surface_load_store) - self._device_id = _get_current_device_id() + self._device_id = device_id return self diff --git a/cuda_core/cuda/core/texture/_surface.pyi b/cuda_core/cuda/core/texture/_surface.pyi index b153ff31fed..c363b63b05e 100644 --- a/cuda_core/cuda/core/texture/_surface.pyi +++ b/cuda_core/cuda/core/texture/_surface.pyi @@ -1,5 +1,8 @@ # This file was generated by stubgen-pyx v0.2.19 from cuda_core/cuda/core/texture/_surface.pyx +from cuda.core._context import Context + + class SurfaceObject: """A bindless surface handle for kernel-side typed load/store. @@ -39,8 +42,8 @@ class SurfaceObject: def __exit__(self, exc_type, exc, tb): ... def __repr__(self): ... -def _create_surface_object(resource): - """Create a :class:`SurfaceObject` on the current device. +def _create_surface_object(resource, ctx: Context, device_id: int): + """Create a :class:`SurfaceObject` on the specified device. Backs :meth:`cuda.core.Device.create_surface_object`. ``resource`` must be a :class:`ResourceDescriptor` wrapping an :class:`OpaqueArray` allocated with diff --git a/cuda_core/cuda/core/texture/_surface.pyx b/cuda_core/cuda/core/texture/_surface.pyx index 074f438ad47..790ce048ecd 100644 --- a/cuda_core/cuda/core/texture/_surface.pyx +++ b/cuda_core/cuda/core/texture/_surface.pyx @@ -7,19 +7,19 @@ from __future__ import annotations from libc.string cimport memset from cuda.bindings cimport cydriver +from cuda.core._context cimport Context from cuda.core.texture._array cimport OpaqueArray, OpaqueArray_check_open from cuda.core._resource_handles cimport ( + ContextHandle, SurfObjectHandle, as_cu, as_intptr, create_surf_object_handle, + get_array_context, get_last_error, ) from cuda.core.texture._texture import ResourceDescriptor -from cuda.core._utils.cuda_utils cimport ( - HANDLE_RETURN, - _get_current_device_id, -) +from cuda.core._utils.cuda_utils cimport HANDLE_RETURN cdef class SurfaceObject: @@ -85,8 +85,8 @@ cdef class SurfaceObject: return f"SurfaceObject(handle=0x{as_intptr(self._handle):x})" -def _create_surface_object(resource): - """Create a :class:`SurfaceObject` on the current device. +def _create_surface_object(resource, Context ctx, int device_id): + """Create a :class:`SurfaceObject` on the specified device. Backs :meth:`cuda.core.Device.create_surface_object`. ``resource`` must be a :class:`ResourceDescriptor` wrapping an :class:`OpaqueArray` allocated with @@ -106,6 +106,14 @@ def _create_surface_object(resource): cdef OpaqueArray arr = resource.source OpaqueArray_check_open(arr) + if arr._device_id != device_id: + raise ValueError( + f"resource belongs to device {arr._device_id}, " + f"but surface creation was requested on device {device_id}" + ) + cdef ContextHandle resource_context = get_array_context(arr._handle) + if resource_context and as_cu(resource_context) != as_cu(ctx._h_context): + raise ValueError("resource is not compatible with this Device object") if not arr.is_surface_load_store: raise ValueError( "OpaqueArray must be created with is_surface_load_store=True to be " @@ -117,12 +125,13 @@ def _create_surface_object(resource): res_desc.resType = cydriver.CU_RESOURCE_TYPE_ARRAY res_desc.res.array.hArray = as_cu(arr._handle) - cdef SurfObjectHandle h = create_surf_object_handle(res_desc, arr._handle) + cdef SurfObjectHandle h = create_surf_object_handle( + ctx._h_context, res_desc, arr._handle) if not h: HANDLE_RETURN(get_last_error()) cdef SurfaceObject self = SurfaceObject.__new__(SurfaceObject) self._handle = h self._source_ref = resource - self._device_id = _get_current_device_id() + self._device_id = device_id return self diff --git a/cuda_core/cuda/core/texture/_texture.pyi b/cuda_core/cuda/core/texture/_texture.pyi index 16508003091..f210be8ac89 100644 --- a/cuda_core/cuda/core/texture/_texture.pyi +++ b/cuda_core/cuda/core/texture/_texture.pyi @@ -3,6 +3,7 @@ from dataclasses import dataclass from cuda.bindings import cydriver +from cuda.core._context import Context from cuda.core.typing import AddressModeType, FilterModeType, ReadModeType _TRSF_READ_AS_INTEGER = 1 @@ -215,8 +216,8 @@ def _normalize_enum(name, value, enum_type): def _normalize_address_modes(address_mode): """Return a 3-tuple of :class:`AddressModeType` values from a scalar or 1-3 tuple. Individual entries may be plain strings.""" -def _create_texture_object(resource, options): - """Create a :class:`TextureObject` on the current device. +def _create_texture_object(resource, options, ctx: Context, device_id: int): + """Create a :class:`TextureObject` on the specified device. Backs :meth:`cuda.core.Device.create_texture_object`. ``resource`` is a :class:`ResourceDescriptor`; ``options`` is a :class:`TextureObjectOptions` diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index ef63c01972a..f1c29490b8a 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -8,6 +8,7 @@ from libc.stdint cimport intptr_t from libc.string cimport memset from cuda.bindings cimport cydriver +from cuda.core._context cimport Context from cuda.core.texture._array cimport OpaqueArray, OpaqueArray_check_open from cuda.core.texture._array import ( _ARRAYFORMAT_TO_CU, @@ -19,18 +20,18 @@ from cuda.core._memory._buffer cimport Buffer, Buffer_check_open from cuda.core.texture._mipmapped_array cimport MipmappedArray, MipmappedArray_check_open from cuda.core.texture._mipmapped_array import MipmappedArray as _PyMipmappedArray from cuda.core._resource_handles cimport ( + ContextHandle, TexObjectHandle, as_cu, as_intptr, create_tex_object_handle_array, create_tex_object_handle_linear, create_tex_object_handle_mipmap, + get_array_context, get_last_error, + get_mipmapped_array_context, ) -from cuda.core._utils.cuda_utils cimport ( - HANDLE_RETURN, - _get_current_device_id, -) +from cuda.core._utils.cuda_utils cimport HANDLE_RETURN from cuda.core.typing import AddressModeType, FilterModeType, ReadModeType @@ -474,8 +475,9 @@ cdef class TextureObject: return f"TextureObject(handle=0x{as_intptr(self._handle):x})" -def _create_texture_object(resource, options): - """Create a :class:`TextureObject` on the current device. +def _create_texture_object( + resource, options, Context ctx, int device_id): + """Create a :class:`TextureObject` on the specified device. Backs :meth:`cuda.core.Device.create_texture_object`. ``resource`` is a :class:`ResourceDescriptor`; ``options`` is a :class:`TextureObjectOptions` @@ -500,19 +502,26 @@ def _create_texture_object(resource, options): cdef MipmappedArray mip cdef Buffer buf cdef intptr_t devptr + cdef ContextHandle resource_context + cdef int resource_device_id if resource.kind == "array": arr = resource.source OpaqueArray_check_open(arr) + resource_context = get_array_context(arr._handle) + resource_device_id = arr._device_id res_desc.resType = cydriver.CU_RESOURCE_TYPE_ARRAY res_desc.res.array.hArray = as_cu(arr._handle) elif resource.kind == "mipmapped_array": mip = resource.source MipmappedArray_check_open(mip) + resource_context = get_mipmapped_array_context(mip._handle) + resource_device_id = mip._device_id res_desc.resType = cydriver.CU_RESOURCE_TYPE_MIPMAPPED_ARRAY res_desc.res.mipmap.hMipmappedArray = as_cu(mip._handle) elif resource.kind == "linear": buf = resource.source Buffer_check_open(buf) + resource_device_id = buf.device_id devptr = int(buf.handle) res_desc.resType = cydriver.CU_RESOURCE_TYPE_LINEAR res_desc.res.linear.devPtr = devptr @@ -522,6 +531,7 @@ def _create_texture_object(resource, options): elif resource.kind == "pitch2d": buf = resource.source Buffer_check_open(buf) + resource_device_id = buf.device_id devptr = int(buf.handle) res_desc.resType = cydriver.CU_RESOURCE_TYPE_PITCH2D res_desc.res.pitch2D.devPtr = devptr @@ -534,6 +544,13 @@ def _create_texture_object(resource, options): raise NotImplementedError( f"ResourceDescriptor kind {resource.kind!r} is not yet supported" ) + if resource_device_id >= 0 and resource_device_id != device_id: + raise ValueError( + f"resource belongs to device {resource_device_id}, " + f"but texture creation was requested on device {device_id}" + ) + if resource_context and as_cu(resource_context) != as_cu(ctx._h_context): + raise ValueError("resource is not compatible with this Device object") # --- Texture descriptor --- # filter_mode/read_mode/mipmap_filter_mode are normalized to their @@ -585,11 +602,14 @@ def _create_texture_object(resource, options): cdef TexObjectHandle h if resource.kind == "array": - h = create_tex_object_handle_array(res_desc, tex_desc, arr._handle) + h = create_tex_object_handle_array( + ctx._h_context, res_desc, tex_desc, arr._handle) elif resource.kind == "mipmapped_array": - h = create_tex_object_handle_mipmap(res_desc, tex_desc, mip._handle) + h = create_tex_object_handle_mipmap( + ctx._h_context, res_desc, tex_desc, mip._handle) else: # linear or pitch2d — both backed by a device Buffer - h = create_tex_object_handle_linear(res_desc, tex_desc, buf._h_ptr) + h = create_tex_object_handle_linear( + ctx._h_context, res_desc, tex_desc, buf._h_ptr) if not h: HANDLE_RETURN(get_last_error()) @@ -597,5 +617,5 @@ def _create_texture_object(resource, options): self._handle = h self._source_ref = resource self._options = opts - self._device_id = _get_current_device_id() + self._device_id = device_id return self diff --git a/cuda_core/docs/source/interoperability.rst b/cuda_core/docs/source/interoperability.rst index 87347eb9d25..33d11d540c7 100644 --- a/cuda_core/docs/source/interoperability.rst +++ b/cuda_core/docs/source/interoperability.rst @@ -26,6 +26,10 @@ Conversely, if any GPU library already sets a device (or context) to current, th method ensures that the same device/context is picked up by and shared with ``cuda.core``. +Other :class:`Device` methods do not change the current context. For example, +``dev1.sync()`` synchronizes device 1 and leaves the current context unchanged, +even when another device is current. + ``__cuda_stream__`` protocol ---------------------------- diff --git a/cuda_core/docs/source/release/1.2.0-notes.rst b/cuda_core/docs/source/release/1.2.0-notes.rst index f96a205d1e8..9591cc3a083 100644 --- a/cuda_core/docs/source/release/1.2.0-notes.rst +++ b/cuda_core/docs/source/release/1.2.0-notes.rst @@ -40,6 +40,13 @@ New features Fixes and enhancements ---------------------- +- :class:`Device` methods that create resources or synchronize now act on that + device, even when another device is current. They do not change which device + is current. For example, ``dev1.sync()`` synchronizes device 1 even when + device 0 is current. :meth:`Device.set_current` now always returns a + :class:`Context` with the correct device ID. + (`#2311 `__) + - A :class:`Buffer` is now freed correctly even when the CUDA context current at teardown is not the one it was allocated in, or when no context is current at all. This happens routinely when a buffer is released by the garbage diff --git a/cuda_core/tests/conftest.py b/cuda_core/tests/conftest.py index ff4cddcb28f..0466277bf6c 100644 --- a/cuda_core/tests/conftest.py +++ b/cuda_core/tests/conftest.py @@ -212,6 +212,15 @@ def deinit_cuda(): _ = _device_unset_current() +@pytest.fixture +def device_x2(deinit_cuda): + """Provide two CUDA devices, or skip when fewer are available.""" + devices = Device.get_all_devices() + if len(devices) < 2: + pytest.skip("Test requires at least 2 CUDA devices") + return devices[:2] + + @pytest.fixture def deinit_all_contexts_function(): def pop_all_contexts(): diff --git a/cuda_core/tests/helpers/contexts.py b/cuda_core/tests/helpers/contexts.py new file mode 100644 index 00000000000..9818875c4da --- /dev/null +++ b/cuda_core/tests/helpers/contexts.py @@ -0,0 +1,90 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +from contextlib import contextmanager + +from cuda.core._utils.cuda_utils import driver, handle_return + +__all__ = [ + "assert_device_operations_use_bound_context", + "current_context_handle", + "no_current_context", + "use_context", +] + + +def current_context_handle(): + """Return the current CUDA context handle, or zero if none is current.""" + return int(handle_return(driver.cuCtxGetCurrent())) + + +def assert_device_operations_use_bound_context(device): + """Check that Device operations use its bound context and preserve the ambient context.""" + bound_context = device.context + ambient_context_handle = current_context_handle() + stream = event = builder = None + + try: + stream = device.create_stream() + assert stream.context == bound_context + assert current_context_handle() == ambient_context_handle + + event = device.create_event() + assert event.context == bound_context + assert current_context_handle() == ambient_context_handle + + builder = device.create_graph_builder() + assert builder.stream.context == bound_context + assert current_context_handle() == ambient_context_handle + + device.sync() + assert current_context_handle() == ambient_context_handle + + builder.close() + builder = None + assert current_context_handle() == ambient_context_handle + + event.close() + event = None + assert current_context_handle() == ambient_context_handle + + stream.close() + stream = None + assert current_context_handle() == ambient_context_handle + finally: + if builder is not None: + builder.close() + if event is not None: + event.close() + if stream is not None: + stream.close() + + +@contextmanager +def no_current_context(): + """Temporarily remove the calling thread's sole current CUDA context.""" + if current_context_handle() == 0: + raise RuntimeError("no_current_context requires a current CUDA context") + + previous = handle_return(driver.cuCtxPopCurrent()) + try: + if current_context_handle() != 0: + raise RuntimeError("no_current_context requires exactly one stacked CUDA context") + yield + finally: + handle_return(driver.cuCtxPushCurrent(previous)) + + +@contextmanager +def use_context(device, context): + """Temporarily make a context current and restore the previous context.""" + if current_context_handle() == 0: + raise RuntimeError("use_context requires a current CUDA context to restore") + + previous = device.set_current(context) + if previous is None: + raise RuntimeError("Device.set_current() did not return the previous CUDA context") + try: + yield + finally: + device.set_current(previous) diff --git a/cuda_core/tests/memory_ipc/test_peer_access.py b/cuda_core/tests/memory_ipc/test_peer_access.py index 992e01aa540..a82690d46b5 100644 --- a/cuda_core/tests/memory_ipc/test_peer_access.py +++ b/cuda_core/tests/memory_ipc/test_peer_access.py @@ -96,7 +96,11 @@ def test_main(self, ipc_mempool_device_x2, grant_access_in_parent): buffer.close() # TODO(seberg): 2026-06: mr close may be unsafe with incomplete `buf.close()` + # Make dev0 current; Device.sync() must act on dev1 and leave dev0 current. + dev0.set_current() + assert Device().device_id == dev0.device_id dev1.sync() + assert Device().device_id == dev0.device_id mr.close() def child_main(self, mr, buffer): diff --git a/cuda_core/tests/test_device.py b/cuda_core/tests/test_device.py index 0d2e5e00952..48f5b3cf484 100644 --- a/cuda_core/tests/test_device.py +++ b/cuda_core/tests/test_device.py @@ -2,12 +2,18 @@ # SPDX-License-Identifier: Apache-2.0 import contextlib +from concurrent.futures import ThreadPoolExecutor import pytest +from helpers.contexts import ( + assert_device_operations_use_bound_context, + current_context_handle, + no_current_context, +) import cuda.core from cuda.bindings import driver, runtime -from cuda.core import Device +from cuda.core import Device, StreamOptions from cuda.core._utils.cuda_utils import ComputeCapability, handle_return from cuda.core._utils.version import driver_version @@ -100,6 +106,80 @@ def test_device_create_event(init_cuda): assert event.handle +@pytest.mark.agent_authored(model="gpt-5.6") +def test_device_operations_target_receiver_and_restore_current(device_x2): + dev0, dev1 = device_x2 + dev0.set_current() + dev1.set_current() + assert_device_operations_use_bound_context(dev0) + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_device_operations_restore_no_current_context(deinit_cuda): + device = Device(0) + device.set_current() + + with no_current_context(): + assert_device_operations_use_bound_context(device) + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_device_create_stream_restores_context_after_failure(device_x2): + dev0, dev1 = device_x2 + dev0.set_current() + dev1.set_current() + ctx1_handle = current_context_handle() + + with pytest.raises(ValueError, match="priority=.*out of range"): + dev0.create_stream(options=StreamOptions(priority=2**30)) + assert current_context_handle() == ctx1_handle + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_set_current_returns_previous_context_with_owning_device(device_x2): + dev0, dev1 = device_x2 + dev0.set_current() + ctx0 = dev0.context + dev1.set_current() + + previous = dev0.set_current(ctx0) + assert previous.handle == dev1.context.handle + dev1.set_current(previous) + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_device_receiver_switching_is_thread_local(device_x2): + dev0, dev1 = device_x2 + dev0.set_current() + main_context = current_context_handle() + + def worker(): + worker_dev0 = Device(dev0.device_id) + worker_dev0.set_current() + target_context = current_context_handle() + worker_dev1 = Device(dev1.device_id) + worker_dev1.set_current() + foreign_context = current_context_handle() + + stream = None + try: + stream = worker_dev0.create_stream() + resource_context = int(stream.context.handle) + finally: + if stream is not None: + stream.close() + + return resource_context, target_context, current_context_handle(), foreign_context + + with ThreadPoolExecutor(max_workers=1) as executor: + worker_result = executor.submit(worker).result() + + resource_context, target_context, restored_context, foreign_context = worker_result + assert resource_context == target_context + assert restored_context == foreign_context + assert current_context_handle() == main_context + + def test_pci_bus_id(): device = Device() bus_id = handle_return(runtime.cudaDeviceGetPCIBusId(13, device.device_id)) diff --git a/cuda_core/tests/test_green_context.py b/cuda_core/tests/test_green_context.py index 52fc372a27b..5b88083d22b 100644 --- a/cuda_core/tests/test_green_context.py +++ b/cuda_core/tests/test_green_context.py @@ -2,11 +2,9 @@ # # SPDX-License-Identifier: Apache-2.0 - -import contextlib - import numpy as np import pytest +from helpers.contexts import assert_device_operations_use_bound_context, use_context from cuda.core import ( ContextOptions, @@ -152,16 +150,6 @@ def _find_backfill_only_two_group_split(sm): return None -@contextlib.contextmanager -def _use_green_ctx(dev, ctx): - """Context manager: set green ctx current, restore previous on exit.""" - prev = dev.set_current(ctx) - try: - yield - finally: - dev.set_current(prev) - - @pytest.mark.agent_authored(model="gpt-5.6") def test_memory_node_updates_preserve_green_context( init_cuda, @@ -173,7 +161,7 @@ def test_memory_node_updates_preserve_green_context( memory_resource = LegacyPinnedMemoryResource() src = memory_resource.allocate(4) dst = memory_resource.allocate(4) - with _use_green_ctx(init_cuda, green_ctx): + with use_context(init_cuda, green_ctx): graph_def = GraphDefinition() memset_node = graph_def.memset(dst, 0, 4) memcpy_node = graph_def.memcpy(dst, src, 4) @@ -528,16 +516,48 @@ def test_stream_and_event_track_green_context(self, green_ctx): stream.sync() event.sync() + @pytest.mark.agent_authored(model="gpt-5.6") + def test_device_receiver_targets_stored_green_context(self, init_cuda, green_ctx): + primary_ctx = init_cuda.context + + with use_context(init_cuda, green_ctx): + handle_return(driver.cuCtxSetCurrent(primary_ctx.handle)) + assert_device_operations_use_bound_context(init_cuda) + + @pytest.mark.agent_authored(model="gpt-5.6") + def test_texture_rejects_resource_from_other_context(self, init_cuda, green_ctx): + from cuda.core.texture import ( + OpaqueArrayOptions, + ResourceDescriptor, + ) + from cuda.core.typing import ArrayFormatType + + array = init_cuda.create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + ) + ) + try: + with ( + use_context(init_cuda, green_ctx), + pytest.raises(ValueError, match="resource is not compatible with this Device object"), + ): + init_cuda.create_texture_object(resource=ResourceDescriptor.from_opaque_array(array)) + finally: + array.close() + def test_close_while_current_raises(self, init_cuda, green_ctx): """close() on a current context raises — test via set_current.""" dev = init_cuda - with _use_green_ctx(dev, green_ctx), pytest.raises(RuntimeError, match="while it is current"): + with use_context(dev, green_ctx), pytest.raises(RuntimeError, match="while it is current"): green_ctx.close() def test_set_current_swap_regression(self, init_cuda, green_ctx): """set_current still works (backward compat) and preserves identity.""" dev = init_cuda - with _use_green_ctx(dev, green_ctx): + with use_context(dev, green_ctx): pass # just verify push/pop works # Swap again and check identity round-trip prev = dev.set_current(green_ctx) diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 2ab766cc2f0..2ee783016f2 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -24,7 +24,7 @@ StreamOptions, launch, ) -from cuda.core._memory._legacy import _SynchronousMemoryResource +from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource from cuda.core._utils.cuda_utils import CUDAError from cuda.core.typing import ObjectCodeFormatType, SourceCodeType diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index 44227d4b1c5..b6b219e002e 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -23,6 +23,7 @@ thread_unsafe_on_windows, ) from helpers.constants import POOL_SIZE +from helpers.contexts import current_context_handle, no_current_context from helpers.memory import ( create_managed_memory_resource_or_skip, create_pinned_memory_resource_or_xfail, @@ -739,14 +740,10 @@ def test_close_with_default_stream_requires_context(): # Use a real stream at creation so _init succeeds without a current context later. buf = Buffer.from_handle(1, 1024, mr=mr, stream=stream) - previous = handle_return(driver.cuCtxPopCurrent()) - assert int(previous) != 0 - try: - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + with no_current_context(): + assert current_context_handle() == 0 with pytest.raises(RuntimeError, match="no CUDA context is current"): buf.close(stream=default_stream()) - finally: - handle_return(driver.cuCtxSetCurrent(previous)) buf.close() # clean up using the recorded stream (which carries a context) @@ -758,14 +755,10 @@ def test_from_handle_mr_default_stream_requires_context(buffer_type): device = Device() device.set_current() mr = StubMemoryResource(device) - previous = handle_return(driver.cuCtxPopCurrent()) - assert int(previous) != 0 - try: - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + with no_current_context(): + assert current_context_handle() == 0 with pytest.raises(RuntimeError, match="no CUDA context is current"): buffer_type.from_handle(1, 1024, mr=mr) - finally: - handle_return(driver.cuCtxSetCurrent(previous)) @pytest.mark.agent_authored(model="gpt-5.6") @@ -777,15 +770,11 @@ def test_from_handle_mr_explicit_stream_without_current_context(buffer_type): stream = device.create_stream() CapturingMR, telemetry = make_instrumented_memory_resource(record_streams=True) mr = CapturingMR(device) - previous = handle_return(driver.cuCtxPopCurrent()) - assert int(previous) != 0 - try: - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + with no_current_context(): + assert current_context_handle() == 0 buf = buffer_type.from_handle(1, 1024, mr=mr, stream=stream) buf.close() - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 - finally: - handle_return(driver.cuCtxSetCurrent(previous)) + assert current_context_handle() == 0 assert telemetry["deallocations"][-1]["stream"].handle == stream.handle @@ -814,39 +803,31 @@ def test_mr_deallocation_without_current_context(init_cuda, capsys, replace_stre stream = init_cuda.create_stream() if replace_stream else None assert len(telemetry["active"]) == 1 - previous = handle_return(driver.cuCtxPopCurrent()) - assert int(previous) != 0 - try: - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + with no_current_context(): + assert current_context_handle() == 0 buf.close(stream) assert len(telemetry["active"]) == 0 - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + assert current_context_handle() == 0 assert "mr.deallocate() failed" not in capsys.readouterr().err - finally: - handle_return(driver.cuCtxSetCurrent(previous)) @pytest.mark.agent_authored(model="cursor-grok-4.5") @pytest.mark.parametrize("replace_stream", [False, True]) -def test_mr_deallocation_with_foreign_context(capsys, replace_stream): +def test_mr_deallocation_with_foreign_context(device_x2, capsys, replace_stream): """MR-backed Buffer teardown switches away from an unrelated current context.""" - if len(Device.get_all_devices()) < 2: - pytest.skip("Test requires at least 2 GPUs") - - alloc_dev = Device(0) + alloc_dev, foreign_dev = device_x2 alloc_dev.set_current() TrackingMR, telemetry = make_instrumented_memory_resource(DummyDeviceMemoryResource, track_active=True) mr = TrackingMR(alloc_dev) buf = mr.allocate(1024) stream = alloc_dev.create_stream() if replace_stream else None assert len(telemetry["active"]) == 1 - alloc_ctx = int(handle_return(driver.cuCtxGetCurrent())) + alloc_ctx = current_context_handle() - foreign_dev = Device(1) foreign_dev.set_current() - foreign_ctx = int(handle_return(driver.cuCtxGetCurrent())) + foreign_ctx = current_context_handle() assert foreign_ctx != 0 assert foreign_ctx != alloc_ctx @@ -854,7 +835,7 @@ def test_mr_deallocation_with_foreign_context(capsys, replace_stream): buf.close(stream) assert len(telemetry["active"]) == 0 - assert int(handle_return(driver.cuCtxGetCurrent())) == foreign_ctx + assert current_context_handle() == foreign_ctx assert "mr.deallocate() failed" not in capsys.readouterr().err finally: alloc_dev.set_current() @@ -886,21 +867,17 @@ def test_pool_buffer_deallocates_without_current_context(mempool_device, capfd): stream.sync() used_after_alloc = mr.attributes.used_mem_current - previous = handle_return(driver.cuCtxPopCurrent()) - assert int(previous) != 0 - try: - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + with no_current_context(): + assert current_context_handle() == 0 buf.close() stream.sync() assert mr.attributes.used_mem_current < used_after_alloc - assert int(handle_return(driver.cuCtxGetCurrent())) == 0 + assert current_context_handle() == 0 err = capfd.readouterr().err - assert "failed during resource destruction" not in err + assert "cuMemFreeAsync failed" not in err assert "mr.deallocate() failed" not in err - finally: - handle_return(driver.cuCtxSetCurrent(previous)) @pytest.mark.agent_authored(model="cursor-grok-4.5") @@ -914,16 +891,16 @@ def test_pool_buffer_deallocates_with_foreign_context(mempool_device_x2, capfd): buf = mr.allocate(size, stream=stream) stream.sync() used_after_alloc = mr.attributes.used_mem_current - alloc_ctx = int(handle_return(driver.cuCtxGetCurrent())) + alloc_ctx = current_context_handle() foreign_dev.set_current() - foreign_ctx = int(handle_return(driver.cuCtxGetCurrent())) + foreign_ctx = current_context_handle() assert foreign_ctx != 0 assert foreign_ctx != alloc_ctx try: buf.close() - assert int(handle_return(driver.cuCtxGetCurrent())) == foreign_ctx + assert current_context_handle() == foreign_ctx # Observe the free on the allocation device, then restore the foreign context. alloc_dev.set_current() @@ -932,7 +909,7 @@ def test_pool_buffer_deallocates_with_foreign_context(mempool_device_x2, capfd): foreign_dev.set_current() err = capfd.readouterr().err - assert "failed during resource destruction" not in err + assert "cuMemFreeAsync failed" not in err finally: alloc_dev.set_current() @@ -2189,7 +2166,7 @@ def test_legacy_pinned_device_id_raises(): def test_synchronous_memory_resource_basic(init_cuda): """_SynchronousMemoryResource exercises properties and allocate paths (zero, non-zero, with-stream).""" - from cuda.core._memory._legacy import _SynchronousMemoryResource + from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource dev = Device() mr = _SynchronousMemoryResource(dev.device_id) @@ -2222,7 +2199,7 @@ def test_synchronous_memory_resource_basic(init_cuda): def test_synchronous_memory_resource_deallocate_accepts_stream(init_cuda): """_SynchronousMemoryResource.deallocate accepts an explicit stream.""" - from cuda.core._memory._legacy import _SynchronousMemoryResource + from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource dev = Device() mr = _SynchronousMemoryResource(dev.device_id) @@ -2232,6 +2209,60 @@ def test_synchronous_memory_resource_deallocate_accepts_stream(init_cuda): stream.close() +@pytest.mark.agent_authored(model="gpt-5.6") +def test_synchronous_memory_resource_uses_its_context(device_x2): + """Synchronous allocation targets its stored context and restores the current one.""" + from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource + + alloc_dev, current_dev = device_x2 + alloc_dev.set_current() + stream = alloc_dev.create_stream() + mr = _SynchronousMemoryResource(alloc_dev.device_id, alloc_dev.context) + + current_dev.set_current() + current_context = current_context_handle() + + buf = mr.allocate(64, stream=stream) + try: + pointer_context = handle_return( + driver.cuPointerGetAttribute( + driver.CUpointer_attribute.CU_POINTER_ATTRIBUTE_CONTEXT, + int(buf.handle), + ) + ) + pointer_device = handle_return( + driver.cuPointerGetAttribute( + driver.CUpointer_attribute.CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL, + int(buf.handle), + ) + ) + assert int(pointer_context) == int(alloc_dev.context.handle) + assert pointer_device == alloc_dev.device_id + assert current_context_handle() == current_context + finally: + buf.close(stream=stream) + stream.close() + + assert current_context_handle() == current_context + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_synchronous_memory_resource_restores_context_after_failure(device_x2): + """A failed synchronous allocation restores the context that was current.""" + from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource + + alloc_dev, current_dev = device_x2 + alloc_dev.set_current() + mr = _SynchronousMemoryResource(alloc_dev.device_id, alloc_dev.context) + current_dev.set_current() + current_context = current_context_handle() + + with pytest.raises(CUDAError): + mr.allocate(sys.maxsize) + + assert current_context_handle() == current_context + + @pytest.mark.parametrize( ("method", "spec", "match"), [ diff --git a/cuda_core/tests/test_stream.py b/cuda_core/tests/test_stream.py index 6309c134ccc..bd909734e09 100644 --- a/cuda_core/tests/test_stream.py +++ b/cuda_core/tests/test_stream.py @@ -190,7 +190,7 @@ class MyStream(Stream): dev = Device() dev.set_current() - stream = MyStream._init(options=StreamOptions(), device_id=dev.device_id) + stream = MyStream._init(options=StreamOptions(), device_id=dev.device_id, ctx=dev.context) assert isinstance(stream, MyStream) diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 7436b99647f..254eb2faf4d 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -5,6 +5,7 @@ import numpy as np import pytest +from helpers.contexts import current_context_handle, no_current_context import cuda.core from cuda.core import ( @@ -45,6 +46,110 @@ def test_resource_descriptor_init_disabled(): ResourceDescriptor() +@pytest.mark.agent_authored(model="gpt-5.6") +def test_texture_resources_target_receiver_context(device_x2): + dev0, dev1 = device_x2 + dev0.set_current() + ctx0 = dev0.context + dev1.set_current() + ctx1_handle = current_context_handle() + + array = dev0.create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + is_surface_load_store=True, + ) + ) + assert array.device == dev0 + assert current_context_handle() == ctx1_handle + + mipmap = dev0.create_mipmapped_array( + MipmappedArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + num_levels=2, + ) + ) + assert mipmap.device == dev0 + assert current_context_handle() == ctx1_handle + + level = mipmap.get_level(0) + assert level.device == dev0 + assert current_context_handle() == ctx1_handle + + resource = ResourceDescriptor.from_opaque_array(array) + texture = dev0.create_texture_object(resource=resource, options=TextureObjectOptions()) + surface = dev0.create_surface_object(resource=resource) + assert texture.device == dev0 + assert surface.device == dev0 + assert current_context_handle() == ctx1_handle + + surface.close() + texture.close() + level.close() + mipmap.close() + array.close() + assert current_context_handle() == ctx1_handle + assert int(ctx0.handle) != ctx1_handle + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_texture_resources_restore_no_current_context(deinit_cuda): + device = Device(0) + device.set_current() + + array = texture = None + try: + with no_current_context(): + array = device.create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + ) + ) + texture = device.create_texture_object(resource=ResourceDescriptor.from_opaque_array(array)) + assert current_context_handle() == 0 + + texture.close() + texture = None + array.close() + array = None + assert current_context_handle() == 0 + finally: + if texture is not None: + texture.close() + if array is not None: + array.close() + + +@pytest.mark.agent_authored(model="gpt-5.6") +def test_texture_creation_rejects_mismatched_receiver(device_x2): + dev0, dev1 = device_x2 + dev0.set_current() + array = dev0.create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + is_surface_load_store=True, + ) + ) + resource = ResourceDescriptor.from_opaque_array(array) + dev1.set_current() + ctx1_handle = current_context_handle() + + with pytest.raises(ValueError, match="resource belongs to device 0"): + dev1.create_texture_object(resource=resource) + with pytest.raises(ValueError, match="resource belongs to device 0"): + dev1.create_surface_object(resource=resource) + assert current_context_handle() == ctx1_handle + array.close() + + def test_array_2d_create_and_properties(init_cuda): arr = Device().create_opaque_array( OpaqueArrayOptions(shape=(32, 16), format=ArrayFormatType.FLOAT32, num_channels=1) From 51ede5119b24282eb98a08eb302d33e03f810b0a Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Fri, 4 Sep 2026 11:17:26 -0700 Subject: [PATCH 2/9] cuda.core: address review of #2750 - Fix _SynchronousMemoryResource to record a deallocation token bound to its own context, so Buffer teardown enters the right context regardless of what is current, and works with no context current. Move the class to its own module and resolve the primary context lazily. - LegacyPinnedMemoryResource.device_id returns -1 as documented; texture creation over a pinned buffer works again. - Device.set_current delegates a foreign-device context to its owning device instead of raising, so the save/restore idiom round-trips across devices. - Guard empty context handles once in invoke_in_context(_or_undo); warn when an undo is skipped because the target context is gone. - context_synchronize and context_get_stream_priority_range release the GIL like the other helpers. Rename array/mipmap box accessors to get_box. - Tests: query the driver (cuStreamGetCtx, cross-context cuEventRecord) instead of comparing cached metadata; add sync(), set_current round-trip, pinned texture, and synchronous-resource teardown tests; register device_x2 with the parallel-test plugin. - Move the release note to 1.3.0 and describe sync() as acting on the bound context. Co-Authored-By: Claude Fable 5.1 --- cuda_core/cuda/core/_cpp/resource_handles.cpp | 42 ++++-- cuda_core/cuda/core/_cpp/resource_handles.hpp | 10 +- cuda_core/cuda/core/_device.pyi | 13 +- cuda_core/cuda/core/_device.pyx | 29 +++-- cuda_core/cuda/core/_memory/_buffer.pyi | 2 +- cuda_core/cuda/core/_memory/_buffer.pyx | 2 +- .../core/_memory/_device_memory_resource.pyi | 17 --- .../core/_memory/_device_memory_resource.pyx | 78 +----------- cuda_core/cuda/core/_memory/_legacy.py | 4 +- .../_memory/_synchronous_memory_resource.pyi | 23 ++++ .../_memory/_synchronous_memory_resource.pyx | 115 +++++++++++++++++ cuda_core/cuda/core/_resource_handles.pxd | 2 + cuda_core/cuda/core/_resource_handles.pyx | 2 + cuda_core/cuda/core/_stream.pyx | 4 +- cuda_core/cuda/core/texture/_texture.pyx | 4 +- cuda_core/docs/source/release/1.2.0-notes.rst | 7 - cuda_core/docs/source/release/1.3.0-notes.rst | 29 +++++ cuda_core/tests/conftest.py | 17 ++- cuda_core/tests/helpers/contexts.py | 38 ++++++ cuda_core/tests/test_device.py | 52 ++++++++ cuda_core/tests/test_green_context.py | 27 ++-- cuda_core/tests/test_launcher.py | 2 +- cuda_core/tests/test_memory.py | 56 ++++++-- cuda_core/tests/test_texture_surface.py | 120 +++++++++--------- 24 files changed, 471 insertions(+), 224 deletions(-) create mode 100644 cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi create mode 100644 cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx create mode 100644 cuda_core/docs/source/release/1.3.0-notes.rst diff --git a/cuda_core/cuda/core/_cpp/resource_handles.cpp b/cuda_core/cuda/core/_cpp/resource_handles.cpp index ee4aef8e2e9..ec9c033829e 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.cpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.cpp @@ -270,6 +270,9 @@ ContextHandle deallocation_context(const DeallocationStream& stream) noexcept { template CUresult invoke_in_context(const ContextHandle& h_context, Fn&& operation, Args&&... args) noexcept { ASSERT_NOTHROW_INVOCABLE(Fn&&, Args&&...); + if (!h_context) { + return CUDA_ERROR_INVALID_CONTEXT; + } CUcontext previous = nullptr; int changed = 0; CUresult status = enter_context(h_context, &previous, &changed); @@ -288,6 +291,9 @@ CUresult invoke_in_context_or_undo(const ContextHandle& h_context, Fn&& operatio Undo&& undo, bool undo_requires_target_context) noexcept { ASSERT_NOTHROW_INVOCABLE(Fn&&); ASSERT_NOTHROW_INVOCABLE(Undo&&); + if (!h_context) { + return CUDA_ERROR_INVALID_CONTEXT; + } CUcontext previous = nullptr; int changed = 0; CUresult status = enter_context(h_context, &previous, &changed); @@ -305,6 +311,11 @@ 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( + "cuCtxSetCurrent (restoring the caller's context)", composite, + "failed; cleanup of the new resource skipped because its context " + "is no longer current (resource leaked)"); } } return composite; @@ -398,9 +409,7 @@ const WarnOnFailure pw_cuSurfObjectDestroy{"cuSurfObjectD // Synchronize the provided context. CUresult context_synchronize(const ContextHandle& h_context) noexcept { - if (!h_context) { - return CUDA_ERROR_INVALID_CONTEXT; - } + GILReleaseGuard gil; return invoke_in_context(h_context, []() noexcept { return p_cuCtxSynchronize(); }); @@ -410,9 +419,7 @@ CUresult context_synchronize(const ContextHandle& h_context) noexcept { CUresult context_get_stream_priority_range(const ContextHandle& h_context, int* least_priority, int* greatest_priority) noexcept { - if (!h_context) { - return CUDA_ERROR_INVALID_CONTEXT; - } + GILReleaseGuard gil; return invoke_in_context(h_context, [&]() noexcept { return p_cuCtxGetStreamPriorityRange(least_priority, greatest_priority); }); @@ -954,6 +961,16 @@ StreamHandle get_per_thread_stream() { return handle; } +StreamHandle create_context_bound_legacy_stream(const ContextHandle& h_context) { + if (!h_context) { + return {}; + } + // Default deleter: this handle never owns CU_STREAM_LEGACY, so nothing + // needs to run when the last reference is released. + auto box = std::make_shared(StreamBox{CU_STREAM_LEGACY, h_context}); + return StreamHandle(box, &box->resource); +} + // ============================================================================ // Deallocation streams // @@ -1301,9 +1318,6 @@ DevicePtrHandle deviceptr_alloc_async(size_t size, const StreamHandle& h_stream) // Allocate device memory synchronously with the provided context current. CUresult deviceptr_alloc_raw(CUdeviceptr* ptr, size_t size, const ContextHandle& h_context) noexcept { - if (!h_context) { - return CUDA_ERROR_INVALID_CONTEXT; - } GILReleaseGuard gil; return invoke_in_context_or_undo( h_context, @@ -2837,14 +2851,14 @@ struct SurfObjectBox { }; // Recover an array's owning box from its aliased resource pointer. -const ArrayBox* get_array_box(const OpaqueArrayHandle& h) noexcept { +const ArrayBox* get_box(const OpaqueArrayHandle& h) noexcept { const CUarray* p = h.get(); return reinterpret_cast( reinterpret_cast(p) - offsetof(ArrayBox, resource)); } // Recover a mipmapped array's owning box from its aliased resource pointer. -const MipmappedArrayBox* get_mipmapped_array_box(const MipmappedArrayHandle& h) noexcept { +const MipmappedArrayBox* get_box(const MipmappedArrayHandle& h) noexcept { const CUmipmappedArray* p = h.get(); return reinterpret_cast( reinterpret_cast(p) @@ -2897,13 +2911,13 @@ OpaqueArrayHandle create_array_handle_owning(CUarray arr) { // Return the context retained by an array handle. ContextHandle get_array_context(const OpaqueArrayHandle& h) noexcept { - return h ? get_array_box(h)->h_context : ContextHandle{}; + return h ? get_box(h)->h_context : ContextHandle{}; } OpaqueArrayHandle create_array_level_handle(const MipmappedArrayHandle& h_mip, unsigned int level) { GILReleaseGuard gil; CUarray arr; - ContextHandle h_context = h_mip ? get_mipmapped_array_box(h_mip)->h_context : ContextHandle{}; + ContextHandle h_context = h_mip ? get_box(h_mip)->h_context : ContextHandle{}; if (CUDA_SUCCESS != (err = p_cuMipmappedArrayGetLevel(&arr, as_cu(h_mip), level))) { return {}; } @@ -2942,7 +2956,7 @@ MipmappedArrayHandle create_mipmapped_array_handle(const ContextHandle& h_contex // Return the context retained by a mipmapped array handle. ContextHandle get_mipmapped_array_context(const MipmappedArrayHandle& h) noexcept { - return h ? get_mipmapped_array_box(h)->h_context : ContextHandle{}; + return h ? get_box(h)->h_context : ContextHandle{}; } namespace { diff --git a/cuda_core/cuda/core/_cpp/resource_handles.hpp b/cuda_core/cuda/core/_cpp/resource_handles.hpp index 57ac9a244d6..525100cb0db 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.hpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.hpp @@ -251,7 +251,7 @@ ContextHandle get_primary_context(int device_id); // Returns empty handle if no context is current (caller must check) ContextHandle get_current_context(); -// Synchronize the provided context. +// Synchronize the provided context. Releases the GIL around the driver call. // Returns CUDA_ERROR_INVALID_CONTEXT for an empty handle. CUresult context_synchronize(const ContextHandle& h_context) noexcept; @@ -303,6 +303,14 @@ StreamHandle get_legacy_stream(); // Note: Per-thread stream has no specific context dependency. StreamHandle get_per_thread_stream(); +// Wrap CU_STREAM_LEGACY with an explicit context, bypassing the "bind to +// whatever is current" resolution that a bare default-stream token uses (see +// make_deallocation_stream). Lets a resource that always operates in one +// known context (e.g. a synchronous, non-pooled allocator) record a correct +// deallocation context without requiring that context to be current when the +// token is created. Returns an empty handle for an empty h_context. +StreamHandle create_context_bound_legacy_stream(const ContextHandle& h_context); + // ============================================================================ // Event handle functions // ============================================================================ diff --git a/cuda_core/cuda/core/_device.pyi b/cuda_core/cuda/core/_device.pyi index 39ff7d3a26e..4f656a0d27f 100644 --- a/cuda_core/cuda/core/_device.pyi +++ b/cuda_core/cuda/core/_device.pyi @@ -606,6 +606,12 @@ class Device: Providing a `ctx` causes the previous set context to be popped and returned. + If `ctx` was created on a different device than this receiver, the call + is delegated to that device's own :meth:`set_current`. This keeps the + owning device's bookkeeping consistent and lets a context this method + handed out for a foreign device be pushed back through any ``Device`` + object, matching the CUDA context stack's own thread-wide semantics. + Parameters ---------- ctx : :obj:`~_context.Context`, optional @@ -719,7 +725,12 @@ class Device: """ def sync(self) -> None: - """Synchronize this device. + """Synchronize this device's bound context. + + Waits for all preceding work in this device's bound :obj:`~_context.Context` + to complete. Only that context is synchronized, not the device as a + whole; work queued in a different context on the same device (e.g. a + green context) is unaffected. Note ---- diff --git a/cuda_core/cuda/core/_device.pyx b/cuda_core/cuda/core/_device.pyx index 78245c6107d..a7e7d59e04a 100644 --- a/cuda_core/cuda/core/_device.pyx +++ b/cuda_core/cuda/core/_device.pyx @@ -1206,7 +1206,7 @@ class Device: from cuda.core._memory import DeviceMemoryResource self._memory_resource = DeviceMemoryResource(self._device_id) else: - from cuda.core._memory._device_memory_resource import ( + from cuda.core._memory._synchronous_memory_resource import ( _SynchronousMemoryResource, ) self._memory_resource = _SynchronousMemoryResource( @@ -1263,6 +1263,12 @@ class Device: Providing a `ctx` causes the previous set context to be popped and returned. + If `ctx` was created on a different device than this receiver, the call + is delegated to that device's own :meth:`set_current`. This keeps the + owning device's bookkeeping consistent and lets a context this method + handed out for a foreign device be pushed back through any ``Device`` + object, matching the CUDA context stack's own thread-wide semantics. + Parameters ---------- ctx : :obj:`~_context.Context`, optional @@ -1296,10 +1302,11 @@ class Device: assert_type(ctx, Context) Context_check_open(ctx) if ctx._device_id != self._device_id: - raise RuntimeError( - "the provided context was created on the device with" - f" id={ctx._device_id}, which is different from the target id={self._device_id}" - ) + # The CUDA context stack is per-thread, not per-Device-object, + # so pushing/popping a foreign-device context is delegated to + # the device that owns it; its own bookkeeping (_context, + # _has_inited) is what should track this push, not ours. + return Device(ctx._device_id).set_current(ctx) if self._has_inited and self._context is not None: prev_owned = self._context curr_ctx = as_cu(ctx._h_context) @@ -1478,7 +1485,12 @@ class Device: return self.memory_resource.allocate(size, stream=stream) def sync(self) -> None: - """Synchronize this device. + """Synchronize this device's bound context. + + Waits for all preceding work in this device's bound :obj:`~_context.Context` + to complete. Only that context is synchronized, not the device as a + whole; work queued in a different context on the same device (e.g. a + green context) is unaffected. Note ---- @@ -1487,10 +1499,7 @@ class Device: """ self._check_context_initialized() cdef Context ctx = self._context - cdef cydriver.CUresult status - with nogil: - status = context_synchronize(ctx._h_context) - HANDLE_RETURN(status) + HANDLE_RETURN(context_synchronize(ctx._h_context)) def create_graph_builder(self) -> GraphBuilder: """Create a new :obj:`~graph.GraphBuilder` on this device. diff --git a/cuda_core/cuda/core/_memory/_buffer.pyi b/cuda_core/cuda/core/_memory/_buffer.pyi index 616d62060da..3c02e731f48 100644 --- a/cuda_core/cuda/core/_memory/_buffer.pyi +++ b/cuda_core/cuda/core/_memory/_buffer.pyi @@ -250,7 +250,7 @@ class Buffer: def __release_buffer__(self, buffer: memoryview, /) -> None: ... @property def device_id(self) -> int: - """Return the device ordinal of this buffer.""" + """Return the device ordinal of this buffer, or -1 for memory not bound to a device.""" @property def handle(self) -> int: """Return the buffer handle object. diff --git a/cuda_core/cuda/core/_memory/_buffer.pyx b/cuda_core/cuda/core/_memory/_buffer.pyx index 316f07f2912..2484ad82b00 100644 --- a/cuda_core/cuda/core/_memory/_buffer.pyx +++ b/cuda_core/cuda/core/_memory/_buffer.pyx @@ -666,7 +666,7 @@ cdef class Buffer: @property def device_id(self) -> int: - """Return the device ordinal of this buffer.""" + """Return the device ordinal of this buffer, or -1 for memory not bound to a device.""" Buffer_check_open(self) if self._memory_resource is not None: return self._memory_resource.device_id diff --git a/cuda_core/cuda/core/_memory/_device_memory_resource.pyi b/cuda_core/cuda/core/_memory/_device_memory_resource.pyi index b74c344d2cd..897e1d03302 100644 --- a/cuda_core/cuda/core/_memory/_device_memory_resource.pyi +++ b/cuda_core/cuda/core/_memory/_device_memory_resource.pyi @@ -4,13 +4,9 @@ import uuid from dataclasses import dataclass from cuda.core._device import Device -from cuda.core._memory._buffer import Buffer, MemoryResource from cuda.core._memory._ipc import IPCAllocationHandle from cuda.core._memory._memory_pool import _MemPool from cuda.core._memory._peer_access_utils import PeerAccessibleBySetProxy -from cuda.core._stream import Stream -from cuda.core.graph import GraphBuilder -from cuda.core.typing import DevicePointerType __all__ = ['DeviceMemoryResource', 'DeviceMemoryResourceOptions'] @@ -32,19 +28,6 @@ class DeviceMemoryResourceOptions: ipc_enabled: bool = False max_size: int = 0 -class _SynchronousMemoryResource(MemoryResource): - __slots__ = ('_context', '_device_id') - - def __init__(self, device_id: int, context=None) -> None: ... - def allocate(self, size: int, *, stream: Stream | GraphBuilder | None=None) -> Buffer: ... - def deallocate(self, ptr: DevicePointerType, size: int, *, stream: Stream | GraphBuilder | None=None) -> None: ... - @property - def is_device_accessible(self) -> bool: ... - @property - def is_host_accessible(self) -> bool: ... - @property - def device_id(self) -> int: ... - class DeviceMemoryResource(_MemPool): """ A device memory resource managing a stream-ordered memory pool. diff --git a/cuda_core/cuda/core/_memory/_device_memory_resource.pyx b/cuda_core/cuda/core/_memory/_device_memory_resource.pyx index 426ddb6d8dc..62dc4f9e747 100644 --- a/cuda_core/cuda/core/_memory/_device_memory_resource.pyx +++ b/cuda_core/cuda/core/_memory/_device_memory_resource.pyx @@ -4,26 +4,14 @@ from __future__ import annotations -from libc.stdint cimport uintptr_t - from cuda.bindings cimport cydriver -from cuda.core._context cimport Context -from cuda.core._memory._buffer cimport Buffer, MemoryResource from cuda.core._memory._location cimport cumemlocation_from_id from cuda.core._memory._memory_pool cimport ( _MemPool, MP_check_open, MP_init_create_pool, MP_raise_release_threshold, ) from cuda.core._memory cimport _ipc from cuda.core._memory._ipc cimport IPCAllocationHandle -from cuda.core._resource_handles cimport ( - ContextHandle, - as_cu, - deviceptr_alloc_raw, - get_device_mempool, - get_last_error, - get_primary_context, -) -from cuda.core._stream cimport Stream, Stream_accept +from cuda.core._resource_handles cimport as_cu, get_device_mempool, get_last_error from cuda.core._utils.cuda_utils cimport ( check_or_create_options, HANDLE_RETURN, @@ -42,8 +30,6 @@ from typing import TYPE_CHECKING if TYPE_CHECKING: from cuda.core._device import Device - from cuda.core.graph import GraphBuilder - from cuda.core.typing import DevicePointerType __all__ = ['DeviceMemoryResource', 'DeviceMemoryResourceOptions'] @@ -67,68 +53,6 @@ cdef class DeviceMemoryResourceOptions: max_size : int = 0 -class _SynchronousMemoryResource(MemoryResource): - __slots__ = ("_context", "_device_id") - - def __init__(self, device_id: int, context=None) -> None: - cdef ContextHandle h_context - from .._device import Device - - self._device_id = Device(device_id).device_id - if context is None: - h_context = get_primary_context(self._device_id) - if not h_context: - HANDLE_RETURN(get_last_error()) - context = Context._from_handle( - Context, h_context, self._device_id) - self._context = context - - def allocate( - self, - size_t size, - *, - stream: Stream | GraphBuilder | None = None, - ) -> Buffer: - # cuMemAlloc is synchronous; stream is accepted (and validated) - # for interface conformance but not used. - if stream is not None: - Stream_accept(stream) - - cdef Context context = self._context - cdef cydriver.CUdeviceptr ptr = 0 - if size: - with nogil: - HANDLE_RETURN(deviceptr_alloc_raw(&ptr, size, context._h_context)) - return Buffer._init(ptr, size, self) - - def deallocate( - self, - ptr: DevicePointerType, - size_t size, - *, - stream: Stream | GraphBuilder | None = None, - ) -> None: - if stream is not None: - Stream_accept(stream).sync() - cdef cydriver.CUdeviceptr devptr - if size: - devptr = int(ptr) - with nogil: - HANDLE_RETURN(cydriver.cuMemFree(devptr)) - - @property - def is_device_accessible(self) -> bool: - return True - - @property - def is_host_accessible(self) -> bool: - return False - - @property - def device_id(self) -> int: - return self._device_id - - cdef class DeviceMemoryResource(_MemPool): """ A device memory resource managing a stream-ordered memory pool. diff --git a/cuda_core/cuda/core/_memory/_legacy.py b/cuda_core/cuda/core/_memory/_legacy.py index a2a8843a448..f3dff33a133 100644 --- a/cuda_core/cuda/core/_memory/_legacy.py +++ b/cuda_core/cuda/core/_memory/_legacy.py @@ -94,5 +94,5 @@ def is_host_accessible(self) -> bool: @property def device_id(self) -> int: - """This memory resource is not bound to any GPU.""" - raise RuntimeError("a pinned memory resource is not bound to any GPU") + """Return -1. Pinned memory is host memory and is not bound to a specific device.""" + return -1 diff --git a/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi b/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi new file mode 100644 index 00000000000..4dfc17f10f2 --- /dev/null +++ b/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi @@ -0,0 +1,23 @@ +# This file was generated by stubgen-pyx v0.2.19 from cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx + +from cuda.core._context import Context +from cuda.core._memory._buffer import Buffer, MemoryResource +from cuda.core._stream import Stream +from cuda.core.graph import GraphBuilder +from cuda.core.typing import DevicePointerType + +__all__ = [] + +class _SynchronousMemoryResource(MemoryResource): + __slots__ = ('_context', '_device_id') + + def __init__(self, device_id: int, context=None) -> None: ... + def _resolve_context(self) -> Context: ... + def allocate(self, size: int, *, stream: Stream | GraphBuilder | None=None) -> Buffer: ... + def deallocate(self, ptr: DevicePointerType, size: int, *, stream: Stream | GraphBuilder | None=None) -> None: ... + @property + def is_device_accessible(self) -> bool: ... + @property + def is_host_accessible(self) -> bool: ... + @property + def device_id(self) -> int: ... diff --git a/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx b/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx new file mode 100644 index 00000000000..f02f38f69b1 --- /dev/null +++ b/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx @@ -0,0 +1,115 @@ +# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: Apache-2.0 + +from __future__ import annotations + +from libc.stdint cimport uintptr_t + +from cuda.bindings cimport cydriver +from cuda.core._context cimport Context +from cuda.core._memory._buffer cimport Buffer, MemoryResource +from cuda.core._resource_handles cimport ( + ContextHandle, + create_context_bound_legacy_stream, + deviceptr_alloc_raw, + get_last_error, + get_primary_context, +) +from cuda.core._stream cimport Stream, Stream_accept, Stream_is_default_token +from cuda.core._utils.cuda_utils cimport HANDLE_RETURN + +from typing import TYPE_CHECKING + +if TYPE_CHECKING: + from cuda.core.graph import GraphBuilder + from cuda.core.typing import DevicePointerType + +__all__ = [] + + +class _SynchronousMemoryResource(MemoryResource): + __slots__ = ("_context", "_device_id") + + def __init__(self, device_id: int, context=None) -> None: + from .._device import Device + + self._device_id = Device(device_id).device_id + # Resolved lazily (in _resolve_context) so that construction with + # context=None does no CUDA work; the primary context is retained + # only once actually needed, on the first allocate()/deallocate(). + self._context = context + + def _resolve_context(self) -> Context: + cdef ContextHandle h_context + if self._context is None: + h_context = get_primary_context(self._device_id) + if not h_context: + HANDLE_RETURN(get_last_error()) + self._context = Context._from_handle( + Context, h_context, self._device_id) + return self._context + + def allocate( + self, + size_t size, + *, + stream: Stream | GraphBuilder | None = None, + ) -> Buffer: + # cuMemAlloc/cuMemFree are synchronous; a caller-supplied stream is + # accepted (and validated) for interface conformance and, if it is a + # real stream, recorded as the stream that orders deallocation. + cdef Context context = self._resolve_context() + cdef Stream dealloc_stream = None + if stream is not None: + dealloc_stream = Stream_accept(stream) + if dealloc_stream is None or Stream_is_default_token(dealloc_stream): + # A default-stream token carries no context of its own; Buffer._init + # would bind it to whichever context is current when it records the + # deallocation stream (and fail if none is). Bind it to this + # resource's context instead, so Buffer teardown frees in the right + # context no matter what is current then. Always the legacy token: + # a per-thread token would also arm the cross-thread PTDS warning, + # which is noise for a synchronous resource. + dealloc_stream = Stream._from_handle( + Stream, create_context_bound_legacy_stream(context._h_context)) + + cdef cydriver.CUdeviceptr ptr = 0 + if size: + with nogil: + HANDLE_RETURN(deviceptr_alloc_raw(&ptr, size, context._h_context)) + return Buffer._init(ptr, size, self, stream=dealloc_stream) + + def deallocate( + self, + ptr: DevicePointerType, + size_t size, + *, + stream: Stream | GraphBuilder | None = None, + ) -> None: + if stream is not None: + Stream_accept(stream).sync() + # No context switch here, by design (settled in the review of #2750): + # cuMemFree does not need a current context. The driver resolves the + # allocation's owning context from the pointer through unified + # addressing and frees it there (cuapiMemFree_common: "a current + # context is not required to free the device memory"). On the Buffer + # teardown path the C++ deleter has additionally already made the + # recorded deallocation context, bound by allocate() above, current. + cdef cydriver.CUdeviceptr devptr + if size: + devptr = int(ptr) + with nogil: + HANDLE_RETURN(cydriver.cuMemFree(devptr)) + + @property + def is_device_accessible(self) -> bool: + return True + + @property + def is_host_accessible(self) -> bool: + return False + + @property + def device_id(self) -> int: + return self._device_id diff --git a/cuda_core/cuda/core/_resource_handles.pxd b/cuda_core/cuda/core/_resource_handles.pxd index acf10e0fa2c..f8dea004831 100644 --- a/cuda_core/cuda/core/_resource_handles.pxd +++ b/cuda_core/cuda/core/_resource_handles.pxd @@ -195,6 +195,8 @@ cdef void retry_deferred_cleanup() noexcept cdef ContextHandle get_stream_context(const StreamHandle& h) noexcept nogil cdef StreamHandle get_legacy_stream() except+ nogil cdef StreamHandle get_per_thread_stream() except+ nogil +cdef StreamHandle create_context_bound_legacy_stream( + const ContextHandle& h_context) except+ nogil # Event handles cdef EventHandle create_event_handle( diff --git a/cuda_core/cuda/core/_resource_handles.pyx b/cuda_core/cuda/core/_resource_handles.pyx index beecb4b745a..da4a5a20ff3 100644 --- a/cuda_core/cuda/core/_resource_handles.pyx +++ b/cuda_core/cuda/core/_resource_handles.pyx @@ -73,6 +73,8 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": const StreamHandle& h) noexcept nogil StreamHandle get_legacy_stream "cuda_core::get_legacy_stream" () except+ nogil StreamHandle get_per_thread_stream "cuda_core::get_per_thread_stream" () except+ nogil + StreamHandle create_context_bound_legacy_stream "cuda_core::create_context_bound_legacy_stream" ( + const ContextHandle& h_context) except+ nogil # Event handles (note: _create_event_handle* are internal due to C++ overloading) EventHandle create_event_handle "cuda_core::create_event_handle" ( diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index 9c51c488a6c..ed473b977d3 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -160,9 +160,7 @@ cdef class Stream: # TODO: we might want to consider memoizing high/low per CUDA context and avoid this call cdef int high, low cdef cydriver.CUresult res_code - with nogil: - res_code = context_get_stream_priority_range( - context._h_context, &high, &low) + res_code = context_get_stream_priority_range(context._h_context, &high, &low) HANDLE_RETURN(res_code) cdef int prio if priority is not None: diff --git a/cuda_core/cuda/core/texture/_texture.pyx b/cuda_core/cuda/core/texture/_texture.pyx index f1c29490b8a..28ddf2d6aa8 100644 --- a/cuda_core/cuda/core/texture/_texture.pyx +++ b/cuda_core/cuda/core/texture/_texture.pyx @@ -521,7 +521,7 @@ def _create_texture_object( elif resource.kind == "linear": buf = resource.source Buffer_check_open(buf) - resource_device_id = buf.device_id + resource_device_id = buf.device_id # -1 for memory not bound to a device devptr = int(buf.handle) res_desc.resType = cydriver.CU_RESOURCE_TYPE_LINEAR res_desc.res.linear.devPtr = devptr @@ -531,7 +531,7 @@ def _create_texture_object( elif resource.kind == "pitch2d": buf = resource.source Buffer_check_open(buf) - resource_device_id = buf.device_id + resource_device_id = buf.device_id # -1 for memory not bound to a device devptr = int(buf.handle) res_desc.resType = cydriver.CU_RESOURCE_TYPE_PITCH2D res_desc.res.pitch2D.devPtr = devptr diff --git a/cuda_core/docs/source/release/1.2.0-notes.rst b/cuda_core/docs/source/release/1.2.0-notes.rst index 9591cc3a083..f96a205d1e8 100644 --- a/cuda_core/docs/source/release/1.2.0-notes.rst +++ b/cuda_core/docs/source/release/1.2.0-notes.rst @@ -40,13 +40,6 @@ New features Fixes and enhancements ---------------------- -- :class:`Device` methods that create resources or synchronize now act on that - device, even when another device is current. They do not change which device - is current. For example, ``dev1.sync()`` synchronizes device 1 even when - device 0 is current. :meth:`Device.set_current` now always returns a - :class:`Context` with the correct device ID. - (`#2311 `__) - - A :class:`Buffer` is now freed correctly even when the CUDA context current at teardown is not the one it was allocated in, or when no context is current at all. This happens routinely when a buffer is released by the garbage diff --git a/cuda_core/docs/source/release/1.3.0-notes.rst b/cuda_core/docs/source/release/1.3.0-notes.rst new file mode 100644 index 00000000000..a76fb8409a2 --- /dev/null +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -0,0 +1,29 @@ +.. SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +.. SPDX-License-Identifier: Apache-2.0 + +.. currentmodule:: cuda.core + +``cuda.core`` 1.3.0 Release Notes +================================== + +Fixes and enhancements +---------------------- + +- :class:`Device` methods that create resources or synchronize now act on that + device's bound context, even when another device is current. They do not + change which device is current. For example, ``dev1.sync()`` synchronizes + device 1's bound context even when device 0 is current, and no longer + touches other contexts on device 1 (such as a green context). :meth:`Device.set_current` + now always returns a :class:`Context` with the correct device ID; passing a + context created on a different device than the receiver now delegates to + that device's own :meth:`~Device.set_current` instead of raising, so a + context this method returns can always be pushed back through any + :class:`Device` object. + (`#2311 `__) + +- :attr:`LegacyPinnedMemoryResource.device_id` now returns ``-1``, as + documented for memory that is not bound to a device and as + :class:`PinnedMemoryResource` already does, instead of raising + ``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. diff --git a/cuda_core/tests/conftest.py b/cuda_core/tests/conftest.py index 0466277bf6c..91265fc47b8 100644 --- a/cuda_core/tests/conftest.py +++ b/cuda_core/tests/conftest.py @@ -116,6 +116,8 @@ def wrapper(*args, **kwargs): kwargs["mempool_device_x2"] = _mempool_device_impl(2) if "mempool_device_x3" in kwargs: kwargs["mempool_device_x3"] = _mempool_device_impl(3) + if "device_x2" in kwargs: + kwargs["device_x2"] = _device_x2_impl() # These are used by test_green_context.py. The original fixtures include # pytest.skip() but that should have correctly fired by this time. @@ -212,15 +214,24 @@ def deinit_cuda(): _ = _device_unset_current() -@pytest.fixture -def device_x2(deinit_cuda): - """Provide two CUDA devices, or skip when fewer are available.""" +def _device_x2_impl(): devices = Device.get_all_devices() if len(devices) < 2: pytest.skip("Test requires at least 2 CUDA devices") return devices[:2] +@pytest.fixture +def device_x2(init_cuda): + """Provide two CUDA devices, or skip when fewer are available. + + Depends on ``init_cuda`` so that, under pytest-run-parallel, the test is + wrapped by ``_wrap_worker_cuda_test`` and the devices are re-fetched on the + worker thread (Device objects are thread-local). + """ + return _device_x2_impl() + + @pytest.fixture def deinit_all_contexts_function(): def pop_all_contexts(): diff --git a/cuda_core/tests/helpers/contexts.py b/cuda_core/tests/helpers/contexts.py index 9818875c4da..a2227c0f737 100644 --- a/cuda_core/tests/helpers/contexts.py +++ b/cuda_core/tests/helpers/contexts.py @@ -18,20 +18,58 @@ def current_context_handle(): return int(handle_return(driver.cuCtxGetCurrent())) +def _assert_event_record_rejected_from_ambient_context(event): + """Assert that recording ``event`` into a stream from the ambient context fails. + + An event and the stream it records must belong to the same context; + otherwise cuEventRecord fails with CUDA_ERROR_INVALID_HANDLE. cuStreamCreate + creates the probe stream in whatever context is currently ambient, so this + is a live check that ``event`` was not actually recorded from there. + """ + err, ambient_stream = driver.cuStreamCreate(0) + handle_return(err) + try: + (record_status,) = driver.cuEventRecord(event.handle, ambient_stream) + assert record_status == driver.CUresult.CUDA_ERROR_INVALID_HANDLE, ( + "Recording an event into a stream from a different context should fail with " + f"CUDA_ERROR_INVALID_HANDLE, got {record_status!r}" + ) + finally: + handle_return(driver.cuStreamDestroy(ambient_stream)) + + def assert_device_operations_use_bound_context(device): """Check that Device operations use its bound context and preserve the ambient context.""" bound_context = device.context ambient_context_handle = current_context_handle() + assert int(bound_context.handle) != ambient_context_handle, ( + "Precondition failed: the device's bound context must not be the current (ambient) context." + ) stream = event = builder = None try: stream = device.create_stream() assert stream.context == bound_context assert current_context_handle() == ambient_context_handle + # Live check: query the driver directly, rather than comparing cached + # metadata, to confirm the stream was actually created in the bound + # context rather than whatever was ambient. + driver_stream_ctx = handle_return(driver.cuStreamGetCtx(stream.handle)) + assert int(driver_stream_ctx) == int(bound_context.handle), ( + "cuStreamGetCtx reports a context other than the one the stream was created in." + ) event = device.create_event() assert event.context == bound_context assert current_context_handle() == ambient_context_handle + # Only exercised when the ambient context belongs to a different + # physical device: cuEventRecord's cross-context rejection is + # guaranteed distinct there. Two contexts on the *same* device (e.g. + # a green context vs. the primary context) may resolve to the same + # underlying device context for this check, so skip rather than + # assert unverified driver behavior. + if ambient_context_handle and handle_return(driver.cuCtxGetDevice()) != device.device_id: + _assert_event_record_rejected_from_ambient_context(event) builder = device.create_graph_builder() assert builder.stream.context == bound_context diff --git a/cuda_core/tests/test_device.py b/cuda_core/tests/test_device.py index 48f5b3cf484..dbfb9fed5a9 100644 --- a/cuda_core/tests/test_device.py +++ b/cuda_core/tests/test_device.py @@ -10,6 +10,7 @@ current_context_handle, no_current_context, ) +from helpers.nanosleep_kernel import NanosleepKernel import cuda.core from cuda.bindings import driver, runtime @@ -147,6 +148,57 @@ def test_set_current_returns_previous_context_with_owning_device(device_x2): dev1.set_current(previous) +@pytest.mark.agent_authored(model="claude-sonnet-5") +def test_set_current_round_trips_through_a_different_device(device_x2): + """The pre-#2311 idiom `prev = dev.set_current(ctx); ...; dev.set_current(prev)` + must keep working even when `prev` belongs to a different device than + `dev`: set_current() delegates to the context's owning device instead of + raising, so restoring through the original Device handle round-trips.""" + dev0, dev1 = device_x2 + dev0.set_current() + ctx0 = dev0.context + + dev1.set_current() + ctx1 = dev1.context + + dev0.set_current() # dev0 current again; dev1.context is still ctx1 + + prev = dev1.set_current(ctx1) + assert prev.handle == ctx0.handle + assert current_context_handle() == int(ctx1.handle) + + restored = dev1.set_current(prev) + assert restored.handle == ctx1.handle + assert current_context_handle() == int(ctx0.handle) + + +@pytest.mark.agent_authored(model="claude-sonnet-5") +def test_device_sync_waits_for_bound_context_work(device_x2): + """dev0.sync() must wait for work queued on dev0's bound context even + while dev1 is ambient, not just preserve the ambient context (#2311).""" + dev0, dev1 = device_x2 + if dev0.compute_capability.major < 7: + pytest.skip("__nanosleep is only available starting Volta (sm70)") + dev0.set_current() + stream = dev0.create_stream() + nanosleep = NanosleepKernel(dev0, sleep_duration_ms=20) + event = None + try: + nanosleep.launch(stream) + event = stream.record() + + dev1.set_current() + ambient_context_handle = current_context_handle() + + dev0.sync() + assert event.is_done + assert current_context_handle() == ambient_context_handle + finally: + if event is not None: + event.close() + stream.close() + + @pytest.mark.agent_authored(model="gpt-5.6") def test_device_receiver_switching_is_thread_local(device_x2): dev0, dev1 = device_x2 diff --git a/cuda_core/tests/test_green_context.py b/cuda_core/tests/test_green_context.py index 5b88083d22b..5f1954c6b58 100644 --- a/cuda_core/tests/test_green_context.py +++ b/cuda_core/tests/test_green_context.py @@ -532,21 +532,18 @@ def test_texture_rejects_resource_from_other_context(self, init_cuda, green_ctx) ) from cuda.core.typing import ArrayFormatType - array = init_cuda.create_opaque_array( - OpaqueArrayOptions( - shape=(8, 8), - format=ArrayFormatType.UINT8, - num_channels=4, - ) - ) - try: - with ( - use_context(init_cuda, green_ctx), - pytest.raises(ValueError, match="resource is not compatible with this Device object"), - ): - init_cuda.create_texture_object(resource=ResourceDescriptor.from_opaque_array(array)) - finally: - array.close() + with ( + init_cuda.create_opaque_array( + OpaqueArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + ) + ) as array, + use_context(init_cuda, green_ctx), + pytest.raises(ValueError, match="resource is not compatible with this Device object"), + ): + init_cuda.create_texture_object(resource=ResourceDescriptor.from_opaque_array(array)) def test_close_while_current_raises(self, init_cuda, green_ctx): """close() on a current context raises — test via set_current.""" diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py index 2ee783016f2..083bdcbee8a 100644 --- a/cuda_core/tests/test_launcher.py +++ b/cuda_core/tests/test_launcher.py @@ -24,7 +24,7 @@ StreamOptions, launch, ) -from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource +from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource from cuda.core._utils.cuda_utils import CUDAError from cuda.core.typing import ObjectCodeFormatType, SourceCodeType diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index b6b219e002e..769d780b36b 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -2157,16 +2157,15 @@ def test_legacy_pinned_allocate_zero_size(init_cuda): assert int(buf.handle) == 0 -def test_legacy_pinned_device_id_raises(): - """LegacyPinnedMemoryResource.device_id raises; pinned memory is not bound to a GPU.""" +def test_legacy_pinned_device_id_is_not_applicable(): + """LegacyPinnedMemoryResource.device_id is -1, as documented for memory not bound to a device.""" mr = LegacyPinnedMemoryResource() - with pytest.raises(RuntimeError, match="not bound to any GPU"): - _ = mr.device_id + assert mr.device_id == -1 def test_synchronous_memory_resource_basic(init_cuda): """_SynchronousMemoryResource exercises properties and allocate paths (zero, non-zero, with-stream).""" - from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource + from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource dev = Device() mr = _SynchronousMemoryResource(dev.device_id) @@ -2199,7 +2198,7 @@ def test_synchronous_memory_resource_basic(init_cuda): def test_synchronous_memory_resource_deallocate_accepts_stream(init_cuda): """_SynchronousMemoryResource.deallocate accepts an explicit stream.""" - from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource + from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource dev = Device() mr = _SynchronousMemoryResource(dev.device_id) @@ -2212,7 +2211,7 @@ def test_synchronous_memory_resource_deallocate_accepts_stream(init_cuda): @pytest.mark.agent_authored(model="gpt-5.6") def test_synchronous_memory_resource_uses_its_context(device_x2): """Synchronous allocation targets its stored context and restores the current one.""" - from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource + from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource alloc_dev, current_dev = device_x2 alloc_dev.set_current() @@ -2249,7 +2248,7 @@ def test_synchronous_memory_resource_uses_its_context(device_x2): @pytest.mark.agent_authored(model="gpt-5.6") def test_synchronous_memory_resource_restores_context_after_failure(device_x2): """A failed synchronous allocation restores the context that was current.""" - from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource + from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource alloc_dev, current_dev = device_x2 alloc_dev.set_current() @@ -2263,6 +2262,47 @@ def test_synchronous_memory_resource_restores_context_after_failure(device_x2): assert current_context_handle() == current_context +@pytest.mark.agent_authored(model="claude-sonnet-5") +def test_synchronous_memory_resource_default_stream_deallocates_in_own_context(device_x2, capsys): + """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 + + alloc_dev, current_dev = device_x2 + alloc_dev.set_current() + mr = _SynchronousMemoryResource(alloc_dev.device_id, alloc_dev.context) + + current_dev.set_current() + current_context = current_context_handle() + + 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 + assert current_context_handle() == current_context + assert capsys.readouterr().err == "" + + +@pytest.mark.agent_authored(model="claude-sonnet-5") +def test_synchronous_memory_resource_allocate_without_current_context(device_x2, capsys): + """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 + + alloc_dev, current_dev = device_x2 + alloc_dev.set_current() + mr = _SynchronousMemoryResource(alloc_dev.device_id, alloc_dev.context) + current_dev.set_current() + + with no_current_context(): + 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"), [ diff --git a/cuda_core/tests/test_texture_surface.py b/cuda_core/tests/test_texture_surface.py index 254eb2faf4d..42ae716e8f3 100644 --- a/cuda_core/tests/test_texture_surface.py +++ b/cuda_core/tests/test_texture_surface.py @@ -10,6 +10,7 @@ import cuda.core from cuda.core import ( Device, + LegacyPinnedMemoryResource, ) from cuda.core.texture import ( MipmappedArrayOptions, @@ -54,44 +55,37 @@ def test_texture_resources_target_receiver_context(device_x2): dev1.set_current() ctx1_handle = current_context_handle() - array = dev0.create_opaque_array( + with dev0.create_opaque_array( OpaqueArrayOptions( shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True, ) - ) - assert array.device == dev0 - assert current_context_handle() == ctx1_handle - - mipmap = dev0.create_mipmapped_array( - MipmappedArrayOptions( - shape=(8, 8), - format=ArrayFormatType.UINT8, - num_channels=4, - num_levels=2, - ) - ) - assert mipmap.device == dev0 - assert current_context_handle() == ctx1_handle - - level = mipmap.get_level(0) - assert level.device == dev0 - assert current_context_handle() == ctx1_handle - - resource = ResourceDescriptor.from_opaque_array(array) - texture = dev0.create_texture_object(resource=resource, options=TextureObjectOptions()) - surface = dev0.create_surface_object(resource=resource) - assert texture.device == dev0 - assert surface.device == dev0 - assert current_context_handle() == ctx1_handle - - surface.close() - texture.close() - level.close() - mipmap.close() - array.close() + ) as array: + assert array.device == dev0 + assert current_context_handle() == ctx1_handle + with dev0.create_mipmapped_array( + MipmappedArrayOptions( + shape=(8, 8), + format=ArrayFormatType.UINT8, + num_channels=4, + num_levels=2, + ) + ) as mipmap: + assert mipmap.device == dev0 + assert current_context_handle() == ctx1_handle + with mipmap.get_level(0) as level: + assert level.device == dev0 + assert current_context_handle() == ctx1_handle + resource = ResourceDescriptor.from_opaque_array(array) + with ( + dev0.create_texture_object(resource=resource, options=TextureObjectOptions()) as texture, + dev0.create_surface_object(resource=resource) as surface, + ): + assert texture.device == dev0 + assert surface.device == dev0 + assert current_context_handle() == ctx1_handle assert current_context_handle() == ctx1_handle assert int(ctx0.handle) != ctx1_handle @@ -101,53 +95,57 @@ def test_texture_resources_restore_no_current_context(deinit_cuda): device = Device(0) device.set_current() - array = texture = None - try: - with no_current_context(): - array = device.create_opaque_array( + with no_current_context(): + with ( + device.create_opaque_array( OpaqueArrayOptions( shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, ) - ) - texture = device.create_texture_object(resource=ResourceDescriptor.from_opaque_array(array)) + ) as array, + device.create_texture_object(resource=ResourceDescriptor.from_opaque_array(array)) as texture, + ): + assert array.device == device + assert texture.device == device assert current_context_handle() == 0 - - texture.close() - texture = None - array.close() - array = None - assert current_context_handle() == 0 - finally: - if texture is not None: - texture.close() - if array is not None: - array.close() + assert current_context_handle() == 0 @pytest.mark.agent_authored(model="gpt-5.6") def test_texture_creation_rejects_mismatched_receiver(device_x2): dev0, dev1 = device_x2 dev0.set_current() - array = dev0.create_opaque_array( + with dev0.create_opaque_array( OpaqueArrayOptions( shape=(8, 8), format=ArrayFormatType.UINT8, num_channels=4, is_surface_load_store=True, ) - ) - resource = ResourceDescriptor.from_opaque_array(array) - dev1.set_current() - ctx1_handle = current_context_handle() - - with pytest.raises(ValueError, match="resource belongs to device 0"): - dev1.create_texture_object(resource=resource) - with pytest.raises(ValueError, match="resource belongs to device 0"): - dev1.create_surface_object(resource=resource) - assert current_context_handle() == ctx1_handle - array.close() + ) as array: + resource = ResourceDescriptor.from_opaque_array(array) + dev1.set_current() + ctx1_handle = current_context_handle() + with pytest.raises(ValueError, match="resource belongs to device 0"): + dev1.create_texture_object(resource=resource) + with pytest.raises(ValueError, match="resource belongs to device 0"): + dev1.create_surface_object(resource=resource) + assert current_context_handle() == ctx1_handle + + +@pytest.mark.agent_authored(model="claude-sonnet-5") +def test_texture_linear_accepts_pinned_buffer(init_cuda): + """Pinned memory is device-accessible and not bound to any device, so a + pinned buffer is a valid linear texture backing whose device check is + skipped (device_id == -1) rather than failed (#2311).""" + mr = LegacyPinnedMemoryResource() + with mr.allocate(256) as buf: + assert buf.device_id == -1 + + resource = ResourceDescriptor.from_linear(buf, format=ArrayFormatType.UINT8, num_channels=1) + with init_cuda.create_texture_object(resource=resource) as texture: + assert texture.device == init_cuda def test_array_2d_create_and_properties(init_cuda): From 85986ee757619bfc6a0b856d6b7a3d3540c5231e Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Fri, 4 Sep 2026 12:17:44 -0700 Subject: [PATCH 3/9] cuda.core tests: fix driver-call plumbing in the bound-context helper handle_return() takes the whole result tuple, and cuCtxGetDevice() returns a CUdevice that never compares equal to an int; both made the cross-context event-record check raise instead of run (or skip) as intended. Co-Authored-By: Claude Fable 5.1 --- cuda_core/tests/helpers/contexts.py | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/cuda_core/tests/helpers/contexts.py b/cuda_core/tests/helpers/contexts.py index a2227c0f737..7ee01bb255f 100644 --- a/cuda_core/tests/helpers/contexts.py +++ b/cuda_core/tests/helpers/contexts.py @@ -26,8 +26,7 @@ def _assert_event_record_rejected_from_ambient_context(event): creates the probe stream in whatever context is currently ambient, so this is a live check that ``event`` was not actually recorded from there. """ - err, ambient_stream = driver.cuStreamCreate(0) - handle_return(err) + ambient_stream = handle_return(driver.cuStreamCreate(0)) try: (record_status,) = driver.cuEventRecord(event.handle, ambient_stream) assert record_status == driver.CUresult.CUDA_ERROR_INVALID_HANDLE, ( @@ -68,7 +67,7 @@ def assert_device_operations_use_bound_context(device): # a green context vs. the primary context) may resolve to the same # underlying device context for this check, so skip rather than # assert unverified driver behavior. - if ambient_context_handle and handle_return(driver.cuCtxGetDevice()) != device.device_id: + if ambient_context_handle and int(handle_return(driver.cuCtxGetDevice())) != device.device_id: _assert_event_record_rejected_from_ambient_context(event) builder = device.create_graph_builder() From cdb75ceda82cc86629127343333392d3238ebc58 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Fri, 4 Sep 2026 12:24:35 -0700 Subject: [PATCH 4/9] cuda.core: regenerate stub with stubgen-pyx 0.2.22 Co-Authored-By: Claude Fable 5.1 --- cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi b/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi index 4dfc17f10f2..73136b896b9 100644 --- a/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi +++ b/cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyi @@ -1,4 +1,4 @@ -# This file was generated by stubgen-pyx v0.2.19 from cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx +# This file was generated by stubgen-pyx v0.2.22 from cuda_core/cuda/core/_memory/_synchronous_memory_resource.pyx from cuda.core._context import Context from cuda.core._memory._buffer import Buffer, MemoryResource From fb683822371eaf9054e3bae4615a7df4e85ffa25 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Fri, 4 Sep 2026 12:58:31 -0700 Subject: [PATCH 5/9] cuda.core: create stream-ordering events in the recorded stream's context The uniform empty-context guard added for review item 6 broke create_event_handle_noctx, which relied on an empty handle meaning "current context". That helper was itself a #2311-class bug: Stream.wait(stream) and the foreign-array/tensor import paths created their temporary ordering event in whatever context was current, and cuEventRecord rejects an event from a different context than the stream it is recorded on, so cross-device waits failed unless the right device happened to be current. Replace it with create_event_handle_for_stream, which resolves the stream's owning context via cuStreamGetCtx and creates the event there. Callers now check the returned handle and surface the real creation error. With that, every creation helper requires a context and no exception remains. Co-Authored-By: Claude Fable 5.1 --- cuda_core/cuda/core/_cpp/resource_handles.cpp | 20 +++++++++++++-- cuda_core/cuda/core/_cpp/resource_handles.hpp | 10 +++++--- cuda_core/cuda/core/_memoryview.pyx | 10 ++++++-- cuda_core/cuda/core/_resource_handles.pxd | 3 ++- cuda_core/cuda/core/_resource_handles.pyx | 8 +++--- cuda_core/cuda/core/_stream.pyx | 11 +++++--- cuda_core/cuda/core/_tensor_bridge.pyx | 11 +++++--- cuda_core/docs/source/release/1.3.0-notes.rst | 8 ++++++ cuda_core/tests/test_stream.py | 25 +++++++++++++++++++ 9 files changed, 89 insertions(+), 17 deletions(-) diff --git a/cuda_core/cuda/core/_cpp/resource_handles.cpp b/cuda_core/cuda/core/_cpp/resource_handles.cpp index ec9c033829e..46f7b019379 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.cpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.cpp @@ -54,6 +54,7 @@ decltype(&cuGreenCtxStreamCreate) p_cuGreenCtxStreamCreate = nullptr; decltype(&cuStreamCreateWithPriority) p_cuStreamCreateWithPriority = nullptr; decltype(&cuStreamDestroy) p_cuStreamDestroy = nullptr; +decltype(&cuStreamGetCtx) p_cuStreamGetCtx = nullptr; decltype(&cuEventCreate) p_cuEventCreate = nullptr; decltype(&cuEventDestroy) p_cuEventDestroy = nullptr; @@ -1099,8 +1100,23 @@ EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, return h; } -EventHandle create_event_handle_noctx(unsigned int flags) { - return create_event_handle(ContextHandle{}, flags, false, false, false, -1); +EventHandle create_event_handle_for_stream(CUstream stream, unsigned int flags) { + // Resolve the stream's owning context (for default-stream tokens this is + // the current context, per cuStreamGetCtx) and create the event there, so + // it can be recorded on `stream` no matter which context is current. + CUcontext ctx = nullptr; + { + GILReleaseGuard gil; + err = p_cuStreamGetCtx(stream, &ctx); + } + if (err != CUDA_SUCCESS) { + return {}; + } + if (!ctx) { + err = CUDA_ERROR_INVALID_CONTEXT; + return {}; + } + return create_event_handle(create_context_handle_ref(ctx), flags, false, false, false, -1); } EventHandle create_event_handle_ref(CUevent event) { diff --git a/cuda_core/cuda/core/_cpp/resource_handles.hpp b/cuda_core/cuda/core/_cpp/resource_handles.hpp index 525100cb0db..419710ea0d9 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.hpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.hpp @@ -82,6 +82,7 @@ extern decltype(&cuGreenCtxStreamCreate) p_cuGreenCtxStreamCreate; extern decltype(&cuStreamCreateWithPriority) p_cuStreamCreateWithPriority; extern decltype(&cuStreamDestroy) p_cuStreamDestroy; +extern decltype(&cuStreamGetCtx) p_cuStreamGetCtx; extern decltype(&cuEventCreate) p_cuEventCreate; extern decltype(&cuEventDestroy) p_cuEventDestroy; @@ -324,11 +325,14 @@ EventHandle create_event_handle(const ContextHandle& h_ctx, unsigned int flags, bool timing_enabled, bool is_blocking_sync, bool ipc_enabled, int device_id); -// Create an owning event handle without context dependency. -// Use for temporary events that are created and destroyed in the same scope. +// Create an owning event in the context that owns `stream`, so it can be +// recorded on that stream regardless of which context is current. Default- +// stream tokens resolve to the current context (cuStreamGetCtx semantics). +// Use for temporary ordering events that are created and destroyed in the +// same scope; the handle carries no device id. // When the last reference is released, cuEventDestroy is called automatically. // Returns empty handle on error (caller must check). -EventHandle create_event_handle_noctx(unsigned int flags); +EventHandle create_event_handle_for_stream(CUstream stream, unsigned int flags); // Create an owning event handle from an IPC handle. // The originating process owns the event and its context. diff --git a/cuda_core/cuda/core/_memoryview.pyx b/cuda_core/cuda/core/_memoryview.pyx index 6b287f68c1a..c4a49d76946 100644 --- a/cuda_core/cuda/core/_memoryview.pyx +++ b/cuda_core/cuda/core/_memoryview.pyx @@ -26,8 +26,9 @@ import numpy from cuda.bindings cimport cydriver from cuda.core._resource_handles cimport ( EventHandle, - create_event_handle_noctx, + create_event_handle_for_stream, as_cu, + get_last_error, ) from cuda.core._utils.cuda_utils import handle_return, driver @@ -1227,7 +1228,12 @@ cpdef StridedMemoryView view_as_cai(obj, stream_ptr, view=None): # establish stream order if producer_s != consumer_s: with nogil: - h_event = create_event_handle_noctx(cydriver.CUevent_flags.CU_EVENT_DISABLE_TIMING) + # The event must belong to the producer stream's context to + # be recorded on it, whatever context is current here. + h_event = create_event_handle_for_stream( + producer_s, cydriver.CUevent_flags.CU_EVENT_DISABLE_TIMING) + if not h_event: + HANDLE_RETURN(get_last_error()) HANDLE_RETURN(cydriver.cuEventRecord( as_cu(h_event), producer_s)) HANDLE_RETURN(cydriver.cuStreamWaitEvent( diff --git a/cuda_core/cuda/core/_resource_handles.pxd b/cuda_core/cuda/core/_resource_handles.pxd index f8dea004831..fe075b6414b 100644 --- a/cuda_core/cuda/core/_resource_handles.pxd +++ b/cuda_core/cuda/core/_resource_handles.pxd @@ -203,7 +203,8 @@ cdef EventHandle create_event_handle( const ContextHandle& h_ctx, unsigned int flags, bint timing_enabled, bint is_blocking_sync, bint ipc_enabled, int device_id) except+ nogil -cdef EventHandle create_event_handle_noctx(unsigned int flags) except+ nogil +cdef EventHandle create_event_handle_for_stream( + cydriver.CUstream stream, unsigned int flags) except+ nogil cdef EventHandle create_event_handle_ref(cydriver.CUevent event) except+ nogil cdef EventHandle create_event_handle_ipc( const cydriver.CUipcEventHandle& ipc_handle, bint is_blocking_sync) except+ nogil diff --git a/cuda_core/cuda/core/_resource_handles.pyx b/cuda_core/cuda/core/_resource_handles.pyx index da4a5a20ff3..0f8d6e15cde 100644 --- a/cuda_core/cuda/core/_resource_handles.pyx +++ b/cuda_core/cuda/core/_resource_handles.pyx @@ -81,8 +81,8 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": const ContextHandle& h_ctx, unsigned int flags, bint timing_enabled, bint is_blocking_sync, bint ipc_enabled, int device_id) except+ nogil - EventHandle create_event_handle_noctx "cuda_core::create_event_handle_noctx" ( - unsigned int flags) except+ nogil + EventHandle create_event_handle_for_stream "cuda_core::create_event_handle_for_stream" ( + cydriver.CUstream stream, unsigned int flags) except+ nogil EventHandle create_event_handle_ref "cuda_core::create_event_handle_ref" ( cydriver.CUevent event) except+ nogil EventHandle create_event_handle_ipc "cuda_core::create_event_handle_ipc" ( @@ -332,6 +332,7 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": # Stream void* p_cuStreamCreateWithPriority "reinterpret_cast(cuda_core::p_cuStreamCreateWithPriority)" void* p_cuStreamDestroy "reinterpret_cast(cuda_core::p_cuStreamDestroy)" + void* p_cuStreamGetCtx "reinterpret_cast(cuda_core::p_cuStreamGetCtx)" # Event void* p_cuEventCreate "reinterpret_cast(cuda_core::p_cuEventCreate)" @@ -434,7 +435,7 @@ cdef void _init_driver_fn_pointers() noexcept: global p_cuCtxSetCurrent, p_cuCtxSynchronize, p_cuCtxGetStreamPriorityRange global p_cuGreenCtxCreate, p_cuGreenCtxDestroy, p_cuCtxFromGreenCtx global p_cuDevResourceGenerateDesc, p_cuGreenCtxStreamCreate - global p_cuStreamCreateWithPriority, p_cuStreamDestroy + global p_cuStreamCreateWithPriority, p_cuStreamDestroy, p_cuStreamGetCtx global p_cuEventCreate, p_cuEventDestroy, p_cuIpcOpenEventHandle global p_cuDeviceGetCount global p_cuMemPoolSetAccess, p_cuMemPoolDestroy, p_cuMemPoolCreate @@ -477,6 +478,7 @@ cdef void _init_driver_fn_pointers() noexcept: # Stream p_cuStreamCreateWithPriority = _get_driver_fn("cuStreamCreateWithPriority") p_cuStreamDestroy = _get_driver_fn("cuStreamDestroy") + p_cuStreamGetCtx = _get_driver_fn("cuStreamGetCtx") # Event p_cuEventCreate = _get_driver_fn("cuEventCreate") diff --git a/cuda_core/cuda/core/_stream.pyx b/cuda_core/cuda/core/_stream.pyx index b5f1bc6a167..e662d67c87f 100644 --- a/cuda_core/cuda/core/_stream.pyx +++ b/cuda_core/cuda/core/_stream.pyx @@ -32,7 +32,7 @@ from cuda.core._resource_handles cimport ( EventHandle, StreamHandle, create_context_handle_ref, - create_event_handle_noctx, + create_event_handle_for_stream, create_stream_handle, create_stream_handle_with_owner, context_get_stream_priority_range, @@ -367,9 +367,14 @@ cdef class Stream: f" got {type(event_or_stream)}" ) from e - # Wait on stream via temporary event + # Wait on stream via a temporary event created in that stream's own + # context; an event from the current context would be rejected by + # cuEventRecord when the streams live on different devices. with nogil: - h_event = create_event_handle_noctx(cydriver.CUevent_flags.CU_EVENT_DISABLE_TIMING) + h_event = create_event_handle_for_stream( + as_cu(stream._h_stream), cydriver.CUevent_flags.CU_EVENT_DISABLE_TIMING) + if not h_event: + HANDLE_RETURN(get_last_error()) HANDLE_RETURN(cydriver.cuEventRecord(as_cu(h_event), as_cu(stream._h_stream))) # TODO: support flags other than 0? HANDLE_RETURN(cydriver.cuStreamWaitEvent(as_cu(self._h_stream), as_cu(h_event), 0)) diff --git a/cuda_core/cuda/core/_tensor_bridge.pyx b/cuda_core/cuda/core/_tensor_bridge.pyx index c7a6743213c..ae7a6794507 100644 --- a/cuda_core/cuda/core/_tensor_bridge.pyx +++ b/cuda_core/cuda/core/_tensor_bridge.pyx @@ -56,8 +56,9 @@ from cuda.core._layout cimport _StridedLayout from cuda.bindings cimport cydriver from cuda.core._resource_handles cimport ( EventHandle, - create_event_handle_noctx, + create_event_handle_for_stream, as_cu, + get_last_error, ) from cuda.core._utils.cuda_utils cimport HANDLE_RETURN @@ -318,8 +319,12 @@ cpdef int sync_torch_stream(int32_t device_index, b"aoti_torch_get_current_cuda_stream") if producer_s != consumer_s: with nogil: - h_event = create_event_handle_noctx( - cydriver.CUevent_flags.CU_EVENT_DISABLE_TIMING) + # The event must belong to the producer stream's context to be + # recorded on it, whatever context is current here. + h_event = create_event_handle_for_stream( + producer_s, cydriver.CUevent_flags.CU_EVENT_DISABLE_TIMING) + if not h_event: + HANDLE_RETURN(get_last_error()) HANDLE_RETURN(cydriver.cuEventRecord( as_cu(h_event), producer_s)) HANDLE_RETURN(cydriver.cuStreamWaitEvent( 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 a76fb8409a2..37ff06b34e5 100644 --- a/cuda_core/docs/source/release/1.3.0-notes.rst +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -21,6 +21,14 @@ Fixes and enhancements :class:`Device` object. (`#2311 `__) +- :meth:`Stream.wait` given a stream now works when that stream belongs to a + device that is not current. The temporary ordering event is created in the + waited-on stream's context rather than the current one, which + ``cuEventRecord`` rejects when the two differ. The stream ordering applied + when importing foreign arrays and tensors uses the producer stream's context + the same way. + (`#2311 `__) + - :attr:`LegacyPinnedMemoryResource.device_id` now returns ``-1``, as documented for memory that is not bound to a device and as :class:`PinnedMemoryResource` already does, instead of raising diff --git a/cuda_core/tests/test_stream.py b/cuda_core/tests/test_stream.py index bd909734e09..02e817cf85b 100644 --- a/cuda_core/tests/test_stream.py +++ b/cuda_core/tests/test_stream.py @@ -74,6 +74,31 @@ def test_stream_wait_event(init_cuda): s2.sync() +@pytest.mark.agent_authored(model="claude-fable-5-1") +def test_stream_wait_stream_on_other_device(device_x2): + """Stream.wait(other_stream) must work when the streams live on different + devices and neither device is necessarily current: the temporary ordering + event has to be created in the *recorded* stream's context, since + cuEventRecord rejects an event from another context (#2311).""" + from helpers.contexts import current_context_handle + + dev0, dev1 = device_x2 + dev0.set_current() + s0 = dev0.create_stream() + dev1.set_current() + s1 = dev1.create_stream() + ambient = current_context_handle() + try: + s1.wait(s0) # dev1 current: s0's device is not current + s0.wait(s1) # dev1 current: self's device is not current + s0.sync() + s1.sync() + assert current_context_handle() == ambient + finally: + s0.close() + s1.close() + + def test_stream_wait_invalid_event(init_cuda): stream = Device().create_stream(options=StreamOptions()) with pytest.raises(ValueError): From fcc88d5565d32ad88d189084644f549452392052 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Wed, 2 Sep 2026 17:14:12 -0700 Subject: [PATCH 6/9] cuda.core: define the error handling policy and report failures that cannot be raised Write down how cuda.core handles CUDA failures (docs/source/error_handling.rst for users, a "Failure handling" section in AGENTS.md and _cpp/DESIGN.md for contributors) and bring the code into line with it: - Add cuda.core.CUDAWarning, emitted for CUDA errors that cannot be raised (destructors, CUDA callbacks, cleanup after an earlier failure). The C++ handle layer reports through one helper that uses the Python warnings machinery when the interpreter is usable, delivers an escalated warning as an unraisable exception, and falls back to stderr otherwise. CUDA_ERROR_DEINITIALIZED is not reported. - Wrap every destroy call made from a deleter (pw_*) so its failure is reported instead of discarded, including memory pools, green contexts, graphs, graph execs, graphics resources, the linker, user objects, the NVRTC/NVVM/nvJitLink handles and file descriptors; release the GIL around the compiler-handle destroys like the CUDA ones. - When the caller's context cannot be restored after a successful operation, undo the creation and raise a CUDAError that says which context is current; report the same failure as a warning in deleters; report a skipped context-sensitive undo instead of leaking silently. - Add context_get_device and graph_node_set_params so Stream_get_ctx_device and _set_definition_node_params stop hand-rolling cuCtxPush/Pop/SetCurrent. The node update now publishes its attachment before raising a restoration failure, closing a window that left the node referencing released owners. - Device.set_current(ctx) switches with a single cuCtxSetCurrent, so a failure leaves the previous context current and the call works without one. - Report failed cuStreamEndCapture in GraphBuilder.__dealloc__ and failed child-graph rollbacks; warn from _mr_dealloc_callback instead of printing. - Add a test hook that makes the next context restoration fail, tests for the policy, and release notes for 1.3.0. Co-Authored-By: Claude Fable 5.1 --- cuda_core/AGENTS.md | 56 +++ cuda_core/cuda/core/__init__.py | 2 + cuda_core/cuda/core/_cpp/DESIGN.md | 31 ++ cuda_core/cuda/core/_cpp/resource_handles.cpp | 334 ++++++++++++++---- cuda_core/cuda/core/_cpp/resource_handles.hpp | 52 +++ cuda_core/cuda/core/_device.pyx | 8 +- cuda_core/cuda/core/_memory/_buffer.pyx | 23 +- cuda_core/cuda/core/_resource_handles.pxd | 16 + cuda_core/cuda/core/_resource_handles.pyi | 8 + cuda_core/cuda/core/_resource_handles.pyx | 35 ++ cuda_core/cuda/core/_stream.pyx | 16 +- cuda_core/cuda/core/_utils/cuda_utils.pyi | 20 ++ cuda_core/cuda/core/_utils/cuda_utils.pyx | 45 ++- cuda_core/cuda/core/graph/_graph_builder.pyx | 15 +- cuda_core/cuda/core/graph/_graph_node.pyx | 7 + cuda_core/cuda/core/graph/_subclasses.pyx | 30 +- cuda_core/docs/source/api.rst | 15 + cuda_core/docs/source/error_handling.rst | 127 +++++++ cuda_core/docs/source/index.rst | 1 + cuda_core/docs/source/release/1.3.0-notes.rst | 51 +++ cuda_core/tests/helpers/contexts.py | 18 + cuda_core/tests/test_error_handling.py | 207 +++++++++++ cuda_core/tests/test_memory.py | 48 +-- 23 files changed, 1033 insertions(+), 132 deletions(-) create mode 100644 cuda_core/docs/source/error_handling.rst create mode 100644 cuda_core/tests/test_error_handling.py diff --git a/cuda_core/AGENTS.md b/cuda_core/AGENTS.md index 9d80ab74aaa..1e7c8da1077 100644 --- a/cuda_core/AGENTS.md +++ b/cuda_core/AGENTS.md @@ -101,6 +101,62 @@ 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__`, CUDA callbacks and cleanup after a failure 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. +- **Rollback failure**: the original exception propagates; the failed rollback + is reported out of band (or chained with `raise ... from` when a second + exception must be raised). Bare `except:` is acceptable only for + rollback-then-`raise` blocks. +- **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`). +- **Aborting**: `std::abort` (or any process termination) is reserved for an + internal invariant violation where continuing could corrupt memory or produce + silently wrong results *and* no leak-based fallback exists. A failed CUDA + call, including a failed context restoration, never qualifies: raise or + report instead. There is currently no such path; if one is ever needed it + must go through a single helper that writes a diagnostic (call, CUDA error, + invariant, "please report") to stderr before aborting, must never trigger + during interpreter finalization or for driver-shutdown errors, and must be + called out in the docs and release notes. 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..149fe327f12 100644 --- a/cuda_core/cuda/core/__init__.py +++ b/cuda_core/cuda/core/__init__.py @@ -102,8 +102,10 @@ 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 CUDAWarning __all__ = [ + "CUDAWarning", *_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..8e1a55a34a3 100644 --- a/cuda_core/cuda/core/_cpp/DESIGN.md +++ b/cuda_core/cuda/core/_cpp/DESIGN.md @@ -275,6 +275,37 @@ Related functions: - `peek_last_error()`: Returns the error without clearing it - `clear_last_error()`: Clears the error state +Some functions return a `CUresult` directly instead of a handle (for example +`context_synchronize`, `context_get_device`, `graph_node_set_params`). Their +callers `HANDLE_RETURN` the value. + +### 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. When restoration fails after the operation succeeded, the +creation is undone and the restoration status is returned; the helper also +records a thread-local detail (`take_last_error_detail()`) that the Cython error +path appends to the raised `CUDAError`, so the user learns that the caller's +context was not restored and which context is current. When both the operation +and the restoration fail, the operation status is returned and the restoration +failure is reported out of band. Tests inject restoration failures with +`set_context_restore_fault_for_testing()`. + +### Reporting from non-propagating paths + +Deleters, CUDA callbacks and cleanup-after-failure cannot raise. They report +through `report_cuda_error()` / `report_message()` (the `pw_*` wrappers +decorate destroy calls with it), 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 terminates +the process; see `docs/source/error_handling.rst` and the "Failure handling" +section of `AGENTS.md` for the policy. + ## 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 46f7b019379..75d2d27bc7d 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,7 +199,120 @@ 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 (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[256] = {0}; +thread_local char taken_error_detail[256] = {0}; + +// 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)); + } +} + +} // 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); +} + +// 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); +} + +const char* take_last_error_detail() noexcept { + if (!last_error_detail[0]) { + return nullptr; + } + std::memcpy(taken_error_detail, last_error_detail, sizeof(taken_error_detail)); + last_error_detail[0] = 0; + return taken_error_detail; +} + +void clear_last_error_detail() noexcept { + last_error_detail[0] = 0; +} + +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 @@ -205,6 +320,7 @@ void warn_on_cuda_error(const char* operation, CUresult status, const char* deta 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,17 +336,49 @@ CUresult enter_context(const ContextHandle& h_context, CUcontext* previous, int* return status; } -// Restore the previous context and preserve an earlier operation error. +// 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; + } + GILReleaseGuard gil; + return p_cuCtxSetCurrent(previous); +} + +// Record why the CUresult about to be returned should be explained further +// when it is raised as a CUDAError: the caller's context was not restored. +void note_context_not_restored(CUcontext previous) noexcept { + CUcontext current = nullptr; + if (p_cuCtxGetCurrent(¤t) != CUDA_SUCCESS) { + current = nullptr; + } + std::snprintf(last_error_detail, sizeof(last_error_detail), + "the calling thread's CUDA context (%#llx) could not be restored; " + "context %#llx is now current. Call Device.set_current() before issuing " + "further CUDA work on this thread", + static_cast(reinterpret_cast(previous)), + static_cast(reinterpret_cast(current))); +} + +// Restore the previous context and preserve an earlier operation error. The +// operation error, if any, is returned; a restoration failure is then reported +// out of band. Otherwise the restoration status is returned, annotated for the +// eventual CUDAError. 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); + CUresult restore_status = changed ? restore_context(previous) : CUDA_SUCCESS; + if (restore_status == CUDA_SUCCESS) { + return operation_status; } - if (operation_status != CUDA_SUCCESS && restore_status != CUDA_SUCCESS) { - warn_on_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restore_status); + if (operation_status != CUDA_SUCCESS) { + report_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restore_status); + return operation_status; } - return operation_status != CUDA_SUCCESS ? operation_status : restore_status; + note_context_not_restored(previous); + return restore_status; } // Require a callable to be invocable without throwing. @@ -257,12 +405,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 +460,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,32 +469,6 @@ 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 @@ -360,39 +481,50 @@ CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, int changed = 0; CUresult status = enter_context(h_context, &previous, &changed); if (status != CUDA_SUCCESS) { - warn_on_cuda_error(name, status, + report_cuda_error(name, status, "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); + report_cuda_error(name, status); } } 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"); + report_cuda_error(name, restore, "failed while restoring the caller's context"); } return status; } #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. 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); - } + auto operator()(Args&&... args) const noexcept { + auto status = Function(std::forward(args)...); + report(status); return status; } private: + void report(CUresult status) const noexcept { + report_cuda_error(operation_, status); + } + + template + void report(Status status) const noexcept { + if (static_cast(status) != 0) { + report_status_code(operation_, static_cast(status)); + } + } + const char* operation_; }; @@ -405,6 +537,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 +570,52 @@ 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, a restoration failure is reported +// out of band 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; + } + if (status != CUDA_SUCCESS) { + report_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restored); + return status; + } + note_context_not_restored(previous); + *restore_status = restored; + return status; +} + // ============================================================================ // CUDA user-object deferred cleanup // @@ -767,7 +957,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 +1377,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; } ); @@ -1353,7 +1543,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; } ); @@ -1909,7 +2099,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); } } @@ -1955,7 +2145,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; @@ -2187,7 +2377,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; } } @@ -2318,7 +2508,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(); @@ -2338,7 +2528,7 @@ GraphExecHandle make_graph_exec_handle( ~RawGraphExecGuard() noexcept { if (resource) { GILReleaseGuard gil; - p_cuGraphExecDestroy(resource); + pw_cuGraphExecDestroy(resource); } retry_deferred_cleanup(); } @@ -2360,7 +2550,8 @@ struct ExecAttachmentStaging { ExecAttachments* accumulator = nullptr; ~ExecAttachmentStaging() noexcept { - release(); + report_cuda_error("cuGraphReleaseUserObject", release(), + "failed while dropping a staged graph attachment"); } CUresult release() noexcept { @@ -2406,7 +2597,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; } } @@ -2676,7 +2867,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; } ); @@ -2699,8 +2890,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; } ); @@ -2730,7 +2923,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; } @@ -2761,7 +2955,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; } @@ -2789,9 +2984,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; } @@ -2814,7 +3009,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 } diff --git a/cuda_core/cuda/core/_cpp/resource_handles.hpp b/cuda_core/cuda/core/_cpp/resource_handles.hpp index 419710ea0d9..069e3ec8319 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.hpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.hpp @@ -57,6 +57,41 @@ 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, CUDA callbacks, __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. See docs/source/error_handling.rst. +// ============================================================================ + +// 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; + +// 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 appends it to the raised CUDAError. Thread-local; take_ returns +// the detail (valid until the next take on this thread) and clears it, or +// nullptr when none is recorded. +const char* take_last_error_detail() 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 +108,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 +300,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 2484ad82b00..a24035c68d4 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,25 +60,33 @@ 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: - print( - "Warning: no deallocation stream was recorded; falling back to " + warnings.warn( + "no deallocation stream was recorded; falling back to " "the default stream for mr.deallocate() during Buffer " "destruction. This is an internal cuda-core error; please " "report it with your CUDA driver, CUDA Toolkit, and " "cuda-python versions.", - file=sys.stderr, + CUDAWarning, + stacklevel=2, ) stream = default_stream() else: 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() failed during Buffer destruction; 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..37f4bcb0435 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,16 @@ 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 const char* take_last_error_detail() 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 +195,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..c44bcd46a03 100644 --- a/cuda_core/cuda/core/_resource_handles.pyi +++ b/cuda_core/cuda/core/_resource_handles.pyi @@ -41,3 +41,11 @@ 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. + """ diff --git a/cuda_core/cuda/core/_resource_handles.pyx b/cuda_core/cuda/core/_resource_handles.pyx index 0f8d6e15cde..09692f9d9a2 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,19 @@ 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 + const char* take_last_error_detail "cuda_core::take_last_error_detail" () 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 +71,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 +342,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 +454,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 +491,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 +579,16 @@ 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) + # ============================================================================= # 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..b190e5967d0 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyi +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyi @@ -14,6 +14,26 @@ _fork_warning_checked = False class CUDAError(Exception): ... +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, or while a context switch is undone + after the requested operation already succeeded. 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 ComputeCapability(NamedTuple): diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyx b/cuda_core/cuda/core/_utils/cuda_utils.pyx index ce75746de56..38fc42027df 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyx +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyx @@ -27,6 +27,10 @@ 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 @@ -36,6 +40,31 @@ class CUDAError(Exception): pass +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, or while a context switch is undone + after the requested operation already succeeded. 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 @@ -135,19 +164,23 @@ 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() 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: + suffix = f" ({detail.decode()})" if detail != NULL else "" # 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()}") + raise CUDAError(f"{name.decode()}: {expl}{suffix}") + if desc_err != cydriver.CUresult.CUDA_SUCCESS: + raise CUDAError(f"{name.decode()}{suffix}") + raise CUDAError(f"{name.decode()}: {desc.decode()}{suffix}") 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..7033f90b239 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 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; the failed rollback is + # reported out of band (error handling policy). + report_cuda_error( + 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..831c0f30b69 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 report_cuda_error 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; the failed rollback is + # reported out of band (error handling policy). + report_cuda_error( + 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..140b69c0cc8 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -298,6 +298,21 @@ DLPack zero-copy interop. Data is moved in and out only by copying — use TextureObject SurfaceObject + +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/ + + CUDAWarning + :template: dataclass.rst OpaqueArrayOptions diff --git a/cuda_core/docs/source/error_handling.rst b/cuda_core/docs/source/error_handling.rst new file mode 100644 index 00000000000..a6761aadd24 --- /dev/null +++ b/cuda_core/docs/source/error_handling.rst @@ -0,0 +1,127 @@ +.. 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 (``CUDAError`` for driver and runtime failures) whose message contains +the CUDA error name and its description. 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. A failure that would otherwise be +hidden, for example because it occurred while another exception was already +propagating, is reported as described in the next section. + +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; +- callbacks invoked by CUDA; +- cleanup performed after an operation has already failed, such as rolling back + a partially built graph node or restoring the caller's CUDA context. + +A CUDA error in one of these places is reported as a :class:`CUDAWarning`. The +message names the failed driver call and the CUDA error. 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 whose message states that the caller's +context could not be restored and which context is now current. 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`` does not abort the process in response to a CUDA error, including +errors that cannot be raised, and including failures to restore the caller's +context. Aborting is reserved for an internal invariant violation where +continuing could corrupt memory, and no such code path exists in this release. 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 37ff06b34e5..3e94b19a77b 100644 --- a/cuda_core/docs/source/release/1.3.0-notes.rst +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -6,6 +6,18 @@ ``cuda.core`` 1.3.0 Release Notes ================================== +New features +------------ + +- 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. + Fixes and enhancements ---------------------- @@ -35,3 +47,42 @@ 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. ``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 states 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. 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``. + +- Failed rollbacks of a partially embedded child graph node and failed + ``cuStreamEndCapture`` calls made when a still-building + :class:`~graph.GraphBuilder` is garbage collected are 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..e716b001f66 --- /dev/null +++ b/cuda_core/tests/test_error_handling.py @@ -0,0 +1,207 @@ +# 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; and a failure to restore the +caller's context is raised (or reported) with an explanation rather than +swallowed or turned into a process abort. 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 +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._resource_handles import _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) + +thread_unsafe_context_fault = pytest.mark.thread_unsafe( + reason="injects a thread-local restoration fault and mutates the CUDA context stack" +) + + +@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) + + +@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, match="could not be restored") as excinfo: + dev.create_stream() + message = str(excinfo.value) + assert "CUDA_ERROR_INVALID_CONTEXT" in message + assert "Device.set_current()" in message + # 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, match="could not be restored"): + dev.sync() + 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()) + with no_context_with_restore_fault(): + with pytest.warns(CUDAWarning, match="restoring the caller's context") as records: + buf.close() + assert any("CUDA_ERROR_INVALID_CONTEXT" in str(record.message) for record in records) + + +@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, match="could not be restored"): + node.update(dst=replacement, value=0x22) + 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() diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index 769d780b36b..661e540c308 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -23,7 +23,7 @@ thread_unsafe_on_windows, ) 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, @@ -33,6 +33,7 @@ from cuda.core import ( Buffer, + CUDAWarning, Device, DeviceMemoryResource, DeviceMemoryResourceOptions, @@ -779,23 +780,25 @@ def test_from_handle_mr_explicit_stream_without_current_context(buffer_type): assert telemetry["deallocations"][-1]["stream"].handle == stream.handle +@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\(\) 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) @@ -806,16 +809,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() @@ -832,11 +836,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() @@ -856,8 +860,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() @@ -870,18 +875,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() @@ -899,7 +903,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. @@ -907,9 +912,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() From 2f45e0bcbefa24cea234181b9f9fadb0fcd05ac0 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Thu, 3 Sep 2026 07:07:33 -0700 Subject: [PATCH 7/9] cuda.core: keep the texture autosummary contiguous in api.rst The "Errors and warnings" section was inserted between the texture classes and the texture option dataclasses, which moved OpaqueArrayOptions, MipmappedArrayOptions and TextureObjectOptions under cuda.core in the docs index and failed test_api_docs_consistency on every CI platform. Place the section after the texture section instead. Co-Authored-By: Claude Fable 5.1 --- cuda_core/docs/source/api.rst | 26 +++++++++++++------------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/cuda_core/docs/source/api.rst b/cuda_core/docs/source/api.rst index 140b69c0cc8..9fb6a8eb718 100644 --- a/cuda_core/docs/source/api.rst +++ b/cuda_core/docs/source/api.rst @@ -298,6 +298,19 @@ DLPack zero-copy interop. Data is moved in and out only by copying — use TextureObject SurfaceObject + :template: dataclass.rst + + OpaqueArrayOptions + MipmappedArrayOptions + TextureObjectOptions + +The associated enumerations — +:class:`~cuda.core.typing.ArrayFormatType`, +:class:`~cuda.core.typing.AddressModeType`, +:class:`~cuda.core.typing.FilterModeType`, and +:class:`~cuda.core.typing.ReadModeType` — live in :mod:`cuda.core.typing` +alongside the other ``cuda.core`` enumerations. + Errors and warnings ------------------- @@ -313,19 +326,6 @@ reported as a warning instead. CUDAWarning - :template: dataclass.rst - - OpaqueArrayOptions - MipmappedArrayOptions - TextureObjectOptions - -The associated enumerations — -:class:`~cuda.core.typing.ArrayFormatType`, -:class:`~cuda.core.typing.AddressModeType`, -:class:`~cuda.core.typing.FilterModeType`, and -:class:`~cuda.core.typing.ReadModeType` — live in :mod:`cuda.core.typing` -alongside the other ``cuda.core`` enumerations. - CUDA process checkpointing -------------------------- From a9499d485982ac96cc5c4c50a78ca41e9f1fd758 Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Thu, 3 Sep 2026 09:39:53 -0700 Subject: [PATCH 8/9] cuda.core: attach secondary failures to the propagating exception as notes Review follow-ups on the error-handling policy: - A failure that happens while an exception is being raised is no longer reported out of band. When both an operation and the restoration of the caller's context fail, the operation's CUDAError is raised with the restoration failure attached; when only the restoration fails, its error is raised with the context explanation attached. The attachment is a PEP 678 note on Python 3.11+ and is appended to the message on 3.10. The thread-local detail is keyed to the status it was recorded for, so it cannot attach to an unrelated error if that status is never raised. - A failed rollback inside a Cython `except` block is attached to the exception being handled through note_or_report_cuda_error(), which falls back to a CUDAWarning when nothing is being handled or notes are unavailable. - Reporting stays reserved for destructors and CUDA callbacks; CUDAWarning's docstring and the docs say so. - DESIGN.md explains the two status conventions of the C++ layer (handle factories use thread-local err, everything else returns CUresult) and the abort-helper guidance in AGENTS.md asks for a faulthandler-style traceback. - Drop the release-relative "in this release" wording from the stable docs. Co-Authored-By: Claude Fable 5.1 --- cuda_core/AGENTS.md | 16 ++- cuda_core/cuda/core/_cpp/DESIGN.md | 39 ++++-- cuda_core/cuda/core/_cpp/resource_handles.cpp | 114 +++++++++++++----- cuda_core/cuda/core/_cpp/resource_handles.hpp | 16 ++- cuda_core/cuda/core/_resource_handles.pxd | 4 +- cuda_core/cuda/core/_resource_handles.pyi | 7 ++ cuda_core/cuda/core/_resource_handles.pyx | 20 ++- cuda_core/cuda/core/_utils/cuda_utils.pyi | 7 +- cuda_core/cuda/core/_utils/cuda_utils.pyx | 32 +++-- cuda_core/cuda/core/graph/_graph_builder.pyx | 8 +- cuda_core/cuda/core/graph/_graph_node.pyx | 8 +- cuda_core/docs/source/error_handling.rst | 26 ++-- cuda_core/docs/source/release/1.3.0-notes.rst | 19 +-- cuda_core/tests/test_error_handling.py | 101 ++++++++++++++-- 14 files changed, 314 insertions(+), 103 deletions(-) diff --git a/cuda_core/AGENTS.md b/cuda_core/AGENTS.md index 1e7c8da1077..fa1b2ce5c51 100644 --- a/cuda_core/AGENTS.md +++ b/cuda_core/AGENTS.md @@ -127,16 +127,18 @@ below are for contributors. Reviewers and agents should flag violations. 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__`, CUDA callbacks and cleanup after a failure report - through one channel, `report_cuda_error()` / `report_message()` in C++ (the + 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. - **Rollback failure**: the original exception propagates; the failed rollback - is reported out of band (or chained with `raise ... from` when a second - exception must be raised). Bare `except:` is acceptable only for - rollback-then-`raise` blocks. + is attached to it with `note_or_report_cuda_error()` (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`). @@ -146,7 +148,9 @@ below are for contributors. Reviewers and agents should flag violations. call, including a failed context restoration, never qualifies: raise or report instead. There is currently no such path; if one is ever needed it must go through a single helper that writes a diagnostic (call, CUDA error, - invariant, "please report") to stderr before aborting, must never trigger + invariant, "please report") and a Python traceback of all threads to stderr + before aborting (as `faulthandler` does, via the GIL-free + `_Py_DumpTracebackThreads`; no Python-object work), must never trigger during interpreter finalization or for driver-shutdown errors, and must be called out in the docs and release notes. An *implicit* abort (an exception escaping a `noexcept` function or a deleter, including `std::bad_alloc` from diff --git a/cuda_core/cuda/core/_cpp/DESIGN.md b/cuda_core/cuda/core/_cpp/DESIGN.md index 8e1a55a34a3..d0f2fdcc799 100644 --- a/cuda_core/cuda/core/_cpp/DESIGN.md +++ b/cuda_core/cuda/core/_cpp/DESIGN.md @@ -275,9 +275,16 @@ Related functions: - `peek_last_error()`: Returns the error without clearing it - `clear_last_error()`: Clears the error state -Some functions return a `CUresult` directly instead of a handle (for example -`context_synchronize`, `context_get_device`, `graph_node_set_params`). Their -callers `HANDLE_RETURN` the value. +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 @@ -285,19 +292,22 @@ 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. When restoration fails after the operation succeeded, the -creation is undone and the restoration status is returned; the helper also -records a thread-local detail (`take_last_error_detail()`) that the Cython error -path appends to the raised `CUDAError`, so the user learns that the caller's -context was not restored and which context is current. When both the operation -and the restoration fail, the operation status is returned and the restoration -failure is reported out of band. Tests inject restoration failures with +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 keeps it from attaching +to an unrelated error if the caller never raises that status; `enter_context` +clears any stale detail. Tests inject restoration failures with `set_context_restore_fault_for_testing()`. ### Reporting from non-propagating paths -Deleters, CUDA callbacks and cleanup-after-failure cannot raise. They report -through `report_cuda_error()` / `report_message()` (the `pw_*` wrappers -decorate destroy calls with it), which emit a `cuda.core.CUDAWarning` through +Deleters and CUDA callbacks cannot raise. They report through +`report_cuda_error()` / `report_message()` (the `pw_*` wrappers decorate +destroy calls with it), 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` @@ -306,6 +316,11 @@ discarded silently anywhere in this layer, and nothing in this layer terminates 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: `note_or_report_cuda_error()` 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. + ## 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 75d2d27bc7d..fc05d376487 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.cpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.cpp @@ -211,11 +211,12 @@ class GILAcquireGuard { // Warning category registered by _resource_handles.pyx (cuda.core.CUDAWarning). std::atomic warning_category{nullptr}; -// Thread-local detail attached to the next raised CUDAError (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[256] = {0}; -thread_local char taken_error_detail[256] = {0}; +// 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; @@ -295,17 +296,62 @@ void report_cuda_error(const char* operation, CUresult status, const char* detai report_message(message); } -const char* take_last_error_detail() noexcept { - if (!last_error_detail[0]) { +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 note_or_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); + 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)); - last_error_detail[0] = 0; + 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 { @@ -349,36 +395,47 @@ CUresult restore_context(CUcontext previous) noexcept { return p_cuCtxSetCurrent(previous); } -// Record why the CUresult about to be returned should be explained further -// when it is raised as a CUDAError: the caller's context was not restored. -void note_context_not_restored(CUcontext previous) noexcept { +// 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; " + "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; a restoration failure is then reported -// out of band. Otherwise the restoration status is returned, annotated for the -// eventual CUDAError. +// 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; } - if (operation_status != CUDA_SUCCESS) { - report_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restore_status); - return operation_status; - } - note_context_not_restored(previous); - return restore_status; + note_context_not_restored(previous, operation_status, restore_status); + return operation_status != CUDA_SUCCESS ? operation_status : restore_status; } // Require a callable to be invocable without throwing. @@ -491,7 +548,10 @@ CUresult cleanup_in_context(const ContextHandle& h_context, const char* name, } CUresult restore = exit_context(previous, changed, CUDA_SUCCESS); if (restore != CUDA_SUCCESS) { + // Nothing is raised here, so the detail exit_context recorded has no + // exception to attach to: report it and drop the detail. report_cuda_error(name, restore, "failed while restoring the caller's context"); + clear_last_error_detail(); } return status; } @@ -581,8 +641,8 @@ CUresult context_get_device(const ContextHandle& h_context, CUdevice* device) no // 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, a restoration failure is reported -// out of band and *restore_status is CUDA_SUCCESS. +// 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 { @@ -607,12 +667,10 @@ CUresult graph_node_set_params(CUgraphNode node, CUgraphNodeParams* params, if (restored == CUDA_SUCCESS) { return status; } - if (status != CUDA_SUCCESS) { - report_cuda_error("cuCtxSetCurrent (restoring the caller's context)", restored); - return status; + note_context_not_restored(previous, status, restored); + if (status == CUDA_SUCCESS) { + *restore_status = restored; } - note_context_not_restored(previous); - *restore_status = restored; return status; } diff --git a/cuda_core/cuda/core/_cpp/resource_handles.hpp b/cuda_core/cuda/core/_cpp/resource_handles.hpp index 069e3ec8319..5ad65143659 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.hpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.hpp @@ -80,12 +80,20 @@ 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 note_or_report_cuda_error(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 appends it to the raised CUDAError. Thread-local; take_ returns -// the detail (valid until the next take on this thread) and clears it, or -// nullptr when none is recorded. -const char* take_last_error_detail() noexcept; +// 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 diff --git a/cuda_core/cuda/core/_resource_handles.pxd b/cuda_core/cuda/core/_resource_handles.pxd index 37f4bcb0435..339e2610b0e 100644 --- a/cuda_core/cuda/core/_resource_handles.pxd +++ b/cuda_core/cuda/core/_resource_handles.pxd @@ -176,7 +176,9 @@ 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 const char* take_last_error_detail() noexcept nogil +cdef void note_or_report_cuda_error( + 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 diff --git a/cuda_core/cuda/core/_resource_handles.pyi b/cuda_core/cuda/core/_resource_handles.pyi index c44bcd46a03..882ccdc483d 100644 --- a/cuda_core/cuda/core/_resource_handles.pyi +++ b/cuda_core/cuda/core/_resource_handles.pyi @@ -49,3 +49,10 @@ def _set_context_restore_fault_for_testing(status: int): injected failure leaves the target context current, exactly as a failing ``cuCtxSetCurrent`` would, so callers must restore the context themselves. """ +def _note_or_report_cuda_error_for_testing(status: int): + """Attach a failed CUDA call to the exception being handled, or report it. + + Test hook for ``note_or_report_cuda_error()``. 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 09692f9d9a2..692ae7368e5 100644 --- a/cuda_core/cuda/core/_resource_handles.pyx +++ b/cuda_core/cuda/core/_resource_handles.pyx @@ -45,7 +45,14 @@ cdef extern from "_cpp/resource_handles.hpp" namespace "cuda_core": 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 - const char* take_last_error_detail "cuda_core::take_last_error_detail" () noexcept nogil + void note_or_report_cuda_error "cuda_core::note_or_report_cuda_error" ( + 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 _note_or_report_cuda_error_local "cuda_core::note_or_report_cuda_error" ( + 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 @@ -589,6 +596,17 @@ def _set_context_restore_fault_for_testing(int status): """ set_context_restore_fault_for_testing(status) + +def _note_or_report_cuda_error_for_testing(int status): + """Attach a failed CUDA call to the exception being handled, or report it. + + Test hook for ``note_or_report_cuda_error()``. Called inside an ``except`` + block it adds a note to the exception being handled (Python 3.11+); anywhere + else it emits a ``CUDAWarning``. + """ + _note_or_report_cuda_error_local( + b"cuTestOperation", status, b"failed while testing") + # ============================================================================= # NVRTC function pointer initialization # ============================================================================= diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyi b/cuda_core/cuda/core/_utils/cuda_utils.pyi index b190e5967d0..565200277ba 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyi +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyi @@ -19,9 +19,10 @@ class CUDAWarning(RuntimeWarning): ``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, or while a context switch is undone - after the requested operation already succeeded. Those failures are reported - as this warning instead, and the affected resource may have leaked. + 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:: diff --git a/cuda_core/cuda/core/_utils/cuda_utils.pyx b/cuda_core/cuda/core/_utils/cuda_utils.pyx index 38fc42027df..b6f33112953 100644 --- a/cuda_core/cuda/core/_utils/cuda_utils.pyx +++ b/cuda_core/cuda/core/_utils/cuda_utils.pyx @@ -45,9 +45,10 @@ class CUDAWarning(RuntimeWarning): ``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, or while a context switch is undone - after the requested operation already succeeded. Those failures are reported - as this warning instead, and the affected resource may have leaked. + 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:: @@ -160,6 +161,16 @@ 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 @@ -167,20 +178,23 @@ cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil: 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() + 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: - suffix = f" ({detail.decode()})" if detail != NULL else "" # 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}{suffix}") - if desc_err != cydriver.CUresult.CUDA_SUCCESS: - raise CUDAError(f"{name.decode()}{suffix}") - raise CUDAError(f"{name.decode()}: {desc.decode()}{suffix}") + 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 7033f90b239..98faf09eec3 100644 --- a/cuda_core/cuda/core/graph/_graph_builder.pyx +++ b/cuda_core/cuda/core/graph/_graph_builder.pyx @@ -20,7 +20,7 @@ from cuda.core.graph._subclasses cimport ( ExecutableGraphNode, create_executable_node_view, ) -from cuda.core._resource_handles cimport report_cuda_error +from cuda.core._resource_handles cimport note_or_report_cuda_error, report_cuda_error from cuda.core._resource_handles cimport ( GraphExecHandle, GraphHandle, @@ -862,9 +862,9 @@ cdef class GraphBuilder: invalidate_child_graph_state( self._h_graph, c_new_node) else: - # The original exception propagates; the failed rollback is - # reported out of band (error handling policy). - report_cuda_error( + # The original exception propagates with the failed rollback + # attached as a note (error handling policy). + note_or_report_cuda_error( b"cuGraphDestroyNode", rollback_status, b"failed while rolling back a child graph node; the node remains in the graph") raise diff --git a/cuda_core/cuda/core/graph/_graph_node.pyx b/cuda_core/cuda/core/graph/_graph_node.pyx index 831c0f30b69..d072971615d 100644 --- a/cuda_core/cuda/core/graph/_graph_node.pyx +++ b/cuda_core/cuda/core/graph/_graph_node.pyx @@ -45,7 +45,7 @@ from cuda.core.graph._subclasses cimport ( SwitchNode, WhileNode, ) -from cuda.core._resource_handles cimport report_cuda_error +from cuda.core._resource_handles cimport note_or_report_cuda_error from cuda.core._resource_handles cimport ( GraphHandle, GraphNodeHandle, @@ -1116,9 +1116,9 @@ cdef inline ChildGraphNode GN_embed(GraphNode self, GraphDefinition child_def): if rollback_status == cydriver.CUDA_SUCCESS: invalidate_child_graph_state(h_graph, new_node) else: - # The original exception propagates; the failed rollback is - # reported out of band (error handling policy). - report_cuda_error( + # The original exception propagates with the failed rollback + # attached as a note (error handling policy). + note_or_report_cuda_error( b"cuGraphDestroyNode", rollback_status, b"failed while rolling back a child graph node; the node remains in the graph") raise diff --git a/cuda_core/docs/source/error_handling.rst b/cuda_core/docs/source/error_handling.rst index a6761aadd24..ad37c76422e 100644 --- a/cuda_core/docs/source/error_handling.rst +++ b/cuda_core/docs/source/error_handling.rst @@ -35,9 +35,14 @@ When a ``cuda.core`` call raises, the following hold: documentation says so where it applies (for example the graph mutation methods that add several driver edges). -``cuda.core`` does not swallow driver errors. A failure that would otherwise be -hidden, for example because it occurred while another exception was already -propagating, is reported as described in the next section. +``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 ------------------------------ @@ -45,10 +50,9 @@ 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; -- callbacks invoked by CUDA; -- cleanup performed after an operation has already failed, such as rolling back - a partially built graph node or restoring the caller's CUDA context. + 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 and the CUDA error. The warning means the @@ -86,8 +90,10 @@ 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 whose message states that the caller's -context could not be restored and which context is now current. Call +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` @@ -124,4 +130,4 @@ Process termination ``cuda.core`` does not abort the process in response to a CUDA error, including errors that cannot be raised, and including failures to restore the caller's context. Aborting is reserved for an internal invariant violation where -continuing could corrupt memory, and no such code path exists in this release. +continuing could corrupt memory. 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 3e94b19a77b..2eb19e5b22a 100644 --- a/cuda_core/docs/source/release/1.3.0-notes.rst +++ b/cuda_core/docs/source/release/1.3.0-notes.rst @@ -63,10 +63,12 @@ Fixes and enhancements - 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 states 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. A - restoration failure during resource cleanup is reported as + 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 @@ -79,10 +81,11 @@ Fixes and enhancements current instead of leaving the thread with no context. It also works when no context is current, returning ``None``. -- Failed rollbacks of a partially embedded child graph node and failed - ``cuStreamEndCapture`` calls made when a still-building - :class:`~graph.GraphBuilder` is garbage collected are now reported as - :class:`CUDAWarning`; both were silent. +- 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/test_error_handling.py b/cuda_core/tests/test_error_handling.py index e716b001f66..b0dbc10d9df 100644 --- a/cuda_core/tests/test_error_handling.py +++ b/cuda_core/tests/test_error_handling.py @@ -4,15 +4,17 @@ """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; and a failure to restore the -caller's context is raised (or reported) with an explanation rather than -swallowed or turned into a process abort. 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. +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 from contextlib import contextmanager import pytest @@ -26,17 +28,32 @@ DeviceMemoryResourceOptions, LegacyPinnedMemoryResource, ) -from cuda.core._resource_handles import _set_context_restore_fault_for_testing +from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource +from cuda.core._resource_handles import ( + _note_or_report_cuda_error_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 @@ -69,11 +86,16 @@ 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, match="could not be restored") as excinfo: + with assert_no_cuda_warning(), pytest.raises(CUDAError) as excinfo: dev.create_stream() - message = str(excinfo.value) - assert "CUDA_ERROR_INVALID_CONTEXT" in message - assert "Device.set_current()" in message + 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) @@ -85,8 +107,29 @@ 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, match="could not be restored"): + 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) @@ -184,8 +227,9 @@ def test_memset_update_keeps_new_owners_alive_when_context_cannot_be_restored(de other_dev.set_current() _set_context_restore_fault_for_testing(INVALID_CONTEXT) try: - with pytest.raises(CUDAError, match="could not be restored"): + 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() @@ -205,3 +249,34 @@ def as_bytes(buffer): 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(): + _note_or_report_cuda_error_for_testing(INVALID_VALUE) + else: + with pytest.warns(CUDAWarning, match="cuTestOperation failed while testing"): + _note_or_report_cuda_error_for_testing(INVALID_VALUE) + with assert_no_cuda_warning(): + _note_or_report_cuda_error_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"): + _note_or_report_cuda_error_for_testing(INVALID_VALUE) From 9773124876b589981896be5f83c9356eeec18cad Mon Sep 17 00:00:00 2001 From: Andy Jost Date: Fri, 4 Sep 2026 11:30:08 -0700 Subject: [PATCH 9/9] cuda.core: follow the review of #2750 and flush the stderr fallback Rebased onto the reviewed head of #2750. Adjustments the rebase needed: - The review's warning for an undo skipped after a failed context restoration is routed through report_cuda_error(), so it carries the CUDA status and becomes a CUDAWarning like every other non-raising report. - invoke_in_context and invoke_in_context_or_undo now reject empty handles themselves, so context_get_device drops its own guard like the other helpers did; enter_context's no-op for empty handles is documented as used only by graph_node_set_params. - _SynchronousMemoryResource moved to its own module; the error-handling test imports it from there. The review's two teardown tests asserted that stderr stayed empty; under the policy a teardown failure is a CUDAWarning, so they assert that no CUDAWarning is issued instead (and are marked thread_unsafe because warning capture is process-global). - report_message() flushes stderr after its last-resort fprintf, so the text is not lost if the process dies right after (review comment). Co-Authored-By: Claude Fable 5.1 --- cuda_core/cuda/core/_cpp/resource_handles.cpp | 6 +++++- cuda_core/tests/test_error_handling.py | 2 +- cuda_core/tests/test_memory.py | 14 +++++++------- 3 files changed, 13 insertions(+), 9 deletions(-) diff --git a/cuda_core/cuda/core/_cpp/resource_handles.cpp b/cuda_core/cuda/core/_cpp/resource_handles.cpp index fc05d376487..8e455c73854 100644 --- a/cuda_core/cuda/core/_cpp/resource_handles.cpp +++ b/cuda_core/cuda/core/_cpp/resource_handles.cpp @@ -270,6 +270,7 @@ void report_message(const char* message) noexcept { } } std::fprintf(stderr, "%s\n", message); + std::fflush(stderr); } // Report a failed non-CUDA call (NVRTC, NVVM, nvJitLink) from a path that @@ -362,7 +363,10 @@ 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; diff --git a/cuda_core/tests/test_error_handling.py b/cuda_core/tests/test_error_handling.py index b0dbc10d9df..caed0beea66 100644 --- a/cuda_core/tests/test_error_handling.py +++ b/cuda_core/tests/test_error_handling.py @@ -28,7 +28,7 @@ DeviceMemoryResourceOptions, LegacyPinnedMemoryResource, ) -from cuda.core._memory._device_memory_resource import _SynchronousMemoryResource +from cuda.core._memory._synchronous_memory_resource import _SynchronousMemoryResource from cuda.core._resource_handles import ( _note_or_report_cuda_error_for_testing, _set_context_restore_fault_for_testing, diff --git a/cuda_core/tests/test_memory.py b/cuda_core/tests/test_memory.py index 661e540c308..ccf064df767 100644 --- a/cuda_core/tests/test_memory.py +++ b/cuda_core/tests/test_memory.py @@ -2264,8 +2264,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 @@ -2280,13 +2281,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 @@ -2296,14 +2298,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"),