fix(example): report stream faults from callback-based CUDA awaitables

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.

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<int>(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<std::size_t>
             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};

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 <mechanism>` 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
 

example/cuda/notification-strategies/notification_strategies.cu

@@ -22,9 +22,12 @@
 #include <boost/capy/concept/io_awaitable.hpp>
 #include <boost/capy/ex/thread_pool.hpp>
 
+#include <chrono>
 #include <cstdlib>
+#include <cstring>
 #include <iostream>
 #include <latch>
+#include <semaphore>
 #include <system_error>
 #include <type_traits>
 #include <vector>
@@ -53,13 +56,36 @@ 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,
     poll,
     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<std::error_code>
+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)
     {

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);
     }
 };