From V4L2 MMAP to CUDA-Mapped USERPTR: Reducing Capture-Side GPU Overhead in Holoscan SDK 5

RidgeRun improved the performance of an NVIDIA HSDK inference pipeline by optimizing how video data moves between system components.

In edge-AI applications, efficient frame movement can matter as much as model execution. By reducing unnecessary transfers between camera capture and GPU processing, RidgeRun lowered capture-side GPU overhead and created a more efficient pipeline.

Why High GPU Utilization Doesn't Always Mean an Inference Bottleneck

High GPU utilization does not always mean your inference model is the bottleneck.

For teams developing embedded AI and computer vision applications, optimizing inference is only part of the performance equation. Inefficient memory transfers and camera integration can introduce unnecessary GPU overhead before a frame even reaches the model.

An edge-AI pipeline can show high GPU utilization even when the inference model is not the only source of the workload. Camera capture, image preparation, memory transfers, and synchronization can also consume GPU time before the model receives a frame.

RidgeRun investigated high GPU utilization in a HSDK inference pipeline using V4L2 MMAP capture. Profiling showed that the application copied each captured frame from host memory to a CUDA tensor before preprocessing and inference.

The optimization replaced that path with a CUDA-mapped V4L2 USERPTR source. V4L2 writes each frame into pinned memory allocated with CUDA mapping enabled, while CUDA accesses the same allocation through its device alias. The application can then publish the captured memory to HSDK without an explicit full-frame host-to-device copy for every frame.

The result is a more direct capture path and a useful example of how profiling can guide an HSDK optimization. The change does not remove the Bayer preprocessing or TensorRT workload. It removes the repeated capture-side transfer that the profile identified as the dominant memory-operation cost.

This case study demonstrates why identifying the source of GPU overhead is essential before making changes to an inference pipeline.

The Challenge: Understanding High GPU Utilization in HSDK

The investigation began with an HSDK inference run that used a V4L2 MMAP source. The run reported an average GPU utilization of 54.34%. That value was the symptom, not an explanation. The application still needed to separate the work performed by capture, memory transfers, Bayer preprocessing, and TensorRT inference.

The pipeline under investigation was:

The test used 1920 × 1080 input at 60 FPS and a 600-frame run.

Figure 1. Initial GPU utilization measured with the HSDK MMAP capture path.

GPU utilization alone cannot identify which stage is responsible for the load. A CUDA kernel summary can show preprocessing and inference activity, but it does not explain every memory transfer. Nsight Systems provides the timeline needed to correlate memory operations with captured frames.

For development teams, this distinction is important. Without visibility into individual pipeline stages, optimization efforts may target inference while leaving capture-side inefficiencies unresolved.

Identifying the Bottleneck with NVIDIA Nsight Systems

RidgeRun analyzed the application with NVIDIA Nsight Systems using the corresponding profile.nsys-rep report. The analysis followed four steps:

1. Inspect the CUDA GPU Memory and CUDA API sections.

2. Search for CUDA memcpy Host-to-Device and record the operation count, total size, and total time.

3. Inspect the CUDA GPU Kernel Summary to separate Bayer preprocessing from TensorRT execution.

4. Correlate the memory operations with the application timeline and the number of captured frames.

Figure 2. Nsight Systems summary for the HSDK MMAP run.

The MMAP trace contained the following CUDA memory operations:

Operation	Time	Total Time	Count	Average Time
CUDA memcpy Host-to-Device	99.8%	1,925.117 ms	602	3.198 ms
CUDA memset	0.2%	3.552 ms	4,207	0.844 microseconds
CUDA memcpy Device-to-Device	0.0%	0.036 ms	4	8.968 microseconds
CUDA memcpy Device-to-Host	0.0%	0.023 ms	4	5.808 microseconds

The host-to-device transfers accounted for 99.8% of the CUDA memory-operation time in the trace. The application recorded 602 H2D operations, which is consistent with one capture-to-CUDA transfer per frame during the 600-frame run, plus a small amount of setup activity.

The timeline showed the same pattern at the frame level: one host-to-device copy appeared for each captured frame.

Figure 3. Nsight Systems timeline for events 596 to 600 in the MMAP run.

This profile identified the repeated frame copy as the main memory-transfer cost in the capture path. The next step was to inspect how the MMAP source produced the tensor consumed by the rest of the graph.

The findings provided a clear optimization target: reduce unnecessary frame transfers without altering the preprocessing and inference workloads.

Understanding the V4L2 MMAP Capture Path

The original operator uses V4L2 MMAP buffers for capture. It allocates a CUDA tensor, copies each captured frame from the host buffer into that tensor, and sends the resulting tensor to the next operator in the HSDK graph.

auto writer = loan->write(runtime_stream);
cudaMemcpy2DAsync(writer->data(),
	row_bytes_,
	mapped.address,
	bytes_per_line_,
	row_bytes_,
	height_,
	cudaMemcpyHostToDevice,
	stream);

// The driver may overwrite an MMAP buffer after QBUF.
cudaStreamSynchronize(stream);

The relevant memory objects have different owners:

• mapped.address is the V4L2 MMAP buffer.

• writer->data() is the CUDA or device tensor published to HSDK.

• cudaMemcpy2DAsync(..., cudaMemcpyHostToDevice, ...) transfers every frame.

• cudaStreamSynchronize blocks the calling thread until all operations in the CUDA stream are complete. This prevents the driver from reusing the MMAP buffer while CUDA is still reading it.

  
  

This implementation protects the MMAP buffer from being reused too early. It also adds a full-frame host-to-device transfer and a synchronization point to the capture path. The code is correct for the ownership model, but the profile shows the cost of that model when the application processes hundreds of frames.

Optimizing Camera Capture with CUDA-Mapped USERPTR

The alternative source uses V4L2 USERPTR buffers backed by CUDA-mapped pinned memory. The operator allocates the memory, gives the host address to V4L2, and keeps the corresponding CUDA device alias.

void* host_address = nullptr;
cudaHostAlloc(&host_address,
              slot->length,
              cudaHostAllocMapped | cudaHostAllocPortable);

void* device_address = nullptr;
cudaHostGetDevicePointer(&device_address, host_address, 0);

buffer.memory = V4L2_MEMORY_USERPTR;
buffer.m.userptr = reinterpret_cast<unsigned long>(host_address);
buffer.length = slot->length;
ioctl(fd, VIDIOC_QBUF, &buffer);

cudaHostAlloc creates pinned host memory with CUDA mapping enabled. cudaHostGetDevicePointer returns the device alias for the same allocation. The V4L2 buffer is then configured with V4L2_MEMORY_USERPTR, so the driver writes captured data into that memory.

After V4L2 fills the buffer, the operator wraps the CUDA alias directly in the HSDK tensor:

auto memory_reference =
	std::make_shared<UserPtrTensorMemoryReference>(slot);

auto tensor = holoscan::make_tensor(
    shape,
    kUInt8Dtype,
    std::nullopt,
    holoscan::MemoryStorageType::kDevice,
    slot->device_address,
    std::move(memory_reference),
    cuda_device_);

The frame path no longer calls cudaMemcpy2DAsync for every captured image. V4L2 writes to pinned system memory, and CUDA reads that allocation through its mapped device alias. The UserPtrTensorMemoryReference keeps the USERPTR slot alive while downstream operators use the tensor. Once those operators release the tensor, the slot can be returned to V4L2.

This lifetime relationship is important. The driver must not receive a buffer back while downstream CUDA work still depends on it. The memory reference connects the V4L2 buffer lifetime to the HSDK tensor lifetime.

The approach illustrates how memory ownership and operator design can influence pipeline efficiency without requiring changes to the underlying inference model.

Validating the Optimization: From 602 H2D Transfers to 2

The new USERPTR implementation was profiled with the same type of Nsight inspection. The trace contained two host-to-device operations instead of one full-frame transfer for each captured frame.

Figure 4. Nsight Systems summary for the HSDK USERPTR run.

The USERPTR trace contained these operations:

Operation	Count	Total Size	Total Time
CUDA memcpy Host-to-Device	2	5.341 MB	0.737 ms
CUDA memcpy Device-to-Device	4	0.004 MB	0.037 ms
CUDA memcpy Device-to-Host	4	0.002 MB	0.023 ms

The two traces show the change clearly:

Capture Path	H2D Operations	H2D Data Transferred
HSDK MMAP	602	Approximately 2.32 GB
HSDK USERPTR	2	5.341 MB

The per-frame capture copy disappeared from the USERPTR profile. The remaining GPU work is expected pipeline work. For example, the trace includes raw_bayer_to_yolo_kernel with 600 instances and 124.839 ms of total time. This kernel preprocesses the RAW Bayer input for inference. It is not evidence of a capture-memory copy.

The key improvement was not eliminating GPU computation, but removing redundant data movement from the capture path.

GPU Utilization Results: MMAP vs. USERPTR

The copy optimization was compared with the initial HSDK MMAP measurement result:

Configuration	GPU Utilization
Holoscan SDK 5 MMAP, initial measurement	54.34%
Holoscan SDK 5 USERPTR	34.95%

Figure 5. GPU utilization after switching from the MMAP path to CUDA-mapped USERPTR.

The initial 54.34% value motivated the investigation. The controlled comparison uses the USERPTR path and reports 34.95% in the same benchmark.

These values describe the tested capture, model, resolution, frame rate, and sampling configuration. They should be used to understand the effect of the capture-memory path, not as a universal performance ranking between Holoscan SDK versions.

For embedded AI development, the results highlight the value of investigating data movement before introducing more complex changes to a processing pipeline. Understanding where GPU resources are consumed helps engineering teams prioritize optimization efforts based on measurable bottlenecks.

Interpreting Nsight Results: Memory Transfers vs. GPU Processing

Nsight observations should be interpreted together with the operator design:

Nsight Observation	Interpretation
One H2D transfer per frame	The source publishes a separate device tensor through a copy-based path.
No repeated full-frame H2D transfers	The capture buffer is being consumed through the CUDA-mapped USERPTR path.
raw_bayer_to_yolo_kernel remains	RAW Bayer preprocessing is still required and runs on the GPU.
High GPU utilization remains	The load can come from preprocessing, TensorRT, memory initialization, or shorter execution time. It is not automatically caused by capture copies.

A useful profiling sequence is to inspect memory transfers first, then GPU kernels, CUDA API calls, and the end-to-end Holoscan SDK flow. Flow data helps explain operator and message latency. Nsight Systems confirms whether physical CUDA memory transfers occur and how they line up with captured frames.

This profiling methodology is particularly valuable when multiple operators, memory types, and processing stages interact. It provides a measurable basis for deciding which part of the pipeline requires optimization.

Why Capture-Side Optimization Matters for Embedded AI Applications

The optimization addresses a common integration problem: a camera pipeline can appear to be inference-bound when the capture path is also moving every frame across the host-to-device boundary.

The MMAP path remains useful when driver-owned buffers and an explicit copy provide the ownership behavior required by the application. The CUDA-mapped USERPTR path is useful when the system can allocate the capture buffers and keep them valid for downstream CUDA processing.

Choosing between the paths requires checking more than the average GPU percentage. The implementation must account for V4L2 buffer ownership, CUDA memory visibility, tensor lifetime, preprocessing, inference, and the point at which the driver can reuse a buffer.

For companies building camera-based AI products, these decisions can influence how efficiently capture, preprocessing, and inference components work together. Addressing data movement at the integration level can be an important part of preparing an embedded vision pipeline for production requirements.

RidgeRun can help teams profile these paths, implement custom HSDK operators, integrate camera drivers with CUDA processing, and validate the resulting pipeline on the target hardware. The same workflow applies to other camera configurations and embedded platforms when the capture and memory-ownership requirements are understood.

Key Takeaways: Building More Efficient HSDK Pipelines

The investigation started with a high GPU utilization value in a HSDK MMAP pipeline. Nsight Systems showed that the main memory-operation cost came from one host-to-device transfer per captured frame. A CUDA-mapped V4L2 USERPTR source removed that repeated copy by allowing V4L2 and CUDA to access the same mapped allocation.

The USERPTR trace reduced H2D activity from 602 operations and approximately 2.32 GB of transferred data to two operations totaling 5.341 MB. The remaining GPU activity belongs to the actual processing pipeline, including RAW Bayer preprocessing and TensorRT inference.

The broader lesson is that effective embedded AI optimization starts with understanding the complete data path, not just the inference stage. Identifying and addressing unnecessary memory transfers can simplify GPU workloads while preserving the processing functionality required by the application.

Optimize Your Embedded AI Pipeline with RidgeRun

Are memory transfers, camera integration, or GPU overhead affecting your embedded AI development?

RidgeRun provides engineering services for Holoscan SDK integration, custom operators, CUDA processing, and performance profiling, helping development teams identify bottlenecks and implement optimizations tailored to their hardware and application requirements.

Ready to optimize your Edge AI applications?

Contact RidgeRun to discuss your Holoscan SDK application and explore opportunities to optimize your pipeline.