Virtual Memory Management & Ring Buffers
Relevant source files
The following files were used as context for generating this wiki page:
Virtual Memory Management (VMM) and Ring Buffers in the gr-cuda project provide a high-performance, zero-copy pipeline for GPU-accelerated signal processing. By leveraging CUDA VMM and POSIX shared memory mechanisms, the system implements circular buffers that appear as contiguous virtual address ranges. This architecture mirrors the circular buffer "trick" used in GNU Radio on the CPU, allowing wrap-around accesses to be handled transparently without the need for data compaction.
The core of this system is the cuda_buffer, which manages both host and device memory. It uses double-mapped 2N virtual regions to ensure that any contiguous span up to N items starting at any index is valid. This eliminates the overhead of managing buffer wraps in signal processing kernels and enables efficient data movement between blocks with minimal synchronization overhead.
Sources: include/gnuradio/cuda/cuda_buffer.h:23-45, README.md:12-20
Architecture and Components
The VMM system is divided into host-side and device-side implementations, both following the principle of double-mapping a physical allocation into a contiguous virtual address space.
Device VMM Ring Buffer
The device_vmm_ring class manages GPU memory using the CUDA Virtual Memory Management API. It reserves a virtual address range of 2N bytes and maps a physical allocation of N bytes twice into that range.
- Allocation: Uses
cuMemAddressReservefor the 2N range andcuMemCreatefor the N-byte physical allocation. - Mapping: The physical allocation is mapped twice via
cuMemMapto make the buffer circular. - Granularity: The system queries the minimum allocation granularity for the current device to ensure proper alignment.
Sources: lib/detail/device_vmm_ring.h:24-52, include/gnuradio/cuda/cuda_buffer.h:26-30
Host MMAP Ring Buffer
The host_mmap_ring class manages RAM-backed pages on the host side. It uses POSIX memfd_create and mmap to achieve the same 2N virtual mapping for host memory.
- DMA Support: Optionally registers the 2N range for asynchronous CUDA DMA using
cudaHostRegister. - Transparency: Provides a standard
char*pointer to the base of the virtual region.
Sources: lib/detail/host_mmap_ring.h:26-52, include/gnuradio/cuda/cuda_buffer.h:32-34
Virtual Address Mapping Layout
The following diagram illustrates how the physical N-byte buffer is mapped into the 2N virtual address space.
This double-mapping ensures that a read or write operation crossing the N-item boundary automatically wraps back to the beginning of the physical buffer via hardware address translation.
Sources: include/gnuradio/cuda/cuda_buffer.h:26-34, lib/detail/device_vmm_ring.h:31-35
Synchronization Model
The cuda_buffer utilizes three primary CUDA events to coordinate data movement and processing between producers and consumers without blocking the CPU thread.
| Event | Purpose | Implementation Detail |
|---|---|---|
| Device-Ready | Signals data is ready on GPU | Recorded by producer after writing; waited on by consumer before reading. |
| Host-Ready | Signals D2H copy complete | Recorded after async Device-to-Host copy; used by _read_pointer for CPU waits. |
| Read-Done | Signals consumers finished | Tracks when all consumers have finished reading so the producer can overwrite. |
Sources: include/gnuradio/cuda/cuda_buffer.h:47-58
Buffer Management Flow
The flow of synchronization during a standard work cycle is managed by wait_for_work() and mark_work_done().
Implementation Details
Key Data Structures and Functions
| Component | Description |
|---|---|
cuda_buffer | The primary buffer class inheriting from buffer_single_mapped. |
space_available() | Returns items the writer may produce using circular semantics. |
post_work(nitems) | Orchestrates H2D, D2H, or D2D transfers after block execution. |
index_sub(a, b) | Corrected modular distance calculation (a - b) mod d_bufsize. |
query_vmm_granularity_for_current_device() | Helper to determine device-specific alignment requirements. |
Sources: include/gnuradio/cuda/cuda_buffer.h:70-134, lib/detail/device_vmm_ring.h:20-22
Configuration Options
Users can configure the synchronization behavior via the GNU Radio configuration files.
- [cuda_buffer] gpu_busy_wait:
true(default): UsescudaEventSynchronize(spin-poll) for lowest latency.false: UsescudaEventBlockingSyncto put the thread to sleep, freeing the CPU core.
Sources: include/gnuradio/cuda/cuda_buffer.h:60-68
Integration with Block Helpers
The cuda_block_helper.h provides inline functions that simplify buffer interactions for block developers:
// Wait for upstream data and ensure output buffers are safe to write
gr::cuda::wait_for_work(detail(), d_stream);
// Launch GPU kernel
my_kernel<<<grid, block, 0, d_stream>>>(in, out, n);
// Signal completion to downstream and upstream blocks
gr::cuda::mark_work_done(detail(), d_stream);
Sources: include/gnuradio/cuda/cuda_block_helper.h:28-100, include/gnuradio/cuda/cuda_block.h:58-90
Summary
The Virtual Memory Management and Ring Buffer system in gr-cuda provides a robust foundation for high-throughput GPU DSP. By abstracting the complexities of circular buffer management and CUDA event synchronization into the cuda_buffer and associated detail classes, the project enables developers to write GPU blocks that interact with data pointers as if they were simple, linear arrays, while maintaining the performance benefits of a zero-copy circular pipeline.