diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 2c9d96f3e..b3bb11d2f 100644 --- a/example/CMakeLists.txt +++ b/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) diff --git a/example/README.md b/example/README.md index 7b26155e5..0b999bda7 100644 --- a/example/README.md +++ b/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 diff --git a/example/cuda/batched-write/CMakeLists.txt b/example/cuda/batched-write/CMakeLists.txt new file mode 100644 index 000000000..d311cc1a2 --- /dev/null +++ b/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) diff --git a/example/cuda/batched-write/README.md b/example/cuda/batched-write/README.md new file mode 100644 index 000000000..15cc19239 --- /dev/null +++ b/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 +``` diff --git a/example/cuda/batched-write/batched_write.cu b/example/cuda/batched-write/batched_write.cu new file mode 100644 index 000000000..7e4f4d9a2 --- /dev/null +++ b/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 +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace capy = boost::capy; +namespace ex = capy::example; + +capy::io_task +write_batch(capy::any_write_stream& dest, + std::array 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 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 r; + std::latch done{1}; + capy::run_async(pool.get_executor(), + [&](capy::io_result ir) { r = ir; done.count_down(); })( + write_batch(dest, bufs)); + done.wait(); + auto const& [ec, n] = r; + + std::vector 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; +} diff --git a/example/cuda/datamovement/README.md b/example/cuda/datamovement/README.md index adb8200ee..2e7d919c2 100644 --- a/example/cuda/datamovement/README.md +++ b/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: diff --git a/example/cuda/datamovement/cuda_datamovement.hpp b/example/cuda/datamovement/cuda_datamovement.hpp index 6a540f418..d14305593 100644 --- a/example/cuda/datamovement/cuda_datamovement.hpp +++ b/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)}; } };