Writing C++ CUDA Blocks

Relevant source files

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

• include/gnuradio/cuda/cuda_block.h
• include/gnuradio/cuda/cuda_block_helper.h
• README.md
• lib/CMakeLists.txt
• lib/null_source_impl.h
• include/gnuradio/cuda/load.h

Writing C++ CUDA blocks in the gr-cuda module allows for high-performance signal processing by leveraging the GPU while maintaining a seamless integration with the GNU Radio scheduler. The core of this integration is the cuda_block base class, which provides managed CUDA streams and works in tandem with cuda_buffer to eliminate manual host-to-device (H2D) and device-to-host (D2H) memory copies.

The system utilizes a zero-copy GPU pipeline where GPU kernels run back-to-back. Synchronization between blocks is handled via specialized helper functions that manage CUDA events on a per-stream basis, ensuring that data is ready before a kernel executes and that memory is safe to overwrite after execution.

Sources: README.md:12-16, include/gnuradio/cuda/cuda_block.h:16-24

Architecture and the Base Class

C++ GPU blocks must use multiple inheritance, inheriting from both a standard GNU Radio block class (like gr::sync_block) and the gr::cuda::cuda_block mixin. The cuda_block class manages the lifecycle of a non-blocking CUDA stream used for all asynchronous operations within the block.

Key Class Members

The cuda_block class provides several protected members to assist with kernel launches and execution:

Member	Type	Description
d_stream	cudaStream_t	A non-blocking CUDA stream created in the constructor and destroyed in the destructor.
d_min_grid_size	int	The minimum grid size required for full GPU occupancy.
d_block_size	int	The thread block size for kernel launches.

Sources: include/gnuradio/cuda/cuda_block.h:32-38, include/gnuradio/cuda/cuda_block.h:103-107

Class Hierarchy

The following diagram illustrates the relationship between the standard GNU Radio architecture and the CUDA extension.

Sources: include/gnuradio/cuda/cuda_block.h:32-38, lib/null_source_impl.h:20-22

Synchronization Workflow

Synchronization in gr-cuda is explicit but managed through two primary helper functions: wait_for_work and mark_work_done. These functions perform GPU-side event synchronization using cudaStreamWaitEvent and do not block the calling CPU thread.

wait_for_work()

This function must be called at the start of the work() or general_work() method. It ensures:

1. Input Readiness: All input buffers' data is ready on the device (upstream producer finished writing).
2. Output Safety: All output buffers are safe to write (downstream consumers finished reading).

mark_work_done()

This function must be called after all GPU kernels have been enqueued. It:

1. Signals Consumers: Records the DEV_READY event on output buffers so downstream blocks can proceed.
2. Signals Producers: Records the READ_DONE event on input buffers so upstream blocks know they can overwrite the data.

Sources: include/gnuradio/cuda/cuda_block_helper.h:35-50, include/gnuradio/cuda/cuda_block_helper.h:80-95

Execution Flow Diagram

Implementation Patterns

Standard work() Implementation

For blocks where the number of input items equals the number of output items, the work() method follows this pattern:

int my_block_impl::work(int noutput_items,
                        gr_vector_const_void_star& input_items,
                        gr_vector_void_star& output_items)
{
    // 1. Add GPU-side dependencies
    gr::cuda::wait_for_work(detail(), d_stream);

    // 2. Access device pointers and launch kernel
    auto in  = static_cast<const float*>(input_items[0]);
    auto out = static_cast<float*>(output_items[0]);
    my_kernel<<<grid, block, 0, d_stream>>>(in, out, noutput_items);

    // 3. Record events for synchronization
    gr::cuda::mark_work_done(detail(), d_stream);
    return noutput_items;
}

Sources: include/gnuradio/cuda/cuda_block.h:60-76

General Work and Explicit Produce/Consume

In general_work() blocks that call produce() or consume() explicitly, mark_work_done() must be called before produce() or consume(). This is because produce() immediately advances the write pointer; if the event is not recorded first, downstream blocks might see stale data.

Sources: README.md:105-115, include/gnuradio/cuda/cuda_block_helper.h:92-96

Build System Integration

To include C++ CUDA blocks in the build, the source files (both .cc and .cu) must be added to the cuda_sources list in the library's CMakeLists.txt. The project uses project(gr-cuda CXX C CUDA) to enable the CUDA compiler.

Relevant build targets include:

• gnuradio-cuda: The shared library containing the implementations.
• CUDA::cudart, CUDA::cuda_driver, and CUDA::cufft: Linked libraries required for CUDA functionality.

Sources: CMakeLists.txt:16-21, lib/CMakeLists.txt:17-50

Conclusion

Writing C++ CUDA blocks requires adhering to the specific synchronization pattern involving wait_for_work() and mark_work_done(). By inheriting from cuda_block, developers gain access to a managed stream that allows GPU operations to overlap with CPU processing and other GPU memory transfers, achieving maximum throughput in complex flowgraphs.