Zero-Host-Copy Inference: Keeping the CPU Out of the Hot Path in a C++/CUDA Pipeline (165 FPS)

A developer published a bare-metal C++/CUDA video inference pipeline, ZeroHostCopyInference, on GitHub that sustains 165 FPS by keeping frames entirely in GPU memory and eliminating all runtime calls to cudaMalloc, cudaFree, new and delete from the hot path. The InferencePipeline class splits work across three asynchronous worker threads — Source, Detector and Sink — connected by a bounded SafeQueue built on std::mutex and std::condition_variable that applies backpressure by blocking the Source thread when the queue is full, preventing out-of-memory crashes when a decoder pushing 300 FPS outpaces a 165 FPS inference engine. The author chose the custom implementation over NVIDIA DeepStream to avoid GStreamer dependencies and gain full control over the threading model.

When modern INT8-quantized neural networks execute in single-digit milliseconds, the system bottleneck shifts from raw compute to I/O and pipeline infrastructure. On discrete GPUs, constantly bouncing high-resolution video frames and intermediate payloads like bounding box tensors across the PCIe bus severely bottlenecks overall throughput. In a Zero-Host-Copy architecture, the data life cycle stays strictly on the GPU. From the initial pre-decode disk read to the final post-encode disk write, frames reside exclusively in device memory, heavily minimizing H2D Host-to-Device and D2H transfers. Furthermore, a strict Zero-Copy pipeline passes payloads between stages strictly via pointers, completely eliminating deep copies. However, redundant PCIe transfers aren’t the only silent performance killers; dynamic memory allocation during runtime is just as fatal. Calling cudaMalloc triggers implicit synchronization, blocking the host CPU on the GPU and wrecking pipeline asynchrony. To max out throughput, a high-performance system requires a strict three-phase life cycle: init, warm-up, and the hot path. During initialization, the application pre-allocates all memory pools, buffers, and queues to their target depths. The warm-up phase then pushes a full processing queue, triggering any lazy allocations or buffer expansions needed for real-world workloads. Once the pipeline drops into the hot path after flushing a few initial frames, dynamic allocation must hit absolute zero. This architecture strictly forbids runtime calls to cudaMalloc or new, relying entirely on recycled memory blocks to guarantee a flat, stable frame rate. Why didn’t I choose NVIDIA DeepStream? DeepStream is great, but building custom GStreamer plugins in C++ is notoriously heavy/inflexible. I wanted to implement a bare-metal pipeline with zero overhead, full control over the threading model, and no GStreamer dependencies to maximize performance and portability. You can view the full bare-metal pipeline implementation on GitHub here: ZeroHostCopyInference https://github.com/Igkho/ZeroHostCopyInference Fast individual components aren’t enough for a high-performance video pipeline. If the hardware decoder pushes 300 FPS but the inference engine pulls 165 FPS, a naive pipeline will OOM Out-Of-Memory and crash the system. To throttle the data flow and prevent memory exhaustion, I designed the InferencePipeline class. It splits the workload across three asynchronous worker threads—Source, Detector, and Sink—operating as a chained producer-consumer: To safely bridge these stages, I built a custom SafeQueue backed by std::mutex and std::condition variable. Standard thread-safe queues typically grow unbounded. To guarantee system stability, my queue enforces strict backpressure by hard-capping the queue depth. If the Detector lags behind the Source, the bounded queue fills up and blocks the Source until slots free up. The OS suspends the Source thread entirely, preventing CPU busy-waiting spin-locking , unbounded queue growth, and memory exhaustion. template <typename T class SafeQueue { // ... void Push T value { std::unique lock<std::mutex lock mutex ; // Suspend Producer thread if queue is full Backpressure if capacity 0 { cond full .wait lock, this { return queue .size < capacity ; } ; } queue .push std::move value ; cond empty .notify one ; } // ...}; Sustaining 165 FPS requires absolutely zero calls to cudaMalloc, cudaFree, new, or delete in the hot path. Runtime memory allocation fragments device memory and stalls the GPU pipeline via implicit synchronization. The fix is aggressive memory pooling via recycle queues. InferencePipeline implements a closed-loop memory manager. Payloads move through the execution queues preProcessQueue , postProcessQueue , but once the Sink worker finishes flushing a frame, we don't free the memory. Instead, the pipeline pushes the exhausted BatchData and BatchDetections buffers back into recycleQueue and resultRecycleQueue . When the producer thread needs to fetch a new frame, it simply pops a pre-allocated buffer from the pool: CudaError InferencePipeline::Run { // ... thread initialization ... while process && running { BatchData batch; // Try to recycle an old batch buffer to save allocation recycleQueue .TryPop batch, std::chrono::milliseconds 0 ; // If recycle queue is empty 'batch' is safely initialized once CUDA TRY source - GetNextBatch batch, batchSize , process ; // ... if process { // ... preProcessQueue .Push std::move batch ; } } // ...} Takeaway: Combining strict backpressure with cyclic memory pools guarantees all cudaMalloc calls occur during the warm-up phase. In the hot path, the pipeline strictly recycles pre-allocated memory blocks, keeping both discrete GPUs and Jetson Orin modules saturated at peak throughput. The fundamental architectural split handled by this project is the difference between a discrete GPU tethered via PCIe and the Jetson Orin’s Unified Memory Architecture UMA . On discrete setups, the CPU and GPU have physically isolated memory banks. Crossing the PCIe bus is an expensive operation requiring careful DMA orchestration. On the Jetson UMA, the ARM cores and the GPU share the exact same physical RAM chips. To abstract these differences, I centralized memory behavior in a custom Block.h allocator. It maps logical memory aliases based on data locality—where the data lives and who needs it—to the correct hardware-specific allocation strategy. This yields three universal memory types: enum class MemoryType { Device, // Standard VRAM — fastest for pure GPU compute Pinned, // Page-locked host memory staging buffer for PCIe transfers on RTX ZeroCopy // Mapped memory shared pointer for CPU/GPU — optimal for Jetson I/O }; ifdef PLATFORM JETSON// --- Jetson Implementation Unified Memory Architecture / UMA ---constexpr MemoryType BoundaryMemType = MemoryType::ZeroCopy;constexpr MemoryType HostStagingMemType = MemoryType::ZeroCopy; else// --- x64 Implementation Discrete GPU over PCIe ---constexpr MemoryType BoundaryMemType = MemoryType::Device;constexpr MemoryType HostStagingMemType = MemoryType::Pinned; endif True zero-host-copy inference on x64 requires keeping video frames out of standard host RAM entirely. In FFmpegSource.cpp, I initialize an AV HWDEVICE TYPE CUDA hardware context and explicitly bind it to our primary CUDA execution context. // Get the current primary CUDA contextCUDA TRY cudaFree 0 ;CUcontext primary ctx = nullptr;CUresult cuRes = cuCtxGetCurrent &primary ctx ;// ...// Allocate the FFmpeg hardware device contexthwDeviceCtx = av hwdevice ctx alloc AV HWDEVICE TYPE CUDA ;// Manually override the FFmpeg context with our primary CUDA contextAVCUDADeviceContext internal ctx = AVCUDADeviceContext AVHWDeviceContext hwDeviceCtx - data - hwctx;internal ctx- cuda ctx = primary ctx;//...// Force the decoder to use hardware acceleration tied to our CUDA context,// guaranteeing output frames land directly in GPU memorydecCtx - hw device ctx = av buffer ref hwDeviceCtx ; In the hot path, avcodec receive frame yields an AVFrame gpuFrame formatted as AV PIX FMT CUDA. The data 0 and data 1 pointers are already valid device pointers pointing to the NV12 planes. We async-copy these directly into our VRAM accumulation batch buffers. Once the batch fills, a custom CUDA kernel BatchedNV12ToRGBPlanar handles the color space conversion strictly in-device. int ret = avcodec receive frame decCtx , gpuFrame ;//...CUDA TRY cudaMemcpyAsync device nv12 y buf idx .data + framesCollected y plane size , gpuFrame - data 0 , y plane size, cudaMemcpyDeviceToDevice, 0 ;CUDA TRY cudaMemcpyAsync device nv12 uv buf idx .data + framesCollected uv plane size , gpuFrame - data 1 , uv plane size, cudaMemcpyDeviceToDevice, 0 ;//...CUDA TRY BatchedNV12ToRGBPlanar device nv12 y buf idx .data , device nv12 uv buf idx .data , framePitch, outBatch.deviceData.data , framesCollected, actualW, actualH, outW, outH, cuda stream ; While standard H.264 streams cap out around 165 FPS, switching the source to H.265 HEVC pushes throughput past 190 FPS and significantly flattens performance dips at larger batch sizes e.g., batch 16 . Why? Even though HEVC is mathematically heavier, the Ampere NVDEC hardware block processes 64x64 CTUs Coding Tree Units compared to H.264’s 16x16 macroblocks. For the same resolution, NVDEC processes drastically fewer blocks. This sharply reduces internal hardware scheduling overhead, allowing the hardware decoder to run perfectly async alongside TensorRT, even at deep batch sizes. Since FFmpeg hardware decoding lacks native support on Jetson UMA, the pipeline falls back to reading directories of JPEG frames -i ./frames/ . On discrete GPUs, this fallback is implemented via NVJpegSource and NVJpegSink using the nvJPEG library. The CPU parses the bitstream nvjpegDecodeJpegHost , and the payload transfers asynchronously to the GPU for final decoding nvjpegDecodeJpegTransferToDevice / nvjpegDecodeJpegDevice . However, Jetson SOCs feature dedicated JPEG hardware engines that nvJPEG doesn't fully exploit. To redline the Orin's throughput, I wired up the Jetson Multimedia API MMAPI via MMAPIJpegSource and MMAPIJpegSink. MMAPI demands rigorous memory lifecycle management. To saturate the silicon, this class defines a logical pool of hardware decoders and encoders, dispatching concurrent decode/encode threads via std::async. struct DecodeResource { static constexpr int HW DECODERS PER BUFFER = 2; // ...}; To overlap hardware decoding with inference processing, I use a strict ping-pong double-buffering scheme. Because of this, all staging structures are two-element arrays. class MMAPIJpegSource : public ISource { //... std::unique ptr<CudaEvent dma complete event 2 ; //... std::vector<DecodeResource resource pool 2 ; std::vector<std::unique ptr<NvJPEGDecoder hw decoders 2 ; std::vector<std::future<std::vector<MMAPIDecodeStatus futures 2 ; int frames in buffer 2 = {0, 0}; //...}; The MMAPI Memory Lifecycle: CudaError MMAPIJpegSource::MapAndConvert DecodeResource& res, ... { //.. cudaEglFrame eglFrame; // Ultra-fast transfer: simply map device pointers from the already registered resource CUDA TRY cudaGraphicsResourceGetMappedEglFrame &eglFrame, res.egl cuda resource, 0, 0 ; // Extract standard Pitch-Linear pointers const uint8 t srcY = static cast<const uint8 t eglFrame.frame.pPitch 0 .ptr ; const uint8 t srcUV = static cast<const uint8 t eglFrame.frame.pPitch 1 .ptr ; int srcPitch = eglFrame.frame.pPitch 0 .pitch; // Read directly from the surface CUDA TRY NV12ToRGBPlanar srcY, srcUV, srcPitch, ... ; //...} Ultimately, porting this double-buffering architecture and logical decoder pool over to the x64 NVJpegSource and NVJpegSink made the JPEG scenario the outright performance winner, peaking at 215 FPS. This happens because JPEG files are independent, intra-coded frames. A CPU thread pool can parse them and blast them across the PCIe bus concurrently, completely hiding transfer latency. Unlike FFmpeg video streams—which enforce sequential decoding due to inter-frame dependencies—ingesting a directory of JPEGs parallelizes flawlessly. Many modern inference pipelines fall into a common trap: they successfully run neural network compute on the GPU, but immediately copy the metadata bounding boxes, class IDs, probabilities back to host memory to run trackers like SORT or ByteTrack. This host-side processing introduces massive synchronization overhead and stalls the pipeline. To hit 165 FPS, I designed the tracking logic to operate purely in VRAM, eliminating device-to-host transfers. The ObjectTracker class uses a TypedBlock<TrackState residing strictly in device memory. This ensures track history positions, velocities, class probabilities never crosses the PCIe bus. Detection outputs that explicitly require host access are passed via BoundaryMemType arrays. These default to device memory on x64 architectures and leverage ZeroCopy pinned/mapped memory on Jetson UMA architectures, ensuring optimal data placement across both hardware targets. Initially, using standard libraries like thrust::sort on the hot path introduced implicit cudaMalloc calls. To completely eliminate runtime allocations, I replaced this with cub::DeviceMergeSort::SortKeys, passing a pre-allocated sortWorkspace buffer managed by a FrameResources memory pool. This sort feeds into the Non-Maximum Suppression NMS pass to filter raw detector logits before tracking begins. The tracking pipeline itself is split into discrete CUDA kernels, executed asynchronously and sequentially per frame in the batch. Before matching, the tracker must predict existing object locations. This kernel simply iterates over active tracks and applies their current velocity vectors vx, vy to the bounding box coordinates. It also increments the timeSinceUpdate and missedFrames counters—these reset to zero upon a successful match. Instead of a heavy, sequential Hungarian algorithm, this kernel leverages massive GPU parallelism to greedily match detections to predicted track states. Box IoU calculation happens directly in-kernel. To prevent false matches during transient occlusion, the kernel applies a dynamic threshold: if a track missed previous frames, the IoU requirement to confirm a match tightens from 0.30 to 0.45. In dense scenes, N-to-1 data races occur when multiple independent detections simultaneously claim the same track. To resolve these collisions strictly in device memory without host sync , I implemented an atomic claim mechanism. The kernel packs the IoU high 32 bits and the detection index low 32 bits into a single 64-bit integer. It then uses atomicMax to lock the claim in a dedicated trackClaims buffer. // Pack claim: high 32 bits = IoU, low 32 bits = indexunsigned long long iouBits = unsigned long long float as uint bestIOU ;unsigned long long claim = iouBits << 32 | unsigned int idx;atomicMax &trackClaims bestTrackIdx , claim ; This stage starts by validating track ownership. Threads read the final value from the trackClaims buffer and compare their thread index against the winning index. If a thread loses the race another overlapping object had a higher IoU , it resets its match index trackIdx = -1 and falls through to the new track creation logic. This completely prevents dropping detections during heavy occlusion. Next, the kernel applies real-world physical constraints to the winning tracks. Bounding box velocity estimation is notoriously noisy, especially when objects enter or exit the frame. When a box hits the screen edge, the object gets partially cropped, shifting its center of mass and causing severe velocity spikes. To mitigate this, the kernel analyzes margins. If an object touches the left edge det.x - det.w 0.5f <= margin , the kernel ignores the left edge and relies entirely on the right edge to calculate the instantaneous velocity raw vx . Once computed, this raw velocity passes through an exponential moving average EMA filter using TRACK VELOCITY FILTER RATIO to stabilize the trajectory: // Adaptive velocity logic Y-axis example bool touchingTop = det.y - det.h 0.5f <= margin;bool touchingBottom = det.y + det.h 0.5f = heightF - margin ;if touchingTop && touchingBottom { // Top edge clipped. Trust the BOTTOM edge. float det y b = det.y + det.h 0.5f; float t y b = t.y + t.h 0.5f; raw vy = det y b - t y b + t.vy;}// ... logic repeats for the other edge/center ...// Apply EMA filtert.vy += t.vy == 0 ? raw vy : TRACK VELOCITY FILTER RATIO raw vy - t.vy ; If an object is temporarily occluded, we don’t want its ID to thrash. The system maintains “ghost tracks.” If an existing track is missed for several frames, the kernel dynamically injects a ghost detection back into the pipeline. This ensures the renderer keeps drawing the predicted bounding box while the object coasts behind an obstacle. However, to prune stale tracks correctly, the system must filter out outliers by computing the global average velocity of all established tracks in the current frame. Relying on standard atomicAdd to shared memory here creates massive thread contention. Instead, the kernel uses a two-tier reduction architecture paired with a grid-stride loop: threads first accumulate velocities in registers using the shfl down sync warp primitive. inline device float warp reduce sum float val { pragma unroll for int s = 16; s = 1; s = 1 { val += shfl down sync 0xFFFFFFFF, val, s ; } return val;} The grid-stride loop ensures the kernel scales seamlessly and safely processes bounding box arrays exceeding the hard limit of 1024 threads per block. int idx = threadIdx.x + blockDim.x blockIdx.x;// ...// The grid-stride loopint sidx = idx;while sidx < totalTracks / ... / { // ... sidx += gridDim.x blockDim.x;} Lane 0 in each warp writes its partial sum to a dynamically allocated shared memory block extern shared float s reduce . Warp 0 then performs the final block-level reduction. Trajectories deviating too far from this global mean are immediately flagged for deletion t.age = -999 . Finally, we must strip dead tracks from the array. Standard functions like thrust::remove if introduce heavy synchronization barriers. To bypass this entirely, I wrote a custom CompactTracksKernel. While the earlier tracking stages run per-frame, CompactTracksKernel runs once per batch. The kernel evaluates active tracks and safely packs the surviving elements into a pre-allocated auxiliary buffer tempTracks using an atomic counter. The host finalizes compaction via an O 1 pointer swap tracksBuffer.raw .swap tempTracksBuffer.raw . This instantly readies the state for the next step without a single device-to-device memcpy. You can’t optimize what you can’t measure in high-performance C++/CUDA systems. To isolate bottlenecks without introducing synchronization overhead, I built a modular profiling system in PerformanceTimer.h: Because the pipeline runs asynchronously across three dedicated threads, simply summing the component runtimes yields a total exceeding actual wall-clock time. To provide an accurate representation of system load, PipelineStats::PrintReport calculates the relative active work share by dividing each stage's execution time by the total active component time. Here is a sample report from a real run: ======================================================= INFERENCE PIPELINE REPORT =======================================================Total frames processed: 122Total batches processed: 81. Initialization: 745 ms2. Total pipeline uptime: 570 ms 214.04 FPS 3. Decode Source : - Total: 280.83 ms - Avg/batch: 35.10 ms - Avg/frame: 2.30 ms - Share: 30.23% active work 4. Inference Detector : - Total: 465.24 ms - Avg/batch: 58.15 ms - Avg/frame: 3.81 ms - Share: 50.07% active work 5. Data Sink: - Total: 183.05 ms - Avg/batch: 22.88 ms - Avg/frame: 1.50 ms - Share: 19.70% active work ======================================================= For deep profiling and isolating specific bottleneck kernels or operations, I rely heavily on NVIDIA Nsight Systems. The trace shows that using NVJpegSource achieves the core goal of the asynchronous architecture: dense, virtually gapless packing of TensorRT execution contexts orange blocks . Thanks to double buffering and a pool of logical decoders in NVJpegSource, the GPU never starves waiting for host data. This pushes the primary bottleneck exactly where it belongs: the neural network inference stage. I benchmarked the pipeline on an NVIDIA RTX 3060 Ti and an NVIDIA Jetson Orin Nano. To measure the pipeline’s architectural overhead I/O latency , I used a stub detector that acts as a pure passthrough, forwarding data without running any compute. Running YOLOv8m with explicit INT8 quantization alongside full object tracking and serialization via NVJpegSink. Scenario A: Video Stream FFmpegSource Scenario B: Image Directory NVJpegSource Load Distribution Directory Images : For image directories NVJpegSource utilizing pure intra-frame compression, parallelization scales perfectly, distributing the load concurrently across multiple CPU threads and hardware JPEG decoders. Video streams FFmpegSource are inherently sequential, bottlenecking on the single hardware NVDEC block. The pipeline supports both TensorRT for bare-metal performance and ONNX Runtime for universal model compatibility . Scenario: 1440x1440 input JPEG directory . Model: YOLOv8 Medium YOLOv8m 1024x1024. Note: ONNX Runtime is unsupported. Why Fast Models Are Only Half the Equation The Computer Vision industry traditionally fixates on the neural network itself: tweaking architectures, applying INT8 quantization, and exporting to TensorRT. However, in practice, once inference time drops to single-digit milliseconds, the surrounding infrastructure becomes the absolute bottleneck. To achieve a stable 165+ FPS on a consumer RTX 3060 Ti and squeeze 38 FPS out of an edge Jetson Orin Nano, you must optimize the entire data pipeline end-to-end: I built this project as a deep dive into systems-level C++ and high-performance GPU inference. I am currently open to new opportunities and looking to join a strong engineering team tackling high-load computer vision systems, High-Performance Computing HPC , or distributed inference e.g., TensorRT-LLM, multi-GPU pipelines . I am targeting remote Senior C++ / CUDA Developer or Performance Engineer roles. Full source code GitHub : ZeroHostCopyInference https://github.com/Igkho/ZeroHostCopyInference Architectural analysis of my repositories: app.igkho-systems.dev https://app.igkho-systems.dev/ Contacts: Zero-Host-Copy Inference: Keeping the CPU Out of the Hot Path in a C++/CUDA Pipeline 165 FPS https://pub.towardsai.net/zero-host-copy-inference-keeping-the-cpu-out-of-the-hot-path-in-a-c-cuda-pipeline-165-fps-6da07bfdea9d was originally published in Towards AI https://pub.towardsai.net on Medium, where people are continuing the conversation by highlighting and responding to this story.