Zero-Copy GPU Decoding on Every Platform

A 4K RGB16 frame is about 50 megabytes. Decoding thousands of them from cinema footage means the copies between CPU and GPU memory are a large share of the time.

The FrameQuery decode pipeline avoids those copies where the platform allows it and removes them on Apple Silicon. Each platform has a different memory architecture, so each gets a different strategy.

The memory problem

The simple approach: allocate a regular CPU buffer, decode a frame into it, allocate a GPU buffer, copy the data over, run the GPU kernel, copy the result back. That is two full-frame copies per frame, each hitting the PCIe bus.

For a 4K RGB16 frame at 50MB, those two copies take about 1.5ms on PCIe 4.0. At 24fps that is 36ms per second of transfers. At 8K it is four times that.

The fix is different on every platform. CUDA has pinned memory. Metal has shared memory. Vulkan has compute shaders that can keep everything on-device. We use all three.

CUDA pinned buffers on Windows

On machines with NVIDIA GPUs we use CUDA page-locked (pinned) memory for host-to-device transfers. Normal malloc memory can be paged out by the OS at any time, which forces the CUDA driver to stage it through an internal pinned buffer before transferring to the GPU. Pinning the memory ourselves skips that intermediate copy.

We wrote a PinnedBuffer<T> RAII wrapper around cudaMallocHost that cleans up with cudaFreeHost on drop:

struct PinnedBuffer<T> {
    ptr: *mut T,
    capacity: usize,
    len: usize,
}

impl<T> PinnedBuffer<T> {
    fn new(capacity: usize) -> Result<Self, String> {
        let size_bytes = capacity * std::mem::size_of::<T>();
        let mut ptr: *mut std::ffi::c_void = ptr::null_mut();
        unsafe {
            let err = cudaMallocHost(&mut ptr, size_bytes);
            if err != 0 {
                return Err(format!("cudaMallocHost failed: {}", err));
            }
        }
        Ok(Self { ptr: ptr as *mut T, capacity, len: capacity })
    }
}
unsafe impl<T> Send for PinnedBuffer<T> {}
unsafe impl<T> Sync for PinnedBuffer<T> {}

We mark PinnedBuffer as Send + Sync because CUDA pinned memory is accessible from any CPU thread. Buffers are shared across the decode pipeline without a lock.

The pipeline uses two buffer slots (double buffering). While one buffer is being filled by the CPU decoder, the other is being transferred to the GPU for indexing. The stages look like this:

1. CPU decodes raw sensor data into pinned buffer A
2. Async H2D copy transfers buffer A to GPU memory
3. GPU processing kernel runs on the transferred data
4. D2H copy pulls the result back to pinned buffer B

In step 4, a cudaMemcpy with cudaMemcpyDeviceToHost implicitly synchronises with the GPU. We do not need an explicit cudaDeviceSynchronize during normal operation. The only place we call cudaDeviceSynchronize is at cleanup, to make sure everything has finished before freeing resources. There are no other sync points.

We also use resolution-adaptive decoding. For on-device frame extraction and search indexing, there is no reason to decode at full resolution. The vendor SDKs expose multiple decode resolutions (full, half, quarter, eighth, sixteenth). Our pipeline picks the smallest resolution that still provides enough pixels for the final output, then lets FFmpeg scale to the exact target size. Decoding at quarter resolution cuts decode time and the bandwidth through every later stage.

Metal zero-copy on Apple Silicon

Apple Silicon has a unified memory architecture. The CPU and GPU share the same physical memory. Metal's storageModeShared lets us allocate a buffer that both processors can access directly, with no copy at all.

We wrap these shared buffers in an RAII struct:

struct MetalSharedBuffer {
    buffer: *mut std::ffi::c_void,  // id<MTLBuffer> retained
    contents: *mut u8,              // .contents pointer (CPU-accessible)
    size: usize,
}

The buffer field holds a retained MTLBuffer for passing to GPU kernels. The contents field is the CPU-accessible pointer from .contents. Because the memory is shared, these point to the same physical pages. The CPU writes raw sensor data directly into the buffer, the GPU reads it for processing, the GPU writes output, and the CPU reads it back. The data never moves.

On a discrete GPU, pinned memory still leaves the PCIe bus as the limit. On unified memory the limit is compute.

The same RAII pattern applies here. When the buffer wrapper is dropped, the underlying MTLBuffer is released via the vendor SDK's deallocation function, which handles the Objective-C runtime calls behind a C interface that Rust can link against.

GPU initialisation and fallback

At startup, FrameQuery probes for available GPU backends using compile-time platform detection:

#[cfg(not(target_os = "macos"))]
let gpu_backend = GpuBackend::Cuda;

#[cfg(target_os = "macos")]
let gpu_backend = GpuBackend::Metal;

The vendor SDKs accept a flag indicating which GPU backend to initialise. If the GPU is not available (no NVIDIA GPU on a Windows machine, or an older Mac without Metal support), the decoder falls back to CPU silently. The same API returns the same frames on either path, at different speeds.

ProRes RAW via Vulkan compute

ProRes RAW is a different problem. Apple designed it, and the only first-party decode path is through AVFoundation on macOS. Windows has no native option.

We built a cross-platform ProRes RAW pipeline using FFmpeg 8.1 and Vulkan compute shaders. Vulkan runs on every major platform, so there is one code path.

The pipeline uses a three-stage cascade with automatic fallback:

1. Full GPU: -hwaccel vulkan with scale_vulkan and h264_vulkan. Everything stays on the GPU. Zero CPU transfer.
2. Hybrid: -hwaccel vulkan with hwdownload and libx264. The GPU handles decode, the CPU handles encode. This kicks in when the full GPU path hits a driver limitation.
3. Software: CPU decode with libx264. The last resort when Vulkan is not available at all.

The Vulkan compute shaders handle Bayer demosaicing directly. FFmpeg has no software demosaic for ProRes RAW, so without Vulkan the only decoder is AVFoundation on macOS.

ProRes RAW also requires HDR tone mapping to produce correct colors. The raw Bayer output has a wider dynamic range than SDR displays can show, and the RGGB Bayer pattern has a 2:1 ratio of green pixels to red and blue. Without tone mapping the output is tinted green. The chain:

zscale=t=linear:npl=100,format=gbrpf32le,zscale=p=bt709,tonemap=hable:desat=0,zscale=t=bt709:m=bt709:r=tv

This converts from linear light to BT.709, applies Hable tone mapping with no desaturation, and outputs in TV range. The intermediate gbrpf32le format gives us 32-bit float precision for the tone mapping maths, which prevents banding in gradients.

Earlier FFmpeg versions had format negotiation bugs between hwdownload and Vulkan surfaces that broke the hybrid pipeline on some drivers. FFmpeg 8.1 fixed them upstream, so we do not carry a patched fork.

The same patterns, different SDKs

The BRAW decoder follows the same architecture. It uses the same PinnedBuffer type for CUDA pinned memory on Windows, and the same Metal shared buffers on macOS.

The BRAW pipeline adds frame-level parallelism through mpsc channels with a capacity of 4. Three pooled buffers rotate through the pipeline: one is being decoded, one is sitting in the channel waiting to be consumed, and one is being written to the output encoder. All three stages (decode, transfer, encode) stay busy and no buffers are allocated per frame.

Trade-offs

cudaMallocHost allocates page-locked memory that the OS cannot swap out. If you pin too much, you reduce the memory available to everything else on the system. We keep the pinned allocation to exactly what the double-buffer pipeline needs: two frame-sized buffers per active decoder.

The Vulkan ProRes RAW path adds FFmpeg as a dependency, which increases the binary size and introduces version sensitivity. We pin to FFmpeg 8.1 and ship it bundled rather than relying on system installs.

With Metal shared memory the CPU and GPU contend for the same memory bus. For this workload the transfer cost removed is larger than the contention added. A purely GPU-bound workload would be faster on dedicated GPU memory.

In the app

R3D, BRAW and ProRes RAW files decode on the fastest path the hardware supports. There are no settings and no GPU toolkit to install. Thousands of clips from a multi-day shoot index in hours rather than a weekend.

Download FrameQuery to try FrameQuery with your own camera originals.