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
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
Architectural analysis of my repositories: app.igkho-systems.dev
Contacts:
Zero-Host-Copy Inference: Keeping the CPU Out of the Hot Path in a C++/CUDA Pipeline (165 FPS) was originally published in Towards AI on Medium, where people are continuing the conversation by highlighting and responding to this story.