evo_cuda_utils is a header-only C++17 library with utilities for CUDA runtime
code developed by Evocargo. The library is mainly useful for code that works
asynchronously with CUDA streams, but it also provides synchronous variants for
simpler use cases.
It wraps common CUDA resources, reports CUDA failures via GpuError, and
provides CudaStorage: a resizable GPU buffer with vector-like operations.
CudaStorage is parameterized by a memory strategy. Use CudaAsyncStorage for
stream-bound asynchronous pipelines, and CudaSyncStorage when blocking CUDA
operations are sufficient.
Public headers are located in include/evo_cuda_utils/*_inl.h.
- Header-only library for easy integration
- RAII wrappers for CUDA streams and events
- Synchronous and stream-ordered asynchronous GPU storage
- Pinned host-memory allocators, vectors, and values
- Exception-based CUDA error reporting with
GpuError
- C++17-compatible compiler
- CUDA Runtime (tested on CUDA 12+)
- CMake (optional)
Because evo_cuda_utils is header-only, it does not have to be built or
installed. You can copy the include/evo_cuda_utils directory into your project
and include its headers directly:
#include "your_dir/evo_cuda_utils/cuda_storage_inl.h"
#include "your_dir/evo_cuda_utils/cuda_stream_inl.h"The library can also be distributed as a DEB package or another package format supported by CPack.
From the repository root, configure and build the package:
cmake -B build
cpack -G DEB --config build/cpack_config/evo_cuda_utils_Config.cmakeTo generate another package format, replace DEB with the required CPack
generator.
After installing the generated package through the appropriate system package manager, you can use the library from CMake.
Installed headers will be placed in include/evo_cuda_utils.
find_package(evo_cuda_utils REQUIRED)
find_package(CUDAToolkit REQUIRED)
target_link_libraries(my_cuda_target
PUBLIC
evo_cuda_utils::evo_cuda_utils
CUDA::cudart
)To configure, build, and run the tests from the repository root:
cmake -B build -DBUILD_TESTS=ON
cmake --build build --parallel
ctest --test-dir build --verboseevo_cuda_utils provides several groups of CUDA utilities:
| Component | Purpose |
|---|---|
CudaStream |
RAII wrapper for cudaStream_t. |
CudaHostPinnedAllocator<T> |
STL-compatible allocator based on pinned host memory. Can be used with standard containers. |
CudaHostPinnedVector<T> |
std::vector<T, CudaHostPinnedAllocator<T>>; useful as a host staging buffer for async CUDA transfers. |
CudaHostPinnedValue<T> |
Single value stored in pinned host memory; useful for scalar parameters, counters, sizes, or flags. |
CudaAsyncMemoryStrategy<T> |
Memory strategy based on cudaMallocAsync, cudaMemcpyAsync, cudaMemsetAsync, and cudaFreeAsync on a fixed stream. |
CudaSyncMemoryStrategy<T> |
Memory strategy based on blocking CUDA runtime memory operations. |
CudaStorage<T, MemoryStrategy> |
Resizable device buffer with vector-like operations. The memory strategy defines whether operations are synchronous or stream-ordered asynchronous. |
CudaAsyncStorage<T> |
Alias for CudaStorage with CudaAsyncMemoryStrategy; use it for buffers bound to a CUDA stream. |
CudaSyncStorage<T> |
Alias for CudaStorage with CudaSyncMemoryStrategy; use it when blocking CUDA operations are sufficient. |
CudaEvent |
RAII wrapper for cudaEvent_t; used for completion points, synchronization, and timing. |
CudaEventProtected<T> |
Host-side value protected by a CUDA event; useful when host data is reused across asynchronous CUDA work and synchronous CPU access. |
GpuError |
Exception type used by the library. |
get_multi_processor_count() |
Returns the SM count of the current CUDA device. |
In asynchronous code, evo_cuda_utils is usually used around one explicitly
owned CUDA stream.
A typical pipeline looks like this:
- Create or receive a
CudaStream. - Create
CudaAsyncStoragebuffers bound to that stream. - Upload host data from pinned memory.
- Queue kernels and storage operations on the same stream.
- Download results into pinned host memory.
- Synchronize the stream or a recorded event before reading results on the CPU.
When the CPU needs access to a specific intermediate result, CudaEvent can be
used to wait for a selected point in the stream instead of synchronizing all
queued work. For host-side data that is reused between asynchronous CUDA work
and synchronous CPU access, CudaEventProtected can be used to record this
boundary automatically.
The important rule is that asynchronous storage is stream-bound. Operations on a
CudaAsyncStorage object, kernels that use its data(), and manual CUDA calls
that touch the same memory should be ordered on the same stream, unless the
caller provides explicit synchronization.
For synchronous code, use CudaSyncStorage. Its memory operations are blocking
and do not require a stream contract.
A CUDA stream is an ordered queue of GPU work. In asynchronous code, a stream usually defines the execution order for memory operations, kernel launches, and synchronization points.
CudaStream is an RAII wrapper around cudaStream_t: create on construction,
destroy on destruction. It is movable, non-copyable, and implicitly
convertible to cudaStream_t.
synchronize_stream blocks until all work previously queued on the given stream
is complete. CudaStream::synchronize() provides the same operation for the
owned stream.
- After move, the source holds a null handle (
is_null() == true). Do not use it for new work. CallingCudaStream::synchronize()on a moved-from stream object throwsGpuError. synchronize_stream(nullptr)invokescudaStreamSynchronize(nullptr)(default stream semantics in CUDA). Prefer an explicitCudaStreamobject in new asynchronous code.
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void pipeline_with_owned_stream()
{
evo::cuda_utils::CudaStream stream{};
// The storage queues memory operations on this stream but does not own it.
evo::cuda_utils::CudaAsyncStorage<float> buffer{stream};
buffer.resize(1024);
my_kernel<<<blocks, threads, 0, stream>>>(buffer.data(), buffer.size());
// Wait for resize and kernel execution before using results or leaving the sync boundary.
stream.synchronize();
}CudaHostPinnedAllocator<T> is an STL-compatible allocator based on pinned host
memory (cudaMallocHost / cudaFreeHost). CudaHostPinnedAllocator<T> can be
used with any STL-compatible container when pinned host allocation is needed.
CudaHostPinnedVector<T> is an alias for
std::vector<T, CudaHostPinnedAllocator<T>>. It is useful as a resizable host
staging buffer for asynchronous host-device transfers, for example with
CudaAsyncStorage::load_from_vector, load_to_vector, load_to_host_memory,
or direct CUDA Runtime API calls.
Pinned host memory does not synchronize with GPU work by itself. The caller is responsible for ordering asynchronous transfers on a CUDA stream and synchronizing the stream or an event before reading data written by the GPU.
- Failed pinned memory allocation throws
GpuError.
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void pinned_vector_upload_download(evo::cuda_utils::CudaStream& stream)
{
evo::cuda_utils::CudaHostPinnedVector<float> input(1024, 1.0f);
evo::cuda_utils::CudaHostPinnedVector<float> output;
evo::cuda_utils::CudaAsyncStorage<float> device_buffer{stream};
device_buffer.load_from_vector(input); // async H2D from pinned host memory
process_kernel<<<grid, block, 0, stream>>>(device_buffer.data(), device_buffer.size());
device_buffer.load_to_vector(output); // async D2H into pinned host memory
// Wait before reading output on the CPU.
stream.synchronize();
}CudaHostPinnedValue<T> owns a single value in pinned host memory.
Pinned host memory is useful for asynchronous CUDA transfers. Typical use cases
are scalar parameters, counters, sizes, flags, or small POD-like structs that
are copied between host and device with cudaMemcpyAsync or storage helper
methods.
CudaHostPinnedValue<T> does not synchronize with GPU work by itself. The
caller is responsible for ordering transfers and kernels on a CUDA stream and
synchronizing the stream or an event before reading data written by the GPU.
Tmust not be an array type.- The object is copyable; move operations fall back to copy. Each instance owns its own pinned allocation.
CudaHostPinnedValue is useful when a kernel produces a single scalar result,
such as a count, status, flag, or computed parameter. The value can be copied
from device memory to pinned host memory asynchronously and read on the CPU
after synchronization.
#include <evo_cuda_utils/cuda_host_pinned_value_inl.h>
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void compute_count_async(evo::cuda_utils::CudaStream& stream)
{
evo::cuda_utils::CudaAsyncStorage<int> input{stream};
evo::cuda_utils::CudaAsyncStorage<std::size_t> device_count{stream};
evo::cuda_utils::CudaHostPinnedValue<std::size_t> host_count{0};
device_count.resize_for_overwrite(1);
evo::cuda_utils::CudaHostPinnedVector<int> source;
fill_input(source);
input.load_from_vector(source);
compute_count_kernel<<<grid, block, 0, stream>>>(input.data(), input.size(), device_count.data());
// Copy one scalar result into pinned host memory.
device_count.load_to_host_memory(host_count.pointer(), 1);
stream.synchronize();
const std::size_t count = host_count.get();
// ...
}CudaStorage<T, MemoryStrategy> owns a resizable buffer in GPU memory. It
provides vector-like operations such as resize, reserve, clear,
host-device copies, device-device copies, and append. When capacity grows, the
new capacity is at least 1.5x the previous capacity, or the requested size if it
is larger. CudaStorage does not define by itself whether memory operations are
blocking or asynchronous. Allocation, copy, memset, and free operations are
delegated to a MemoryStrategy.
This allows the same storage API to be used in two execution modes:
| Strategy | CUDA operations | Behavior |
|---|---|---|
CudaSyncMemoryStrategy<T> |
cudaMalloc, cudaMemcpy, cudaMemset, cudaFree |
Blocking memory operations |
CudaAsyncMemoryStrategy<T> |
cudaMallocAsync, cudaMemcpyAsync, cudaMemsetAsync, cudaFreeAsync on a fixed stream |
Stream-ordered memory operations |
template <typename T>
using CudaSyncStorage = CudaStorage<T, CudaSyncMemoryStrategy<T>>;
template <typename T>
using CudaAsyncStorage = CudaStorage<T, CudaAsyncMemoryStrategy<T>>;Use CudaAsyncStorage<T> when the buffer is part of an asynchronous CUDA
pipeline. Use CudaSyncStorage<T> when blocking memory operations are
sufficient.
CudaAsyncStorage<T> is bound to one CUDA stream through its
CudaAsyncMemoryStrategy<T>.
Storage operations, kernels that use data(), and manual CUDA calls that touch
the same memory should be ordered on the same stream, unless the caller provides
explicit synchronization.
Copy assignment, append, and device-device copies between two asynchronous
storages require compatible memory strategies. For CudaAsyncStorage<T>, this
means that both storages must use the same CUDA stream.
CudaAsyncStorage<T> and CudaAsyncMemoryStrategy<T> do not own the CUDA
stream. The stream lifetime must be managed externally and must exceed the
lifetime of every storage object that uses it. In particular, the stream must
still be valid when CudaAsyncStorage<T> is destroyed, because destruction may
queue an asynchronous free on that stream.
This example shows the basic CudaAsyncStorage usage pattern: host data is
uploaded to the device, kernels are queued on the same stream, and the result is
downloaded back to host memory. The host result must not be read before the
stream is synchronized.
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void process_frame_async()
{
evo::cuda_utils::CudaStream stream{};
evo::cuda_utils::CudaAsyncStorage<float> device_buffer{stream};
evo::cuda_utils::CudaHostPinnedVector<float> host_frame(1024, 0.0f);
device_buffer.load_from_vector(host_frame); // async H2D, queued on stream
my_preprocess_kernel<<<grid, block, 0, stream>>>(device_buffer.data(), device_buffer.size());
my_main_kernel<<<grid, block, 0, stream>>>(device_buffer.data(), device_buffer.size());
evo::cuda_utils::CudaHostPinnedVector<float> result;
device_buffer.load_to_vector(result); // async D2H, queued on the same stream
stream.synchronize();
// It is now safe to read result on the CPU.
}operator+= appends one device buffer to another using a device-to-device copy.
For asynchronous storage, both buffers must be bound to the same stream.
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void append_on_same_stream(evo::cuda_utils::CudaStream& stream)
{
evo::cuda_utils::CudaAsyncStorage<int> a{stream};
evo::cuda_utils::CudaAsyncStorage<int> b{stream};
evo::cuda_utils::CudaHostPinnedVector<int> src_a{1, 2, 3};
evo::cuda_utils::CudaHostPinnedVector<int> src_b{4, 5};
a.load_from_vector(src_a);
b.load_from_vector(src_b);
// Appends b to a by queuing a device-to-device copy on the shared stream.
a += b;
stream.synchronize();
// a now contains: 1, 2, 3, 4, 5
}reserve can be used to allocate capacity in advance. resize allocates if
necessary, changes the logical size, and initializes new elements with zero.
resize_for_overwrite changes the logical size without preserving old contents
when reallocation is needed, so the buffer must be filled before reading.
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void sizing_async(evo::cuda_utils::CudaStream& stream)
{
evo::cuda_utils::CudaAsyncStorage<float> buffer{stream};
buffer.reserve(50); // May queue allocation and device-to-device copy.
buffer.resize(100); // May grow capacity again. New elements are zero-filled.
// After resize_for_overwrite, treat the buffer contents as uninitialized.
buffer.resize_for_overwrite(200);
// Next operation must initialize buffer.data()[0..buffer.size() - 1].
init_kernel<<<grid, block, 0, stream>>>(buffer.data(), buffer.size());
stream.synchronize();
}Two device buffers use different streams and copy their results into different parts of one pinned host buffer. Each stream is synchronized before reading the corresponding part of host memory.
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void download_halves_in_parallel()
{
evo::cuda_utils::CudaStream stream_a{};
evo::cuda_utils::CudaStream stream_b{};
// Each async storage is bound to its own stream.
evo::cuda_utils::CudaAsyncStorage<float> gpu_a{stream_a};
evo::cuda_utils::CudaAsyncStorage<float> gpu_b{stream_b};
evo::cuda_utils::CudaHostPinnedVector<float> upload_a(50, 1.0f);
evo::cuda_utils::CudaHostPinnedVector<float> upload_b(50, 2.0f);
gpu_a.load_from_vector(upload_a);
gpu_b.load_from_vector(upload_b);
my_kernel_a<<<grid, block, 0, stream_a>>>(gpu_a.data(), gpu_a.size());
my_kernel_b<<<grid, block, 0, stream_b>>>(gpu_b.data(), gpu_b.size());
evo::cuda_utils::CudaHostPinnedVector<float> host(gpu_a.size() + gpu_b.size());
// Copy each device buffer into its own range of the same host buffer.
gpu_a.load_to_host_memory(host.data(), gpu_a.size());
gpu_b.load_to_host_memory(host.data() + gpu_a.size(), gpu_b.size());
// Wait for both async D2H copies before reading host.
stream_a.synchronize();
stream_b.synchronize();
// host[0..49] contains gpu_a result, host[50..99] contains gpu_b result.
}CudaSyncStorage uses blocking CUDA memory operations. It is useful when the
code does not need stream-ordered asynchronous execution.
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
void sync_upload_download()
{
evo::cuda_utils::CudaSyncStorage<double> gpu;
const evo::cuda_utils::CudaHostPinnedVector<double> host{1.0, 2.0, 3.0};
gpu.load_from_vector(host); // blocking H2D
evo::cuda_utils::CudaHostPinnedVector<double> back;
gpu.load_to_vector(back); // blocking D2H
gpu.clear(); // size = 0, capacity is retained
}Events mark points in a stream’s timeline. Use them to:
- check whether previously queued work has finished (
synchronize,is_ready); - coordinate multiple streams;
- measure elapsed GPU time (with timing-enabled flags).
CudaEvent owns cudaEvent_t (RAII, movable, non-copyable). record(stream)
inserts a marker after all work previously queued on that stream. The event
becomes ready when the stream reaches that marker.
synchronize_event blocks until the given event is complete.
- Default construction uses
cudaEventDisableTiming; usecudaEventDefaultor other timing-enabled flags if you needcudaEventElapsedTime. - After move, the source holds a null handle (
is_null() == true). Do not use it for new work. Callingrecord,synchronize, oris_readyon a moved-from event object throwsGpuError. synchronize_event(nullptr)throwsGpuError.
The stream first computes the output size and copies it to pinned host memory, then records an event. Heavier GPU work is queued after the event and can continue while the CPU does unrelated work.
The CPU synchronizes on the event later, only when it needs to read the size and resize buffers. Filling the output and the full device-to-host copy are queued after that. The stream is synchronized at the end.
#include <evo_cuda_utils/cuda_event_inl.h>
#include <evo_cuda_utils/cuda_host_pinned_value_inl.h>
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void pipeline_with_early_size_event()
{
evo::cuda_utils::CudaStream stream{};
evo::cuda_utils::CudaEvent size_ready{};
evo::cuda_utils::CudaAsyncStorage<int> input{stream};
evo::cuda_utils::CudaAsyncStorage<int> output{stream};
evo::cuda_utils::CudaAsyncStorage<std::size_t> device_size{stream};
evo::cuda_utils::CudaHostPinnedValue<std::size_t> host_size{0};
evo::cuda_utils::CudaHostPinnedVector<int> source(1024);
fill_input(source);
input.load_from_vector(source);
device_size.resize_for_overwrite(1);
compute_output_size_kernel<<<grid, block, 0, stream>>>(
input.data(),
input.size(),
device_size.data());
device_size.load_to_host_memory(host_size.pointer(), 1);
// The event marks the point where host_size becomes available.
// Work queued after this point is not waited for by size_ready.synchronize().
size_ready.record(stream);
preprocess_input_kernel<<<grid, block, 0, stream>>>(input.data(), input.size());
prepare_host_side_resources();
// Wait only for input upload, size computation, and size D2H copy.
size_ready.synchronize();
const std::size_t count = host_size.get();
output.resize_for_overwrite(count);
fill_output_kernel<<<grid, block, 0, stream>>>(
input.data(),
input.size(),
output.data(),
output.size());
evo::cuda_utils::CudaHostPinnedVector<int> host_result;
output.load_to_vector(host_result);
// Wait for preprocessing, output resize, output fill, and final D2H copy.
stream.synchronize();
}CudaEventProtected<T> stores a host-side value and protects access to it when
the value is reused between asynchronous CUDA operations and ordinary CPU code.
It is intended for one CUDA stream and one host thread. Internally, it uses a
CudaEvent to separate asynchronous access sessions from synchronous host
access:
get_async_access_guard(stream)— returns a guard for the protected value; when the guard is destroyed, an event is recorded on the stream.get_for_sync_access()— waits for the recorded event, then returns the value for ordinary CPU use without synchronizing the whole stream.
This is useful when a host-side object, for example CudaHostPinnedVector, is
reused between pipeline iterations while previously queued CUDA transfers may
still be using it.
- Single host thread only.
- Single CUDA stream per protected object.
- Access the value only through
get_for_sync_access()orget_async_access_guard(). - At most one
AsyncAccessGuardfor the same protected object should exist in a scope. - The asynchronous access scope starts when
get_async_access_guard()returns the guard and ends when the guard is destroyed. - Do not call
get_for_sync_access()inside a scope where anAsyncAccessGuardfor the same object still exists. - Do not call
get_async_access_guard()while using a reference returned byget_for_sync_access(). - References and pointers obtained through
AsyncAccessGuardmust not outlive the guard. - The guard must stay alive until all asynchronous operations using the value have been submitted on that stream.
CudaEventProtected is useful when the same host-side buffer is reused. The CPU
fills the buffer, submits an async upload, and later waits only for the upload
boundary before writing new input into the same buffer again.
#include <evo_cuda_utils/cuda_event_protected_inl.h>
#include <evo_cuda_utils/cuda_host_pinned_vector_inl.h>
#include <evo_cuda_utils/cuda_storage_inl.h>
#include <evo_cuda_utils/cuda_stream_inl.h>
void process_input_stream(evo::cuda_utils::CudaStream& stream)
{
evo::cuda_utils::CudaEventProtected<evo::cuda_utils::CudaHostPinnedVector<float>> host_input{1024};
evo::cuda_utils::CudaAsyncStorage<float> device_input{stream};
evo::cuda_utils::CudaAsyncStorage<float> device_output{stream};
for (std::size_t iteration = 0; iteration < 100; ++iteration) {
{
auto& input = host_input.get_for_sync_access();
fill_input(input, iteration);
}
{
auto input_guard = host_input.get_async_access_guard(stream);
device_input.load_from_vector(*input_guard);
}
// The next sync access waits only for the async use of host_input, not for later stream work.
device_output.resize_for_overwrite(device_input.size());
process_kernel<<<grid, block, 0, stream>>>(device_input.data(), device_input.size(), device_output.data());
}
stream.synchronize();
}GpuError is the exception type used by evo_cuda_utils.
It derives from std::runtime_error and is used to report errors detected by
this module, including failed CUDA Runtime calls.
GpuError can be constructed from cudaError_t, from a custom description, or
from a custom description together with cudaError_t.
get_multi_processor_count() returns the streaming multiprocessor count of the
current CUDA device.
This can be useful when choosing launch parameters relative to the device size, for example when computing the number of blocks from blocks-per-SM.
The function throws GpuError if the device information cannot be read.