diff --git a/examples/camera_livekit_bridge/CMakeLists.txt b/examples/camera_livekit_bridge/CMakeLists.txt new file mode 100644 index 0000000000..0c8bb954fa --- /dev/null +++ b/examples/camera_livekit_bridge/CMakeLists.txt @@ -0,0 +1,70 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +cmake_minimum_required(VERSION 3.24) +project(camera_livekit_bridge LANGUAGES CXX CUDA) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +if(NOT CMAKE_CUDA_ARCHITECTURES) + set(CMAKE_CUDA_ARCHITECTURES native) +endif() +if(NOT CMAKE_BUILD_TYPE) + set(CMAKE_BUILD_TYPE Release) +endif() + +find_package(CUDAToolkit REQUIRED) +find_package(nlohmann_json REQUIRED) +find_package(Threads REQUIRED) + +include(FetchContent) + +# ============================================================================== +# NVENC headers (ffnvcodec): headers only, loaded dynamically at run time +# ============================================================================== +# Point FETCHCONTENT_SOURCE_DIR_NV_CODEC_HEADERS at an unpacked copy to build offline. +FetchContent_Declare( + nv_codec_headers + URL https://github.com/FFmpeg/nv-codec-headers/archive/refs/tags/n12.2.72.0.tar.gz + URL_HASH SHA256=dbeaec433d93b850714760282f1d0992b1254fc3b5a6cb7d76fc1340a1e47563 + SOURCE_SUBDIR no-cmake-here +) +FetchContent_MakeAvailable(nv_codec_headers) +add_library(nv_codec_headers INTERFACE) +target_include_directories(nv_codec_headers SYSTEM INTERFACE "${nv_codec_headers_SOURCE_DIR}/include") + +# ============================================================================== +# LiveKit C++ SDK: a prebuilt package (headers and shared libraries) +# ============================================================================== +# Point LIVEKIT_SDK_DIR at an unpacked copy to use another, or to build offline. +if(NOT LIVEKIT_SDK_DIR) + set(LIVEKIT_SDK_UBUNTU "24.04" CACHE STRING "Ubuntu release the prebuilt SDK was built for") + if(LIVEKIT_SDK_UBUNTU STREQUAL "22.04") + set(_livekit_hash "9e568ee6b2bb3c7b0d07faba694930767096a14d0e09422ba01f34077ffb03cb") + else() + set(_livekit_hash "5e0feb2e519c6ff131dd2d1c45149115057ae4e21a324fdb249e1789d5614e7b") + endif() + FetchContent_Declare( + livekit_sdk + URL https://github.com/livekit/client-sdk-cpp/releases/download/v1.12.4/livekit-sdk-ubuntu-${LIVEKIT_SDK_UBUNTU}-x64-1.12.4.tar.gz + URL_HASH SHA256=${_livekit_hash} + ) + FetchContent_MakeAvailable(livekit_sdk) + set(LIVEKIT_SDK_DIR "${livekit_sdk_SOURCE_DIR}") +endif() +find_package(LiveKit CONFIG REQUIRED PATHS "${LIVEKIT_SDK_DIR}" NO_DEFAULT_PATH) +message(STATUS "LiveKit SDK=${LIVEKIT_SDK_DIR}") + +add_subdirectory(video_types) +add_subdirectory(ipc_client) +add_subdirectory(gpu_ops) +add_subdirectory(nvenc) +add_subdirectory(livekit_publisher) +add_subdirectory(probe) +add_subdirectory(bridge) + +include(CTest) +if(BUILD_TESTING) + add_subdirectory(tests) +endif() diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md new file mode 100644 index 0000000000..f14ffae051 --- /dev/null +++ b/examples/camera_livekit_bridge/README.md @@ -0,0 +1,95 @@ + + +# camera_livekit_bridge + +A downstream consumer of `camera_streamer`'s CUDA IPC output that publishes the +camera streams to LiveKit as pre-encoded H.264 tracks. It talks to +`camera_receiver` only through its IPC socket and imports nothing from egocap or +the Isaac Capture runtime; egocap joins the LiveKit room on its own. + +```text +camera_receiver ──CUDA IPC──> bridge ──H.264──> LiveKit room ──> egocap page / recorder + (GPU NV12) GPU copy or shrink, NVENC +``` + +Each stream runs on its own thread. A frame is copied (or box-averaged down) on +the GPU from the IPC slot into an NVENC input surface, encoded, and submitted as +a pre-encoded access unit. Pixels never reach host memory. The frame's source +sequence is the track's `frame_id`, the sensor start time (when the receiver +derived one) is the `user_timestamp` in Unix microseconds, and the sensor +counter is the 8-byte `user_data`. The newest picture is encoded and older +queued ones are skipped. The receiver overwrites a slot without notice once its ring wraps, so after the GPU copy 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 (counted as `slot overruns`). The session also answers clock-sync requests that +arrive as LiveKit data messages. + +## Build + +```bash +cmake -S . -B build && cmake --build build -j +``` + +Needs the CUDA toolkit and nlohmann_json. The NVENC headers and the LiveKit C++ +SDK are fetched at configure time; pass `-DLIVEKIT_SDK_DIR=` and +`-DFETCHCONTENT_SOURCE_DIR_NV_CODEC_HEADERS=` to build +offline. The build is standalone and is not part of the repository's top-level +build. + +## Run + +With `camera_receiver` running (for example fed by `camera_streamer/mock_sender`) +and a LiveKit server reachable: + +```bash +build/bridge/camera_livekit_bridge \ + --livekit-url ws://127.0.0.1:7880 --livekit-token-file publisher-token.txt \ + --socket /tmp/camera_receiver_avs1_cuda_ipc.sock \ + --stream head=sensing-head:2 --stream wrist1=sensing-wrist1 \ + --fps 30 --bitrate-mbps 8 +``` + +`--stream NAME=TRACK[:DIV]` publishes the receiver's stream `NAME` as the +LiveKit track `TRACK`, shrunk by the integer `DIV` on the GPU (the size must +divide to even dimensions). Track names, the join token's identity and the +published sizes must match the streams file egocap is given. The token needs +publish permission. Every 5 s it prints each track's frame rate, keyframes, +skipped, rejected and slot-overrun frames and mean encode time. + +`--timestamp-basis` sets what each frame's `user_timestamp` means, and the consumer's streams file must name +the same basis: + +| Basis | Value | +| --- | --- | +| `capture_unix_us` (default) | The sensor start-of-frame as Unix time on the sender's clock. | +| `monotonic_us` | This host's monotonic clock at the instant the picture is ready to encode (before encoding): frame-ready time, later than the sensor capture. | + +`capture_unix_us` is used only when the receiver reports a realtime sender clock; a frame with no usable sensor +start-of-frame (missing, in the future, older than 5 s, or on another clock domain) carries no `user_timestamp`, so +a consumer that needs a capture time says so instead of using a guess. `capture_monotonic_us` is refused: it would +need a measured sender-to-host clock mapping with a stated uncertainty, which does not exist yet. + +`--bitrate-mbps` is the maximum. The encoder starts at 2 Mbps and follows the target bitrate that LiveKit's +congestion control reports for each track (with 10% headroom, never above the maximum or below 0.5 Mbps), so a +fresh connection does not fill the sender's pacing queue. A fixed 8 Mbps start made the first frames 1-2 s old and +growing; with the adaptive start they are about 60-150 ms old from the first frame and settle near 80 ms while +the rate climbs to the maximum within roughly 20 s. The `[stats]` line shows the current rate, the transport's +target and how many times the rate changed. + +Measured with the mock sender on an RTX PRO 5000 against egocap's receiver and recording-feed checks (30 fps, head +and one wrist, local LiveKit server): the head at 2560x1984 and 8 Mbps and at 1280x992 both ran with no missing +frames and a median capture-to-receive time of 65-85 ms. A fixed 20 Mbps head built a standing backlog of about +1.2 s, so keep the maximum near 8 Mbps. + +## Probe + +```bash +build/probe/camera_ipc_probe --socket /tmp/camera_receiver_avs1_cuda_ipc.sock \ + --stream head --frames 30 +``` + +It prints each stream's plane layout, then per frame the slot, generation, +stream sequence, session id, sensor start time and a checksum of the first +pixels. Slots are overwritten without notice when the publisher's ring wraps. diff --git a/examples/camera_livekit_bridge/bridge/CMakeLists.txt b/examples/camera_livekit_bridge/bridge/CMakeLists.txt new file mode 100644 index 0000000000..c647bca7ef --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/CMakeLists.txt @@ -0,0 +1,20 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +add_executable(camera_livekit_bridge + main.cpp + stream_worker.cpp +) + +target_link_libraries(camera_livekit_bridge + PRIVATE + camera_livekit_bridge::ipc_client + camera_livekit_bridge::gpu_ops + camera_livekit_bridge::nvenc + camera_livekit_bridge::livekit_publisher + camera_livekit_bridge::video_types + CUDA::cudart + Threads::Threads +) + +target_compile_options(camera_livekit_bridge PRIVATE -Wall -Wextra) diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp new file mode 100644 index 0000000000..45af62f354 --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -0,0 +1,226 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +// Publishes camera_receiver's CUDA IPC streams to a LiveKit room as pre-encoded H.264 tracks. + +#include "stream_worker.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace +{ + +std::atomic g_stop{ false }; + +void on_signal(int) +{ + g_stop = true; +} + +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"; +} + +std::string read_token(const std::string& path) +{ + std::ifstream file(path); + std::string token; + std::getline(file, token); + while (!token.empty() && (token.back() == ' ' || token.back() == '\r' || token.back() == '\t')) + token.pop_back(); + if (token.empty()) + throw std::runtime_error("cannot read a token from " + path); + return token; +} + +} // namespace + +int main(int argc, char** argv) +{ + using namespace camera_bridge; + + std::string url; + std::string token_file; + std::string socket_path = "/tmp/camera_receiver_avs1_cuda_ipc.sock"; + struct StreamSpec + { + std::string name; + std::string track; + uint32_t scale_div = 1; + }; + std::vector streams; + WorkerConfig defaults; + int gpu = 0; + for (int i = 1; i < argc; ++i) + { + const std::string arg = argv[i]; + const auto next = [&]() -> std::string + { + if (i + 1 >= argc) + { + usage(); + std::exit(2); + } + return argv[++i]; + }; + if (arg == "--livekit-url") + url = next(); + else if (arg == "--livekit-token-file") + token_file = next(); + else if (arg == "--socket") + socket_path = next(); + else if (arg == "--fps") + defaults.fps = static_cast(std::atoi(next().c_str())); + else if (arg == "--bitrate-mbps") + defaults.bitrate_bps = static_cast(std::atof(next().c_str()) * 1e6); + else if (arg == "--idr-period") + defaults.idr_period = static_cast(std::atoi(next().c_str())); + else if (arg == "--timestamp-basis") + { + const auto basis = parse_timestamp_basis(next()); + if (!basis) + { + std::cerr << "--timestamp-basis must be monotonic_us, capture_monotonic_us or capture_unix_us\n"; + return 2; + } + if (*basis == TimestampBasis::CaptureMonotonicUs) + { + std::cerr + << "--timestamp-basis capture_monotonic_us is not available: it needs a measured sender-to-host clock " + "mapping, which is not implemented. Use monotonic_us (frame-ready time) or capture_unix_us.\n"; + return 2; + } + defaults.timestamp_basis = *basis; + } + else if (arg == "--gpu") + gpu = std::atoi(next().c_str()); + else if (arg == "--stream") + { + std::string spec = next(); + StreamSpec parsed; + const auto colon = spec.rfind(':'); + if (colon != std::string::npos) + { + parsed.scale_div = static_cast(std::atoi(spec.substr(colon + 1).c_str())); + spec.erase(colon); + } + const auto eq = spec.find('='); + if (eq == std::string::npos || eq == 0 || eq + 1 == spec.size() || parsed.scale_div == 0) + { + std::cerr << "--stream needs NAME=TRACK or NAME=TRACK:DIV, got '" << spec << "'\n"; + return 2; + } + parsed.name = spec.substr(0, eq); + parsed.track = spec.substr(eq + 1); + streams.push_back(parsed); + } + else + { + usage(); + return 2; + } + } + if (url.empty() || token_file.empty() || streams.empty() || defaults.fps == 0) + { + usage(); + return 2; + } + + std::signal(SIGINT, on_signal); + std::signal(SIGTERM, on_signal); + + try + { + CudaContext context(gpu); + auto session = LiveKitSession::connect(url, read_token(token_file)); + std::vector> workers; + for (const auto& spec : streams) + { + WorkerConfig config = defaults; + config.stream = spec.name; + config.track = spec.track; + config.scale_div = spec.scale_div; + config.socket_path = socket_path; + workers.push_back(std::make_unique(config, context, *session)); + } + for (auto& worker : workers) + worker->start(); + std::cerr << "[info] publishing " << workers.size() << " stream(s) on " << context.device_name() << "\n"; + + auto last_report = std::chrono::steady_clock::now(); + std::vector last_frames(workers.size(), 0); + int exit_code = 0; + while (!g_stop && exit_code == 0) + { + std::this_thread::sleep_for(std::chrono::milliseconds(200)); + if (session->failed()) + { + std::cerr << "[error] " << session->failure() << "\n"; + exit_code = 1; + } + for (const auto& worker : workers) + { + const std::string failure = worker->failure(); + if (!failure.empty()) + { + std::cerr << "[error] " << worker->config().stream << ": " << failure << "\n"; + exit_code = 1; + } + } + const auto now = std::chrono::steady_clock::now(); + if (now - last_report >= std::chrono::seconds(5)) + { + const double seconds = std::chrono::duration(now - last_report).count(); + for (std::size_t i = 0; i < workers.size(); ++i) + { + const auto& stats = workers[i]->stats(); + const uint64_t frames = stats.frames.load(); + const double mean_ms = + frames == 0 ? 0.0 : static_cast(stats.encode_us_total.load()) / frames / 1000.0; + std::cerr << "[stats] " << workers[i]->config().track << ": " << (frames - last_frames[i]) / seconds + << " fps, " << frames << " frames, " << stats.keyframes.load() << " keyframes, " + << stats.skipped.load() << " skipped, " << stats.rejected.load() << " rejected, " + << stats.slot_overruns.load() << " slot overruns, encode " << mean_ms + << " ms mean, bitrate " << stats.bitrate_bps.load() / 1000 << " kbps (transport target " + << stats.target_bps.load() / 1000 << " kbps, " << stats.rate_changes.load() << " changes)"; + const uint64_t ready_samples = stats.capture_to_ready_samples.load(); + if (ready_samples > 0) + { + std::cerr << ", capture-to-ready " + << static_cast(stats.capture_to_ready_us_total.load()) / ready_samples / 1000.0 + << " ms mean " << stats.capture_to_ready_max_us.load() / 1000.0 << " ms max"; + } + std::cerr << "\n"; + last_frames[i] = frames; + } + last_report = now; + } + } + for (auto& worker : workers) + worker->stop(); + // Tracks unpublish when their workers are destroyed, which must happen before the session leaves the room. + workers.clear(); + return exit_code; + } + catch (const std::exception& error) + { + std::cerr << "[error] " << error.what() << "\n"; + return 1; + } +} diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp new file mode 100644 index 0000000000..44ebba2a33 --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -0,0 +1,347 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "stream_worker.hpp" + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +namespace +{ + +struct PlaneLayout +{ + size_t y_offset = 0; + size_t y_pitch = 0; + size_t uv_offset = 0; + size_t uv_pitch = 0; +}; + +PlaneLayout read_layout(const IpcStream& stream) +{ + PlaneLayout layout; + bool have_y = false; + bool have_uv = false; + for (const auto& plane : stream.layout.value("planes", nlohmann::json::array())) + { + const std::string name = plane.value("name", std::string()); + if (name == "y") + { + layout.y_offset = plane.at("offset").get(); + layout.y_pitch = plane.at("pitch").get(); + have_y = true; + } + else if (name == "uv") + { + layout.uv_offset = plane.at("offset").get(); + layout.uv_pitch = plane.at("pitch").get(); + have_uv = true; + } + } + if (stream.pixel_format != "nv12" || !have_y || !have_uv) + throw std::runtime_error("stream '" + stream.name + "' is not NV12 with y and uv planes; only NV12 is supported"); + return layout; +} + +constexpr uint32_t kStartBitrateBps = 2'000'000; + +uint64_t unsigned_field(const nlohmann::json& metadata, const char* key) +{ + const auto it = metadata.find(key); + return (it != metadata.end() && it->is_number_unsigned()) ? it->get() : 0; +} + +} // namespace + +struct StreamWorker::Impl +{ + Impl(const WorkerConfig& cfg, CudaContext& ctx, LiveKitSession& session) + : config(cfg), context(ctx), client(cfg.socket_path, { cfg.stream }) + { + client.connect(); + if (client.streams().size() != 1) + throw std::runtime_error("the IPC publisher has no stream named '" + cfg.stream + "'"); + 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 }; + + PublishSpec spec; + spec.name = cfg.track; + spec.width = geometry.width; + spec.height = geometry.height; + spec.max_bitrate_bps = cfg.bitrate_bps; + spec.max_fps = cfg.fps; + sink = session.publish_video(spec); + } + + ~Impl() + { + running = false; + if (thread.joinable()) + thread.join(); + if (copy_stream != nullptr) + cudaStreamDestroy(copy_stream); + } + + void run(WorkerStats& stats) + { + try + { + loop(stats); + } + catch (const std::exception& error) + { + std::lock_guard lock(failure_mutex); + failure = error.what(); + } + } + + void loop(WorkerStats& stats) + { + context.make_current(); + if (cudaStreamCreateWithFlags(©_stream, cudaStreamNonBlocking) != cudaSuccess) + throw std::runtime_error("cannot create a CUDA stream"); + + EncoderConfig encoder_config; + encoder_config.geometry = geometry; + encoder_config.fps_numerator = config.fps; + // Start low: a fixed high rate sent into a connection whose bandwidth is not yet known queues behind the + // sender's pacer and ages every frame. The transport's target raises it toward the configured maximum. + encoder_config.bitrate_bps = std::min(config.bitrate_bps, kStartBitrateBps); + encoder_config.idr_period = config.idr_period; + NvencEncoder encoder(context, encoder_config); + + const IpcStream& stream = client.streams().front(); + bool need_idr = true; + uint32_t own_frame_id = 0; + uint64_t last_session = 0; + FrameNotice notice; + std::optional pending; // a newer notice seen while checking the last slot; used next + while (running) + { + if (pending) + { + notice = std::move(*pending); + pending.reset(); + } + else if (!client.next_frame(notice, std::chrono::milliseconds(200))) + continue; + // Encode the newest picture: anything queued behind it is older, and a late frame is worse than a skipped + // one. + FrameNotice newer; + while (client.next_frame(newer, std::chrono::milliseconds(0))) + { + notice = std::move(newer); + ++stats.skipped; + } + + const IpcSlot& slot = client.slot(notice); + const auto* base = static_cast(slot.device_ptr); + 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); + }(); + if (copy_status != cudaSuccess) + { + encoder.release(surface); + throw std::runtime_error(std::string("GPU copy from the IPC slot failed: ") + + cudaGetErrorString(copy_status)); + } + + // The publisher overwrites a slot without notice once its ring wraps. The copy must have finished before + // that, and a newer notice for the same slot, or as many newer notices as the ring has other slots, means + // it may not have. Such a picture is dropped, never encoded. + if (cudaStreamSynchronize(copy_stream) != cudaSuccess) + { + encoder.release(surface); + throw std::runtime_error("GPU copy from the IPC slot did not complete"); + } + { + bool overwritten = false; + std::size_t newer_count = 0; + FrameNotice newer; + while (client.next_frame(newer, std::chrono::milliseconds(0))) + { + ++newer_count; + if (newer.slot == notice.slot && newer.slot_generation > notice.slot_generation) + overwritten = true; + pending = std::move(newer); + } + if (stream.slots.size() > 1 && newer_count + 1 >= stream.slots.size()) + overwritten = true; + if (overwritten) + { + encoder.release(surface); + ++stats.slot_overruns; + continue; + } + } + + const uint64_t session_id = unsigned_field(notice.metadata, "session_id"); + if (session_id != last_session) + { + // A restarted source: start the track's picture sequence over at an IDR. + need_idr = true; + last_session = session_id; + } + const auto feedback = sink->take_feedback(); + const bool force_idr = need_idr || feedback.keyframe_requested; + if (feedback.target_bitrate_bps) + { + // The congestion controller's target for this track, with headroom. A fixed rate above what the + // connection can carry only fills the sender's pacing queue and ages every frame behind it. Never + // above the configured rate, never below a floor that keeps the picture usable. + constexpr uint32_t kMinBitrateBps = 500'000; + stats.target_bps = *feedback.target_bitrate_bps; + const uint64_t shaded = static_cast(*feedback.target_bitrate_bps) * 9 / 10; + const auto wanted = + static_cast(std::clamp(shaded, kMinBitrateBps, config.bitrate_bps)); + const uint32_t current = encoder.bitrate(); + // Only a real change reconfigures: each one resets rate control inside the encoder. + if (wanted > current + current / 10 || wanted + wanted / 10 < current) + { + encoder.set_bitrate(wanted); + ++stats.rate_changes; + } + } + stats.bitrate_bps = encoder.bitrate(); + + // The frame-ready instant: the picture is on the encoder's input, before any encoding delay. Both clocks + // are read together so the capture bases can move the sensor time between them. + timespec ready_realtime{}; + clock_gettime(CLOCK_REALTIME, &ready_realtime); + const int64_t ready_realtime_ns = + static_cast(ready_realtime.tv_sec) * 1'000'000'000 + ready_realtime.tv_nsec; + const int64_t ready_monotonic_ns = monotonic_ns(); + + const auto begin = std::chrono::steady_clock::now(); + EncodedPacket packet = encoder.encode(surface, copy_stream, 0, force_idr); + stats.encode_us_total += static_cast( + std::chrono::duration_cast(std::chrono::steady_clock::now() - begin).count()); + need_idr = false; + + const uint64_t sequence = unsigned_field(notice.metadata, "stream_sequence"); + EncodedVideoFrame frame; + frame.data = packet.data.data(); + frame.size = packet.data.size(); + frame.keyframe = packet.keyframe; + frame.timestamp_us = static_cast(ready_monotonic_ns / 1000); + // The sensor start-of-frame is wall-clock time on the sender; only a realtime (Unix) sender clock may be + // used as Unix time. Any other domain, or none, means the publisher cannot vouch for a capture time. + const auto wall_domain = notice.metadata.find("sender_wall_clock"); + const bool wall_is_realtime = wall_domain != notice.metadata.end() && wall_domain->is_string() && + wall_domain->get() == "realtime"; + const uint64_t sensor_sof_ns = wall_is_realtime ? unsigned_field(notice.metadata, "sensor_sof_ns") : 0; + frame.user_timestamp_us = + user_timestamp_us(config.timestamp_basis, sensor_sof_ns, ready_realtime_ns, ready_monotonic_ns); + if (sensor_sof_ns != 0) + { + const int64_t age_ns = ready_realtime_ns - static_cast(sensor_sof_ns); + if (age_ns >= 0 && age_ns <= kMaxCaptureAgeNs) + { + const auto age_us = static_cast(age_ns / 1000); + stats.capture_to_ready_us_total += age_us; + ++stats.capture_to_ready_samples; + uint32_t seen = stats.capture_to_ready_max_us.load(); + while (age_us > seen && !stats.capture_to_ready_max_us.compare_exchange_weak(seen, age_us)) + { + } + } + } + // The source's own sequence keeps frames lost before the encoder visible to the consumer; a source without + // one is numbered here. + frame.frame_id = sequence != 0 ? static_cast(sequence) : ++own_frame_id; + frame.capture_tsc = unsigned_field(notice.metadata, "sof_tsc_ns"); + + ++stats.frames; + if (packet.keyframe) + ++stats.keyframes; + if (!sink->submit(frame)) + { + // The SDK refused a unit, which can break the reference chain: follow with a keyframe. + need_idr = true; + ++stats.rejected; + } + } + (void)stream; + } + + WorkerConfig config; + CudaContext& context; + CudaIpcClient client; + PlaneLayout layout; + Geometry source; + Geometry geometry; + std::unique_ptr sink; + cudaStream_t copy_stream = nullptr; + std::atomic running{ false }; + std::thread thread; + mutable std::mutex failure_mutex; + std::string failure; +}; + +StreamWorker::StreamWorker(const WorkerConfig& config, CudaContext& context, LiveKitSession& session) + : m_config(config), m_impl(std::make_unique(config, context, session)) +{ +} + +StreamWorker::~StreamWorker() = default; + +void StreamWorker::start() +{ + m_impl->running = true; + m_impl->thread = std::thread([this] { m_impl->run(m_stats); }); +} + +void StreamWorker::stop() +{ + m_impl->running = false; + if (m_impl->thread.joinable()) + m_impl->thread.join(); +} + +std::string StreamWorker::failure() const +{ + std::lock_guard lock(m_impl->failure_mutex); + return m_impl->failure; +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp new file mode 100644 index 0000000000..f1879cbb04 --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -0,0 +1,94 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +struct WorkerConfig +{ + /// The stream's name in camera_receiver's IPC catalog. + std::string stream; + /// The LiveKit track name egocap's streams file expects. + std::string track; + std::string socket_path; + uint32_t fps = 30; + uint32_t bitrate_bps = 8'000'000; + 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; + /// What each frame's user_timestamp means; the consumer's streams file must name the same basis. + TimestampBasis timestamp_basis = TimestampBasis::CaptureUnixUs; +}; + +struct WorkerStats +{ + std::atomic frames{ 0 }; + std::atomic skipped{ 0 }; + std::atomic keyframes{ 0 }; + std::atomic rejected{ 0 }; + /// Pictures dropped because the publisher may have overwritten the IPC slot during the copy. + std::atomic slot_overruns{ 0 }; + std::atomic encode_us_total{ 0 }; + /// The bitrate the encoder is configured for now, and how many times congestion feedback changed it. + std::atomic bitrate_bps{ 0 }; + std::atomic rate_changes{ 0 }; + /// Sensor start-of-frame to frame-ready, from frames whose sensor time is on a realtime sender clock. Only as + /// good as the agreement between the sender's wall clock and this host's. + std::atomic capture_to_ready_us_total{ 0 }; + std::atomic capture_to_ready_samples{ 0 }; + std::atomic capture_to_ready_max_us{ 0 }; + /// The latest target the transport reported (0 until it reports one). + std::atomic target_bps{ 0 }; +}; + +/** + * @brief One camera stream from CUDA IPC to a LiveKit track on its own thread. + * + * Each frame is copied on the GPU from the IPC slot into an NVENC input surface, encoded to H.264 and submitted as a + * pre-encoded access unit. The pixels never reach host memory. Streams share nothing but the CUDA context and the + * LiveKit room, so one stalling does not delay another. + */ +class StreamWorker +{ +public: + /// Connects to the IPC socket and publishes the track. Throws std::runtime_error with the reason on failure. + StreamWorker(const WorkerConfig& config, CudaContext& context, LiveKitSession& session); + ~StreamWorker(); + + StreamWorker(const StreamWorker&) = delete; + StreamWorker& operator=(const StreamWorker&) = delete; + + void start(); + void stop(); + + const WorkerConfig& config() const + { + return m_config; + } + const WorkerStats& stats() const + { + return m_stats; + } + /// The reason the worker ended, or empty while it runs or after a requested stop. + std::string failure() const; + +private: + struct Impl; + WorkerConfig m_config; + WorkerStats m_stats; + std::unique_ptr m_impl; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt b/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt new file mode 100644 index 0000000000..b7c31026cc --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt @@ -0,0 +1,17 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# 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) + +target_include_directories(gpu_ops + INTERFACE + $ +) + +target_link_libraries(gpu_ops + PUBLIC + CUDA::cudart +) + +add_library(camera_livekit_bridge::gpu_ops ALIAS gpu_ops) diff --git a/examples/camera_livekit_bridge/gpu_ops/box_downscale.cu b/examples/camera_livekit_bridge/gpu_ops/box_downscale.cu new file mode 100644 index 0000000000..6529979e43 --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/box_downscale.cu @@ -0,0 +1,102 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/gpu_ops/box_downscale.hpp" + +namespace camera_bridge +{ + +namespace +{ + +constexpr uint32_t kMaxDivisor = 8; + +__global__ void downscale_luma(const uint8_t* src, + size_t src_pitch, + uint8_t* dst, + size_t dst_pitch, + uint32_t dst_width, + uint32_t dst_height, + uint32_t divisor) +{ + const uint32_t x = blockIdx.x * blockDim.x + threadIdx.x; + const uint32_t y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst_width || y >= dst_height) + return; + uint32_t sum = 0; + for (uint32_t j = 0; j < divisor; ++j) + { + const uint8_t* row = src + static_cast(y * divisor + j) * src_pitch + x * divisor; + for (uint32_t i = 0; i < divisor; ++i) + sum += row[i]; + } + dst[static_cast(y) * dst_pitch + x] = + static_cast((sum + divisor * divisor / 2) / (divisor * divisor)); +} + +// One thread per destination chroma sample pair (U and V), averaging the same block of source pairs. +__global__ void downscale_chroma(const uint8_t* src, + size_t src_pitch, + uint8_t* dst, + size_t dst_pitch, + uint32_t dst_pairs_x, + uint32_t dst_rows, + uint32_t divisor) +{ + const uint32_t x = blockIdx.x * blockDim.x + threadIdx.x; + const uint32_t y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst_pairs_x || y >= dst_rows) + return; + uint32_t sum_u = 0; + uint32_t sum_v = 0; + for (uint32_t j = 0; j < divisor; ++j) + { + const uint8_t* row = + src + static_cast(y * divisor + j) * src_pitch + static_cast(x) * divisor * 2; + for (uint32_t i = 0; i < divisor; ++i) + { + sum_u += row[2 * i]; + sum_v += row[2 * i + 1]; + } + } + const uint32_t count = divisor * divisor; + uint8_t* out = dst + static_cast(y) * dst_pitch + static_cast(x) * 2; + out[0] = static_cast((sum_u + count / 2) / count); + out[1] = static_cast((sum_v + count / 2) / count); +} + +} // namespace + +bool box_downscale_supported(uint32_t width, uint32_t height, uint32_t divisor) +{ + if (divisor == 0 || divisor > kMaxDivisor) + return false; + if (width % divisor != 0 || height % divisor != 0) + return false; + return (width / divisor) % 2 == 0 && (height / divisor) % 2 == 0 && width / divisor > 0 && height / divisor > 0; +} + +cudaError_t box_downscale_nv12(const Nv12Planes& source, + uint32_t width, + uint32_t height, + const Nv12PlanesOut& destination, + uint32_t divisor, + cudaStream_t stream) +{ + if (!box_downscale_supported(width, height, divisor)) + return cudaErrorInvalidValue; + const uint32_t dst_width = width / divisor; + const uint32_t dst_height = height / divisor; + const dim3 block(32, 8); + const dim3 luma_grid((dst_width + block.x - 1) / block.x, (dst_height + block.y - 1) / block.y); + downscale_luma<<>>( + source.y, source.y_pitch, destination.y, destination.y_pitch, dst_width, dst_height, divisor); + const uint32_t pairs = dst_width / 2; + const uint32_t rows = dst_height / 2; + const dim3 chroma_grid((pairs + block.x - 1) / block.x, (rows + block.y - 1) / block.y); + downscale_chroma<<>>( + source.uv, source.uv_pitch, destination.uv, destination.uv_pitch, pairs, rows, divisor); + return cudaGetLastError(); +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/box_downscale.hpp b/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/box_downscale.hpp new file mode 100644 index 0000000000..a47a2f9516 --- /dev/null +++ b/examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/box_downscale.hpp @@ -0,0 +1,42 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +namespace camera_bridge +{ + +/// A device NV12 image: a luma plane and one interleaved chroma plane at half resolution in both directions. +struct Nv12Planes +{ + const uint8_t* y = nullptr; + const uint8_t* uv = nullptr; + size_t y_pitch = 0; + size_t uv_pitch = 0; +}; + +struct Nv12PlanesOut +{ + uint8_t* y = nullptr; + uint8_t* uv = nullptr; + size_t y_pitch = 0; + size_t uv_pitch = 0; +}; + +/// True when an image of this size shrinks by `divisor` to a valid NV12 size (even width and height). +bool box_downscale_supported(uint32_t width, uint32_t height, uint32_t divisor); + +/// Shrinks an NV12 image by an integer `divisor` (1 copies), averaging each divisor x divisor block of luma and of +/// chroma, on `stream`. The destination is width/divisor x height/divisor. Returns the CUDA status of the launch. +cudaError_t box_downscale_nv12(const Nv12Planes& source, + uint32_t width, + uint32_t height, + const Nv12PlanesOut& destination, + uint32_t divisor, + cudaStream_t stream); + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/ipc_client/CMakeLists.txt b/examples/camera_livekit_bridge/ipc_client/CMakeLists.txt new file mode 100644 index 0000000000..490db23f5c --- /dev/null +++ b/examples/camera_livekit_bridge/ipc_client/CMakeLists.txt @@ -0,0 +1,20 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +add_library(ipc_client STATIC cuda_ipc_client.cpp) + +target_include_directories(ipc_client + INTERFACE + $ +) + +target_link_libraries(ipc_client + PUBLIC + nlohmann_json::nlohmann_json + PRIVATE + CUDA::cudart +) + +target_compile_options(ipc_client PRIVATE -Wall -Wextra) + +add_library(camera_livekit_bridge::ipc_client ALIAS ipc_client) diff --git a/examples/camera_livekit_bridge/ipc_client/cuda_ipc_client.cpp b/examples/camera_livekit_bridge/ipc_client/cuda_ipc_client.cpp new file mode 100644 index 0000000000..e91668df5f --- /dev/null +++ b/examples/camera_livekit_bridge/ipc_client/cuda_ipc_client.cpp @@ -0,0 +1,246 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/ipc_client/cuda_ipc_client.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +namespace +{ + +constexpr std::size_t kMaxMessageBytes = 1024 * 1024; + +[[noreturn]] void fail(const std::string& what) +{ + throw std::runtime_error(what); +} + +void check_cuda(cudaError_t status, const char* what) +{ + if (status != cudaSuccess) + fail(std::string(what) + ": " + cudaGetErrorString(status)); +} + +std::vector base64_decode(const std::string& text) +{ + static const std::string alphabet = "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"; + std::vector out; + unsigned int accumulator = 0; + int bits = 0; + for (const char c : text) + { + if (c == '=') + break; + const auto position = alphabet.find(c); + if (position == std::string::npos) + fail("invalid base64 in an IPC handle"); + accumulator = (accumulator << 6) | static_cast(position); + bits += 6; + if (bits >= 8) + { + bits -= 8; + out.push_back(static_cast((accumulator >> bits) & 0xFF)); + } + } + return out; +} + +} // namespace + +CudaIpcClient::CudaIpcClient(std::string socket_path, std::vector streams) + : m_socket_path(std::move(socket_path)), m_requested(std::move(streams)) +{ +} + +CudaIpcClient::~CudaIpcClient() +{ + close_all(); +} + +void CudaIpcClient::close_all() +{ + for (auto& stream : m_streams) + { + for (auto& slot : stream.slots) + { + if (slot.device_ptr != nullptr) + { + cudaSetDevice(stream.gpu_id); + cudaIpcCloseMemHandle(slot.device_ptr); + slot.device_ptr = nullptr; + } + } + } + m_streams.clear(); + if (m_fd >= 0) + { + ::close(m_fd); + m_fd = -1; + } +} + +void CudaIpcClient::send_message(const nlohmann::json& message) +{ + const std::string line = message.dump() + "\n"; + std::size_t sent = 0; + while (sent < line.size()) + { + const ssize_t count = ::send(m_fd, line.data() + sent, line.size() - sent, MSG_NOSIGNAL); + if (count < 0) + fail(std::string("cannot send to the IPC publisher: ") + std::strerror(errno)); + sent += static_cast(count); + } +} + +nlohmann::json CudaIpcClient::read_message(std::chrono::milliseconds timeout, bool& timed_out) +{ + timed_out = false; + const auto deadline = std::chrono::steady_clock::now() + timeout; + for (;;) + { + const auto newline = m_buffer.find('\n'); + if (newline != std::string::npos) + { + const std::string line = m_buffer.substr(0, newline); + m_buffer.erase(0, newline + 1); + auto message = nlohmann::json::parse(line, nullptr, /*allow_exceptions=*/false); + if (message.is_discarded() || !message.is_object()) + fail("the IPC publisher sent a message that is not a JSON object"); + return message; + } + if (m_buffer.size() > kMaxMessageBytes) + fail("the IPC publisher sent an oversized message"); + const auto remaining = + std::chrono::duration_cast(deadline - std::chrono::steady_clock::now()); + if (remaining.count() < 0) + { + timed_out = true; + return {}; + } + pollfd descriptor{ m_fd, POLLIN, 0 }; + const int ready = ::poll(&descriptor, 1, static_cast(remaining.count())); + if (ready < 0 && errno != EINTR) + fail(std::string("poll on the IPC socket failed: ") + std::strerror(errno)); + if (ready <= 0) + continue; + char chunk[65536]; + const ssize_t count = ::recv(m_fd, chunk, sizeof(chunk), 0); + if (count == 0) + fail("the IPC publisher closed the connection"); + if (count < 0) + { + if (errno == EINTR || errno == EAGAIN) + continue; + fail(std::string("cannot read from the IPC publisher: ") + std::strerror(errno)); + } + m_buffer.append(chunk, static_cast(count)); + } +} + +void CudaIpcClient::connect() +{ + close_all(); + m_fd = ::socket(AF_UNIX, SOCK_STREAM, 0); + if (m_fd < 0) + fail(std::string("cannot create a socket: ") + std::strerror(errno)); + sockaddr_un address{}; + address.sun_family = AF_UNIX; + if (m_socket_path.size() >= sizeof(address.sun_path)) + fail("the IPC socket path is too long"); + std::strncpy(address.sun_path, m_socket_path.c_str(), sizeof(address.sun_path) - 1); + if (::connect(m_fd, reinterpret_cast(&address), sizeof(address)) != 0) + fail("cannot connect to the IPC socket " + m_socket_path + ": " + std::strerror(errno)); + + bool timed_out = false; + const auto hello = read_message(std::chrono::seconds(5), timed_out); + if (timed_out || hello.value("type", "") != "hello") + fail("expected a hello message from the IPC publisher"); + send_message({ { "type", "subscribe" }, { "streams", m_requested } }); + const auto subscribed = read_message(std::chrono::seconds(5), timed_out); + if (timed_out) + fail("the IPC publisher did not answer the subscription"); + if (subscribed.value("type", "") == "error") + fail("the IPC subscription was refused: " + subscribed.value("error", std::string("no reason given"))); + if (subscribed.value("type", "") != "subscribed") + fail("expected a subscribed message from the IPC publisher"); + open_streams(subscribed); +} + +void CudaIpcClient::open_streams(const nlohmann::json& subscribed) +{ + for (const auto& entry : subscribed.at("streams")) + { + IpcStream stream; + stream.name = entry.at("name").get(); + stream.width = entry.at("width").get(); + stream.height = entry.at("height").get(); + stream.pixel_format = entry.value("pixel_format", std::string()); + stream.gpu_id = entry.value("gpu_id", 0); + stream.layout = entry.value("layout", nlohmann::json::object()); + check_cuda(cudaSetDevice(stream.gpu_id), "cudaSetDevice"); + for (const auto& item : entry.at("slots")) + { + const auto handle_bytes = base64_decode(item.at("handle").get()); + cudaIpcMemHandle_t handle{}; + if (handle_bytes.size() != sizeof(handle)) + fail("an IPC handle has the wrong size"); + std::memcpy(&handle, handle_bytes.data(), sizeof(handle)); + IpcSlot slot; + check_cuda( + cudaIpcOpenMemHandle(&slot.device_ptr, handle, cudaIpcMemLazyEnablePeerAccess), "cudaIpcOpenMemHandle"); + slot.nbytes = item.at("nbytes").get(); + slot.shape = item.at("shape").get>(); + slot.strides = item.at("strides").get>(); + const auto index = static_cast(item.at("slot").get()); + if (stream.slots.size() <= index) + stream.slots.resize(index + 1); + stream.slots[index] = slot; + } + m_streams.push_back(std::move(stream)); + } +} + +const IpcSlot& CudaIpcClient::slot(const FrameNotice& notice) const +{ + for (const auto& stream : m_streams) + { + if (stream.name == notice.stream) + return stream.slots.at(static_cast(notice.slot)); + } + throw std::out_of_range("unknown IPC stream " + notice.stream); +} + +bool CudaIpcClient::next_frame(FrameNotice& out, std::chrono::milliseconds timeout) +{ + const auto deadline = std::chrono::steady_clock::now() + timeout; + for (;;) + { + bool timed_out = false; + const auto remaining = + std::max(std::chrono::milliseconds(0), + std::chrono::duration_cast(deadline - std::chrono::steady_clock::now())); + auto message = read_message(remaining, timed_out); + if (timed_out) + return false; + if (message.value("type", "") != "frame") + continue; + out.stream = message.at("stream").get(); + out.slot = message.at("slot").get(); + out.slot_generation = message.value("slot_generation", std::uint64_t{ 0 }); + out.metadata = message.value("metadata", nlohmann::json::object()); + return true; + } +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/ipc_client/inc/ipc_client/cuda_ipc_client.hpp b/examples/camera_livekit_bridge/ipc_client/inc/ipc_client/cuda_ipc_client.hpp new file mode 100644 index 0000000000..6b4de91e5d --- /dev/null +++ b/examples/camera_livekit_bridge/ipc_client/inc/ipc_client/cuda_ipc_client.hpp @@ -0,0 +1,90 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include + +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +/// One ring slot of a stream: device memory opened from the publisher's CUDA IPC handle. +struct IpcSlot +{ + void* device_ptr = nullptr; + std::size_t nbytes = 0; + std::vector shape; + std::vector strides; +}; + +struct IpcStream +{ + std::string name; + int width = 0; + int height = 0; + std::string pixel_format; + int gpu_id = 0; + nlohmann::json layout; + std::vector slots; +}; + +/// A frame notification: the slot now holds a new picture. The slot is overwritten without notice when the +/// publisher's ring wraps, so a consumer must finish with it before `slot_generation` advances by the ring size. +struct FrameNotice +{ + std::string stream; + int slot = 0; + std::uint64_t slot_generation = 0; + nlohmann::json metadata; +}; + +/** + * @brief Client of camera_receiver's CUDA IPC publisher (newline-delimited JSON on a unix socket). + * + * Protocol: the publisher sends `hello`; the client answers `subscribe` with stream names (all when empty); the + * publisher replies `subscribed` with one CUDA IPC handle per slot, then sends `frame` notices. Pixels never cross + * the socket. Must run in a different process from the publisher. + */ +class CudaIpcClient +{ +public: + CudaIpcClient(std::string socket_path, std::vector streams); + ~CudaIpcClient(); + + CudaIpcClient(const CudaIpcClient&) = delete; + CudaIpcClient& operator=(const CudaIpcClient&) = delete; + + /// Connects, subscribes and opens every slot. Throws std::runtime_error with the reason on failure. + void connect(); + + const std::vector& streams() const + { + return m_streams; + } + + /// The slot a frame notice refers to; throws std::out_of_range for an unknown stream or slot. + const IpcSlot& slot(const FrameNotice& notice) const; + + /// Waits up to `timeout` for the next frame notice. Returns false on timeout; throws if the publisher went away. + bool next_frame(FrameNotice& out, std::chrono::milliseconds timeout); + +private: + nlohmann::json read_message(std::chrono::milliseconds timeout, bool& timed_out); + void send_message(const nlohmann::json& message); + void open_streams(const nlohmann::json& subscribed); + void close_all(); + + std::string m_socket_path; + std::vector m_requested; + int m_fd = -1; + std::string m_buffer; + std::vector m_streams; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/livekit_publisher/CMakeLists.txt b/examples/camera_livekit_bridge/livekit_publisher/CMakeLists.txt new file mode 100644 index 0000000000..11b1f5e961 --- /dev/null +++ b/examples/camera_livekit_bridge/livekit_publisher/CMakeLists.txt @@ -0,0 +1,22 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# The LiveKit SDK behind a small interface: join a room, publish a video track, submit H.264 access units, and answer +# the PC's clock-sync requests. No LiveKit type appears in an installed header, and nothing here touches CUDA. +add_library(livekit_publisher STATIC livekit_session.cpp) + +target_include_directories(livekit_publisher + INTERFACE + $ +) + +target_link_libraries(livekit_publisher + PUBLIC + camera_livekit_bridge::video_types + PRIVATE + LiveKit::livekit +) + +target_compile_options(livekit_publisher PRIVATE -Wall -Wextra) + +add_library(camera_livekit_bridge::livekit_publisher ALIAS livekit_publisher) diff --git a/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/clock_sync.hpp b/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/clock_sync.hpp new file mode 100644 index 0000000000..8315f59aae --- /dev/null +++ b/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/clock_sync.hpp @@ -0,0 +1,110 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// The publisher's half of the PC's four-timestamp clock exchange: parse a request, build the reply. All times are +// CLOCK_MONOTONIC nanoseconds, the clock the PC maps. Wire format (little-endian), shared with +// examples/egocap/backend/egocap_web/livekit_clock_sync.py and the egocap emulator: +// request, 13 bytes: 'Q', u32 sequence, i64 t1 (the requester's clock) +// reply, 29 bytes: 'R', u32 sequence, i64 t1 (both echoed), i64 t2 (receipt), i64 t3 (reply send) +// No LiveKit type appears here, so the format is testable on its own. + +#pragma once + +#include + +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +inline constexpr const char* kClockSyncTopic = "isaaccapture.clock-sync"; +inline constexpr std::size_t kClockSyncRequestSize = 13; +inline constexpr std::size_t kClockSyncReplySize = 29; +/// Byte offset of t3 in a reply, so it can be written as late as possible, just before the send. +inline constexpr std::size_t kClockSyncReplyT3Offset = 21; + +struct ClockSyncRequest +{ + std::uint32_t sequence = 0; + std::int64_t t1 = 0; +}; + +inline void put_le(std::vector& out, std::size_t at, std::uint64_t value, std::size_t bytes) +{ + for (std::size_t i = 0; i < bytes; ++i) + out[at + i] = static_cast((value >> (8 * i)) & 0xFF); +} + +inline std::uint64_t get_le(const std::vector& in, std::size_t at, std::size_t bytes) +{ + std::uint64_t value = 0; + for (std::size_t i = 0; i < bytes; ++i) + value |= static_cast(in[at + i]) << (8 * i); + return value; +} + +/// The request in `data`, or nullopt for anything that is not exactly a request. +inline std::optional parse_clock_sync_request(const std::vector& data) +{ + if (data.size() != kClockSyncRequestSize || data[0] != 'Q') + return std::nullopt; + ClockSyncRequest request; + request.sequence = static_cast(get_le(data, 1, 4)); + request.t1 = static_cast(get_le(data, 5, 8)); + return request; +} + +inline std::vector build_clock_sync_reply(const ClockSyncRequest& request, std::int64_t t2, std::int64_t t3) +{ + std::vector reply(kClockSyncReplySize); + reply[0] = 'R'; + put_le(reply, 1, request.sequence, 4); + put_le(reply, 5, static_cast(request.t1), 8); + put_le(reply, 13, static_cast(t2), 8); + put_le(reply, kClockSyncReplyT3Offset, static_cast(t3), 8); + return reply; +} + +/// The requester's half of the same exchange, for a peer that sends requests (a PC-side process over its control +/// link). +inline std::vector build_clock_sync_request(const ClockSyncRequest& request) +{ + std::vector out(kClockSyncRequestSize); + out[0] = 'Q'; + put_le(out, 1, request.sequence, 4); + put_le(out, 5, static_cast(request.t1), 8); + return out; +} + +struct ClockSyncReply +{ + std::uint32_t sequence = 0; + std::int64_t t1 = 0; + std::int64_t t2 = 0; + std::int64_t t3 = 0; +}; + +/// The reply in `data`, or nullopt for anything that is not exactly a reply. +inline std::optional parse_clock_sync_reply(const std::vector& data) +{ + if (data.size() != kClockSyncReplySize || data[0] != 'R') + return std::nullopt; + ClockSyncReply reply; + reply.sequence = static_cast(get_le(data, 1, 4)); + reply.t1 = static_cast(get_le(data, 5, 8)); + reply.t2 = static_cast(get_le(data, 13, 8)); + reply.t3 = static_cast(get_le(data, kClockSyncReplyT3Offset, 8)); + return reply; +} + +/// Writes t3 into an already built reply. +inline void set_clock_sync_reply_t3(std::vector& reply, std::int64_t t3) +{ + put_le(reply, kClockSyncReplyT3Offset, static_cast(t3), 8); +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/livekit_session.hpp b/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/livekit_session.hpp new file mode 100644 index 0000000000..dab693ad1b --- /dev/null +++ b/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/livekit_session.hpp @@ -0,0 +1,75 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "video_track_sink.hpp" + +#include +#include +#include + +namespace camera_bridge +{ + +struct LiveKitSessionImpl; + +/** + * @brief Keeps the LiveKit SDK initialised while it is held. + * + * The SDK is process-wide and must be shut down only when nothing in the process uses it, so every + * user holds one of these and the last release shuts it down. + */ +class LiveKitSdkReference +{ +public: + LiveKitSdkReference(); + ~LiveKitSdkReference(); + LiveKitSdkReference(const LiveKitSdkReference&) = delete; + LiveKitSdkReference& operator=(const LiveKitSdkReference&) = delete; +}; + +struct PublishSpec +{ + /// Name subscribers see for the track. + std::string name; + uint32_t width = 0; + uint32_t height = 0; + /// Upper bound the upstream encoder may use; WebRTC adapts below it. + uint64_t max_bitrate_bps = 0; + double max_fps = 30.0; +}; + +/** + * @brief One participant in one LiveKit room. + * + * The room and the participant's identity come from the join token. A terminal loss of the room + * (the server went away, the participant was removed) sets failed(); the SDK's own bounded + * reconnection runs first and is not wrapped in another supervisor. + */ +class LiveKitSession +{ +public: + /// @throws std::runtime_error if the room cannot be joined. The token is never logged. + static std::unique_ptr connect(const std::string& url, const std::string& token); + + /// Unpublish every track (destroy them first), then leave the room. + ~LiveKitSession(); + + LiveKitSession(const LiveKitSession&) = delete; + LiveKitSession& operator=(const LiveKitSession&) = delete; + + /// Publishes an H.264 camera track that carries pre-encoded access units; the SDK does not encode. + /// @throws std::runtime_error if publication fails. + std::unique_ptr publish_video(const PublishSpec& spec); + + /// True once the room has been lost for good. + bool failed() const; + std::string failure() const; + +private: + explicit LiveKitSession(std::shared_ptr impl); + std::shared_ptr m_impl; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/video_track_sink.hpp b/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/video_track_sink.hpp new file mode 100644 index 0000000000..5be2cbdace --- /dev/null +++ b/examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/video_track_sink.hpp @@ -0,0 +1,52 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include + +#include +#include +#include +#include + +namespace camera_bridge +{ + +/// What the receiving side asked of the encoder since the last call. +struct EncoderFeedback +{ + /// The next frame must be an IDR with its parameter sets. + bool keyframe_requested = false; + /// Latest bitrate target from congestion control, if one was requested. + std::optional target_bitrate_bps; +}; + +/** + * @brief One published video track that carries frames already encoded by hardware. + * + * The SDK forwards each access unit without encoding it again. A single thread must drive one sink. + */ +class VideoTrackSink +{ +public: + virtual ~VideoTrackSink() = default; + + /// Size of the encoded pictures this track was published for. + virtual uint32_t width() const = 0; + virtual uint32_t height() const = 0; + + /// Publishes one access unit; the bytes are copied before this returns. + /// @return false when the SDK did not accept the frame, which can break the reference chain: + /// the caller must follow with a keyframe. + /// @throws std::runtime_error when the frame cannot be published at all. + virtual bool submit(const EncodedVideoFrame& frame) = 0; + + /// Keyframe and rate requests accumulated since the previous call. + virtual EncoderFeedback take_feedback() = 0; + + /// One line naming the transport's view of the stream, or empty if unavailable. + virtual std::string describe() = 0; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/livekit_publisher/livekit_session.cpp b/examples/camera_livekit_bridge/livekit_publisher/livekit_session.cpp new file mode 100644 index 0000000000..fb97e951fa --- /dev/null +++ b/examples/camera_livekit_bridge/livekit_publisher/livekit_session.cpp @@ -0,0 +1,389 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/livekit_publisher/livekit_session.hpp" + +#include "inc/livekit_publisher/clock_sync.hpp" +#include "log.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +namespace +{ + +std::mutex g_sdk_mutex; +int g_sdk_users = 0; + +const char* reason_name(livekit::DisconnectReason reason) +{ + using R = livekit::DisconnectReason; + switch (reason) + { + case R::ClientInitiated: + return "client initiated"; + case R::DuplicateIdentity: + return "another participant joined with the same identity"; + case R::ServerShutdown: + return "the server is shutting down"; + case R::ParticipantRemoved: + return "the participant was removed from the room"; + case R::RoomDeleted: + return "the room was deleted"; + default: + return "connection lost"; + } +} + +} // namespace + +LiveKitSdkReference::LiveKitSdkReference() +{ + std::lock_guard lock(g_sdk_mutex); + if (g_sdk_users == 0) + livekit::initialize(livekit::LogLevel::Warn); + ++g_sdk_users; +} + +LiveKitSdkReference::~LiveKitSdkReference() +{ + std::lock_guard lock(g_sdk_mutex); + if (--g_sdk_users == 0) + livekit::shutdown(); +} + +struct LiveKitSessionImpl : livekit::RoomDelegate +{ + ~LiveKitSessionImpl() override + { + stop_clock_sync(); + // The members below are destroyed before the room: stop callbacks reaching them, then leave the + // room while the SDK is still up (its reference is released after the room). + room.setDelegate(nullptr); + room.disconnect(); + } + + // A clock-sync request from the PC or the web page. t2 is read first, before anything else runs; the reply is + // sent from a small thread so this SDK callback never waits on a send, and t3 is written just before it. + void onUserPacketReceived(livekit::Room&, const livekit::UserDataPacketEvent& event) override + { + const std::int64_t t2 = monotonic_ns(); + if (event.topic != kClockSyncTopic || event.participant == nullptr) + return; + const auto request = parse_clock_sync_request(event.data); + if (!request) + return; + { + std::lock_guard lock(sync_mutex); + // Bounded: a flood drops the oldest request, whose requester would discard a late reply anyway. + if (sync_queue.size() >= kMaxPendingSync) + sync_queue.pop_front(); + sync_queue.push_back({ event.participant->identity(), *request, t2 }); + } + sync_wake.notify_one(); + } + + void start_clock_sync() + { + sync_thread = std::thread([this] { run_clock_sync(); }); + } + + void stop_clock_sync() + { + { + std::lock_guard lock(sync_mutex); + sync_stopping = true; + } + sync_wake.notify_all(); + if (sync_thread.joinable()) + sync_thread.join(); + } + + void run_clock_sync() + { + for (;;) + { + PendingSync pending; + { + std::unique_lock lock(sync_mutex); + sync_wake.wait(lock, [this] { return sync_stopping || !sync_queue.empty(); }); + if (sync_stopping) + return; + pending = std::move(sync_queue.front()); + sync_queue.pop_front(); + } + auto reply = build_clock_sync_reply(pending.request, pending.t2, 0); + try + { + const auto local = room.localParticipant().lock(); + if (!local) + continue; + // t3 as late as possible: nothing but the call follows it. + set_clock_sync_reply_t3(reply, monotonic_ns()); + local->publishData(reply, false, { pending.identity }, kClockSyncTopic); + } + catch (const std::exception& e) + { + log_warn(std::string("clock-sync reply failed: ") + e.what()); + } + } + } + + void onDisconnected(livekit::Room&, const livekit::DisconnectedEvent& event) override + { + if (event.reason == livekit::DisconnectReason::ClientInitiated) + return; + fail(std::string("LiveKit room lost: ") + reason_name(event.reason)); + } + + void onReconnecting(livekit::Room&, const livekit::ReconnectingEvent&) override + { + log_warn("LiveKit connection interrupted; the SDK is reconnecting"); + } + + void onReconnected(livekit::Room&, const livekit::ReconnectedEvent&) override + { + log_info("LiveKit connection restored"); + } + + void fail(const std::string& message) + { + std::lock_guard lock(mutex); + if (!failed.load()) + { + failure = message; + failed.store(true); + log_error(message); + } + } + + struct PendingSync + { + std::string identity; + ClockSyncRequest request; + std::int64_t t2 = 0; + }; + static constexpr std::size_t kMaxPendingSync = 16; + std::mutex sync_mutex; + std::condition_variable sync_wake; + std::deque sync_queue; + bool sync_stopping = false; + std::thread sync_thread; + + // Declared before the room, so destroyed after it. + LiveKitSdkReference sdk; + livekit::Room room; + std::atomic failed{ false }; + mutable std::mutex mutex; + std::string failure; +}; + +namespace +{ + +/// A track fed with H.264 access units that were encoded upstream. +class LiveKitVideoTrack : public VideoTrackSink +{ +public: + LiveKitVideoTrack(std::shared_ptr session, const PublishSpec& spec) + : m_session(std::move(session)), m_width(spec.width), m_height(spec.height) + { + m_source = std::make_shared( + livekit::VideoCodec::H264, static_cast(spec.width), static_cast(spec.height)); + m_track = livekit::LocalVideoTrack::createLocalVideoTrack(spec.name, m_source); + + livekit::TrackPublishOptions options; + options.source = livekit::TrackSource::SOURCE_CAMERA; + options.video_codec = livekit::VideoCodec::H264; + // Pre-encoded access units go through as they are: no software or SDK-selected encoder. + options.video_encoder = livekit::VideoEncoderBackend::PreEncoded; + options.simulcast = false; + options.video_encoding = livekit::VideoEncodingOptions{ spec.max_bitrate_bps, spec.max_fps }; + options.frame_metadata_features = livekit::FrameMetadataFeatures{ true, true, true }; + + auto participant = m_session->room.localParticipant().lock(); + if (!participant) + throw std::runtime_error("LiveKit: no local participant to publish '" + spec.name + "'"); + try + { + participant->publishTrack(m_track, options); + } + catch (const std::exception& e) + { + throw std::runtime_error("LiveKit: cannot publish track '" + spec.name + "': " + e.what()); + } + m_participant = participant; + } + + ~LiveKitVideoTrack() override + { + try + { + if (auto participant = m_participant.lock()) + participant->unpublishTrack(m_track->sid()); + } + catch (const std::exception& e) + { + log_warn(std::string("LiveKit: unpublish failed: ") + e.what()); + } + } + + uint32_t width() const override + { + return m_width; + } + uint32_t height() const override + { + return m_height; + } + + bool submit(const EncodedVideoFrame& frame) override + { + if (m_session->failed.load()) + throw std::runtime_error(m_session->failure); + + livekit::EncodedVideoSource::Frame out; + out.is_keyframe = frame.keyframe; + out.width = m_width; + out.height = m_height; + out.timestamp_us = static_cast(frame.timestamp_us); + out.data.assign(frame.data, frame.data + frame.size); + livekit::VideoFrameMetadata metadata; + // Zero means the publisher cannot vouch for a capture time: the frame is sent without one, and a consumer + // that needs a timestamp (a recording) says so instead of using a guessed one. + if (frame.user_timestamp_us != 0) + metadata.user_timestamp_us = frame.user_timestamp_us; + metadata.frame_id = frame.frame_id; + std::vector tsc(8); + for (int i = 0; i < 8; ++i) + tsc[i] = static_cast((frame.capture_tsc >> (8 * i)) & 0xFF); + metadata.user_data = std::move(tsc); + out.metadata = std::move(metadata); + + try + { + return m_source->captureFrame(out); + } + catch (const std::exception& e) + { + throw std::runtime_error(std::string("LiveKit: cannot submit a frame: ") + e.what()); + } + } + + EncoderFeedback take_feedback() override + { + EncoderFeedback feedback; + try + { + const auto pending = m_source->takeFeedback(); + feedback.keyframe_requested = pending.keyframe_requested; + if (pending.rate_control && pending.rate_control->target_bitrate_bps > 0) + feedback.target_bitrate_bps = static_cast(pending.rate_control->target_bitrate_bps); + } + catch (const std::exception& e) + { + throw std::runtime_error(std::string("LiveKit: cannot read encoder feedback: ") + e.what()); + } + return feedback; + } + + // Called from the polling thread, which must not wait on the network: this returns the last answer + // and starts, or collects, a stats request that was asked for earlier. + std::string describe() override + { + try + { + if (m_stats.valid() && m_stats.wait_for(std::chrono::seconds(0)) == std::future_status::ready) + { + for (const auto& entry : m_stats.get()) + { + const auto* out = std::get_if(&entry.stats); + if (!out || out->outbound.frame_width == 0) + continue; + const auto& o = out->outbound; + m_description = "encoder " + o.encoder_implementation + ", " + std::to_string(o.frame_width) + "x" + + std::to_string(o.frame_height) + " @ " + + std::to_string(static_cast(o.frames_per_second)) + " fps, " + + std::to_string(static_cast(o.target_bitrate / 1000)) + " kbps target"; + break; + } + } + if (!m_stats.valid()) + m_stats = m_track->getStats(); + } + catch (const std::exception&) + { + } + return m_description; + } + +private: + std::shared_ptr m_session; + uint32_t m_width; + uint32_t m_height; + std::shared_ptr m_source; + std::shared_ptr m_track; + std::weak_ptr m_participant; + decltype(std::declval().getStats()) m_stats; + std::string m_description; +}; + +} // namespace + +LiveKitSession::LiveKitSession(std::shared_ptr impl) : m_impl(std::move(impl)) +{ +} + +LiveKitSession::~LiveKitSession() = default; + +std::unique_ptr LiveKitSession::connect(const std::string& url, const std::string& token) +{ + auto impl = std::make_shared(); + impl->room.setDelegate(impl.get()); + + livekit::RoomOptions options; + options.auto_subscribe = false; // this process only publishes + options.dynacast = false; + if (!impl->room.connect(url, token, options)) + { + throw std::runtime_error("LiveKit: cannot join the room at " + url + + " (check the URL, that the server is reachable and that the token is valid)"); + } + log_info("Joined LiveKit room at " + url); + impl->start_clock_sync(); + return std::unique_ptr(new LiveKitSession(std::move(impl))); +} + +std::unique_ptr LiveKitSession::publish_video(const PublishSpec& spec) +{ + return std::make_unique(m_impl, spec); +} + +bool LiveKitSession::failed() const +{ + return m_impl->failed.load(); +} + +std::string LiveKitSession::failure() const +{ + std::lock_guard lock(m_impl->mutex); + return m_impl->failure; +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/livekit_publisher/log.hpp b/examples/camera_livekit_bridge/livekit_publisher/log.hpp new file mode 100644 index 0000000000..d3285767d6 --- /dev/null +++ b/examples/camera_livekit_bridge/livekit_publisher/log.hpp @@ -0,0 +1,27 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +namespace camera_bridge +{ + +inline void log_info(const std::string& message) +{ + std::cerr << "[info] " << message << "\n"; +} + +inline void log_warn(const std::string& message) +{ + std::cerr << "[warn] " << message << "\n"; +} + +inline void log_error(const std::string& message) +{ + std::cerr << "[error] " << message << "\n"; +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/nvenc/CMakeLists.txt b/examples/camera_livekit_bridge/nvenc/CMakeLists.txt new file mode 100644 index 0000000000..b5b9ab2f10 --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/CMakeLists.txt @@ -0,0 +1,29 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# H.264 NVENC encoder over the device's primary CUDA context. +add_library(nvenc STATIC + cuda_context.cpp + nvenc_encoder.cpp +) + +set_target_properties(nvenc PROPERTIES POSITION_INDEPENDENT_CODE ON) + +target_include_directories(nvenc + INTERFACE + $ +) + +target_link_libraries(nvenc + PUBLIC + camera_livekit_bridge::video_types + CUDA::cudart + CUDA::cuda_driver + PRIVATE + nv_codec_headers + ${CMAKE_DL_LIBS} +) + +target_compile_options(nvenc PRIVATE -Wall -Wextra) + +add_library(camera_livekit_bridge::nvenc ALIAS nvenc) diff --git a/examples/camera_livekit_bridge/nvenc/cuda_context.cpp b/examples/camera_livekit_bridge/nvenc/cuda_context.cpp new file mode 100644 index 0000000000..cf83ec2430 --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/cuda_context.cpp @@ -0,0 +1,50 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/nvenc/cuda_context.hpp" + +#include + +namespace camera_bridge +{ + +void check_cu(CUresult result, const char* what) +{ + if (result == CUDA_SUCCESS) + return; + const char* message = nullptr; + cuGetErrorString(result, &message); + throw std::runtime_error(std::string(what) + ": " + (message ? message : "unknown CUDA driver error")); +} + +CudaContext::CudaContext(int device_index) +{ + check_cu(cuInit(0), "cuInit"); + check_cu(cuDeviceGet(&m_device, device_index), "cuDeviceGet"); + char name[128] = {}; + check_cu(cuDeviceGetName(name, sizeof(name), m_device), "cuDeviceGetName"); + m_name = name; + check_cu(cuDevicePrimaryCtxRetain(&m_context, m_device), "cuDevicePrimaryCtxRetain"); + make_current(); + // Bind the runtime API to the same primary context. + if (cudaSetDevice(device_index) != cudaSuccess) + throw std::runtime_error("cudaSetDevice failed"); +} + +CudaContext::~CudaContext() +{ + if (m_context) + cuDevicePrimaryCtxRelease(m_device); +} + +void CudaContext::make_current_noexcept() const noexcept +{ + cuCtxSetCurrent(m_context); +} + +void CudaContext::make_current() const +{ + check_cu(cuCtxSetCurrent(m_context), "cuCtxSetCurrent"); +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/nvenc/inc/nvenc/cuda_context.hpp b/examples/camera_livekit_bridge/nvenc/inc/nvenc/cuda_context.hpp new file mode 100644 index 0000000000..36b90316a8 --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/inc/nvenc/cuda_context.hpp @@ -0,0 +1,49 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +namespace camera_bridge +{ + +/// Throws std::runtime_error naming `what` when a CUDA driver call failed. +void check_cu(CUresult result, const char* what); + +/** + * @brief The device's primary CUDA context, shared by decode, the GPU operation and encode. + * + * The runtime API (kernels, streams) and the driver API (NVDEC, NVENC) both use the primary context, so surfaces + * are visible to all of them. Every thread that touches CUDA calls make_current() once. + */ +class CudaContext +{ +public: + explicit CudaContext(int device_index = 0); + ~CudaContext(); + + CudaContext(const CudaContext&) = delete; + CudaContext& operator=(const CudaContext&) = delete; + + void make_current() const; + /// For cleanup paths that must not throw: the result is ignored. + void make_current_noexcept() const noexcept; + CUcontext handle() const + { + return m_context; + } + const std::string& device_name() const + { + return m_name; + } + +private: + CUdevice m_device = 0; + CUcontext m_context = nullptr; + std::string m_name; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp b/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp new file mode 100644 index 0000000000..eb654cc47f --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp @@ -0,0 +1,80 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuda_context.hpp" + +#include + +#include +#include +#include +#include +#include + +namespace camera_bridge +{ + +struct EncoderConfig +{ + Geometry geometry; + uint32_t fps_numerator = 30; + uint32_t fps_denominator = 1; + uint32_t bitrate_bps = 12'000'000; + /// Frames between IDRs; a stream recovers from loss only at an IDR, so keep it short. + uint32_t idr_period = 30; + bool full_range = false; + /// Reusable input surfaces, registered with the encoder once at startup. + uint32_t input_surfaces = 3; +}; + +/** + * @brief NV12 device surfaces to H.264 access units on NVENC, tuned for low latency. + * + * The encoder owns a small pool of CUDA surfaces registered once; the GPU operation writes its output straight + * into one, so there is no copy and no per-frame registration. Low-latency settings: the ultra-low-latency tuning, + * no B-frames, no look-ahead, constant bitrate with a one-frame buffer, parameter sets repeated before every IDR. + * encode() returns the access unit for the picture it was given. + */ +class NvencEncoder +{ +public: + /// A pooled input surface. Write the picture into `view`, then pass the same `slot` to encode(). + struct Surface + { + Nv12View view; + uint32_t slot = 0; + }; + + NvencEncoder(CudaContext& context, const EncoderConfig& config); + ~NvencEncoder(); + + NvencEncoder(const NvencEncoder&) = delete; + NvencEncoder& operator=(const NvencEncoder&) = delete; + + /// The next free input surface, or nullopt-equivalent `false` when all are in flight. + bool acquire(Surface& out); + + /// Give back a surface acquired but not encoded (an error path). Idempotent, and encode() releases its own. + void release(const Surface& surface); + + /// Encode the picture written to `surface` after `stream` has completed the work that wrote it, and release the + /// surface. `key` comes back on the packet. `force_idr` makes this picture an IDR with parameter sets. + /// Throws std::runtime_error on an encoder error. + EncodedPacket encode(const Surface& surface, cudaStream_t stream, uint64_t key, bool force_idr); + + /// Changes the target bitrate (and the one-frame rate-control buffer with it) without restarting the stream or + /// forcing an IDR. Call between encode() calls, from the thread that encodes. Throws std::runtime_error on an + /// encoder error. + void set_bitrate(uint32_t bitrate_bps); + + /// The bitrate the encoder is currently configured for. + uint32_t bitrate() const; + +private: + struct Impl; + std::unique_ptr m_impl; +}; + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp new file mode 100644 index 0000000000..11245481d8 --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp @@ -0,0 +1,327 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/nvenc/nvenc_encoder.hpp" + +#include + +#include +#include +#include + +namespace camera_bridge +{ + +namespace +{ + +void check_nv(NVENCSTATUS status, const char* what) +{ + if (status != NV_ENC_SUCCESS) + throw std::runtime_error(std::string("NvencEncoder: ") + what + " failed with status " + std::to_string(status)); +} + +} // namespace + +struct NvencEncoder::Impl +{ + CudaContext& context; + EncoderConfig config; + + NvencFunctions* nvenc = nullptr; + NV_ENCODE_API_FUNCTION_LIST fn{}; + void* encoder = nullptr; + NV_ENC_OUTPUT_PTR bitstream = nullptr; + + struct Slot + { + CUdeviceptr ptr = 0; + size_t pitch = 0; + NV_ENC_REGISTERED_PTR registered = nullptr; + bool in_use = false; + }; + std::vector slots; + // Kept after configure() so the rate can be changed later with the same settings. + NV_ENC_CONFIG enc_config{}; + NV_ENC_INITIALIZE_PARAMS init_params{}; + uint32_t bitrate_bps = 0; + + Impl(CudaContext& ctx, const EncoderConfig& cfg) : context(ctx), config(cfg) + { + if (config.geometry.width == 0 || config.geometry.height == 0 || (config.geometry.width % 2) != 0 || + (config.geometry.height % 2) != 0) + throw std::invalid_argument("NvencEncoder: NV12 needs non-zero even dimensions"); + try + { + context.make_current(); + if (nvenc_load_functions(&nvenc, nullptr) < 0 || nvenc == nullptr) + throw std::runtime_error("NvencEncoder: libnvidia-encode is not available"); + fn.version = NV_ENCODE_API_FUNCTION_LIST_VER; + check_nv(nvenc->NvEncodeAPICreateInstance(&fn), "NvEncodeAPICreateInstance"); + + NV_ENC_OPEN_ENCODE_SESSION_EX_PARAMS open{}; + open.version = NV_ENC_OPEN_ENCODE_SESSION_EX_PARAMS_VER; + open.deviceType = NV_ENC_DEVICE_TYPE_CUDA; + open.device = context.handle(); + open.apiVersion = NVENCAPI_VERSION; + check_nv(fn.nvEncOpenEncodeSessionEx(&open, &encoder), "nvEncOpenEncodeSessionEx"); + + configure(); + allocate_surfaces(); + + NV_ENC_CREATE_BITSTREAM_BUFFER out{}; + out.version = NV_ENC_CREATE_BITSTREAM_BUFFER_VER; + check_nv(fn.nvEncCreateBitstreamBuffer(encoder, &out), "nvEncCreateBitstreamBuffer"); + bitstream = out.bitstreamBuffer; + } + catch (...) + { + // A constructor that throws does not run the destructor, so a failure part way releases what exists. + shutdown(); + throw; + } + } + + ~Impl() + { + shutdown(); + } + + // Releases the session, the registered surfaces and the loaded library; safe to call on a partly built object. + void shutdown() noexcept + { + context.make_current_noexcept(); + if (encoder) + { + for (Slot& s : slots) + { + if (s.registered) + fn.nvEncUnregisterResource(encoder, s.registered); + if (s.ptr) + cuMemFree(s.ptr); + s = Slot{}; + } + if (bitstream) + fn.nvEncDestroyBitstreamBuffer(encoder, bitstream); + fn.nvEncDestroyEncoder(encoder); + encoder = nullptr; + bitstream = nullptr; + } + if (nvenc) + nvenc_free_functions(&nvenc); + } + + void configure() + { + const GUID codec = NV_ENC_CODEC_H264_GUID; + const GUID preset = NV_ENC_PRESET_P1_GUID; + const NV_ENC_TUNING_INFO tuning = NV_ENC_TUNING_INFO_ULTRA_LOW_LATENCY; + + NV_ENC_PRESET_CONFIG preset_config{}; + preset_config.version = NV_ENC_PRESET_CONFIG_VER; + preset_config.presetCfg.version = NV_ENC_CONFIG_VER; + check_nv(fn.nvEncGetEncodePresetConfigEx(encoder, codec, preset, tuning, &preset_config), + "nvEncGetEncodePresetConfigEx"); + + enc_config = preset_config.presetCfg; + enc_config.version = NV_ENC_CONFIG_VER; + enc_config.gopLength = config.idr_period; + enc_config.frameIntervalP = 1; // no B-frames + enc_config.rcParams.rateControlMode = NV_ENC_PARAMS_RC_CBR; + apply_rate(config.bitrate_bps); + enc_config.rcParams.lookaheadDepth = 0; + NV_ENC_CONFIG_H264& h264 = enc_config.encodeCodecConfig.h264Config; + h264.idrPeriod = config.idr_period; + h264.repeatSPSPPS = 1; + h264.chromaFormatIDC = 1; + h264.h264VUIParameters.videoSignalTypePresentFlag = 1; + h264.h264VUIParameters.videoFormat = NV_ENC_VUI_VIDEO_FORMAT_UNSPECIFIED; + h264.h264VUIParameters.videoFullRangeFlag = config.full_range ? 1 : 0; + + NV_ENC_INITIALIZE_PARAMS& init = init_params; + init = {}; + init.version = NV_ENC_INITIALIZE_PARAMS_VER; + init.encodeGUID = codec; + init.presetGUID = preset; + init.tuningInfo = tuning; + init.encodeWidth = config.geometry.width; + init.encodeHeight = config.geometry.height; + init.darWidth = config.geometry.width; + init.darHeight = config.geometry.height; + init.maxEncodeWidth = config.geometry.width; + init.maxEncodeHeight = config.geometry.height; + init.frameRateNum = config.fps_numerator; + init.frameRateDen = config.fps_denominator; + init.enablePTD = 1; + init.encodeConfig = &enc_config; + check_nv(fn.nvEncInitializeEncoder(encoder, &init), "nvEncInitializeEncoder"); + } + + // Sets the target bitrate, with a buffer of one frame's worth of bits so the output stays close to the + // instantaneous rate. + void apply_rate(uint32_t bps) + { + enc_config.rcParams.averageBitRate = bps; + enc_config.rcParams.maxBitRate = bps; + const uint32_t frame_bits = + static_cast(static_cast(bps) * config.fps_denominator / config.fps_numerator); + enc_config.rcParams.vbvBufferSize = frame_bits; + enc_config.rcParams.vbvInitialDelay = frame_bits; + bitrate_bps = bps; + } + + void allocate_surfaces() + { + slots.resize(config.input_surfaces); + const size_t rows = static_cast(config.geometry.height) * 3 / 2; + for (Slot& s : slots) + { + // One allocation: luma rows then chroma rows, the layout NVENC reads for NV12. + check_cu(cuMemAllocPitch(&s.ptr, &s.pitch, config.geometry.width, rows, 16), "cuMemAllocPitch"); + NV_ENC_REGISTER_RESOURCE reg{}; + reg.version = NV_ENC_REGISTER_RESOURCE_VER; + reg.resourceType = NV_ENC_INPUT_RESOURCE_TYPE_CUDADEVICEPTR; + reg.width = config.geometry.width; + reg.height = config.geometry.height; + reg.pitch = static_cast(s.pitch); + reg.resourceToRegister = reinterpret_cast(s.ptr); + reg.bufferFormat = NV_ENC_BUFFER_FORMAT_NV12; + reg.bufferUsage = NV_ENC_INPUT_IMAGE; + check_nv(fn.nvEncRegisterResource(encoder, ®), "nvEncRegisterResource"); + s.registered = reg.registeredResource; + } + } +}; + +NvencEncoder::NvencEncoder(CudaContext& context, const EncoderConfig& config) + : m_impl(std::make_unique(context, config)) +{ +} + +NvencEncoder::~NvencEncoder() = default; + +void NvencEncoder::release(const Surface& surface) +{ + m_impl->slots.at(surface.slot).in_use = false; +} + +bool NvencEncoder::acquire(Surface& out) +{ + Impl& e = *m_impl; + for (uint32_t i = 0; i < e.slots.size(); ++i) + { + Impl::Slot& s = e.slots[i]; + if (s.in_use) + continue; + s.in_use = true; + out.slot = i; + out.view.geometry = e.config.geometry; + out.view.color.full_range = e.config.full_range; + out.view.y = reinterpret_cast(s.ptr); + out.view.y_pitch = s.pitch; + out.view.uv = reinterpret_cast(s.ptr) + s.pitch * e.config.geometry.height; + out.view.uv_pitch = s.pitch; + return true; + } + return false; +} + +void NvencEncoder::set_bitrate(uint32_t bitrate_bps) +{ + Impl& e = *m_impl; + if (bitrate_bps == 0) + throw std::invalid_argument("NvencEncoder: the bitrate must be positive"); + e.context.make_current(); + const uint32_t previous = e.bitrate_bps; + e.apply_rate(bitrate_bps); + NV_ENC_RECONFIGURE_PARAMS reconfigure{}; + reconfigure.version = NV_ENC_RECONFIGURE_PARAMS_VER; + reconfigure.reInitEncodeParams = e.init_params; + reconfigure.reInitEncodeParams.encodeConfig = &e.enc_config; + reconfigure.resetEncoder = 0; + reconfigure.forceIDR = 0; + try + { + check_nv(e.fn.nvEncReconfigureEncoder(e.encoder, &reconfigure), "nvEncReconfigureEncoder"); + } + catch (...) + { + // The encoder kept its old rate, so the reported rate and the settings used for the next change must too. + e.apply_rate(previous); + throw; + } +} + +uint32_t NvencEncoder::bitrate() const +{ + return m_impl->bitrate_bps; +} + +EncodedPacket NvencEncoder::encode(const Surface& surface, cudaStream_t stream, uint64_t key, bool force_idr) +{ + Impl& e = *m_impl; + Impl::Slot& slot = e.slots.at(surface.slot); + EncodedPacket packet; + try + { + e.context.make_current(); + // NVENC reads the surface directly; wait for this stream's work (the warp) to finish writing it. + if (cudaStreamSynchronize(stream) != cudaSuccess) + throw std::runtime_error("NvencEncoder: the GPU work feeding the encoder failed"); + + NV_ENC_MAP_INPUT_RESOURCE map{}; + map.version = NV_ENC_MAP_INPUT_RESOURCE_VER; + map.registeredResource = slot.registered; + check_nv(e.fn.nvEncMapInputResource(e.encoder, &map), "nvEncMapInputResource"); + // Unmap on every exit path: a registered surface that stays mapped cannot be reused or freed. + struct MapGuard + { + Impl& impl; + NV_ENC_INPUT_PTR mapped; + ~MapGuard() + { + impl.fn.nvEncUnmapInputResource(impl.encoder, mapped); + } + } map_guard{ e, map.mappedResource }; + + NV_ENC_PIC_PARAMS pic{}; + pic.version = NV_ENC_PIC_PARAMS_VER; + pic.inputBuffer = map.mappedResource; + pic.bufferFmt = map.mappedBufferFmt; + pic.inputWidth = e.config.geometry.width; + pic.inputHeight = e.config.geometry.height; + pic.inputPitch = static_cast(slot.pitch); + pic.outputBitstream = e.bitstream; + pic.pictureStruct = NV_ENC_PIC_STRUCT_FRAME; + pic.inputTimeStamp = key; + if (force_idr) + pic.encodePicFlags = NV_ENC_PIC_FLAG_FORCEIDR | NV_ENC_PIC_FLAG_OUTPUT_SPSPPS; + check_nv(e.fn.nvEncEncodePicture(e.encoder, &pic), "nvEncEncodePicture"); + + NV_ENC_LOCK_BITSTREAM lock{}; + lock.version = NV_ENC_LOCK_BITSTREAM_VER; + lock.outputBitstream = e.bitstream; + check_nv(e.fn.nvEncLockBitstream(e.encoder, &lock), "nvEncLockBitstream"); + struct LockGuard + { + Impl& impl; + ~LockGuard() + { + impl.fn.nvEncUnlockBitstream(impl.encoder, impl.bitstream); + } + } lock_guard{ e }; + const auto* bytes = static_cast(lock.bitstreamBufferPtr); + packet.data.assign(bytes, bytes + lock.bitstreamSizeInBytes); + packet.key = lock.outputTimeStamp; + packet.keyframe = lock.pictureType == NV_ENC_PIC_TYPE_IDR; + } + catch (...) + { + slot.in_use = false; + throw; + } + slot.in_use = false; + return packet; +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/probe/CMakeLists.txt b/examples/camera_livekit_bridge/probe/CMakeLists.txt new file mode 100644 index 0000000000..f086c7fd54 --- /dev/null +++ b/examples/camera_livekit_bridge/probe/CMakeLists.txt @@ -0,0 +1,12 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +add_executable(camera_ipc_probe main.cpp) + +target_link_libraries(camera_ipc_probe + PRIVATE + camera_livekit_bridge::ipc_client + CUDA::cudart +) + +target_compile_options(camera_ipc_probe PRIVATE -Wall -Wextra) diff --git a/examples/camera_livekit_bridge/probe/main.cpp b/examples/camera_livekit_bridge/probe/main.cpp new file mode 100644 index 0000000000..b359289fde --- /dev/null +++ b/examples/camera_livekit_bridge/probe/main.cpp @@ -0,0 +1,87 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +// Connects to camera_receiver's CUDA IPC socket and prints per-frame metadata with a small checksum of the frame's +// first pixels, to show that frames reach this process without leaving the GPU. + +#include + +#include +#include +#include +#include +#include +#include + +namespace +{ + +std::uint64_t first_row_checksum(const camera_bridge::IpcSlot& slot) +{ + std::vector host(std::min(4096, slot.nbytes)); + if (host.empty() || cudaMemcpy(host.data(), slot.device_ptr, host.size(), cudaMemcpyDeviceToHost) != cudaSuccess) + return 0; + std::uint64_t sum = 0; + for (const unsigned char byte : host) + sum += byte; + return sum; +} + +} // namespace + +int main(int argc, char** argv) +{ + std::string socket_path = "/tmp/camera_receiver_avs1_cuda_ipc.sock"; + std::vector streams; + int frames = 60; + for (int i = 1; i < argc; ++i) + { + const std::string arg = argv[i]; + if (arg == "--socket" && i + 1 < argc) + socket_path = argv[++i]; + else if (arg == "--stream" && i + 1 < argc) + streams.emplace_back(argv[++i]); + else if (arg == "--frames" && i + 1 < argc) + frames = std::atoi(argv[++i]); + else + { + std::cerr << "usage: camera_ipc_probe [--socket PATH] [--stream NAME]... [--frames N]\n"; + return 2; + } + } + + try + { + camera_bridge::CudaIpcClient client(socket_path, streams); + client.connect(); + for (const auto& stream : client.streams()) + { + std::cout << "stream " << stream.name << " " << stream.width << "x" << stream.height << " " + << stream.pixel_format << " slots=" << stream.slots.size() << " layout=" << stream.layout.dump() + << "\n"; + } + camera_bridge::FrameNotice notice; + for (int count = 0; count < frames;) + { + if (!client.next_frame(notice, std::chrono::seconds(5))) + { + std::cerr << "no frame for 5 s\n"; + return 1; + } + const auto& slot = client.slot(notice); + const auto& meta = notice.metadata; + std::cout << notice.stream << " slot=" << notice.slot << " gen=" << notice.slot_generation + << " stream_sequence=" << meta.value("stream_sequence", -1) + << " session=" << meta.value("session_id", std::uint64_t{ 0 }) + << " sensor_sof_ns=" << (meta.contains("sensor_sof_ns") ? meta["sensor_sof_ns"].dump() : "null") + << " checksum=" << first_row_checksum(slot) << "\n"; + ++count; + } + } + catch (const std::exception& error) + { + std::cerr << "error: " << error.what() << "\n"; + return 1; + } + return 0; +} diff --git a/examples/camera_livekit_bridge/tests/CMakeLists.txt b/examples/camera_livekit_bridge/tests/CMakeLists.txt new file mode 100644 index 0000000000..0e6512e01b --- /dev/null +++ b/examples/camera_livekit_bridge/tests/CMakeLists.txt @@ -0,0 +1,10 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# One test executable per subdirectory (each with its own CMakeLists.txt); Python tests live in tests/python. +file(GLOB _test_dirs LIST_DIRECTORIES true RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*) +foreach(_dir IN LISTS _test_dirs) + if(IS_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/${_dir} AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${_dir}/CMakeLists.txt) + add_subdirectory(${_dir}) + endif() +endforeach() diff --git a/examples/camera_livekit_bridge/tests/clock_sync/CMakeLists.txt b/examples/camera_livekit_bridge/tests/clock_sync/CMakeLists.txt new file mode 100644 index 0000000000..c952e2ea7d --- /dev/null +++ b/examples/camera_livekit_bridge/tests/clock_sync/CMakeLists.txt @@ -0,0 +1,7 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# The header needs no SDK, but linking the publisher target pulls the SDK libraries in at link time. +add_executable(test_clock_sync test_clock_sync.cpp) +target_link_libraries(test_clock_sync PRIVATE camera_livekit_bridge::livekit_publisher) +add_test(NAME clock_sync COMMAND test_clock_sync) diff --git a/examples/camera_livekit_bridge/tests/clock_sync/test_clock_sync.cpp b/examples/camera_livekit_bridge/tests/clock_sync/test_clock_sync.cpp new file mode 100644 index 0000000000..59567f9173 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/clock_sync/test_clock_sync.cpp @@ -0,0 +1,133 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include +#include +#include + +using namespace camera_bridge; + +namespace +{ + +int g_failures = 0; + +#define CHECK(cond) \ + do \ + { \ + if (!(cond)) \ + { \ + std::fprintf(stderr, "FAIL %s:%d: %s\n", __FILE__, __LINE__, #cond); \ + ++g_failures; \ + } \ + } while (0) + +using Bytes = std::vector; + +// sequence 7, t1 123456789 (0x075BCD15) +const Bytes kRequest = { 'Q', 0x07, 0x00, 0x00, 0x00, 0x15, 0xCD, 0x5B, 0x07, 0x00, 0x00, 0x00, 0x00 }; + +// t2 = 0x0102030405060708, t3 = 1000000000000 (0x000000E8D4A51000) +const Bytes kReply = { 'R', 0x07, 0x00, 0x00, 0x00, 0x15, 0xCD, 0x5B, 0x07, 0x00, 0x00, 0x00, 0x00, 0x08, 0x07, + 0x06, 0x05, 0x04, 0x03, 0x02, 0x01, 0x00, 0x10, 0xA5, 0xD4, 0xE8, 0x00, 0x00, 0x00 }; + +void test_parse_request() +{ + CHECK(kRequest.size() == kClockSyncRequestSize); + const auto r = parse_clock_sync_request(kRequest); + CHECK(r.has_value()); + if (r) + { + CHECK(r->sequence == 7); + CHECK(r->t1 == 123456789); + } + + // High bits: negative t1 and a sequence with the top byte set. + Bytes high = { 'Q', 0x78, 0x56, 0x34, 0x82, 0xFF, 0xFF, 0xFF, 0xFF, 0xFF, 0xFF, 0xFF, 0xFF }; + const auto h = parse_clock_sync_request(high); + CHECK(h.has_value()); + if (h) + { + CHECK(h->sequence == 0x82345678u); + CHECK(h->t1 == -1); + } +} + +void test_parse_request_rejects() +{ + CHECK(!parse_clock_sync_request({}).has_value()); + Bytes short_req(kRequest.begin(), kRequest.end() - 1); + CHECK(!parse_clock_sync_request(short_req).has_value()); + Bytes long_req = kRequest; + long_req.push_back(0); + CHECK(!parse_clock_sync_request(long_req).has_value()); + Bytes wrong_tag = kRequest; + wrong_tag[0] = 'R'; + CHECK(!parse_clock_sync_request(wrong_tag).has_value()); + wrong_tag[0] = 'q'; + CHECK(!parse_clock_sync_request(wrong_tag).has_value()); + // A reply-sized buffer is not a request. + CHECK(!parse_clock_sync_request(kReply).has_value()); +} + +void test_build_reply() +{ + const Bytes reply = build_clock_sync_reply(ClockSyncRequest{ 7, 123456789 }, 0x0102030405060708LL, 1000000000000LL); + CHECK(reply.size() == kClockSyncReplySize); + CHECK(reply == kReply); + + const auto parsed = parse_clock_sync_reply(reply); + CHECK(parsed.has_value()); + if (parsed) + { + CHECK(parsed->sequence == 7); + CHECK(parsed->t1 == 123456789); + CHECK(parsed->t2 == 0x0102030405060708LL); + CHECK(parsed->t3 == 1000000000000LL); + } +} + +void test_set_t3() +{ + Bytes reply = build_clock_sync_reply(ClockSyncRequest{ 7, 123456789 }, 0x0102030405060708LL, 0); + const Bytes before = reply; + set_clock_sync_reply_t3(reply, 1000000000000LL); + CHECK(reply == kReply); + // Only the eight bytes at the t3 offset differ. + for (std::size_t i = 0; i < reply.size(); ++i) + if (i < kClockSyncReplyT3Offset || i >= kClockSyncReplyT3Offset + 8) + CHECK(reply[i] == before[i]); + CHECK(kClockSyncReplyT3Offset == 21); +} + +void test_request_round_trip_and_rejects() +{ + CHECK(build_clock_sync_request(ClockSyncRequest{ 7, 123456789 }) == kRequest); + CHECK(!parse_clock_sync_reply(kRequest).has_value()); + Bytes bad = kReply; + bad[0] = 'Q'; + CHECK(!parse_clock_sync_reply(bad).has_value()); + bad = kReply; + bad.pop_back(); + CHECK(!parse_clock_sync_reply(bad).has_value()); +} + +} // namespace + +int main() +{ + test_parse_request(); + test_parse_request_rejects(); + test_build_reply(); + test_set_t3(); + test_request_round_trip_and_rejects(); + if (g_failures != 0) + { + std::fprintf(stderr, "%d check(s) failed\n", g_failures); + return 1; + } + std::printf("clock_sync tests passed\n"); + return 0; +} diff --git a/examples/camera_livekit_bridge/tests/gpu_ops/CMakeLists.txt b/examples/camera_livekit_bridge/tests/gpu_ops/CMakeLists.txt new file mode 100644 index 0000000000..ae05786bb8 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/gpu_ops/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_ops test_box_downscale.cpp) +target_link_libraries(test_gpu_ops PRIVATE camera_livekit_bridge::gpu_ops CUDA::cudart) +add_test(NAME gpu_ops COMMAND test_gpu_ops) +set_tests_properties(gpu_ops PROPERTIES SKIP_RETURN_CODE 77) diff --git a/examples/camera_livekit_bridge/tests/gpu_ops/test_box_downscale.cpp b/examples/camera_livekit_bridge/tests/gpu_ops/test_box_downscale.cpp new file mode 100644 index 0000000000..c639657f7a --- /dev/null +++ b/examples/camera_livekit_bridge/tests/gpu_ops/test_box_downscale.cpp @@ -0,0 +1,228 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include +#include +#include + +using namespace camera_bridge; + +namespace +{ + +int g_failures = 0; + +#define CHECK(cond) \ + do \ + { \ + if (!(cond)) \ + { \ + std::fprintf(stderr, "FAIL %s:%d: %s\n", __FILE__, __LINE__, #cond); \ + ++g_failures; \ + } \ + } while (0) + +struct HostNv12 +{ + uint32_t width = 0; + uint32_t height = 0; + std::vector y; // width * height, tightly packed + std::vector uv; // width * (height / 2), interleaved U,V +}; + +HostNv12 make_pattern(uint32_t width, uint32_t height, bool blocky, uint32_t block) +{ + HostNv12 img; + img.width = width; + img.height = height; + img.y.resize(static_cast(width) * height); + img.uv.resize(static_cast(width) * (height / 2)); + for (uint32_t r = 0; r < height; ++r) + for (uint32_t c = 0; c < width; ++c) + { + uint8_t v; + if (blocky) + { + // Within a block the values are base, base+1, base+2, base+3 so the average rounds up or down. + const uint32_t base = ((r / block) * 37 + (c / block) * 11) % 240; + v = static_cast(base + ((r % block) + (c % block)) % 4); + } + else + v = static_cast((r * 3 + c * 5) % 256); + img.y[static_cast(r) * width + c] = v; + } + for (uint32_t r = 0; r < height / 2; ++r) + for (uint32_t c = 0; c < width; ++c) + img.uv[static_cast(r) * width + c] = + static_cast(c % 2 == 0 ? (r * 7 + c * 3) % 256 : (r * 13 + c * 5 + 90) % 256); + return img; +} + +HostNv12 reference_downscale(const HostNv12& src, uint32_t d) +{ + HostNv12 out; + out.width = src.width / d; + out.height = src.height / d; + out.y.resize(static_cast(out.width) * out.height); + out.uv.resize(static_cast(out.width) * (out.height / 2)); + const uint32_t n = d * d; + for (uint32_t r = 0; r < out.height; ++r) + for (uint32_t c = 0; c < out.width; ++c) + { + uint32_t sum = 0; + for (uint32_t j = 0; j < d; ++j) + for (uint32_t i = 0; i < d; ++i) + sum += src.y[static_cast(r * d + j) * src.width + c * d + i]; + out.y[static_cast(r) * out.width + c] = static_cast((sum + n / 2) / n); + } + for (uint32_t r = 0; r < out.height / 2; ++r) + for (uint32_t c = 0; c < out.width / 2; ++c) + for (uint32_t comp = 0; comp < 2; ++comp) + { + uint32_t sum = 0; + for (uint32_t j = 0; j < d; ++j) + for (uint32_t i = 0; i < d; ++i) + sum += src.uv[static_cast(r * d + j) * src.width + (c * d + i) * 2 + comp]; + out.uv[static_cast(r) * out.width + c * 2 + comp] = static_cast((sum + n / 2) / n); + } + return out; +} + +bool cuda_ok(cudaError_t e, const char* what) +{ + if (e != cudaSuccess) + { + std::fprintf(stderr, "FAIL %s: %s\n", what, cudaGetErrorString(e)); + ++g_failures; + return false; + } + return true; +} + +struct DevicePlane +{ + uint8_t* ptr = nullptr; + size_t pitch = 0; +}; + +DevicePlane alloc_plane(uint32_t row_bytes, uint32_t rows) +{ + DevicePlane p; + // Pad the row so the pitch is always larger than the data, whatever alignment the allocator picks. + cuda_ok(cudaMallocPitch(reinterpret_cast(&p.ptr), &p.pitch, row_bytes + 64, rows), "cudaMallocPitch"); + cuda_ok(cudaMemset2D(p.ptr, p.pitch, 0xEE, p.pitch, rows), "cudaMemset2D"); + return p; +} + +void run_case(const HostNv12& src, uint32_t divisor, const char* label) +{ + const uint32_t dw = src.width / divisor; + const uint32_t dh = src.height / divisor; + DevicePlane sy = alloc_plane(src.width, src.height); + DevicePlane suv = alloc_plane(src.width, src.height / 2); + DevicePlane dy = alloc_plane(dw, dh); + DevicePlane duv = alloc_plane(dw, dh / 2); + CHECK(sy.pitch > src.width && dy.pitch > dw); + + cuda_ok( + cudaMemcpy2D(sy.ptr, sy.pitch, src.y.data(), src.width, src.width, src.height, cudaMemcpyHostToDevice), "up y"); + cuda_ok(cudaMemcpy2D(suv.ptr, suv.pitch, src.uv.data(), src.width, src.width, src.height / 2, cudaMemcpyHostToDevice), + "up uv"); + + const Nv12Planes in{ sy.ptr, suv.ptr, sy.pitch, suv.pitch }; + const Nv12PlanesOut out{ dy.ptr, duv.ptr, dy.pitch, duv.pitch }; + CHECK(box_downscale_nv12(in, src.width, src.height, out, divisor, nullptr) == cudaSuccess); + cuda_ok(cudaDeviceSynchronize(), "sync"); + + std::vector got_y(static_cast(dw) * dh), got_uv(static_cast(dw) * (dh / 2)); + cuda_ok(cudaMemcpy2D(got_y.data(), dw, dy.ptr, dy.pitch, dw, dh, cudaMemcpyDeviceToHost), "down y"); + cuda_ok(cudaMemcpy2D(got_uv.data(), dw, duv.ptr, duv.pitch, dw, dh / 2, cudaMemcpyDeviceToHost), "down uv"); + + const HostNv12 ref = reference_downscale(src, divisor); + const bool y_match = got_y == ref.y; + const bool uv_match = got_uv == ref.uv; + if (!y_match || !uv_match) + std::fprintf(stderr, "case %s divisor %u: y_match=%d uv_match=%d\n", label, divisor, y_match, uv_match); + CHECK(y_match); + CHECK(uv_match); + + // The padding bytes past each destination row must be left alone. + std::vector raw(dy.pitch * dh); + cuda_ok(cudaMemcpy(raw.data(), dy.ptr, raw.size(), cudaMemcpyDeviceToHost), "raw"); + bool pad_intact = true; + for (uint32_t r = 0; r < dh; ++r) + for (size_t c = dw; c < dy.pitch; ++c) + pad_intact = pad_intact && raw[r * dy.pitch + c] == 0xEE; + CHECK(pad_intact); + + cudaFree(sy.ptr); + cudaFree(suv.ptr); + cudaFree(dy.ptr); + cudaFree(duv.ptr); +} + +void test_supported() +{ + CHECK(box_downscale_supported(640, 480, 1)); + CHECK(box_downscale_supported(640, 480, 2)); + CHECK(box_downscale_supported(640, 480, 4)); + CHECK(box_downscale_supported(1280, 720, 4)); + CHECK(box_downscale_supported(64, 64, 8)); + CHECK(!box_downscale_supported(640, 480, 0)); + CHECK(!box_downscale_supported(640, 480, 9)); + CHECK(!box_downscale_supported(640, 480, 16)); + CHECK(!box_downscale_supported(642, 480, 4)); // not divisible + CHECK(!box_downscale_supported(640, 482, 4)); + CHECK(!box_downscale_supported(644, 480, 4)); // divisible, odd result width (161) + CHECK(!box_downscale_supported(640, 484, 4)); // odd result height (121) + CHECK(!box_downscale_supported(641, 480, 1)); // odd source + CHECK(!box_downscale_supported(0, 480, 1)); + CHECK(!box_downscale_supported(640, 0, 1)); +} + +void test_invalid_launch() +{ + DevicePlane y = alloc_plane(64, 64); + DevicePlane uv = alloc_plane(64, 32); + const Nv12Planes in{ y.ptr, uv.ptr, y.pitch, uv.pitch }; + const Nv12PlanesOut out{ y.ptr, uv.ptr, y.pitch, uv.pitch }; + CHECK(box_downscale_nv12(in, 64, 64, out, 0, nullptr) == cudaErrorInvalidValue); + CHECK(box_downscale_nv12(in, 64, 64, out, 9, nullptr) == cudaErrorInvalidValue); + CHECK(box_downscale_nv12(in, 66, 64, out, 4, nullptr) == cudaErrorInvalidValue); + CHECK(box_downscale_nv12(in, 68, 64, out, 4, nullptr) == cudaErrorInvalidValue); + cudaFree(y.ptr); + cudaFree(uv.ptr); +} + +} // namespace + +int main() +{ + test_supported(); + + int count = 0; + if (cudaGetDeviceCount(&count) != cudaSuccess || count == 0) + { + std::fprintf(stderr, "SKIP: no CUDA device; only the CPU-side checks ran\n"); + return g_failures == 0 ? 77 : 1; + } + + test_invalid_launch(); + // 128x96 divides by 1, 2 and 4 into even sizes; 64x32 also exercises divisor 8 to the minimum even size. + for (uint32_t d : { 1u, 2u, 4u }) + { + run_case(make_pattern(128, 96, false, d), d, "gradient"); + run_case(make_pattern(128, 96, true, d), d, "blocky"); + } + run_case(make_pattern(64, 32, false, 8), 8, "gradient"); + + if (g_failures != 0) + { + std::fprintf(stderr, "%d check(s) failed\n", g_failures); + return 1; + } + std::printf("gpu_ops tests passed\n"); + return 0; +} diff --git a/examples/camera_livekit_bridge/tests/nvenc/CMakeLists.txt b/examples/camera_livekit_bridge/tests/nvenc/CMakeLists.txt new file mode 100644 index 0000000000..771ae60aac --- /dev/null +++ b/examples/camera_livekit_bridge/tests/nvenc/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 and libnvidia-encode; the test exits 77 (skipped) without them. +add_executable(test_nvenc test_nvenc_encoder.cpp) +target_link_libraries(test_nvenc PRIVATE camera_livekit_bridge::nvenc) +add_test(NAME nvenc COMMAND test_nvenc) +set_tests_properties(nvenc PROPERTIES SKIP_RETURN_CODE 77) diff --git a/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp b/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp new file mode 100644 index 0000000000..6f03097785 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp @@ -0,0 +1,376 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +using namespace camera_bridge; + +namespace +{ + +int g_failures = 0; + +#define CHECK(cond) \ + do \ + { \ + if (!(cond)) \ + { \ + std::fprintf(stderr, "FAIL %s:%d: %s\n", __FILE__, __LINE__, #cond); \ + ++g_failures; \ + } \ + } while (0) + +// NAL unit types found after Annex-B start codes (00 00 01). +std::vector nal_types(const std::vector& d) +{ + std::vector types; + for (size_t i = 0; i + 3 < d.size(); ++i) + if (d[i] == 0 && d[i + 1] == 0 && d[i + 2] == 1) + types.push_back(d[i + 3] & 0x1F); + return types; +} + +bool has_type(const std::vector& t, int type) +{ + for (int v : t) + if (v == type) + return true; + return false; +} + +bool starts_with_start_code(const std::vector& d) +{ + return (d.size() >= 3 && d[0] == 0 && d[1] == 0 && d[2] == 1) || + (d.size() >= 4 && d[0] == 0 && d[1] == 0 && d[2] == 0 && d[3] == 1); +} + +void fill_surface(const NvencEncoder::Surface& s, uint8_t luma) +{ + const uint32_t w = s.view.geometry.width; + const uint32_t h = s.view.geometry.height; + CHECK(cudaMemset2D(s.view.y, s.view.y_pitch, luma, w, h) == cudaSuccess); + CHECK(cudaMemset2D(s.view.uv, s.view.uv_pitch, 128, w, h / 2) == cudaSuccess); + CHECK(cudaDeviceSynchronize() == cudaSuccess); +} + +void test_invalid_geometry(CudaContext& ctx) +{ + for (Geometry g : { Geometry{ 641, 480 }, Geometry{ 640, 479 }, Geometry{ 0, 480 }, Geometry{ 640, 0 } }) + { + EncoderConfig cfg; + cfg.geometry = g; + bool threw = false; + try + { + NvencEncoder bad(ctx, cfg); + } + catch (const std::invalid_argument&) + { + threw = true; + } + catch (const std::exception& e) + { + std::fprintf(stderr, "unexpected exception type for %ux%u: %s\n", g.width, g.height, e.what()); + } + CHECK(threw); + } +} + +void test_pool(NvencEncoder& enc, uint32_t pool_size) +{ + // The pool hands out exactly input_surfaces distinct slots, then reports false until one is released. + std::vector held(pool_size); + for (uint32_t i = 0; i < pool_size; ++i) + { + CHECK(enc.acquire(held[i])); + for (uint32_t j = 0; j < i; ++j) + CHECK(held[i].slot != held[j].slot); + } + NvencEncoder::Surface extra; + CHECK(!enc.acquire(extra)); + enc.release(held[0]); + enc.release(held[0]); // idempotent + CHECK(enc.acquire(extra)); + CHECK(extra.slot == held[0].slot); + enc.release(extra); + for (uint32_t i = 1; i < pool_size; ++i) + enc.release(held[i]); +} + +bool ffmpeg_available() +{ + return std::system("command -v ffmpeg >/dev/null 2>&1") == 0; +} + +void test_decodable(const std::vector& idr) +{ + if (!ffmpeg_available()) + return; + const char* tmp = std::getenv("TMPDIR"); + const std::string path = std::string(tmp ? tmp : "/tmp") + "/test_nvenc_idr_" + std::to_string(::getpid()) + ".h264"; + { + std::ofstream f(path, std::ios::binary); + f.write(reinterpret_cast(idr.data()), static_cast(idr.size())); + } + const int rc = std::system(("ffmpeg -v error -i " + path + " -f null - luma(static_cast(w) * h); + uint32_t x = seed * 2654435761u + 1u; + for (uint8_t& v : luma) + { + x ^= x << 13; + x ^= x >> 17; + x ^= x << 5; + v = static_cast(x >> 24); + } + CHECK(cudaMemcpy2D(s.view.y, s.view.y_pitch, luma.data(), w, w, h, cudaMemcpyHostToDevice) == cudaSuccess); + CHECK(cudaMemset2D(s.view.uv, s.view.uv_pitch, 128, w, h / 2) == cudaSuccess); + CHECK(cudaDeviceSynchronize() == cudaSuccess); +} + +EncodedPacket encode_noise(NvencEncoder& enc, uint32_t seed, uint64_t key, bool force_idr) +{ + NvencEncoder::Surface s; + CHECK(enc.acquire(s)); + fill_noise(s, seed); + return enc.encode(s, nullptr, key, force_idr); +} + +EncodedPacket encode_flat(NvencEncoder& enc, uint64_t key, bool force_idr) +{ + NvencEncoder::Surface s; + CHECK(enc.acquire(s)); + fill_surface(s, 100); + return enc.encode(s, nullptr, key, force_idr); +} + +void test_bitrate_getter(NvencEncoder& enc, uint32_t configured) +{ + CHECK(enc.bitrate() == configured); + enc.set_bitrate(1'000'000); + CHECK(enc.bitrate() == 1'000'000); + enc.set_bitrate(6'000'000); + CHECK(enc.bitrate() == 6'000'000); +} + +void test_zero_bitrate_rejected(NvencEncoder& enc) +{ + const uint32_t before = enc.bitrate(); + bool threw = false; + try + { + enc.set_bitrate(0); + } + catch (const std::invalid_argument&) + { + threw = true; + } + CHECK(threw); + CHECK(enc.bitrate() == before); +} + +void test_output_after_reconfigure(NvencEncoder& enc) +{ + // Start from an IDR so the following pictures are P-frames. + CHECK(encode_flat(enc, 500, true).keyframe); + enc.set_bitrate(2'000'000); + for (int i = 0; i < 4; ++i) + { + const uint64_t key = 501 + static_cast(i); + EncodedPacket p = encode_flat(enc, key, false); + CHECK(!p.data.empty()); + CHECK(starts_with_start_code(p.data)); + CHECK(p.key == key); + CHECK(!p.keyframe); + const std::vector types = nal_types(p.data); + CHECK(has_type(types, 1)); + CHECK(!has_type(types, 5)); + } + EncodedPacket idr = encode_flat(enc, 600, true); + const std::vector types = nal_types(idr.data); + CHECK(idr.keyframe); + CHECK(idr.key == 600); + CHECK(has_type(types, 7) && has_type(types, 8) && has_type(types, 5)); +} + +double mean_noise_size(NvencEncoder& enc, int frames, int skip, uint32_t& seed) +{ + double total = 0; + int counted = 0; + for (int i = 0; i < frames; ++i) + { + EncodedPacket p = encode_noise(enc, seed++, 700 + static_cast(seed), false); + CHECK(!p.data.empty()); + if (i >= skip) + { + total += static_cast(p.data.size()); + ++counted; + } + } + return counted ? total / counted : 0.0; +} + +void test_rate_changes_output_size(NvencEncoder& enc) +{ + uint32_t seed = 1; + CHECK(encode_noise(enc, seed++, 650, true).keyframe); + enc.set_bitrate(6'000'000); + const double high = mean_noise_size(enc, 30, 5, seed); + enc.set_bitrate(500'000); + const double low = mean_noise_size(enc, 30, 5, seed); + std::printf("noise mean packet size: high=%.0f B low=%.0f B ratio=%.3f\n", high, low, high > 0 ? low / high : 0.0); + CHECK(high > 0); + CHECK(low < 0.3 * high); +} + +void test_decodable_across_reconfigure(NvencEncoder& enc) +{ + if (!ffmpeg_available()) + return; + std::vector stream; + auto append = [&stream](const EncodedPacket& p) { stream.insert(stream.end(), p.data.begin(), p.data.end()); }; + uint32_t seed = 5000; + enc.set_bitrate(4'000'000); + append(encode_noise(enc, seed++, 800, true)); + for (int i = 0; i < 5; ++i) + append(encode_noise(enc, seed++, 801 + static_cast(i), false)); + enc.set_bitrate(800'000); + for (int i = 0; i < 6; ++i) + append(encode_noise(enc, seed++, 810 + static_cast(i), false)); + enc.set_bitrate(3'000'000); + for (int i = 0; i < 6; ++i) + append(encode_noise(enc, seed++, 820 + static_cast(i), false)); + + const char* tmp = std::getenv("TMPDIR"); + const std::string path = + std::string(tmp ? tmp : "/tmp") + "/test_nvenc_reconfig_" + std::to_string(::getpid()) + ".h264"; + { + std::ofstream f(path, std::ios::binary); + f.write(reinterpret_cast(stream.data()), static_cast(stream.size())); + } + const int rc = std::system(("ffmpeg -v error -i " + path + " -f null - ctx; + std::unique_ptr enc; + try + { + ctx = std::make_unique(0); + enc = std::make_unique(*ctx, cfg); + } + catch (const std::exception& e) + { + std::fprintf(stderr, "SKIP: cannot create an NVENC encoder here: %s\n", e.what()); + return 77; + } + + test_invalid_geometry(*ctx); + test_pool(*enc, kPool); + test_bitrate_getter(*enc, cfg.bitrate_bps); + test_zero_bitrate_rejected(*enc); + + std::vector first_idr; + constexpr int kFrames = 8; + for (int i = 0; i < kFrames; ++i) + { + NvencEncoder::Surface s; + CHECK(enc->acquire(s)); + CHECK(s.view.geometry == cfg.geometry); + CHECK(s.view.y_pitch >= 640); + fill_surface(s, 100); + const uint64_t key = 1000 + static_cast(i); + EncodedPacket p = enc->encode(s, nullptr, key, i == 0); + + CHECK(!p.data.empty()); + CHECK(starts_with_start_code(p.data)); + CHECK(p.key == key); + const std::vector types = nal_types(p.data); + if (i == 0) + { + CHECK(p.keyframe); + CHECK(has_type(types, 7)); + CHECK(has_type(types, 8)); + CHECK(has_type(types, 5)); + first_idr = p.data; + } + else + { + CHECK(!p.keyframe); + CHECK(!has_type(types, 5)); + CHECK(has_type(types, 1)); + } + } + + // encode() released its surface: the whole pool is free again. + { + std::vector all(kPool); + for (auto& s : all) + CHECK(enc->acquire(s)); + for (auto& s : all) + enc->release(s); + } + + // A later forced IDR carries parameter sets again. + { + NvencEncoder::Surface s; + CHECK(enc->acquire(s)); + fill_surface(s, 100); + EncodedPacket p = enc->encode(s, nullptr, 77, true); + const std::vector types = nal_types(p.data); + CHECK(p.keyframe); + CHECK(p.key == 77); + CHECK(has_type(types, 7) && has_type(types, 8) && has_type(types, 5)); + } + + test_decodable(first_idr); + test_output_after_reconfigure(*enc); + test_rate_changes_output_size(*enc); + test_decodable_across_reconfigure(*enc); + + if (g_failures != 0) + { + std::fprintf(stderr, "%d check(s) failed\n", g_failures); + return 1; + } + std::printf("nvenc tests passed\n"); + return 0; +} diff --git a/examples/camera_livekit_bridge/tests/python/test_bridge_pipeline.py b/examples/camera_livekit_bridge/tests/python/test_bridge_pipeline.py new file mode 100644 index 0000000000..d605f905c5 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/python/test_bridge_pipeline.py @@ -0,0 +1,478 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""End to end: mock sender -> camera_receiver -> camera_livekit_bridge -> LiveKit -> subscriber. + +Run: python3 test_bridge_pipeline.py -v (about 40 s; needs a CUDA device). + +Environment overrides: LIVEKIT_SERVER_BIN, CAMERA_LIVEKIT_BRIDGE_BIN, +LIVEKIT_PYTHON (interpreter with the `livekit` SDK and numpy, used for the +subscriber). The test is skipped, never faked, when a prerequisite is missing. +""" + +from __future__ import annotations + +import base64 +import hashlib +import hmac +import json +import os +import shutil +import signal +import socket +import subprocess +import sys +import tempfile +import time +import unittest +import uuid +from pathlib import Path + +HERE = Path(__file__).resolve().parent +EXAMPLES = HERE.parents[2] +MOCK_DIR = EXAMPLES / "camera_streamer" / "mock_sender" +RECEIVER_DIR = EXAMPLES / "camera_streamer" / "camera_receiver" + +SERVER_BIN = Path( + os.environ.get( + "LIVEKIT_SERVER_BIN", + shutil.which("livekit-server") + or Path.home() / "github/livekit-server-1.13.10/livekit-server", + ) +) +BRIDGE_BIN = Path( + os.environ.get( + "CAMERA_LIVEKIT_BRIDGE_BIN", + EXAMPLES / "camera_livekit_bridge/build/bridge/camera_livekit_bridge", + ) +) +RECEIVER_PY = RECEIVER_DIR / ".venv" / "bin" / "python" +SUBSCRIBER_PY = Path(os.environ.get("LIVEKIT_PYTHON", sys.executable)) + +API_KEY, API_SECRET = "devkey", "secret" +SOURCE_FPS = 30 +# name, source width, source height, track name, bridge divisor +STREAMS = [ + ("head", 640, 496, "head_track", 2), + ("wrist", 320, 240, "wrist_track", 1), +] +COLLECT_S = 8.0 +PUBLISHER_IDENTITY = "bridge-pub" + +# Runs under LIVEKIT_PYTHON; prints one JSON document with per-track samples. +SUBSCRIBER_SCRIPT = r""" +import asyncio, json, sys, time +from livekit import rtc + +url, token_file, n_tracks, collect_s = sys.argv[1], sys.argv[2], int(sys.argv[3]), float(sys.argv[4]) +token = open(token_file).read().strip() + + +async def main(): + room = rtc.Room() + tracks = {} + tasks = [] + + async def read(name, identity, track): + t = tracks.setdefault(name, {"identity": identity, "frames": []}) + first = None + async for event in rtc.VideoStream(track): + now = time.time() + first = first or now + if now - first > collect_s: + break + m = event.metadata + t["frames"].append({ + "t": now, + "w": event.frame.width, + "h": event.frame.height, + "id": m.frame_id if m is not None and m.HasField("frame_id") else None, + "ts": m.user_timestamp if m is not None and m.HasField("user_timestamp") else None, + "ud": bytes(m.user_data).hex() if m is not None and m.HasField("user_data") else None, + }) + + @room.on("track_subscribed") + def on_track(track, publication, participant): + if track.kind == rtc.TrackKind.KIND_VIDEO: + tasks.append(asyncio.create_task(read(publication.name, participant.identity, track))) + + await room.connect(url, token) + deadline = time.time() + 25 + collect_s + while time.time() < deadline: + await asyncio.sleep(0.25) + if len(tasks) >= n_tracks and all(t.done() for t in tasks): + break + keyframe_api = any(hasattr(rtc.VideoStream, n) for n in ("request_keyframe", "send_pli")) + await room.disconnect() + print("RESULT " + json.dumps({"tracks": tracks, "keyframe_api": keyframe_api})) + + +asyncio.run(main()) +""" + + +def _b64(data: bytes) -> str: + return base64.urlsafe_b64encode(data).rstrip(b"=").decode() + + +def make_token(room: str, identity: str, publish: bool) -> str: + header = _b64(json.dumps({"alg": "HS256", "typ": "JWT"}).encode()) + now = int(time.time()) + claims = { + "iss": API_KEY, + "sub": identity, + "nbf": now, + "exp": now + 3600, + "name": identity, + "video": { + "room": room, + "roomJoin": True, + "canPublish": publish, + "canSubscribe": not publish, + "canPublishData": publish, + }, + } + body = _b64(json.dumps(claims).encode()) + sig = hmac.new(API_SECRET.encode(), f"{header}.{body}".encode(), hashlib.sha256) + return f"{header}.{body}.{_b64(sig.digest())}" + + +def free_port(kind: int) -> int: + with socket.socket(socket.AF_INET, kind) as s: + s.bind(("127.0.0.1", 0)) + return s.getsockname()[1] + + +def _python_has(python: Path, module: str) -> bool: + if not python.exists(): + return False + try: + return ( + subprocess.run( + [str(python), "-c", f"import {module}"], + capture_output=True, + check=False, + timeout=60, + ).returncode + == 0 + ) + except (OSError, subprocess.TimeoutExpired): + return False + + +def _cuda_present() -> bool: + smi = shutil.which("nvidia-smi") + if smi is None: + return False + try: + out = subprocess.run( + [smi, "-L"], capture_output=True, text=True, timeout=30, check=False + ) + except (OSError, subprocess.TimeoutExpired): + return False + return out.returncode == 0 and "GPU" in out.stdout + + +def _skip_reason() -> str | None: + if not SERVER_BIN.is_file(): + return f"livekit-server not found ({SERVER_BIN}); set LIVEKIT_SERVER_BIN" + if not BRIDGE_BIN.is_file(): + return f"camera_livekit_bridge not built ({BRIDGE_BIN}); set CAMERA_LIVEKIT_BRIDGE_BIN" + if shutil.which("ffmpeg") is None: + return "ffmpeg is required by the mock sender" + if not RECEIVER_PY.is_file(): + return f"receiver venv missing ({RECEIVER_PY})" + if not _cuda_present(): + return "no CUDA device (nvidia-smi -L lists no GPU)" + if not _python_has(SUBSCRIBER_PY, "livekit.rtc") or not _python_has( + SUBSCRIBER_PY, "numpy" + ): + return f"{SUBSCRIBER_PY} lacks the livekit SDK or numpy; set LIVEKIT_PYTHON" + return None + + +class _Proc: + """A child process with its output captured to a file and scanned for markers.""" + + def __init__(self, name: str, argv: list[str], logdir: Path, cwd=None): + self.name = name + self.log_path = logdir / f"{name}.log" + self._log = open(self.log_path, "wb") + self.proc = subprocess.Popen( + argv, + stdout=self._log, + stderr=subprocess.STDOUT, + cwd=cwd, + start_new_session=True, + ) + + def text(self) -> str: + self._log.flush() + return self.log_path.read_text(errors="replace") + + def wait_for(self, marker: str, timeout: float) -> None: + end = time.monotonic() + timeout + while time.monotonic() < end: + if marker in self.text(): + return + if self.proc.poll() is not None: + raise AssertionError( + f"{self.name} exited ({self.proc.returncode}) before {marker!r}:\n" + f"{self.text()[-3000:]}" + ) + time.sleep(0.2) + raise AssertionError( + f"{self.name}: no {marker!r} within {timeout}s:\n{self.text()[-3000:]}" + ) + + def stop(self) -> None: + if self.proc.poll() is None: + self.proc.send_signal(signal.SIGTERM) + try: + self.proc.wait(timeout=10) + except subprocess.TimeoutExpired: + self.proc.kill() + self.proc.wait(timeout=10) + self._log.close() + + +@unittest.skipIf(_skip_reason() is not None, _skip_reason() or "") +class BridgePipelineTest(unittest.TestCase): + result: dict + + @classmethod + def setUpClass(cls): + cls._procs: list[_Proc] = [] + cls._tmp = Path(tempfile.mkdtemp(prefix="bridge_pipeline_")) + try: + cls.result = cls._run_pipeline() + except BaseException: + cls._stop_all() + raise + cls._stop_all() + + @classmethod + def _stop_all(cls): + for proc in reversed(cls._procs): + proc.stop() + cls._procs.clear() + + @classmethod + def _spawn(cls, name, argv, cwd=None) -> _Proc: + proc = _Proc(name, argv, cls._tmp, cwd) + cls._procs.append(proc) + return proc + + @classmethod + def _run_pipeline(cls) -> dict: + tmp = cls._tmp + room = f"bridge-pipeline-{uuid.uuid4().hex[:10]}" + lk_port, lk_tcp, lk_udp = ( + free_port(socket.SOCK_STREAM), + free_port(socket.SOCK_STREAM), + free_port(socket.SOCK_DGRAM), + ) + mock_port = free_port(socket.SOCK_STREAM) + sock = tmp / "receiver.sock" + url = f"ws://127.0.0.1:{lk_port}" + + server = cls._spawn( + "livekit", + [ + str(SERVER_BIN), + "--dev", + "--bind", + "127.0.0.1", + "--node-ip", + "127.0.0.1", + "--port", + str(lk_port), + "--rtc.tcp_port", + str(lk_tcp), + "--udp-port", + str(lk_udp), + ], + ) + server.wait_for("starting LiveKit server", 20) + deadline = time.monotonic() + 20 + while time.monotonic() < deadline: + try: + socket.create_connection(("127.0.0.1", lk_port), timeout=1).close() + break + except OSError: + time.sleep(0.2) + else: + raise AssertionError("livekit-server never accepted connections") + + mock_cmd = [ + sys.executable, + str(MOCK_DIR / "mock_sender.py"), + "--port", + str(mock_port), + "--duration", + "120", + "--cache-dir", + str(Path(tempfile.gettempdir()) / "camera_livekit_bridge_test_cache"), + ] + for name, w, h, _, _ in STREAMS: + mock_cmd += ["--stream", f"{name}:{w}x{h}@{SOURCE_FPS}"] + mock = cls._spawn("mock_sender", mock_cmd) + mock.wait_for("mock sender listening", 120) + + cfg = tmp / "receiver.yaml" + lines = [ + "version: 1", + "receiver:", + " host: 127.0.0.1", + f" port: {mock_port}", + " gpu_id: 0", + " stats_log_interval_s: 5", + "ipc:", + f" socket_path: {sock}", + f" owner_name: bridge_pipeline_{os.getpid()}", + " queue_size: 1024", + "streams:", + ] + for i, (name, w, h, _, _) in enumerate(STREAMS): + lines += [ + f" - index: {i}", + f" name: {name}", + " type: mono", + f" width: {w}", + f" height: {h}", + " output_format: nv12", + " ring_depth: 6", + " ipc:", + f" mono: {name}", + ] + cfg.write_text("\n".join(lines) + "\n") + receiver = cls._spawn( + "camera_receiver", + [ + str(RECEIVER_PY), + str(RECEIVER_DIR / "camera_receiver.py"), + "--config", + str(cfg), + "--port", + str(mock_port), + "--socket", + str(sock), + ], + cwd=RECEIVER_DIR, + ) + receiver.wait_for("camera_receiver running", 60) + receiver.wait_for("connected to AVS1 sender", 30) + + pub_token = tmp / "pub.token" + pub_token.write_text(make_token(room, PUBLISHER_IDENTITY, True)) + bridge_cmd = [ + str(BRIDGE_BIN), + "--livekit-url", + url, + "--livekit-token-file", + str(pub_token), + "--socket", + str(sock), + "--fps", + str(SOURCE_FPS), + ] + for name, _, _, track, div in STREAMS: + bridge_cmd += ["--stream", f"{name}={track}:{div}"] + bridge = cls._spawn("bridge", bridge_cmd) + bridge.wait_for("publishing", 60) + + sub_token = tmp / "sub.token" + sub_token.write_text(make_token(room, "test-subscriber", False)) + script = tmp / "subscriber.py" + script.write_text(SUBSCRIBER_SCRIPT) + try: + out = subprocess.run( + [ + str(SUBSCRIBER_PY), + str(script), + url, + str(sub_token), + str(len(STREAMS)), + str(COLLECT_S), + ], + capture_output=True, + check=False, + text=True, + timeout=COLLECT_S + 60, + ) + except subprocess.TimeoutExpired as exc: + raise AssertionError(f"subscriber timed out: {exc.stdout}") from exc + marker = [ln for ln in out.stdout.splitlines() if ln.startswith("RESULT ")] + if not marker: + raise AssertionError( + f"subscriber produced no result (rc={out.returncode}):\n" + f"{out.stdout[-2000:]}\n{out.stderr[-2000:]}\nbridge:\n{bridge.text()[-2000:]}" + ) + return json.loads(marker[-1][len("RESULT ") :]) + + def _frames(self, track: str) -> list[dict]: + tracks = self.result["tracks"] + self.assertIn(track, tracks, f"track {track} missing; got {sorted(tracks)}") + return tracks[track]["frames"] + + def test_tracks_and_sizes(self): + self.assertEqual(sorted(self.result["tracks"]), sorted(s[3] for s in STREAMS)) + for _name, w, h, track, div in STREAMS: + info = self.result["tracks"][track] + self.assertEqual(info["identity"], PUBLISHER_IDENTITY) + frames = self._frames(track) + self.assertTrue(frames, f"{track} delivered no frames") + sizes = {(f["w"], f["h"]) for f in frames} + # The first frames can be pre-adaptation; the steady state must match. + steady = {(f["w"], f["h"]) for f in frames[len(frames) // 4 :]} + self.assertEqual(steady, {(w // div, h // div)}, f"{track}: {sizes}") + + def test_frame_ids_and_rate(self): + for _, _, _, track, _ in STREAMS: + frames = self._frames(track) + ids = [f["id"] for f in frames] + self.assertNotIn(None, ids, f"{track}: frame without frame_id") + steps = [b - a for a, b in zip(ids, ids[1:])] + self.assertTrue(all(s > 0 for s in steps), f"{track}: ids not increasing") + consecutive = sum(1 for s in steps if s == 1) + self.assertGreaterEqual( + consecutive / max(len(steps), 1), 0.8, f"{track}: steps {steps[:60]}" + ) + span = frames[-1]["t"] - frames[0]["t"] + rate = (len(frames) - 1) / span + self.assertTrue( + 0.7 * SOURCE_FPS <= rate <= 1.15 * SOURCE_FPS, + f"{track}: {rate:.1f} fps over {span:.1f}s", + ) + + def test_user_timestamp_is_capture_time(self): + for _, _, _, track, _ in STREAMS: + frames = self._frames(track) + t0 = frames[0]["t"] + late = [f for f in frames if f["t"] - t0 >= 1.0] + self.assertGreater(len(late), 30, f"{track}: too few frames after 1 s") + for f in late: + self.assertIsNotNone(f["ts"], f"{track}: frame without user_timestamp") + age_ms = (f["t"] * 1e6 - f["ts"]) / 1000.0 + self.assertTrue( + 0 <= age_ms <= 1500, + f"{track}: frame {f['id']} age {age_ms:.1f} ms", + ) + + def test_user_data_is_sensor_counter(self): + for _, _, _, track, _ in STREAMS: + for f in self._frames(track): + self.assertIsNotNone(f["ud"], f"{track}: frame without user_data") + raw = bytes.fromhex(f["ud"]) + self.assertEqual(len(raw), 8, f"{track}: user_data {raw!r}") + self.assertNotEqual(int.from_bytes(raw, "little"), 0) + + def test_subscriber_keyframe_request(self): + if not self.result["keyframe_api"]: + self.skipTest( + "the livekit Python SDK exposes no subscriber keyframe request" + ) + + +if __name__ == "__main__": + unittest.main() diff --git a/examples/camera_livekit_bridge/tests/python/test_bridge_timing.py b/examples/camera_livekit_bridge/tests/python/test_bridge_timing.py new file mode 100644 index 0000000000..d74d8e8fa8 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/python/test_bridge_timing.py @@ -0,0 +1,472 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Bridge timing against egocap's own clock-sync and recording code. + +Mock sender -> camera_receiver -> camera_livekit_bridge -> LiveKit -> egocap's +ClockSyncClient / LiveKitRecordingFeed (run in a subprocess under egocap's +interpreter). The bridge is restarted per `--timestamp-basis`; everything +else is shared. + +Run: python3 test_bridge_timing.py -v (about 40 s; needs a CUDA device). + +Environment overrides: LIVEKIT_SERVER_BIN, CAMERA_LIVEKIT_BRIDGE_BIN, +EGOCAP_DIR (an egocap checkout with the timing code, the directory holding +backend/ and capture_server/), EGOCAP_PYTHON (interpreter with egocap's +dependencies and the livekit SDK). Skipped, never faked, when one is missing. +""" + +from __future__ import annotations + +import hashlib +import hmac +import json +import os +import shutil +import socket +import statistics +import subprocess +import sys +import tempfile +import time +import unittest +import uuid +from pathlib import Path + +from test_bridge_pipeline import ( + BRIDGE_BIN, + MOCK_DIR, + RECEIVER_DIR, + RECEIVER_PY, + SERVER_BIN, + SOURCE_FPS, + _cuda_present, + _Proc, + free_port, + API_KEY, + API_SECRET, + _b64, + make_token, +) + +# name, source width, source height, track name, bridge divisor +STREAMS = [ + ("head", 640, 496, "head_track", 2), + ("wrist", 320, 240, "wrist_track", 1), +] +BASES = ("monotonic_us",) +COLLECT_S = 10.0 +MAX_AGE_MS = 1500.0 +OFFSET_TOLERANCE_NS = 2_000_000 + +# Defaults: the egocap example in this repository and the current interpreter; override with the environment. +EGOCAP_DIR = Path( + os.environ.get("EGOCAP_DIR", Path(__file__).resolve().parents[3] / "egocap") +) +EGOCAP_PY = Path(os.environ.get("EGOCAP_PYTHON", sys.executable)) + +# Runs under EGOCAP_PY with PYTHONPATH at egocap's packages; prints one RESULT line. +DRIVER_SCRIPT = r""" +import asyncio, json, sys, time +from pathlib import Path +from types import SimpleNamespace +from capture_server.livekit_streams import load_livekit_streams +from egocap_web.livekit_receiver import LiveKitFrameReceiver +from egocap_web.livekit_recording import ( + LiveKitRecordingFeed, unrecordable_reason, uncertainty_budget_ns) + +yaml_path, token_file, seconds = sys.argv[1], sys.argv[2], float(sys.argv[3]) + + +class Ingress: + # Keeps timestamps only: holding decoded frames would pin the SDK's buffers. + def __init__(self): + self.samples = {} + + def offer(self, sample): + self.samples.setdefault(sample.source, []).append( + (sample.capture_time, sample.receive_time)) + return True + + +async def main(): + streams = load_livekit_streams(Path(yaml_path)) + reason = unrecordable_reason(SimpleNamespace(streams=streams), "http://127.0.0.1:8765") + receiver = LiveKitFrameReceiver( + streams, open(token_file).read().strip(), clock_sync=True) + await receiver.start() + client = receiver.clock_sync + exchanges = {"added": 0, "rejected_by_estimator": 0} + real_add = client.estimator.add + + def counting_add(exchange): + exchanges["added"] += 1 + why = real_add(exchange) + if why is not None: + exchanges["rejected_by_estimator"] += 1 + return why + + client.estimator.add = counting_add + ingress = Ingress() + # As LiveKitRecordingSource.prepare() does: the feed starts only once a mapping exists. + t0 = time.monotonic() + while client.mapping() is None and time.monotonic() - t0 < 10: + await asyncio.sleep(0.05) + first_mapping_s = time.monotonic() - t0 + receiver.clear_queued() + feed = LiveKitRecordingFeed( + receiver, streams, ingress, {s.stream_id: s.name for s in streams}) + feed.start() + maps, end = [], time.monotonic() + seconds + while time.monotonic() < end and feed.failure is None: + await asyncio.sleep(0.5) + m = client.mapping() + maps.append(None if m is None else [m.epoch, m.offset_ns, m.rtt_ns, m.uncertainty_ns]) + out = { + "first_mapping_s": first_mapping_s, + "unrecordable_reason": reason, + "budget_ns": uncertainty_budget_ns(), + "feed_failure": feed.failure, + "loss": feed.loss(), + "exchanges": exchanges, + "client_rejected": client.rejected, + "discontinuity": client.discontinuity, + "epochs": sorted({h.epoch for h in client.estimator.history}), + "mappings": maps, + "ages_ms": {n: [(r - c) * 1000 for c, r in s] for n, s in ingress.samples.items()}, + } + await feed.stop() + await receiver.stop() + print("RESULT " + json.dumps(out)) + + +asyncio.run(main()) +""" + +# Runs under EGOCAP_PY: only the loader and recordability check, no network. +REFUSE_SCRIPT = r""" +import json, sys +from pathlib import Path +from types import SimpleNamespace +from capture_server.livekit_streams import load_livekit_streams +from egocap_web.livekit_recording import recordable_livekit_streams, unrecordable_reason + +streams = load_livekit_streams(Path(sys.argv[1])) +started = SimpleNamespace(streams=streams) +print("RESULT " + json.dumps({ + "reason": unrecordable_reason(started, "http://127.0.0.1:8765"), + "recordable": len(recordable_livekit_streams(started)), +})) +""" + + +def make_receiver_token(room: str, identity: str) -> str: + """Subscribe and publish-data: clock sync sends its requests as data messages.""" + header = _b64(json.dumps({"alg": "HS256", "typ": "JWT"}).encode()) + now = int(time.time()) + claims = { + "iss": API_KEY, "sub": identity, "nbf": now, "exp": now + 3600, "name": identity, + "video": { + "room": room, "roomJoin": True, "canPublish": False, + "canSubscribe": True, "canPublishData": True, + }, + } # fmt: skip + body = _b64(json.dumps(claims).encode()) + sig = hmac.new(API_SECRET.encode(), f"{header}.{body}".encode(), hashlib.sha256) + return f"{header}.{body}.{_b64(sig.digest())}" + + +def _egocap_env() -> dict[str, str]: + env = dict(os.environ) + parts = [ + "backend", + "capture_server", + "recording", + "../../src/plugins/manus_calibration/python", + ] + env["PYTHONPATH"] = os.pathsep.join(str(EGOCAP_DIR / p) for p in parts) + return env + + +def _skip_reason() -> str | None: + if not SERVER_BIN.is_file(): + return f"livekit-server not found ({SERVER_BIN}); set LIVEKIT_SERVER_BIN" + if not BRIDGE_BIN.is_file(): + return f"camera_livekit_bridge not built ({BRIDGE_BIN}); set CAMERA_LIVEKIT_BRIDGE_BIN" + if shutil.which("ffmpeg") is None: + return "ffmpeg is required by the mock sender" + if not RECEIVER_PY.is_file(): + return f"receiver venv missing ({RECEIVER_PY})" + if not _cuda_present(): + return "no CUDA device (nvidia-smi -L lists no GPU)" + if not (EGOCAP_DIR / "backend/egocap_web/livekit_clock_sync.py").is_file(): + return f"egocap timing checkout not found ({EGOCAP_DIR}); set EGOCAP_DIR" + if not EGOCAP_PY.is_file(): + return f"egocap interpreter missing ({EGOCAP_PY}); set EGOCAP_PYTHON" + probe = subprocess.run( + [str(EGOCAP_PY), "-c", "import livekit.rtc, egocap_web.livekit_recording"], + capture_output=True, + env=_egocap_env(), + check=False, + timeout=120, + ) + if probe.returncode != 0: + return ( + "egocap code or livekit SDK not importable under EGOCAP_PYTHON: " + + (probe.stderr.decode(errors="replace")[-300:]) + ) + return None + + +def _streams_yaml(url: str, room: str, publisher: str, basis: str) -> str: + lines = [ + "livekit:", + f" url: {url}", + f" room: {room}", + f" publisher: {publisher}", + f" timestamp_basis: {basis}", + " streams:", + ] + for name, w, h, track, div in STREAMS: + lines.append( + f" - {{name: {name}, track: {track}, width: {w // div}, height: {h // div}, fps: {SOURCE_FPS}}}" + ) + return "\n".join(lines) + "\n" + + +def _result(stdout: str, what: str, extra: str = "") -> dict: + marker = [ln for ln in stdout.splitlines() if ln.startswith("RESULT ")] + if not marker: + raise AssertionError(f"{what} produced no result:\n{stdout[-2000:]}\n{extra}") + return json.loads(marker[-1][len("RESULT ") :]) + + +def _summ(ages: list[float]) -> str: + return f"n={len(ages)} min={min(ages):.1f} median={statistics.median(ages):.1f} max={max(ages):.1f} ms" + + +def _report(r: dict) -> str: + maps = [m for m in r["mappings"] if m is not None] + last = maps[-1] if maps else None + ages = {n: _summ(a) for n, a in r["ages_ms"].items() if a} + return ( + f"first_mapping_s={r['first_mapping_s']:.2f} replies={r['exchanges']} mappings={len(maps)} last(epoch,offset_ns,rtt_ns,unc_ns)={last} " + f"discontinuity={r['discontinuity']} feed_failure={r['feed_failure']!r} ages={ages}" + ) + + +@unittest.skipIf(_skip_reason() is not None, _skip_reason() or "") +class BridgeTimingTest(unittest.TestCase): + results: dict[str, dict] + refuse: dict + + @classmethod + def setUpClass(cls): + cls._procs: list[_Proc] = [] + cls._tmp = Path(tempfile.mkdtemp(prefix="bridge_timing_")) + cls.results = {} + try: + cls._run_all() + except BaseException: + cls._stop_all() + raise + cls._stop_all() + + @classmethod + def _stop_all(cls): + for proc in reversed(cls._procs): + proc.stop() + cls._procs.clear() + + @classmethod + def _spawn(cls, name, argv, cwd=None) -> _Proc: + proc = _Proc(name, argv, cls._tmp, cwd) + cls._procs.append(proc) + return proc + + @classmethod + def _start_chain(cls, room: str) -> tuple[str, Path]: + tmp = cls._tmp + lk_port, lk_tcp, lk_udp = ( + free_port(socket.SOCK_STREAM), + free_port(socket.SOCK_STREAM), + free_port(socket.SOCK_DGRAM), + ) + mock_port = free_port(socket.SOCK_STREAM) + sock = tmp / "receiver.sock" + url = f"ws://127.0.0.1:{lk_port}" + server = cls._spawn( + "livekit", + [ + str(SERVER_BIN), "--dev", "--bind", "127.0.0.1", "--node-ip", "127.0.0.1", + "--port", str(lk_port), "--rtc.tcp_port", str(lk_tcp), "--udp-port", str(lk_udp), + ], + ) # fmt: skip + server.wait_for("starting LiveKit server", 20) + deadline = time.monotonic() + 20 + while time.monotonic() < deadline: + try: + socket.create_connection(("127.0.0.1", lk_port), timeout=1).close() + break + except OSError: + time.sleep(0.2) + else: + raise AssertionError("livekit-server never accepted connections") + + mock_cmd = [ + sys.executable, str(MOCK_DIR / "mock_sender.py"), + "--port", str(mock_port), "--duration", "400", + "--cache-dir", str(Path(tempfile.gettempdir()) / "camera_livekit_bridge_test_cache"), + ] # fmt: skip + for name, w, h, _, _ in STREAMS: + mock_cmd += ["--stream", f"{name}:{w}x{h}@{SOURCE_FPS}"] + mock = cls._spawn("mock_sender", mock_cmd) + mock.wait_for("mock sender listening", 120) + + lines = [ + "version: 1", "receiver:", " host: 127.0.0.1", f" port: {mock_port}", " gpu_id: 0", + " stats_log_interval_s: 5", "ipc:", f" socket_path: {sock}", + f" owner_name: bridge_timing_{os.getpid()}", " queue_size: 1024", "streams:", + ] # fmt: skip + for i, (name, w, h, _, _) in enumerate(STREAMS): + lines += [ + f" - index: {i}", f" name: {name}", " type: mono", f" width: {w}", + f" height: {h}", " output_format: nv12", " ring_depth: 6", " ipc:", + f" mono: {name}", + ] # fmt: skip + cfg = tmp / "receiver.yaml" + cfg.write_text("\n".join(lines) + "\n") + receiver = cls._spawn( + "camera_receiver", + [ + str(RECEIVER_PY), str(RECEIVER_DIR / "camera_receiver.py"), + "--config", str(cfg), "--port", str(mock_port), "--socket", str(sock), + ], + cwd=RECEIVER_DIR, + ) # fmt: skip + receiver.wait_for("camera_receiver running", 60) + receiver.wait_for("connected to AVS1 sender", 30) + return url, sock + + @classmethod + def _run_basis(cls, url: str, sock: Path, room: str, basis: str) -> dict: + tmp = cls._tmp + identity = f"bridge-{basis}" + pub_token = tmp / f"pub-{basis}.token" + pub_token.write_text(make_token(room, identity, True)) + cmd = [ + str(BRIDGE_BIN), "--livekit-url", url, "--livekit-token-file", str(pub_token), + "--socket", str(sock), "--fps", str(SOURCE_FPS), "--timestamp-basis", basis, + ] # fmt: skip + for name, _, _, track, div in STREAMS: + cmd += ["--stream", f"{name}={track}:{div}"] + n_before = len(cls._procs) + bridge = cls._spawn(f"bridge_{basis}", cmd) + try: + bridge.wait_for("publishing", 60) + yaml_path = tmp / f"streams-{basis}.yaml" + yaml_path.write_text(_streams_yaml(url, room, identity, basis)) + # Clock sync sends data messages, so the receiver needs publish-data as well as subscribe. + token = tmp / f"recv-{basis}.token" + token.write_text(make_receiver_token(room, f"timing-{basis}")) + out = subprocess.run( + [str(EGOCAP_PY), "-c", DRIVER_SCRIPT, str(yaml_path), str(token), str(COLLECT_S)], + capture_output=True, text=True, check=False, env=_egocap_env(), + timeout=COLLECT_S + 90, + ) # fmt: skip + result = _result( + out.stdout, + f"egocap driver ({basis})", + out.stderr[-2000:] + bridge.text()[-2000:], + ) + finally: + # One identity at a time: stop only this bridge before the next starts. + for proc in reversed(cls._procs[n_before:]): + proc.stop() + del cls._procs[n_before:] + return result + + @classmethod + def _run_all(cls): + room = f"bridge-timing-{uuid.uuid4().hex[:10]}" + url, sock = cls._start_chain(room) + for basis in BASES: + cls.results[basis] = cls._run_basis(url, sock, room, basis) + print(f"\n[{basis}] {_report(cls.results[basis])}", file=sys.stderr) + unix_yaml = cls._tmp / "streams-unix.yaml" + unix_yaml.write_text(_streams_yaml(url, room, "bridge-unix", "capture_unix_us")) + out = subprocess.run( + [str(EGOCAP_PY), "-c", REFUSE_SCRIPT, str(unix_yaml)], + capture_output=True, text=True, check=False, env=_egocap_env(), timeout=120, + ) # fmt: skip + cls.refuse = _result(out.stdout, "refuse check", out.stderr[-2000:]) + + def _check_feed(self, basis: str): + r = self.results[basis] + self.assertIsNone( + r["feed_failure"], f"{basis}: feed failed: {r['feed_failure']}" + ) + self.assertEqual(sorted(r["ages_ms"]), sorted(s[0] for s in STREAMS)) + for name, ages in r["ages_ms"].items(): + self.assertGreater( + len(ages), 30, f"{basis}/{name}: too few recorded frames" + ) + self.assertTrue( + all(0 <= a <= MAX_AGE_MS for a in ages), + f"{basis}/{name}: age out of range; {_summ(ages)}; " + f"over limit: {[round(a) for a in ages if a > MAX_AGE_MS][:10]}", + ) + + def _check_clock(self, basis: str): + r = self.results[basis] + self.assertGreaterEqual( + r["exchanges"]["added"], 5, f"{basis}: few replies: {r['exchanges']}" + ) + maps = [m for m in r["mappings"] if m is not None] + self.assertTrue(maps, f"{basis}: estimator never produced a mapping") + self.assertFalse(r["discontinuity"], f"{basis}: clock discontinuity") + self.assertEqual(r["epochs"], [1], f"{basis}: epochs {r['epochs']}") + budget = r["budget_ns"] + for epoch, offset, rtt, unc in maps: + self.assertEqual(epoch, 1) + self.assertLess( + unc, budget, f"{basis}: uncertainty {unc} ns over budget {budget} ns" + ) + self.assertLess( + abs(offset), + OFFSET_TOLERANCE_NS, + f"{basis}: offset {offset} ns (rtt {rtt} ns)", + ) + + def test_monotonic_us_clock_and_feed(self): + self._check_clock("monotonic_us") + self._check_feed("monotonic_us") + + def test_capture_monotonic_us_is_refused_by_bridge(self): + # Needs a measured sender-to-PC clock mapping the bridge does not have, so it must fail at startup. + token = self._tmp / "refuse.token" + token.write_text(make_token("unused", "bridge-refuse", True)) + cmd = [ + str(BRIDGE_BIN), "--livekit-url", "ws://127.0.0.1:1", "--livekit-token-file", str(token), + "--socket", str(self._tmp / "none.sock"), "--stream", "head=head_track", + "--timestamp-basis", "capture_monotonic_us", + ] # fmt: skip + started = time.monotonic() + out = subprocess.run( + cmd, capture_output=True, text=True, check=False, timeout=30 + ) + self.assertLess(time.monotonic() - started, 15) + self.assertEqual(out.returncode, 2, out.stdout + out.stderr) + self.assertIn("monotonic_us", out.stderr) + + def test_capture_unix_us_is_refused_by_egocap(self): + reason = self.refuse["reason"] + self.assertIsNotNone( + reason, "egocap accepted a capture_unix_us stream for recording" + ) + self.assertIn("monotonic clock", reason) + self.assertEqual(self.refuse["recordable"], 0) + + +if __name__ == "__main__": + unittest.main() diff --git a/examples/camera_livekit_bridge/tests/python/test_ipc_client_probe.py b/examples/camera_livekit_bridge/tests/python/test_ipc_client_probe.py new file mode 100755 index 0000000000..440cb0f51f --- /dev/null +++ b/examples/camera_livekit_bridge/tests/python/test_ipc_client_probe.py @@ -0,0 +1,263 @@ +#!/usr/bin/env python3 +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""camera_ipc_probe must read GPU frames published by the real CUDA IPC publisher. + +A separate publisher process (CUDA IPC handles cannot be opened by the process +that exported them) serves two NV12 streams through ``CudaIpcFramePublisher``; +the probe binary is run against its socket and its stdout is checked. +""" + +from __future__ import annotations + +import os +import re +import signal +import subprocess +import sys +import tempfile +import time +import unittest +from pathlib import Path + +_HERE = Path(__file__).resolve() +_EXAMPLES = _HERE.parents[3] +_RECEIVER_DIR = _EXAMPLES / "camera_streamer" / "camera_receiver" +_DEFAULT_PROBE = _HERE.parents[2] / "build" / "probe" / "camera_ipc_probe" + +# name -> (width, height). Each NV12 buffer must hold at least the 4096 bytes +# the probe checksums. +_STREAMS = {"head": (64, 48), "wrist": (96, 48)} +_RING = 3 +_SESSION_ID = 424242 +_SOF_BASE_NS = 1_000_000_000 +_SOF_STEP_NS = 33_000_000 +_CHECKSUM_BYTES = 4096 +_WAIT_S = 20.0 + +_FRAME_RE = re.compile( + r"^(?P\S+) slot=(?P\d+) gen=(?P\d+) " + r"stream_sequence=(?P-?\d+) session=(?P\d+) " + r"sensor_sof_ns=(?P\S+) checksum=(?P\d+)$" +) +_STREAM_RE = re.compile( + r"^stream (?P\S+) (?P\d+)x(?P\d+) (?P\S+) " + r"slots=(?P\d+) layout=(?P.*)$" +) + + +def _slot_value(slot: int) -> int: + return 10 * slot + 3 + + +def _expected_checksum(slot: int) -> int: + return _CHECKSUM_BYTES * _slot_value(slot) + + +def _probe_path() -> Path: + return Path(os.environ.get("CAMERA_IPC_PROBE", str(_DEFAULT_PROBE))) + + +def _skip_reason() -> str | None: + probe = _probe_path() + if not probe.is_file() or not os.access(probe, os.X_OK): + return f"probe binary not found or not executable: {probe}" + try: + import cupy as cp + except ImportError: + return "CuPy is not importable" + try: + if cp.cuda.runtime.getDeviceCount() < 1: + return "no CUDA device" + except Exception as exc: # CUDA runtime errors vary by driver state + return f"CUDA unavailable: {exc}" + return None + + +def _publisher_main(socket_path: str) -> None: + """Child process: serve the streams and publish frames until terminated.""" + import cupy as cp + + sys.path.insert(0, str(_RECEIVER_DIR)) + from cuda_ipc import CudaIpcFramePublisher + + publisher = CudaIpcFramePublisher(socket_path=socket_path, owner_name="test") + rings = {} + for name, (width, height) in _STREAMS.items(): + y_bytes = width * height + layout = { + "planes": [ + { + "name": "y", + "offset": 0, + "width": width, + "height": height, + "pitch": width, + }, + { + "name": "uv", + "offset": y_bytes, + "width": width, + "height": height // 2, + "pitch": width, + }, + ] + } + buffers = [ + cp.full((y_bytes * 3 // 2,), _slot_value(i), dtype=cp.uint8) + for i in range(_RING) + ] + rings[name] = buffers + publisher.register_stream(name, width, height, "nv12", 0, buffers, layout) + cp.cuda.Device(0).synchronize() + publisher.start() + + stopping = False + + def _stop(*_args: object) -> None: + nonlocal stopping + stopping = True + + signal.signal(signal.SIGTERM, _stop) + print("READY", flush=True) + + seq = 0 + while not stopping: + for name in _STREAMS: + sof_ns = _SOF_BASE_NS + seq * _SOF_STEP_NS + publisher.publish( + name, + slot=seq % _RING, + slot_generation=seq // _RING + 1, + seq=seq, + timestamp_ns=sof_ns, + metadata={ + "session_id": _SESSION_ID, + "stream_name": name, + "stream_sequence": seq, + "sensor_sof_ns": sof_ns, + "sof_tsc_ns": sof_ns - 5, + "sof_tsc_valid": True, + }, + ) + seq += 1 + time.sleep(0.02) + publisher.stop() + + +def _terminate(proc: subprocess.Popen) -> None: + if proc.poll() is None: + proc.send_signal(signal.SIGTERM) + try: + proc.wait(timeout=10) + except subprocess.TimeoutExpired: + proc.kill() + proc.wait(timeout=5) + + +@unittest.skipIf(_skip_reason() is not None, _skip_reason() or "") +class IpcClientProbeTest(unittest.TestCase): + @classmethod + def setUpClass(cls) -> None: + cls._tmp = tempfile.TemporaryDirectory(prefix="ipc_probe_test_") + cls.socket_path = os.path.join(cls._tmp.name, "ipc.sock") + cls._publisher = subprocess.Popen( + [sys.executable, str(_HERE), "--publisher", cls.socket_path], + stdout=subprocess.PIPE, + stderr=subprocess.PIPE, + text=True, + ) + cls.addClassCleanup(cls._tmp.cleanup) + cls.addClassCleanup(_terminate, cls._publisher) + deadline = time.monotonic() + _WAIT_S + line = "" + while time.monotonic() < deadline: + line = cls._publisher.stdout.readline().strip() + if line == "READY" or cls._publisher.poll() is not None: + break + if line != "READY": + _terminate(cls._publisher) + err = cls._publisher.stderr.read() + raise RuntimeError(f"publisher did not become ready: {err}") + + def run_probe(self, *args: str, socket_path: str | None = None, frames: int = 8): + cmd = [ + str(_probe_path()), + "--socket", + socket_path or self.socket_path, + "--frames", + str(frames), + *args, + ] + return subprocess.run( + cmd, capture_output=True, text=True, timeout=_WAIT_S, check=False + ) + + def parse(self, stdout: str): + streams, frames = {}, [] + for line in stdout.splitlines(): + if line.startswith("stream "): + m = _STREAM_RE.match(line) + self.assertIsNotNone(m, line) + streams[m["name"]] = m + else: + m = _FRAME_RE.match(line) + self.assertIsNotNone(m, f"unparseable frame line: {line!r}") + frames.append(m) + return streams, frames + + def check_frames(self, frames) -> None: + for m in frames: + seq = int(m["seq"]) + self.assertEqual(int(m["slot"]), seq % _RING, m.string) + self.assertEqual(int(m["gen"]), seq // _RING + 1, m.string) + self.assertEqual(int(m["session"]), _SESSION_ID, m.string) + self.assertEqual(int(m["sof"]), _SOF_BASE_NS + seq * _SOF_STEP_NS, m.string) + self.assertEqual( + int(m["sum"]), _expected_checksum(int(m["slot"])), m.string + ) + + def test_all_streams(self) -> None: + result = self.run_probe(frames=12) + self.assertEqual(result.returncode, 0, result.stderr) + streams, frames = self.parse(result.stdout) + self.assertEqual(set(streams), set(_STREAMS)) + for name, (w, h) in _STREAMS.items(): + m = streams[name] + self.assertEqual((int(m["w"]), int(m["h"])), (w, h)) + self.assertEqual(m["fmt"], "nv12") + self.assertEqual(int(m["slots"]), _RING) + self.assertIn('"uv"', m["layout"]) + self.assertIn('"y"', m["layout"]) + self.assertEqual(len(frames), 12) + self.assertEqual({m["name"] for m in frames}, set(_STREAMS)) + self.check_frames(frames) + + def test_single_stream_subscription(self) -> None: + result = self.run_probe("--stream", "head", frames=6) + self.assertEqual(result.returncode, 0, result.stderr) + streams, frames = self.parse(result.stdout) + self.assertEqual(set(streams), {"head"}) + self.assertEqual(len(frames), 6) + self.assertEqual({m["name"] for m in frames}, {"head"}) + self.check_frames(frames) + + def test_unknown_stream_fails_clearly(self) -> None: + result = self.run_probe("--stream", "nonexistent") + self.assertNotEqual(result.returncode, 0) + self.assertIn("nonexistent", result.stderr) + self.assertNotRegex(result.stdout, _FRAME_RE.pattern) + + def test_missing_socket_fails_clearly(self) -> None: + missing = os.path.join(self._tmp.name, "no_such.sock") + result = self.run_probe(socket_path=missing) + self.assertNotEqual(result.returncode, 0) + self.assertIn(missing, result.stderr) + + +if __name__ == "__main__": + if len(sys.argv) == 3 and sys.argv[1] == "--publisher": + _publisher_main(sys.argv[2]) + else: + unittest.main() diff --git a/examples/camera_livekit_bridge/tests/timestamp_basis/CMakeLists.txt b/examples/camera_livekit_bridge/tests/timestamp_basis/CMakeLists.txt new file mode 100644 index 0000000000..b2dc28cc11 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/timestamp_basis/CMakeLists.txt @@ -0,0 +1,13 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +add_executable(test_timestamp_basis test_timestamp_basis.cpp) + +target_link_libraries(test_timestamp_basis + PRIVATE + camera_livekit_bridge::video_types +) + +target_compile_options(test_timestamp_basis PRIVATE -Wall -Wextra) + +add_test(NAME timestamp_basis COMMAND test_timestamp_basis) diff --git a/examples/camera_livekit_bridge/tests/timestamp_basis/test_timestamp_basis.cpp b/examples/camera_livekit_bridge/tests/timestamp_basis/test_timestamp_basis.cpp new file mode 100644 index 0000000000..0ec75b4c02 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/timestamp_basis/test_timestamp_basis.cpp @@ -0,0 +1,76 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include + +#include +#include + +namespace +{ + +int g_failures = 0; + +void check(bool condition, const char* what) +{ + if (!condition) + { + std::cerr << "FAILED: " << what << "\n"; + ++g_failures; + } +} + +} // namespace + +int main() +{ + using camera_bridge::TimestampBasis; + using camera_bridge::user_timestamp_us; + + check(camera_bridge::parse_timestamp_basis("monotonic_us") == TimestampBasis::MonotonicUs, "parse monotonic_us"); + check(camera_bridge::parse_timestamp_basis("capture_monotonic_us") == TimestampBasis::CaptureMonotonicUs, + "parse capture_monotonic_us"); + check(camera_bridge::parse_timestamp_basis("capture_unix_us") == TimestampBasis::CaptureUnixUs, + "parse capture_unix_us"); + check(!camera_bridge::parse_timestamp_basis("processing_unix_us").has_value(), "reject an unsupported basis"); + check(!camera_bridge::parse_timestamp_basis("").has_value(), "reject an empty basis"); + + const int64_t realtime = 1'800'000'000'000'000'000; + const int64_t monotonic = 5'000'000'000'000; + + // Frame-ready time ignores the sensor time and is the monotonic clock in microseconds. + check(user_timestamp_us(TimestampBasis::MonotonicUs, 0, realtime, monotonic) == 5'000'000'000, + "monotonic_us without sensor time"); + check(user_timestamp_us(TimestampBasis::MonotonicUs, 123, realtime, monotonic) == 5'000'000'000, + "monotonic_us ignores sensor time"); + + // Capture time 80 ms ago on the sender's wall clock. + const uint64_t sof = static_cast(realtime - 80'000'000); + check(user_timestamp_us(TimestampBasis::CaptureUnixUs, sof, realtime, monotonic) == sof / 1000, + "capture_unix_us is the sensor time"); + check(user_timestamp_us(TimestampBasis::CaptureMonotonicUs, sof, realtime, monotonic) == + (5'000'000'000'000 - 80'000'000) / 1000, + "capture_monotonic_us is now minus the age"); + + // No sensor time, a future one and an implausibly old one are not vouched for. + for (const auto basis : { TimestampBasis::CaptureUnixUs, TimestampBasis::CaptureMonotonicUs }) + { + check(user_timestamp_us(basis, 0, realtime, monotonic) == 0, "no sensor time gives 0"); + check(user_timestamp_us(basis, static_cast(realtime + 1'000'000), realtime, monotonic) == 0, + "future sensor time gives 0"); + check(user_timestamp_us(basis, static_cast(realtime - camera_bridge::kMaxCaptureAgeNs - 1), realtime, + monotonic) == 0, + "too old sensor time gives 0"); + check(user_timestamp_us( + basis, static_cast(realtime - camera_bridge::kMaxCaptureAgeNs), realtime, monotonic) != 0, + "sensor time at the age limit is accepted"); + } + + // An age larger than the monotonic clock itself cannot be placed on it. + check(user_timestamp_us(TimestampBasis::CaptureMonotonicUs, sof, realtime, 10'000'000) == 0, + "age beyond the monotonic clock gives 0"); + + if (g_failures == 0) + std::cout << "timestamp_basis: all checks passed\n"; + return g_failures == 0 ? 0 : 1; +} diff --git a/examples/camera_livekit_bridge/video_types/CMakeLists.txt b/examples/camera_livekit_bridge/video_types/CMakeLists.txt new file mode 100644 index 0000000000..e7ce2b0fdd --- /dev/null +++ b/examples/camera_livekit_bridge/video_types/CMakeLists.txt @@ -0,0 +1,11 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +add_library(video_types INTERFACE) + +target_include_directories(video_types + INTERFACE + $ +) + +add_library(camera_livekit_bridge::video_types ALIAS video_types) diff --git a/examples/camera_livekit_bridge/video_types/inc/video_types/timestamp_basis.hpp b/examples/camera_livekit_bridge/video_types/inc/video_types/timestamp_basis.hpp new file mode 100644 index 0000000000..e2cca3ea5c --- /dev/null +++ b/examples/camera_livekit_bridge/video_types/inc/video_types/timestamp_basis.hpp @@ -0,0 +1,65 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +namespace camera_bridge +{ + +/// What a published frame's application timestamp means; consumers must be told the same value. +enum class TimestampBasis +{ + /// The publisher's CLOCK_MONOTONIC when the frame is published (frame-ready time). + MonotonicUs, + /// The sensor's capture event on the publisher's CLOCK_MONOTONIC. The sensor time arrives as sender wall-clock + /// time, so this assumes the sender's clock agrees with this host's wall clock and carries no bound on that error. + /// The bridge does not offer it until a measured sender-to-host clock mapping exists. + CaptureMonotonicUs, + /// The sensor's capture event as Unix time on the sender's clock. + CaptureUnixUs, +}; + +inline std::optional parse_timestamp_basis(const std::string& name) +{ + if (name == "monotonic_us") + return TimestampBasis::MonotonicUs; + if (name == "capture_monotonic_us") + return TimestampBasis::CaptureMonotonicUs; + if (name == "capture_unix_us") + return TimestampBasis::CaptureUnixUs; + return std::nullopt; +} + +/// A sensor time older than this, or in the future, is not a capture time this host can vouch for. +inline constexpr int64_t kMaxCaptureAgeNs = 5'000'000'000; + +/** + * @brief The frame's `user_timestamp` in microseconds on `basis`, or 0 when the publisher cannot vouch for one. + * + * `sensor_sof_ns` is the sensor start of frame as Unix nanoseconds on the sender's clock (0 when unknown); + * `realtime_now_ns` and `monotonic_now_ns` are this host's CLOCK_REALTIME and CLOCK_MONOTONIC read together. + * For the capture bases a missing, future or implausibly old sensor time gives 0, never a guess. + */ +inline uint64_t user_timestamp_us(TimestampBasis basis, + uint64_t sensor_sof_ns, + int64_t realtime_now_ns, + int64_t monotonic_now_ns) +{ + if (basis == TimestampBasis::MonotonicUs) + return monotonic_now_ns > 0 ? static_cast(monotonic_now_ns) / 1000 : 0; + if (sensor_sof_ns == 0) + return 0; + const int64_t age_ns = realtime_now_ns - static_cast(sensor_sof_ns); + if (age_ns < 0 || age_ns > kMaxCaptureAgeNs) + return 0; + if (basis == TimestampBasis::CaptureUnixUs) + return sensor_sof_ns / 1000; + const int64_t capture_ns = monotonic_now_ns - age_ns; + return capture_ns > 0 ? static_cast(capture_ns) / 1000 : 0; +} + +} // namespace camera_bridge diff --git a/examples/camera_livekit_bridge/video_types/inc/video_types/video_types.hpp b/examples/camera_livekit_bridge/video_types/inc/video_types/video_types.hpp new file mode 100644 index 0000000000..f1492fcb13 --- /dev/null +++ b/examples/camera_livekit_bridge/video_types/inc/video_types/video_types.hpp @@ -0,0 +1,76 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include +#include + +namespace camera_bridge +{ + +struct Geometry +{ + uint32_t width = 0; + uint32_t height = 0; + + bool operator==(const Geometry& other) const + { + return width == other.width && height == other.height; + } +}; + +/// One access unit out of the encoder. +struct EncodedPacket +{ + std::vector data; + uint64_t key = 0; + bool keyframe = false; +}; + +/// Quantisation range of the picture; it decides the value of a border fill (black is Y=16 limited, Y=0 full). +struct ColorInfo +{ + bool full_range = false; +}; + +/// A device-memory NV12 image: a luma plane followed by one interleaved chroma (U,V) plane at half resolution in both +/// directions. Pitches are in bytes and may exceed the row size. A view does not own its memory. +struct Nv12View +{ + uint8_t* y = nullptr; + uint8_t* uv = nullptr; + size_t y_pitch = 0; + size_t uv_pitch = 0; + Geometry geometry; + ColorInfo color; +}; + +struct EncodedVideoFrame +{ + /// Annex-B bytes: an IDR carries its SPS/PPS in front of it. + const uint8_t* data = nullptr; + size_t size = 0; + /// The unit is an IDR, decodable on its own. + bool keyframe = false; + /// CLOCK_MONOTONIC microseconds when the frame was published: the host timeline, not sensor time. + uint64_t timestamp_us = 0; + /// The frame's application timestamp in microseconds, on the clock its stream declares. Zero means the publisher + /// cannot vouch for a time, and the frame carries none. + uint64_t user_timestamp_us = 0; + /// The frame's sequence on its track, from 1. + uint32_t frame_id = 0; + /// The sensor's capture time on its own clock, carried as user data. + uint64_t capture_tsc = 0; +}; + +inline int64_t monotonic_ns() +{ + timespec ts{}; + clock_gettime(CLOCK_MONOTONIC, &ts); + return static_cast(ts.tv_sec) * 1'000'000'000 + ts.tv_nsec; +} + +} // namespace camera_bridge