diff --git a/examples/camera_livekit_bridge/GPU_PROCESSING.md b/examples/camera_livekit_bridge/GPU_PROCESSING.md new file mode 100644 index 0000000000..f25985c4ae --- /dev/null +++ b/examples/camera_livekit_bridge/GPU_PROCESSING.md @@ -0,0 +1,148 @@ + + +# GPU processing in the camera bridge + +Every picture the bridge publishes can pass through a series of GPU steps on its way from the receiver to the +encoder. This page describes how that works and how to add a step, for example lens rectification. + +None of this runs by default. With no `--gpu-step` the picture is only copied (or shrunk, if the stream has `:DIV`). + +## Where it sits + +```text +camera_receiver (NVDEC) camera_livekit_bridge, one thread per stream + decoded NV12 in a GPU ring --CUDA IPC--> slot (GPU) --> [step 1] --> [step 2] --> ... --> NVENC input surface + (slot notices on a socket) all on one CUDA stream, no host copy + --> H.264 --> LiveKit +``` + +- The receiver decodes into a ring of GPU buffers and tells the bridge which slot holds the newest picture. The bridge + maps the buffers once with CUDA IPC; pixels never cross the socket and never reach host memory. +- The bridge owns the code between the slot and the encoder input surface. That code is a `GpuPipeline` + (`gpu_ops/inc/gpu_ops/gpu_step.hpp`): an ordered list of `GpuStep`s. With no steps it is one plain copy. +- Each stream has its own pipeline, so streams can have different steps and different calibrations. +- The encoder is created for the size the last step produces, so a step may change the picture size. + +## Using steps + +```bash +camera_livekit_bridge ... --stream head=sensing-head:2 \ + --gpu-step sine-warp=14,90 --gpu-step another-step=ARGS +``` + +- `--gpu-step NAME[=ARGS]` adds a step to every stream. It is repeatable. Steps run in the order given. +- Steps run at the source size. The stream's own `:DIV` downscale is added last, so it shrinks the result of the steps. + Run corrections such as rectification at full resolution and let the downscale come after. +- A step that cannot run on a stream (an odd size, a calibration for another size, a size NV12 cannot hold) stops the + bridge at start with the stream, the step and the reason. + +Known steps: + +| Step | Arguments | Purpose | +| --- | --- | --- | +| `sine-warp` | `AMPLITUDE_PX,WAVELENGTH_PX,PHASE_RAD`, all optional | An example horizontal ripple. It shows the shape of a step; it is not calibrated rectification. | + +## The step contract + +A step derives from `GpuStep` and implements three things: + +| Method | Called | Does | +| --- | --- | --- | +| `prepare(input_size)` | once per picture size, before the first frame | Returns an empty string if the step can run on that size, otherwise the reason. Allocate and upload here: lookup tables, calibration, scratch. Remember the size; `enqueue()` receives planes only. | +| `output_size(input_size)` | at prepare time | The size this step produces. Default: unchanged. Return a different size for a crop or scale. | +| `enqueue(in, out, stream)` | every frame | Launch kernels on `stream` and return the CUDA status. | + +Rules for `enqueue()`: + +- Only queue work. Do not allocate, synchronise the device or the stream, or throw. The bridge waits for the stream + once per frame, after the whole pipeline, and then checks the receiver's slot (see below). +- `in` and `out` are different buffers. Work is not in place. +- Pixels only. Frame number, timestamps and metadata travel with the frame; a step never touches them. +- NV12 layout: a full-size luma plane and one interleaved chroma plane at half size in both directions. Planes are + pitched; use the pitch, not the width, to step between rows. Sizes must stay even. +- The output must be fully written, including positions the step cannot compute: use the video-range fill (luma 16, + chroma 128). + +The pipeline owns the buffers between steps (allocated in `GpuPipeline::prepare()`) and gives the last step the encoder +input surface directly, so a step never allocates its own output. + +## Adding a step + +1. Add a class deriving from `GpuStep`. Put the kernels in a `.cu` file next to it in `gpu_ops/`, add the file to + `gpu_ops/CMakeLists.txt`, and keep the public class in `gpu_ops/inc/gpu_ops/`. `sine_warp.hpp` and `sine_warp.cu` + are a complete small example to copy. +2. Register it with one branch in `make_gpu_step()` in `gpu_ops/gpu_pipeline.cpp`, which maps `NAME[=ARGS]` to the + step and parses its arguments (throw `std::invalid_argument` for bad ones). +3. Add a case to `tests/gpu_pipeline/test_gpu_pipeline.cpp`. A good first test is a case where the answer is known: an + identity setting that must equal a plain copy, and one setting that visibly changes the picture. +4. Run it: `cmake --build build && ctest --test-dir build`, then a live run with `--gpu-step` and a saved frame. + +Skeleton: + +```cpp +class MyStep final : public GpuStep +{ +public: + explicit MyStep(Params params) : m_params(params) {} + const char* name() const override { return "my-step"; } + + std::string prepare(const ImageSize& input) override + { + m_size = input; + if (input.width % 2 != 0 || input.height % 2 != 0) + return "needs an even picture size"; + // upload tables here + return {}; + } + + cudaError_t enqueue(const Nv12Planes& in, const Nv12PlanesOut& out, cudaStream_t stream) override + { + my_kernel_launch(in, out, m_size, m_params, stream); // luma and chroma + return cudaGetLastError(); + } + +private: + Params m_params; + ImageSize m_size; +}; +``` + +## Writing lens rectification as a step + +Rectification is a remap: for each pixel of the corrected picture, find where it comes from in the camera image and +sample there. The sine warp is the same operation with a made-up mapping. + +- **The mapping.** For output pixel `(u, v)` build the ray of the ideal pinhole camera (the new camera matrix), distort + it with the lens model (fisheye or radial-tangential coefficients from the camera's intrinsics), and project it with + the real intrinsics. The result is the source position `(sx, sy)`. The sine warp uses + `sx = x + A*sin(2*pi*y/L + phase), sy = y` in its place. +- **Compute it once.** The mapping depends on the calibration and the picture size, never on the pixels. In + `prepare()`, compute `(sx, sy)` for every output pixel on the host, `cudaMalloc` a `float2` table of that size, and + upload it. Free it in the destructor. The per-frame kernel then reads the table entry and bilinear-samples the + source, which costs about the same as the sine warp. +- **Chroma.** The chroma plane is half size. Either keep a second half-size table, or evaluate the luma table at each + chroma sample's luma position (`(2*cx, 2*cy + 0.5)` with the usual siting) and halve the result. +- **Outside the image.** A source position outside the image gets the fill (luma 16, chroma 128). Test for it with a + positive predicate so a NaN also takes the fill path. +- **Size.** Rectification often crops or rescales the field of view. Return that from `output_size()`; the encoder is + created for the last step's output. +- **Calibration input.** Pass a calibration file path as the step's argument, parse it in `make_gpu_step()`, and + refuse in `prepare()` a picture size the calibration was not made for. If cameras differ, the per-stream pipelines + allow a different calibration per stream; the step spec would need a stream name or a per-stream argument. +- **Check it.** A known calibration of a synthetic grid with straight lines is the natural test: after the step the + lines are straight within a pixel. + +## What the bridge does around the steps + +- The pipeline runs on one CUDA stream and the bridge waits for that stream once per frame. That wait is required: + the receiver's ring slot can be overwritten, so the copy out of it must be finished before the slot may be trusted. +- After the wait, the bridge checks for a newer notice for the same slot, or as many newer notices as the ring has + other slots, and drops the picture instead of encoding it (the `slot overruns` count in the stats line). This is a + best-effort check, not exclusive ownership of the slot; see `README.md`. +- Steps add GPU time to every frame. Keep a step's kernels light and measure with the stats line (`encode` and the + capture-to-ready latency) before and after. +- The encoder surface is copied into by the last step, so a step that leaves part of its output unwritten publishes + whatever that surface held before. diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md index 6a52b9c0b1..43a2b16bcf 100644 --- a/examples/camera_livekit_bridge/README.md +++ b/examples/camera_livekit_bridge/README.md @@ -26,6 +26,24 @@ bridge checks for a newer notice for the same slot, or as many newer notices as picture instead of encoding it (counted as `slot overruns`). The session also answers clock-sync requests that arrive as LiveKit data messages. +## GPU steps (per-frame image operations) + +Between the receiver's IPC buffer and the encoder the picture can pass through a series of GPU steps, with no host +copy. None run by default. `--gpu-step NAME[=ARGS]` adds one, repeatable; steps run in the order given, at the source +size, and a stream's own `:DIV` downscale shrinks the result afterwards. + +| Step | Arguments | Purpose | +| --- | --- | --- | +| `sine-warp` | `AMPLITUDE_PX,WAVELENGTH_PX,PHASE_RAD` (all optional) | An EXAMPLE horizontal ripple. It is not calibrated rectification; it exists to show the shape of a step. | + +The system and how to add a step (with a walk-through of writing lens rectification) are described in +[`GPU_PROCESSING.md`](GPU_PROCESSING.md). A step is a class deriving from `GpuStep` (`gpu_ops/inc/gpu_ops/gpu_step.hpp`) with `prepare()` (once per picture size: +allocate and upload) and `enqueue()` (per frame: launch kernels on the given stream). To add one, write the class and +add a line to `make_gpu_step()` in `gpu_ops/gpu_pipeline.cpp`. `SineWarpStep` (`gpu_ops/inc/gpu_ops/sine_warp.hpp`) is +the template, and its comment explains how lens rectification would be written the same way: a remap table computed +once from the camera calibration and uploaded in `prepare()`, and a per-frame kernel that samples the source through it +(luma and chroma), with the fill for positions outside the image. + ## Build ```bash diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index e9fbef9f5f..1ecc277c98 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -34,8 +34,10 @@ void usage() { std::cerr << "usage: camera_livekit_bridge --livekit-url URL --livekit-token-file FILE --stream NAME=TRACK[:DIV] [--stream ...]\n" - " [--socket PATH] [--fps N] [--bitrate-mbps N] [--idr-period N] [--gpu N]\n [--timestamp-basis monotonic_us|capture_monotonic_us|capture_unix_us]\n" - " NAME is a stream in camera_receiver's IPC catalog; TRACK is the LiveKit track name egocap expects;\n DIV shrinks the picture on the GPU by that integer factor before encoding.\n"; + " [--socket PATH] [--fps N] [--bitrate-mbps N] [--idr-period N] [--gpu N]\n [--gpu-step NAME[=ARGS] ...] [--timestamp-basis monotonic_us|capture_monotonic_us|capture_unix_us]\n" + " NAME is a stream in camera_receiver's IPC catalog; TRACK is the LiveKit track name egocap expects;\n DIV shrinks the picture on the GPU by that integer factor before encoding.\n" + " --gpu-step adds a per-frame GPU image operation, repeatable, run in the order given before the downscale;\n" + " none run by default. Known: sine-warp[=AMPLITUDE_PX,WAVELENGTH_PX,PHASE_RAD] (an example, not rectification).\n"; } std::string read_token(const std::string& path) @@ -120,6 +122,8 @@ int main(int argc, char** argv) } else if (arg == "--clock-uncertainty-budget-ms") defaults.clock_budget_ns = static_cast(std::atof(next().c_str()) * 1e6); + else if (arg == "--gpu-step") + defaults.gpu_steps.push_back(next()); else if (arg == "--gpu") gpu = std::atoi(next().c_str()); else if (arg == "--stream") diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 62a542cc9d..a947d04aca 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -3,7 +3,7 @@ #include "stream_worker.hpp" -#include +#include #include #include #include @@ -78,11 +78,16 @@ struct StreamWorker::Impl const IpcStream& stream = client.streams().front(); layout = read_layout(stream); source = { static_cast(stream.width), static_cast(stream.height) }; - if (!box_downscale_supported(source.width, source.height, cfg.scale_div)) - throw std::runtime_error("stream '" + cfg.stream + "' is " + std::to_string(source.width) + "x" + - std::to_string(source.height) + " and cannot be shrunk by " + - std::to_string(cfg.scale_div) + " to an even size"); - geometry = { source.width / cfg.scale_div, source.height / cfg.scale_div }; + // The steps asked for run in order at the source size, then the stream's own downscale shrinks the result. + for (const auto& step_spec : cfg.gpu_steps) + pipeline.add(make_gpu_step(step_spec)); + if (cfg.scale_div > 1) + pipeline.add(make_box_downscale_step(cfg.scale_div)); + const std::string unusable = pipeline.prepare({ source.width, source.height }); + if (!unusable.empty()) + throw std::runtime_error("stream '" + cfg.stream + "' (" + std::to_string(source.width) + "x" + + std::to_string(source.height) + "): " + unusable); + geometry = { pipeline.output_size().width, pipeline.output_size().height }; PublishSpec spec; spec.name = cfg.track; @@ -160,26 +165,12 @@ struct StreamWorker::Impl NvencEncoder::Surface surface; if (!encoder.acquire(surface)) throw std::runtime_error("no free encoder input surface"); - // The picture goes from the IPC slot to the encoder input on the GPU: a plain copy at full size, or a block - // average that shrinks it. The encode below waits for this stream. - const auto copy_status = [&] - { - if (config.scale_div > 1) - { - const Nv12Planes in{ base + layout.y_offset, base + layout.uv_offset, layout.y_pitch, - layout.uv_pitch }; - const Nv12PlanesOut out{ surface.view.y, surface.view.uv, surface.view.y_pitch, - surface.view.uv_pitch }; - return box_downscale_nv12(in, source.width, source.height, out, config.scale_div, copy_stream); - } - const auto y = - cudaMemcpy2DAsync(surface.view.y, surface.view.y_pitch, base + layout.y_offset, layout.y_pitch, - source.width, source.height, cudaMemcpyDeviceToDevice, copy_stream); - if (y != cudaSuccess) - return y; - return cudaMemcpy2DAsync(surface.view.uv, surface.view.uv_pitch, base + layout.uv_offset, layout.uv_pitch, - source.width, source.height / 2, cudaMemcpyDeviceToDevice, copy_stream); - }(); + // The picture goes from the IPC slot to the encoder input on the GPU, through the configured series of + // GPU steps (a plain copy when there are none). This is the place a per-frame image operation such as lens + // rectification runs: see GpuStep in gpu_ops/gpu_step.hpp. The encode below waits for this stream. + const Nv12Planes in{ base + layout.y_offset, base + layout.uv_offset, layout.y_pitch, layout.uv_pitch }; + const Nv12PlanesOut out{ surface.view.y, surface.view.uv, surface.view.y_pitch, surface.view.uv_pitch }; + const cudaError_t copy_status = pipeline.enqueue(in, out, copy_stream); if (copy_status != cudaSuccess) { encoder.release(surface); @@ -333,6 +324,7 @@ struct StreamWorker::Impl CudaContext& context; CudaIpcClient client; PlaneLayout layout; + GpuPipeline pipeline; Geometry source; Geometry geometry; std::unique_ptr sink; diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp index 361fb58df5..08b8194e21 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.hpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -14,6 +14,7 @@ #include #include #include +#include namespace camera_bridge { @@ -30,6 +31,9 @@ struct WorkerConfig uint32_t idr_period = 30; /// Integer factor the picture is shrunk by on the GPU before encoding (1 keeps the source size). uint32_t scale_div = 1; + /// GPU steps run, in order, on every picture before the downscale (see gpu_ops/gpu_step.hpp). Empty by default: + /// the picture is only copied or shrunk. + std::vector gpu_steps; /// What each frame's user_timestamp means; the consumer's streams file must name the same basis. TimestampBasis timestamp_basis = TimestampBasis::CaptureUnixUs; /// The sender's clock, measured over its management channel; required for capture_monotonic_us. Shared by every diff --git a/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt b/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt index b7c31026cc..f2a995fef2 100644 --- a/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt +++ b/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt @@ -2,7 +2,7 @@ # SPDX-License-Identifier: Apache-2.0 # GPU operations on NV12 pictures, run between the IPC slot and the encoder input. -add_library(gpu_ops STATIC box_downscale.cu) +add_library(gpu_ops STATIC box_downscale.cu sine_warp.cu gpu_pipeline.cpp) target_include_directories(gpu_ops INTERFACE diff --git a/examples/camera_livekit_bridge/gpu_ops/gpu_pipeline.cpp b/examples/camera_livekit_bridge/gpu_ops/gpu_pipeline.cpp new file mode 100644 index 0000000000..4553c8746c --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/gpu_pipeline.cpp @@ -0,0 +1,182 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/gpu_ops/gpu_step.hpp" +#include "inc/gpu_ops/sine_warp.hpp" + +#include +#include +#include +#include + +namespace camera_bridge +{ + +namespace +{ + +/// A plain copy: what the pipeline does when it has no steps. +cudaError_t copy_nv12(const Nv12Planes& in, const Nv12PlanesOut& out, const ImageSize& size, cudaStream_t stream) +{ + const cudaError_t y = cudaMemcpy2DAsync( + out.y, out.y_pitch, in.y, in.y_pitch, size.width, size.height, cudaMemcpyDeviceToDevice, stream); + if (y != cudaSuccess) + return y; + return cudaMemcpy2DAsync( + out.uv, out.uv_pitch, in.uv, in.uv_pitch, size.width, size.height / 2, cudaMemcpyDeviceToDevice, stream); +} + +class BoxDownscaleStep final : public GpuStep +{ +public: + explicit BoxDownscaleStep(uint32_t divisor) : m_divisor(divisor) + { + } + + const char* name() const override + { + return "box-downscale"; + } + + ImageSize output_size(const ImageSize& input) const override + { + return { input.width / m_divisor, input.height / m_divisor }; + } + + std::string prepare(const ImageSize& input) override + { + m_size = input; + return box_downscale_supported(input.width, input.height, m_divisor) ? + std::string() : + "a " + std::to_string(input.width) + "x" + std::to_string(input.height) + + " picture cannot be shrunk by " + std::to_string(m_divisor) + " to an even size"; + } + + cudaError_t enqueue(const Nv12Planes& in, const Nv12PlanesOut& out, cudaStream_t stream) override + { + return box_downscale_nv12(in, m_size.width, m_size.height, out, m_divisor, stream); + } + +private: + uint32_t m_divisor; + ImageSize m_size; +}; + +std::vector parse_floats(const std::string& text) +{ + std::vector values; + std::stringstream stream(text); + std::string item; + while (std::getline(stream, item, ',')) + { + char* end = nullptr; + const float value = std::strtof(item.c_str(), &end); + if (item.empty() || end == nullptr || *end != '\0') + throw std::invalid_argument("not a number: '" + item + "'"); + values.push_back(value); + } + return values; +} + +} // namespace + +GpuPipeline::~GpuPipeline() +{ + for (auto& scratch : m_scratch) + { + cudaFree(scratch.y); + cudaFree(scratch.uv); + } +} + +void GpuPipeline::add(std::unique_ptr step) +{ + m_steps.push_back(std::move(step)); +} + +std::string GpuPipeline::prepare(const ImageSize& source) +{ + m_source = source; + m_inputs.clear(); + ImageSize size = source; + for (const auto& step : m_steps) + { + m_inputs.push_back(size); + const std::string reason = step->prepare(size); + if (!reason.empty()) + return std::string(step->name()) + ": " + reason; + size = step->output_size(size); + if (size.width == 0 || size.height == 0 || size.width % 2 != 0 || size.height % 2 != 0) + return std::string(step->name()) + ": it would produce a " + std::to_string(size.width) + "x" + + std::to_string(size.height) + " picture, and NV12 needs even sizes"; + } + m_output = size; + // One buffer for the output of every step except the last, which writes into the caller's destination. + for (std::size_t i = 0; i + 1 < m_steps.size(); ++i) + { + const ImageSize next = m_steps[i]->output_size(m_inputs[i]); + Scratch scratch; + if (cudaMallocPitch(reinterpret_cast(&scratch.y), &scratch.y_pitch, next.width, next.height) != + cudaSuccess || + cudaMallocPitch(reinterpret_cast(&scratch.uv), &scratch.uv_pitch, next.width, next.height / 2) != + cudaSuccess) + { + cudaFree(scratch.y); + return "cannot allocate the buffer between " + std::string(m_steps[i]->name()) + " and the next step"; + } + m_scratch.push_back(scratch); + } + return {}; +} + +cudaError_t GpuPipeline::enqueue(const Nv12Planes& source, const Nv12PlanesOut& destination, cudaStream_t stream) +{ + if (m_steps.empty()) + return copy_nv12(source, destination, m_source, stream); + Nv12Planes current = source; + for (std::size_t i = 0; i < m_steps.size(); ++i) + { + const bool last = i + 1 == m_steps.size(); + const Nv12PlanesOut out = + last ? destination : + Nv12PlanesOut{ m_scratch[i].y, m_scratch[i].uv, m_scratch[i].y_pitch, m_scratch[i].uv_pitch }; + const cudaError_t status = m_steps[i]->enqueue(current, out, stream); + if (status != cudaSuccess) + return status; + current = { out.y, out.uv, out.y_pitch, out.uv_pitch }; + } + return cudaSuccess; +} + +std::unique_ptr make_box_downscale_step(uint32_t divisor) +{ + return std::make_unique(divisor); +} + +// The one place that maps a step name to a step. To add a step, add a branch here. +std::unique_ptr make_gpu_step(const std::string& spec) +{ + const auto equals = spec.find('='); + const std::string name = spec.substr(0, equals); + const std::string args = equals == std::string::npos ? std::string() : spec.substr(equals + 1); + if (name == "sine-warp") + { + // sine-warp[=AMPLITUDE_PX[,WAVELENGTH_PX[,PHASE_RAD]]]: an EXAMPLE ripple, not rectification. + SineWarpParams params; + const auto values = parse_floats(args); + if (values.size() > 3) + throw std::invalid_argument("sine-warp takes at most AMPLITUDE_PX,WAVELENGTH_PX,PHASE_RAD"); + if (values.size() > 0) + params.amplitude_px = values[0]; + if (values.size() > 1) + params.wavelength_px = values[1]; + if (values.size() > 2) + params.phase_rad = values[2]; + return std::make_unique(params); + } + // Add steps here, for example: + // if (name == "rectify") return std::make_unique(load_calibration(args)); + throw std::invalid_argument("unknown GPU step '" + name + "' (known: sine-warp)"); +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/gpu_step.hpp b/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/gpu_step.hpp new file mode 100644 index 0000000000..a26b35b099 --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/gpu_step.hpp @@ -0,0 +1,117 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "box_downscale.hpp" + +#include +#include +#include +#include + +namespace camera_bridge +{ + +struct ImageSize +{ + uint32_t width = 0; + uint32_t height = 0; + bool operator==(const ImageSize& other) const + { + return width == other.width && height == other.height; + } +}; + +/** + * @brief One GPU image operation on an NV12 picture: planes in, planes out, queued on a CUDA stream. + * + * The bridge runs a series of steps between the receiver's IPC buffer and the encoder input, all on the GPU, with no + * host copy. Steps run in the order they are added and each reads the previous step's output. With no steps the + * picture is only copied. A step is a good place for anything that changes pixels per frame: rectification, a crop, + * a color change, an overlay. + * + * To add a step: + * 1. Write a class deriving from GpuStep (see SineWarpStep in sine_warp.hpp for a complete, small example). + * 2. Add one line to make_gpu_step() in gpu_pipeline.cpp so `--gpu-step NAME[=ARGS]` can create it. + * 3. Pass `--gpu-step NAME=ARGS` once per step, in the order they should run. Nothing runs unless asked for. + * + * Contract: + * - prepare() runs once per input size, before the first frame. Allocate and upload here (lookup tables, + * calibration, scratch). enqueue() must not allocate, synchronise the device or the stream, or throw. + * - enqueue() only queues work on `stream` and returns the CUDA status of the launches. The output is not written + * when it returns; the caller waits for the stream. `in` and `out` are different buffers (no in-place work). + * - Pixels only: frame identity, timestamps and metadata travel with the frame and are not a step's business. + * - A step may change the size (output_size()); later steps and the encoder see the new size. + */ +class GpuStep +{ +public: + virtual ~GpuStep() = default; + + virtual const char* name() const = 0; + + /// The size of the picture this step produces from one of `input` size. Default: unchanged. + virtual ImageSize output_size(const ImageSize& input) const + { + return input; + } + + /// Called once for the size of the picture this step will receive. Returns an empty string when the step can run + /// on it, otherwise the reason it cannot (for example an odd size, or a calibration made for another size). + virtual std::string prepare(const ImageSize& input) = 0; + + virtual cudaError_t enqueue(const Nv12Planes& in, const Nv12PlanesOut& out, cudaStream_t stream) = 0; +}; + +/// A series of steps. Owns the intermediate buffers between them; the last step writes straight into the caller's +/// destination (the encoder input surface). +class GpuPipeline +{ +public: + GpuPipeline() = default; + ~GpuPipeline(); + GpuPipeline(const GpuPipeline&) = delete; + GpuPipeline& operator=(const GpuPipeline&) = delete; + + void add(std::unique_ptr step); + + /// Prepares every step for a picture of `source` size and allocates the intermediate buffers. Returns an empty + /// string on success, otherwise which step cannot run and why. + std::string prepare(const ImageSize& source); + + /// The size of the picture the last step produces (the source size when there are no steps). + ImageSize output_size() const + { + return m_output; + } + + /// Queues every step on `stream`, source to destination. With no steps, a plain copy. + cudaError_t enqueue(const Nv12Planes& source, const Nv12PlanesOut& destination, cudaStream_t stream); + +private: + struct Scratch + { + uint8_t* y = nullptr; + uint8_t* uv = nullptr; + size_t y_pitch = 0; + size_t uv_pitch = 0; + }; + + std::vector> m_steps; + std::vector m_inputs; // the size each step receives + std::vector m_scratch; // the output of every step except the last + ImageSize m_source; + ImageSize m_output; +}; + +/** + * @brief Creates a step from a command-line spec `NAME` or `NAME=ARG[,ARG...]`; throws std::invalid_argument for an + * unknown name or bad arguments. The single place that maps names to steps: add new steps here. + */ +std::unique_ptr make_gpu_step(const std::string& spec); + +/// A box downscale by an integer divisor, as a step. +std::unique_ptr make_box_downscale_step(uint32_t divisor); + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/sine_warp.hpp b/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/sine_warp.hpp new file mode 100644 index 0000000000..928be18e4a --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/sine_warp.hpp @@ -0,0 +1,96 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "gpu_step.hpp" + +namespace camera_bridge +{ + +/// Parameters of the example warp. The defaults are a visible but modest ripple. +struct SineWarpParams +{ + /// Peak horizontal displacement in luma pixels. Zero makes the operation an identity. + float amplitude_px = 8.0F; + /// Rows per full sine period. + float wavelength_px = 128.0F; + /// Phase in radians at row 0. + float phase_rad = 0.0F; +}; + +/// True when the parameters are finite with a positive wavelength and the size is a valid NV12 size. +bool sine_warp_supported(uint32_t width, uint32_t height, const SineWarpParams& params); + +/** + * @brief EXAMPLE of a per-frame GPU operation, not calibrated rectification: a horizontal sine ripple. + * + * Output pixel (x, y) is the bilinear sample of the source at (x + amplitude * sin(2*pi*y/wavelength + phase), y); + * a position outside the image gives the video-range black fill. The output has the source's size. Chroma is sampled + * at the mapped position of its own luma-space location (even columns, half way between rows). + */ +cudaError_t sine_warp_nv12(const Nv12Planes& source, + uint32_t width, + uint32_t height, + const Nv12PlanesOut& destination, + const SineWarpParams& params, + cudaStream_t stream); + +/** + * @brief The sine warp as a GpuStep: the template to copy for a real step. + * + * HOW LENS RECTIFICATION WOULD BE WRITTEN AS A STEP (this class is the same shape, with a different mapping): + * + * - Every pixel of the output picture asks "where in the distorted camera image does this pixel come from?". For + * rectification that answer comes from the camera calibration: for the output pixel (u, v) build the ray of the + * ideal pinhole camera (the new camera matrix), distort it with the lens model (for example the fisheye or + * Brown-Conrady coefficients stored with the camera's intrinsics), project it with the real intrinsics, and you + * have the source position (sx, sy). The sine warp replaces exactly this mapping with + * `sx = x + A*sin(2*pi*y/L + phase), sy = y`. + * - The mapping depends only on the calibration and the picture size, never on the pixels, so do it ONCE: + * in prepare(), compute (sx, sy) for every output pixel on the host, cudaMalloc a float2 table of that size + * and upload it. Keep the table as a member and free it in the destructor. The per-frame kernel is then only + * "read the table entry, bilinear-sample the source" and costs about the same as the sine warp. + * - Chroma: the NV12 chroma plane is half size. Either keep a second half-size table, or evaluate the table of + * the luma position of each chroma sample ((2*cx, 2*cy + 0.5) here) and halve the result. Samples whose source + * position falls outside the image get the fill: video-range black (luma 16, chroma 128). + * - Calibration is per camera and per resolution. Give the step its calibration as arguments (a file path, via + * make_gpu_step), refuse in prepare() a picture size the calibration was not made for, and remember that a + * stream shrunk by `--stream NAME=TRACK:DIV` is shrunk AFTER the steps, so a step always sees the source size. + * Rectify at full resolution, then let the downscale shrink the result. + * - Rectification usually crops or scales the field of view: that is output_size() (a different size than the + * input is allowed; the encoder is created for the last step's output size). + * - Several cameras: one pipeline exists per stream, so each stream can get its own calibration; pass the stream + * name or a per-stream argument if the steps differ. + */ +class SineWarpStep final : public GpuStep +{ +public: + explicit SineWarpStep(const SineWarpParams& params = {}) : m_params(params) + { + } + + const char* name() const override + { + return "sine-warp"; + } + + std::string prepare(const ImageSize& input) override + { + m_size = input; // remembered: enqueue() gets planes only, the size is known here + return sine_warp_supported(input.width, input.height, m_params) ? + std::string() : + "the picture size must be even and the warp parameters finite with a positive wavelength"; + } + + cudaError_t enqueue(const Nv12Planes& in, const Nv12PlanesOut& out, cudaStream_t stream) override + { + return sine_warp_nv12(in, m_size.width, m_size.height, out, m_params, stream); + } + +private: + SineWarpParams m_params; + ImageSize m_size; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/gpu_ops/sine_warp.cu b/examples/camera_livekit_bridge/gpu_ops/sine_warp.cu new file mode 100644 index 0000000000..9a731aba06 --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/sine_warp.cu @@ -0,0 +1,121 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/gpu_ops/sine_warp.hpp" + +#include + +namespace camera_bridge +{ + +namespace +{ + +constexpr float kTwoPi = 6.28318530717958647692F; +constexpr unsigned kBlockX = 32; +constexpr unsigned kBlockY = 8; + +struct WarpArgs +{ + float amplitude; + float cycles_per_row; + float phase; +}; + +__device__ __forceinline__ float displaced_x(float x, float y, const WarpArgs& a) +{ + return x + a.amplitude * sinf(kTwoPi * y * a.cycles_per_row + a.phase); +} + +// Bilinear sample of one byte channel; `stride` is the byte distance between neighbouring samples of the channel +// (1 for luma, 2 for interleaved chroma). The caller has checked that sx lies inside [0, w-1]; sy is a row. +__device__ __forceinline__ float bilinear(const uint8_t* plane, size_t pitch, int stride, int w, int h, float sx, float sy) +{ + const int x0 = static_cast(floorf(sx)); + const int y0 = static_cast(floorf(sy)); + const int x1 = min(x0 + 1, w - 1); + const int y1 = min(y0 + 1, h - 1); + const float fx = sx - static_cast(x0); + const float fy = sy - static_cast(y0); + const uint8_t* r0 = plane + static_cast(y0) * pitch; + const uint8_t* r1 = plane + static_cast(y1) * pitch; + const float top = (1.0F - fx) * r0[x0 * stride] + fx * r0[x1 * stride]; + const float bottom = (1.0F - fx) * r1[x0 * stride] + fx * r1[x1 * stride]; + return (1.0F - fy) * top + fy * bottom; +} + +__device__ __forceinline__ uint8_t to_byte(float v) +{ + return static_cast(fminf(fmaxf(v + 0.5F, 0.0F), 255.0F)); +} + +__global__ void warp_luma(const uint8_t* in, size_t in_pitch, uint8_t* out, size_t out_pitch, int w, int h, WarpArgs args) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= w || y >= h) + return; + const float sx = displaced_x(static_cast(x), static_cast(y), args); + // A positive predicate: a NaN coordinate takes the fill path instead of reaching the sampler. + const bool inside = sx >= 0.0F && sx <= static_cast(w - 1); + out[static_cast(y) * out_pitch + x] = + inside ? to_byte(bilinear(in, in_pitch, 1, w, h, sx, static_cast(y))) : 16; +} + +// Chroma sample (cx, cy) sits at luma position (2*cx, 2*cy + 0.5); the mapping is applied there and converted back +// to chroma-plane coordinates. +__global__ void warp_chroma(const uint8_t* in, size_t in_pitch, uint8_t* out, size_t out_pitch, int cw, int ch, WarpArgs args) +{ + const int cx = blockIdx.x * blockDim.x + threadIdx.x; + const int cy = blockIdx.y * blockDim.y + threadIdx.y; + if (cx >= cw || cy >= ch) + return; + const float luma_y = 2.0F * static_cast(cy) + 0.5F; + const float scx = 0.5F * displaced_x(2.0F * static_cast(cx), luma_y, args); + const float scy = static_cast(cy); + uint8_t* dst = out + static_cast(cy) * out_pitch + 2 * cx; + if (!(scx >= 0.0F && scx <= static_cast(cw - 1))) + { + dst[0] = 128; + dst[1] = 128; + return; + } + dst[0] = to_byte(bilinear(in, in_pitch, 2, cw, ch, scx, scy)); + dst[1] = to_byte(bilinear(in + 1, in_pitch, 2, cw, ch, scx, scy)); +} + +} // namespace + +bool sine_warp_supported(uint32_t width, uint32_t height, const SineWarpParams& params) +{ + const bool finite = + std::isfinite(params.amplitude_px) && std::isfinite(params.wavelength_px) && std::isfinite(params.phase_rad); + return finite && params.wavelength_px > 0.0F && std::isfinite(1.0F / params.wavelength_px) && width != 0 && + height != 0 && width % 2 == 0 && height % 2 == 0; +} + +cudaError_t sine_warp_nv12(const Nv12Planes& source, + uint32_t width, + uint32_t height, + const Nv12PlanesOut& destination, + const SineWarpParams& params, + cudaStream_t stream) +{ + if (!sine_warp_supported(width, height, params)) + return cudaErrorInvalidValue; + const int w = static_cast(width); + const int h = static_cast(height); + const WarpArgs args{ params.amplitude_px, 1.0F / params.wavelength_px, params.phase_rad }; + const dim3 block(kBlockX, kBlockY); + warp_luma<<>>( + source.y, source.y_pitch, destination.y, destination.y_pitch, w, h, args); + if (const cudaError_t status = cudaGetLastError(); status != cudaSuccess) + return status; + const int cw = w / 2; + const int ch = h / 2; + warp_chroma<<>>( + source.uv, source.uv_pitch, destination.uv, destination.uv_pitch, cw, ch, args); + return cudaGetLastError(); +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/tests/gpu_pipeline/CMakeLists.txt b/examples/camera_livekit_bridge/tests/gpu_pipeline/CMakeLists.txt new file mode 100644 index 0000000000..56d9d4540b --- /dev/null +++ b/examples/camera_livekit_bridge/tests/gpu_pipeline/CMakeLists.txt @@ -0,0 +1,8 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# Needs a CUDA device; the test exits 77 (skipped) without one. +add_executable(test_gpu_pipeline test_gpu_pipeline.cpp) +target_link_libraries(test_gpu_pipeline PRIVATE camera_livekit_bridge::gpu_ops CUDA::cudart) +add_test(NAME gpu_pipeline COMMAND test_gpu_pipeline) +set_tests_properties(gpu_pipeline PROPERTIES SKIP_RETURN_CODE 77) diff --git a/examples/camera_livekit_bridge/tests/gpu_pipeline/test_gpu_pipeline.cpp b/examples/camera_livekit_bridge/tests/gpu_pipeline/test_gpu_pipeline.cpp new file mode 100644 index 0000000000..b1a856bb95 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/gpu_pipeline/test_gpu_pipeline.cpp @@ -0,0 +1,165 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +// The GPU step pipeline: no steps copies, steps run in series in the order given, the sine warp is an identity at +// zero amplitude and a horizontal ripple otherwise, and bad specs are refused. + +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace +{ + +using camera_bridge::GpuPipeline; +using camera_bridge::ImageSize; +using camera_bridge::Nv12Planes; +using camera_bridge::Nv12PlanesOut; + +int g_failures = 0; + +void check(bool condition, const std::string& what) +{ + if (!condition) + { + std::cerr << "FAILED: " << what << "\n"; + ++g_failures; + } +} + +struct Image +{ + ImageSize size; + uint8_t* device = nullptr; // Y then UV, tightly packed (pitch = width) + explicit Image(const ImageSize& s) : size(s) + { + cudaMalloc(reinterpret_cast(&device), static_cast(s.width) * s.height * 3 / 2); + } + ~Image() + { + cudaFree(device); + } + Nv12Planes in() const + { + return { device, device + static_cast(size.width) * size.height, size.width, size.width }; + } + Nv12PlanesOut out() + { + return { device, device + static_cast(size.width) * size.height, size.width, size.width }; + } + std::vector download() const + { + std::vector host(static_cast(size.width) * size.height * 3 / 2); + cudaMemcpy(host.data(), device, host.size(), cudaMemcpyDeviceToHost); + return host; + } + void upload(const std::vector& host) + { + cudaMemcpy(device, host.data(), host.size(), cudaMemcpyHostToDevice); + } +}; + +std::vector gradient(const ImageSize& s) +{ + std::vector host(static_cast(s.width) * s.height * 3 / 2); + for (uint32_t y = 0; y < s.height; ++y) + for (uint32_t x = 0; x < s.width; ++x) + host[static_cast(y) * s.width + x] = static_cast(16 + (x * 3 + y) % 200); + for (size_t i = static_cast(s.width) * s.height; i < host.size(); ++i) + host[i] = static_cast(100 + i % 50); + return host; +} + +std::vector run(const std::vector& steps, const ImageSize& size, ImageSize& result) +{ + GpuPipeline pipeline; + for (const auto& spec : steps) + pipeline.add(camera_bridge::make_gpu_step(spec)); + const std::string reason = pipeline.prepare(size); + check(reason.empty(), "prepare: " + reason); + result = pipeline.output_size(); + Image source(size); + source.upload(gradient(size)); + Image destination(result); + check(pipeline.enqueue(source.in(), destination.out(), nullptr) == cudaSuccess, "enqueue"); + cudaDeviceSynchronize(); + return destination.download(); +} + +} // namespace + +int main() +{ + int devices = 0; + if (cudaGetDeviceCount(&devices) != cudaSuccess || devices == 0) + { + std::cout << "SKIP: no CUDA device\n"; + return 77; + } + const ImageSize size{ 128, 96 }; + ImageSize out; + + // No steps: a plain copy. + const auto copied = run({}, size, out); + check(out == size && copied == gradient(size), "no steps copies the picture"); + + // Zero amplitude is an identity (the sampler lands exactly on pixel centres). + const auto identity = run({ "sine-warp=0" }, size, out); + check(identity == gradient(size), "sine-warp with zero amplitude is an identity"); + + // A real ripple moves pixels sideways but keeps rows' overall content: the picture changes, same size. + const auto rippled = run({ "sine-warp=8,32" }, size, out); + check(out == size && rippled != gradient(size), "sine-warp changes the picture and keeps its size"); + + // Steps run in series: warp(0) then a downscale equals the downscale alone; two warps differ from one. + ImageSize shrunk; + GpuPipeline plain; + plain.add(camera_bridge::make_box_downscale_step(2)); + check(plain.prepare(size).empty() && plain.output_size() == ImageSize{ 64, 48 }, "downscale output size"); + const auto in_series = run({ "sine-warp=0" }, size, shrunk); // series of one, then compare chain of two below + GpuPipeline chain; + chain.add(camera_bridge::make_gpu_step("sine-warp=0")); + chain.add(camera_bridge::make_box_downscale_step(2)); + check(chain.prepare(size).empty() && chain.output_size() == ImageSize{ 64, 48 }, "chain output size"); + Image source(size); + source.upload(gradient(size)); + Image chained_out({ 64, 48 }); + Image direct_out({ 64, 48 }); + check(chain.enqueue(source.in(), chained_out.out(), nullptr) == cudaSuccess, "chain enqueue"); + check(plain.enqueue(source.in(), direct_out.out(), nullptr) == cudaSuccess, "plain enqueue"); + cudaDeviceSynchronize(); + check(chained_out.download() == direct_out.download(), "identity warp then downscale equals the downscale alone"); + (void)in_series; + + // Refusals. + bool threw = false; + try + { + camera_bridge::make_gpu_step("no-such-step"); + } + catch (const std::invalid_argument&) + { + threw = true; + } + check(threw, "an unknown step name is refused"); + GpuPipeline bad; + bad.add(camera_bridge::make_gpu_step("sine-warp=4,0")); + check(!bad.prepare(size).empty(), "a zero wavelength is refused at prepare"); + GpuPipeline odd; + odd.add(camera_bridge::make_gpu_step("sine-warp")); + check(!odd.prepare({ 127, 96 }).empty(), "an odd picture size is refused at prepare"); + + if (g_failures != 0) + { + std::cerr << g_failures << " check(s) failed\n"; + return 1; + } + std::cout << "all gpu_pipeline checks passed\n"; + return 0; +}