cuda-interop

# CUDA Interop ## Overview ovstream consumes raw video frames as **BGRA8 CUDA buffers** directly from device memory. No CPU copy, no upload — the SDK reads the GPU pointer you hand it and encodes via NVENC. The frame descriptor (`VideoFrame` / `ovstream_video_frame_t`) holds: - `buffer` — a CUDA device pointer. - `width`, `height` — pixel dimensions. - `pitch_bytes` — row stride in bytes (≥ `width * 4`). ## Allocate with `cudaMallocPitch` NVENC tolerates the buffer being row-padded for DMA alignment. Let CUDA pick the pitch: > **Source:** `examples/c/basic_stream/main.cu` snippet `cuda-buffer-alloc` `cudaMallocPitch(&buf, &pitch, widthBytes, height)` returns a device pointer plus the actual pitch CUDA picked (often `width * 4` rounded up to a 256 / 512-byte multiple). Pass that exact `pitch` into the `VideoFrame.pitch_bytes` field. ## Python: ctypes-bound cudart The Python wheel bundles `cudart64_*.dll` / `libcudart.so` alongside `ovstream.dll`, so a Python app can use ctypes to call `cudaMalloc` / `cudaMemset` / `cudaFree` directly without any extra CUDA Python package. The example discovers the bundled CUDA runtime via `ovstream._bindings._find_library()` and loads it with ctypes. > **Source:** `examples/python/basic_stream/main.py` snippet `bundled-cudart` If you'd rather use a higher-level CUDA Python package, [`CuPy`](https://cupy.dev), [`PyCUDA`](https://documen.tician.de/pycuda/), or [`NVIDIA Warp`](https://developer.nvidia.com/warp-python) all expose `__cuda_array_interface__` or a `.data_ptr()` accessor that yields a raw device pointer — pass that as `VideoFrame.buffer`. See the `warp_stream` example for a Warp-based producer, and `ovrtx_stream` for handing a renderer's mapped CUDA pointer straight to `VideoFrame`. ## BGRA, not RGBA ovstream's CUDA video format is BGRA8 (blue first), not RGBA8. If your producer outputs RGBA, swizzle on the GPU before submitting. Sketch: ```cpp // CUDA kernel sketch uint8_t* px = buffer + y * pitch + x * 4; px[0] = rgba[2]; // B px[1] = rgba[1]; // G px[2] = rgba[0]; // R px[3] = rgba[3]; // A ``` ovrtx's `LdrColor` render var is RGBA8, but ovrtx 0.5 can publish a BGRA8 render var from its SPG `Swizzle` node (`spg:rtx.spg.stdlib/Swizzle`, `swizzle = "zyxw"`), so the swizzle happens renderer-side and no application kernel is needed. The `ovrtx_stream` example authors that render var in its `stream_scene.usda`. Note that with ovrtx 0.5 an extra render var on a RenderProduct disables DLSS Frame Generation, so an application that needs both keeps the ovrtx output RGBA8 and swizzles the frame it composes for streaming instead. ## Lifetime / re-use pattern `stream_video` stages the frame data into server-owned memory before returning, so the caller may reuse or free the buffer immediately after the call returns — except with `webrtc.resizable`, which disables frame copying and requires a persistent buffer preallocated at the maximum resolution (see the `streaming-frames` skill). The recommended pattern is **one long-lived CUDA buffer that you re-fill in place** each frame — it avoids per-frame allocator churn and is the only supported pattern in resizable mode — while outside that mode per-frame `cudaMalloc` / `cudaFree` is also safe (the encoder never reads through the caller's pointer after return). ## Synchronization `stream_video` does **not** synchronize internally by default — the caller is responsible for making sure the buffer's pixels are visible to the encoder when the call lands. Two patterns: - **Caller pre-syncs.** Leave `frame.sync` zero-initialized (the default). The contract is "the buffer is safe to read on entry to `stream_video`". `cudaDeviceSynchronize` between your render kernel and `stream_video` is sufficient (and is what `basic_stream` does for simplicity). - **Caller hands a sync hint.** Set `frame.sync.wait_event` (a `cudaEvent_t` recorded on your stream after the producer kernel) or `frame.sync.stream` (the CUDA stream the kernel ran on). ovstream chains on the event/stream without a global device sync. SHM and CUDASHM use `cudaStreamWaitEvent` and avoid any host block; RTSP / WebRTC host-block on the event (or stream) before handing the buffer to the encoder. `wait_event` takes precedence when both are set. The sync hint is the right pattern for production render loops that already manage their own CUDA streams — it avoids the global `cudaDeviceSynchronize` cost between producer and ovstream. ## Multi-server zero-copy When you feed multiple servers from one buffer (e.g. WebRTC + RTSP simultaneously, see `basic_stream` example), each server reads the same device pointer. No extra copies — NVENC encodes once per server, but the source data is shared. ## Multi-GPU device selection On a multi-GPU host the GPU your frames live on is often **not** the GPU the backend uses by default: WebRTC/native default to StreamSDK's display adapter, RTSP to the encoder element's registered device, and SHM/CUDASHM to the calling thread's current device. A mismatch yields a connected client with no decodable video (WebRTC/native) or an encoder failure at the first client (RTSP). SHM and CUDASHM copy the frame across devices through the CUDA runtime, which works without peer access but is slower than a same-device copy and reports nothing, so a wrong `cuda_device` there costs bandwidth rather than failing. Set `ServerConfig.cuda_device` (C: `config.cuda_device`) to the ordinal your frame buffers are allocated on — every CUDA-input backend honors it. **For WebRTC/native, also pass `cuda_context`** — the producer's `CUcontext` as an int. This is **required** whenever the producer allocates frames in its own CUDA context, which Warp and most real renderers do. StreamSDK encodes from the device pointer through that context; with `cuda_context=0` it falls back to the device's *primary* context, which cannot read a buffer allocated in a different context and fails the encode (`CUDA error invalid argument: Copying device buffer to VideoFrameResourceCuda`). `cuda_context=0` is only safe when the frames live in the primary context already, e.g. a plain `cudaMalloc` buffer. An ovrtx 0.5 render output mapped to CUDA is readable that way only while ovrtx sees a single GPU (see below). ```python # Warp producer (own CUDA context): pass both. cfg = ovstream.ServerConfig( width=W, height=H, cuda_device=0, cuda_context=int(wp.get_device("cuda:0").context), ) # Plain cudaMalloc producer (primary context): device alone is enough. cfg = ovstream.ServerConfig(width=W, height=H, cuda_device=0) ``` Left at the default `-1`, `cuda_device` preserves each backend's prior behavior, so single-GPU apps need not set either field. **ovrtx on multi-GPU hosts.** With more than one GPU visible to ovrtx 0.5, a RenderProduct without an authored `deviceIds` is assigned to a GPU of ovrtx's choosing, and the mapped render var may then not be readable from the primary context of the device you pass as `cuda_device`: every backend rejects the frame (`cudaMemcpy2DAsync failed: invalid argument` from SHM, CUDASHM and RTSP; a connected client with no decodable video from WebRTC/native). ovrtx exposes no CUDA context for the mapping, so there is nothing to pass as `cuda_context`. Author `uint[] deviceIds = [<ordinal>]` on the RenderProduct and pass the same ordinal as `cuda_device` (the `ovrtx_stream` example does this with 0). If frames are still rejected, restrict the process to one GPU with `CUDA_VISIBLE_DEVICES=<ordinal>` and pass `cuda_device=0`, since the only exposed GPU is then runtime ordinal 0 whatever its physical index. This guidance will change if ovrtx starts exposing the mapping's context or allocates the output differently. ## Key Types / Functions | Python | C | |--------|---| | `ovstream.VideoFrame(buffer=ptr, width=w, height=h, pitch_bytes=p)` | `ovstream_video_frame_t { .buffer = ptr, .width = w, .height = h, .pitch_bytes = p }` | | Any ctypes / CuPy / Warp device pointer (int) | `void*` device pointer | ## Common Pitfalls - **BGRA, not RGBA.** This is the #1 source of "my stream looks blue-tinted" reports. - **Don't pass a host pointer.** `buffer` must be a CUDA *device* pointer. ovstream does not check (passing a host pointer crashes inside NVENC). - **Freeing the buffer mid-stream is safe (since v0.4).** `stream_video` stages each frame into server-owned memory before returning, so `cudaFree` immediately after a `stream_video` call doesn't race the encoder. Earlier ovstream versions held the source by reference and required the buffer to outlive the encoder read; if you're targeting both, keep the buffer alive until `stop()`. - **`pitch * height`, not `width * 4 * height`.** If your frame producer asserts on `pitch == width * 4`, you may need to allocate with plain `cudaMalloc` instead. That works too, but you lose the alignment benefit. - **Multi-GPU: set `cuda_device`.** The backend's default encode/copy GPU is often not where your frames live (the display GPU for WebRTC/native). Set `ServerConfig.cuda_device` to your buffers' device — see *Multi-GPU device selection* above. Single-GPU apps can ignore it.