cppalliance · m-fila · Jun 22, 2026 · Jul 13, 2026
diff --git a/source/2026-07-july/d4251-ioawaitable-cuda.md b/source/2026-07-july/d4251-ioawaitable-cuda.md
@@ -129,13 +129,15 @@ The full execution model built on this protocol is specified in [P4003R3](https:
 
 CUDA streams are in-order queues where operations execute sequentially.<sup>[14]</sup> When GPU work completes, the host needs notification. Three mechanisms exist, and the IoAwaitable protocol is independent of which one a given awaitable uses:
 
-- **Polling**: a service thread loops `cudaEventQuery` on a recorded event.<sup>[15]</sup> Costs a spinning thread, but stays stable as the number of worker threads grows.
-- **Blocking**: a service thread runs `cudaStreamSynchronize`.<sup>[16]</sup> Costs one parked thread per outstanding wait, but keeps the worker threads free.
+- **Polling**: a thread periodically calls `cudaStreamQuery`.<sup>[15]</sup> The polling can be implemented with a dedicated thread, but it can also be integrated into an existing work loop by interleaving completion checks with other work items. This avoids blocking threads, but requires periodic polling activity.
+- **Blocking**: a service thread runs `cudaStreamSynchronize`.<sup>[15]</sup> Costs one parked thread per outstanding wait, but keeps the worker threads free.
 - **Callback**: `cudaLaunchHostFunc` enqueues a host function into the stream.<sup>[7]</sup> No busy-wait and the simplest to wire up, but a single CUDA-internal worker services every callback across all streams, so it scales poorly as the number of worker threads grows.
 
-`cudaLaunchHostFunc` is the recommended replacement for the deprecated `cudaStreamAddCallback`.<sup>[14]</sup> Its host function fires on a dedicated internal CPU thread created by the CUDA driver, not the application thread.<sup>[17]</sup><sup>[18]</sup> It cannot call CUDA APIs and must not create transitive dependencies on outstanding CUDA work.
+`cudaLaunchHostFunc` is the recommended replacement for the deprecated `cudaStreamAddCallback`.<sup>[14]</sup> Its host function fires on a dedicated internal CPU thread created by the CUDA driver, not the application thread.<sup>[16]</sup><sup>[17]</sup> It cannot call CUDA APIs and must not create transitive dependencies on outstanding CUDA work.
 
-The choice among the three is a scaling tradeoff, not a correctness one. All three satisfy `IoAwaitable` and, driving the same GPU pipeline, produce identical results at runtime; the accompanying notification-strategies example<sup>[19]</sup> demonstrates this directly. A CHEP 2026 report on trigger scheduling<sup>[20]</sup> finds that the callback handler scales poorly as the worker-thread count grows, while event polling and deferred synchronization remain stable. In a multi-threaded framework, prefer polling or deferred synchronization; reach for the callback for its simplicity in low-concurrency settings.
+For cases where waiting for the entire stream is too coarse, CUDA events provide finer-grained completion points. An event can be recorded at a specific position in a stream, and the host can then wait for that event instead of the entire stream. The same polling and blocking approaches apply: `cudaEventQuery` can be used to poll event completion, and `cudaEventSynchronize` can be used to block until the event is complete.<sup>[18]</sup> Unlike streams, events do not support callback-based notification.
+
+The choice among the three  mechanisms is a scaling tradeoff, not a correctness one. All three satisfy `IoAwaitable` and, driving the same GPU pipeline, produce identical results at runtime; the accompanying notification-strategies example<sup>[19]</sup> demonstrates this directly. For example, a report from CERN Next Generation Triggers<sup>[20]</sup> finds that the callback handler scales poorly as the worker-thread count grows, while event polling and deferred synchronization remain stable. In a multi-threaded framework, prefer polling or deferred synchronization; reach for the callback for its simplicity in low-concurrency settings.
 
 This is the same structural pattern as epoll, IOCP, or io_uring completions arriving on arbitrary threads. In all cases, an async operation completes on a thread that is not the application's, and the application must dispatch the result to the correct execution context. This is the exact problem that Capy's executor-affinity dispatch was designed to solve.
 
@@ -335,7 +337,7 @@ public:
 
 The `resume_ctx` is a pre-allocated member of `cuda_stream`, not heap-allocated per operation. This is safe because the coroutine suspends on each `co_await`, so only one operation is in-flight per `cuda_stream` at a time. The CUDA Programming Guide<sup>[14]</sup> confirms that operations in a stream execute in enqueue order, and the CUDA Runtime API documentation<sup>[7]</sup> states that `cudaLaunchHostFunc` callbacks block later work in the stream until they return.<sup>[21]</sup> The pre-allocated `resume_ctx` is never accessed concurrently. This is the same one-at-a-time invariant that Capy's sockets rely on for their pre-allocated op states in the networking domain.
 
-`cudaLaunchHostFunc` has documented constraints that production code must respect. The callback must not call CUDA APIs or synchronize on outstanding CUDA work.<sup>[7]</sup> A single CUDA-internal worker thread may service all callbacks across all streams; on loaded systems, OS scheduling can starve this thread, producing latency spikes up to 12ms between callback completion and stream resumption.<sup>[22]</sup> If the callback blocks on a user lock while the CUDA launch queue is full, the enqueuing thread blocks too, producing deadlock.<sup>[23]</sup> Notification is unidirectional: `cudaLaunchHostFunc` provides stream-to-CPU notification only and cannot make the stream wait for a CPU-side signal.<sup>[24]</sup> These constraints apply equally to any pattern that uses `cudaLaunchHostFunc` for completion notification, including the hand-rolled awaitable in Section 6 and any sender-based wrapper that uses the same mechanism. They do not invalidate the pattern but they bound its applicability in high-throughput pipelines. They are specific to the callback mechanism: the polling and deferred-synchronization awaitables of Section 5 sidestep all four. A CHEP 2026 scaling measurement<sup>[20]</sup> favors those alternatives as the worker-thread count grows, and CERN's traccc port<sup>[25]</sup> implements all three strategies over this one protocol so the mechanism can be selected per deployment. The IoAwaitable protocol is the same in every case; only the notification source changes.
+`cudaLaunchHostFunc` has documented constraints that production code must respect. The callback must not call CUDA APIs or synchronize on outstanding CUDA work.<sup>[7]</sup> A single CUDA-internal worker thread may service all callbacks across all streams; on loaded systems, OS scheduling can starve this thread, producing latency spikes up to 12ms between callback completion and stream resumption.<sup>[22]</sup> If the callback blocks on a user lock while the CUDA launch queue is full, the enqueuing thread blocks too, producing deadlock.<sup>[23]</sup> Notification is unidirectional: `cudaLaunchHostFunc` provides stream-to-CPU notification only and cannot make the stream wait for a CPU-side signal.<sup>[24]</sup> These constraints apply equally to any pattern that uses `cudaLaunchHostFunc` for completion notification, including the hand-rolled awaitable in Section 6 and any sender-based wrapper that uses the same mechanism. They do not invalidate the pattern but they bound its applicability in high-throughput pipelines. They are specific to the callback mechanism: the polling and deferred-synchronization awaitables of Section 5 sidestep all these limitations. CERN's traccc port<sup>[25]</sup> implements all three strategies through the same IoAwaitable protocol, allowing the notification mechanism to be selected per deployment. Scaling measurements<sup>[20]</sup> show that polling and deferred synchronization remain stable as the worker-thread count grows, whereas the callback-based approach scales less well. In all cases, the IoAwaitable protocol remains unchanged; only the underlying notification source differs.
 
 One caveat: `cudaMemcpyAsync` is only truly asynchronous with pinned (page-locked) memory.<sup>[26]</sup> With pageable memory allocated via `malloc` or `new`, the call blocks the host thread despite the `Async` suffix.<sup>[27]</sup> For multi-gigabyte model weight transfers, this distinction matters.
 
@@ -649,11 +651,11 @@ The gap between networking ambition and deployed evidence suggests that data mov
 
 ## 14. Independent Validation
 
-Several independent projects have arrived at the same design: coroutine-based async completion for GPU and HPC data movement. The notification mechanism that bridges GPU completion to coroutine resumption varies - a host-function callback (`cudaLaunchHostFunc`, or its driver-level equivalent `cuLaunchHostFunc`), event polling, or deferred stream synchronization - but the coroutine completion model is common to all of them. The callback is the most frequently chosen bridge in the projects below because it is the simplest; it is not the only one in use.
+Several independent projects have arrived at the same design: coroutine-based async completion for GPU and HPC data movement. The notification mechanism that bridges GPU completion to coroutine resumption varies - a host-function callback (`cudaLaunchHostFunc`, or its driver-level equivalent `cuLaunchHostFunc`), event or stream polling, or deferred synchronization - but the coroutine completion model is common to all of them. The callback is the most frequently chosen bridge in the projects below because it is the simplest; it is not the only one in use.
 
 **cuda-oxide (NVIDIA Labs, Rust).**<sup>[36]</sup> NVIDIA's own research lab implemented the same mechanism in Rust. Their `DeviceFuture` submits GPU work, enqueues a `cuLaunchHostFunc` callback that sets an `AtomicBool` and wakes a Tokio `Waker`, and the async runtime resumes the task on the next poll. Zero busy-wait. The three-state machine (Idle, Executing, Complete) is structurally identical to a network socket future. When NVIDIA's own research lab arrives at the same `cudaLaunchHostFunc`-to-async-runtime pattern independently, in a different language, the convergence is a data point about where the pattern fits naturally.
 
-**CERN wp1.7-traccc.**<sup>[25]</sup> As part of its wp1.7 work package evaluating C++20 coroutines for task scheduling, CERN ported the traccc GPU track-reconstruction pipeline from stdexec to Capy. It implements its CUDA completion strategies behind a single `await_strategy` selector - among them a `cudaLaunchHostFunc` callback, event polling, and deferred `cudaStreamSynchronize` - each an awaitable with the signature `await_suspend(std::coroutine_handle<>, boost::capy::io_env const*)` that posts the coroutine handle back to `env->executor`. That a real reconstruction workload exercises all three notification mechanisms over one protocol is the most concrete evidence in this survey that the coroutine model is not bound to the callback.
+**CERN wp1.7-traccc.**<sup>[25]</sup> As part of its evaluation of C++20 coroutines for task scheduling, CERN Next Generation Triggers project ported the traccc GPU track-reconstruction pipeline from stdexec to Capy. It implements its CUDA completion strategies behind a single `await_strategy` selector - among them a `cudaLaunchHostFunc` callback, event polling, and deferred `cudaStreamSynchronize` - each an awaitable with the signature `await_suspend(std::coroutine_handle<>, boost::capy::io_env const*)` that posts the coroutine handle back to `env->executor`. That a real reconstruction workload exercises all three notification mechanisms over one protocol is the most concrete evidence in this survey that the coroutine model is not bound to the callback.
 
 **Taro (University of Wisconsin-Madison).**<sup>[37]</sup> A C++20 coroutine task-graph system for CPU-GPU workloads. GPU tasks suspend the CPU thread via coroutines when waiting for GPU completion, allowing other tasks to run. Uses `cudaLaunchHostFunc` for the callback. Published at Euro-Par 2024 (as TaroRTL) and presented at CppCon 2023. TaroRTL reported a 40-80% speedup over RTLflow, a state-of-the-art GPU-accelerated RTL simulator.
 
@@ -673,7 +675,7 @@ These projects span GPU compute, molecular dynamics, high-energy physics, RDMA n
 
 Sender pipelines provide compile-time `operation_state` fusion. [P3425R1](https://www.open-std.org/jtc1/sc22/wg21/docs/papers/2025/p3425r1.html)<sup>[44]</sup> documents 8 bytes saved per nesting level via constant pointer offsets. This is real.
 
-CUDA Graphs<sup>[45]</sup> provide GPU-side work-graph optimization at the driver level. The driver sees SM count, memory bandwidth, occupancy, and hardware topology. Stream capture<sup>[16]</sup> records kernel DAGs:
+CUDA Graphs<sup>[45]</sup> provide GPU-side work-graph optimization at the driver level. The driver sees SM count, memory bandwidth, occupancy, and hardware topology. Stream capture<sup>[15]</sup> records kernel DAGs:
 
 ```c
 cudaStreamBeginCapture(stream,
@@ -846,7 +848,7 @@ The preceding sections present convergent findings. This section addresses fores
 
 ## 19. Conclusion
 
-A protocol handler compiled once links against TCP, RDMA, or GPU device memory without recompilation. This is possible because byte-oriented data movement - host-device memcpy, inter-GPU collectives over NVLink, RDMA transfers between nodes, and TCP sockets - shares a common async completion model that the IoAwaitable protocol captures with zero per-operation allocation. The CUDA Programming Guide confirms that single-stream callbacks are strictly serialized,<sup>[14]</sup> enabling the same pre-allocated op-state pattern that networking sockets use. Independent projects at NVIDIA Research (cuda-oxide),<sup>[36]</sup> CERN,<sup>[25]</sup> the University of Wisconsin-Madison (Taro),<sup>[37]</sup> and Schr&ouml;dinger (Desmond)<sup>[39]</sup> have converged on coroutine-based completion for data movement without coordination. The notification mechanism is a free variable the protocol does not fix: CERN's traccc port<sup>[25]</sup> implements the callback, event polling, and deferred synchronization as interchangeable awaitables, and a CHEP 2026 measurement<sup>[20]</sup> finds the callback the worst-scaling of the three under many worker threads.
+A protocol handler compiled once links against TCP, RDMA, or GPU device memory without recompilation. This is possible because byte-oriented data movement - host-device memcpy, inter-GPU collectives over NVLink, RDMA transfers between nodes, and TCP sockets - shares a common async completion model that the IoAwaitable protocol captures with zero per-operation allocation. The CUDA Programming Guide confirms that single-stream callbacks are strictly serialized,<sup>[14]</sup> enabling the same pre-allocated op-state pattern that networking sockets use. Independent projects at NVIDIA Research (cuda-oxide),<sup>[36]</sup> CERN,<sup>[25]</sup> the University of Wisconsin-Madison (Taro),<sup>[37]</sup> and Schr&ouml;dinger (Desmond)<sup>[39]</sup> have converged on coroutine-based completion for data movement without coordination. The notification mechanism is a free variable the protocol does not fix: CERN's traccc port<sup>[25]</sup> implements the callback, event polling, and deferred synchronization as interchangeable awaitables, and the measurements<sup>[20]</sup> finds the callback the worst-scaling of the three under many worker threads.
 
 `cudaLaunchHostFunc` has documented limitations (Section 7) that bound the applicability of the callback mechanism in high-throughput GPU pipelines. Those limitations are specific to the callback: the protocol equally admits event polling and deferred synchronization, which sidestep them where they bite.
 
@@ -864,7 +866,7 @@ Eric Niebler, Micha&lstrok; Dominiak, Lewis Baker, Lucian Radu Teodorescu, Lee H
 
 Richard Smith and Gor Nishanov for P0981R0 (HALO analysis). Chuanqi Xu for the `[[clang::coro_await_elidable]]` attribute and P2477R3 (coroutine allocation elision). Dietmar K&uuml;hl and Maikel Nadolski for P3552R3 (`std::execution::task`). Lewis Baker for cppcoro, the operator `co_await` and symmetric transfer blog posts, and P3425R1 (operation-state sizes). Michael Wong for P4029R0 (SG14 priority list).
 
-Michael Garland and the NVIDIA stdexec team for the nvexec GPU schedulers and the Maxwell FDTD benchmark. The CERN wp1.7 team for their C++20 coroutine task-scheduling experiments and the Capy IoAwaitable integration. Dian-Lun Lin (University of Wisconsin-Madison) for Taro and its CppCon 2023 presentation. The NVIDIA Labs team for cuda-oxide. Jiqun Tu (NVIDIA) and Ellery Russell (Schr&ouml;dinger) for the Desmond coroutine integration presented at GTC 2024. The TTG/PaRSEC team for demonstrating coroutine-based heterogeneous GPU dispatch at DOE Exascale scale.
+Michael Garland and the NVIDIA stdexec team for the nvexec GPU schedulers and the Maxwell FDTD benchmark. The CERN Next Generation Triggers project for their C++20 coroutine task-scheduling experiments and the Capy IoAwaitable integration. Dian-Lun Lin (University of Wisconsin-Madison) for Taro and its CppCon 2023 presentation. The NVIDIA Labs team for cuda-oxide. Jiqun Tu (NVIDIA) and Ellery Russell (Schr&ouml;dinger) for the Desmond coroutine integration presented at GTC 2024. The TTG/PaRSEC team for demonstrating coroutine-based heterogeneous GPU dispatch at DOE Exascale scale.
 
 This paper was generated with AI assistance (Claude, via Cursor).
 
@@ -898,13 +900,13 @@ This paper was generated with AI assistance (Claude, via Cursor).
 
 [14] [CUDA Programming Guide: Asynchronous Concurrent Execution](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/asynchronous-execution.html) (NVIDIA, 2024).
 
-[15] [CUDA Runtime API: Event Management](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html) (NVIDIA, 2024).
+[15] [CUDA Runtime API: Stream Management](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html) (NVIDIA, 2024).
 
-[16] [CUDA Runtime API: Stream Management](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html) (NVIDIA, 2024).
+[16] [CUDA Handbook: Stream Callbacks](https://www.cudahandbook.com/2012/09/stream-callbacks/) (Nicholas Wilt, 2012).
 
-[17] [CUDA Handbook: Stream Callbacks](https://www.cudahandbook.com/2012/09/stream-callbacks/) (Nicholas Wilt, 2012).
+[17] [Stack Overflow: Exception Handling in cudaLaunchHostFunc Callbacks](https://stackoverflow.com/questions/75145603/catching-an-exception-thrown-from-a-callback-in-cudalaunchhostfunc) (2023).
 
-[18] [Stack Overflow: Exception Handling in cudaLaunchHostFunc Callbacks](https://stackoverflow.com/questions/75145603/catching-an-exception-thrown-from-a-callback-in-cudalaunchhostfunc) (2023).
+[18] [CUDA Runtime API: Event Management](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html) (NVIDIA, 2024).
 
 [19] [Accompanying examples](https://github.com/cppalliance/capy/tree/a226b793a3409f07723d2e90dd154e7461fffe89/example) - the compileable demonstrations for this paper, pinned at commit `a226b79` of the official repository (C++ Alliance). Section 5 (the three notification mechanisms, `callback_awaitable`, `poll_awaitable`, `deferred_sync_awaitable`): [`example/cuda/notification-strategies`](https://github.com/cppalliance/capy/tree/a226b793a3409f07723d2e90dd154e7461fffe89/example/cuda/notification-strategies). Sections 6-8 and 14 (`cuda_stream`, `cuda_device_stream`, CUDA Graphs): [`example/cuda/datamovement`](https://github.com/cppalliance/capy/tree/a226b793a3409f07723d2e90dd154e7461fffe89/example/cuda/datamovement). Section 16 (the `await_sender` bridge, `handle_request`): [`example/cuda/pipeline/cuda_pipeline.cu`](https://github.com/cppalliance/capy/blob/a226b793a3409f07723d2e90dd154e7461fffe89/example/cuda/pipeline/cuda_pipeline.cu). Sections 10-11 (compound results and HPC-fabric signatures): [`example/fabrics/fabrics.cpp`](https://github.com/cppalliance/capy/blob/a226b793a3409f07723d2e90dd154e7461fffe89/example/fabrics/fabrics.cpp).