CUDA OpenGL interop for 3D camera image display

A stereo or structured-light 3D camera can hand you frames that already live in GPU memory. The default habit — convert the SDK image into a CPU buffer, then push it to the screen — drags every frame across the PCIe bus twice and leaves your render thread waiting on both transfers. There is a shorter road: register the window's OpenGL texture with CUDA once, copy each grab device-to-device, and let the GPU display its own data. You get the per-frame call sequence below, the arithmetic that says whether the change pays off at your frame size and frame rate, the cases where the CPU path is still right, and the camera classes the fast path fits.

The decision, stated first

If the SDK's image object can hand you a device pointer, register the display texture once and copy device-to-device on every grab. A GPU-resident frame routed through the CPU pays two PCIe trips per frame; the interop copy is one transfer inside GPU memory that finishes in tens of microseconds. If the SDK only delivers host memory, skip interop: a frame born in host RAM has no download to lose, and a single upload into the texture is already the one-trip minimum.

That is the same rule that governs the rest of a GPU pipeline, including how you feed GPU tensors to a vision model without a CPU round trip: every stage that can stay on the device should. Display is one more consumer that never needs to see the CPU.

How the interop path works

Registration happens once, after both the CUDA and the OpenGL context exist. Create a GL_TEXTURE_2D texture at the display size, hand it to cudaGraphicsGLRegisterImage, and pass the write-discard flag so the driver knows each frame overwrites the whole texture and keeps no read-back bookkeeping. From then on, every grab runs the same four steps: map the resource, fetch the mapped array, copy the SDK's device buffer into it, unmap. Draw only after the unmap — OpenGL must not touch the texture while CUDA holds it.

// Once, after the GL and CUDA contexts both exist:
GLuint glTex = 0;                       // GL_TEXTURE_2D, GL_RGBA8, sized width × height
cudaGraphicsResource_t glRes = nullptr; // handle to the registered texture
cudaGraphicsGLRegisterImage(&glRes, glTex, GL_TEXTURE_2D,
                            cudaGraphicsRegisterFlagsWriteDiscard);

// Per grab: devPtr and pitch come from the SDK's GPU image object.
cudaGraphicsMapResources(1, &glRes, 0);
cudaArray_t texArray = nullptr;
cudaGraphicsSubResourceGetMappedArray(&texArray, glRes, 0, 0);
cudaMemcpy2DToArray(texArray, 0, 0, devPtr, pitch,
                    width * bytesPerPixel, height, cudaMemcpyDeviceToDevice);
cudaGraphicsUnmapResources(1, &glRes, 0);
// GL now samples glTex as an ordinary texture; draw the quad.

In the sketch, width, height, bytesPerPixel, devPtr and pitch are illustrative inputs: devPtr is the device pointer the SDK's GPU image object exposes and pitch is its row stride, typed as that SDK documents them. The remaining calls are documented CUDA runtime functions. Prove the result two ways: time a hundred grabs with CUDA events and confirm the copy sits in the tens of microseconds, then read one frame back with glGetTexImage after the unmap and diff it against the SDK viewer's snapshot for the same timestamp. Identical bytes mean the texture holds what the camera captured.

What each route costs per frame

The rates, with their basis: PCIe 3.0 with four lanes carries 3.94 GB/s raw per direction, and pinned host-device copies sustain about 3 GB/s, so 3.0 GB/s is the working number. Device memory on a mid-range desktop GPU moves 200-400 GB/s; the estimate stays conservative at 200 GB/s, and device-to-device copies pay no bus protocol overhead.

Worked example: a 2K stereo pair at 30 fps

The inputs are illustrative — substitute your own resolution and format. Take a rectified mono 8-bit image from a 2K-class stereo pair, 2208 × 1242 pixels at 30 fps:

1. Frame bytes: 2208 × 1242 × 1 B = 2,742,336 B ≈ 2.74 MB.
2. One PCIe trip: 2,742,336 B / 3.0 × 10⁹ B/s = 0.914 ms ≈ 0.9 ms.
3. The CPU detour, download then upload: 0.9 + 0.9 = 1.8 ms per frame.
4. The interop copy: 2,742,336 B / 2.0 × 10¹¹ B/s = 13.7 µs; with map and unmap overhead, call it 20-50 µs.
5. Per second at 30 fps, the detour holds two transfers on the critical path for 1.8 ms of every 33 ms frame; the interop copy adds under 0.05 ms.
6. Scaling: 16-bit depth doubles the frame to 5.48 MB, so one trip costs 1.8 ms and the detour 3.7 ms; an 8-bit RGBA image at the same resolution is 10.97 MB, or 3.7 ms per trip.

The three routes side by side

Bandwidth is not the binding limit here: two trips × 2.74 MB × 30 fps = 164 MB/s against a 3.9 GB/s link. The cost is latency and synchronization — the picture cannot reach the screen until both DMA transfers and their driver calls finish. On an older PCIe 2.0 ×4 slot the sustained rate falls to roughly 1.5 GB/s and the detour doubles to about 3.7 ms per frame, which starts to bite at modest frame rates.

Keep inspection data separate from display pixels. False colour, contrast stretching, even inverting a depth image for viewing belong in the shader after the interop copy, not in the buffer you measure or export. If the on-screen picture and the exported data ever disagree, why exported point clouds differ from the viewer render walks through the usual causes.

Where the interop path is the wrong tool

• SDK delivers host memory only. Skip interop; the one-trip upload is already optimal.
• Processing still lives on the CPU. The download is then a necessary cost: upload the processed result and you still save the second trip. When a stage can move to the GPU, move it — diagnosing slow GPU image processing shows where GPU time actually goes.
• Frame rate is low. Below roughly 10 fps, a set-up screen or a debug viewer, 1.8 ms is invisible; take the CPU path and keep the code simple.
• No desktop GL context. Headless Linux units need a real or virtual X server before registration can succeed; the FAQ below has the details.
• Multi-GPU node. Interop does not cross devices: the CUDA context and the GL context must sit on the same physical GPU. Run the display on the capture GPU, or copy through host memory and pay for it.
• The destination is a panel, not your window. An operator HMI is a different pipeline; camera integration for HMI image display over a PLC link covers that route.

Kernels and the display copy share the same device memory bandwidth. A 14 µs ask is small next to full-frame processing, but if you stream every camera line to the GPU, budget it deliberately — how FPGA and CUDA FFT pipelines sustain 250 kHz camera lines shows what such a budget buys on live camera data.

Which cameras the pipeline fits

SDKs for stereo and structured-light 3D sensors are usually CUDA-based, so the GPU-resident image object and this interop loop go together naturally: if the image object offers a GPU memory type, the fast path applies. The MV-DPS200P-02 is a projected structured-light 3D camera of that class.

The MV-DLS300P-04 is a laser galvanometer stereo camera of the same family.

Depth and ToF cameras behave the same way: once the SDK hands the depth frame over as a GPU image, the loop is identical. The 3D sensors category gathers the range, from laser profilers to volume cameras:

GigE and USB3 area scan cameras deliver frames into host memory through their SDKs, so the one-trip upload is the best display path out of the box; GPU-direct capture is a separate project. For the 2D texture leg, the LS-CMR-0053 is a 12 MP GigE area scan camera in that group.

Filter the whole camera range by interface and resolution in the camera finder, or start the system configurator from the camera step. If you are unsure which side of the host/GPU line a specific SDK falls on, send the SDK name and your frame rate through the contact page and we will walk the delivery path with you.

Checks before you write the display loop

• Buffer location. Look for a memory-type or GPU accessor on the SDK's image object; if every accessor returns host memory, plan the upload path instead.
• One GPU. CUDA device and GL context on the same physical GPU; registration fails across devices.
• A desktop context. GLX on Linux, WGL on Windows; an EGL-only headless context is refused.
• Texture format. One to four components at 8 or 16 bits per component; compressed and depth formats are refused at registration.
• Ordering. Map, copy, unmap, draw — never GL access while mapped.
• Context lifetime. Register after both contexts exist and keep the camera's CUDA context alive as long as the resource lives; CUDA context invalid after camera close: manage GPU buffers shows what happens when that ordering slips.
• Measurement. Wrap a hundred grabs in CUDA events, quote the steady-state number, and check end-to-end against camera buffer timestamp latency so queue time is not hiding inside the copy.

The wider 3D set — placement, calibration, capture formats — lives on the 3D vision topic page.