diff --git a/source/2026-07-july/d4251-ioawaitable-cuda.md b/source/2026-07-july/d4251-ioawaitable-cuda.md index 35562f4..7a7f419 100644 --- 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.[14] 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.[15] Costs a spinning thread, but stays stable as the number of worker threads grows. -- **Blocking**: a service thread runs `cudaStreamSynchronize`.[16] Costs one parked thread per outstanding wait, but keeps the worker threads free. +- **Polling**: a thread periodically calls `cudaStreamQuery`.[15] 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`.[15] Costs one parked thread per outstanding wait, but keeps the worker threads free. - **Callback**: `cudaLaunchHostFunc` enqueues a host function into the stream.[7] 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`.[14] Its host function fires on a dedicated internal CPU thread created by the CUDA driver, not the application thread.[17][18] It cannot call CUDA APIs and must not create transitive dependencies on outstanding CUDA work. +`cudaLaunchHostFunc` is the recommended replacement for the deprecated `cudaStreamAddCallback`.[14] Its host function fires on a dedicated internal CPU thread created by the CUDA driver, not the application thread.[16][17] 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[19] demonstrates this directly. A CHEP 2026 report on trigger scheduling[20] 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.[18] 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[19] demonstrates this directly. For example, a report from CERN Next Generation Triggers[20] 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[14] confirms that operations in a stream execute in enqueue order, and the CUDA Runtime API documentation[7] states that `cudaLaunchHostFunc` callbacks block later work in the stream until they return.[21] 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.[7] 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.[22] If the callback blocks on a user lock while the CUDA launch queue is full, the enqueuing thread blocks too, producing deadlock.[23] Notification is unidirectional: `cudaLaunchHostFunc` provides stream-to-CPU notification only and cannot make the stream wait for a CPU-side signal.[24] 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[20] favors those alternatives as the worker-thread count grows, and CERN's traccc port[25] 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.[7] 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.[22] If the callback blocks on a user lock while the CUDA launch queue is full, the enqueuing thread blocks too, producing deadlock.[23] Notification is unidirectional: `cudaLaunchHostFunc` provides stream-to-CPU notification only and cannot make the stream wait for a CPU-side signal.[24] 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[25] implements all three strategies through the same IoAwaitable protocol, allowing the notification mechanism to be selected per deployment. Scaling measurements[20] 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.[26] With pageable memory allocated via `malloc` or `new`, the call blocks the host thread despite the `Async` suffix.[27] 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).**[36] 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.**[25] 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.**[25] 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).**[37] 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)[44] documents 8 bytes saved per nesting level via constant pointer offsets. This is real. -CUDA Graphs[45] provide GPU-side work-graph optimization at the driver level. The driver sees SM count, memory bandwidth, occupancy, and hardware topology. Stream capture[16] records kernel DAGs: +CUDA Graphs[45] provide GPU-side work-graph optimization at the driver level. The driver sees SM count, memory bandwidth, occupancy, and hardware topology. Stream capture[15] 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,[14] enabling the same pre-allocated op-state pattern that networking sockets use. Independent projects at NVIDIA Research (cuda-oxide),[36] CERN,[25] the University of Wisconsin-Madison (Taro),[37] and Schrödinger (Desmond)[39] 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[25] implements the callback, event polling, and deferred synchronization as interchangeable awaitables, and a CHEP 2026 measurement[20] 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,[14] enabling the same pre-allocated op-state pattern that networking sockets use. Independent projects at NVIDIA Research (cuda-oxide),[36] CERN,[25] the University of Wisconsin-Madison (Taro),[37] and Schrödinger (Desmond)[39] 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[25] implements the callback, event polling, and deferred synchronization as interchangeable awaitables, and the measurements[20] 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ł 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ü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ö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ö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).