From 0cf43632a8a729aafca39d383bab0587d179227d Mon Sep 17 00:00:00 2001 From: Steve Gerbino Date: Wed, 26 Aug 2026 19:51:48 +0200 Subject: [PATCH] fix(example): report stream faults from callback-based CUDA awaitables cudaLaunchHostFunc passes no completion status to its host function, so a callback-based IoAwaitable could not observe a stream fault and resumed with success. Query the stream with cudaStreamQuery in await_resume, back on a worker thread where CUDA calls are permitted, and surface the sticky error through the awaitable's normal error path. Apply this to callback_awaitable in notification-strategies and to cuda_stream / cuda_device_stream in datamovement. Add a --fault mode to notification-strategies that launches a null-pointer kernel and awaits it via one mechanism per process under a watchdog; all three now resume with cudaErrorIllegalAddress, where the callback previously resumed with success. Poll and deferred-sync were already correct because cudaEventQuery and cudaStreamSynchronize return the status. --- example/cuda/datamovement/README.md | 5 +- .../cuda/datamovement/cuda_datamovement.hpp | 33 +++++-- .../cuda/notification-strategies/README.md | 42 +++++++- .../notification_strategies.cu | 98 ++++++++++++++++++- .../notification_strategies.hpp | 15 ++- 5 files changed, 177 insertions(+), 16 deletions(-) diff --git a/example/cuda/datamovement/README.md b/example/cuda/datamovement/README.md index e645722cd..adb8200ee 100644 --- a/example/cuda/datamovement/README.md +++ b/example/cuda/datamovement/README.md @@ -11,7 +11,10 @@ What is validated: - `cuda_stream_awaiter`: the io_env-less baseline. Asserted to be a standard awaitable but **not** an `IoAwaitable`. - `cuda_stream`: `memcpy_h2d` / `memcpy_d2h` / `synchronize` return - `IoAwaitable`s. + `IoAwaitable`s. Because `cudaLaunchHostFunc` passes no status to its + host function, `await_resume` calls `cudaStreamQuery` after + resumption so a stream fault surfaces as an error instead of success + (see the `--fault` probe in `../notification-strategies`). - NCCL interop: `ncclAllReduce` on `cuda_stream::native_handle()` followed by `co_await synchronize()`. Built only when NCCL is found at configure time. diff --git a/example/cuda/datamovement/cuda_datamovement.hpp b/example/cuda/datamovement/cuda_datamovement.hpp index 165bbc260..6a540f418 100644 --- a/example/cuda/datamovement/cuda_datamovement.hpp +++ b/example/cuda/datamovement/cuda_datamovement.hpp @@ -58,6 +58,19 @@ inline std::error_code make_cuda_error(cudaError_t e) noexcept return std::error_code(static_cast(e), cuda_category()); } +/// Return the stream's sticky error, or success if it is idle or busy. +/// +/// `cudaLaunchHostFunc` passes no completion status to its host function, +/// so a callback-based awaitable queries the stream after resumption, +/// back on a worker thread where CUDA calls are permitted. +inline std::error_code stream_error(cudaStream_t s) noexcept +{ + auto st = cudaStreamQuery(s); + if(st == cudaSuccess || st == cudaErrorNotReady) + return {}; + return make_cuda_error(st); +} + /// A minimal hand-rolled CUDA-completion awaitable (no executor /// affinity, cancellation, or frame allocator). Resumes on the CUDA /// driver callback thread. @@ -153,9 +166,11 @@ class cuda_stream void await_resume() { + if(! self->error_) + self->error_ = stream_error(self->stream_); if(self->error_) - throw std::system_error(self->error_); - self->error_ = {}; + throw std::system_error( + std::exchange(self->error_, {})); } }; @@ -185,9 +200,11 @@ class cuda_stream void await_resume() { + if(! self->error_) + self->error_ = stream_error(self->stream_); if(self->error_) - throw std::system_error(self->error_); - self->error_ = {}; + throw std::system_error( + std::exchange(self->error_, {})); } }; @@ -331,12 +348,10 @@ class cuda_device_stream io_result await_resume() { + if(! self->error_) + self->error_ = stream_error(self->stream_); if(self->error_) - { - auto ec = self->error_; - self->error_ = {}; - return {ec, 0}; - } + return {std::exchange(self->error_, {}), 0}; auto n = buf.size(); self->offset_ += n; return {std::error_code(), n}; diff --git a/example/cuda/notification-strategies/README.md b/example/cuda/notification-strategies/README.md index 0a69765ca..5b80fc85b 100644 --- a/example/cuda/notification-strategies/README.md +++ b/example/cuda/notification-strategies/README.md @@ -21,10 +21,44 @@ Each awaitable captures the executor and posts the continuation through it, so the coroutine always resumes on a worker thread, never on a CUDA or service thread. -The callback mechanism is the only one that cannot report a stream error -through its host function; `cudaLaunchHostFunc` does not pass completion -status to the callback, so `callback_awaitable` always resumes with -success — this is an inherent limitation of the API. +The callback mechanism is the only one that cannot observe a stream +error from inside its notification: `cudaLaunchHostFunc` passes no +completion status to the host function (unlike the deprecated +`cudaStreamAddCallback`, whose callback received a `cudaError_t`). +`callback_awaitable` compensates by calling `cudaStreamQuery` in +`await_resume`, after the coroutine is back on a worker thread where +CUDA calls are permitted, and reports the stream's sticky error from +there. Polling and deferred synchronization get the status for free from +`cudaEventQuery` and `cudaStreamSynchronize`. + +### Fault probe + +`--fault ` launches a kernel that writes through a null +pointer (a sticky `cudaErrorIllegalAddress`) and awaits the stream via +one mechanism. The context is dead afterwards, so each mechanism is +probed in its own process; a watchdog reports a coroutine that never +resumes. + +``` +for m in callback poll deferred-sync; do + ./build-cuda/example/cuda/notification-strategies/capy_example_cuda_notification_strategies --fault $m +done +``` + +Observed on an RTX 4060 (CUDA 13.3, clang 22): + +``` +callback: resumed with error: an illegal memory access was encountered +poll: resumed with error: an illegal memory access was encountered +deferred-sync: resumed with error: an illegal memory access was encountered +``` + +Before the `cudaStreamQuery` in `await_resume`, the callback line read +`resumed with success despite the fault`. On this driver the host +function still fires after a sticky error, contrary to the +`cudaLaunchHostFunc` documentation's statement that it "will not be +called in the event of an error in the CUDA context"; the watchdog path +exists for drivers where it does not. ### Service lifetime diff --git a/example/cuda/notification-strategies/notification_strategies.cu b/example/cuda/notification-strategies/notification_strategies.cu index 733422bf0..1753fa101 100644 --- a/example/cuda/notification-strategies/notification_strategies.cu +++ b/example/cuda/notification-strategies/notification_strategies.cu @@ -22,9 +22,12 @@ #include #include +#include #include +#include #include #include +#include #include #include #include @@ -53,6 +56,15 @@ fill_kernel(int* p, int n, int v) p[i] = v; } +// Writes through a null pointer: a sticky cudaErrorIllegalAddress that +// poisons the context. Used by --fault to see which mechanisms still +// deliver the error to the awaiting coroutine. +__global__ void +fault_kernel(int* p) +{ + p[threadIdx.x] = 1; +} + enum class notify { callback, @@ -60,6 +72,20 @@ enum class notify deferred_sync }; +bool +parse_notify(char const* s, notify& out) noexcept +{ + if(std::strcmp(s, "callback") == 0) + out = notify::callback; + else if(std::strcmp(s, "poll") == 0) + out = notify::poll; + else if(std::strcmp(s, "deferred-sync") == 0) + out = notify::deferred_sync; + else + return false; + return true; +} + char const* name_of(notify how) noexcept { @@ -149,11 +175,81 @@ run_one(capy::thread_pool& pool, return result; } +// Launch a faulting kernel and await the stream via `how`. Returns the +// error the mechanism reported, or success if it reported none. +capy::task +run_fault(ex::cuda_stream& stream, + ex::cuda_event& event, + notify how, + ex::poll_service& poll_svc, + ex::sync_service& sync_svc) +{ + auto s = stream.native_handle(); + fault_kernel<<<1, 32, 0, s>>>(nullptr); + event.record(s); + co_return co_await wait(stream, event, how, poll_svc, sync_svc); +} + +// The faulted context is dead afterwards, so each mechanism is probed +// in its own process. A mechanism whose coroutine never resumes is +// reported as such after the watchdog expires; the process is then +// exited without teardown because the coroutine is still suspended. +int +fault_main(notify how) +{ + capy::thread_pool pool(4); + ex::poll_service poll_svc; + ex::sync_service sync_svc; + ex::cuda_stream stream; + ex::cuda_event event; + + std::error_code ec; + std::binary_semaphore done{0}; + capy::run_async(pool.get_executor(), + [&](std::error_code e) { ec = e; done.release(); })( + run_fault(stream, event, how, poll_svc, sync_svc)); + + std::cout << name_of(how) << ": "; + if(! done.try_acquire_for(std::chrono::seconds(3))) + { + std::cout << "coroutine never resumed (watchdog expired)\n"; + std::cout.flush(); + std::_Exit(EXIT_FAILURE); + } + if(! ec) + { + std::cout << "resumed with success despite the fault\n"; + std::cout.flush(); + std::_Exit(EXIT_FAILURE); + } + std::cout << "resumed with error: " << ec.message() << "\n"; + std::cout.flush(); + std::_Exit(EXIT_SUCCESS); +} + } // namespace int -main() +main(int argc, char** argv) { + if(argc == 3 && std::strcmp(argv[1], "--fault") == 0) + { + notify how; + if(! parse_notify(argv[2], how)) + { + std::cerr << "usage: " << argv[0] + << " [--fault callback|poll|deferred-sync]\n"; + return EXIT_FAILURE; + } + return fault_main(how); + } + if(argc != 1) + { + std::cerr << "usage: " << argv[0] + << " [--fault callback|poll|deferred-sync]\n"; + return EXIT_FAILURE; + } + int device_count = 0; if(cudaGetDeviceCount(&device_count) != cudaSuccess || device_count == 0) { diff --git a/example/cuda/notification-strategies/notification_strategies.hpp b/example/cuda/notification-strategies/notification_strategies.hpp index 5c0956e68..e154b4892 100644 --- a/example/cuda/notification-strategies/notification_strategies.hpp +++ b/example/cuda/notification-strategies/notification_strategies.hpp @@ -301,6 +301,10 @@ class cuda_event continuation through the captured executor instead. One operation is in flight at a time, so the resume context is a member rather than a per-operation allocation. + + The host function receives no completion status, so `await_resume` + queries the stream after resumption to report a fault that occurred + before the callback ran. */ struct callback_awaitable { @@ -345,9 +349,18 @@ struct callback_awaitable return std::noop_coroutine(); } + // cudaLaunchHostFunc hands the host function no completion status, + // so a stream fault is invisible from inside on_complete. Back on the + // worker thread, CUDA calls are permitted again and cudaStreamQuery + // exposes the stream's sticky error. std::error_code await_resume() const noexcept { - return ec_; + if(ec_) + return ec_; + auto s = cudaStreamQuery(stream); + if(s == cudaSuccess || s == cudaErrorNotReady) + return {}; + return make_cuda_error(s); } };