Optimizing Interprocess GPU Frame Transport in Holoscan SDK
A technical review of how the Interprocess camera application evolved from a functional multi-process prototype into a handle-based HoloMQ transport path on NVIDIA Jetson.
Introduction
Moving a video frame between two GPU applications is not automatically zero-copy, even when both use CUDA. Intermediate tensor allocations, transport staging, receiver-side memory copies, and CUDA synchronization can introduce unnecessary overhead, increasing latency and resource consumption in real-time embedded vision pipelines.
These operations are not always visible at the application level. A seemingly straightforward pipeline:
The Holoscan SDK Interprocess application was developed to transfer live IMX274 camera frames between two independent processes on NVIDIA Jetson. The sender captures RG10 frames, converts them to BGRA using GPU processing, and publishes them through HoloMQ. The receiver imports the CUDA tensor and displays it using X11DisplayOp.
Although the initial implementation demonstrated functional interprocess communication, Nsight Systems profiling revealed redundant device-to-device (D2D) memory copies. To address these inefficiencies, RidgeRun introduced four key improvements:
Direct GPU memory export: The sender writes converted frames directly into HoloMQ-exportable CUDA allocations, eliminating intermediate staging copies.
Retained tensor ownership: The receiver forwards imported tensors while preserving their HoloMQ read claims, avoiding additional memory copies and synchronization.
Optimized white balance: Fixed gains eliminate per-frame accumulator resets and reduction operations.
Simplified communication: A single video endpoint replaces the previous endpoint-switching mechanism.
Together, these optimizations eliminate frame-payload copies at the interprocess boundary, significantly reducing glass-to-glass latency, GPU utilization, and CPU overhead. While the transport path achieves zero-copy communication, the X11 display stage still requires a device-to-host transfer.
Application Architecture
Interprocess consists of two Holoscan executables.
Two small control endpoints complement the video path:
receiver_ready prevents the sender from consuming its frame budget before the receiver has attached.
sender_done allows the receiver to terminate cleanly when the sender reaches its frame limit.
Only the sender accesses the camera and performs Bayer conversion. The receiver receives an already converted BGRA tensor.
The Original Memory Path
The original application used CUDA-aware components throughout the graph, but CUDA visibility alone did not prevent redundant transfers.
Camera capture
V4L2UserPtrSourceOp allocated pinned USERPTR buffers and exposed their CUDA device aliases. The IMX274 driver wrote RG10 data into those buffers, and the conversion operator read them from CUDA without first copying the complete camera frame from host to device memory.
Bayer conversion output
RawBayerToBgraOp allocated a normal CUDA tensor from its Holoscan runtime output pool. Its CUDA kernel performed optical-black correction, white balance, Bayer interpolation, gain, gamma encoding, and BGRA generation into that tensor.
Sender-side export staging
The exported topic originally used:
foreign_emit_policy = InterfaceTensorEmitPolicy::kCopyIfNeeded; |
The BGRA tensor belonged to the conversion operator's runtime pool rather than the HoloMQ export pool. When a subscriber was active, the publisher could not export that allocation directly, so it staged the frame into transport-owned exportable CUDA memory. That staging operation introduced a full-frame device-to-device copy.
Receiver-side bridge copy
HoloMQ transported the export allocation through CUDA handles and metadata. The payload itself was not copied while crossing the process boundary.
However, the original InterpipeSrcOp treated the imported allocation as a temporary external claim. It performed the following operations for every frame:
Allocate a receiver-local CUDA tensor.
Copy the imported tensor with cudaMemcpyAsync(..., cudaMemcpyDeviceToDevice).
Synchronize the CUDA stream.
Release the HoloMQ read claim.
Emit the local tensor to the display operator.
The synchronization was necessary because the imported allocation could not be released before the copy completed. This protected correctness but added another full-frame D2D transfer and a per-frame synchronization point.
Dynamic white balance
The original Bayer operator calculated gray-world gains for every frame. It reset three channel accumulators with cudaMemsetAsync, ran a reduction over the Bayer image, calculated the gains, and then used them during demosaicing.
The accumulator was small, so the memset did not move a full video frame. Nevertheless, it created one visible CUDA memory operation per frame and added reduction work to the conversion path.
Original copy model
The effective path began with the camera writing the RG10 frame through DMA into a mapped USERPTR buffer. The Bayer kernel then generated the BGRA image in a runtime CUDA tensor.
Because that tensor was not directly exportable, the sender performed a device-to-device copy into the HoloMQ export pool. HoloMQ transported the CUDA allocation handle and tensor metadata across the process boundary, after which the receiver performed another device-to-device copy into a local CUDA tensor. Finally, the X11 display path copied the image from device memory to host memory for presentation.
HoloMQ itself was not copying the payload between processes. The redundant copies existed immediately before and after the transport boundary because the producer and consumer did not preserve the export allocation's ownership.
What Nsight Systems Revealed
An early 600-frame Nsight Systems capture reported:
CUDA memory operation | Count |
Device-to-device memcpy | 1,203 |
CUDA memset | 610 |
Device-to-host memcpy | 1 |
The approximately two D2D operations per frame were consistent with sender-side export staging and the receiver-side bridge copy. The approximately one memset per frame matched the dynamic white-balance accumulator reset.
The display was not active for the complete early capture, which explains why that report contained only one D2H operation. Therefore, this run was useful for diagnosing memory behavior but should not be used as a complete display-performance comparison.
An intermediate implementation retained the imported receiver tensor but still staged the sender output. Its 600-frame report showed:
CUDA memory operation | Count |
Device-to-host memcpy | 602 |
Device-to-device memcpy | 600 |
CUDA memset | 10 |
Host-to-device memcpy | 1 |
This was an important validation point. Removing the receiver bridge copy reduced the path from approximately two D2D copies per frame to one. The remaining 600 D2D operations identified the sender-to-export-pool staging path as the next target.
Optimization 1: Fixed White-Balance Gains
The conversion operator was extended to accept fixed red, green, and blue gains. Interprocess uses:
R = 0.833
G = 1.000
B = 0.833
The gains are copied to device memory once during startup. The per-frame accumulator reset, channel reduction, and gain calculation are skipped. This change was appropriate for a controlled benchmark with a fixed camera and lighting setup. It also makes the amount of GPU work more deterministic from frame to frame.
Fixed gains are not universally better than dynamic white balance. A product that must adapt to changing illumination may still need automatic gain calculation. The optimization is a workload decision, not a claim that dynamic white balance is unnecessary in every camera pipeline.
Optimization 2: Produce Directly Into Exportable Memory
The central sender optimization was to change where RawBayerToBgraOp writes its output.
Instead of generating BGRA in a normal runtime tensor and asking the publisher to stage it later, the operator can acquire a tensor from an ExportableHoloMQPool. The pool contains persistent CUDA device allocations created by HoloMQ and exposed to Holoscan as tensors.
The conversion kernel writes directly into the selected exportable allocation:
RG10 mapped input
-> Bayer/BGRA CUDA kernel
-> HoloMQ-exportable CUDA buffer
The topic contract was also changed from kCopyIfNeeded to:
foreign_emit_policy = InterfaceTensorEmitPolicy::kRequireExportable; |
This policy is important. kCopyIfNeeded allows the graph to preserve functionality by inserting a staging copy when the allocation is not exportable. kRequireExportable makes a non-exportable tensor an error instead of silently reintroducing the copy.
The export pool uses persistent slots rather than allocating a new CUDA frame buffer on every activation. A slot remains busy while any tensor handle still refers to it. It becomes reusable only after the imported generation has retired.
CUDA stream readiness is attached to the tensor before it is emitted. HoloMQ can then export the allocation with the correct producer-stream dependency instead of forcing a CPU-side synchronization after the conversion kernel. This design removes the sender-side full-frame D2D staging operation.
Optimization 3: Retain the Imported Tensor Claim
The receiver optimization required a change in ownership, not only the removal of a cudaMemcpyAsync call.
An imported HoloMQ tensor is valid only while its read claim remains alive. If InterprocessSrcOp emitted a raw pointer and destroyed the received sample at the end of compute(), downstream operators could access memory that the publisher had already reclaimed.
The retained-import mode solves that lifetime problem by moving the complete sample to the output:
output.emit(std::move(sample)); |
The tensor's memory reference retains the HoloMQ read claim while the sample is present in graph queues or held by downstream consumers. The claim is released only when the last reference disappears.
Consequently, InterprocessSrcOp no longer needs to:
Allocate a local frame tensor.
Copy the imported payload device-to-device.
Synchronize the stream before returning from compute().
The receiver still declares a runtime output-pool obligation so Holoscan can materialize the port and CUDA-device binding, but retained samples are not allocated from or copied into that pool.
This is the key distinction between forwarding a pointer and forwarding ownership. The optimized implementation preserves the transport claim for the full downstream lifetime.
Final Copy Behavior
The final pipeline handles each frame as follows:
The camera writes RG10 data into a CUDA-mapped USERPTR buffer.
The Bayer kernel reads the mapped input and generates BGRA directly in a HoloMQ-exportable CUDA allocation.
HoloMQ transfers the allocation handle and tensor metadata to the receiver without copying the frame payload.
InterprocessSrcOp forwards the imported sample and preserves its read claim.
X11DisplayOp performs the D2H transfer required by its host-backed X11 presentation path.
A final 600-frame Nsight report contained:
CUDA memory operation | Count |
Device-to-host memcpy | 600 |
Device-to-device memcpy | 0 |
CUDA memset | 21 |
Host-to-device memcpy | 1 |
The one H2D operation corresponds to startup-sized state such as the fixed white-balance gains rather than a frame payload. The remaining memsets are initialization or runtime bookkeeping and no longer occur once per frame. The 600 D2H operations correspond to one X11 presentation transfer for each displayed frame.
The result is zero-copy across the interprocess transport boundary, but it is not end-to-end zero-copy. X11 still needs a host-visible image. A display path based on a directly importable GPU or DRM allocation would be a separate optimization with different ownership and presentation requirements.
Before And After
Stage | Original implementation | Final implementation |
Camera capture | V4L2 USERPTR mapped for CUDA | Same efficient USERPTR path |
White balance | Gray-world reduction and memset every frame | Fixed gains uploaded at startup |
Bayer output | Normal Holoscan runtime CUDA tensor | HoloMQ-exportable CUDA tensor |
Export policy | kCopyIfNeeded | kRequireExportable |
Sender transport preparation | D2D staging into export pool | Conversion kernel writes into export pool |
Process boundary | CUDA handle and metadata | CUDA handle and metadata |
Receiver ownership | Temporary imported claim | Claim retained through tensor lifetime |
Receiver output | Local allocation plus D2D copy and synchronization | Imported sample emitted directly |
Endpoint policy | Two endpoints with switching | One video endpoint |
X11 presentation | D2H copy | D2H copy remains |
Holoscan Interprocess Performance Results
The original and final Interprocess implementations produced the following project measurements:
The optimized implementation reduced mean G2G latency by 8.55 ms while also lowering the measured CPU, GPU, and resident-memory usage. CPU utilization showed the largest relative reduction, falling by 64.52%. GPU utilization decreased by 8.56 percentage points, and RSS decreased slightly by 1.08 MB. The standard deviation was also lower for all four metrics, indicating more stable measurements in the final configuration, particularly for GPU utilization, whose deviation fell from 10.25 to 0.81 percentage points.
These improvements are consistent with removing the per-frame D2D staging and receiver bridge copies, the receiver-side synchronization, and the dynamic white-balance work. They should still be interpreted as whole-configuration results rather than as the isolated cost of a single code change.
Key Engineering Insights
Zero-copy is an ownership property
Sharing a CUDA pointer is not sufficient. The producer allocation must be exportable, and its lifetime must extend until every receiver-side consumer releases it.
Fallback policies can hide copies
kCopyIfNeeded is useful for compatibility, but it can obscure transport staging. A strict exportable-memory policy makes the intended memory behavior enforceable.
The producer should write into the final transport allocation
Eliminating a copy is easier when the processing kernel generates its result in the memory that the next boundary already owns. Copying into the correct pool after processing is still a copy.
Imported claims must follow downstream references
The receiver bridge copy originally existed because it was the simplest safe way to end the external claim. Retaining the sample converts that safety requirement into explicit lifetime management.
Conclusion
The Interprocess optimization did not replace HoloMQ or remove the source operator that demonstrates interprocess transport. It changed how the application participates in HoloMQ's memory model.
On the sender, RawBayerToBgraOp now produces directly into persistent exportable CUDA storage. At the topic boundary, kRequireExportable prevents a fallback staging copy. On the receiver, InterprocessSrcOp keeps the imported claim alive and emits the same tensor handle instead of creating a local payload copy.
Nsight Systems captured the progression clearly: the initial implementation showed approximately two D2D transfers per frame, the intermediate implementation showed one, and the final 600-frame run showed none. The only full-frame transfer left in the measured path was the D2H copy required for X11 presentation.
These improvements are particularly relevant to embedded vision systems composed of independent processing stages. Industrial inspection, robotics, intelligent cameras, and real-time video analytics can benefit from minimizing memory transfers and preserving GPU resources for image processing and inference.
Although the measured results are specific to this implementation, the underlying principles may be applied to other applications. Efficient interprocess GPU communication depends on where the memory was allocated, strict transfer policies, stream readiness and syncing, and explicit ownership lifetime. Once these elements are aligned, process boundaries can exchange handles and metadata while the frame payload remains in the same CUDA allocation.
In conclusion, the application demonstrates how exportable CUDA memory, explicit ownership management, and profiling-driven development can improve pipeline efficiency without requiring a complete rewrite of the app.
Ready to optimize your Edge AI applications?
Contact RidgeRun to discuss your project.
