fix(example): batch a buffer sequence in cuda_device_stream::write_some

example/CMakeLists.txt

@@ -29,6 +29,7 @@ if(BOOST_CAPY_BUILD_P2300_EXAMPLES)
 endif()
 
 if(BOOST_CAPY_BUILD_CUDA_EXAMPLES)
+    add_subdirectory(cuda/batched-write)
     add_subdirectory(cuda/datamovement)
     add_subdirectory(cuda/notification-strategies)
     add_subdirectory(cuda/pipeline)

example/README.md

@@ -46,6 +46,12 @@ Measures the `exec::any_sender` operation state and the heap allocation its
 `connect` performs, against the concrete operation state for the same
 pipeline. Requires stdexec (`BOOST_CAPY_BUILD_P2300_EXAMPLES=ON`).
 
+### cuda/batched-write/
+
+Runs the batched `write_some` of the CUDA device stream: three buffers, one
+await, device holds the concatenation. Requires CUDA
+(`BOOST_CAPY_BUILD_CUDA_EXAMPLES=ON`).
+
 ## Building
 
 ### CMake

example/cuda/batched-write/CMakeLists.txt

@@ -0,0 +1,38 @@
+#
+# Copyright (c) 2026 Steve Gerbino
+#
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+#
+# Official repository: https://github.com/cppalliance/capy
+#
+
+if(NOT CMAKE_CUDA_COMPILER)
+    message(FATAL_ERROR
+        "example/cuda/batched-write requires CUDA; "
+        "did you set BOOST_CAPY_BUILD_CUDA_EXAMPLES?")
+endif()
+
+file(GLOB_RECURSE PFILES CONFIGURE_DEPENDS
+    *.cu *.cuh *.hpp
+    CMakeLists.txt
+    README.md)
+
+source_group(TREE ${CMAKE_CURRENT_SOURCE_DIR} PREFIX "" FILES ${PFILES})
+
+add_executable(capy_example_cuda_batched_write ${PFILES})
+
+set_target_properties(capy_example_cuda_batched_write PROPERTIES
+    FOLDER "examples"
+    CUDA_STANDARD 20
+    CUDA_STANDARD_REQUIRED ON
+    CUDA_SEPARABLE_COMPILATION OFF)
+
+target_compile_features(capy_example_cuda_batched_write PRIVATE cxx_std_20)
+
+target_link_libraries(capy_example_cuda_batched_write PRIVATE
+    Boost::capy
+    CUDA::cudart)
+
+add_test(NAME capy_example_cuda_batched_write
+    COMMAND capy_example_cuda_batched_write)

example/cuda/batched-write/README.md

@@ -0,0 +1,15 @@
+# CUDA batched write example (P4251R0)
+
+Runs the batched `write_some` of `cuda_device_stream` from
+`example/cuda/datamovement`. Three host buffers are gathered through
+`any_write_stream` in a single `write_some`: every buffer is enqueued as
+its own `cudaMemcpyAsync`, one `cudaLaunchHostFunc` follows the last, and
+the coroutine suspends once. The program checks that the returned count
+is the sum of the buffer sizes and that the device holds the buffers'
+concatenation, and exits non-zero otherwise.
+
+Observed on an RTX 4060 (CUDA 13.3, clang 22):
+
+```
+batched write_some: 3 buffers, one await, 9 of 9 bytes, device holds "abcdefghi": ok
+```

example/cuda/batched-write/batched_write.cu

@@ -0,0 +1,86 @@
+//
+// Copyright (c) 2026 Steve Gerbino
+//
+// Distributed under the Boost Software License, Version 1.0. (See accompanying
+// file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+//
+// Official repository: https://github.com/cppalliance/capy
+//
+
+// Runs the batched write path of cuda_device_stream: three host buffers
+// go through any_write_stream in one write_some, so a single co_await
+// covers the whole sequence, and the device is checked to hold their
+// concatenation.
+
+#include "../datamovement/cuda_datamovement.hpp"
+
+#include <boost/capy.hpp>
+#include <boost/capy/ex/thread_pool.hpp>
+
+#include <array>
+#include <cstddef>
+#include <cstring>
+#include <iostream>
+#include <latch>
+#include <string>
+#include <string_view>
+#include <system_error>
+#include <vector>
+
+namespace capy = boost::capy;
+namespace ex   = capy::example;
+
+capy::io_task<std::size_t>
+write_batch(capy::any_write_stream& dest,
+    std::array<capy::const_buffer, 3> const& bufs)
+{
+    co_return co_await dest.write_some(bufs);
+}
+
+int main()
+{
+    constexpr std::string_view parts[] = {"abc", "defg", "hi"};
+    std::string expected;
+    for(auto p : parts)
+        expected += p;
+
+    std::byte* d_ptr = nullptr;
+    if(cudaMalloc(&d_ptr, expected.size()) != cudaSuccess)
+    {
+        std::cout << "cudaMalloc failed; no device available\n";
+        return 1;
+    }
+    cudaStream_t s = nullptr;
+    cudaStreamCreate(&s);
+
+    ex::cuda_device_stream gpu(s, d_ptr);
+    capy::any_write_stream dest(&gpu);
+    std::array<capy::const_buffer, 3> bufs{
+        capy::make_buffer(parts[0].data(), parts[0].size()),
+        capy::make_buffer(parts[1].data(), parts[1].size()),
+        capy::make_buffer(parts[2].data(), parts[2].size())};
+
+    capy::thread_pool pool(2);
+    capy::io_result<std::size_t> r;
+    std::latch done{1};
+    capy::run_async(pool.get_executor(),
+        [&](capy::io_result<std::size_t> ir) { r = ir; done.count_down(); })(
+            write_batch(dest, bufs));
+    done.wait();
+    auto const& [ec, n] = r;
+
+    std::vector<char> back(expected.size());
+    cudaMemcpy(back.data(), d_ptr, back.size(), cudaMemcpyDeviceToHost);
+    cudaStreamDestroy(s);
+    cudaFree(d_ptr);
+
+    bool const ok = ! ec && n == expected.size()
+        && std::memcmp(back.data(), expected.data(), expected.size()) == 0;
+    std::cout << "batched write_some: " << bufs.size() << " buffers, "
+        << "one await, " << n << " of " << expected.size()
+        << " bytes, device holds \""
+        << std::string_view(back.data(), back.size()) << "\": "
+        << (ok ? "ok" : "FAILED")
+        << (ec ? " (" + ec.message() + ")" : "") << "\n";
+    return ok ? 0 : 1;
+}

example/cuda/datamovement/README.md

@@ -1,10 +1,11 @@
 # CUDA data-movement example (P4251R0)
 
 Validation that the CUDA data-movement listings from
-P4251R0 "IoAwaitables for GPU Data Movement" are type-correct against the
+P4251R0 "Coroutine Completion for GPU Data Movement: Convergent Findings"
+are type-correct against the
 real `boost::capy` API and CUDA. The paper flags this code as AI-generated
 and unverified; this target proves it compiles. Nothing here is executed
-at runtime.
+at runtime; `example/cuda/batched-write` runs the batched `write_some`.
 
 What is validated:
 

example/cuda/datamovement/cuda_datamovement.hpp

@@ -271,9 +271,12 @@ class cuda_stream
 /// GPU device memory exposed as a WriteStream.
 ///
 /// Reshapes the `cuda_stream` memcpy pattern to satisfy `WriteStream`, so device
-/// memory can hide behind `any_write_stream`. Because `cudaMemcpyAsync`
-/// transfers the whole buffer in one operation, `write_some` never
-/// performs a partial write. Errors are delivered via `io_result`
+/// memory can hide behind `any_write_stream`. A buffer sequence is one
+/// batch: every buffer is enqueued as its own `cudaMemcpyAsync` and a
+/// single host function follows the last, so the stream keeps its queue
+/// depth and the coroutine suspends once per `write_some`. Because
+/// `cudaMemcpyAsync` transfers each buffer in one operation, `write_some`
+/// never performs a partial write. Errors are delivered via `io_result`
 /// rather than exceptions. Does not own `stream_`; the caller is
 /// responsible for the stream's lifetime.
 class cuda_device_stream
@@ -312,7 +315,8 @@ class cuda_device_stream
         struct awaitable
         {
             cuda_device_stream* self;
-            const_buffer buf;
+            Buffers buffers;
+            std::size_t total = 0;
 
             bool await_ready() const noexcept
             {
@@ -322,20 +326,27 @@ class cuda_device_stream
             std::coroutine_handle<>
             await_suspend(std::coroutine_handle<> h, io_env const* env)
             {
-                auto n = buf.size();
-                auto err = cudaMemcpyAsync(
-                    self->d_ptr_ + self->offset_,
-                    buf.data(), n,
-                    cudaMemcpyHostToDevice,
-                    self->stream_);
-                if(err != cudaSuccess)
+                // Enqueue the whole sequence before the host function so
+                // the stream runs the batch back to back.
+                auto const end = capy::end(buffers);
+                for(auto it = capy::begin(buffers); it != end; ++it)
                 {
-                    self->error_ = make_cuda_error(err);
-                    return h;
+                    const_buffer b = *it;
+                    auto err = cudaMemcpyAsync(
+                        self->d_ptr_ + self->offset_ + total,
+                        b.data(), b.size(),
+                        cudaMemcpyHostToDevice,
+                        self->stream_);
+                    if(err != cudaSuccess)
+                    {
+                        self->error_ = make_cuda_error(err);
+                        return h;
+                    }
+                    total += b.size();
                 }
                 self->cont_.h = h;
                 self->ctx_ = resume_ctx{env->executor, &self->cont_};
-                err = cudaLaunchHostFunc(
+                auto err = cudaLaunchHostFunc(
                     self->stream_, &on_complete, &self->ctx_);
                 if(err != cudaSuccess)
                 {
@@ -352,12 +363,11 @@ class cuda_device_stream
                     self->error_ = stream_error(self->stream_);
                 if(self->error_)
                     return {std::exchange(self->error_, {}), 0};
-                auto n = buf.size();
-                self->offset_ += n;
-                return {std::error_code(), n};
+                self->offset_ += total;
+                return {std::error_code(), total};
             }
         };
-        return awaitable{this, *capy::begin(buffers)};
+        return awaitable{this, std::move(buffers)};
     }
 };