Zero-copy GPU Pipeline Architecture

Relevant source files

The following files were used as context for generating this wiki page:

• README.md
• include/gnuradio/cuda/api.h
• include/gnuradio/cuda/cuda_block.h
• include/gnuradio/cuda/cuda_block_helper.h
• include/gnuradio/cuda/cuda_buffer.h
• lib/CMakeLists.txt

The Zero-copy GPU Pipeline Architecture in gr-cuda is designed to enable high-performance signal processing by minimizing data movement between the CPU and GPU. The architecture centers around the cuda_buffer, which manages host-to-device (H2D) and device-to-host (D2H) transfers automatically. This allows GPU kernels to execute back-to-back on the device without "round-tripping" through host memory, effectively creating a continuous processing pipeline on the GPU hardware.

This system leverages CUDA Virtual Memory Management (VMM) to implement circular buffers on the GPU, mirroring the traditional buffer handling mechanisms used by GNU Radio on the CPU. By using asynchronous CUDA events and streams, the architecture ensures that data dependencies are respected without blocking the calling CPU threads, enabling maximum overlap of computation and data transfer on supported hardware.

Sources: README.md:12-20, include/gnuradio/cuda/cuda_block.h:73-82

Core Components and Mechanisms

cuda_buffer and Virtual Memory Management

The cuda_buffer is the foundational element that facilitates the zero-copy behavior. It uses the CUDA VMM API to implement circular buffers on the device. This allows the GPU to handle wrapped memory accesses similarly to how the CPU handles them via mmap in standard GNU Radio blocks.

Sources: README.md:18-20, lib/CMakeLists.txt:20-21

CUDA Stream Management

Every GPU-enabled block inherits from the cuda_block base class, which provides a dedicated, non-blocking CUDA stream (d_stream). This stream is used to enqueue all GPU kernels and memory operations, ensuring that tasks are executed in the correct order relative to each other while remaining asynchronous to the host CPU.

Sources: include/gnuradio/cuda/cuda_block.h:32-41, include/gnuradio/cuda/cuda_block.h:98-102

Component	Description
d_stream	Non-blocking CUDA stream for enqueuing GPU work.
d_min_grid_size	Populated via cudaOccupancyMaxPotentialBlockSize for optimal occupancy.
d_block_size	Thread block size determined for the specific kernel launch.
cuda_buffer	Manages GPU memory, VMM-based circularity, and synchronization events.

Sources: include/gnuradio/cuda/cuda_block.h:35-42, README.md:12-20

Synchronization Flow

The architecture utilizes a two-step synchronization process for every work() or general_work() call to handle data hazards without manual cudaMemcpy calls.

Pre-Kernel: wait_for_work

Before a block launches its kernel, it must call wait_for_work(). This function performs two critical GPU-side checks:

1. Input Buffers: It waits for "device-ready" events from upstream blocks, ensuring input data is fully transferred to or produced on the GPU.
2. Output Buffers: It waits for "read-done" events from downstream consumers, ensuring that previous data has been read and the buffer is safe to overwrite.

Sources: include/gnuradio/cuda/cuda_block_helper.h:38-66

Post-Kernel: mark_work_done

After the kernel launch, mark_work_done() is called to record new events:

1. Outputs: Records a "device-ready" event for downstream blocks.
2. Inputs: Records a "read-done" event to notify the upstream producer that the memory can be reused.

Sources: include/gnuradio/cuda/cuda_block_helper.h:88-112

Synchronization Logic Flow

The following diagram illustrates the sequence of synchronization between the CPU host thread and the GPU device stream during a standard block execution.

Data Pipeline Architecture

In a chain of GPU blocks, data stays on the device. The runtime only triggers H2D transfers at the start of the GPU chain and D2H transfers at the end.

Integration Patterns

C++ Block Implementation

C++ blocks must manually invoke the synchronization helpers within their work method. For general_work() blocks, mark_work_done() must be called specifically before produce() and consume() to ensure the GPU event is recorded before the write pointer advances.

int my_block_impl::work(int noutput_items, ...)
{
    // GPU-side wait (non-blocking for CPU)
    gr::cuda::wait_for_work(detail(), d_stream);

    // Launch kernels on managed stream
    my_kernel<<<grid, block, 0, d_stream>>>(in, out, noutput_items);

    // Signal completion to pipeline
    gr::cuda::mark_work_done(detail(), d_stream);
    return noutput_items;
}

Sources: include/gnuradio/cuda/cuda_block.h:60-85, include/gnuradio/cuda/cuda_block_helper.h:81-86

Python/CuPy Block Implementation

Python blocks inheriting from cuda.sync_block or cuda.basic_block have this synchronization handled automatically by the wrapper. The wrapper provides CuPy arrays representing device memory directly to the work() function.

Sources: README.md:46-70

Performance Significance

The zero-copy architecture allows CuPy-based Python blocks to achieve performance within 5% of hand-written C++ CUDA blocks. By avoiding synchronous cudaMemcpy and CPU-side blocking, the pipeline maximizes the utilization of the GPU's copy engines and compute units.

Sources: README.md:14-16, README.md:103-115