From e9939d2cde71813b5415fa8d36e485214309f87b Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 21:06:51 -0500 Subject: [PATCH 01/10] examples/camera_livekit_bridge: add the CUDA IPC client and a probe Start a downstream consumer of camera_receiver's CUDA IPC output: a client that connects to the publisher's socket, subscribes, opens each slot's IPC handle and reads frame notices with their metadata, and a probe that prints them with a checksum of the first pixels. Pixels stay on the GPU. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/CMakeLists.txt | 18 ++ examples/camera_livekit_bridge/README.md | 36 +++ .../ipc_client/CMakeLists.txt | 20 ++ .../ipc_client/cuda_ipc_client.cpp | 246 ++++++++++++++++++ .../inc/ipc_client/cuda_ipc_client.hpp | 90 +++++++ .../probe/CMakeLists.txt | 12 + examples/camera_livekit_bridge/probe/main.cpp | 86 ++++++ 7 files changed, 508 insertions(+) create mode 100644 examples/camera_livekit_bridge/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/README.md create mode 100644 examples/camera_livekit_bridge/ipc_client/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/ipc_client/cuda_ipc_client.cpp create mode 100644 examples/camera_livekit_bridge/ipc_client/inc/ipc_client/cuda_ipc_client.hpp create mode 100644 examples/camera_livekit_bridge/probe/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/probe/main.cpp diff --git a/examples/camera_livekit_bridge/CMakeLists.txt b/examples/camera_livekit_bridge/CMakeLists.txt new file mode 100644 index 0000000000..c98f252ee1 --- /dev/null +++ b/examples/camera_livekit_bridge/CMakeLists.txt @@ -0,0 +1,18 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +cmake_minimum_required(VERSION 3.20) +project(camera_livekit_bridge LANGUAGES CXX) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +if(NOT CMAKE_BUILD_TYPE) + set(CMAKE_BUILD_TYPE Release) +endif() + +find_package(CUDAToolkit REQUIRED) +find_package(nlohmann_json REQUIRED) + +add_subdirectory(ipc_client) +add_subdirectory(probe) diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md new file mode 100644 index 0000000000..faecd6d6df --- /dev/null +++ b/examples/camera_livekit_bridge/README.md @@ -0,0 +1,36 @@ + + +# camera_livekit_bridge + +A downstream consumer of `camera_streamer`'s CUDA IPC output that publishes the +camera streams to LiveKit. 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. + +Status: the IPC client and a probe exist. Encoding and publishing are not +implemented yet. + +## Build + +```bash +cmake -S . -B build && cmake --build build -j +``` + +Needs the CUDA toolkit and nlohmann_json. The build is standalone and is not +part of the repository's top-level build. + +## Probe + +With `camera_receiver` running (for example fed by `camera_streamer/mock_sender`): + +```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/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/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..826c1e2e7c --- /dev/null +++ b/examples/camera_livekit_bridge/probe/main.cpp @@ -0,0 +1,86 @@ +// 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 + +namespace +{ + +std::uint64_t first_row_checksum(const camera_bridge::IpcSlot& slot) +{ + std::vector host(4096); + if (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; +} From 1139410aad7e0b2576c149a789b36019f9bda180 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 21:14:03 -0500 Subject: [PATCH 02/10] examples/camera_livekit_bridge: encode the IPC frames and publish them to LiveKit Add the rest of the bridge: each camera stream's frames are copied on the GPU from the CUDA IPC slot into an NVENC input surface, encoded to H.264 and submitted to a LiveKit track as pre-encoded access units, with the frame's sequence, its sensor start time as the application timestamp and the sensor counter as user data. Frames reach the encoder without leaving the GPU. The NVENC encoder and the LiveKit session are small modules; the session also answers the clock-sync requests that arrive as data messages. Each stream runs on its own thread and encodes the newest picture, skipping older ones. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/CMakeLists.txt | 45 +- .../bridge/CMakeLists.txt | 19 + .../camera_livekit_bridge/bridge/main.cpp | 183 ++++++++ .../bridge/stream_worker.cpp | 242 +++++++++++ .../bridge/stream_worker.hpp | 77 ++++ .../livekit_publisher/CMakeLists.txt | 22 + .../inc/livekit_publisher/clock_sync.hpp | 110 +++++ .../inc/livekit_publisher/livekit_session.hpp | 75 ++++ .../livekit_publisher/video_track_sink.hpp | 52 +++ .../livekit_publisher/livekit_session.cpp | 389 ++++++++++++++++++ .../livekit_publisher/log.hpp | 27 ++ .../nvenc/CMakeLists.txt | 29 ++ .../nvenc/cuda_context.cpp | 50 +++ .../nvenc/inc/nvenc/cuda_context.hpp | 49 +++ .../nvenc/inc/nvenc/nvenc_encoder.hpp | 72 ++++ .../nvenc/nvenc_encoder.cpp | 284 +++++++++++++ .../video_types/CMakeLists.txt | 11 + .../inc/video_types/video_types.hpp | 76 ++++ 18 files changed, 1811 insertions(+), 1 deletion(-) create mode 100644 examples/camera_livekit_bridge/bridge/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/bridge/main.cpp create mode 100644 examples/camera_livekit_bridge/bridge/stream_worker.cpp create mode 100644 examples/camera_livekit_bridge/bridge/stream_worker.hpp create mode 100644 examples/camera_livekit_bridge/livekit_publisher/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/clock_sync.hpp create mode 100644 examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/livekit_session.hpp create mode 100644 examples/camera_livekit_bridge/livekit_publisher/inc/livekit_publisher/video_track_sink.hpp create mode 100644 examples/camera_livekit_bridge/livekit_publisher/livekit_session.cpp create mode 100644 examples/camera_livekit_bridge/livekit_publisher/log.hpp create mode 100644 examples/camera_livekit_bridge/nvenc/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/nvenc/cuda_context.cpp create mode 100644 examples/camera_livekit_bridge/nvenc/inc/nvenc/cuda_context.hpp create mode 100644 examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp create mode 100644 examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp create mode 100644 examples/camera_livekit_bridge/video_types/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/video_types/inc/video_types/video_types.hpp diff --git a/examples/camera_livekit_bridge/CMakeLists.txt b/examples/camera_livekit_bridge/CMakeLists.txt index c98f252ee1..ecde04fa23 100644 --- a/examples/camera_livekit_bridge/CMakeLists.txt +++ b/examples/camera_livekit_bridge/CMakeLists.txt @@ -1,7 +1,7 @@ # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -cmake_minimum_required(VERSION 3.20) +cmake_minimum_required(VERSION 3.24) project(camera_livekit_bridge LANGUAGES CXX) set(CMAKE_CXX_STANDARD 17) @@ -13,6 +13,49 @@ 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(nvenc) +add_subdirectory(livekit_publisher) add_subdirectory(probe) +add_subdirectory(bridge) diff --git a/examples/camera_livekit_bridge/bridge/CMakeLists.txt b/examples/camera_livekit_bridge/bridge/CMakeLists.txt new file mode 100644 index 0000000000..efe1e02193 --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/CMakeLists.txt @@ -0,0 +1,19 @@ +# 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::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..6646c702f0 --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -0,0 +1,183 @@ +// 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 [--stream ...]\n" + " [--socket PATH] [--fps N] [--bitrate-mbps N] [--idr-period N] [--gpu N]\n" + " NAME is a stream in camera_receiver's IPC catalog; TRACK is the LiveKit track name egocap expects.\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"; + 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 == "--gpu") + gpu = std::atoi(next().c_str()); + else if (arg == "--stream") + { + const std::string spec = next(); + const auto eq = spec.find('='); + if (eq == std::string::npos || eq == 0 || eq + 1 == spec.size()) + { + std::cerr << "--stream needs NAME=TRACK, got '" << spec << "'\n"; + return 2; + } + streams.emplace_back(spec.substr(0, eq), spec.substr(eq + 1)); + } + 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& [name, track] : streams) + { + WorkerConfig config = defaults; + config.stream = name; + config.track = track; + 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, encode " + << mean_ms << " ms mean\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..2e28fbb7bd --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -0,0 +1,242 @@ +// 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 + +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; +} + +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); + geometry = { static_cast(stream.width), static_cast(stream.height) }; + + 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; + encoder_config.bitrate_bps = config.bitrate_bps; + 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; + while (running) + { + 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"); + const auto copy_status = [&] + { + const auto y = + cudaMemcpy2DAsync(surface.view.y, surface.view.y_pitch, base + layout.y_offset, layout.y_pitch, + geometry.width, geometry.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, + geometry.width, geometry.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)); + } + + 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; + + 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(monotonic_ns() / 1000); + // Capture time: the sensor start-of-frame in Unix nanoseconds when the receiver could derive one; zero + // tells consumers the frame has no capture time they can rely on. + frame.user_timestamp_us = unsigned_field(notice.metadata, "sensor_sof_ns") / 1000; + // 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 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..c8edab43ab --- /dev/null +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -0,0 +1,77 @@ +// 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 + +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; +}; + +struct WorkerStats +{ + std::atomic frames{ 0 }; + std::atomic skipped{ 0 }; + std::atomic keyframes{ 0 }; + std::atomic rejected{ 0 }; + std::atomic encode_us_total{ 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/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..18e5f45a3f --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp @@ -0,0 +1,72 @@ +// 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); + +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..2e165c1c33 --- /dev/null +++ b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp @@ -0,0 +1,284 @@ +// 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; + + 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"); + + NV_ENC_CONFIG 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; + enc_config.rcParams.averageBitRate = config.bitrate_bps; + enc_config.rcParams.maxBitRate = config.bitrate_bps; + // A buffer of one frame's worth of bits keeps the output close to the instantaneous rate. + const uint32_t frame_bits = static_cast(static_cast(config.bitrate_bps) * + config.fps_denominator / config.fps_numerator); + enc_config.rcParams.vbvBufferSize = frame_bits; + enc_config.rcParams.vbvInitialDelay = frame_bits; + 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.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"); + } + + 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; +} + +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/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/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 From f8b0d686c5c26fc59ecb83b987895836613b7084 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 21:25:29 -0500 Subject: [PATCH 03/10] examples/camera_livekit_bridge: shrink a stream on the GPU before encoding Add a per-stream integer divisor (--stream NAME=TRACK:DIV) that averages each block of the NV12 picture on the GPU between the IPC slot and the encoder input, so a large camera stream can be published at a size the consumers can handle without the pixels leaving the GPU. This is also the point where a GPU operation on the picture belongs. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/CMakeLists.txt | 6 +- .../bridge/CMakeLists.txt | 1 + .../camera_livekit_bridge/bridge/main.cpp | 36 +++++-- .../bridge/stream_worker.cpp | 23 +++- .../bridge/stream_worker.hpp | 2 + .../gpu_ops/CMakeLists.txt | 17 +++ .../gpu_ops/box_downscale.cu | 102 ++++++++++++++++++ .../gpu_ops/inc/gpu_ops/box_downscale.hpp | 42 ++++++++ 8 files changed, 215 insertions(+), 14 deletions(-) create mode 100644 examples/camera_livekit_bridge/gpu_ops/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/gpu_ops/box_downscale.cu create mode 100644 examples/camera_livekit_bridge/gpu_ops/inc/gpu_ops/box_downscale.hpp diff --git a/examples/camera_livekit_bridge/CMakeLists.txt b/examples/camera_livekit_bridge/CMakeLists.txt index ecde04fa23..0901251d52 100644 --- a/examples/camera_livekit_bridge/CMakeLists.txt +++ b/examples/camera_livekit_bridge/CMakeLists.txt @@ -2,11 +2,14 @@ # SPDX-License-Identifier: Apache-2.0 cmake_minimum_required(VERSION 3.24) -project(camera_livekit_bridge LANGUAGES CXX) +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() @@ -55,6 +58,7 @@ 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) diff --git a/examples/camera_livekit_bridge/bridge/CMakeLists.txt b/examples/camera_livekit_bridge/bridge/CMakeLists.txt index efe1e02193..c647bca7ef 100644 --- a/examples/camera_livekit_bridge/bridge/CMakeLists.txt +++ b/examples/camera_livekit_bridge/bridge/CMakeLists.txt @@ -9,6 +9,7 @@ add_executable(camera_livekit_bridge 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 diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index 6646c702f0..1c858f882c 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -32,9 +32,9 @@ void on_signal(int) void usage() { std::cerr - << "usage: camera_livekit_bridge --livekit-url URL --livekit-token-file FILE --stream NAME=TRACK [--stream ...]\n" + << "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" - " NAME is a stream in camera_receiver's IPC catalog; TRACK is the LiveKit track name egocap expects.\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) @@ -58,7 +58,13 @@ int main(int argc, char** argv) std::string url; std::string token_file; std::string socket_path = "/tmp/camera_receiver_avs1_cuda_ipc.sock"; - std::vector> streams; + 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) @@ -89,14 +95,23 @@ int main(int argc, char** argv) gpu = std::atoi(next().c_str()); else if (arg == "--stream") { - const std::string spec = next(); + 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()) + if (eq == std::string::npos || eq == 0 || eq + 1 == spec.size() || parsed.scale_div == 0) { - std::cerr << "--stream needs NAME=TRACK, got '" << spec << "'\n"; + std::cerr << "--stream needs NAME=TRACK or NAME=TRACK:DIV, got '" << spec << "'\n"; return 2; } - streams.emplace_back(spec.substr(0, eq), spec.substr(eq + 1)); + parsed.name = spec.substr(0, eq); + parsed.track = spec.substr(eq + 1); + streams.push_back(parsed); } else { @@ -118,11 +133,12 @@ int main(int argc, char** argv) CudaContext context(gpu); auto session = LiveKitSession::connect(url, read_token(token_file)); std::vector> workers; - for (const auto& [name, track] : streams) + for (const auto& spec : streams) { WorkerConfig config = defaults; - config.stream = name; - config.track = track; + 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)); } diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 2e28fbb7bd..6e28eda19a 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -3,6 +3,7 @@ #include "stream_worker.hpp" +#include #include #include #include @@ -71,7 +72,12 @@ struct StreamWorker::Impl throw std::runtime_error("the IPC publisher has no stream named '" + cfg.stream + "'"); const IpcStream& stream = client.streams().front(); layout = read_layout(stream); - geometry = { static_cast(stream.width), static_cast(stream.height) }; + 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; @@ -140,15 +146,25 @@ struct StreamWorker::Impl NvencEncoder::Surface surface; if (!encoder.acquire(surface)) throw std::runtime_error("no free encoder input surface"); + // The picture goes from the IPC slot to the encoder input on the GPU: a plain copy at full size, or a block + // average that shrinks it. The encode below waits for this stream. const auto copy_status = [&] { + if (config.scale_div > 1) + { + const Nv12Planes in{ base + layout.y_offset, base + layout.uv_offset, layout.y_pitch, + layout.uv_pitch }; + const Nv12PlanesOut out{ surface.view.y, surface.view.uv, surface.view.y_pitch, + surface.view.uv_pitch }; + return box_downscale_nv12(in, source.width, source.height, out, config.scale_div, copy_stream); + } const auto y = cudaMemcpy2DAsync(surface.view.y, surface.view.y_pitch, base + layout.y_offset, layout.y_pitch, - geometry.width, geometry.height, cudaMemcpyDeviceToDevice, copy_stream); + 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, - geometry.width, geometry.height / 2, cudaMemcpyDeviceToDevice, copy_stream); + source.width, source.height / 2, cudaMemcpyDeviceToDevice, copy_stream); }(); if (copy_status != cudaSuccess) { @@ -204,6 +220,7 @@ struct StreamWorker::Impl CudaContext& context; CudaIpcClient client; PlaneLayout layout; + Geometry source; Geometry geometry; std::unique_ptr sink; cudaStream_t copy_stream = nullptr; diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp index c8edab43ab..c3f91d9cf6 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.hpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -25,6 +25,8 @@ struct WorkerConfig 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; }; struct WorkerStats 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 From 77e0de44d02f6059b0a7dfc8aa7ebe3c4756d11d Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 21:39:08 -0500 Subject: [PATCH 04/10] examples/camera_livekit_bridge: document the bridge Describe how the bridge moves frames from the IPC slot to a LiveKit track, how to build and run it, the per-stream shrink option, and the capacity measured against egocap's receiver and recording checks. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/README.md | 57 ++++++++++++++++++++---- 1 file changed, 48 insertions(+), 9 deletions(-) diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md index faecd6d6df..5998b7fa64 100644 --- a/examples/camera_livekit_bridge/README.md +++ b/examples/camera_livekit_bridge/README.md @@ -6,12 +6,23 @@ SPDX-License-Identifier: Apache-2.0 # camera_livekit_bridge A downstream consumer of `camera_streamer`'s CUDA IPC output that publishes the -camera streams to LiveKit. 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. +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. -Status: the IPC client and a probe exist. Encoding and publishing are not -implemented yet. +```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 session also answers clock-sync requests that +arrive as LiveKit data messages. ## Build @@ -19,12 +30,40 @@ implemented yet. cmake -S . -B build && cmake --build build -j ``` -Needs the CUDA toolkit and nlohmann_json. The build is standalone and is not -part of the repository's top-level build. +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. -## Probe +## 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 and rejected frames and mean encode time. -With `camera_receiver` running (for example fed by `camera_streamer/mock_sender`): +Measured with the mock sender on an RTX PRO 5000 against egocap's receiver and +recording-feed checks (30 fps, head and one wrist): 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. The head at 20 Mbps built a backlog of +about 1.2 s, so keep the bitrate near 8 Mbps. The first frames after a consumer +subscribes can be 1-2 s old. + +## Probe ```bash build/probe/camera_ipc_probe --socket /tmp/camera_receiver_avs1_cuda_ipc.sock \ From 823ad590d3a3336a7bbbd2fbbfc656b1d0939ed5 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 22:06:01 -0500 Subject: [PATCH 05/10] examples/camera_livekit_bridge: add tests for the bridge modules and the pipeline Add unit tests for the GPU downscale (exact against a CPU reference), the clock-sync wire format (golden bytes) and the NVENC encoder (Annex-B output, SPS/PPS and IDR, pool behaviour, decodable by ffmpeg), a test that reads known GPU patterns through the IPC client and probe from a separate publisher process, and an end-to-end test of mock sender, receiver, bridge, a local LiveKit server and a subscriber that checks track names, sizes, frame ids, capture-time age and user data. Each skips with a stated reason when a prerequisite is missing. The probe's checksum now reads at most the slot's size instead of always 4096 bytes. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/CMakeLists.txt | 5 + examples/camera_livekit_bridge/probe/main.cpp | 5 +- .../tests/CMakeLists.txt | 10 + .../tests/clock_sync/CMakeLists.txt | 7 + .../tests/clock_sync/test_clock_sync.cpp | 133 +++++ .../tests/gpu_ops/CMakeLists.txt | 8 + .../tests/gpu_ops/test_box_downscale.cpp | 228 +++++++++ .../tests/nvenc/CMakeLists.txt | 8 + .../tests/nvenc/test_nvenc_encoder.cpp | 228 +++++++++ .../tests/python/test_bridge_pipeline.py | 478 ++++++++++++++++++ .../tests/python/test_ipc_client_probe.py | 263 ++++++++++ 11 files changed, 1371 insertions(+), 2 deletions(-) create mode 100644 examples/camera_livekit_bridge/tests/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/tests/clock_sync/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/tests/clock_sync/test_clock_sync.cpp create mode 100644 examples/camera_livekit_bridge/tests/gpu_ops/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/tests/gpu_ops/test_box_downscale.cpp create mode 100644 examples/camera_livekit_bridge/tests/nvenc/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp create mode 100644 examples/camera_livekit_bridge/tests/python/test_bridge_pipeline.py create mode 100755 examples/camera_livekit_bridge/tests/python/test_ipc_client_probe.py diff --git a/examples/camera_livekit_bridge/CMakeLists.txt b/examples/camera_livekit_bridge/CMakeLists.txt index 0901251d52..0c8bb954fa 100644 --- a/examples/camera_livekit_bridge/CMakeLists.txt +++ b/examples/camera_livekit_bridge/CMakeLists.txt @@ -63,3 +63,8 @@ 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/probe/main.cpp b/examples/camera_livekit_bridge/probe/main.cpp index 826c1e2e7c..b359289fde 100644 --- a/examples/camera_livekit_bridge/probe/main.cpp +++ b/examples/camera_livekit_bridge/probe/main.cpp @@ -6,6 +6,7 @@ #include +#include #include #include #include @@ -17,8 +18,8 @@ namespace std::uint64_t first_row_checksum(const camera_bridge::IpcSlot& slot) { - std::vector host(4096); - if (cudaMemcpy(host.data(), slot.device_ptr, host.size(), cudaMemcpyDeviceToHost) != cudaSuccess) + 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) 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..c43f26b541 --- /dev/null +++ b/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.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 +#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 - 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); + + 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); + + 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_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() From 6db0055627ced063bd753b0a786eeaf6380dbb6d Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 22:17:04 -0500 Subject: [PATCH 06/10] examples/camera_livekit_bridge: choose what the frame timestamp means Add --timestamp-basis: capture_unix_us (default, the sensor start-of-frame as Unix time), capture_monotonic_us (the same instant on this host's monotonic clock, derived as now minus the frame's age) and monotonic_us (the time the frame is published). A frame with no usable sensor start-of-frame carries no timestamp on the capture bases. The conversion is a small header-only helper with its own test. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/README.md | 12 +++ .../camera_livekit_bridge/bridge/main.cpp | 12 ++- .../bridge/stream_worker.cpp | 11 ++- .../bridge/stream_worker.hpp | 3 + .../tests/timestamp_basis/CMakeLists.txt | 13 ++++ .../timestamp_basis/test_timestamp_basis.cpp | 76 +++++++++++++++++++ .../inc/video_types/timestamp_basis.hpp | 64 ++++++++++++++++ 7 files changed, 187 insertions(+), 4 deletions(-) create mode 100644 examples/camera_livekit_bridge/tests/timestamp_basis/CMakeLists.txt create mode 100644 examples/camera_livekit_bridge/tests/timestamp_basis/test_timestamp_basis.cpp create mode 100644 examples/camera_livekit_bridge/video_types/inc/video_types/timestamp_basis.hpp diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md index 5998b7fa64..fb603f04e0 100644 --- a/examples/camera_livekit_bridge/README.md +++ b/examples/camera_livekit_bridge/README.md @@ -56,6 +56,18 @@ 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 and rejected 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. | +| `capture_monotonic_us` | The sensor start-of-frame on this host's monotonic clock, derived as now minus the frame's age, so it assumes the sender's wall clock agrees with this host's. | +| `monotonic_us` | This host's monotonic clock when the frame is published (frame-ready time). | + +For the capture bases a frame with no usable sensor start-of-frame (missing, in the future, or older than 5 s) +carries no `user_timestamp`, so a consumer that needs a capture time says so instead of using a guess. + Measured with the mock sender on an RTX PRO 5000 against egocap's receiver and recording-feed checks (30 fps, head and one wrist): the head at 2560x1984 and 8 Mbps and at 1280x992 both ran with no missing frames and a median diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index 1c858f882c..04102c4e4c 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -33,7 +33,7 @@ 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" + " [--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"; } @@ -91,6 +91,16 @@ int main(int argc, char** argv) 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; + } + defaults.timestamp_basis = *basis; + } else if (arg == "--gpu") gpu = std::atoi(next().c_str()); else if (arg == "--stream") diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 6e28eda19a..1cd5ad2f1b 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -9,6 +9,7 @@ #include #include +#include #include #include #include @@ -195,9 +196,13 @@ struct StreamWorker::Impl frame.size = packet.data.size(); frame.keyframe = packet.keyframe; frame.timestamp_us = static_cast(monotonic_ns() / 1000); - // Capture time: the sensor start-of-frame in Unix nanoseconds when the receiver could derive one; zero - // tells consumers the frame has no capture time they can rely on. - frame.user_timestamp_us = unsigned_field(notice.metadata, "sensor_sof_ns") / 1000; + // The application timestamp on the configured basis; zero tells consumers the frame has no time the + // publisher can vouch for (for the capture bases, no usable sensor start-of-frame). + timespec realtime{}; + clock_gettime(CLOCK_REALTIME, &realtime); + const int64_t realtime_ns = static_cast(realtime.tv_sec) * 1'000'000'000 + realtime.tv_nsec; + frame.user_timestamp_us = user_timestamp_us( + config.timestamp_basis, unsigned_field(notice.metadata, "sensor_sof_ns"), realtime_ns, monotonic_ns()); // 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; diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp index c3f91d9cf6..cfa1099be4 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.hpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -5,6 +5,7 @@ #include #include +#include #include #include @@ -27,6 +28,8 @@ struct WorkerConfig uint32_t idr_period = 30; /// Integer factor the picture is shrunk by on the GPU before encoding (1 keeps the source size). uint32_t scale_div = 1; + /// What each frame's user_timestamp means; the consumer's streams file must name the same basis. + TimestampBasis timestamp_basis = TimestampBasis::CaptureUnixUs; }; struct WorkerStats 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/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..bc97f72c50 --- /dev/null +++ b/examples/camera_livekit_bridge/video_types/inc/video_types/timestamp_basis.hpp @@ -0,0 +1,64 @@ +// 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. + 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 From 359c19c7b5b68a396215657dcfd9549d16a0b63c Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 22:22:34 -0500 Subject: [PATCH 07/10] examples/camera_livekit_bridge: stamp frame-ready before encoding, gate unqualified capture time Take the frame-ready instant when the picture is on the encoder's input, before encoding, so the monotonic timestamp does not include encode delay. Use the sender's sensor start-of-frame only when the receiver reports a realtime (Unix) sender clock; any other domain publishes no capture time instead of labelling it Unix. Refuse --timestamp-basis capture_monotonic_us: it moved the sensor time onto the host's monotonic clock by assuming the sender's wall clock agrees with the host's, with no bound on that error. It needs a measured sender-to-host clock mapping, which does not exist yet. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/README.md | 9 ++++--- .../camera_livekit_bridge/bridge/main.cpp | 7 ++++++ .../bridge/stream_worker.cpp | 25 +++++++++++++------ .../inc/video_types/timestamp_basis.hpp | 3 ++- 4 files changed, 31 insertions(+), 13 deletions(-) diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md index fb603f04e0..4fecdae3ef 100644 --- a/examples/camera_livekit_bridge/README.md +++ b/examples/camera_livekit_bridge/README.md @@ -62,11 +62,12 @@ the same basis: | Basis | Value | | --- | --- | | `capture_unix_us` (default) | The sensor start-of-frame as Unix time on the sender's clock. | -| `capture_monotonic_us` | The sensor start-of-frame on this host's monotonic clock, derived as now minus the frame's age, so it assumes the sender's wall clock agrees with this host's. | -| `monotonic_us` | This host's monotonic clock when the frame is published (frame-ready time). | +| `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. | -For the capture bases a frame with no usable sensor start-of-frame (missing, in the future, or older than 5 s) -carries no `user_timestamp`, so a consumer that needs a capture time says so instead of using a guess. +`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. Measured with the mock sender on an RTX PRO 5000 against egocap's receiver and recording-feed checks (30 fps, head and one wrist): the head at 2560x1984 and diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index 04102c4e4c..c53a769d34 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -99,6 +99,13 @@ int main(int argc, char** argv) 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") diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 1cd5ad2f1b..0ce9981b51 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -184,6 +184,14 @@ struct StreamWorker::Impl const auto feedback = sink->take_feedback(); const bool force_idr = need_idr || feedback.keyframe_requested; + // 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( @@ -195,14 +203,15 @@ struct StreamWorker::Impl frame.data = packet.data.data(); frame.size = packet.data.size(); frame.keyframe = packet.keyframe; - frame.timestamp_us = static_cast(monotonic_ns() / 1000); - // The application timestamp on the configured basis; zero tells consumers the frame has no time the - // publisher can vouch for (for the capture bases, no usable sensor start-of-frame). - timespec realtime{}; - clock_gettime(CLOCK_REALTIME, &realtime); - const int64_t realtime_ns = static_cast(realtime.tv_sec) * 1'000'000'000 + realtime.tv_nsec; - frame.user_timestamp_us = user_timestamp_us( - config.timestamp_basis, unsigned_field(notice.metadata, "sensor_sof_ns"), realtime_ns, monotonic_ns()); + 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); // 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; 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 index bc97f72c50..e2cca3ea5c 100644 --- 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 @@ -16,7 +16,8 @@ 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. + /// 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, From 4d7028f61a5ebd4660a9106330401c2d84963756 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 22:27:58 -0500 Subject: [PATCH 08/10] examples/camera_livekit_bridge: follow the transport's target bitrate Start each stream's encoder at 2 Mbps and apply the target bitrate that the LiveKit congestion controller reports for the track (10% headroom, capped at --bitrate-mbps, floor 0.5 Mbps) by reconfiguring NVENC without a restart or an IDR. A fixed high rate sent into a connection whose bandwidth is not yet known queued frames behind the sender's pacer, making the first frames 1-2 s old and growing. With the adaptive start they are about 60-150 ms old and settle near 80 ms. The stats line reports the current rate, the transport target and the number of changes. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/README.md | 17 ++++--- .../camera_livekit_bridge/bridge/main.cpp | 4 +- .../bridge/stream_worker.cpp | 26 +++++++++- .../bridge/stream_worker.hpp | 5 ++ .../nvenc/inc/nvenc/nvenc_encoder.hpp | 8 +++ .../nvenc/nvenc_encoder.cpp | 51 +++++++++++++++---- 6 files changed, 94 insertions(+), 17 deletions(-) diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md index 4fecdae3ef..9860324bb1 100644 --- a/examples/camera_livekit_bridge/README.md +++ b/examples/camera_livekit_bridge/README.md @@ -69,12 +69,17 @@ start-of-frame (missing, in the future, older than 5 s, or on another clock doma 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. -Measured with the mock sender on an RTX PRO 5000 against egocap's receiver and -recording-feed checks (30 fps, head and one wrist): 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. The head at 20 Mbps built a backlog of -about 1.2 s, so keep the bitrate near 8 Mbps. The first frames after a consumer -subscribes can be 1-2 s old. +`--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 diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index c53a769d34..45de70db70 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -196,7 +196,9 @@ int main(int argc, char** argv) 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, encode " - << mean_ms << " ms mean\n"; + << mean_ms << " ms mean, bitrate " << stats.bitrate_bps.load() / 1000 + << " kbps (transport target " << stats.target_bps.load() / 1000 << " kbps, " + << stats.rate_changes.load() << " changes)\n"; last_frames[i] = frames; } last_report = now; diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 0ce9981b51..80574b7be6 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -8,6 +8,7 @@ #include #include +#include #include #include #include @@ -55,6 +56,8 @@ PlaneLayout read_layout(const IpcStream& stream) 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); @@ -120,7 +123,9 @@ struct StreamWorker::Impl EncoderConfig encoder_config; encoder_config.geometry = geometry; encoder_config.fps_numerator = config.fps; - encoder_config.bitrate_bps = config.bitrate_bps; + // 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); @@ -183,6 +188,25 @@ struct StreamWorker::Impl } 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. diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp index cfa1099be4..a64d178d3b 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.hpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -39,6 +39,11 @@ struct WorkerStats std::atomic keyframes{ 0 }; std::atomic rejected{ 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 }; + /// The latest target the transport reported (0 until it reports one). + std::atomic target_bps{ 0 }; }; /** diff --git a/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp b/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp index 18e5f45a3f..eb654cc47f 100644 --- a/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp +++ b/examples/camera_livekit_bridge/nvenc/inc/nvenc/nvenc_encoder.hpp @@ -64,6 +64,14 @@ class NvencEncoder /// 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; diff --git a/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp index 2e165c1c33..4597782249 100644 --- a/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp +++ b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp @@ -41,6 +41,10 @@ struct NvencEncoder::Impl 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) { @@ -119,18 +123,12 @@ struct NvencEncoder::Impl check_nv(fn.nvEncGetEncodePresetConfigEx(encoder, codec, preset, tuning, &preset_config), "nvEncGetEncodePresetConfigEx"); - NV_ENC_CONFIG enc_config = preset_config.presetCfg; + 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; - enc_config.rcParams.averageBitRate = config.bitrate_bps; - enc_config.rcParams.maxBitRate = config.bitrate_bps; - // A buffer of one frame's worth of bits keeps the output close to the instantaneous rate. - const uint32_t frame_bits = static_cast(static_cast(config.bitrate_bps) * - config.fps_denominator / config.fps_numerator); - enc_config.rcParams.vbvBufferSize = frame_bits; - enc_config.rcParams.vbvInitialDelay = frame_bits; + apply_rate(config.bitrate_bps); enc_config.rcParams.lookaheadDepth = 0; NV_ENC_CONFIG_H264& h264 = enc_config.encodeCodecConfig.h264Config; h264.idrPeriod = config.idr_period; @@ -140,7 +138,8 @@ struct NvencEncoder::Impl h264.h264VUIParameters.videoFormat = NV_ENC_VUI_VIDEO_FORMAT_UNSPECIFIED; h264.h264VUIParameters.videoFullRangeFlag = config.full_range ? 1 : 0; - NV_ENC_INITIALIZE_PARAMS init{}; + NV_ENC_INITIALIZE_PARAMS& init = init_params; + init = {}; init.version = NV_ENC_INITIALIZE_PARAMS_VER; init.encodeGUID = codec; init.presetGUID = preset; @@ -158,6 +157,19 @@ struct NvencEncoder::Impl 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); @@ -214,6 +226,27 @@ bool NvencEncoder::acquire(Surface& out) 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(); + 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; + check_nv(e.fn.nvEncReconfigureEncoder(e.encoder, &reconfigure), "nvEncReconfigureEncoder"); +} + +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; From f4e3e7a60950a2dae08f4d91005cbe68f6bb7e2e Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 22:31:22 -0500 Subject: [PATCH 09/10] examples/camera_livekit_bridge: report capture-to-ready latency, keep the rate honest, add timing tests Report the sensor start-of-frame to frame-ready time in the stats line, from frames whose sensor time is on a realtime sender clock, so the latency before a consumer's own frame-ready clock starts can be added to what it reports. If NVENC rejects a bitrate change the encoder keeps its old rate, so the reported rate and the settings for the next change are restored too. Add tests: the bitrate setter and getter, zero rejected, no forced IDR after a reconfigure, smaller packets at a lower rate, and a stream that stays decodable across changes; and a timing test that runs egocap's clock-sync client and recording feed against the bridge on monotonic_us, refuses capture_monotonic_us, and checks that egocap's loader rejects capture_unix_us under its timing branch. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- .../camera_livekit_bridge/bridge/main.cpp | 10 +- .../bridge/stream_worker.cpp | 14 + .../bridge/stream_worker.hpp | 5 + .../nvenc/nvenc_encoder.cpp | 12 +- .../tests/nvenc/test_nvenc_encoder.cpp | 148 ++++++ .../tests/python/test_bridge_timing.py | 472 ++++++++++++++++++ 6 files changed, 659 insertions(+), 2 deletions(-) create mode 100644 examples/camera_livekit_bridge/tests/python/test_bridge_timing.py diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index 45de70db70..1a7a00d2a7 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -198,7 +198,15 @@ int main(int argc, char** argv) << stats.skipped.load() << " skipped, " << stats.rejected.load() << " rejected, 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)\n"; + << 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; diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 80574b7be6..06f35f1a24 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -236,6 +236,20 @@ struct StreamWorker::Impl 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; diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp index a64d178d3b..f480dba514 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.hpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -42,6 +42,11 @@ struct WorkerStats /// 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 }; }; diff --git a/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp index 4597782249..11245481d8 100644 --- a/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp +++ b/examples/camera_livekit_bridge/nvenc/nvenc_encoder.cpp @@ -232,6 +232,7 @@ void NvencEncoder::set_bitrate(uint32_t bitrate_bps) 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; @@ -239,7 +240,16 @@ void NvencEncoder::set_bitrate(uint32_t bitrate_bps) reconfigure.reInitEncodeParams.encodeConfig = &e.enc_config; reconfigure.resetEncoder = 0; reconfigure.forceIDR = 0; - check_nv(e.fn.nvEncReconfigureEncoder(e.encoder, &reconfigure), "nvEncReconfigureEncoder"); + 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 diff --git a/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp b/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp index c43f26b541..6f03097785 100644 --- a/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp +++ b/examples/camera_livekit_bridge/tests/nvenc/test_nvenc_encoder.cpp @@ -127,6 +127,149 @@ void test_decodable(const std::vector& idr) std::remove(path.c_str()); } +void fill_noise(const NvencEncoder::Surface& s, uint32_t seed) +{ + const uint32_t w = s.view.geometry.width; + const uint32_t h = s.view.geometry.height; + std::vector 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 - first_idr; constexpr int kFrames = 8; @@ -217,6 +362,9 @@ int main() } test_decodable(first_idr); + test_output_after_reconfigure(*enc); + test_rate_changes_output_size(*enc); + test_decodable_across_reconfigure(*enc); if (g_failures != 0) { 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() From d950e1fd47d59b90c316ff2912eabed7cfaed9f6 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Tue, 10 Nov 2026 23:27:46 -0500 Subject: [PATCH 10/10] examples/camera_livekit_bridge: drop a picture whose IPC slot may have been overwritten The receiver overwrites a slot without notice once its ring wraps. After the GPU copy completes, the bridge now 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 a possibly torn one. Dropped pictures are counted as slot overruns in the stats line. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- examples/camera_livekit_bridge/README.md | 6 ++- .../camera_livekit_bridge/bridge/main.cpp | 8 ++-- .../bridge/stream_worker.cpp | 38 ++++++++++++++++++- .../bridge/stream_worker.hpp | 2 + 4 files changed, 47 insertions(+), 7 deletions(-) diff --git a/examples/camera_livekit_bridge/README.md b/examples/camera_livekit_bridge/README.md index 9860324bb1..f14ffae051 100644 --- a/examples/camera_livekit_bridge/README.md +++ b/examples/camera_livekit_bridge/README.md @@ -21,7 +21,9 @@ 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 session also answers clock-sync requests that +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 @@ -54,7 +56,7 @@ 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 and rejected frames and mean encode time. +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: diff --git a/examples/camera_livekit_bridge/bridge/main.cpp b/examples/camera_livekit_bridge/bridge/main.cpp index 1a7a00d2a7..45af62f354 100644 --- a/examples/camera_livekit_bridge/bridge/main.cpp +++ b/examples/camera_livekit_bridge/bridge/main.cpp @@ -195,10 +195,10 @@ int main(int argc, char** argv) 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, 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)"; + << 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) { diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.cpp b/examples/camera_livekit_bridge/bridge/stream_worker.cpp index 06f35f1a24..44ebba2a33 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.cpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.cpp @@ -13,6 +13,7 @@ #include #include #include +#include #include #include @@ -134,9 +135,15 @@ struct StreamWorker::Impl 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 (!client.next_frame(notice, std::chrono::milliseconds(200))) + 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. @@ -179,6 +186,35 @@ struct StreamWorker::Impl 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) { diff --git a/examples/camera_livekit_bridge/bridge/stream_worker.hpp b/examples/camera_livekit_bridge/bridge/stream_worker.hpp index f480dba514..f1879cbb04 100644 --- a/examples/camera_livekit_bridge/bridge/stream_worker.hpp +++ b/examples/camera_livekit_bridge/bridge/stream_worker.hpp @@ -38,6 +38,8 @@ struct WorkerStats 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 };