From 97dfdf8483b85ca874c9695b333f5232f51c4b6b Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sat, 24 Oct 2026 17:26:19 -0400 Subject: [PATCH 01/10] Add a mock backend so the SENSING plugin runs without a Jetson Add SENSING_BACKEND (AUTO, JETSON, MOCK). AUTO keeps JETSON on aarch64, so a Jetson missing the SDK still gets the setup instructions, and falls back to MOCK elsewhere. The mock replaces only SiplCamera and JetsonEncoder. Everything above them is the code that runs on the rig: SensingCamera, the sinks, the YUV-to-RGBA and RGBA-to-NV12 kernels, and the CUDA IPC publisher. - SiplCamera: answers query() from the vendor JSON, including the renumbering of sensors to pipeline indices and the per-deserializer link masks; paints a synthetic NV12 surface and runs the real conversion on it; keeps the three-slot mailbox with its reader lease and a capture TSC shared by every sensor. Fault injection and scaling through SENSING_MOCK_* variables. - JetsonEncoder: runs the real RGBA-to-NV12 kernel and the device-to-host copy, then emits Annex-B shaped units (not decodable video) with the real encoder's contract: six input buffers that make submit() refuse frames when behind, a warm-up delay, microsecond-resolution stamps, and a concatenated tail at end of stream. Tests (sensing_mock_tests) cover the camera, the encoder, SensingCamera end to end, and a CUDA IPC consumer that maps the shared allocation and reads the published frames. The existing test_sipl_query now also runs against the mock. One test is marked [!shouldfail]: the capture TSC of an encoded stream is never recovered, because the encoder returns stamps at microsecond resolution and SensingCamera matches them for nanosecond equality. Nothing here has been run against SIPL or the V4L2 encoder. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/CMakeLists.txt | 34 +- src/plugins/sensing/README.md | 42 +++ src/plugins/sensing/core/CMakeLists.txt | 92 +++-- .../sensing/core/mock_jetson_encoder.cpp | 296 ++++++++++++++++ src/plugins/sensing/core/mock_nv12_source.cu | 83 +++++ src/plugins/sensing/core/mock_nv12_source.cuh | 34 ++ .../sensing/core/mock_platform_config.cpp | 107 ++++++ .../sensing/core/mock_platform_config.hpp | 40 +++ src/plugins/sensing/core/mock_sipl_camera.cpp | 334 ++++++++++++++++++ tests/cpp/plugins/CMakeLists.txt | 4 + tests/cpp/plugins/sensing_mock/CMakeLists.txt | 29 ++ .../sensing_mock/mock_test_support.hpp | 146 ++++++++ .../sensing_mock/test_cuda_ipc_consumer.cpp | 285 +++++++++++++++ .../plugins/sensing_mock/test_mock_camera.cpp | 200 +++++++++++ .../sensing_mock/test_mock_encoder.cpp | 179 ++++++++++ .../sensing_mock/test_mock_pipeline.cpp | 216 +++++++++++ 16 files changed, 2095 insertions(+), 26 deletions(-) create mode 100644 src/plugins/sensing/core/mock_jetson_encoder.cpp create mode 100644 src/plugins/sensing/core/mock_nv12_source.cu create mode 100644 src/plugins/sensing/core/mock_nv12_source.cuh create mode 100644 src/plugins/sensing/core/mock_platform_config.cpp create mode 100644 src/plugins/sensing/core/mock_platform_config.hpp create mode 100644 src/plugins/sensing/core/mock_sipl_camera.cpp create mode 100644 tests/cpp/plugins/sensing_mock/CMakeLists.txt create mode 100644 tests/cpp/plugins/sensing_mock/mock_test_support.hpp create mode 100644 tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp create mode 100644 tests/cpp/plugins/sensing_mock/test_mock_camera.cpp create mode 100644 tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp create mode 100644 tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp diff --git a/src/plugins/sensing/CMakeLists.txt b/src/plugins/sensing/CMakeLists.txt index a168711c0c..3acff5b1bd 100644 --- a/src/plugins/sensing/CMakeLists.txt +++ b/src/plugins/sensing/CMakeLists.txt @@ -4,7 +4,12 @@ # ============================================================================== # SENSING Camera Plugin (Jetson SIPL capture + V4L2 H.264 encode) # ============================================================================== -# Jetson-only: SIPL and the V4L2 M2M encoder exist only on L4T. +# Two backends sit behind SiplCamera and JetsonEncoder: +# JETSON SIPL capture and the V4L2 M2M encoder, which exist only on L4T. +# MOCK synthetic frames and a stand-in encoder, for any machine with an +# NVIDIA GPU and the CUDA toolkit. Everything above those two classes +# (SensingCamera, sinks, CUDA IPC, the app) is the same code. +# SENSING_BACKEND=AUTO picks JETSON on aarch64 and MOCK elsewhere. # # Four leaf targets, one per directory (cmake/cmake-structure.md rule 1): # cuda_ipc/ sensing_cuda_ipc — CUDA-only IPC publisher + wire format @@ -12,14 +17,16 @@ # app/ camera_plugin_sensing # tools/ sensing_ipc_testsrc — links cuda_ipc only, no SIPL, no encoder # -# The Jetson SDK checks below gate the whole plugin, tools/ included. +# The Jetson SDK checks below gate the whole plugin, tools/ included, when the +# backend is JETSON. # # Prerequisites: # src/plugins/sensing/setup_container.sh # fetches the Multimedia API, # # locates jetson_sipl_api # # Configure with: -# cmake -B build -DBUILD_PLUGIN_SENSING=ON +# cmake -B build -DBUILD_PLUGIN_SENSING=ON # Jetson, or mock elsewhere +# cmake -B build -DBUILD_PLUGIN_SENSING=ON -DSENSING_BACKEND=MOCK # force the mock # ============================================================================== enable_language(CUDA) @@ -31,12 +38,32 @@ endif() find_package(CUDAToolkit REQUIRED) +set(SENSING_BACKEND "AUTO" CACHE STRING "SENSING backend: AUTO, JETSON (SIPL + V4L2 NVENC) or MOCK (synthetic)") +set_property(CACHE SENSING_BACKEND PROPERTY STRINGS AUTO JETSON MOCK) +if(SENSING_BACKEND STREQUAL "AUTO") + # A Jetson that lacks the SDK should still get the setup instructions below, + # not a silent mock, so only non-aarch64 hosts fall back. + if(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|arm64") + set(SENSING_BACKEND_RESOLVED "JETSON" CACHE INTERNAL "") + else() + set(SENSING_BACKEND_RESOLVED "MOCK" CACHE INTERNAL "") + endif() +elseif(SENSING_BACKEND STREQUAL "JETSON" OR SENSING_BACKEND STREQUAL "MOCK") + set(SENSING_BACKEND_RESOLVED "${SENSING_BACKEND}" CACHE INTERNAL "") +else() + message(FATAL_ERROR "SENSING_BACKEND must be AUTO, JETSON or MOCK, got '${SENSING_BACKEND}'") +endif() +# Cache, not a plain variable, so tests/ (a sibling directory) can gate on it too. +message(STATUS "SENSING plugin backend: ${SENSING_BACKEND_RESOLVED}") + if(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|arm64") set(_tegra_lib_dirs "/usr/lib/aarch64-linux-gnu/nvidia" "/usr/lib/aarch64-linux-gnu/tegra") else() set(_tegra_lib_dirs "/usr/lib/x86_64-linux-gnu/nvidia") endif() +if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") +# Jetson-only SDK checks. # ============================================================================== # SIPL # ============================================================================== @@ -159,6 +186,7 @@ message(STATUS "SENSING plugin: V4L2=${V4L2_LIBRARY}") # No libnvjpeg here on purpose: on R39 libnvbufsurface has no libjpeg # dependency at all (readelf -d lists only EGL, cuda, dl, pthread, rt, c), and # the tegra lib dir ships no libnvjpeg to link against. +endif() # ============================================================================== # Targets diff --git a/src/plugins/sensing/README.md b/src/plugins/sensing/README.md index 3c96b01d73..dcc68f423d 100644 --- a/src/plugins/sensing/README.md +++ b/src/plugins/sensing/README.md @@ -24,6 +24,7 @@ That guide owns the operating instructions; this README maps the source tree. - [`docker/`](docker/): example image, image-side group configuration, and host-side container setup. - [`configs/`](configs/): the platform configuration and its provenance. +- [`core/mock_*`](core/): the mock backend (see below). ## Platform configuration provenance @@ -44,3 +45,44 @@ Keep the file byte-identical to the vendor copy. Compare new vendor drops before replacing it; do not hand-edit the board wiring. Its CSI ports, I2C buses, deserializer addresses, GPIOs, and link modes must match the installed SIPL drivers. + +## Mock backend: run the plugin without a Jetson + +`SENSING_BACKEND` selects what sits behind `SiplCamera` and `JetsonEncoder`: + +| Value | Capture | Encoder | Needs | +| --- | --- | --- | --- | +| `JETSON` | SIPL | V4L2 NVENC | L4T SDKs | +| `MOCK` | synthetic frames | stand-in | NVIDIA GPU, CUDA toolkit | +| `AUTO` (default) | `JETSON` on aarch64, `MOCK` elsewhere | | | + +```bash +cmake -B build -DBUILD_PLUGIN_SENSING=ON -DSENSING_BACKEND=MOCK +cmake --build build --target camera_plugin_sensing sensing_mock_tests +./build/src/plugins/sensing/app/camera_plugin_sensing --list-sensors +``` + +Everything above those two classes is the code that runs on the rig: `SensingCamera`, the +sinks, the YUV-to-RGBA and RGBA-to-NV12 kernels, and the CUDA IPC publisher. The mock +resolves `--list-sensors` and the link masks from the vendor JSON, paints a synthetic NV12 +surface (a 16-bit frame counter top-left, chroma tinted per sensor), and runs the real +conversion on it. Both sensors share a capture TSC, as the rig's fsync does. + +The mock encoder's output is **not decodable video**: Annex-B start codes, SPS/PPS/IDR and +non-IDR NAL types, and a payload naming the frame. It reproduces what callers depend on: a +fixed pool of input buffers that makes `submit()` refuse frames when the encoder falls behind, +a warm-up delay before the first unit, and capture stamps that come back at microsecond +resolution. + +Environment, read once at construction: + +| Variable | Effect | +| --- | --- | +| `SENSING_MOCK_FPS` | override the platform config's frame rate | +| `SENSING_MOCK_SCALE_DIV` | divide the reported width and height (kept even); useful for tests | +| `SENSING_MOCK_FAIL_AFTER` | every sensor fails after N frames, as a link drop would | +| `SENSING_MOCK_ENCODER_LATENCY_MS` | per-frame encode time | +| `SENSING_MOCK_ENCODER_DEPTH` | frames held back before the first unit (default 2) | + +A pass on the mock says nothing about SIPL, the ISP output format, NvSci buffer lifetime, the +V4L2 encoder or real frame rates. Those are only checked on an Orin with SENSING cameras. diff --git a/src/plugins/sensing/core/CMakeLists.txt b/src/plugins/sensing/core/CMakeLists.txt index 3b4c856fd1..f303e7c7ea 100644 --- a/src/plugins/sensing/core/CMakeLists.txt +++ b/src/plugins/sensing/core/CMakeLists.txt @@ -1,17 +1,51 @@ # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -# SIPL capture, the V4L2 M2M encoder wrapper, and the output sinks. Everything -# that needs the Jetson SDKs lives here, so tools/ stays buildable without them. -add_library(sensing_core STATIC +# Capture, encoder and output sinks. The JETSON backend needs the Jetson SDKs, +# so everything that does lives in this directory and tools/ stays buildable +# without them. The MOCK backend replaces SiplCamera and JetsonEncoder with +# synthetic ones that need only CUDA. +set(SENSING_COMMON_SOURCES sensing_camera.cpp frame_sink.cpp rawdata_writer.cpp - sipl_camera.cpp yuv_to_rgba.cu - jetson_encoder.cpp rgba_to_nv12.cu - ${JETSON_MMAPI_SOURCES} +) + +if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") + set(SENSING_BACKEND_SOURCES + sipl_camera.cpp + jetson_encoder.cpp + ${JETSON_MMAPI_SOURCES} + ) +else() + # Parses the vendor JSON itself, so a MOCK build has no SIPL query library. + find_package(nlohmann_json 3 QUIET) + if(NOT nlohmann_json_FOUND) + include(FetchContent) + FetchContent_Declare( + nlohmann_json + GIT_REPOSITORY https://github.com/nlohmann/json.git + GIT_TAG v3.11.3 + GIT_SHALLOW TRUE + ) + set(JSON_BuildTests OFF CACHE INTERNAL "") + isaac_teleop_third_party_scope_begin() + FetchContent_MakeAvailable(nlohmann_json) + isaac_teleop_third_party_scope_end() + endif() + set(SENSING_BACKEND_SOURCES + mock_sipl_camera.cpp + mock_platform_config.cpp + mock_nv12_source.cu + mock_jetson_encoder.cpp + ) +endif() + +add_library(sensing_core STATIC + ${SENSING_COMMON_SOURCES} + ${SENSING_BACKEND_SOURCES} ) add_library(sensing::core ALIAS sensing_core) @@ -27,19 +61,6 @@ set_target_properties(sensing_core PROPERTIES target_include_directories(sensing_core PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/.. - PRIVATE - # sipl_camera.hpp pimpls the SIPL and NvSci types away, so these stay - # private: they drag in NvMedia and X11-style macros that collide with - # mcap's StatusCode, which is what sensing_types.hpp exists to avoid. - ${SIPL_INCLUDE_DIRS} - ${EGL_INCLUDE_DIRS} - ${JETSON_MMAPI_DIR}/include -) - -# The Multimedia API reference sources carry legacy warnings. -target_compile_options(sensing_core PRIVATE - $<$:-Wno-reorder -Wno-uninitialized -Wno-unused-variable - -Wno-unused-parameter -Wno-missing-field-initializers -Wno-sign-compare> ) target_link_libraries(sensing_core @@ -48,10 +69,35 @@ target_link_libraries(sensing_core sensing::cuda_ipc isaacteleop_schema PRIVATE - ${SIPL_LIBRARIES} - ${EGL_LIBRARIES} - ${NVBUFSURFACE_LIBRARY} - ${V4L2_LIBRARY} mcap::mcap oxr::oxr_utils ) + +if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") + target_include_directories(sensing_core + PRIVATE + # sipl_camera.hpp pimpls the SIPL and NvSci types away, so these stay + # private: they drag in NvMedia and X11-style macros that collide with + # mcap's StatusCode, which is what sensing_types.hpp exists to avoid. + ${SIPL_INCLUDE_DIRS} + ${EGL_INCLUDE_DIRS} + ${JETSON_MMAPI_DIR}/include + ) + + # The Multimedia API reference sources carry legacy warnings. + target_compile_options(sensing_core PRIVATE + $<$:-Wno-reorder -Wno-uninitialized -Wno-unused-variable + -Wno-unused-parameter -Wno-missing-field-initializers -Wno-sign-compare> + ) + + target_link_libraries(sensing_core + PRIVATE + ${SIPL_LIBRARIES} + ${EGL_LIBRARIES} + ${NVBUFSURFACE_LIBRARY} + ${V4L2_LIBRARY} + ) +else() + target_link_libraries(sensing_core PRIVATE nlohmann_json::nlohmann_json) + target_compile_definitions(sensing_core PUBLIC SENSING_MOCK_BACKEND=1) +endif() diff --git a/src/plugins/sensing/core/mock_jetson_encoder.cpp b/src/plugins/sensing/core/mock_jetson_encoder.cpp new file mode 100644 index 0000000000..76ed50cd54 --- /dev/null +++ b/src/plugins/sensing/core/mock_jetson_encoder.cpp @@ -0,0 +1,296 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Mock backend for JetsonEncoder: no V4L2 and no NVENC. It runs the real +// RGBA-to-NV12 kernel and the same device-to-host copy as the V4L2 path, then +// a worker thread turns each frame into an Annex-B *shaped* unit. +// +// The units are NOT decodable video. Each carries start codes, plausible NAL +// types (SPS/PPS/IDR on the first frame of every GOP, non-IDR slices between) +// and a payload that records the frame, so sinks and tests can follow a frame +// through. What it reproduces from the real encoder is the behaviour callers +// depend on: +// * a fixed pool of input buffers, so submit() returns false when the encoder +// falls behind rather than blocking; +// * warm-up: a unit appears only after more frames have been submitted; +// * the capture stamp comes back at microsecond resolution (a timeval), not +// the nanoseconds it went in with; +// * end_of_stream() returns the remaining tail as one concatenated unit. +// +// Environment (read once, at construction): +// SENSING_MOCK_ENCODER_LATENCY_MS=N per-frame encode time (default 0) +// SENSING_MOCK_ENCODER_DEPTH=N frames held back before output (default 2) + +#include "jetson_encoder.hpp" +#include "rgba_to_nv12.cuh" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace plugins +{ +namespace sensing +{ + +namespace +{ + +constexpr uint32_t kInputBuffers = 6; + +void check_cuda(cudaError_t err, const char* what) +{ + if (err != cudaSuccess) + throw std::runtime_error(std::string("mock JetsonEncoder: ") + what + ": " + cudaGetErrorString(err)); +} + +uint32_t env_uint(const char* name, uint32_t fallback) +{ + const char* value = std::getenv(name); + return value ? static_cast(std::strtoul(value, nullptr, 10)) : fallback; +} + +void append_start_code(std::vector& out) +{ + for (uint8_t byte : { 0, 0, 0, 1 }) + out.push_back(byte); +} + +/// One byte becomes two printable bytes, so a payload never contains the zero +/// bytes that would read as a start code. +void append_hex(std::vector& out, uint32_t value, int nibbles) +{ + for (int shift = (nibbles - 1) * 4; shift >= 0; shift -= 4) + out.push_back(static_cast('A' + ((value >> shift) & 0xF))); +} + +} // namespace + +struct JetsonEncoder::Impl +{ + struct Job + { + uint32_t buffer = 0; + uint32_t frame = 0; + int64_t timestamp_ns = 0; + }; + + explicit Impl(const EncoderConfig& cfg) : config(cfg) + { + gop = config.gop ? config.gop : (config.fps ? config.fps * 5 : 150); + latency = std::chrono::milliseconds(env_uint("SENSING_MOCK_ENCODER_LATENCY_MS", 0)); + depth = env_uint("SENSING_MOCK_ENCODER_DEPTH", 2); + + check_cuda(cudaMallocPitch(reinterpret_cast(&d_y), &d_y_pitch, config.width, config.height), + "cudaMallocPitch(Y)"); + check_cuda(cudaMallocPitch(reinterpret_cast(&d_uv), &d_uv_pitch, config.width, config.height / 2), + "cudaMallocPitch(UV)"); + + // Host staging, standing in for the V4L2 output plane's mapped buffers. + const size_t nv12_bytes = static_cast(config.width) * config.height * 3 / 2; + for (uint32_t i = 0; i < kInputBuffers; ++i) + { + host_buffers.emplace_back(nv12_bytes); + free_buffers.push_back(i); + } + worker = std::thread([this] { run(); }); + } + + ~Impl() + { + { + std::lock_guard lock(mutex); + closing = true; + } + wake.notify_all(); + if (worker.joinable()) + worker.join(); + cudaFree(d_y); + cudaFree(d_uv); + } + + std::vector build_unit(const Job& job, uint32_t checksum) const + { + std::vector unit; + if (job.frame % gop == 0) + { + append_start_code(unit); + unit.push_back(0x67); // SPS + append_hex(unit, config.width, 4); + append_hex(unit, config.height, 4); + append_hex(unit, config.fps, 4); + append_start_code(unit); + unit.push_back(0x68); // PPS + append_hex(unit, 0, 2); + } + const bool idr = job.frame % gop == 0; + append_start_code(unit); + unit.push_back(idr ? 0x65 : 0x41); + for (char c : { 'M', 'O', 'C', 'K' }) + unit.push_back(static_cast(c)); + append_hex(unit, job.frame, 8); + append_hex(unit, checksum, 8); + + // Size follows the bitrate budget (capped), so throughput looks plausible. + const size_t budget = config.bitrate_bps / 8 / std::max(config.fps, 1); + const size_t target = std::min(std::max(budget, unit.size()), 64 * 1024); + unit.resize(target, 'U'); + return unit; + } + + void run() + { + for (;;) + { + Job job; + { + std::unique_lock lock(mutex); + wake.wait(lock, [&] { return closing || !jobs.empty(); }); + if (jobs.empty()) + return; + job = jobs.front(); + jobs.pop_front(); + busy = true; + } + + if (latency.count() > 0) + std::this_thread::sleep_for(latency); + + // Cheap content fingerprint: a sparse sum over the NV12 bytes. + uint32_t checksum = 0; + const auto& nv12 = host_buffers[job.buffer]; + for (size_t i = 0; i < nv12.size(); i += 97) + checksum = checksum * 31u + nv12[i]; + + EncodedUnit unit; + unit.data = build_unit(job, checksum); + // A V4L2 timeval holds microseconds; the stamp does not survive finer. + unit.timestamp_ns = (job.timestamp_ns / 1000) * 1000; + + { + std::lock_guard lock(mutex); + free_buffers.push_back(job.buffer); + held.push_back(std::move(unit)); + while (held.size() > depth) + { + ready.push_back(std::move(held.front())); + held.pop_front(); + } + busy = false; + } + idle.notify_all(); + } + } + + EncoderConfig config; + uint32_t gop = 150; + uint32_t depth = 2; + std::chrono::milliseconds latency{ 0 }; + + uint8_t* d_y = nullptr; + uint8_t* d_uv = nullptr; + size_t d_y_pitch = 0; + size_t d_uv_pitch = 0; + + std::vector> host_buffers; + uint32_t submitted = 0; + + std::mutex mutex; + std::condition_variable wake; + std::condition_variable idle; + std::deque free_buffers; + std::deque jobs; + std::deque held; // encoded, still inside the warm-up window + std::deque ready; // what poll() hands out + bool busy = false; + bool closing = false; + std::thread worker; +}; + +JetsonEncoder::JetsonEncoder(const EncoderConfig& config) +{ + if (config.width == 0 || config.height == 0) + throw std::runtime_error("JetsonEncoder: width/height must be non-zero"); + if ((config.width % 2) != 0 || (config.height % 2) != 0) + throw std::runtime_error("JetsonEncoder: width/height must be even for NV12"); + m_impl = std::make_unique(config); +} + +JetsonEncoder::~JetsonEncoder() = default; + +bool JetsonEncoder::submit(uintptr_t rgba_device_ptr, std::size_t row_pitch_bytes, int64_t timestamp_ns) +{ + auto& impl = *m_impl; + const auto& config = impl.config; + + uint32_t buffer = 0; + { + std::lock_guard lock(impl.mutex); + if (impl.free_buffers.empty()) + return false; // encoder still busy; drop this frame rather than block + buffer = impl.free_buffers.front(); + impl.free_buffers.pop_front(); + } + + launch_rgba_to_nv12(reinterpret_cast(rgba_device_ptr), static_cast(row_pitch_bytes), impl.d_y, + static_cast(impl.d_y_pitch), impl.d_uv, static_cast(impl.d_uv_pitch), + static_cast(config.width), static_cast(config.height), config.full_range, nullptr); + check_cuda(cudaGetLastError(), "rgba_to_nv12 launch"); + check_cuda(cudaStreamSynchronize(nullptr), "rgba_to_nv12 sync"); + + uint8_t* host = impl.host_buffers[buffer].data(); + check_cuda( + cudaMemcpy2D(host, config.width, impl.d_y, impl.d_y_pitch, config.width, config.height, cudaMemcpyDeviceToHost), + "cudaMemcpy2D(Y)"); + check_cuda(cudaMemcpy2D(host + static_cast(config.width) * config.height, config.width, impl.d_uv, + impl.d_uv_pitch, config.width, config.height / 2, cudaMemcpyDeviceToHost), + "cudaMemcpy2D(UV)"); + + { + std::lock_guard lock(impl.mutex); + impl.jobs.push_back({ buffer, impl.submitted++, timestamp_ns }); + } + impl.wake.notify_one(); + return true; +} + +EncodedUnit JetsonEncoder::poll() +{ + std::lock_guard lock(m_impl->mutex); + if (m_impl->ready.empty()) + return {}; + auto out = std::move(m_impl->ready.front()); + m_impl->ready.pop_front(); + return out; +} + +EncodedUnit JetsonEncoder::end_of_stream() +{ + auto& impl = *m_impl; + std::unique_lock lock(impl.mutex); + impl.idle.wait(lock, [&] { return impl.jobs.empty() && !impl.busy; }); + + // The warm-up window empties at EOS, then the tail is concatenated, so it + // carries the stamp of the last unit in it. + EncodedUnit out; + for (auto* queue : { &impl.ready, &impl.held }) + { + for (auto& chunk : *queue) + { + out.data.insert(out.data.end(), chunk.data.begin(), chunk.data.end()); + out.timestamp_ns = chunk.timestamp_ns; + } + queue->clear(); + } + return out; +} + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_nv12_source.cu b/src/plugins/sensing/core/mock_nv12_source.cu new file mode 100644 index 0000000000..81d757083c --- /dev/null +++ b/src/plugins/sensing/core/mock_nv12_source.cu @@ -0,0 +1,83 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "mock_nv12_source.cuh" + +namespace plugins +{ +namespace sensing +{ + +namespace +{ + +constexpr int kCounterCell = 16; +constexpr int kCounterBits = 16; + +__global__ void paint_luma(uint8_t* y_plane, int y_pitch, int width, int height, uint32_t frame_index) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) + return; + + const float u = static_cast(x) / width; + const float v = static_cast(y) / height; + const float t = static_cast(frame_index % 240u) / 240.0f; + + // Diagonal gradient scrolling with the frame index. + float luma = 255.0f * (0.25f + 0.5f * fabsf(__sinf(6.2831f * (u + v + t)))); + + // A bright square orbiting the centre: a repeated frame freezes it. + const float cx = 0.5f + 0.3f * __cosf(6.2831f * t); + const float cy = 0.5f + 0.3f * __sinf(6.2831f * t); + if (fabsf(u - cx) < 0.04f && fabsf(v - cy) < 0.04f * width / height) + luma = 255.0f; + + if (y < kCounterCell && x < kCounterBits * kCounterCell) + { + const unsigned bit = static_cast(x / kCounterCell); + luma = ((frame_index >> (kCounterBits - 1u - bit)) & 1u) ? 255.0f : 0.0f; + } + + y_plane[static_cast(y) * y_pitch + x] = static_cast(luma); +} + +__global__ void paint_chroma(uint8_t* uv_plane, int uv_pitch, int width, int height, uint32_t sensor_id) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; // chroma sample column + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width / 2 || y >= height / 2) + return; + + // Even sensors lean blue-green, odd sensors red-orange. The counter strip stays neutral so its + // cells decode as pure black and white. + const bool odd = (sensor_id & 1u) != 0; + const bool counter = y < kCounterCell / 2 && x < kCounterBits * kCounterCell / 2; + uint8_t* sample = uv_plane + static_cast(y) * uv_pitch + static_cast(x) * 2; + sample[0] = counter ? 128 : (odd ? 100 : 160); // U + sample[1] = counter ? 128 : (odd ? 180 : 100); // V +} + +} // namespace + +void launch_mock_nv12_pattern(uint8_t* y_plane, + int y_pitch, + uint8_t* uv_plane, + int uv_pitch, + int width, + int height, + uint32_t sensor_id, + uint32_t frame_index, + cudaStream_t stream) +{ + const dim3 block(16, 16); + const dim3 luma_grid((width + block.x - 1) / block.x, (height + block.y - 1) / block.y); + paint_luma<<>>(y_plane, y_pitch, width, height, frame_index); + + const dim3 chroma_grid((width / 2 + block.x - 1) / block.x, (height / 2 + block.y - 1) / block.y); + paint_chroma<<>>(uv_plane, uv_pitch, width, height, sensor_id); +} + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_nv12_source.cuh b/src/plugins/sensing/core/mock_nv12_source.cuh new file mode 100644 index 0000000000..5db70f07ec --- /dev/null +++ b/src/plugins/sensing/core/mock_nv12_source.cuh @@ -0,0 +1,34 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +namespace plugins +{ +namespace sensing +{ + +/** + * @brief Paint one synthetic ISP0 frame: full-range NV12 in pitch-linear device memory. + * + * Stands in for the buffer SIPL would fill, so the real YUV-to-RGBA kernel runs + * on it unchanged. A 16-bit frame counter is drawn top-left as 16 px cells (white + * for a set bit, black for clear, most significant bit first), so a test can tell + * a fresh frame from a stale or torn one by reading pixels. Chroma is tinted by + * sensor so the two eyes are distinguishable. + */ +void launch_mock_nv12_pattern(uint8_t* y_plane, + int y_pitch, + uint8_t* uv_plane, + int uv_pitch, + int width, + int height, + uint32_t sensor_id, + uint32_t frame_index, + cudaStream_t stream); + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_platform_config.cpp b/src/plugins/sensing/core/mock_platform_config.cpp new file mode 100644 index 0000000000..c4305b66b0 --- /dev/null +++ b/src/plugins/sensing/core/mock_platform_config.cpp @@ -0,0 +1,107 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "mock_platform_config.hpp" + +#include + +#include +#include +#include +#include + +namespace plugins +{ +namespace sensing +{ + +namespace +{ + +uint32_t scale_divisor() +{ + const char* value = std::getenv("SENSING_MOCK_SCALE_DIV"); + if (!value) + return 1; + const long parsed = std::strtol(value, nullptr, 10); + return parsed > 1 ? static_cast(parsed) : 1; +} + +uint32_t scaled(uint32_t extent, uint32_t divisor) +{ + return std::max(2, (extent / divisor) & ~1u); +} + +} // namespace + +std::vector mock_query_platform_config(const std::string& platform_config_json, + const std::string& platform_config_name, + const std::vector& link_masks) +{ + std::ifstream file(platform_config_json); + if (!file) + throw std::runtime_error("ParseJsonFile failed: cannot open " + platform_config_json); + + nlohmann::json root; + try + { + root = nlohmann::json::parse(file); + } + catch (const nlohmann::json::exception& e) + { + throw std::runtime_error("ParseJsonFile failed: " + std::string(e.what())); + } + + const auto& configs = root.at("cameraConfigs"); + auto find_config = [&](const std::string& name) -> const nlohmann::json* + { + for (const auto& config : configs) + { + if (config.value("name", "") == name) + return &config; + } + return nullptr; + }; + + const auto* platform = find_config(platform_config_name); + if (!platform) + throw std::runtime_error("GetSensorSystemConfig failed: no camera config named '" + platform_config_name + "'"); + + const uint32_t divisor = scale_divisor(); + std::vector sensors; + for (const auto& module : platform->at("cameraModules")) + { + const uint32_t link = module.at("linkIndex"); + const uint32_t transport = module.value("transportIndex", 0u); + if (!link_masks.empty()) + { + const bool enabled = transport < link_masks.size() && ((link_masks[transport] >> (4 * link)) & 0xF) != 0; + if (!enabled) + continue; + } + + const std::string module_name = module.at("name"); + const auto* module_config = find_config(module_name); + if (!module_config) + throw std::runtime_error("GetSensorSystemConfig failed: no camera config named '" + module_name + "'"); + + for (const auto& sensor : module_config->at("sensorInfo")) + { + SensorInfo info; + info.id = static_cast(sensors.size()); + info.name = module_name; + const auto& channels = sensor.at("virtualChannels"); + if (!channels.empty()) + { + info.width = scaled(channels.front().at("width"), divisor); + info.height = scaled(channels.front().at("height"), divisor); + info.fps = channels.front().at("fps"); + } + sensors.push_back(std::move(info)); + } + } + return sensors; +} + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_platform_config.hpp b/src/plugins/sensing/core/mock_platform_config.hpp new file mode 100644 index 0000000000..5c0ac44d52 --- /dev/null +++ b/src/plugins/sensing/core/mock_platform_config.hpp @@ -0,0 +1,40 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Mock backend only. Answers SiplCamera::query from the vendor platform JSON +// itself, so the pipeline-index renumbering and link-mask behaviour can be +// exercised with no SIPL library. + +#pragma once + +#include "sipl_camera.hpp" + +#include +#include +#include + +namespace plugins +{ +namespace sensing +{ + +/** + * @brief Resolve a platform config to its sensors, as INvSIPLCameraQuery would. + * + * Module `m` of the named config is enabled when its transport's mask has a + * non-zero hex digit at the module's link index, i.e. + * `(masks[transportIndex] >> (4 * linkIndex)) & 0xF`. An empty mask list keeps + * every module. `SensorInfo::id` is the position among the enabled sensors, not + * the JSON sensorInfo.id. + * + * SENSING_MOCK_SCALE_DIV=N divides the reported width and height (kept even), + * which keeps tests fast without touching the vendor JSON. + * + * @throws std::runtime_error when the file or the named config is missing. + */ +std::vector mock_query_platform_config(const std::string& platform_config_json, + const std::string& platform_config_name, + const std::vector& link_masks); + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_sipl_camera.cpp b/src/plugins/sensing/core/mock_sipl_camera.cpp new file mode 100644 index 0000000000..1150802472 --- /dev/null +++ b/src/plugins/sensing/core/mock_sipl_camera.cpp @@ -0,0 +1,334 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Mock backend for SiplCamera: synthetic frames from a CUDA pattern generator, +// no SIPL, NvSci or EGL. It keeps the contract the rest of the plugin relies on +// -- per-sensor capture threads, a three-slot RGBA mailbox with a reader lease, +// a TSC stamp shared by every sensor, and failures surfacing from latest() -- +// and runs the real YUV-to-RGBA kernel on the synthetic NV12, so the conversion +// code is exercised too. +// +// Environment (read once, at construction): +// SENSING_MOCK_FPS=N override the platform config's frame rate +// SENSING_MOCK_FAIL_AFTER=N every sensor fails after N frames, as a link drop would +// SENSING_MOCK_SCALE_DIV=N shrink the reported geometry (see mock_platform_config.hpp) + +#include "mock_nv12_source.cuh" +#include "mock_platform_config.hpp" +#include "sipl_camera.hpp" +#include "yuv_to_rgba.cuh" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace plugins +{ +namespace sensing +{ + +namespace +{ + +// Arbitrary offset so nothing can mistake the TSC for CLOCK_MONOTONIC. +constexpr uint64_t kMockTscBaseNs = 1'000'000'000'000'000ull; + +uint64_t monotonic_ns() +{ + return static_cast( + std::chrono::duration_cast(std::chrono::steady_clock::now().time_since_epoch()).count()); +} + +double env_number(const char* name, double fallback) +{ + const char* value = std::getenv(name); + return value ? std::strtod(value, nullptr) : fallback; +} + +void check_cuda(cudaError_t err, const char* what) +{ + if (err != cudaSuccess) + throw std::runtime_error(std::string("mock SIPL: ") + what + ": " + cudaGetErrorString(err)); +} + +} // namespace + +struct SiplCamera::Impl +{ + struct DeviceBuffer + { + uint8_t* ptr = nullptr; + size_t pitch = 0; + }; + + struct Sensor + { + SensorInfo info; + /// Synthetic ISP0 surface, in place of the NvSciBuf SIPL would fill. + DeviceBuffer luma; + DeviceBuffer chroma; + cudaStream_t stream = nullptr; + + /// Triple-buffered RGBA8 mailbox. A reader leases the published slot + /// until its next latest() call, so the producer never overwrites it. + std::array buffers{}; + mutable std::mutex publish_mutex; + int publish_idx = -1; + int lease_idx = -1; + uint64_t published_sequence = 0; + uint64_t consumed_sequence = 0; + uint64_t published_timestamp_ns = 0; + uint64_t published_capture_tsc_ns = 0; + + std::thread frame_thread; + }; + + explicit Impl(const SiplConfig& cfg) : config(cfg) + { + if (config.link_masks.empty()) + throw std::runtime_error("link masks are required for a GMSL platform config"); + + sensor_infos = + mock_query_platform_config(config.platform_config_json, config.platform_config_name, config.link_masks); + if (sensor_infos.empty()) + { + throw std::runtime_error("no modules left in '" + config.platform_config_name + + "' after applying the link masks -- check the mask against the config"); + } + + fps_override = env_number("SENSING_MOCK_FPS", 0.0); + fail_after_frames = static_cast(env_number("SENSING_MOCK_FAIL_AFTER", 0.0)); + + check_cuda(cudaSetDevice(config.gpu_id), "cudaSetDevice"); + for (const auto& info : sensor_infos) + { + auto sensor = std::make_unique(); + sensor->info = info; + allocate(*sensor); + sensors.push_back(std::move(sensor)); + } + } + + ~Impl() + { + stop(); + for (auto& sensor : sensors) + { + cudaFree(sensor->luma.ptr); + cudaFree(sensor->chroma.ptr); + for (auto& buffer : sensor->buffers) + cudaFree(buffer.ptr); + if (sensor->stream) + cudaStreamDestroy(sensor->stream); + } + } + + void allocate(Sensor& sensor) + { + const size_t width = sensor.info.width; + const size_t height = sensor.info.height; + check_cuda(cudaMallocPitch(reinterpret_cast(&sensor.luma.ptr), &sensor.luma.pitch, width, height), + "cudaMallocPitch(Y)"); + check_cuda(cudaMallocPitch(reinterpret_cast(&sensor.chroma.ptr), &sensor.chroma.pitch, width, height / 2), + "cudaMallocPitch(UV)"); + for (auto& buffer : sensor.buffers) + { + check_cuda(cudaMallocPitch(reinterpret_cast(&buffer.ptr), &buffer.pitch, width * 4, height), + "cudaMallocPitch(RGBA)"); + } + check_cuda(cudaStreamCreateWithFlags(&sensor.stream, cudaStreamNonBlocking), "cudaStreamCreate"); + } + + Sensor* find(uint32_t id) + { + for (auto& sensor : sensors) + { + if (sensor->info.id == id) + return sensor.get(); + } + return nullptr; + } + + void set_failure(const std::string& message) + { + std::lock_guard guard(failure_mutex); + if (failure.empty()) + failure = message; + } + + void throw_if_failed() + { + std::lock_guard guard(failure_mutex); + if (!failure.empty()) + throw std::runtime_error(failure); + } + + uint32_t pick_write_index(const Sensor& sensor) const + { + std::lock_guard guard(sensor.publish_mutex); + if (sensor.publish_idx < 0) + return 0; + // Skip the published slot and any slot still leased to a reader. With + // three slots one is always free, so this cannot spin. + uint32_t idx = static_cast((sensor.publish_idx + 1) % 3); + if (static_cast(idx) == sensor.lease_idx) + idx = (idx + 1) % 3; + return idx; + } + + void publish(Sensor& sensor, uint32_t write_idx, uint64_t timestamp_ns, uint64_t capture_tsc_ns) + { + std::lock_guard guard(sensor.publish_mutex); + sensor.publish_idx = static_cast(write_idx); + sensor.published_timestamp_ns = timestamp_ns; + sensor.published_capture_tsc_ns = capture_tsc_ns; + ++sensor.published_sequence; + } + + void capture_one(Sensor& sensor, uint64_t frame_index, uint64_t tsc_ns) + { + const uint32_t write_idx = pick_write_index(sensor); + const DeviceBuffer& dest = sensor.buffers[write_idx]; + const int width = static_cast(sensor.info.width); + const int height = static_cast(sensor.info.height); + + launch_mock_nv12_pattern(sensor.luma.ptr, static_cast(sensor.luma.pitch), sensor.chroma.ptr, + static_cast(sensor.chroma.pitch), width, height, sensor.info.id, + static_cast(frame_index), sensor.stream); + + // Same call, layout and colour matrix as the SIPL path: full-range + // BT.709 semi-planar, from the surface the ISP would have written. + const YuvLayout layout = config.swap_uv ? YuvLayout::YVU420SemiPlanar : YuvLayout::YUV420SemiPlanar; + launch_yuv420_pitch_to_rgba(sensor.luma.ptr, sensor.chroma.ptr, nullptr, static_cast(sensor.luma.pitch), + static_cast(sensor.chroma.pitch), 0, width, height, dest.ptr, + static_cast(dest.pitch), layout, true, true, sensor.stream); + check_cuda(cudaGetLastError(), "mock frame kernels"); + check_cuda(cudaStreamSynchronize(sensor.stream), "cudaStreamSynchronize"); + + publish(sensor, write_idx, monotonic_ns(), tsc_ns); + } + + void frame_loop(Sensor& sensor, std::chrono::steady_clock::time_point epoch) + { + try + { + check_cuda(cudaSetDevice(config.gpu_id), "cudaSetDevice frame_loop"); + const double fps = fps_override > 0.0 ? fps_override : std::max(sensor.info.fps, 1.0); + const double period_ns = 1e9 / fps; + for (uint64_t n = 0; running.load(); ++n) + { + std::this_thread::sleep_until(epoch + std::chrono::nanoseconds(static_cast(n * period_ns))); + if (!running.load()) + break; + if (fail_after_frames != 0 && n >= fail_after_frames) + throw std::runtime_error("injected link failure (SENSING_MOCK_FAIL_AFTER)"); + // Every sensor derives its TSC from the same epoch, as the + // deserializer's shared fsync does: frame n of each sensor + // carries the same capture time. + capture_one(sensor, n, kMockTscBaseNs + static_cast(n * period_ns)); + } + } + catch (const std::exception& e) + { + std::ostringstream oss; + oss << "SIPL sensor " << sensor.info.id << " capture error: " << e.what(); + std::cerr << "[sipl-mock] " << oss.str() << std::endl; + set_failure(oss.str()); + running.store(false); + } + } + + void start() + { + if (running.exchange(true)) + return; + const auto epoch = std::chrono::steady_clock::now(); + for (auto& sensor : sensors) + sensor->frame_thread = std::thread([this, s = sensor.get(), epoch] { frame_loop(*s, epoch); }); + } + + void stop() + { + running.store(false); + for (auto& sensor : sensors) + { + if (sensor->frame_thread.joinable()) + sensor->frame_thread.join(); + } + } + + SiplConfig config; + std::vector sensor_infos; + std::vector> sensors; + std::atomic running{ false }; + double fps_override = 0.0; + uint64_t fail_after_frames = 0; + std::mutex failure_mutex; + std::string failure; +}; + +SiplCamera::SiplCamera(const SiplConfig& config) : m_impl(std::make_unique(config)) +{ +} + +SiplCamera::~SiplCamera() = default; + +std::vector SiplCamera::query(const std::string& platform_config_json, + const std::string& platform_config_name, + const std::vector& link_masks) +{ + return mock_query_platform_config(platform_config_json, platform_config_name, link_masks); +} + +const std::vector& SiplCamera::sensors() const +{ + return m_impl->sensor_infos; +} + +void SiplCamera::start() +{ + m_impl->start(); +} + +void SiplCamera::stop() +{ + m_impl->stop(); +} + +std::optional SiplCamera::latest(uint32_t sensor_id) +{ + m_impl->throw_if_failed(); + Impl::Sensor* sensor = m_impl->find(sensor_id); + if (!sensor) + throw std::runtime_error("no SIPL pipeline with index " + std::to_string(sensor_id)); + + std::lock_guard guard(sensor->publish_mutex); + if (sensor->publish_idx < 0 || sensor->consumed_sequence == sensor->published_sequence) + return std::nullopt; + sensor->consumed_sequence = sensor->published_sequence; + + // The caller now holds this slot; pick_write_index() skips it until the + // next latest() moves the lease on. + sensor->lease_idx = sensor->publish_idx; + + FrameView view; + view.ptr = reinterpret_cast(sensor->buffers[sensor->publish_idx].ptr); + view.pitch = sensor->buffers[sensor->publish_idx].pitch; + view.width = sensor->info.width; + view.height = sensor->info.height; + view.timestamp_ns = sensor->published_timestamp_ns; + view.capture_tsc_ns = sensor->published_capture_tsc_ns; + view.sequence = sensor->published_sequence; + return view; +} + +} // namespace sensing +} // namespace plugins diff --git a/tests/cpp/plugins/CMakeLists.txt b/tests/cpp/plugins/CMakeLists.txt index 7ae2c32d67..b3d5925a3a 100644 --- a/tests/cpp/plugins/CMakeLists.txt +++ b/tests/cpp/plugins/CMakeLists.txt @@ -16,3 +16,7 @@ endif() if(BUILD_PLUGIN_SENSING) add_subdirectory(sensing) endif() +# Needs the MOCK backend; the JETSON backend is exercised on a rig, not here. +if(BUILD_PLUGIN_SENSING AND SENSING_BACKEND_RESOLVED STREQUAL "MOCK") + add_subdirectory(sensing_mock) +endif() diff --git a/tests/cpp/plugins/sensing_mock/CMakeLists.txt b/tests/cpp/plugins/sensing_mock/CMakeLists.txt new file mode 100644 index 0000000000..fbca027d6a --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/CMakeLists.txt @@ -0,0 +1,29 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# Exercises the plugin on the MOCK backend: synthetic frames and a stand-in +# encoder, but the real SensingCamera, sinks, CUDA kernels and CUDA IPC. Needs an +# NVIDIA GPU, not a Jetson. +find_package(CUDAToolkit REQUIRED) + +add_executable(sensing_mock_tests + mock_test_support.hpp + test_mock_camera.cpp + test_mock_encoder.cpp + test_mock_pipeline.cpp + test_cuda_ipc_consumer.cpp +) + +target_link_libraries(sensing_mock_tests + PRIVATE + Catch2::Catch2WithMain + sensing::core + CUDA::cudart + CUDA::cuda_driver +) + +target_compile_definitions(sensing_mock_tests PRIVATE + SENSING_PLATFORM_CONFIG="${CMAKE_SOURCE_DIR}/src/plugins/sensing/configs/shw5g.json" +) + +catch_discover_tests(sensing_mock_tests ADD_TAGS_AS_LABELS) diff --git a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp new file mode 100644 index 0000000000..6038dd7c33 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -0,0 +1,146 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace sensing_test +{ + +/// Sets an environment variable for one scope; the mock reads its knobs at construction. +class ScopedEnv +{ +public: + ScopedEnv(const char* name, const char* value) : m_name(name) + { + ::setenv(name, value, 1); + } + ~ScopedEnv() + { + ::unsetenv(m_name.c_str()); + } + ScopedEnv(const ScopedEnv&) = delete; + ScopedEnv& operator=(const ScopedEnv&) = delete; + +private: + std::string m_name; +}; + +template +bool wait_for(Pred done, std::chrono::milliseconds limit = std::chrono::seconds(5)) +{ + const auto deadline = std::chrono::steady_clock::now() + limit; + while (!done()) + { + if (std::chrono::steady_clock::now() >= deadline) + return false; + std::this_thread::sleep_for(std::chrono::milliseconds(2)); + } + return true; +} + +inline plugins::sensing::SiplConfig make_sipl_config() +{ + plugins::sensing::SiplConfig config; + config.platform_config_json = SENSING_PLATFORM_CONFIG; + config.platform_config_name = "SHW5G_2"; + config.link_masks = { 0x0000, 0x1100 }; + return config; +} + +inline std::string unique_path(const std::string& stem, const std::string& extension) +{ + static int counter = 0; + return "/tmp/" + stem + "_" + std::to_string(::getpid()) + "_" + std::to_string(counter++) + extension; +} + +/// Copies a device RGBA frame to the host as tightly packed rows. +inline std::vector download_rgba(uintptr_t device_ptr, size_t pitch, uint32_t width, uint32_t height) +{ + std::vector host(static_cast(width) * height * 4); + const cudaError_t err = + cudaMemcpy2D(host.data(), static_cast(width) * 4, reinterpret_cast(device_ptr), pitch, + static_cast(width) * 4, height, cudaMemcpyDeviceToHost); + if (err != cudaSuccess) + throw std::runtime_error(std::string("cudaMemcpy2D: ") + cudaGetErrorString(err)); + return host; +} + +/// Reads the 16-bit frame counter the mock paints top-left: 16 px cells, most significant bit first. +inline uint32_t read_counter(const std::vector& rgba, uint32_t width) +{ + uint32_t value = 0; + for (int bit = 0; bit < 16; ++bit) + { + const size_t x = 8 + 16 * static_cast(bit); + const size_t y = 8; + const uint8_t* px = &rgba[(y * width + x) * 4]; + const bool on = px[0] > 127 && px[1] > 127 && px[2] > 127; + value = (value << 1) | (on ? 1u : 0u); + } + return value; +} + +/// One NAL unit of the mock encoder's Annex-B output. +struct Nal +{ + uint8_t type = 0; + std::vector payload; +}; + +inline std::vector split_annex_b(const std::vector& data) +{ + std::vector nals; + for (size_t i = 0; i + 4 < data.size();) + { + if (data[i] == 0 && data[i + 1] == 0 && data[i + 2] == 0 && data[i + 3] == 1) + { + Nal nal; + nal.type = data[i + 4]; + size_t end = i + 5; + while (end + 3 < data.size() && + !(data[end] == 0 && data[end + 1] == 0 && data[end + 2] == 0 && data[end + 3] == 1)) + ++end; + if (end + 3 >= data.size()) + end = data.size(); + nal.payload.assign(data.begin() + i + 5, data.begin() + end); + nals.push_back(std::move(nal)); + i = end; + } + else + { + ++i; + } + } + return nals; +} + +/// The frame number the mock records after the "MOCK" magic, as 8 hex-like letters. +inline uint32_t frame_of(const Nal& slice) +{ + REQUIRE(slice.payload.size() >= 12); + REQUIRE(std::string(slice.payload.begin(), slice.payload.begin() + 4) == "MOCK"); + uint32_t value = 0; + for (int i = 0; i < 8; ++i) + value = (value << 4) | static_cast(slice.payload[4 + i] - 'A'); + return value; +} + +inline bool is_slice(const Nal& nal) +{ + return nal.type == 0x65 || nal.type == 0x41; +} + +} // namespace sensing_test diff --git a/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp b/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp new file mode 100644 index 0000000000..6ee83488eb --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp @@ -0,0 +1,285 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Acts as a CUDA IPC consumer of the plugin: connects to its socket, receives the +// memory handle as a file descriptor, maps the shared allocation through the CUDA +// virtual memory API, and reads the frames it publishes. This is the path any +// out-of-process consumer, such as a JPEG-encoding server, would take. + +#include "mock_test_support.hpp" + +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +using namespace plugins::sensing; +using namespace sensing_test; +using namespace std::chrono_literals; + +namespace +{ + +void check_cu(CUresult result, const char* what) +{ + if (result != CUDA_SUCCESS) + { + const char* name = nullptr; + cuGetErrorName(result, &name); + throw std::runtime_error(std::string(what) + " failed: " + (name ? name : "unknown")); + } +} + +class IpcConsumer +{ +public: + explicit IpcConsumer(const std::string& socket_path) + { + m_socket = ::socket(AF_UNIX, SOCK_STREAM, 0); + REQUIRE(m_socket >= 0); + timeval timeout{ 2, 0 }; + ::setsockopt(m_socket, SOL_SOCKET, SO_RCVTIMEO, &timeout, sizeof(timeout)); + + sockaddr_un address{}; + address.sun_family = AF_UNIX; + std::snprintf(address.sun_path, sizeof(address.sun_path), "%s", socket_path.c_str()); + REQUIRE(::connect(m_socket, reinterpret_cast(&address), sizeof(address)) == 0); + + int fd = receive_hello(); + import_and_map(fd); + ::close(fd); // the import keeps what it needs + } + + ~IpcConsumer() + { + if (m_ptr) + { + cuMemUnmap(m_ptr, m_hello.total_bytes); + cuMemAddressFree(m_ptr, m_hello.total_bytes); + } + if (m_handle) + cuMemRelease(m_handle); + if (m_socket >= 0) + ::close(m_socket); + } + + IpcConsumer(const IpcConsumer&) = delete; + IpcConsumer& operator=(const IpcConsumer&) = delete; + + const ipc::Hello& hello() const + { + return m_hello; + } + + /// Waits up to `limit` for the next published frame. + bool next_frame(ipc::FrameReady& out, std::chrono::milliseconds limit) + { + pollfd pfd{ m_socket, POLLIN, 0 }; + if (::poll(&pfd, 1, static_cast(limit.count())) <= 0) + return false; + const ssize_t got = ::recv(m_socket, &out, sizeof(out), MSG_WAITALL); + REQUIRE(got == static_cast(sizeof(out))); + REQUIRE(out.magic == ipc::kFrameMagic); + REQUIRE(out.slot < m_hello.slot_count); + return true; + } + + std::vector read_slot(uint32_t slot) const + { + std::vector host(static_cast(m_hello.pitch) * m_hello.height); + check_cu(cuMemcpyDtoH(host.data(), m_ptr + slot * m_hello.slot_stride, host.size()), "cuMemcpyDtoH"); + return host; + } + + void release(uint32_t slot) + { + ipc::SlotRelease message{ ipc::kReleaseMagic, slot }; + REQUIRE(::send(m_socket, &message, sizeof(message), MSG_NOSIGNAL) == static_cast(sizeof(message))); + } + +private: + int receive_hello() + { + char control[CMSG_SPACE(sizeof(int))] = {}; + iovec iov{ &m_hello, sizeof(m_hello) }; + msghdr msg{}; + msg.msg_iov = &iov; + msg.msg_iovlen = 1; + msg.msg_control = control; + msg.msg_controllen = sizeof(control); + REQUIRE(::recvmsg(m_socket, &msg, 0) == static_cast(sizeof(m_hello))); + + cmsghdr* header = CMSG_FIRSTHDR(&msg); + REQUIRE(header != nullptr); + REQUIRE(header->cmsg_type == SCM_RIGHTS); + int fd = -1; + std::memcpy(&fd, CMSG_DATA(header), sizeof(fd)); + return fd; + } + + void import_and_map(int fd) + { + REQUIRE(m_hello.magic == ipc::kHelloMagic); + REQUIRE(m_hello.version == ipc::kProtocolVersion); + + CUdevice device = 0; + CUcontext context = nullptr; + check_cu(cuInit(0), "cuInit"); + check_cu(cuDeviceGet(&device, static_cast(m_hello.device_id)), "cuDeviceGet"); + check_cu(cuDevicePrimaryCtxRetain(&context, device), "cuDevicePrimaryCtxRetain"); + check_cu(cuCtxSetCurrent(context), "cuCtxSetCurrent"); + + check_cu(cuMemImportFromShareableHandle(&m_handle, reinterpret_cast(static_cast(fd)), + CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR), + "cuMemImportFromShareableHandle"); + check_cu(cuMemAddressReserve(&m_ptr, m_hello.total_bytes, 0, 0, 0), "cuMemAddressReserve"); + check_cu(cuMemMap(m_ptr, m_hello.total_bytes, 0, m_handle, 0), "cuMemMap"); + + CUmemAccessDesc access{}; + access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + access.location.id = static_cast(m_hello.device_id); + access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; + check_cu(cuMemSetAccess(m_ptr, m_hello.total_bytes, &access, 1), "cuMemSetAccess"); + } + + int m_socket = -1; + ipc::Hello m_hello{}; + CUmemGenericAllocationHandle m_handle = 0; + CUdeviceptr m_ptr = 0; +}; + +/// Runs SensingCamera::update() on its own thread, as the plugin's main loop would. +class UpdateLoop +{ +public: + explicit UpdateLoop(SensingCamera& camera) + : m_thread( + [this, &camera] + { + while (!m_stop.load()) + { + camera.update(); + ++m_ticks; + std::this_thread::sleep_for(1ms); + } + }) + { + } + + ~UpdateLoop() + { + m_stop.store(true); + m_thread.join(); + } + + uint64_t ticks() const + { + return m_ticks.load(); + } + +private: + std::atomic m_stop{ false }; + std::atomic m_ticks{ 0 }; + std::thread m_thread; +}; + +SensingConfig make_config() +{ + SensingConfig config; + config.platform_config_json = SENSING_PLATFORM_CONFIG; + return config; +} + +} // namespace + +TEST_CASE("a consumer receives a described, mappable allocation", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + std::vector streams(1); + streams[0].sensor_id = 1; + streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); + + SensingCamera camera(make_config(), streams, std::make_unique(streams)); + UpdateLoop loop(camera); + IpcConsumer consumer(streams[0].ipc_socket_path); + + const auto& hello = consumer.hello(); + CHECK(hello.width == 320); + CHECK(hello.height == 248); + CHECK(hello.format == static_cast(ipc::PixelFormat::Rgba8)); + CHECK(hello.sensor_id == 1); + CHECK(hello.slot_count >= 3); + CHECK(hello.pitch >= static_cast(hello.width) * 4); + CHECK(hello.slot_stride >= hello.pitch * hello.height); + CHECK(hello.total_bytes >= hello.slot_stride * hello.slot_count); +} + +TEST_CASE("published frames are fresh pixels in order", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + std::vector streams(1); + streams[0].sensor_id = 0; + streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); + + SensingCamera camera(make_config(), streams, std::make_unique(streams)); + UpdateLoop loop(camera); + IpcConsumer consumer(streams[0].ipc_socket_path); + + uint64_t last_sequence = 0; + uint64_t last_stamp = 0; + int64_t last_counter = -1; + for (int i = 0; i < 20; ++i) + { + ipc::FrameReady frame{}; + REQUIRE(consumer.next_frame(frame, 2000ms)); + const auto pixels = consumer.read_slot(frame.slot); + const int64_t counter = read_counter(pixels, consumer.hello().width); + consumer.release(frame.slot); + + CHECK(frame.sequence > last_sequence); + CHECK(frame.timestamp_ns > last_stamp); + // Each published frame is a new capture, never a repeat of the previous one. + CHECK(counter > last_counter); + last_sequence = frame.sequence; + last_stamp = frame.timestamp_ns; + last_counter = counter; + } +} + +TEST_CASE("a consumer that holds its slots cannot stall the producer", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + std::vector streams(1); + streams[0].sensor_id = 0; + streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); + + SensingCamera camera(make_config(), streams, std::make_unique(streams)); + UpdateLoop loop(camera); + IpcConsumer consumer(streams[0].ipc_socket_path); + + // Take frames without releasing any: the producer can publish only as many as there are slots. + std::vector held; + ipc::FrameReady frame{}; + while (consumer.next_frame(frame, 500ms)) + held.push_back(frame.slot); + CHECK(held.size() == consumer.hello().slot_count); + + // The capture loop kept running the whole time. + const uint64_t ticks_before = loop.ticks(); + std::this_thread::sleep_for(100ms); + CHECK(loop.ticks() > ticks_before); + + // Releasing the slots lets frames flow again. + for (uint32_t slot : held) + consumer.release(slot); + CHECK(consumer.next_frame(frame, 2000ms)); +} diff --git a/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp b/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp new file mode 100644 index 0000000000..e202bc6600 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp @@ -0,0 +1,200 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "mock_test_support.hpp" + +#include +#include +#include + +#include +#include + +using namespace plugins::sensing; +using namespace sensing_test; +using namespace std::chrono_literals; + +TEST_CASE("mock query renumbers sensors and honours the link masks", "[unit][gpu][sensing]") +{ + const std::string json = SENSING_PLATFORM_CONFIG; + + const auto both = SiplCamera::query(json, "SHW5G_2", { 0x0000, 0x1100 }); + REQUIRE(both.size() == 2); + // The JSON names these sensors 2 and 3; the pipeline indices are 0 and 1. + CHECK(both[0].id == 0); + CHECK(both[1].id == 1); + CHECK(both[0].name == "SHW5G_MIX_L2"); + CHECK(both[1].name == "SHW5G_MIX_L3"); + CHECK(both[0].width == 2560); + CHECK(both[0].height == 1984); + CHECK(both[0].fps == 60.0); + + // One link enabled leaves one sensor, which is renumbered to 0. + const auto link2 = SiplCamera::query(json, "SHW5G_2", { 0x0000, 0x0100 }); + REQUIRE(link2.size() == 1); + CHECK(link2[0].id == 0); + CHECK(link2[0].name == "SHW5G_MIX_L2"); + + const auto link3 = SiplCamera::query(json, "SHW5G_2", { 0x0000, 0x1000 }); + REQUIRE(link3.size() == 1); + CHECK(link3[0].id == 0); + CHECK(link3[0].name == "SHW5G_MIX_L3"); + + // The mask for the wrong deserializer enables nothing. + CHECK(SiplCamera::query(json, "SHW5G_2", { 0x1100, 0x0000 }).empty()); + CHECK(SiplCamera::query(json, "SHW5G_2", { 0x0000, 0x0000 }).empty()); + + // No mask list keeps every module. + CHECK(SiplCamera::query(json, "SHW5G_2", {}).size() == 2); +} + +TEST_CASE("mock query rejects a missing file or config", "[unit][gpu][sensing]") +{ + CHECK_THROWS_AS(SiplCamera::query("/nonexistent/platform.json", "SHW5G_2", {}), std::runtime_error); + CHECK_THROWS_AS(SiplCamera::query(SENSING_PLATFORM_CONFIG, "NO_SUCH_CONFIG", {}), std::runtime_error); +} + +TEST_CASE("SENSING_MOCK_SCALE_DIV shrinks the reported geometry and keeps it even", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + const auto sensors = SiplCamera::query(SENSING_PLATFORM_CONFIG, "SHW5G_2", { 0x0000, 0x1100 }); + REQUIRE(sensors.size() == 2); + CHECK(sensors[0].width == 320); + CHECK(sensors[0].height == 248); + CHECK(sensors[0].width % 2 == 0); + CHECK(sensors[0].height % 2 == 0); +} + +TEST_CASE("camera construction validates the masks", "[unit][gpu][sensing]") +{ + auto config = make_sipl_config(); + + config.link_masks = {}; + CHECK_THROWS_AS(SiplCamera(config), std::runtime_error); + + config.link_masks = { 0x0000, 0x0000 }; + CHECK_THROWS_WITH(SiplCamera(config), Catch::Matchers::ContainsSubstring("no modules left")); +} + +TEST_CASE("frames arrive with a fresh counter, rising sequence and rising stamps", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + SiplCamera camera(make_sipl_config()); + REQUIRE(camera.sensors().size() == 2); + camera.start(); + + for (const auto& sensor : camera.sensors()) + { + INFO("sensor " << sensor.id); + uint64_t last_sequence = 0; + uint64_t last_stamp = 0; + int seen = 0; + REQUIRE(wait_for( + [&] + { + auto view = camera.latest(sensor.id); + if (!view) + return false; + CHECK(view->width == sensor.width); + CHECK(view->height == sensor.height); + CHECK(view->ptr != 0); + CHECK(view->sequence > last_sequence); + CHECK(view->timestamp_ns > last_stamp); + last_sequence = view->sequence; + last_stamp = view->timestamp_ns; + + // The mock numbers frames from 0 and publishes the sequence from 1. + const auto pixels = download_rgba(view->ptr, view->pitch, view->width, view->height); + CHECK(read_counter(pixels, view->width) == ((view->sequence - 1) & 0xFFFF)); + return ++seen >= 6; + })); + } + camera.stop(); +} + +TEST_CASE("both sensors carry the same capture TSC for the same frame", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + SiplCamera camera(make_sipl_config()); + camera.start(); + + std::map tsc[2]; + REQUIRE(wait_for( + [&] + { + for (uint32_t id = 0; id < 2; ++id) + { + if (auto view = camera.latest(id)) + { + // Not on the CLOCK_MONOTONIC timebase, so nothing can confuse the two. + CHECK(view->capture_tsc_ns != view->timestamp_ns); + tsc[id][view->sequence] = view->capture_tsc_ns; + } + } + int common = 0; + for (const auto& [sequence, value] : tsc[0]) + { + auto other = tsc[1].find(sequence); + if (other != tsc[1].end()) + { + CHECK(other->second == value); + ++common; + } + } + return common >= 3; + })); + camera.stop(); +} + +TEST_CASE("a leased slot is not overwritten while the producer keeps running", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + SiplCamera camera(make_sipl_config()); + camera.start(); + + std::optional held; + REQUIRE(wait_for([&] { return (held = camera.latest(0)).has_value(); })); + const uint32_t counter = read_counter(download_rgba(held->ptr, held->pitch, held->width, held->height), held->width); + + // Several frame periods pass without another latest(): the lease must hold. + std::this_thread::sleep_for(150ms); + CHECK(read_counter(download_rgba(held->ptr, held->pitch, held->width, held->height), held->width) == counter); + camera.stop(); +} + +TEST_CASE("an injected link failure surfaces from latest()", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fail("SENSING_MOCK_FAIL_AFTER", "3"); + SiplCamera camera(make_sipl_config()); + camera.start(); + + std::string message; + REQUIRE(wait_for( + [&] + { + try + { + camera.latest(0); + } + catch (const std::runtime_error& e) + { + message = e.what(); + return true; + } + return false; + })); + CHECK(message.find("capture error") != std::string::npos); + camera.stop(); +} + +TEST_CASE("an unknown sensor index is an error, and stop is idempotent", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + SiplCamera camera(make_sipl_config()); + CHECK_THROWS_AS(camera.latest(7), std::runtime_error); + camera.start(); + camera.start(); + camera.stop(); + camera.stop(); +} diff --git a/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp b/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp new file mode 100644 index 0000000000..041e7a7363 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp @@ -0,0 +1,179 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "mock_test_support.hpp" + +#include +#include + +#include +#include +#include + +using namespace plugins::sensing; +using namespace sensing_test; +using namespace std::chrono_literals; + +namespace +{ + +constexpr uint32_t kWidth = 64; +constexpr uint32_t kHeight = 48; + +class Frame +{ +public: + Frame() + { + cudaMallocPitch(reinterpret_cast(&m_ptr), &m_pitch, kWidth * 4, kHeight); + cudaMemset2D(m_ptr, m_pitch, 0x80, kWidth * 4, kHeight); + } + ~Frame() + { + cudaFree(m_ptr); + } + uintptr_t ptr() const + { + return reinterpret_cast(m_ptr); + } + size_t pitch() const + { + return m_pitch; + } + +private: + uint8_t* m_ptr = nullptr; + size_t m_pitch = 0; +}; + +EncoderConfig make_config() +{ + EncoderConfig config; + config.width = kWidth; + config.height = kHeight; + config.fps = 30; + config.bitrate_bps = 1'000'000; + return config; +} + +} // namespace + +TEST_CASE("encoder rejects zero or odd geometry", "[unit][gpu][sensing]") +{ + auto config = make_config(); + config.width = 0; + CHECK_THROWS_AS(JetsonEncoder(config), std::runtime_error); + config = make_config(); + config.height = 47; + CHECK_THROWS_AS(JetsonEncoder(config), std::runtime_error); +} + +TEST_CASE("encoder holds frames back during warm-up, then emits them in order", "[unit][gpu][sensing]") +{ + Frame frame; + JetsonEncoder encoder(make_config()); + + REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000)); + std::this_thread::sleep_for(100ms); + CHECK(encoder.poll().empty()); // nothing yet: the encoder wants more input first + + REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'033'000'000)); + REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'066'000'000)); + REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'100'000'000)); + + std::vector units; + REQUIRE(wait_for( + [&] + { + for (auto unit = encoder.poll(); !unit.empty(); unit = encoder.poll()) + units.push_back(std::move(unit)); + return units.size() >= 2; + })); + + // The first unit opens the stream: SPS, PPS, IDR slice. + const auto first = split_annex_b(units[0].data); + REQUIRE(first.size() == 3); + CHECK(first[0].type == 0x67); + CHECK(first[1].type == 0x68); + CHECK(first[2].type == 0x65); + CHECK(frame_of(first[2]) == 0); + + const auto second = split_annex_b(units[1].data); + REQUIRE(second.size() == 1); + CHECK(second[0].type == 0x41); + CHECK(frame_of(second[0]) == 1); +} + +TEST_CASE("encoder returns the capture stamp at microsecond resolution", "[unit][gpu][sensing]") +{ + Frame frame; + JetsonEncoder encoder(make_config()); + + // A V4L2 timeval carries microseconds, so the sub-microsecond part cannot come back. + const int64_t stamp = 1'234'567'891'234; + REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), stamp)); + const auto tail = encoder.end_of_stream(); + REQUIRE_FALSE(tail.empty()); + CHECK(tail.timestamp_ns == 1'234'567'891'000); + CHECK(tail.timestamp_ns != stamp); +} + +TEST_CASE("encoder marks an IDR at every GOP boundary", "[unit][gpu][sensing]") +{ + Frame frame; + auto config = make_config(); + config.gop = 3; + JetsonEncoder encoder(config); + // Seven frames against six input buffers: submit() may refuse one until the encoder frees a buffer. + for (int i = 0; i < 7; ++i) + REQUIRE(wait_for([&] { return encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i * 33'000'000); })); + + const auto nals = split_annex_b(encoder.end_of_stream().data); + std::vector idr_frames; + for (const auto& nal : nals) + { + if (nal.type == 0x65) + idr_frames.push_back(frame_of(nal)); + } + CHECK(idr_frames == std::vector{ 0, 3, 6 }); +} + +TEST_CASE("encoder refuses frames when every input buffer is in flight", "[unit][gpu][sensing]") +{ + ScopedEnv latency("SENSING_MOCK_ENCODER_LATENCY_MS", "300"); + Frame frame; + JetsonEncoder encoder(make_config()); + + int accepted = 0; + int rejected = 0; + for (int i = 0; i < 10; ++i) + (encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i) ? accepted : rejected)++; + + // Six input buffers, and the encoder is too slow to return any in the loop. + CHECK(accepted == 6); + CHECK(rejected == 4); + + // Once the encoder frees one, submission works again. + CHECK(wait_for([&] { return encoder.submit(frame.ptr(), frame.pitch(), 2'000'000'000); }, 2s)); +} + +TEST_CASE("end of stream returns the tail as one unit carrying the last stamp", "[unit][gpu][sensing]") +{ + Frame frame; + JetsonEncoder encoder(make_config()); + for (int i = 0; i < 4; ++i) + REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i * 1'000'000)); + + const auto tail = encoder.end_of_stream(); + std::vector frames; + for (const auto& nal : split_annex_b(tail.data)) + { + if (is_slice(nal)) + frames.push_back(frame_of(nal)); + } + CHECK(frames == std::vector{ 0, 1, 2, 3 }); + CHECK(tail.timestamp_ns == 1'003'000'000); + + CHECK(encoder.poll().empty()); + CHECK(encoder.end_of_stream().empty()); +} diff --git a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp new file mode 100644 index 0000000000..895bd53eb3 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp @@ -0,0 +1,216 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// SensingCamera end to end on the mock backend: capture, convert, encode, sink. + +#include "mock_test_support.hpp" + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +using namespace plugins::sensing; +using namespace sensing_test; +using namespace std::chrono_literals; + +namespace +{ + +struct Recorded +{ + uint32_t sensor_id; + uint64_t sequence_number; + uint64_t capture_tsc_ns; + int64_t host_ns; +}; + +/// Collects what would have gone to MCAP. +struct Recorder +{ + std::mutex mutex; + std::vector records; +}; + +class RecordingPusher : public IMetadataPusher +{ +public: + explicit RecordingPusher(std::shared_ptr recorder) : m_recorder(std::move(recorder)) + { + } + + void on_frame_metadata(const core::FrameMetadataSensingT& metadata, + int64_t sample_time_local_common_clock_ns, + int64_t sample_time_raw_device_clock_ns) override + { + (void)sample_time_raw_device_clock_ns; + std::lock_guard lock(m_recorder->mutex); + m_recorder->records.push_back({ metadata.sensor_id, metadata.sequence_number, metadata.capture_tsc_ns, + sample_time_local_common_clock_ns }); + } + +private: + std::shared_ptr m_recorder; +}; + +SensingConfig make_config() +{ + SensingConfig config; + config.platform_config_json = SENSING_PLATFORM_CONFIG; + config.bitrate_bps = 2'000'000; + return config; +} + +std::vector read_file(const std::string& path) +{ + std::ifstream file(path, std::ios::binary); + return { std::istreambuf_iterator(file), std::istreambuf_iterator() }; +} + +std::vector records_for(const Recorder& recorder, uint32_t sensor) +{ + std::vector out; + for (const auto& r : recorder.records) + { + if (r.sensor_id == sensor) + out.push_back(r); + } + return out; +} + +// Runs the plugin's update loop for `duration`, then flushes. +void run_for(SensingCamera& camera, std::chrono::milliseconds duration) +{ + const auto end = std::chrono::steady_clock::now() + duration; + while (std::chrono::steady_clock::now() < end) + { + camera.update(); + std::this_thread::sleep_for(2ms); + } + camera.flush(); +} + +} // namespace + +TEST_CASE("an unknown sensor index is reported with the available pipelines", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + std::vector streams(1); + streams[0].sensor_id = 5; + streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); + CHECK_THROWS_WITH(SensingCamera(make_config(), streams, std::make_unique(streams)), + Catch::Matchers::ContainsSubstring("not a pipeline")); +} + +TEST_CASE("an encoded stream writes Annex-B data and one metadata record per unit", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + const std::string output = unique_path("sensing_mock", ".h264"); + + std::vector streams(1); + streams[0].sensor_id = 0; + streams[0].output_path = output; + + auto recorder = std::make_shared(); + { + SensingCamera camera( + make_config(), streams, std::make_unique(streams, std::make_unique(recorder))); + run_for(camera, 1000ms); + } + + const auto bytes = read_file(output); + std::remove(output.c_str()); + REQUIRE(bytes.size() > 4); + + const auto nals = split_annex_b(bytes); + REQUIRE_FALSE(nals.empty()); + CHECK(nals[0].type == 0x67); // the stream opens with its parameter sets + std::vector frames; + for (const auto& nal : nals) + { + if (is_slice(nal)) + frames.push_back(frame_of(nal)); + } + REQUIRE(frames.size() >= 10); + // The encoder preserves submission order; the mailbox may skip frames but never reorders. + for (size_t i = 1; i < frames.size(); ++i) + CHECK(frames[i] > frames[i - 1]); + + const auto records = records_for(*recorder, 0); + REQUIRE(records.size() >= 5); + for (size_t i = 0; i < records.size(); ++i) + CHECK(records[i].sequence_number == i); // contiguous from zero, one per dispatched unit + // The final flush folds the warm-up tail into one unit, so units can be fewer than frames. + CHECK(records.size() <= frames.size()); +} + +TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + std::vector streams(2); + streams[0].sensor_id = 0; + streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); + streams[1].sensor_id = 1; + streams[1].ipc_socket_path = unique_path("sensing_mock", ".sock"); + + auto recorder = std::make_shared(); + { + SensingCamera camera( + make_config(), streams, std::make_unique(streams, std::make_unique(recorder))); + run_for(camera, 700ms); + } + + std::vector per_sensor[2] = { records_for(*recorder, 0), records_for(*recorder, 1) }; + for (uint32_t sensor = 0; sensor < 2; ++sensor) + { + INFO("sensor " << sensor); + REQUIRE(per_sensor[sensor].size() >= 5); + for (size_t i = 0; i < per_sensor[sensor].size(); ++i) + { + CHECK(per_sensor[sensor][i].sequence_number == i); + CHECK(per_sensor[sensor][i].capture_tsc_ns != 0); + } + } + + // The same frame carries the same TSC on both sensors, so the two eyes can be paired on it. + std::set left; + for (const auto& r : per_sensor[0]) + left.insert(r.capture_tsc_ns); + int shared = 0; + for (const auto& r : per_sensor[1]) + shared += static_cast(left.count(r.capture_tsc_ns)); + CHECK(shared >= 3); +} + +// The encoder returns each unit's stamp at microsecond resolution, while SensingCamera remembers the +// stamp it submitted in nanoseconds and matches them for equality, so the capture TSC of an encoded +// stream is never found. Remove [!shouldfail] when that is fixed. +TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit][gpu][sensing][!shouldfail]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + const std::string output = unique_path("sensing_mock", ".h264"); + + std::vector streams(1); + streams[0].sensor_id = 0; + streams[0].output_path = output; + + auto recorder = std::make_shared(); + { + SensingCamera camera( + make_config(), streams, std::make_unique(streams, std::make_unique(recorder))); + run_for(camera, 700ms); + } + std::remove(output.c_str()); + + const auto records = records_for(*recorder, 0); + REQUIRE(records.size() >= 5); + // The last record is the flushed tail, whose stamp is not matched either way. + for (size_t i = 0; i + 1 < records.size(); ++i) + CHECK(records[i].capture_tsc_ns != 0); +} From 4f94973b10935bdd30831f0d0d93ac2e858fa046 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sat, 24 Oct 2026 17:29:04 -0400 Subject: [PATCH 02/10] Match encoder stamps at microsecond resolution The V4L2 encoder carries the capture stamp on a timeval, so a unit comes back with the stamp truncated to microseconds. SensingCamera remembers the stamp it submitted in nanoseconds and looked the unit up by equality, so the match essentially never succeeded: every metadata record of an encoded stream had capture_tsc_ns == 0, and each failed lookup also discarded the pending stamps ahead of it. Compare at microsecond resolution. The mock encoder reproduces the truncation, and the test that recovers the capture TSC of an encoded stream no longer needs [!shouldfail]. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/core/sensing_camera.cpp | 6 +++++- tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp | 5 +---- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/src/plugins/sensing/core/sensing_camera.cpp b/src/plugins/sensing/core/sensing_camera.cpp index 67a0937298..03a60c0f31 100644 --- a/src/plugins/sensing/core/sensing_camera.cpp +++ b/src/plugins/sensing/core/sensing_camera.cpp @@ -113,11 +113,15 @@ int64_t SensingCamera::take_capture_tsc(Stream& stream, int64_t timestamp_ns) // The encoder returns units in submission order, so the match is at the // front. Anything older than the match was dropped by the encoder and is // discarded with it, which is what keeps this bounded. + // + // The stamp rides a V4L2 timeval, which holds microseconds, so the unit + // comes back with the submitted stamp truncated. Compare at that resolution. + constexpr int64_t kNsPerUs = 1000; while (!stream.pending_stamps.empty()) { auto entry = stream.pending_stamps.front(); stream.pending_stamps.pop_front(); - if (entry.first == timestamp_ns) + if (entry.first / kNsPerUs == timestamp_ns / kNsPerUs) return entry.second; } return 0; diff --git a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp index 895bd53eb3..b320e01330 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp @@ -188,10 +188,7 @@ TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", " CHECK(shared >= 3); } -// The encoder returns each unit's stamp at microsecond resolution, while SensingCamera remembers the -// stamp it submitted in nanoseconds and matches them for equality, so the capture TSC of an encoded -// stream is never found. Remove [!shouldfail] when that is fixed. -TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit][gpu][sensing][!shouldfail]") +TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit][gpu][sensing]") { ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); const std::string output = unique_path("sensing_mock", ".h264"); From 560558e3c18adcceb9845d1a71ee733f8bf39c90 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sat, 24 Oct 2026 17:44:09 -0400 Subject: [PATCH 03/10] Tidy the mock backend after review - Free a mock sensor's device memory from its own destructor, so a constructor that throws part-way through allocating does not leak. - Abandon queued frames when the mock encoder is torn down instead of encoding them, which made teardown take as long as the queue was deep. - Treat a link index of 8 or more as disabled; a mask holds eight hex digits. - Fail configure with an install hint when the EGL headers cudaEGL.h needs are missing, and name them in the README. - Remove a compile definition nothing used, and a stale comment in tools/. - Remove the tests' temp files on every exit path. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/README.md | 2 +- src/plugins/sensing/core/CMakeLists.txt | 8 ++++++- .../sensing/core/mock_jetson_encoder.cpp | 3 ++- .../sensing/core/mock_platform_config.cpp | 4 +++- src/plugins/sensing/core/mock_sipl_camera.cpp | 21 +++++++++------- src/plugins/sensing/tools/CMakeLists.txt | 4 ++-- .../sensing_mock/mock_test_support.hpp | 24 +++++++++++++++++++ .../sensing_mock/test_mock_pipeline.cpp | 9 ++++--- 8 files changed, 55 insertions(+), 20 deletions(-) diff --git a/src/plugins/sensing/README.md b/src/plugins/sensing/README.md index dcc68f423d..d34de184c2 100644 --- a/src/plugins/sensing/README.md +++ b/src/plugins/sensing/README.md @@ -53,7 +53,7 @@ SIPL drivers. | Value | Capture | Encoder | Needs | | --- | --- | --- | --- | | `JETSON` | SIPL | V4L2 NVENC | L4T SDKs | -| `MOCK` | synthetic frames | stand-in | NVIDIA GPU, CUDA toolkit | +| `MOCK` | synthetic frames | stand-in | NVIDIA GPU, CUDA toolkit, EGL headers (`libegl-dev`) | | `AUTO` (default) | `JETSON` on aarch64, `MOCK` elsewhere | | | ```bash diff --git a/src/plugins/sensing/core/CMakeLists.txt b/src/plugins/sensing/core/CMakeLists.txt index f303e7c7ea..0af64a3bd3 100644 --- a/src/plugins/sensing/core/CMakeLists.txt +++ b/src/plugins/sensing/core/CMakeLists.txt @@ -20,6 +20,13 @@ if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") ${JETSON_MMAPI_SOURCES} ) else() + # yuv_to_rgba.cu includes cudaEGL.h, which needs the EGL headers even though + # nothing here opens an EGL display. + find_path(SENSING_EGL_INCLUDE_DIR EGL/egl.h) + if(NOT SENSING_EGL_INCLUDE_DIR) + message(FATAL_ERROR "The SENSING mock backend needs the EGL headers (cudaEGL.h includes them). " + "Install libegl-dev.") + endif() # Parses the vendor JSON itself, so a MOCK build has no SIPL query library. find_package(nlohmann_json 3 QUIET) if(NOT nlohmann_json_FOUND) @@ -99,5 +106,4 @@ if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") ) else() target_link_libraries(sensing_core PRIVATE nlohmann_json::nlohmann_json) - target_compile_definitions(sensing_core PUBLIC SENSING_MOCK_BACKEND=1) endif() diff --git a/src/plugins/sensing/core/mock_jetson_encoder.cpp b/src/plugins/sensing/core/mock_jetson_encoder.cpp index 76ed50cd54..facd1a521e 100644 --- a/src/plugins/sensing/core/mock_jetson_encoder.cpp +++ b/src/plugins/sensing/core/mock_jetson_encoder.cpp @@ -153,7 +153,8 @@ struct JetsonEncoder::Impl { std::unique_lock lock(mutex); wake.wait(lock, [&] { return closing || !jobs.empty(); }); - if (jobs.empty()) + // Frames still queued at teardown are abandoned, as the hardware encoder would. + if (closing || jobs.empty()) return; job = jobs.front(); jobs.pop_front(); diff --git a/src/plugins/sensing/core/mock_platform_config.cpp b/src/plugins/sensing/core/mock_platform_config.cpp index c4305b66b0..bf9cfd24d8 100644 --- a/src/plugins/sensing/core/mock_platform_config.cpp +++ b/src/plugins/sensing/core/mock_platform_config.cpp @@ -75,7 +75,9 @@ std::vector mock_query_platform_config(const std::string& platform_c const uint32_t transport = module.value("transportIndex", 0u); if (!link_masks.empty()) { - const bool enabled = transport < link_masks.size() && ((link_masks[transport] >> (4 * link)) & 0xF) != 0; + // A mask is eight hex digits, one per link; a link beyond that cannot be enabled. + const bool enabled = + link < 8 && transport < link_masks.size() && ((link_masks[transport] >> (4 * link)) & 0xF) != 0; if (!enabled) continue; } diff --git a/src/plugins/sensing/core/mock_sipl_camera.cpp b/src/plugins/sensing/core/mock_sipl_camera.cpp index 1150802472..521dbd9ab3 100644 --- a/src/plugins/sensing/core/mock_sipl_camera.cpp +++ b/src/plugins/sensing/core/mock_sipl_camera.cpp @@ -89,6 +89,18 @@ struct SiplCamera::Impl uint64_t published_capture_tsc_ns = 0; std::thread frame_thread; + + // Owns its device memory, so a constructor that throws part-way through + // allocating does not leak what it had already allocated. + ~Sensor() + { + cudaFree(luma.ptr); + cudaFree(chroma.ptr); + for (auto& buffer : buffers) + cudaFree(buffer.ptr); + if (stream) + cudaStreamDestroy(stream); + } }; explicit Impl(const SiplConfig& cfg) : config(cfg) @@ -120,15 +132,6 @@ struct SiplCamera::Impl ~Impl() { stop(); - for (auto& sensor : sensors) - { - cudaFree(sensor->luma.ptr); - cudaFree(sensor->chroma.ptr); - for (auto& buffer : sensor->buffers) - cudaFree(buffer.ptr); - if (sensor->stream) - cudaStreamDestroy(sensor->stream); - } } void allocate(Sensor& sensor) diff --git a/src/plugins/sensing/tools/CMakeLists.txt b/src/plugins/sensing/tools/CMakeLists.txt index 15634fd338..4a805c9748 100644 --- a/src/plugins/sensing/tools/CMakeLists.txt +++ b/src/plugins/sensing/tools/CMakeLists.txt @@ -3,8 +3,8 @@ # Animated test pattern on the IPC socket, so a consumer can be exercised with # no camera attached. Links sensing::cuda_ipc only, so the binary needs CUDA at -# runtime and neither SIPL nor the encoder. Configuring this directory still -# goes through the plugin's Jetson SDK checks in ../CMakeLists.txt. +# runtime and neither SIPL nor the encoder. With the JETSON backend, configuring +# this directory still goes through the SDK checks in ../CMakeLists.txt. add_executable(sensing_ipc_testsrc ipc_testsrc.cu) set_target_properties(sensing_ipc_testsrc PROPERTIES diff --git a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp index 6038dd7c33..fa7239607f 100644 --- a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -8,6 +8,7 @@ #include #include +#include #include #include #include @@ -66,6 +67,29 @@ inline std::string unique_path(const std::string& stem, const std::string& exten return "/tmp/" + stem + "_" + std::to_string(::getpid()) + "_" + std::to_string(counter++) + extension; } +/// A path that is removed when the scope ends, however the test exits. +class TempFile +{ +public: + explicit TempFile(const std::string& extension) : m_path(unique_path("sensing_mock", extension)) + { + } + ~TempFile() + { + std::remove(m_path.c_str()); + } + TempFile(const TempFile&) = delete; + TempFile& operator=(const TempFile&) = delete; + + const std::string& path() const + { + return m_path; + } + +private: + std::string m_path; +}; + /// Copies a device RGBA frame to the host as tightly packed rows. inline std::vector download_rgba(uintptr_t device_ptr, size_t pitch, uint32_t width, uint32_t height) { diff --git a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp index b320e01330..b0371f5d77 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp @@ -10,7 +10,6 @@ #include #include -#include #include #include #include @@ -111,7 +110,8 @@ TEST_CASE("an unknown sensor index is reported with the available pipelines", "[ TEST_CASE("an encoded stream writes Annex-B data and one metadata record per unit", "[unit][gpu][sensing]") { ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - const std::string output = unique_path("sensing_mock", ".h264"); + TempFile output_file(".h264"); + const std::string& output = output_file.path(); std::vector streams(1); streams[0].sensor_id = 0; @@ -125,7 +125,6 @@ TEST_CASE("an encoded stream writes Annex-B data and one metadata record per uni } const auto bytes = read_file(output); - std::remove(output.c_str()); REQUIRE(bytes.size() > 4); const auto nals = split_annex_b(bytes); @@ -191,7 +190,8 @@ TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", " TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit][gpu][sensing]") { ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - const std::string output = unique_path("sensing_mock", ".h264"); + TempFile output_file(".h264"); + const std::string& output = output_file.path(); std::vector streams(1); streams[0].sensor_id = 0; @@ -203,7 +203,6 @@ TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit] make_config(), streams, std::make_unique(streams, std::make_unique(recorder))); run_for(camera, 700ms); } - std::remove(output.c_str()); const auto records = records_for(*recorder, 0); REQUIRE(records.size() >= 5); From a6316d56c07348bc41de94b1592758bab93ac798 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sat, 24 Oct 2026 18:07:16 -0400 Subject: [PATCH 04/10] Show animated GIFs on the mock camera, with a bundled default Every mock sensor shows an animated GIF, by default the bundled robot. Set SENSING_MOCK_GIF for every sensor, or SENSING_MOCK_GIF_ for one (for example the bundled waving hand on a wrist camera). A GIF that cannot be loaded, the bundled one included, is an error; there is no fallback. The GIF is decoded once, held on the GPU, and played on the capture clock, with each frame's delay clamped to at least 20 ms and the animation looping. Transparent pixels are composited onto a plain gray. Frames are scaled to cover the sensor's configured size and centre-cropped, so the aspect ratio is never distorted whatever the GIF's own size: bilinear when enlarging and area-averaged when reducing. The result is full-range BT.709 NV12, the inverse of the matrix the real conversion applies, so it goes through that conversion unchanged. Sequence numbers and stamps advance on every capture even while the GIF shows one image. The bundled GIFs are Animated Noto Emoji (robot U+1F916, waving hand U+1F44B), copyright Google, licensed CC BY 4.0, included unmodified with their licence, source and attribution in mock_assets/. LICENSES/CC-BY-4.0.txt is the canonical text, and the root README notes the exception. Decoding uses stb_image, a public-domain header fetched at a pinned commit and checked against its hash at configure time. The mock backend no longer paints a test pattern. Tests use small generated GIFs of known colours and check sequence, stamps, whole (untorn) frames, the reader lease and the scaling. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- LICENSES/CC-BY-4.0.txt | 156 +++++++ README.md | 3 +- src/plugins/sensing/README.md | 18 +- src/plugins/sensing/core/CMakeLists.txt | 14 + src/plugins/sensing/core/mock_gif_source.cpp | 109 +++++ src/plugins/sensing/core/mock_gif_source.hpp | 73 ++++ src/plugins/sensing/core/mock_nv12_source.cu | 173 +++++--- src/plugins/sensing/core/mock_nv12_source.cuh | 34 +- src/plugins/sensing/core/mock_sipl_camera.cpp | 58 ++- src/plugins/sensing/mock_assets/README.md | 21 + .../sensing/mock_assets/noto-robot.gif | 3 + .../mock_assets/noto-robot.gif.license | 2 + .../sensing/mock_assets/noto-waving-hand.gif | 3 + .../mock_assets/noto-waving-hand.gif.license | 2 + tests/cpp/plugins/sensing_mock/CMakeLists.txt | 3 + .../cpp/plugins/sensing_mock/gif_fixture.hpp | 171 ++++++++ .../sensing_mock/mock_test_support.hpp | 107 ++++- .../sensing_mock/test_cuda_ipc_consumer.cpp | 28 +- .../plugins/sensing_mock/test_mock_camera.cpp | 59 ++- .../plugins/sensing_mock/test_mock_gif.cpp | 386 ++++++++++++++++++ .../sensing_mock/test_mock_pipeline.cpp | 8 +- 21 files changed, 1309 insertions(+), 122 deletions(-) create mode 100644 LICENSES/CC-BY-4.0.txt create mode 100644 src/plugins/sensing/core/mock_gif_source.cpp create mode 100644 src/plugins/sensing/core/mock_gif_source.hpp create mode 100644 src/plugins/sensing/mock_assets/README.md create mode 100644 src/plugins/sensing/mock_assets/noto-robot.gif create mode 100644 src/plugins/sensing/mock_assets/noto-robot.gif.license create mode 100644 src/plugins/sensing/mock_assets/noto-waving-hand.gif create mode 100644 src/plugins/sensing/mock_assets/noto-waving-hand.gif.license create mode 100644 tests/cpp/plugins/sensing_mock/gif_fixture.hpp create mode 100644 tests/cpp/plugins/sensing_mock/test_mock_gif.cpp diff --git a/LICENSES/CC-BY-4.0.txt b/LICENSES/CC-BY-4.0.txt new file mode 100644 index 0000000000..13ca539f37 --- /dev/null +++ b/LICENSES/CC-BY-4.0.txt @@ -0,0 +1,156 @@ +Creative Commons Attribution 4.0 International + + Creative Commons Corporation (“Creative Commons”) is not a law firm and does not provide legal services or legal advice. Distribution of Creative Commons public licenses does not create a lawyer-client or other relationship. Creative Commons makes its licenses and related information available on an “as-is” basis. Creative Commons gives no warranties regarding its licenses, any material licensed under their terms and conditions, or any related information. Creative Commons disclaims all liability for damages resulting from their use to the fullest extent possible. + +Using Creative Commons Public Licenses + +Creative Commons public licenses provide a standard set of terms and conditions that creators and other rights holders may use to share original works of authorship and other material subject to copyright and certain other rights specified in the public license below. The following considerations are for informational purposes only, are not exhaustive, and do not form part of our licenses. + +Considerations for licensors: Our public licenses are intended for use by those authorized to give the public permission to use material in ways otherwise restricted by copyright and certain other rights. Our licenses are irrevocable. Licensors should read and understand the terms and conditions of the license they choose before applying it. Licensors should also secure all rights necessary before applying our licenses so that the public can reuse the material as expected. Licensors should clearly mark any material not subject to the license. This includes other CC-licensed material, or material used under an exception or limitation to copyright. More considerations for licensors. + +Considerations for the public: By using one of our public licenses, a licensor grants the public permission to use the licensed material under specified terms and conditions. If the licensor’s permission is not necessary for any reason–for example, because of any applicable exception or limitation to copyright–then that use is not regulated by the license. Our licenses grant only permissions under copyright and certain other rights that a licensor has authority to grant. Use of the licensed material may still be restricted for other reasons, including because others have copyright or other rights in the material. A licensor may make special requests, such as asking that all changes be marked or described. Although not required by our licenses, you are encouraged to respect those requests where reasonable. More considerations for the public. + +Creative Commons Attribution 4.0 International Public License + +By exercising the Licensed Rights (defined below), You accept and agree to be bound by the terms and conditions of this Creative Commons Attribution 4.0 International Public License ("Public License"). To the extent this Public License may be interpreted as a contract, You are granted the Licensed Rights in consideration of Your acceptance of these terms and conditions, and the Licensor grants You such rights in consideration of benefits the Licensor receives from making the Licensed Material available under these terms and conditions. + +Section 1 – Definitions. + + a. Adapted Material means material subject to Copyright and Similar Rights that is derived from or based upon the Licensed Material and in which the Licensed Material is translated, altered, arranged, transformed, or otherwise modified in a manner requiring permission under the Copyright and Similar Rights held by the Licensor. For purposes of this Public License, where the Licensed Material is a musical work, performance, or sound recording, Adapted Material is always produced where the Licensed Material is synched in timed relation with a moving image. + + b. Adapter's License means the license You apply to Your Copyright and Similar Rights in Your contributions to Adapted Material in accordance with the terms and conditions of this Public License. + + c. Copyright and Similar Rights means copyright and/or similar rights closely related to copyright including, without limitation, performance, broadcast, sound recording, and Sui Generis Database Rights, without regard to how the rights are labeled or categorized. For purposes of this Public License, the rights specified in Section 2(b)(1)-(2) are not Copyright and Similar Rights. + + d. Effective Technological Measures means those measures that, in the absence of proper authority, may not be circumvented under laws fulfilling obligations under Article 11 of the WIPO Copyright Treaty adopted on December 20, 1996, and/or similar international agreements. + + e. Exceptions and Limitations means fair use, fair dealing, and/or any other exception or limitation to Copyright and Similar Rights that applies to Your use of the Licensed Material. + + f. Licensed Material means the artistic or literary work, database, or other material to which the Licensor applied this Public License. + + g. Licensed Rights means the rights granted to You subject to the terms and conditions of this Public License, which are limited to all Copyright and Similar Rights that apply to Your use of the Licensed Material and that the Licensor has authority to license. + + h. Licensor means the individual(s) or entity(ies) granting rights under this Public License. + + i. Share means to provide material to the public by any means or process that requires permission under the Licensed Rights, such as reproduction, public display, public performance, distribution, dissemination, communication, or importation, and to make material available to the public including in ways that members of the public may access the material from a place and at a time individually chosen by them. + + j. Sui Generis Database Rights means rights other than copyright resulting from Directive 96/9/EC of the European Parliament and of the Council of 11 March 1996 on the legal protection of databases, as amended and/or succeeded, as well as other essentially equivalent rights anywhere in the world. + + k. You means the individual or entity exercising the Licensed Rights under this Public License. Your has a corresponding meaning. + +Section 2 – Scope. + + a. License grant. + + 1. Subject to the terms and conditions of this Public License, the Licensor hereby grants You a worldwide, royalty-free, non-sublicensable, non-exclusive, irrevocable license to exercise the Licensed Rights in the Licensed Material to: + + A. reproduce and Share the Licensed Material, in whole or in part; and + + B. produce, reproduce, and Share Adapted Material. + + 2. Exceptions and Limitations. For the avoidance of doubt, where Exceptions and Limitations apply to Your use, this Public License does not apply, and You do not need to comply with its terms and conditions. + + 3. Term. The term of this Public License is specified in Section 6(a). + + 4. Media and formats; technical modifications allowed. The Licensor authorizes You to exercise the Licensed Rights in all media and formats whether now known or hereafter created, and to make technical modifications necessary to do so. The Licensor waives and/or agrees not to assert any right or authority to forbid You from making technical modifications necessary to exercise the Licensed Rights, including technical modifications necessary to circumvent Effective Technological Measures. For purposes of this Public License, simply making modifications authorized by this Section 2(a)(4) never produces Adapted Material. + + 5. Downstream recipients. + + A. Offer from the Licensor – Licensed Material. Every recipient of the Licensed Material automatically receives an offer from the Licensor to exercise the Licensed Rights under the terms and conditions of this Public License. + + B. No downstream restrictions. You may not offer or impose any additional or different terms or conditions on, or apply any Effective Technological Measures to, the Licensed Material if doing so restricts exercise of the Licensed Rights by any recipient of the Licensed Material. + + 6. No endorsement. Nothing in this Public License constitutes or may be construed as permission to assert or imply that You are, or that Your use of the Licensed Material is, connected with, or sponsored, endorsed, or granted official status by, the Licensor or others designated to receive attribution as provided in Section 3(a)(1)(A)(i). + +b. Other rights. + + 1. Moral rights, such as the right of integrity, are not licensed under this Public License, nor are publicity, privacy, and/or other similar personality rights; however, to the extent possible, the Licensor waives and/or agrees not to assert any such rights held by the Licensor to the limited extent necessary to allow You to exercise the Licensed Rights, but not otherwise. + + 2. Patent and trademark rights are not licensed under this Public License. + + 3. To the extent possible, the Licensor waives any right to collect royalties from You for the exercise of the Licensed Rights, whether directly or through a collecting society under any voluntary or waivable statutory or compulsory licensing scheme. In all other cases the Licensor expressly reserves any right to collect such royalties. + +Section 3 – License Conditions. + +Your exercise of the Licensed Rights is expressly made subject to the following conditions. + + a. Attribution. + + 1. If You Share the Licensed Material (including in modified form), You must: + + A. retain the following if it is supplied by the Licensor with the Licensed Material: + + i. identification of the creator(s) of the Licensed Material and any others designated to receive attribution, in any reasonable manner requested by the Licensor (including by pseudonym if designated); + + ii. a copyright notice; + + iii. a notice that refers to this Public License; + + iv. a notice that refers to the disclaimer of warranties; + + v. a URI or hyperlink to the Licensed Material to the extent reasonably practicable; + + B. indicate if You modified the Licensed Material and retain an indication of any previous modifications; and + + C. indicate the Licensed Material is licensed under this Public License, and include the text of, or the URI or hyperlink to, this Public License. + + 2. You may satisfy the conditions in Section 3(a)(1) in any reasonable manner based on the medium, means, and context in which You Share the Licensed Material. For example, it may be reasonable to satisfy the conditions by providing a URI or hyperlink to a resource that includes the required information. + + 3. If requested by the Licensor, You must remove any of the information required by Section 3(a)(1)(A) to the extent reasonably practicable. + + 4. If You Share Adapted Material You produce, the Adapter's License You apply must not prevent recipients of the Adapted Material from complying with this Public License. + +Section 4 – Sui Generis Database Rights. + +Where the Licensed Rights include Sui Generis Database Rights that apply to Your use of the Licensed Material: + + a. for the avoidance of doubt, Section 2(a)(1) grants You the right to extract, reuse, reproduce, and Share all or a substantial portion of the contents of the database; + + b. if You include all or a substantial portion of the database contents in a database in which You have Sui Generis Database Rights, then the database in which You have Sui Generis Database Rights (but not its individual contents) is Adapted Material; and + + c. You must comply with the conditions in Section 3(a) if You Share all or a substantial portion of the contents of the database. +For the avoidance of doubt, this Section 4 supplements and does not replace Your obligations under this Public License where the Licensed Rights include other Copyright and Similar Rights. + +Section 5 – Disclaimer of Warranties and Limitation of Liability. + + a. Unless otherwise separately undertaken by the Licensor, to the extent possible, the Licensor offers the Licensed Material as-is and as-available, and makes no representations or warranties of any kind concerning the Licensed Material, whether express, implied, statutory, or other. This includes, without limitation, warranties of title, merchantability, fitness for a particular purpose, non-infringement, absence of latent or other defects, accuracy, or the presence or absence of errors, whether or not known or discoverable. Where disclaimers of warranties are not allowed in full or in part, this disclaimer may not apply to You. + + b. To the extent possible, in no event will the Licensor be liable to You on any legal theory (including, without limitation, negligence) or otherwise for any direct, special, indirect, incidental, consequential, punitive, exemplary, or other losses, costs, expenses, or damages arising out of this Public License or use of the Licensed Material, even if the Licensor has been advised of the possibility of such losses, costs, expenses, or damages. Where a limitation of liability is not allowed in full or in part, this limitation may not apply to You. + + c. The disclaimer of warranties and limitation of liability provided above shall be interpreted in a manner that, to the extent possible, most closely approximates an absolute disclaimer and waiver of all liability. + +Section 6 – Term and Termination. + + a. This Public License applies for the term of the Copyright and Similar Rights licensed here. However, if You fail to comply with this Public License, then Your rights under this Public License terminate automatically. + + b. Where Your right to use the Licensed Material has terminated under Section 6(a), it reinstates: + + 1. automatically as of the date the violation is cured, provided it is cured within 30 days of Your discovery of the violation; or + + 2. upon express reinstatement by the Licensor. + + c. For the avoidance of doubt, this Section 6(b) does not affect any right the Licensor may have to seek remedies for Your violations of this Public License. + + d. For the avoidance of doubt, the Licensor may also offer the Licensed Material under separate terms or conditions or stop distributing the Licensed Material at any time; however, doing so will not terminate this Public License. + + e. Sections 1, 5, 6, 7, and 8 survive termination of this Public License. + +Section 7 – Other Terms and Conditions. + + a. The Licensor shall not be bound by any additional or different terms or conditions communicated by You unless expressly agreed. + + b. Any arrangements, understandings, or agreements regarding the Licensed Material not stated herein are separate from and independent of the terms and conditions of this Public License. + +Section 8 – Interpretation. + + a. For the avoidance of doubt, this Public License does not, and shall not be interpreted to, reduce, limit, restrict, or impose conditions on any use of the Licensed Material that could lawfully be made without permission under this Public License. + + b. To the extent possible, if any provision of this Public License is deemed unenforceable, it shall be automatically reformed to the minimum extent necessary to make it enforceable. If the provision cannot be reformed, it shall be severed from this Public License without affecting the enforceability of the remaining terms and conditions. + + c. No term or condition of this Public License will be waived and no failure to comply consented to unless expressly agreed to by the Licensor. + + d. Nothing in this Public License constitutes or may be interpreted as a limitation upon, or waiver of, any privileges and immunities that apply to the Licensor or You, including from the legal processes of any jurisdiction or authority. + +Creative Commons is not a party to its public licenses. Notwithstanding, Creative Commons may elect to apply one of its public licenses to material it publishes and in those instances will be considered the “Licensor.” Except for the limited purpose of indicating that material is shared under a Creative Commons public license or as otherwise permitted by the Creative Commons policies published at creativecommons.org/policies, Creative Commons does not authorize the use of the trademark “Creative Commons” or any other trademark or logo of Creative Commons without its prior written consent including, without limitation, in connection with any unauthorized modifications to any of its public licenses or any other arrangements, understandings, or agreements concerning use of licensed material. For the avoidance of doubt, this paragraph does not form part of the public licenses. + +Creative Commons may be contacted at creativecommons.org. diff --git a/README.md b/README.md index 8f9da8f6b7..2ad335407c 100644 --- a/README.md +++ b/README.md @@ -80,4 +80,5 @@ To get started, please refer to Isaac Lab's [Installation](https://isaac-sim.git The default license is [Apache-2.0](LICENSE.md). EgoPost is licensed under [OpenMDW-1.1](src/python/isaaccapture_egopost/LICENSE.md). -Third-party materials may carry their own licenses. +Third-party materials may carry their own licenses; for example, the sample GIFs under +[`src/plugins/sensing/mock_assets/`](src/plugins/sensing/mock_assets/) are Animated Noto Emoji under CC BY 4.0. diff --git a/src/plugins/sensing/README.md b/src/plugins/sensing/README.md index d34de184c2..e6b63cc16e 100644 --- a/src/plugins/sensing/README.md +++ b/src/plugins/sensing/README.md @@ -53,7 +53,7 @@ SIPL drivers. | Value | Capture | Encoder | Needs | | --- | --- | --- | --- | | `JETSON` | SIPL | V4L2 NVENC | L4T SDKs | -| `MOCK` | synthetic frames | stand-in | NVIDIA GPU, CUDA toolkit, EGL headers (`libegl-dev`) | +| `MOCK` | animated GIF frames | stand-in | NVIDIA GPU, CUDA toolkit, EGL headers (`libegl-dev`) | | `AUTO` (default) | `JETSON` on aarch64, `MOCK` elsewhere | | | ```bash @@ -64,9 +64,9 @@ cmake --build build --target camera_plugin_sensing sensing_mock_tests Everything above those two classes is the code that runs on the rig: `SensingCamera`, the sinks, the YUV-to-RGBA and RGBA-to-NV12 kernels, and the CUDA IPC publisher. The mock -resolves `--list-sensors` and the link masks from the vendor JSON, paints a synthetic NV12 -surface (a 16-bit frame counter top-left, chroma tinted per sensor), and runs the real -conversion on it. Both sensors share a capture TSC, as the rig's fsync does. +resolves `--list-sensors` and the link masks from the vendor JSON, builds a synthetic NV12 +surface from a GIF frame, and runs the real conversion on it. Both sensors share a capture +TSC, as the rig's fsync does. The mock encoder's output is **not decodable video**: Annex-B start codes, SPS/PPS/IDR and non-IDR NAL types, and a payload naming the frame. It reproduces what callers depend on: a @@ -74,10 +74,20 @@ fixed pool of input buffers that makes `submit()` refuse frames when the encoder a warm-up delay before the first unit, and capture stamps that come back at microsecond resolution. +Every sensor shows an animated GIF, by default the bundled robot in [`mock_assets/`](mock_assets/) +(Animated Noto Emoji, CC BY 4.0; see its README for the licence and source). The GIF plays on the +capture clock, looping, scaled to cover the sensor's configured size and centre-cropped so it is +never stretched (bilinear when enlarging, area-averaged when reducing). Transparent pixels show a +plain gray. The sequence number and stamps advance on every capture even while the GIF holds one +image. The GIF's own size and aspect do not matter; the frame size always comes from the platform +config. A GIF that cannot be loaded, the bundled one included, is an error: there is no fallback. + Environment, read once at construction: | Variable | Effect | | --- | --- | +| `SENSING_MOCK_GIF` | path of an animated GIF shown by every sensor (default: the bundled robot) | +| `SENSING_MOCK_GIF_` | a GIF for one sensor, by module name (`SENSING_MOCK_GIF_SHW5G_MIX_L3`); wins over `SENSING_MOCK_GIF`. The bundled `noto-waving-hand.gif` suits a wrist camera | | `SENSING_MOCK_FPS` | override the platform config's frame rate | | `SENSING_MOCK_SCALE_DIV` | divide the reported width and height (kept even); useful for tests | | `SENSING_MOCK_FAIL_AFTER` | every sensor fails after N frames, as a link drop would | diff --git a/src/plugins/sensing/core/CMakeLists.txt b/src/plugins/sensing/core/CMakeLists.txt index 0af64a3bd3..98d06b3bc2 100644 --- a/src/plugins/sensing/core/CMakeLists.txt +++ b/src/plugins/sensing/core/CMakeLists.txt @@ -42,8 +42,20 @@ else() FetchContent_MakeAvailable(nlohmann_json) isaac_teleop_third_party_scope_end() endif() + # Decodes the GIFs the mock camera can show. A single public-domain header, pinned to a commit and + # checked against its hash, so nothing else of the project is cloned. + include(FetchContent) + FetchContent_Declare( + stb_image + URL https://raw.githubusercontent.com/nothings/stb/003501f6d7a8668ce8b6d875e34dd97ad0d5d5b5/stb_image.h + URL_HASH SHA256=594c2fe35d49488b4382dbfaec8f98366defca819d916ac95becf3e75f4200b3 + DOWNLOAD_NO_EXTRACT TRUE + ) + FetchContent_MakeAvailable(stb_image) + set(SENSING_BACKEND_SOURCES mock_sipl_camera.cpp + mock_gif_source.cpp mock_platform_config.cpp mock_nv12_source.cu mock_jetson_encoder.cpp @@ -106,4 +118,6 @@ if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") ) else() target_link_libraries(sensing_core PRIVATE nlohmann_json::nlohmann_json) + target_include_directories(sensing_core SYSTEM PRIVATE ${stb_image_SOURCE_DIR}) + target_compile_definitions(sensing_core PRIVATE SENSING_MOCK_ASSET_DIR="${CMAKE_CURRENT_SOURCE_DIR}/../mock_assets") endif() diff --git a/src/plugins/sensing/core/mock_gif_source.cpp b/src/plugins/sensing/core/mock_gif_source.cpp new file mode 100644 index 0000000000..3f574297de --- /dev/null +++ b/src/plugins/sensing/core/mock_gif_source.cpp @@ -0,0 +1,109 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "mock_gif_source.hpp" + +#include + +#define STBI_ONLY_GIF +#define STBI_NO_STDIO +#define STB_IMAGE_IMPLEMENTATION +#include +#include +#include +#include +#include +#include + +namespace plugins +{ +namespace sensing +{ + +namespace +{ + +constexpr int kMinDelayMs = 20; + +} // namespace + +std::vector gif_frame_ends(const std::vector& delays_ms) +{ + std::vector ends; + ends.reserve(delays_ms.size()); + double total = 0.0; + for (int delay : delays_ms) + { + total += std::max(delay, kMinDelayMs) / 1000.0; + ends.push_back(total); + } + return ends; +} + +size_t gif_frame_index(const std::vector& frame_ends, double seconds) +{ + if (frame_ends.empty()) + return 0; + const double t = std::fmod(std::max(seconds, 0.0), frame_ends.back()); + const auto it = std::upper_bound(frame_ends.begin(), frame_ends.end(), t); + return std::min(static_cast(it - frame_ends.begin()), frame_ends.size() - 1); +} + +MockGifSource::MockGifSource(const std::string& path) +{ + std::ifstream file(path, std::ios::binary); + if (!file) + throw std::runtime_error("mock GIF: cannot open " + path); + const std::vector bytes((std::istreambuf_iterator(file)), std::istreambuf_iterator()); + + int width = 0; + int height = 0; + int frames = 0; + int channels = 0; + int* delays = nullptr; + unsigned char* pixels = stbi_load_gif_from_memory( + bytes.data(), static_cast(bytes.size()), &delays, &width, &height, &frames, &channels, 4); + if (!pixels) + throw std::runtime_error("mock GIF: cannot decode " + path + ": " + stbi_failure_reason()); + + // Composite onto the background, so the kernel only ever sees opaque pixels. + const size_t pixel_count = static_cast(width) * height * frames; + for (size_t i = 0; i < pixel_count; ++i) + { + unsigned char* px = pixels + i * 4; + const int alpha = px[3]; + for (int c = 0; c < 3; ++c) + px[c] = static_cast((px[c] * alpha + kBackground * (255 - alpha) + 127) / 255); + px[3] = 255; + } + + m_width = static_cast(width); + m_height = static_cast(height); + m_frame_ends = gif_frame_ends(std::vector(delays, delays + frames)); + + const size_t bytes_per_frame = static_cast(width) * height * 4; + const cudaError_t alloc = cudaMalloc(reinterpret_cast(&m_device_frames), bytes_per_frame * frames); + const cudaError_t copy = alloc == cudaSuccess ? + cudaMemcpy(m_device_frames, pixels, bytes_per_frame * frames, cudaMemcpyHostToDevice) : + alloc; + stbi_image_free(pixels); + stbi_image_free(delays); + if (copy != cudaSuccess) + { + cudaFree(m_device_frames); + throw std::runtime_error(std::string("mock GIF: cannot upload to the GPU: ") + cudaGetErrorString(copy)); + } +} + +MockGifSource::~MockGifSource() +{ + cudaFree(m_device_frames); +} + +const uint8_t* MockGifSource::frame(size_t index) const +{ + return m_device_frames + index * static_cast(m_width) * m_height * 4; +} + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_gif_source.hpp b/src/plugins/sensing/core/mock_gif_source.hpp new file mode 100644 index 0000000000..5a21edbcf2 --- /dev/null +++ b/src/plugins/sensing/core/mock_gif_source.hpp @@ -0,0 +1,73 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Mock backend only. An animated GIF decoded once and held in device memory, so the +// mock camera can show real moving images. + +#pragma once + +#include +#include +#include +#include + +namespace plugins +{ +namespace sensing +{ + +/// Time in seconds at which each GIF frame stops showing: the running total of the frame delays, +/// each clamped to at least 20 ms, as browsers do for the tiny delays some encoders write. +std::vector gif_frame_ends(const std::vector& delays_ms); + +/// Index of the frame showing `seconds` into the animation, which loops over its whole length. +size_t gif_frame_index(const std::vector& frame_ends, double seconds); + +class MockGifSource +{ +public: + /// Transparent pixels are composited onto this gray, so an emoji or sticker without a + /// background reads as an image rather than as whatever colour sits under its transparency. + static constexpr uint8_t kBackground = 48; + + /// Decodes every frame to opaque RGBA8 and uploads them to the current CUDA device. + /// @throws std::runtime_error if the file cannot be read or is not a GIF. + explicit MockGifSource(const std::string& path); + ~MockGifSource(); + + MockGifSource(const MockGifSource&) = delete; + MockGifSource& operator=(const MockGifSource&) = delete; + + uint32_t width() const + { + return m_width; + } + uint32_t height() const + { + return m_height; + } + size_t frame_count() const + { + return m_frame_ends.size(); + } + /// Row stride of every frame, in bytes. Rows are tightly packed. + size_t pitch() const + { + return static_cast(m_width) * 4; + } + /// Device pointer to one RGBA8 frame. + const uint8_t* frame(size_t index) const; + size_t frame_index_at(double seconds) const + { + return gif_frame_index(m_frame_ends, seconds); + } + +private: + uint32_t m_width = 0; + uint32_t m_height = 0; + std::vector m_frame_ends; + uint8_t* m_device_frames = nullptr; +}; + +} // namespace sensing +} // namespace plugins diff --git a/src/plugins/sensing/core/mock_nv12_source.cu b/src/plugins/sensing/core/mock_nv12_source.cu index 81d757083c..ba0e2c8351 100644 --- a/src/plugins/sensing/core/mock_nv12_source.cu +++ b/src/plugins/sensing/core/mock_nv12_source.cu @@ -11,73 +11,150 @@ namespace sensing namespace { -constexpr int kCounterCell = 16; -constexpr int kCounterBits = 16; - -__global__ void paint_luma(uint8_t* y_plane, int y_pitch, int width, int height, uint32_t frame_index) +// Geometry of the cover resize: the source scaled uniformly until it fills the target, then +// centre-cropped. `scaled_*` is the size of the scaled image and `crop_*` the offset of the +// visible window inside it. +struct CoverGeometry { - const int x = blockIdx.x * blockDim.x + threadIdx.x; - const int y = blockIdx.y * blockDim.y + threadIdx.y; - if (x >= width || y >= height) - return; - - const float u = static_cast(x) / width; - const float v = static_cast(y) / height; - const float t = static_cast(frame_index % 240u) / 240.0f; + int src_width; + int src_height; + int scaled_width; + int scaled_height; + int crop_x; + int crop_y; +}; + +__device__ float3 fetch_rgb(const uint8_t* src, int pitch, int x, int y) +{ + const uint8_t* px = src + static_cast(y) * pitch + static_cast(x) * 4; + return make_float3(px[0], px[1], px[2]); +} - // Diagonal gradient scrolling with the frame index. - float luma = 255.0f * (0.25f + 0.5f * fabsf(__sinf(6.2831f * (u + v + t)))); +// Colour of one pixel of the cropped output. Enlarging samples bilinearly; reducing averages the +// source area the pixel covers, weighted by overlap, as a box filter would. +__device__ float3 sample_cover(const uint8_t* src, int pitch, const CoverGeometry& g, int x, int y) +{ + const float sx = static_cast(g.src_width) / g.scaled_width; + const float sy = static_cast(g.src_height) / g.scaled_height; + const int X = x + g.crop_x; + const int Y = y + g.crop_y; - // A bright square orbiting the centre: a repeated frame freezes it. - const float cx = 0.5f + 0.3f * __cosf(6.2831f * t); - const float cy = 0.5f + 0.3f * __sinf(6.2831f * t); - if (fabsf(u - cx) < 0.04f && fabsf(v - cy) < 0.04f * width / height) - luma = 255.0f; + if (sx <= 1.0f && sy <= 1.0f) + { + const float fx = fminf(fmaxf((X + 0.5f) * sx - 0.5f, 0.0f), g.src_width - 1.0f); + const float fy = fminf(fmaxf((Y + 0.5f) * sy - 0.5f, 0.0f), g.src_height - 1.0f); + const int x0 = static_cast(fx); + const int y0 = static_cast(fy); + const int x1 = min(x0 + 1, g.src_width - 1); + const int y1 = min(y0 + 1, g.src_height - 1); + const float tx = fx - x0; + const float ty = fy - y0; + const float3 a = fetch_rgb(src, pitch, x0, y0); + const float3 b = fetch_rgb(src, pitch, x1, y0); + const float3 c = fetch_rgb(src, pitch, x0, y1); + const float3 d = fetch_rgb(src, pitch, x1, y1); + return make_float3((a.x * (1 - tx) + b.x * tx) * (1 - ty) + (c.x * (1 - tx) + d.x * tx) * ty, + (a.y * (1 - tx) + b.y * tx) * (1 - ty) + (c.y * (1 - tx) + d.y * tx) * ty, + (a.z * (1 - tx) + b.z * tx) * (1 - ty) + (c.z * (1 - tx) + d.z * tx) * ty); + } - if (y < kCounterCell && x < kCounterBits * kCounterCell) + const float left = X * sx; + const float right = (X + 1) * sx; + const float top = Y * sy; + const float bottom = (Y + 1) * sy; + float3 sum = make_float3(0.0f, 0.0f, 0.0f); + float weight = 0.0f; + for (int j = static_cast(top); j < min(static_cast(ceilf(bottom)), g.src_height); ++j) { - const unsigned bit = static_cast(x / kCounterCell); - luma = ((frame_index >> (kCounterBits - 1u - bit)) & 1u) ? 255.0f : 0.0f; + const float wy = fminf(bottom, j + 1.0f) - fmaxf(top, static_cast(j)); + for (int i = static_cast(left); i < min(static_cast(ceilf(right)), g.src_width); ++i) + { + const float w = wy * (fminf(right, i + 1.0f) - fmaxf(left, static_cast(i))); + const float3 rgb = fetch_rgb(src, pitch, i, j); + sum.x += rgb.x * w; + sum.y += rgb.y * w; + sum.z += rgb.z * w; + weight += w; + } } + const float inv = weight > 0.0f ? 1.0f / weight : 0.0f; + return make_float3(sum.x * inv, sum.y * inv, sum.z * inv); +} - y_plane[static_cast(y) * y_pitch + x] = static_cast(luma); +__device__ float luma_709(const float3& rgb) +{ + return fminf(fmaxf(0.2126f * rgb.x + 0.7152f * rgb.y + 0.0722f * rgb.z, 0.0f), 255.0f); } -__global__ void paint_chroma(uint8_t* uv_plane, int uv_pitch, int width, int height, uint32_t sensor_id) +// One thread per 2x2 block: four luma samples and one chroma pair. +__global__ void paint_gif(const uint8_t* src, + int src_pitch, + CoverGeometry geometry, + uint8_t* y_plane, + int y_pitch, + uint8_t* uv_plane, + int uv_pitch, + int width, + int height) { - const int x = blockIdx.x * blockDim.x + threadIdx.x; // chroma sample column - const int y = blockIdx.y * blockDim.y + threadIdx.y; - if (x >= width / 2 || y >= height / 2) + const int cx = blockIdx.x * blockDim.x + threadIdx.x; + const int cy = blockIdx.y * blockDim.y + threadIdx.y; + if (cx >= width / 2 || cy >= height / 2) return; - // Even sensors lean blue-green, odd sensors red-orange. The counter strip stays neutral so its - // cells decode as pure black and white. - const bool odd = (sensor_id & 1u) != 0; - const bool counter = y < kCounterCell / 2 && x < kCounterBits * kCounterCell / 2; - uint8_t* sample = uv_plane + static_cast(y) * uv_pitch + static_cast(x) * 2; - sample[0] = counter ? 128 : (odd ? 100 : 160); // U - sample[1] = counter ? 128 : (odd ? 180 : 100); // V + float3 mean = make_float3(0.0f, 0.0f, 0.0f); + for (int dy = 0; dy < 2; ++dy) + { + for (int dx = 0; dx < 2; ++dx) + { + const int x = 2 * cx + dx; + const int y = 2 * cy + dy; + const float3 rgb = sample_cover(src, src_pitch, geometry, x, y); + y_plane[static_cast(y) * y_pitch + x] = static_cast(luma_709(rgb) + 0.5f); + mean.x += 0.25f * rgb.x; + mean.y += 0.25f * rgb.y; + mean.z += 0.25f * rgb.z; + } + } + + // Inverse of the full-range BT.709 matrix the real conversion applies. + const float luma = luma_709(mean); + const float cb = fminf(fmaxf(128.0f + (mean.z - luma) / 1.8556f, 0.0f), 255.0f); + const float cr = fminf(fmaxf(128.0f + (mean.x - luma) / 1.5748f, 0.0f), 255.0f); + uint8_t* sample = uv_plane + static_cast(cy) * uv_pitch + static_cast(cx) * 2; + sample[0] = static_cast(cb + 0.5f); + sample[1] = static_cast(cr + 0.5f); } } // namespace -void launch_mock_nv12_pattern(uint8_t* y_plane, - int y_pitch, - uint8_t* uv_plane, - int uv_pitch, - int width, - int height, - uint32_t sensor_id, - uint32_t frame_index, - cudaStream_t stream) +void launch_mock_gif_nv12(const uint8_t* src_rgba, + int src_pitch, + int src_width, + int src_height, + uint8_t* y_plane, + int y_pitch, + uint8_t* uv_plane, + int uv_pitch, + int width, + int height, + cudaStream_t stream) { - const dim3 block(16, 16); - const dim3 luma_grid((width + block.x - 1) / block.x, (height + block.y - 1) / block.y); - paint_luma<<>>(y_plane, y_pitch, width, height, frame_index); + const float scale = fmaxf(static_cast(width) / src_width, static_cast(height) / src_height); + CoverGeometry geometry{}; + geometry.src_width = src_width; + geometry.src_height = src_height; + geometry.scaled_width = max(width, static_cast(lrintf(src_width * scale))); + geometry.scaled_height = max(height, static_cast(lrintf(src_height * scale))); + geometry.crop_x = (geometry.scaled_width - width) / 2; + geometry.crop_y = (geometry.scaled_height - height) / 2; - const dim3 chroma_grid((width / 2 + block.x - 1) / block.x, (height / 2 + block.y - 1) / block.y); - paint_chroma<<>>(uv_plane, uv_pitch, width, height, sensor_id); + const dim3 block(16, 16); + const dim3 grid((width / 2 + block.x - 1) / block.x, (height / 2 + block.y - 1) / block.y); + paint_gif<<>>( + src_rgba, src_pitch, geometry, y_plane, y_pitch, uv_plane, uv_pitch, width, height); } + } // namespace sensing } // namespace plugins diff --git a/src/plugins/sensing/core/mock_nv12_source.cuh b/src/plugins/sensing/core/mock_nv12_source.cuh index 5db70f07ec..fb7a2e9698 100644 --- a/src/plugins/sensing/core/mock_nv12_source.cuh +++ b/src/plugins/sensing/core/mock_nv12_source.cuh @@ -12,23 +12,27 @@ namespace sensing { /** - * @brief Paint one synthetic ISP0 frame: full-range NV12 in pitch-linear device memory. + * @brief Paint one synthetic ISP0 frame from a GIF frame already in device memory. * - * Stands in for the buffer SIPL would fill, so the real YUV-to-RGBA kernel runs - * on it unchanged. A 16-bit frame counter is drawn top-left as 16 px cells (white - * for a set bit, black for clear, most significant bit first), so a test can tell - * a fresh frame from a stale or torn one by reading pixels. Chroma is tinted by - * sensor so the two eyes are distinguishable. + * The source is scaled to cover the full width x height and centre-cropped, so its aspect ratio is + * never distorted: a uniform scale of max(width / src_width, height / src_height), bilinear when + * enlarging and area-averaged when reducing. The result is full-range BT.709 NV12, the inverse of + * the matrix the real conversion applies, so the colours come back out of that conversion. + * + * @param src_rgba Device pointer to the GIF frame, src_width x src_height RGBA8. + * @param src_pitch Source row stride in bytes. */ -void launch_mock_nv12_pattern(uint8_t* y_plane, - int y_pitch, - uint8_t* uv_plane, - int uv_pitch, - int width, - int height, - uint32_t sensor_id, - uint32_t frame_index, - cudaStream_t stream); +void launch_mock_gif_nv12(const uint8_t* src_rgba, + int src_pitch, + int src_width, + int src_height, + uint8_t* y_plane, + int y_pitch, + uint8_t* uv_plane, + int uv_pitch, + int width, + int height, + cudaStream_t stream); } // namespace sensing } // namespace plugins diff --git a/src/plugins/sensing/core/mock_sipl_camera.cpp b/src/plugins/sensing/core/mock_sipl_camera.cpp index 521dbd9ab3..2fe1d19780 100644 --- a/src/plugins/sensing/core/mock_sipl_camera.cpp +++ b/src/plugins/sensing/core/mock_sipl_camera.cpp @@ -1,18 +1,28 @@ // SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 // -// Mock backend for SiplCamera: synthetic frames from a CUDA pattern generator, -// no SIPL, NvSci or EGL. It keeps the contract the rest of the plugin relies on -// -- per-sensor capture threads, a three-slot RGBA mailbox with a reader lease, -// a TSC stamp shared by every sensor, and failures surfacing from latest() -- -// and runs the real YUV-to-RGBA kernel on the synthetic NV12, so the conversion -// code is exercised too. +// Mock backend for SiplCamera: frames from an animated GIF, no SIPL, NvSci or EGL. +// It keeps the contract the rest of the plugin relies on -- per-sensor capture +// threads, a three-slot RGBA mailbox with a reader lease, a TSC stamp shared by +// every sensor, and failures surfacing from latest() -- and runs the real +// YUV-to-RGBA kernel on a synthetic NV12 surface made from the GIF, so the +// conversion code is exercised too. // -// Environment (read once, at construction): +// What a sensor shows: an animated GIF, chosen in this order +// SENSING_MOCK_GIF_=PATH one sensor, by its module name +// SENSING_MOCK_GIF=PATH every sensor +// (neither) the bundled robot, in mock_assets/ +// A GIF that cannot be loaded, the bundled one included, is an error. It is scaled to cover the +// sensor's configured size and centre-cropped, never stretched, and follows the capture clock: it +// advances with each frame's time, and the sequence number and stamps advance on every capture even +// while the GIF shows one image. +// +// Other environment (read once, at construction): // SENSING_MOCK_FPS=N override the platform config's frame rate // SENSING_MOCK_FAIL_AFTER=N every sensor fails after N frames, as a link drop would // SENSING_MOCK_SCALE_DIV=N shrink the reported geometry (see mock_platform_config.hpp) +#include "mock_gif_source.hpp" #include "mock_nv12_source.cuh" #include "mock_platform_config.hpp" #include "sipl_camera.hpp" @@ -25,6 +35,8 @@ #include #include #include +#include +#include #include #include #include @@ -88,6 +100,8 @@ struct SiplCamera::Impl uint64_t published_timestamp_ns = 0; uint64_t published_capture_tsc_ns = 0; + std::shared_ptr gif; + std::thread frame_thread; // Owns its device memory, so a constructor that throws part-way through @@ -125,6 +139,7 @@ struct SiplCamera::Impl auto sensor = std::make_unique(); sensor->info = info; allocate(*sensor); + sensor->gif = gif_for(info); sensors.push_back(std::move(sensor)); } } @@ -150,6 +165,23 @@ struct SiplCamera::Impl check_cuda(cudaStreamCreateWithFlags(&sensor.stream, cudaStreamNonBlocking), "cudaStreamCreate"); } + /// The GIF a sensor shows. + std::shared_ptr gif_for(const SensorInfo& info) + { + const char* path = std::getenv(("SENSING_MOCK_GIF_" + info.name).c_str()); + if (!path) + path = std::getenv("SENSING_MOCK_GIF"); + const std::string bundled = std::string(SENSING_MOCK_ASSET_DIR) + "/noto-robot.gif"; + if (!path) + path = bundled.c_str(); + + // Sensors showing the same file share one copy on the device. + auto& slot = gif_cache[path]; + if (!slot) + slot = std::make_shared(path); + return slot; + } + Sensor* find(uint32_t id) { for (auto& sensor : sensors) @@ -196,16 +228,17 @@ struct SiplCamera::Impl ++sensor.published_sequence; } - void capture_one(Sensor& sensor, uint64_t frame_index, uint64_t tsc_ns) + void capture_one(Sensor& sensor, double seconds, uint64_t tsc_ns) { const uint32_t write_idx = pick_write_index(sensor); const DeviceBuffer& dest = sensor.buffers[write_idx]; const int width = static_cast(sensor.info.width); const int height = static_cast(sensor.info.height); - launch_mock_nv12_pattern(sensor.luma.ptr, static_cast(sensor.luma.pitch), sensor.chroma.ptr, - static_cast(sensor.chroma.pitch), width, height, sensor.info.id, - static_cast(frame_index), sensor.stream); + launch_mock_gif_nv12(sensor.gif->frame(sensor.gif->frame_index_at(seconds)), + static_cast(sensor.gif->pitch()), static_cast(sensor.gif->width()), + static_cast(sensor.gif->height()), sensor.luma.ptr, static_cast(sensor.luma.pitch), + sensor.chroma.ptr, static_cast(sensor.chroma.pitch), width, height, sensor.stream); // Same call, layout and colour matrix as the SIPL path: full-range // BT.709 semi-planar, from the surface the ISP would have written. @@ -236,7 +269,7 @@ struct SiplCamera::Impl // Every sensor derives its TSC from the same epoch, as the // deserializer's shared fsync does: frame n of each sensor // carries the same capture time. - capture_one(sensor, n, kMockTscBaseNs + static_cast(n * period_ns)); + capture_one(sensor, n * period_ns / 1e9, kMockTscBaseNs + static_cast(n * period_ns)); } } catch (const std::exception& e) @@ -272,6 +305,7 @@ struct SiplCamera::Impl std::vector sensor_infos; std::vector> sensors; std::atomic running{ false }; + std::map> gif_cache; double fps_override = 0.0; uint64_t fail_after_frames = 0; std::mutex failure_mutex; diff --git a/src/plugins/sensing/mock_assets/README.md b/src/plugins/sensing/mock_assets/README.md new file mode 100644 index 0000000000..470f8f0bde --- /dev/null +++ b/src/plugins/sensing/mock_assets/README.md @@ -0,0 +1,21 @@ + + +# Mock camera images + +The mock backend shows these animated GIFs on its simulated sensors. They are not part of the +Apache-2.0 licensed source: they are Animated Noto Emoji, © Google, licensed under +[CC BY 4.0](https://creativecommons.org/licenses/by/4.0/) (`LICENSES/CC-BY-4.0.txt`). + +| File | Emoji | Source | +| --- | --- | --- | +| `noto-robot.gif` | Robot, U+1F916 | `https://fonts.gstatic.com/s/e/notoemoji/latest/1f916/512.gif` | +| `noto-waving-hand.gif` | Waving hand, U+1F44B | `https://fonts.gstatic.com/s/e/notoemoji/latest/1f44b/512.gif` | + +Licence and artwork: . The files are the unmodified +512 px GIFs. The mock composites their transparent background onto a plain gray when it decodes +them; the files themselves are not changed. + +The mock backend is a development tool. Keep these files out of anything shipped to users. diff --git a/src/plugins/sensing/mock_assets/noto-robot.gif b/src/plugins/sensing/mock_assets/noto-robot.gif new file mode 100644 index 0000000000..92d373cfea --- /dev/null +++ b/src/plugins/sensing/mock_assets/noto-robot.gif @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:2a53e434118916a967f30b9968d1e3fd3cf6a4978d39bdb3551c9fb4c8c1fccf +size 840412 diff --git a/src/plugins/sensing/mock_assets/noto-robot.gif.license b/src/plugins/sensing/mock_assets/noto-robot.gif.license new file mode 100644 index 0000000000..e268adf305 --- /dev/null +++ b/src/plugins/sensing/mock_assets/noto-robot.gif.license @@ -0,0 +1,2 @@ +SPDX-FileCopyrightText: Google LLC +SPDX-License-Identifier: CC-BY-4.0 diff --git a/src/plugins/sensing/mock_assets/noto-waving-hand.gif b/src/plugins/sensing/mock_assets/noto-waving-hand.gif new file mode 100644 index 0000000000..fdb2462241 --- /dev/null +++ b/src/plugins/sensing/mock_assets/noto-waving-hand.gif @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:1025c22c790ada1e071c0ace357817a7745c356b2ace3cb0e508329f67ff9fe0 +size 378319 diff --git a/src/plugins/sensing/mock_assets/noto-waving-hand.gif.license b/src/plugins/sensing/mock_assets/noto-waving-hand.gif.license new file mode 100644 index 0000000000..e268adf305 --- /dev/null +++ b/src/plugins/sensing/mock_assets/noto-waving-hand.gif.license @@ -0,0 +1,2 @@ +SPDX-FileCopyrightText: Google LLC +SPDX-License-Identifier: CC-BY-4.0 diff --git a/tests/cpp/plugins/sensing_mock/CMakeLists.txt b/tests/cpp/plugins/sensing_mock/CMakeLists.txt index fbca027d6a..2e3f445bbb 100644 --- a/tests/cpp/plugins/sensing_mock/CMakeLists.txt +++ b/tests/cpp/plugins/sensing_mock/CMakeLists.txt @@ -7,9 +7,11 @@ find_package(CUDAToolkit REQUIRED) add_executable(sensing_mock_tests + gif_fixture.hpp mock_test_support.hpp test_mock_camera.cpp test_mock_encoder.cpp + test_mock_gif.cpp test_mock_pipeline.cpp test_cuda_ipc_consumer.cpp ) @@ -24,6 +26,7 @@ target_link_libraries(sensing_mock_tests target_compile_definitions(sensing_mock_tests PRIVATE SENSING_PLATFORM_CONFIG="${CMAKE_SOURCE_DIR}/src/plugins/sensing/configs/shw5g.json" + SENSING_BUNDLED_GIF_DIR="${CMAKE_SOURCE_DIR}/src/plugins/sensing/mock_assets" ) catch_discover_tests(sensing_mock_tests ADD_TAGS_AS_LABELS) diff --git a/tests/cpp/plugins/sensing_mock/gif_fixture.hpp b/tests/cpp/plugins/sensing_mock/gif_fixture.hpp new file mode 100644 index 0000000000..35370904b1 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/gif_fixture.hpp @@ -0,0 +1,171 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// A minimal GIF89a writer, so tests can make animated GIFs of known content. + +#pragma once + +#include +#include +#include +#include +#include + +namespace sensing_test +{ + +/// Palette indices a fixture frame can use. +enum GifColor : uint8_t +{ + kGifBlack = 0, + kGifRed = 1, + kGifBlue = 2, + kGifWhite = 3, +}; + +struct GifFrame +{ + /// width * height palette indices, row by row. + std::vector pixels; + int delay_centiseconds = 10; +}; + +namespace detail +{ + +class LzwBits +{ +public: + void put(uint32_t code, int bits) + { + m_buffer |= code << m_count; + m_count += bits; + while (m_count >= 8) + { + m_bytes.push_back(static_cast(m_buffer & 0xFF)); + m_buffer >>= 8; + m_count -= 8; + } + } + std::vector finish() + { + if (m_count > 0) + m_bytes.push_back(static_cast(m_buffer & 0xFF)); + return m_bytes; + } + +private: + std::vector m_bytes; + uint32_t m_buffer = 0; + int m_count = 0; +}; + +inline void put16(std::string& out, int value) +{ + out.push_back(static_cast(value & 0xFF)); + out.push_back(static_cast((value >> 8) & 0xFF)); +} + +/// Valid LZW with no compression: each pixel is a literal 9-bit code, and the table is cleared often +/// enough that the code size never has to grow. +inline std::vector lzw_literal(const std::vector& pixels) +{ + constexpr uint32_t kClear = 256; + constexpr uint32_t kEnd = 257; + constexpr size_t kRun = 250; + LzwBits bits; + bits.put(kClear, 9); + for (size_t i = 0; i < pixels.size(); ++i) + { + bits.put(pixels[i], 9); + if ((i + 1) % kRun == 0) + bits.put(kClear, 9); + } + bits.put(kEnd, 9); + return bits.finish(); +} + +} // namespace detail + +inline void write_gif(const std::string& path, int width, int height, const std::vector& frames) +{ + std::string out = "GIF89a"; + detail::put16(out, width); + detail::put16(out, height); + out.push_back(static_cast(0xF7)); // global colour table, 256 entries + out.push_back(0); + out.push_back(0); + + std::string palette(256 * 3, '\0'); + auto set = [&](int index, int r, int g, int b) + { + palette[index * 3] = static_cast(r); + palette[index * 3 + 1] = static_cast(g); + palette[index * 3 + 2] = static_cast(b); + }; + set(kGifBlack, 0, 0, 0); + set(kGifRed, 255, 0, 0); + set(kGifBlue, 0, 0, 255); + set(kGifWhite, 255, 255, 255); + out += palette; + + for (const auto& frame : frames) + { + out += std::string("\x21\xF9\x04\x00", 4); + detail::put16(out, frame.delay_centiseconds); + out += std::string("\x00\x00", 2); + + out.push_back(0x2C); + detail::put16(out, 0); + detail::put16(out, 0); + detail::put16(out, width); + detail::put16(out, height); + out.push_back(0); + + out.push_back(8); // minimum code size + const auto data = detail::lzw_literal(frame.pixels); + for (size_t offset = 0; offset < data.size(); offset += 255) + { + const size_t n = std::min(255, data.size() - offset); + out.push_back(static_cast(n)); + out.append(reinterpret_cast(data.data()) + offset, n); + } + out.push_back(0); + } + out.push_back(0x3B); + + std::ofstream file(path, std::ios::binary); + file.write(out.data(), static_cast(out.size())); +} + +/// A frame of one colour. +inline GifFrame solid_frame(int width, int height, GifColor color, int delay_centiseconds) +{ + return { std::vector(static_cast(width) * height, color), delay_centiseconds }; +} + +/// Left half `left`, right half `right`. +inline GifFrame split_frame(int width, int height, GifColor left, GifColor right, int delay_centiseconds) +{ + GifFrame frame{ std::vector(static_cast(width) * height), delay_centiseconds }; + for (int y = 0; y < height; ++y) + { + for (int x = 0; x < width; ++x) + frame.pixels[static_cast(y) * width + x] = x < width / 2 ? left : right; + } + return frame; +} + +/// A white square of `side` pixels centred on black. +inline GifFrame square_frame(int width, int height, int side, int delay_centiseconds) +{ + GifFrame frame{ std::vector(static_cast(width) * height, kGifBlack), delay_centiseconds }; + for (int y = (height - side) / 2; y < (height + side) / 2; ++y) + { + for (int x = (width - side) / 2; x < (width + side) / 2; ++x) + frame.pixels[static_cast(y) * width + x] = kGifWhite; + } + return frame; +} + +} // namespace sensing_test diff --git a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp index fa7239607f..5137d64cae 100644 --- a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -3,10 +3,13 @@ #pragma once +#include "gif_fixture.hpp" + #include #include #include +#include #include #include #include @@ -26,7 +29,15 @@ class ScopedEnv public: ScopedEnv(const char* name, const char* value) : m_name(name) { - ::setenv(name, value, 1); + set(value); + } + /// Names the variable now; the value is set later with set(). + explicit ScopedEnv(const char* name) : m_name(name) + { + } + void set(const char* value) + { + ::setenv(m_name.c_str(), value, 1); } ~ScopedEnv() { @@ -90,6 +101,72 @@ class TempFile std::string m_path; }; +/// The mock at test size, showing a short GIF that cycles black, red, blue and white every 40 ms. +/// Small frames keep tests fast, and solid frames of known colours show whether a frame is whole. +class FastMock +{ +public: + FastMock() : m_gif(".gif"), m_scale("SENSING_MOCK_SCALE_DIV", "8"), m_source("SENSING_MOCK_GIF") + { + write_gif(m_gif.path(), 16, 16, + { solid_frame(16, 16, kGifBlack, 4), solid_frame(16, 16, kGifRed, 4), + solid_frame(16, 16, kGifBlue, 4), solid_frame(16, 16, kGifWhite, 4) }); + m_source.set(m_gif.path().c_str()); + } + +private: + TempFile m_gif; + ScopedEnv m_scale; + ScopedEnv m_source; +}; + +struct Rgb +{ + int r; + int g; + int b; +}; + +/// One frame copied to the host. +struct Shot +{ + uint32_t width = 0; + uint32_t height = 0; + uint64_t sequence = 0; + std::vector rgba; + + Rgb at(uint32_t x, uint32_t y) const + { + const uint8_t* px = &rgba[(static_cast(y) * width + x) * 4]; + return { px[0], px[1], px[2] }; + } +}; + +inline bool near(const Rgb& c, int r, int g, int b, int tolerance = 12) +{ + return std::abs(c.r - r) <= tolerance && std::abs(c.g - g) <= tolerance && std::abs(c.b - b) <= tolerance; +} + +/// Black, red, blue or white: the colours of the frames FastMock plays. +inline bool is_cycle_colour(const Rgb& c) +{ + return near(c, 0, 0, 0) || near(c, 255, 0, 0) || near(c, 0, 0, 255) || near(c, 255, 255, 255, 14); +} + +/// True when the corners and centre agree, as in a frame of one solid colour. A frame torn between two +/// captures would show two. +inline bool is_whole(const Shot& shot) +{ + const Rgb centre = shot.at(shot.width / 2, shot.height / 2); + for (const Rgb& corner : { shot.at(0, 0), shot.at(shot.width - 1, 0), shot.at(0, shot.height - 1), + shot.at(shot.width - 1, shot.height - 1) }) + { + if (!near(corner, centre.r, centre.g, centre.b, 14)) + return false; + } + return is_cycle_colour(centre); +} + /// Copies a device RGBA frame to the host as tightly packed rows. inline std::vector download_rgba(uintptr_t device_ptr, size_t pitch, uint32_t width, uint32_t height) { @@ -102,19 +179,23 @@ inline std::vector download_rgba(uintptr_t device_ptr, size_t pitch, ui return host; } -/// Reads the 16-bit frame counter the mock paints top-left: 16 px cells, most significant bit first. -inline uint32_t read_counter(const std::vector& rgba, uint32_t width) +/// Waits for one frame from `sensor` and copies it to the host. +inline Shot grab(plugins::sensing::SiplCamera& camera, uint32_t sensor) { - uint32_t value = 0; - for (int bit = 0; bit < 16; ++bit) - { - const size_t x = 8 + 16 * static_cast(bit); - const size_t y = 8; - const uint8_t* px = &rgba[(y * width + x) * 4]; - const bool on = px[0] > 127 && px[1] > 127 && px[2] > 127; - value = (value << 1) | (on ? 1u : 0u); - } - return value; + Shot shot; + REQUIRE(wait_for( + [&] + { + auto view = camera.latest(sensor); + if (!view) + return false; + shot.width = view->width; + shot.height = view->height; + shot.sequence = view->sequence; + shot.rgba = download_rgba(view->ptr, view->pitch, view->width, view->height); + return true; + })); + return shot; } /// One NAL unit of the mock encoder's Annex-B output. diff --git a/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp b/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp index 6ee83488eb..95cf117abc 100644 --- a/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp +++ b/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp @@ -203,7 +203,7 @@ SensingConfig make_config() TEST_CASE("a consumer receives a described, mappable allocation", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; std::vector streams(1); streams[0].sensor_id = 1; streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); @@ -223,9 +223,9 @@ TEST_CASE("a consumer receives a described, mappable allocation", "[unit][gpu][s CHECK(hello.total_bytes >= hello.slot_stride * hello.slot_count); } -TEST_CASE("published frames are fresh pixels in order", "[unit][gpu][sensing]") +TEST_CASE("published frames are whole, in order, and the picture moves", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; std::vector streams(1); streams[0].sensor_id = 0; streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); @@ -233,31 +233,39 @@ TEST_CASE("published frames are fresh pixels in order", "[unit][gpu][sensing]") SensingCamera camera(make_config(), streams, std::make_unique(streams)); UpdateLoop loop(camera); IpcConsumer consumer(streams[0].ipc_socket_path); + REQUIRE(consumer.hello().pitch == static_cast(consumer.hello().width) * 4); uint64_t last_sequence = 0; uint64_t last_stamp = 0; - int64_t last_counter = -1; + std::vector colours; for (int i = 0; i < 20; ++i) { ipc::FrameReady frame{}; REQUIRE(consumer.next_frame(frame, 2000ms)); - const auto pixels = consumer.read_slot(frame.slot); - const int64_t counter = read_counter(pixels, consumer.hello().width); + Shot shot; + shot.width = consumer.hello().width; + shot.height = consumer.hello().height; + shot.rgba = consumer.read_slot(frame.slot); consumer.release(frame.slot); CHECK(frame.sequence > last_sequence); CHECK(frame.timestamp_ns > last_stamp); - // Each published frame is a new capture, never a repeat of the previous one. - CHECK(counter > last_counter); + CHECK(is_whole(shot)); // not torn between two captures last_sequence = frame.sequence; last_stamp = frame.timestamp_ns; - last_counter = counter; + colours.push_back(shot.at(shot.width / 2, shot.height / 2)); } + + // Twenty frames span several 40 ms GIF frames, so more than one colour appears. + bool changed = false; + for (const Rgb& c : colours) + changed = changed || !near(c, colours[0].r, colours[0].g, colours[0].b, 30); + CHECK(changed); } TEST_CASE("a consumer that holds its slots cannot stall the producer", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; std::vector streams(1); streams[0].sensor_id = 0; streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); diff --git a/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp b/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp index e202bc6600..9fc23779c1 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp @@ -76,9 +76,9 @@ TEST_CASE("camera construction validates the masks", "[unit][gpu][sensing]") CHECK_THROWS_WITH(SiplCamera(config), Catch::Matchers::ContainsSubstring("no modules left")); } -TEST_CASE("frames arrive with a fresh counter, rising sequence and rising stamps", "[unit][gpu][sensing]") +TEST_CASE("frames arrive whole, with rising sequence and stamps, and the picture moves", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; SiplCamera camera(make_sipl_config()); REQUIRE(camera.sensors().size() == 2); camera.start(); @@ -88,7 +88,7 @@ TEST_CASE("frames arrive with a fresh counter, rising sequence and rising stamps INFO("sensor " << sensor.id); uint64_t last_sequence = 0; uint64_t last_stamp = 0; - int seen = 0; + std::vector shots; REQUIRE(wait_for( [&] { @@ -103,18 +103,32 @@ TEST_CASE("frames arrive with a fresh counter, rising sequence and rising stamps last_sequence = view->sequence; last_stamp = view->timestamp_ns; - // The mock numbers frames from 0 and publishes the sequence from 1. - const auto pixels = download_rgba(view->ptr, view->pitch, view->width, view->height); - CHECK(read_counter(pixels, view->width) == ((view->sequence - 1) & 0xFFFF)); - return ++seen >= 6; + Shot shot; + shot.width = view->width; + shot.height = view->height; + shot.sequence = view->sequence; + shot.rgba = download_rgba(view->ptr, view->pitch, view->width, view->height); + CHECK(is_whole(shot)); + shots.push_back(std::move(shot)); + return shots.size() >= 8; })); + + // Eight captures span more than one 40 ms GIF frame, so the colour must have changed. + bool changed = false; + for (const auto& shot : shots) + { + const Rgb first = shots[0].at(shots[0].width / 2, shots[0].height / 2); + const Rgb here = shot.at(shot.width / 2, shot.height / 2); + changed = changed || !near(here, first.r, first.g, first.b, 30); + } + CHECK(changed); } camera.stop(); } TEST_CASE("both sensors carry the same capture TSC for the same frame", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; SiplCamera camera(make_sipl_config()); camera.start(); @@ -148,23 +162,38 @@ TEST_CASE("both sensors carry the same capture TSC for the same frame", "[unit][ TEST_CASE("a leased slot is not overwritten while the producer keeps running", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; SiplCamera camera(make_sipl_config()); camera.start(); std::optional held; REQUIRE(wait_for([&] { return (held = camera.latest(0)).has_value(); })); - const uint32_t counter = read_counter(download_rgba(held->ptr, held->pitch, held->width, held->height), held->width); + const auto before = download_rgba(held->ptr, held->pitch, held->width, held->height); + + // The GIF changes colour every 40 ms and the other two slots are rewritten on every capture. If + // the lease failed, the held slot would soon show a different colour. + const auto end = std::chrono::steady_clock::now() + 250ms; + while (std::chrono::steady_clock::now() < end) + { + CHECK(download_rgba(held->ptr, held->pitch, held->width, held->height) == before); + std::this_thread::sleep_for(5ms); + } - // Several frame periods pass without another latest(): the lease must hold. - std::this_thread::sleep_for(150ms); - CHECK(read_counter(download_rgba(held->ptr, held->pitch, held->width, held->height), held->width) == counter); + // The producer did keep going while the slot was held. + uint64_t latest_sequence = held->sequence; + REQUIRE(wait_for( + [&] + { + if (auto view = camera.latest(1)) + latest_sequence = std::max(latest_sequence, view->sequence); + return latest_sequence >= held->sequence + 5; + })); camera.stop(); } TEST_CASE("an injected link failure surfaces from latest()", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; ScopedEnv fail("SENSING_MOCK_FAIL_AFTER", "3"); SiplCamera camera(make_sipl_config()); camera.start(); @@ -190,7 +219,7 @@ TEST_CASE("an injected link failure surfaces from latest()", "[unit][gpu][sensin TEST_CASE("an unknown sensor index is an error, and stop is idempotent", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; SiplCamera camera(make_sipl_config()); CHECK_THROWS_AS(camera.latest(7), std::runtime_error); camera.start(); diff --git a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp new file mode 100644 index 0000000000..1d1043489e --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp @@ -0,0 +1,386 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "gif_fixture.hpp" +#include "mock_test_support.hpp" + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +using namespace plugins::sensing; +using namespace sensing_test; +using namespace std::chrono_literals; + +namespace +{ + +bool is_white(const Rgb& c) +{ + return near(c, 255, 255, 255, 14); +} + +// Number of consecutive white pixels along a row or column starting at (x, y). +uint32_t white_run_x(const Shot& shot, uint32_t y) +{ + uint32_t count = 0; + for (uint32_t x = 0; x < shot.width; ++x) + count += is_white(shot.at(x, y)) ? 1 : 0; + return count; +} + +uint32_t white_run_y(const Shot& shot, uint32_t x) +{ + uint32_t count = 0; + for (uint32_t y = 0; y < shot.height; ++y) + count += is_white(shot.at(x, y)) ? 1 : 0; + return count; +} + +} // namespace + +TEST_CASE("GIF frame ends clamp tiny delays to 20 ms and accumulate", "[unit][sensing]") +{ + const auto ends = gif_frame_ends({ 0, 10, 100, 20 }); + REQUIRE(ends.size() == 4); + CHECK(ends[0] == 0.02); + CHECK(ends[1] == 0.04); + CHECK(ends[2] == 0.14); + CHECK(ends[3] == 0.16); +} + +TEST_CASE("GIF frame index follows elapsed time and loops", "[unit][sensing]") +{ + const auto ends = gif_frame_ends({ 100, 100, 200 }); // 0.1, 0.2, 0.4 + CHECK(gif_frame_index(ends, 0.0) == 0); + CHECK(gif_frame_index(ends, 0.099) == 0); + CHECK(gif_frame_index(ends, 0.1) == 1); + CHECK(gif_frame_index(ends, 0.25) == 2); + CHECK(gif_frame_index(ends, 0.399) == 2); + CHECK(gif_frame_index(ends, 0.41) == 0); // loops + CHECK(gif_frame_index(ends, 0.65) == 2); + CHECK(gif_frame_index({}, 1.0) == 0); +} + +TEST_CASE("a GIF that is missing or is not a GIF fails construction", "[unit][gpu][sensing]") +{ + CHECK_THROWS_WITH(MockGifSource("/nonexistent/none.gif"), Catch::Matchers::ContainsSubstring("cannot open")); + + TempFile text(".gif"); + std::ofstream(text.path()) << "this is not a GIF"; + CHECK_THROWS_WITH(MockGifSource(text.path()), Catch::Matchers::ContainsSubstring("cannot decode")); + + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv gif("SENSING_MOCK_GIF", "/nonexistent/none.gif"); + CHECK_THROWS_AS(SiplCamera(make_sipl_config()), std::runtime_error); +} + +TEST_CASE("the camera shows the GIF frame for its capture time while sequence and stamps keep advancing", + "[unit][gpu][sensing]") +{ + TempFile gif(".gif"); + // Delays are in centiseconds: 4 is 40 ms. + write_gif(gif.path(), 16, 16, { solid_frame(16, 16, kGifRed, 4), solid_frame(16, 16, kGifBlue, 4) }); + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedEnv fps("SENSING_MOCK_FPS", "50"); // 20 ms per capture: two captures per GIF frame + + SiplCamera camera(make_sipl_config()); + camera.start(); + + std::vector shots; + uint64_t last_sequence = 0; + REQUIRE(wait_for( + [&] + { + if (auto view = camera.latest(0)) + { + Shot shot; + shot.width = view->width; + shot.height = view->height; + shot.sequence = view->sequence; + shot.rgba = download_rgba(view->ptr, view->pitch, view->width, view->height); + CHECK(shot.sequence > last_sequence); + last_sequence = shot.sequence; + shots.push_back(std::move(shot)); + } + return shots.size() >= 14; + })); + camera.stop(); + + // Capture n shows GIF time n / 50 s: red until 0.04 s, blue until 0.08 s, then red again. + int repeats = 0; + int red = 0; + int blue = 0; + for (size_t i = 0; i < shots.size(); ++i) + { + const double t = (shots[i].sequence - 1) / 50.0; + const double phase = std::fmod(t, 0.08); + // A capture exactly on a frame boundary could land on either side by rounding. + if (std::abs(phase - 0.04) < 1e-6 || phase < 1e-6 || phase > 0.08 - 1e-6) + continue; + const bool expect_red = phase < 0.04; + const Rgb centre = shots[i].at(shots[i].width / 2, shots[i].height / 2); + INFO("capture " << shots[i].sequence << " t=" << t); + CHECK((expect_red ? near(centre, 255, 0, 0) : near(centre, 0, 0, 255))); + (expect_red ? red : blue)++; + if (i > 0 && near(centre, 255, 0, 0) == near(shots[i - 1].at(0, 0), 255, 0, 0)) + ++repeats; + } + CHECK(red > 0); + CHECK(blue > 0); + // The same image on consecutive captures is normal; it must not stall the sequence. + CHECK(repeats > 0); +} + +TEST_CASE("an enlarged GIF keeps its proportions on the sensor's geometry", "[unit][gpu][sensing]") +{ + // 32x32 with a centred 16x16 square, shown on 320x248: scale 10 fills the width, the height is + // cropped from 320 to 248, so the square stays square. + TempFile gif(".gif"); + write_gif(gif.path(), 32, 32, { square_frame(32, 32, 16, 100) }); + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot shot = grab(camera, 0); + camera.stop(); + + REQUIRE(shot.width == 320); + REQUIRE(shot.height == 248); + const uint32_t cx = shot.width / 2; + const uint32_t cy = shot.height / 2; + CHECK(is_white(shot.at(cx, cy))); + CHECK(near(shot.at(30, cy), 0, 0, 0)); + CHECK(near(shot.at(cx, 10), 0, 0, 0)); + // 16 px at scale 10 is about 160 px both ways, a little less where bilinear softens the edge. + // Equal extents in x and y mean no independent stretch. + const int run_x = static_cast(white_run_x(shot, cy)); + const int run_y = static_cast(white_run_y(shot, cx)); + CHECK(std::abs(run_x - 160) <= 14); + CHECK(std::abs(run_x - run_y) <= 3); +} + +TEST_CASE("a wide GIF is reduced and cropped to the sensor's aspect", "[unit][gpu][sensing]") +{ + // 640x320 (2:1), left red and right blue, on 320x248 (about 1.29:1): scale 0.775 makes the image + // 496x248, the sides are cropped, and the red/blue boundary stays in the middle. + TempFile gif(".gif"); + write_gif(gif.path(), 640, 320, { split_frame(640, 320, kGifRed, kGifBlue, 100) }); + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot shot = grab(camera, 0); + camera.stop(); + + REQUIRE(shot.width == 320); + REQUIRE(shot.height == 248); + CHECK(near(shot.at(60, 124), 255, 0, 0)); + CHECK(near(shot.at(150, 124), 255, 0, 0)); + CHECK(near(shot.at(172, 124), 0, 0, 255)); + CHECK(near(shot.at(260, 124), 0, 0, 255)); +} + +TEST_CASE("a tall GIF is cropped top and bottom", "[unit][gpu][sensing]") +{ + // 100x400 (1:4) on 320x248: scale 3.2 gives 320x1280, so only the middle 248 rows show. + // The GIF is white in the middle band and red elsewhere, so the visible window is all white. + TempFile gif(".gif"); + GifFrame frame{ std::vector(100 * 400, kGifRed), 100 }; + for (int y = 150; y < 250; ++y) + std::fill_n(frame.pixels.begin() + y * 100, 100, kGifWhite); + write_gif(gif.path(), 100, 400, { frame }); + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot shot = grab(camera, 0); + camera.stop(); + + REQUIRE(shot.height == 248); + CHECK(is_white(shot.at(160, 124))); + CHECK(is_white(shot.at(160, 20))); + CHECK(is_white(shot.at(160, 228))); +} + +TEST_CASE("a large reduction averages the source instead of skipping pixels", "[unit][gpu][sensing]") +{ + // 200x200 reduced to 80x62 (scale 0.4): left red, right blue. + TempFile gif(".gif"); + write_gif(gif.path(), 200, 200, { split_frame(200, 200, kGifRed, kGifBlue, 100) }); + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "32"); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot shot = grab(camera, 0); + camera.stop(); + + REQUIRE(shot.width == 80); + REQUIRE(shot.height == 62); + CHECK(near(shot.at(10, 31), 255, 0, 0)); + CHECK(near(shot.at(70, 31), 0, 0, 255)); +} + +TEST_CASE("without a scale override the GIF is shown at the sensor's configured size", "[unit][gpu][sensing]") +{ + TempFile gif(".gif"); + write_gif(gif.path(), 32, 32, { square_frame(32, 32, 16, 100) }); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedEnv fps("SENSING_MOCK_FPS", "5"); + + SiplCamera camera(make_sipl_config()); + const auto& sensors = camera.sensors(); + REQUIRE(sensors.size() == 2); + camera.start(); + const Shot shot = grab(camera, 0); + camera.stop(); + + // The vendor config's 2560x1984, not the GIF's 32x32. + CHECK(shot.width == sensors[0].width); + CHECK(shot.height == sensors[0].height); + CHECK(shot.width == 2560); + CHECK(shot.height == 1984); + // Scale 80: the square is about 1280 px wide and tall, centred, with soft edges. + CHECK(is_white(shot.at(shot.width / 2, shot.height / 2))); + const int run_x = static_cast(white_run_x(shot, shot.height / 2)); + const int run_y = static_cast(white_run_y(shot, shot.width / 2)); + CHECK(std::abs(run_x - 1280) <= 110); + CHECK(std::abs(run_x - run_y) <= 4); +} + +TEST_CASE("a sensor's GIF can be chosen by its module name", "[unit][gpu][sensing]") +{ + TempFile red(".gif"); + TempFile blue(".gif"); + write_gif(red.path(), 16, 16, { solid_frame(16, 16, kGifRed, 100) }); + write_gif(blue.path(), 16, 16, { solid_frame(16, 16, kGifBlue, 100) }); + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv everyone("SENSING_MOCK_GIF", red.path().c_str()); + ScopedEnv one("SENSING_MOCK_GIF_SHW5G_MIX_L3", blue.path().c_str()); + + SiplCamera camera(make_sipl_config()); + REQUIRE(camera.sensors().size() == 2); + REQUIRE(camera.sensors()[0].name == "SHW5G_MIX_L2"); + REQUIRE(camera.sensors()[1].name == "SHW5G_MIX_L3"); + camera.start(); + const Shot first = grab(camera, 0); + const Shot second = grab(camera, 1); + camera.stop(); + + CHECK(near(first.at(first.width / 2, first.height / 2), 255, 0, 0)); + CHECK(near(second.at(second.width / 2, second.height / 2), 0, 0, 255)); +} + +namespace +{ + +std::string bundled(const char* name) +{ + return std::string(SENSING_BUNDLED_GIF_DIR) + "/" + name; +} + +} // namespace + +TEST_CASE("the bundled GIFs decode, animate and have their transparency composited", "[unit][gpu][sensing]") +{ + for (const auto& [name, frames] : { std::pair{ "noto-robot.gif", 72 }, + std::pair{ "noto-waving-hand.gif", 46 } }) + { + INFO(name); + MockGifSource gif(bundled(name)); + CHECK(gif.width() == 512); + CHECK(gif.height() == 512); + REQUIRE(gif.frame_count() == frames); + + auto pixel = [&](size_t frame, uint32_t x, uint32_t y) + { + std::array px{}; + REQUIRE(cudaMemcpy(px.data(), gif.frame(frame) + (static_cast(y) * gif.width() + x) * 4, 4, + cudaMemcpyDeviceToHost) == cudaSuccess); + return px; + }; + + // The corner is transparent in the file, so it shows the background, fully opaque. + const auto corner = pixel(0, 0, 0); + CHECK(corner[0] == MockGifSource::kBackground); + CHECK(corner[1] == MockGifSource::kBackground); + CHECK(corner[2] == MockGifSource::kBackground); + CHECK(corner[3] == 255); + + // The animation changes the picture: some frame differs from the first at the centre. + const auto first = pixel(0, 256, 256); + bool changed = false; + for (size_t f = 1; f < gif.frame_count() && !changed; ++f) + changed = pixel(f, 256, 256) != first || pixel(f, 200, 300) != pixel(0, 200, 300); + CHECK(changed); + } +} + +TEST_CASE("by default every sensor shows the animated bundled robot", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fps("SENSING_MOCK_FPS", "20"); + SiplCamera camera(make_sipl_config()); + camera.start(); + + // Captures a second apart show different parts of the 2.5 s animation. + std::vector shots; + uint64_t last = 0; + REQUIRE(wait_for( + [&] + { + if (auto view = camera.latest(0)) + { + if (view->sequence >= last + 10 || shots.empty()) + { + Shot shot; + shot.width = view->width; + shot.height = view->height; + shot.sequence = view->sequence; + shot.rgba = download_rgba(view->ptr, view->pitch, view->width, view->height); + shots.push_back(std::move(shot)); + last = view->sequence; + } + } + return shots.size() >= 4; + }, + 10s)); + camera.stop(); + + // Transparent corners of the emoji read as the background gray, and the picture moves. + for (const auto& shot : shots) + CHECK(near(shot.at(2, 2), MockGifSource::kBackground, MockGifSource::kBackground, MockGifSource::kBackground, 8)); + size_t differing = 0; + for (size_t i = 0; i < shots[0].rgba.size(); i += 4) + differing += shots[0].rgba[i] != shots[2].rgba[i] ? 1 : 0; + CHECK(differing > 500); +} + +TEST_CASE("the waving hand can be given to one sensor by module name", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv one("SENSING_MOCK_GIF_SHW5G_MIX_L3", bundled("noto-waving-hand.gif").c_str()); + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot robot = grab(camera, 0); + const Shot hand = grab(camera, 1); + camera.stop(); + + // Different pictures: the robot on one sensor, the hand on the other. + size_t differing = 0; + for (size_t i = 0; i < robot.rgba.size(); i += 4) + differing += std::abs(robot.rgba[i] - hand.rgba[i]) > 20 ? 1 : 0; + CHECK(differing > 2000); +} diff --git a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp index b0371f5d77..df7b3e47a8 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp @@ -99,7 +99,7 @@ void run_for(SensingCamera& camera, std::chrono::milliseconds duration) TEST_CASE("an unknown sensor index is reported with the available pipelines", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; std::vector streams(1); streams[0].sensor_id = 5; streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); @@ -109,7 +109,7 @@ TEST_CASE("an unknown sensor index is reported with the available pipelines", "[ TEST_CASE("an encoded stream writes Annex-B data and one metadata record per unit", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; TempFile output_file(".h264"); const std::string& output = output_file.path(); @@ -151,7 +151,7 @@ TEST_CASE("an encoded stream writes Annex-B data and one metadata record per uni TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; std::vector streams(2); streams[0].sensor_id = 0; streams[0].ipc_socket_path = unique_path("sensing_mock", ".sock"); @@ -189,7 +189,7 @@ TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", " TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit][gpu][sensing]") { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + FastMock fast; TempFile output_file(".h264"); const std::string& output = output_file.path(); From 89abf4d4e3a143d2f02d05f47d9e796e3287db4e Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sun, 25 Oct 2026 00:00:05 -0400 Subject: [PATCH 05/10] Give the mock a three-camera rig with role-correct pictures Add a mock rig config (S56C stereo head camera, two SHF3L wrist cameras) with each camera's real size and aspect. Pictures are chosen by module name, not pipeline index: the head eyes show the robot offset by a fixed stereo disparity, the right wrist is the left wrist's picture flipped. A module with no known picture must be given one by the environment. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/README.md | 27 +- src/plugins/sensing/core/mock_nv12_source.cu | 16 +- src/plugins/sensing/core/mock_nv12_source.cuh | 4 + src/plugins/sensing/core/mock_sipl_camera.cpp | 92 ++++- src/plugins/sensing/mock_assets/mock_rig.json | 335 +++++++++++++++++ tests/cpp/plugins/sensing_mock/CMakeLists.txt | 1 + .../plugins/sensing_mock/test_mock_gif.cpp | 342 ++++++++++++++++-- 7 files changed, 757 insertions(+), 60 deletions(-) create mode 100644 src/plugins/sensing/mock_assets/mock_rig.json diff --git a/src/plugins/sensing/README.md b/src/plugins/sensing/README.md index e6b63cc16e..fc5f86d482 100644 --- a/src/plugins/sensing/README.md +++ b/src/plugins/sensing/README.md @@ -74,20 +74,29 @@ fixed pool of input buffers that makes `submit()` refuse frames when the encoder a warm-up delay before the first unit, and capture stamps that come back at microsecond resolution. -Every sensor shows an animated GIF, by default the bundled robot in [`mock_assets/`](mock_assets/) -(Animated Noto Emoji, CC BY 4.0; see its README for the licence and source). The GIF plays on the -capture clock, looping, scaled to cover the sensor's configured size and centre-cropped so it is -never stretched (bilinear when enlarging, area-averaged when reducing). Transparent pixels show a -plain gray. The sequence number and stamps advance on every capture even while the GIF holds one -image. The GIF's own size and aspect do not matter; the frame size always comes from the platform -config. A GIF that cannot be loaded, the bundled one included, is an error: there is no fallback. +Every sensor shows an animated GIF chosen by its **module name** (never its pipeline index, which the +link mask renumbers), from the bundled Animated Noto Emoji in [`mock_assets/`](mock_assets/) (CC BY 4.0; +see its README): `S56C_HEAD` the robot, `SHF3L_LEFT` the waving hand, `SHF3L_RIGHT` the hand flipped left +to right. The two sensors of `S56C_HEAD` are its left and right eyes; their pictures are offset against +each other by a fixed 2% of the width, as a 60 mm baseline gives. A module not in that table needs +`SENSING_MOCK_GIF` or `SENSING_MOCK_GIF_`. The GIF plays on the capture clock, looping, scaled +to cover the sensor's configured size and centre-cropped so it is never stretched (bilinear when +enlarging, area-averaged when reducing). Transparent pixels show a plain gray. The GIF's own size and +aspect do not matter; the frame size always comes from the platform config. A GIF that cannot be loaded, +the bundled one included, is an error: there is no fallback. + +[`mock_assets/mock_rig.json`](mock_assets/mock_rig.json) describes a three-camera rig for the mock only +(not a SIPL config): one S56C stereo head camera with two 2560x1984 @ 30 eyes, and two SHF3L wrist +cameras at 1920x1536 @ 30. Use `--platform-config=.../mock_rig.json --config-name=MOCK_S56C_2_SHF3L +--link-masks='0x0000 0x0111'`. Environment, read once at construction: | Variable | Effect | | --- | --- | -| `SENSING_MOCK_GIF` | path of an animated GIF shown by every sensor (default: the bundled robot) | -| `SENSING_MOCK_GIF_` | a GIF for one sensor, by module name (`SENSING_MOCK_GIF_SHW5G_MIX_L3`); wins over `SENSING_MOCK_GIF`. The bundled `noto-waving-hand.gif` suits a wrist camera | +| `SENSING_MOCK_GIF` | path of an animated GIF shown by every sensor, as it is (default: by module, above) | +| `SENSING_MOCK_GIF_` | a GIF for one sensor, by module name (`SENSING_MOCK_GIF_SHW5G_MIX_L3`); wins over `SENSING_MOCK_GIF` | +| `SENSING_MOCK_FLIP_` | `1` or `0`: flip that module's picture left to right, or not, whatever its GIF | | `SENSING_MOCK_FPS` | override the platform config's frame rate | | `SENSING_MOCK_SCALE_DIV` | divide the reported width and height (kept even); useful for tests | | `SENSING_MOCK_FAIL_AFTER` | every sensor fails after N frames, as a link drop would | diff --git a/src/plugins/sensing/core/mock_nv12_source.cu b/src/plugins/sensing/core/mock_nv12_source.cu index ba0e2c8351..188f1c5113 100644 --- a/src/plugins/sensing/core/mock_nv12_source.cu +++ b/src/plugins/sensing/core/mock_nv12_source.cu @@ -36,8 +36,8 @@ __device__ float3 sample_cover(const uint8_t* src, int pitch, const CoverGeometr { const float sx = static_cast(g.src_width) / g.scaled_width; const float sy = static_cast(g.src_height) / g.scaled_height; - const int X = x + g.crop_x; - const int Y = y + g.crop_y; + const int X = min(max(x + g.crop_x, 0), g.scaled_width - 1); + const int Y = min(max(y + g.crop_y, 0), g.scaled_height - 1); if (sx <= 1.0f && sy <= 1.0f) { @@ -95,7 +95,9 @@ __global__ void paint_gif(const uint8_t* src, uint8_t* uv_plane, int uv_pitch, int width, - int height) + int height, + bool mirror, + int shift_x) { const int cx = blockIdx.x * blockDim.x + threadIdx.x; const int cy = blockIdx.y * blockDim.y + threadIdx.y; @@ -109,7 +111,9 @@ __global__ void paint_gif(const uint8_t* src, { const int x = 2 * cx + dx; const int y = 2 * cy + dy; - const float3 rgb = sample_cover(src, src_pitch, geometry, x, y); + // A mirrored view samples the picture right to left; the output pixel stays where it is. + const int sx = (mirror ? width - 1 - x : x) - shift_x; + const float3 rgb = sample_cover(src, src_pitch, geometry, sx, y); y_plane[static_cast(y) * y_pitch + x] = static_cast(luma_709(rgb) + 0.5f); mean.x += 0.25f * rgb.x; mean.y += 0.25f * rgb.y; @@ -138,6 +142,8 @@ void launch_mock_gif_nv12(const uint8_t* src_rgba, int uv_pitch, int width, int height, + bool mirror, + int shift_x, cudaStream_t stream) { const float scale = fmaxf(static_cast(width) / src_width, static_cast(height) / src_height); @@ -152,7 +158,7 @@ void launch_mock_gif_nv12(const uint8_t* src_rgba, const dim3 block(16, 16); const dim3 grid((width / 2 + block.x - 1) / block.x, (height / 2 + block.y - 1) / block.y); paint_gif<<>>( - src_rgba, src_pitch, geometry, y_plane, y_pitch, uv_plane, uv_pitch, width, height); + src_rgba, src_pitch, geometry, y_plane, y_pitch, uv_plane, uv_pitch, width, height, mirror, shift_x); } diff --git a/src/plugins/sensing/core/mock_nv12_source.cuh b/src/plugins/sensing/core/mock_nv12_source.cuh index fb7a2e9698..1a79a37a58 100644 --- a/src/plugins/sensing/core/mock_nv12_source.cuh +++ b/src/plugins/sensing/core/mock_nv12_source.cuh @@ -21,6 +21,8 @@ namespace sensing * * @param src_rgba Device pointer to the GIF frame, src_width x src_height RGBA8. * @param src_pitch Source row stride in bytes. + * @param mirror Show the picture flipped left to right. + * @param shift_x Pixels the picture moves right after any flip; the edge pixel repeats into the gap. */ void launch_mock_gif_nv12(const uint8_t* src_rgba, int src_pitch, @@ -32,6 +34,8 @@ void launch_mock_gif_nv12(const uint8_t* src_rgba, int uv_pitch, int width, int height, + bool mirror, + int shift_x, cudaStream_t stream); } // namespace sensing diff --git a/src/plugins/sensing/core/mock_sipl_camera.cpp b/src/plugins/sensing/core/mock_sipl_camera.cpp index 2fe1d19780..160cc4427d 100644 --- a/src/plugins/sensing/core/mock_sipl_camera.cpp +++ b/src/plugins/sensing/core/mock_sipl_camera.cpp @@ -8,14 +8,19 @@ // YUV-to-RGBA kernel on a synthetic NV12 surface made from the GIF, so the // conversion code is exercised too. // -// What a sensor shows: an animated GIF, chosen in this order -// SENSING_MOCK_GIF_=PATH one sensor, by its module name -// SENSING_MOCK_GIF=PATH every sensor -// (neither) the bundled robot, in mock_assets/ -// A GIF that cannot be loaded, the bundled one included, is an error. It is scaled to cover the -// sensor's configured size and centre-cropped, never stretched, and follows the capture clock: it -// advances with each frame's time, and the sequence number and stamps advance on every capture even -// while the GIF shows one image. +// What a sensor shows is chosen by module name, never by pipeline index (which the link mask renumbers): +// S56C_HEAD the robot, the two eyes offset against each other by a fixed disparity +// SHF3L_LEFT the waving hand +// SHF3L_RIGHT the waving hand, flipped left to right +// SHW5G_MIX_L2 / SHW5G_MIX_L3 as SHF3L_LEFT / SHF3L_RIGHT +// A module not listed needs SENSING_MOCK_GIF or SENSING_MOCK_GIF_. Environment, in order: +// SENSING_MOCK_GIF_=PATH one module's GIF +// SENSING_MOCK_GIF=PATH every sensor's GIF +// SENSING_MOCK_FLIP_=1|0 flip one module, whatever its GIF +// A GIF named by the environment is shown as it is, with no eye offset. A GIF that cannot be loaded, +// the bundled one included, is an error. It is scaled to cover the sensor's configured size and +// centre-cropped, never stretched, and follows the capture clock: it advances with each frame's time, +// and the sequence number and stamps advance on every capture even while the GIF shows one image. // // Other environment (read once, at construction): // SENSING_MOCK_FPS=N override the platform config's frame rate @@ -53,6 +58,10 @@ namespace // Arbitrary offset so nothing can mistake the TSC for CLOCK_MONOTONIC. constexpr uint64_t kMockTscBaseNs = 1'000'000'000'000'000ull; +// Total horizontal offset between the two eyes, as a fraction of the width: a 60 mm baseline seen at +// about half a metre through a 140 degree lens. +constexpr double kEyeDisparity = 0.02; + uint64_t monotonic_ns() { return static_cast( @@ -101,6 +110,8 @@ struct SiplCamera::Impl uint64_t published_capture_tsc_ns = 0; std::shared_ptr gif; + bool mirror = false; + int shift_x = 0; std::thread frame_thread; @@ -139,7 +150,7 @@ struct SiplCamera::Impl auto sensor = std::make_unique(); sensor->info = info; allocate(*sensor); - sensor->gif = gif_for(info); + choose_picture(info, *sensor); sensors.push_back(std::move(sensor)); } } @@ -165,21 +176,67 @@ struct SiplCamera::Impl check_cuda(cudaStreamCreateWithFlags(&sensor.stream, cudaStreamNonBlocking), "cudaStreamCreate"); } - /// The GIF a sensor shows. - std::shared_ptr gif_for(const SensorInfo& info) + /// Sets the GIF a sensor shows, whether it is flipped and how far its picture is offset. + void choose_picture(const SensorInfo& info, Sensor& sensor) { + struct Look + { + const char* module; + const char* gif; + bool mirror; + bool stereo; + }; + static constexpr Look kLooks[] = { + { "S56C_HEAD", "noto-robot.gif", false, true }, + { "SHF3L_LEFT", "noto-waving-hand.gif", false, false }, + { "SHF3L_RIGHT", "noto-waving-hand.gif", true, false }, + { "SHW5G_MIX_L2", "noto-waving-hand.gif", false, false }, + { "SHW5G_MIX_L3", "noto-waving-hand.gif", true, false }, + }; + const Look* look = nullptr; + for (const Look& candidate : kLooks) + look = info.name == candidate.module ? &candidate : look; + const char* path = std::getenv(("SENSING_MOCK_GIF_" + info.name).c_str()); if (!path) path = std::getenv("SENSING_MOCK_GIF"); - const std::string bundled = std::string(SENSING_MOCK_ASSET_DIR) + "/noto-robot.gif"; + std::string bundled; if (!path) + { + if (!look) + { + throw std::runtime_error("no mock picture for module '" + info.name + "': set SENSING_MOCK_GIF_" + + info.name + " or SENSING_MOCK_GIF"); + } + bundled = std::string(SENSING_MOCK_ASSET_DIR) + "/" + look->gif; path = bundled.c_str(); + sensor.mirror = look->mirror; + if (look->stereo) + { + // The sensors of one module are its left and right eyes, in that order. The left eye sees + // the scene further right, so its picture moves right and the right eye's moves left. + size_t eye = 0; + for (const SensorInfo& other : sensor_infos) + eye += other.name == info.name && other.id < info.id ? 1 : 0; + const int half = static_cast(info.width * kEyeDisparity / 2); + sensor.shift_x = eye == 0 ? half : -half; + } + } // Sensors showing the same file share one copy on the device. auto& slot = gif_cache[path]; if (!slot) slot = std::make_shared(path); - return slot; + sensor.gif = slot; + + const std::string flip_name = "SENSING_MOCK_FLIP_" + info.name; + if (const char* flip = std::getenv(flip_name.c_str())) + { + const std::string value = flip; + if (value != "0" && value != "1") + throw std::runtime_error(flip_name + " must be 0 or 1, got '" + value + "'"); + sensor.mirror = value == "1"; + } } Sensor* find(uint32_t id) @@ -235,10 +292,11 @@ struct SiplCamera::Impl const int width = static_cast(sensor.info.width); const int height = static_cast(sensor.info.height); - launch_mock_gif_nv12(sensor.gif->frame(sensor.gif->frame_index_at(seconds)), - static_cast(sensor.gif->pitch()), static_cast(sensor.gif->width()), - static_cast(sensor.gif->height()), sensor.luma.ptr, static_cast(sensor.luma.pitch), - sensor.chroma.ptr, static_cast(sensor.chroma.pitch), width, height, sensor.stream); + launch_mock_gif_nv12( + sensor.gif->frame(sensor.gif->frame_index_at(seconds)), static_cast(sensor.gif->pitch()), + static_cast(sensor.gif->width()), static_cast(sensor.gif->height()), sensor.luma.ptr, + static_cast(sensor.luma.pitch), sensor.chroma.ptr, static_cast(sensor.chroma.pitch), width, + height, sensor.mirror, sensor.shift_x, sensor.stream); // Same call, layout and colour matrix as the SIPL path: full-range // BT.709 semi-planar, from the surface the ISP would have written. diff --git a/src/plugins/sensing/mock_assets/mock_rig.json b/src/plugins/sensing/mock_assets/mock_rig.json new file mode 100644 index 0000000000..70c1c84cee --- /dev/null +++ b/src/plugins/sensing/mock_assets/mock_rig.json @@ -0,0 +1,335 @@ +{ + "deserializers": [ + { + "name": "MAX96712_Fusa_nv_AB", + "description": "Maxim MAX96712 GMSL2 Deserializer - CSI-AB", + "i2cAddress": "0x29" + }, + { + "name": "MAX96712_Fusa_nv_CD", + "description": "Maxim MAX96712 GMSL2 Deserializer - CSI-CD", + "i2cAddress": "0x2d" + } + ], + "serializers": [ + { + "name": "MAX9295", + "description": "Maxim MAX9295D Serializer", + "i2cAddress": "0x40", + "useCDIv2API": true + }, + { + "name": "MAX96717F", + "description": "Maxim MAX96717F Serializer", + "i2cAddress": "0x40", + "useCDIv2API": true + }, + { + "name": "MAX96717", + "description": "Maxim MAX96717 Serializer", + "i2cAddress": "0x40", + "useCDIv2API": true + } + ], + "imageSensors": [ + { + "name": "VB1940", + "description": "SHW5G link2/3 VB1940 physical 0x10", + "i2cAddress": "0x10", + "virtualChannels": [ + { + "cfa": "gbrg", + "embeddedTopLines": 1, + "embeddedBottomLines": 0, + "inputFormat": "raw10", + "width": 2560, + "height": 1984, + "fps": 60.0, + "isEmbeddedDataTypeEnabled": true + } + ], + "isTriggerModeEnabled": true, + "isTPGEnabled": false, + "useCDIv2API": true + } + ], + "platformTransportSettings": [ + { + "name": "jetson-thor", + "description": "Thor - Jetson AGX Thor Devkit (p3834)", + "boardIdPrefix": "p3834", + "enableMasks": [ + "0x0000", + "0x1100" + ], + "transportSettings": [ + { + "name": "transportSettings_jetson-thor_AB", + "type": "GMSL", + "description": "GMSL CSI-AB with external MAX96712", + "csiPort": "csi-gh", + "phyMode": "dphy", + "desTxPort": 1, + "deserInfo": { + "name": "Max96712GmslDeserializer", + "i2cAddress": "0x2d" + }, + "powerControlInfo": { + "deserializerInfo": { + "name": "TegraDeserPowerDriver", + "gpioPinIndex": 9043968 + } + }, + "i2cDevice": 10, + "desI2CPort": 0, + "groupInitProg": true + }, + { + "name": "transportSettings_jetson-thor_CD", + "type": "GMSL", + "description": "GMSL CSI-CD with external MAX96712", + "csiPort": "csi-cd", + "phyMode": "dphy", + "desTxPort": 1, + "deserInfo": { + "name": "Max96712GmslDeserializer", + "i2cAddress": "0x29" + }, + "powerControlInfo": { + "deserializerInfo": { + "name": "TegraDeserPowerDriver", + "gpioPinIndex": 3211264 + } + }, + "i2cDevice": 9, + "desI2CPort": 0, + "groupInitProg": true + } + ] + } + ], + "cameraConfigs": [ + { + "name": "MOCK_S56C_2_SHF3L", + "type": "GMSL", + "description": "Mock rig: S56C stereo head camera on link 0, SHF3L left and right wrist cameras on links 1 and 2", + "cameraModules": [ + { + "name": "S56C_HEAD", + "linkIndex": 0, + "transportIndex": 1 + }, + { + "name": "SHF3L_LEFT", + "linkIndex": 1, + "transportIndex": 1 + }, + { + "name": "SHF3L_RIGHT", + "linkIndex": 2, + "transportIndex": 1 + } + ], + "transportSettings": [ + { + "name": "transportSettings_jetson-thor_AB", + "type": "GMSL", + "csiPort": "csi-gh", + "phyMode": "dphy", + "desTxPort": 1, + "deserInfo": { + "name": "Max96712GmslDeserializer", + "i2cAddress": "0x2d" + }, + "powerControlInfo": { + "deserializerInfo": { + "name": "TegraDeserPowerDriver", + "gpioPinIndex": 9043968 + } + }, + "i2cDevice": 10, + "desI2CPort": 0, + "groupInitProg": false + }, + { + "name": "transportSettings_jetson-thor_CD", + "type": "GMSL", + "csiPort": "csi-cd", + "phyMode": "dphy", + "desTxPort": 1, + "deserInfo": { + "name": "Max96712GmslDeserializer", + "i2cAddress": "0x29" + }, + "powerControlInfo": { + "deserializerInfo": { + "name": "TegraDeserPowerDriver", + "gpioPinIndex": 3211264 + } + }, + "i2cDevice": 9, + "desI2CPort": 0, + "groupInitProg": true + } + ] + }, + { + "name": "S56C_HEAD", + "type": "GMSL", + "description": "S56C stereo head camera, left eye then right eye", + "moduleDriverName": "SHW5G", + "serInfo": { + "name": "MAX96717", + "i2cAddress": "0x40" + }, + "linkMode": "LINK_MODE_GMSL2_6GBPS", + "fsyncMode": "osc_manual", + "mipiSettings": { + "dphyRate": 2500000, + "phyMode": "dphy" + }, + "sensorInfo": [ + { + "name": "S56C_LEFT", + "id": 0, + "description": "S56C stereo head camera, left eye then right eye", + "i2cAddress": "0x10", + "numContext": 1, + "isTriggerModeEnabled": true, + "deviceIndex": 0, + "virtualChannels": [ + { + "vcIdSrc": 2, + "vcIdDst": 2, + "cfa": "gbrg", + "embeddedTopLines": 1, + "embeddedBottomLines": 0, + "inputFormat": "raw10", + "width": 2560, + "height": 1984, + "fps": 30.0, + "isEmbeddedDataTypeEnabled": true + } + ] + }, + { + "name": "S56C_RIGHT", + "id": 1, + "description": "S56C stereo head camera, left eye then right eye", + "i2cAddress": "0x10", + "numContext": 1, + "isTriggerModeEnabled": true, + "deviceIndex": 0, + "virtualChannels": [ + { + "vcIdSrc": 2, + "vcIdDst": 2, + "cfa": "gbrg", + "embeddedTopLines": 1, + "embeddedBottomLines": 0, + "inputFormat": "raw10", + "width": 2560, + "height": 1984, + "fps": 30.0, + "isEmbeddedDataTypeEnabled": true + } + ] + } + ], + "eepromInfo": { + "name": "BR24H128", + "i2cAddress": "0x50" + } + }, + { + "name": "SHF3L_LEFT", + "type": "GMSL", + "description": "SHF3L left wrist camera", + "moduleDriverName": "SHW5G", + "serInfo": { + "name": "MAX96717", + "i2cAddress": "0x40" + }, + "linkMode": "LINK_MODE_GMSL2_6GBPS", + "fsyncMode": "osc_manual", + "mipiSettings": { + "dphyRate": 2500000, + "phyMode": "dphy" + }, + "sensorInfo": [ + { + "name": "SHF3L", + "id": 0, + "description": "SHF3L left wrist camera", + "i2cAddress": "0x10", + "numContext": 1, + "isTriggerModeEnabled": true, + "deviceIndex": 0, + "virtualChannels": [ + { + "vcIdSrc": 2, + "vcIdDst": 2, + "cfa": "gbrg", + "embeddedTopLines": 1, + "embeddedBottomLines": 0, + "inputFormat": "raw10", + "width": 1920, + "height": 1536, + "fps": 30.0, + "isEmbeddedDataTypeEnabled": true + } + ] + } + ], + "eepromInfo": { + "name": "BR24H128", + "i2cAddress": "0x50" + } + }, + { + "name": "SHF3L_RIGHT", + "type": "GMSL", + "description": "SHF3L right wrist camera", + "moduleDriverName": "SHW5G", + "serInfo": { + "name": "MAX96717", + "i2cAddress": "0x40" + }, + "linkMode": "LINK_MODE_GMSL2_6GBPS", + "fsyncMode": "osc_manual", + "mipiSettings": { + "dphyRate": 2500000, + "phyMode": "dphy" + }, + "sensorInfo": [ + { + "name": "SHF3L", + "id": 0, + "description": "SHF3L right wrist camera", + "i2cAddress": "0x10", + "numContext": 1, + "isTriggerModeEnabled": true, + "deviceIndex": 0, + "virtualChannels": [ + { + "vcIdSrc": 2, + "vcIdDst": 2, + "cfa": "gbrg", + "embeddedTopLines": 1, + "embeddedBottomLines": 0, + "inputFormat": "raw10", + "width": 1920, + "height": 1536, + "fps": 30.0, + "isEmbeddedDataTypeEnabled": true + } + ] + } + ], + "eepromInfo": { + "name": "BR24H128", + "i2cAddress": "0x50" + } + } + ] +} diff --git a/tests/cpp/plugins/sensing_mock/CMakeLists.txt b/tests/cpp/plugins/sensing_mock/CMakeLists.txt index 2e3f445bbb..96becb0f9b 100644 --- a/tests/cpp/plugins/sensing_mock/CMakeLists.txt +++ b/tests/cpp/plugins/sensing_mock/CMakeLists.txt @@ -26,6 +26,7 @@ target_link_libraries(sensing_mock_tests target_compile_definitions(sensing_mock_tests PRIVATE SENSING_PLATFORM_CONFIG="${CMAKE_SOURCE_DIR}/src/plugins/sensing/configs/shw5g.json" + SENSING_MOCK_RIG_CONFIG="${CMAKE_SOURCE_DIR}/src/plugins/sensing/mock_assets/mock_rig.json" SENSING_BUNDLED_GIF_DIR="${CMAKE_SOURCE_DIR}/src/plugins/sensing/mock_assets" ) diff --git a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp index 1d1043489e..7abec16d3d 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp @@ -12,8 +12,10 @@ #include #include #include +#include #include #include +#include using namespace plugins::sensing; using namespace sensing_test; @@ -328,59 +330,341 @@ TEST_CASE("the bundled GIFs decode, animate and have their transparency composit } } -TEST_CASE("by default every sensor shows the animated bundled robot", "[unit][gpu][sensing]") +namespace { - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv fps("SENSING_MOCK_FPS", "20"); - SiplCamera camera(make_sipl_config()); - camera.start(); - // Captures a second apart show different parts of the 2.5 s animation. - std::vector shots; - uint64_t last = 0; +// Every capture of every sensor, keyed by sequence, until each sensor has `count`. Capture n shows +// the same moment of the animation on every sensor and in every camera, so a sequence number names +// one picture and pictures can be compared across sensors and across cameras. +using Frames = std::map; + +std::vector record(SiplCamera& camera, size_t count) +{ + std::vector frames(camera.sensors().size()); + camera.start(); REQUIRE(wait_for( [&] { - if (auto view = camera.latest(0)) + bool done = true; + for (size_t s = 0; s < frames.size(); ++s) { - if (view->sequence >= last + 10 || shots.empty()) + if (auto view = camera.latest(static_cast(s))) { Shot shot; shot.width = view->width; shot.height = view->height; shot.sequence = view->sequence; shot.rgba = download_rgba(view->ptr, view->pitch, view->width, view->height); - shots.push_back(std::move(shot)); - last = view->sequence; + frames[s][view->sequence] = std::move(shot); } + done = done && frames[s].size() >= count; } - return shots.size() >= 4; + return done; }, - 10s)); + 15s)); camera.stop(); + return frames; +} + +// Mean absolute difference per colour channel between two pictures, optionally reading `b` flipped. +double mean_abs_diff(const Shot& a, const Shot& b, bool mirror_b = false) +{ + double sum = 0; + for (uint32_t y = 0; y < a.height; ++y) + { + for (uint32_t x = 0; x < a.width; ++x) + { + const Rgb pa = a.at(x, y); + const Rgb pb = b.at(mirror_b ? a.width - 1 - x : x, y); + sum += std::abs(pa.r - pb.r) + std::abs(pa.g - pb.g) + std::abs(pa.b - pb.b); + } + } + return sum / (static_cast(a.width) * a.height * 3); +} + +// The pictures two recordings have in common, as (sequence, first, second). +std::vector> shared(const Frames& a, const Frames& b) +{ + std::vector> out; + for (const auto& [sequence, shot] : a) + { + auto it = b.find(sequence); + if (it != b.end()) + out.emplace_back(sequence, &shot, &it->second); + } + return out; +} + +constexpr size_t kFrames = 12; + +} // namespace + +TEST_CASE("by default both sensors show the waving hand, one the mirror image of the other", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fps("SENSING_MOCK_FPS", "20"); + SiplCamera camera(make_sipl_config()); + REQUIRE(camera.sensors().size() == 2); + const auto both = record(camera, kFrames); + + // The two sensors show different pictures, and the second is the first flipped left to right. + const auto pairs = shared(both[0], both[1]); + REQUIRE(pairs.size() >= 8); + for (const auto& [sequence, first, second] : pairs) + { + INFO("capture " << sequence); + CHECK(mean_abs_diff(*second, *first, true) < 3.0); + CHECK(mean_abs_diff(*second, *first, false) > 6.0); + } - // Transparent corners of the emoji read as the background gray, and the picture moves. - for (const auto& shot : shots) + // The picture is the animated waving hand: transparent corners read as the background gray, and + // the picture moves with the clock. + for (const auto& [sequence, shot] : both[0]) CHECK(near(shot.at(2, 2), MockGifSource::kBackground, MockGifSource::kBackground, MockGifSource::kBackground, 8)); - size_t differing = 0; - for (size_t i = 0; i < shots[0].rgba.size(); i += 4) - differing += shots[0].rgba[i] != shots[2].rgba[i] ? 1 : 0; - CHECK(differing > 500); + CHECK(mean_abs_diff(both[0].begin()->second, both[0].rbegin()->second) > 2.0); + + // Same time, same picture as naming the bundled hand explicitly: the first sensor is the plain + // hand, and the explicit setting does not flip the second. + Frames explicit_first; + Frames explicit_second; + { + ScopedEnv hand("SENSING_MOCK_GIF", bundled("noto-waving-hand.gif").c_str()); + SiplCamera named(make_sipl_config()); + auto frames = record(named, kFrames); + explicit_first = std::move(frames[0]); + explicit_second = std::move(frames[1]); + } + const auto first_vs_named = shared(both[0], explicit_first); + REQUIRE(first_vs_named.size() >= 8); + for (const auto& [sequence, shot, named] : first_vs_named) + CHECK(mean_abs_diff(*shot, *named) < 3.0); + const auto second_vs_named = shared(both[1], explicit_second); + REQUIRE(second_vs_named.size() >= 8); + for (const auto& [sequence, shot, named] : second_vs_named) + CHECK(mean_abs_diff(*shot, *named, true) < 3.0); // flipped, as the default second sensor is + + // And it is not the robot. + Frames robot; + { + ScopedEnv face("SENSING_MOCK_GIF", bundled("noto-robot.gif").c_str()); + SiplCamera named(make_sipl_config()); + robot = std::move(record(named, kFrames)[0]); + } + const auto first_vs_robot = shared(both[0], robot); + REQUIRE(first_vs_robot.size() >= 8); + for (const auto& [sequence, shot, face] : first_vs_robot) + CHECK(mean_abs_diff(*shot, *face) > 15.0); } -TEST_CASE("the waving hand can be given to one sensor by module name", "[unit][gpu][sensing]") +TEST_CASE("a GIF named in the environment is shown as it is, on every sensor", "[unit][gpu][sensing]") { + TempFile gif(".gif"); + write_gif(gif.path(), 16, 16, { split_frame(16, 16, kGifRed, kGifBlue, 100) }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv one("SENSING_MOCK_GIF_SHW5G_MIX_L3", bundled("noto-waving-hand.gif").c_str()); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + SiplCamera camera(make_sipl_config()); camera.start(); - const Shot robot = grab(camera, 0); - const Shot hand = grab(camera, 1); + const Shot first = grab(camera, 0); + const Shot second = grab(camera, 1); camera.stop(); - // Different pictures: the robot on one sensor, the hand on the other. - size_t differing = 0; - for (size_t i = 0; i < robot.rgba.size(); i += 4) - differing += std::abs(robot.rgba[i] - hand.rgba[i]) > 20 ? 1 : 0; - CHECK(differing > 2000); + // Red on the left and blue on the right for both: the default flip belongs to the bundled hand. + for (const Shot* shot : { &first, &second }) + { + CHECK(near(shot->at(shot->width / 4, shot->height / 2), 255, 0, 0)); + CHECK(near(shot->at(shot->width * 3 / 4, shot->height / 2), 0, 0, 255)); + } +} + +TEST_CASE("a flip can be asked for on any sensor and any GIF, and taken off the default", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fps("SENSING_MOCK_FPS", "20"); + + SECTION("flipping a named GIF") + { + TempFile gif(".gif"); + write_gif(gif.path(), 16, 16, { split_frame(16, 16, kGifRed, kGifBlue, 100) }); + ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedEnv flip("SENSING_MOCK_FLIP_SHW5G_MIX_L3", "1"); + + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot first = grab(camera, 0); + const Shot second = grab(camera, 1); + camera.stop(); + CHECK(near(first.at(first.width / 4, first.height / 2), 255, 0, 0)); + CHECK(near(second.at(second.width / 4, second.height / 2), 0, 0, 255)); // flipped: blue on the left + } + + SECTION("taking the flip off the default hand leaves both sensors identical") + { + ScopedEnv flip("SENSING_MOCK_FLIP_SHW5G_MIX_L3", "0"); + SiplCamera camera(make_sipl_config()); + const auto both = record(camera, kFrames); + const auto pairs = shared(both[0], both[1]); + REQUIRE(pairs.size() >= 8); + for (const auto& [sequence, first, second] : pairs) + CHECK(mean_abs_diff(*first, *second) < 3.0); + } + + SECTION("choosing which sensor is flipped swaps the pictures") + { + ScopedEnv first_flipped("SENSING_MOCK_FLIP_SHW5G_MIX_L2", "1"); + ScopedEnv second_plain("SENSING_MOCK_FLIP_SHW5G_MIX_L3", "0"); + SiplCamera camera(make_sipl_config()); + const auto both = record(camera, kFrames); + const auto pairs = shared(both[0], both[1]); + REQUIRE(pairs.size() >= 8); + for (const auto& [sequence, first, second] : pairs) + CHECK(mean_abs_diff(*first, *second, true) < 3.0); // now the first is the mirror image + } + + SECTION("a flip must be 0 or 1") + { + ScopedEnv flip("SENSING_MOCK_FLIP_SHW5G_MIX_L2", "yes"); + CHECK_THROWS_WITH(SiplCamera(make_sipl_config()), Catch::Matchers::ContainsSubstring("must be 0 or 1")); + } +} + +TEST_CASE("the robot is for a head camera and is chosen by module name", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fps("SENSING_MOCK_FPS", "20"); + ScopedEnv head("SENSING_MOCK_GIF_SHW5G_MIX_L3", bundled("noto-robot.gif").c_str()); + + SiplCamera camera(make_sipl_config()); + const auto both = record(camera, kFrames); + + // The other sensor still shows the default hand, so the two pictures are plainly different. + const auto pairs = shared(both[0], both[1]); + REQUIRE(pairs.size() >= 8); + for (const auto& [sequence, hand, robot] : pairs) + CHECK(mean_abs_diff(*hand, *robot) > 15.0); +} + +namespace +{ + +// Head stereo pair on link 0, left and right wrist on links 1 and 2. Pipelines: 0 and 1 are the eyes, +// 2 the left wrist, 3 the right wrist. +SiplConfig make_rig_config(std::vector masks = { 0x0000, 0x0111 }) +{ + SiplConfig config; + config.platform_config_json = SENSING_MOCK_RIG_CONFIG; + config.platform_config_name = "MOCK_S56C_2_SHF3L"; + config.link_masks = std::move(masks); + return config; +} + +} // namespace + +TEST_CASE("the mock rig has the sizes and aspects of the real cameras", "[unit][sensing]") +{ + const auto sensors = SiplCamera::query(SENSING_MOCK_RIG_CONFIG, "MOCK_S56C_2_SHF3L", { 0x0000, 0x0111 }); + REQUIRE(sensors.size() == 4); + for (size_t eye = 0; eye < 2; ++eye) + { + CHECK(sensors[eye].name == "S56C_HEAD"); + CHECK(sensors[eye].width == 2560); + CHECK(sensors[eye].height == 1984); + } + for (size_t wrist = 2; wrist < 4; ++wrist) + { + CHECK(sensors[wrist].width == 1920); + CHECK(sensors[wrist].height == 1536); // 5:4, not the head's 2560:1984 + } + CHECK(sensors[2].name == "SHF3L_LEFT"); + CHECK(sensors[3].name == "SHF3L_RIGHT"); +} + +TEST_CASE("each rig camera delivers frames of its configured aspect, and wrists mirror while eyes offset", + "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fps("SENSING_MOCK_FPS", "20"); + SiplCamera camera(make_rig_config()); + REQUIRE(camera.sensors().size() == 4); + const auto frames = record(camera, kFrames); + + for (size_t s = 0; s < 4; ++s) + { + const auto& info = camera.sensors()[s]; + REQUIRE_FALSE(frames[s].empty()); + const Shot& shot = frames[s].begin()->second; + CHECK(shot.width == info.width); + CHECK(shot.height == info.height); + CHECK(shot.width == (s < 2 ? 320u : 240u)); + CHECK(shot.height == (s < 2 ? 248u : 192u)); + } + + // Left wrist plain, right wrist its mirror image. + const auto wrists = shared(frames[2], frames[3]); + REQUIRE(wrists.size() >= 8); + for (const auto& [sequence, left, right] : wrists) + { + CHECK(mean_abs_diff(*right, *left, true) < 3.0); + CHECK(mean_abs_diff(*right, *left, false) > 6.0); + } + + // Head eyes show the same robot offset by the disparity: left(x) is right(x - d), and the eyes differ + // where they are not offset. + const auto eyes = shared(frames[0], frames[1]); + REQUIRE(eyes.size() >= 8); + const uint32_t d = 2 * static_cast(320 * 0.02 / 2); + REQUIRE(d > 0); + for (const auto& [sequence, left, right] : eyes) + { + double sum = 0; + double unshifted = 0; + for (uint32_t y = 0; y < left->height; ++y) + { + for (uint32_t x = d + 4; x < left->width - 4; ++x) + { + const Rgb a = left->at(x, y); + const Rgb b = right->at(x - d, y); + const Rgb c = right->at(x, y); + sum += std::abs(a.r - b.r) + std::abs(a.g - b.g) + std::abs(a.b - b.b); + unshifted += std::abs(a.r - c.r) + std::abs(a.g - c.g) + std::abs(a.b - c.b); + } + } + const double n = static_cast(left->height) * (left->width - d - 8) * 3; + CHECK(sum / n < 3.0); + CHECK(unshifted / n > sum / n + 1.0); + } +} + +TEST_CASE("a rig camera keeps its picture when another camera is masked off", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + ScopedEnv fps("SENSING_MOCK_FPS", "20"); + + SiplCamera full(make_rig_config()); + const auto all = record(full, kFrames); + + // Only link 2, the right wrist, is enabled: it becomes pipeline 0 but is still the mirrored hand. + SiplCamera alone(make_rig_config({ 0x0000, 0x0100 })); + REQUIRE(alone.sensors().size() == 1); + REQUIRE(alone.sensors()[0].name == "SHF3L_RIGHT"); + const auto only = record(alone, kFrames); + + const auto pairs = shared(all[3], only[0]); + REQUIRE(pairs.size() >= 8); + for (const auto& [sequence, from_rig, from_alone] : pairs) + CHECK(mean_abs_diff(*from_rig, *from_alone) < 3.0); +} + +TEST_CASE("a module with no known picture needs one from the environment", "[unit][gpu][sensing]") +{ + ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); + TempFile config(".json"); + std::ofstream(config.path()) << R"({"cameraConfigs":[ + {"name":"P","cameraModules":[{"name":"UNKNOWN_CAM","linkIndex":0,"transportIndex":0}]}, + {"name":"UNKNOWN_CAM","sensorInfo":[{"virtualChannels":[{"width":64,"height":48,"fps":30.0}]}]}]})"; + SiplConfig sipl; + sipl.platform_config_json = config.path(); + sipl.platform_config_name = "P"; + sipl.link_masks = { 0x1 }; + CHECK_THROWS_WITH(SiplCamera(sipl), Catch::Matchers::ContainsSubstring("SENSING_MOCK_GIF_UNKNOWN_CAM")); } From 2fb64b45e13ee268d8f77f14bca1ef1bc7fa2d83 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sun, 25 Oct 2026 00:02:40 -0400 Subject: [PATCH 06/10] Drop the mock's GIF and flip environment variables Pictures are fixed by module name, so the per-module and global GIF overrides and the flip knob only added precedence rules and tests. Tests that need a known picture use an internal hook instead. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/README.md | 7 +- src/plugins/sensing/core/mock_sipl_camera.cpp | 52 +++----- src/plugins/sensing/core/mock_test_hooks.hpp | 20 +++ .../sensing_mock/mock_test_support.hpp | 24 +++- .../plugins/sensing_mock/test_mock_gif.cpp | 118 +++--------------- 5 files changed, 77 insertions(+), 144 deletions(-) create mode 100644 src/plugins/sensing/core/mock_test_hooks.hpp diff --git a/src/plugins/sensing/README.md b/src/plugins/sensing/README.md index fc5f86d482..e36603a632 100644 --- a/src/plugins/sensing/README.md +++ b/src/plugins/sensing/README.md @@ -78,8 +78,8 @@ Every sensor shows an animated GIF chosen by its **module name** (never its pipe link mask renumbers), from the bundled Animated Noto Emoji in [`mock_assets/`](mock_assets/) (CC BY 4.0; see its README): `S56C_HEAD` the robot, `SHF3L_LEFT` the waving hand, `SHF3L_RIGHT` the hand flipped left to right. The two sensors of `S56C_HEAD` are its left and right eyes; their pictures are offset against -each other by a fixed 2% of the width, as a 60 mm baseline gives. A module not in that table needs -`SENSING_MOCK_GIF` or `SENSING_MOCK_GIF_`. The GIF plays on the capture clock, looping, scaled +each other by a fixed 2% of the width (illustrative, not calibrated geometry). A module not in that +table is an error. The GIF plays on the capture clock, looping, scaled to cover the sensor's configured size and centre-cropped so it is never stretched (bilinear when enlarging, area-averaged when reducing). Transparent pixels show a plain gray. The GIF's own size and aspect do not matter; the frame size always comes from the platform config. A GIF that cannot be loaded, @@ -94,9 +94,6 @@ Environment, read once at construction: | Variable | Effect | | --- | --- | -| `SENSING_MOCK_GIF` | path of an animated GIF shown by every sensor, as it is (default: by module, above) | -| `SENSING_MOCK_GIF_` | a GIF for one sensor, by module name (`SENSING_MOCK_GIF_SHW5G_MIX_L3`); wins over `SENSING_MOCK_GIF` | -| `SENSING_MOCK_FLIP_` | `1` or `0`: flip that module's picture left to right, or not, whatever its GIF | | `SENSING_MOCK_FPS` | override the platform config's frame rate | | `SENSING_MOCK_SCALE_DIV` | divide the reported width and height (kept even); useful for tests | | `SENSING_MOCK_FAIL_AFTER` | every sensor fails after N frames, as a link drop would | diff --git a/src/plugins/sensing/core/mock_sipl_camera.cpp b/src/plugins/sensing/core/mock_sipl_camera.cpp index 160cc4427d..491b2c1911 100644 --- a/src/plugins/sensing/core/mock_sipl_camera.cpp +++ b/src/plugins/sensing/core/mock_sipl_camera.cpp @@ -8,19 +8,16 @@ // YUV-to-RGBA kernel on a synthetic NV12 surface made from the GIF, so the // conversion code is exercised too. // -// What a sensor shows is chosen by module name, never by pipeline index (which the link mask renumbers): -// S56C_HEAD the robot, the two eyes offset against each other by a fixed disparity +// What a sensor shows is fixed by its module name, never by pipeline index (which the link mask +// renumbers): +// S56C_HEAD the robot; its two sensors are the left and right eye, offset against each other // SHF3L_LEFT the waving hand // SHF3L_RIGHT the waving hand, flipped left to right // SHW5G_MIX_L2 / SHW5G_MIX_L3 as SHF3L_LEFT / SHF3L_RIGHT -// A module not listed needs SENSING_MOCK_GIF or SENSING_MOCK_GIF_. Environment, in order: -// SENSING_MOCK_GIF_=PATH one module's GIF -// SENSING_MOCK_GIF=PATH every sensor's GIF -// SENSING_MOCK_FLIP_=1|0 flip one module, whatever its GIF -// A GIF named by the environment is shown as it is, with no eye offset. A GIF that cannot be loaded, -// the bundled one included, is an error. It is scaled to cover the sensor's configured size and -// centre-cropped, never stretched, and follows the capture clock: it advances with each frame's time, -// and the sequence number and stamps advance on every capture even while the GIF shows one image. +// A module not listed is an error, and so is a GIF that cannot be loaded. The picture is scaled to cover +// the sensor's configured size and centre-cropped, never stretched, and follows the capture clock: it +// advances with each frame's time, and the sequence number and stamps advance on every capture even +// while the GIF shows one image. // // Other environment (read once, at construction): // SENSING_MOCK_FPS=N override the platform config's frame rate @@ -30,6 +27,7 @@ #include "mock_gif_source.hpp" #include "mock_nv12_source.cuh" #include "mock_platform_config.hpp" +#include "mock_test_hooks.hpp" #include "sipl_camera.hpp" #include "yuv_to_rgba.cuh" @@ -58,8 +56,7 @@ namespace // Arbitrary offset so nothing can mistake the TSC for CLOCK_MONOTONIC. constexpr uint64_t kMockTscBaseNs = 1'000'000'000'000'000ull; -// Total horizontal offset between the two eyes, as a fraction of the width: a 60 mm baseline seen at -// about half a metre through a 140 degree lens. +// Total horizontal offset between the two eyes, as a fraction of the width. Illustrative, not calibrated. constexpr double kEyeDisparity = 0.02; uint64_t monotonic_ns() @@ -80,8 +77,15 @@ void check_cuda(cudaError_t err, const char* what) throw std::runtime_error(std::string("mock SIPL: ") + what + ": " + cudaGetErrorString(err)); } +std::string g_test_picture; + } // namespace +void mock_set_test_picture(const std::string& gif_path) +{ + g_test_picture = gif_path; +} + struct SiplCamera::Impl { struct DeviceBuffer @@ -197,19 +201,12 @@ struct SiplCamera::Impl for (const Look& candidate : kLooks) look = info.name == candidate.module ? &candidate : look; - const char* path = std::getenv(("SENSING_MOCK_GIF_" + info.name).c_str()); - if (!path) - path = std::getenv("SENSING_MOCK_GIF"); - std::string bundled; - if (!path) + std::string path = g_test_picture; + if (path.empty()) { if (!look) - { - throw std::runtime_error("no mock picture for module '" + info.name + "': set SENSING_MOCK_GIF_" + - info.name + " or SENSING_MOCK_GIF"); - } - bundled = std::string(SENSING_MOCK_ASSET_DIR) + "/" + look->gif; - path = bundled.c_str(); + throw std::runtime_error("no mock picture for module '" + info.name + "'"); + path = std::string(SENSING_MOCK_ASSET_DIR) + "/" + look->gif; sensor.mirror = look->mirror; if (look->stereo) { @@ -228,15 +225,6 @@ struct SiplCamera::Impl if (!slot) slot = std::make_shared(path); sensor.gif = slot; - - const std::string flip_name = "SENSING_MOCK_FLIP_" + info.name; - if (const char* flip = std::getenv(flip_name.c_str())) - { - const std::string value = flip; - if (value != "0" && value != "1") - throw std::runtime_error(flip_name + " must be 0 or 1, got '" + value + "'"); - sensor.mirror = value == "1"; - } } Sensor* find(uint32_t id) diff --git a/src/plugins/sensing/core/mock_test_hooks.hpp b/src/plugins/sensing/core/mock_test_hooks.hpp new file mode 100644 index 0000000000..19d96e6591 --- /dev/null +++ b/src/plugins/sensing/core/mock_test_hooks.hpp @@ -0,0 +1,20 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// Mock backend only, for tests: not an operator interface. + +#pragma once + +#include + +namespace plugins +{ +namespace sensing +{ + +/// Every sensor constructed afterwards shows this GIF as it is -- no mirror, no eye offset -- instead of +/// its module's fixed picture. An empty path restores the fixed pictures. +void mock_set_test_picture(const std::string& gif_path); + +} // namespace sensing +} // namespace plugins diff --git a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp index 5137d64cae..a7106ad011 100644 --- a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -6,6 +6,7 @@ #include "gif_fixture.hpp" #include +#include #include #include @@ -14,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -50,6 +52,22 @@ class ScopedEnv std::string m_name; }; +/// Every sensor shows this GIF as it is for one scope. +class ScopedPicture +{ +public: + explicit ScopedPicture(const std::string& path) + { + plugins::sensing::mock_set_test_picture(path); + } + ~ScopedPicture() + { + plugins::sensing::mock_set_test_picture(""); + } + ScopedPicture(const ScopedPicture&) = delete; + ScopedPicture& operator=(const ScopedPicture&) = delete; +}; + template bool wait_for(Pred done, std::chrono::milliseconds limit = std::chrono::seconds(5)) { @@ -106,18 +124,18 @@ class TempFile class FastMock { public: - FastMock() : m_gif(".gif"), m_scale("SENSING_MOCK_SCALE_DIV", "8"), m_source("SENSING_MOCK_GIF") + FastMock() : m_gif(".gif"), m_scale("SENSING_MOCK_SCALE_DIV", "8") { write_gif(m_gif.path(), 16, 16, { solid_frame(16, 16, kGifBlack, 4), solid_frame(16, 16, kGifRed, 4), solid_frame(16, 16, kGifBlue, 4), solid_frame(16, 16, kGifWhite, 4) }); - m_source.set(m_gif.path().c_str()); + m_source = std::make_unique(m_gif.path()); } private: TempFile m_gif; ScopedEnv m_scale; - ScopedEnv m_source; + std::unique_ptr m_source; }; struct Rgb diff --git a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp index 7abec16d3d..946d0e58b1 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp @@ -80,7 +80,7 @@ TEST_CASE("a GIF that is missing or is not a GIF fails construction", "[unit][gp CHECK_THROWS_WITH(MockGifSource(text.path()), Catch::Matchers::ContainsSubstring("cannot decode")); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv gif("SENSING_MOCK_GIF", "/nonexistent/none.gif"); + ScopedPicture gif("/nonexistent/none.gif"); CHECK_THROWS_AS(SiplCamera(make_sipl_config()), std::runtime_error); } @@ -91,7 +91,7 @@ TEST_CASE("the camera shows the GIF frame for its capture time while sequence an // Delays are in centiseconds: 4 is 40 ms. write_gif(gif.path(), 16, 16, { solid_frame(16, 16, kGifRed, 4), solid_frame(16, 16, kGifBlue, 4) }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); ScopedEnv fps("SENSING_MOCK_FPS", "50"); // 20 ms per capture: two captures per GIF frame SiplCamera camera(make_sipl_config()); @@ -149,7 +149,7 @@ TEST_CASE("an enlarged GIF keeps its proportions on the sensor's geometry", "[un TempFile gif(".gif"); write_gif(gif.path(), 32, 32, { square_frame(32, 32, 16, 100) }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); SiplCamera camera(make_sipl_config()); camera.start(); @@ -178,7 +178,7 @@ TEST_CASE("a wide GIF is reduced and cropped to the sensor's aspect", "[unit][gp TempFile gif(".gif"); write_gif(gif.path(), 640, 320, { split_frame(640, 320, kGifRed, kGifBlue, 100) }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); SiplCamera camera(make_sipl_config()); camera.start(); @@ -203,7 +203,7 @@ TEST_CASE("a tall GIF is cropped top and bottom", "[unit][gpu][sensing]") std::fill_n(frame.pixels.begin() + y * 100, 100, kGifWhite); write_gif(gif.path(), 100, 400, { frame }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); SiplCamera camera(make_sipl_config()); camera.start(); @@ -222,7 +222,7 @@ TEST_CASE("a large reduction averages the source instead of skipping pixels", "[ TempFile gif(".gif"); write_gif(gif.path(), 200, 200, { split_frame(200, 200, kGifRed, kGifBlue, 100) }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "32"); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); SiplCamera camera(make_sipl_config()); camera.start(); @@ -239,7 +239,7 @@ TEST_CASE("without a scale override the GIF is shown at the sensor's configured { TempFile gif(".gif"); write_gif(gif.path(), 32, 32, { square_frame(32, 32, 16, 100) }); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); ScopedEnv fps("SENSING_MOCK_FPS", "5"); SiplCamera camera(make_sipl_config()); @@ -262,29 +262,6 @@ TEST_CASE("without a scale override the GIF is shown at the sensor's configured CHECK(std::abs(run_x - run_y) <= 4); } -TEST_CASE("a sensor's GIF can be chosen by its module name", "[unit][gpu][sensing]") -{ - TempFile red(".gif"); - TempFile blue(".gif"); - write_gif(red.path(), 16, 16, { solid_frame(16, 16, kGifRed, 100) }); - write_gif(blue.path(), 16, 16, { solid_frame(16, 16, kGifBlue, 100) }); - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv everyone("SENSING_MOCK_GIF", red.path().c_str()); - ScopedEnv one("SENSING_MOCK_GIF_SHW5G_MIX_L3", blue.path().c_str()); - - SiplCamera camera(make_sipl_config()); - REQUIRE(camera.sensors().size() == 2); - REQUIRE(camera.sensors()[0].name == "SHW5G_MIX_L2"); - REQUIRE(camera.sensors()[1].name == "SHW5G_MIX_L3"); - camera.start(); - const Shot first = grab(camera, 0); - const Shot second = grab(camera, 1); - camera.stop(); - - CHECK(near(first.at(first.width / 2, first.height / 2), 255, 0, 0)); - CHECK(near(second.at(second.width / 2, second.height / 2), 0, 0, 255)); -} - namespace { @@ -428,7 +405,7 @@ TEST_CASE("by default both sensors show the waving hand, one the mirror image of Frames explicit_first; Frames explicit_second; { - ScopedEnv hand("SENSING_MOCK_GIF", bundled("noto-waving-hand.gif").c_str()); + ScopedPicture hand(bundled("noto-waving-hand.gif")); SiplCamera named(make_sipl_config()); auto frames = record(named, kFrames); explicit_first = std::move(frames[0]); @@ -446,7 +423,7 @@ TEST_CASE("by default both sensors show the waving hand, one the mirror image of // And it is not the robot. Frames robot; { - ScopedEnv face("SENSING_MOCK_GIF", bundled("noto-robot.gif").c_str()); + ScopedPicture face(bundled("noto-robot.gif")); SiplCamera named(make_sipl_config()); robot = std::move(record(named, kFrames)[0]); } @@ -456,12 +433,12 @@ TEST_CASE("by default both sensors show the waving hand, one the mirror image of CHECK(mean_abs_diff(*shot, *face) > 15.0); } -TEST_CASE("a GIF named in the environment is shown as it is, on every sensor", "[unit][gpu][sensing]") +TEST_CASE("a test picture is shown as it is, on every sensor", "[unit][gpu][sensing]") { TempFile gif(".gif"); write_gif(gif.path(), 16, 16, { split_frame(16, 16, kGifRed, kGifBlue, 100) }); ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); + ScopedPicture source(gif.path()); SiplCamera camera(make_sipl_config()); camera.start(); @@ -469,7 +446,7 @@ TEST_CASE("a GIF named in the environment is shown as it is, on every sensor", " const Shot second = grab(camera, 1); camera.stop(); - // Red on the left and blue on the right for both: the default flip belongs to the bundled hand. + // Red on the left and blue on the right for both: the mirror belongs to the fixed pictures. for (const Shot* shot : { &first, &second }) { CHECK(near(shot->at(shot->width / 4, shot->height / 2), 255, 0, 0)); @@ -477,73 +454,6 @@ TEST_CASE("a GIF named in the environment is shown as it is, on every sensor", " } } -TEST_CASE("a flip can be asked for on any sensor and any GIF, and taken off the default", "[unit][gpu][sensing]") -{ - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv fps("SENSING_MOCK_FPS", "20"); - - SECTION("flipping a named GIF") - { - TempFile gif(".gif"); - write_gif(gif.path(), 16, 16, { split_frame(16, 16, kGifRed, kGifBlue, 100) }); - ScopedEnv source("SENSING_MOCK_GIF", gif.path().c_str()); - ScopedEnv flip("SENSING_MOCK_FLIP_SHW5G_MIX_L3", "1"); - - SiplCamera camera(make_sipl_config()); - camera.start(); - const Shot first = grab(camera, 0); - const Shot second = grab(camera, 1); - camera.stop(); - CHECK(near(first.at(first.width / 4, first.height / 2), 255, 0, 0)); - CHECK(near(second.at(second.width / 4, second.height / 2), 0, 0, 255)); // flipped: blue on the left - } - - SECTION("taking the flip off the default hand leaves both sensors identical") - { - ScopedEnv flip("SENSING_MOCK_FLIP_SHW5G_MIX_L3", "0"); - SiplCamera camera(make_sipl_config()); - const auto both = record(camera, kFrames); - const auto pairs = shared(both[0], both[1]); - REQUIRE(pairs.size() >= 8); - for (const auto& [sequence, first, second] : pairs) - CHECK(mean_abs_diff(*first, *second) < 3.0); - } - - SECTION("choosing which sensor is flipped swaps the pictures") - { - ScopedEnv first_flipped("SENSING_MOCK_FLIP_SHW5G_MIX_L2", "1"); - ScopedEnv second_plain("SENSING_MOCK_FLIP_SHW5G_MIX_L3", "0"); - SiplCamera camera(make_sipl_config()); - const auto both = record(camera, kFrames); - const auto pairs = shared(both[0], both[1]); - REQUIRE(pairs.size() >= 8); - for (const auto& [sequence, first, second] : pairs) - CHECK(mean_abs_diff(*first, *second, true) < 3.0); // now the first is the mirror image - } - - SECTION("a flip must be 0 or 1") - { - ScopedEnv flip("SENSING_MOCK_FLIP_SHW5G_MIX_L2", "yes"); - CHECK_THROWS_WITH(SiplCamera(make_sipl_config()), Catch::Matchers::ContainsSubstring("must be 0 or 1")); - } -} - -TEST_CASE("the robot is for a head camera and is chosen by module name", "[unit][gpu][sensing]") -{ - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv fps("SENSING_MOCK_FPS", "20"); - ScopedEnv head("SENSING_MOCK_GIF_SHW5G_MIX_L3", bundled("noto-robot.gif").c_str()); - - SiplCamera camera(make_sipl_config()); - const auto both = record(camera, kFrames); - - // The other sensor still shows the default hand, so the two pictures are plainly different. - const auto pairs = shared(both[0], both[1]); - REQUIRE(pairs.size() >= 8); - for (const auto& [sequence, hand, robot] : pairs) - CHECK(mean_abs_diff(*hand, *robot) > 15.0); -} - namespace { @@ -655,7 +565,7 @@ TEST_CASE("a rig camera keeps its picture when another camera is masked off", "[ CHECK(mean_abs_diff(*from_rig, *from_alone) < 3.0); } -TEST_CASE("a module with no known picture needs one from the environment", "[unit][gpu][sensing]") +TEST_CASE("a module with no known picture is an error", "[unit][gpu][sensing]") { ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); TempFile config(".json"); @@ -666,5 +576,5 @@ TEST_CASE("a module with no known picture needs one from the environment", "[uni sipl.platform_config_json = config.path(); sipl.platform_config_name = "P"; sipl.link_masks = { 0x1 }; - CHECK_THROWS_WITH(SiplCamera(sipl), Catch::Matchers::ContainsSubstring("SENSING_MOCK_GIF_UNKNOWN_CAM")); + CHECK_THROWS_WITH(SiplCamera(sipl), Catch::Matchers::ContainsSubstring("no mock picture for module 'UNKNOWN_CAM'")); } From b380899237b947b7bc31d4184822c4af95547149 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sun, 25 Oct 2026 00:03:29 -0400 Subject: [PATCH 07/10] Trim the mock rig config to what the parser reads Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/mock_assets/mock_rig.json | 274 +----------------- 1 file changed, 7 insertions(+), 267 deletions(-) diff --git a/src/plugins/sensing/mock_assets/mock_rig.json b/src/plugins/sensing/mock_assets/mock_rig.json index 70c1c84cee..ce0753682b 100644 --- a/src/plugins/sensing/mock_assets/mock_rig.json +++ b/src/plugins/sensing/mock_assets/mock_rig.json @@ -1,118 +1,7 @@ { - "deserializers": [ - { - "name": "MAX96712_Fusa_nv_AB", - "description": "Maxim MAX96712 GMSL2 Deserializer - CSI-AB", - "i2cAddress": "0x29" - }, - { - "name": "MAX96712_Fusa_nv_CD", - "description": "Maxim MAX96712 GMSL2 Deserializer - CSI-CD", - "i2cAddress": "0x2d" - } - ], - "serializers": [ - { - "name": "MAX9295", - "description": "Maxim MAX9295D Serializer", - "i2cAddress": "0x40", - "useCDIv2API": true - }, - { - "name": "MAX96717F", - "description": "Maxim MAX96717F Serializer", - "i2cAddress": "0x40", - "useCDIv2API": true - }, - { - "name": "MAX96717", - "description": "Maxim MAX96717 Serializer", - "i2cAddress": "0x40", - "useCDIv2API": true - } - ], - "imageSensors": [ - { - "name": "VB1940", - "description": "SHW5G link2/3 VB1940 physical 0x10", - "i2cAddress": "0x10", - "virtualChannels": [ - { - "cfa": "gbrg", - "embeddedTopLines": 1, - "embeddedBottomLines": 0, - "inputFormat": "raw10", - "width": 2560, - "height": 1984, - "fps": 60.0, - "isEmbeddedDataTypeEnabled": true - } - ], - "isTriggerModeEnabled": true, - "isTPGEnabled": false, - "useCDIv2API": true - } - ], - "platformTransportSettings": [ - { - "name": "jetson-thor", - "description": "Thor - Jetson AGX Thor Devkit (p3834)", - "boardIdPrefix": "p3834", - "enableMasks": [ - "0x0000", - "0x1100" - ], - "transportSettings": [ - { - "name": "transportSettings_jetson-thor_AB", - "type": "GMSL", - "description": "GMSL CSI-AB with external MAX96712", - "csiPort": "csi-gh", - "phyMode": "dphy", - "desTxPort": 1, - "deserInfo": { - "name": "Max96712GmslDeserializer", - "i2cAddress": "0x2d" - }, - "powerControlInfo": { - "deserializerInfo": { - "name": "TegraDeserPowerDriver", - "gpioPinIndex": 9043968 - } - }, - "i2cDevice": 10, - "desI2CPort": 0, - "groupInitProg": true - }, - { - "name": "transportSettings_jetson-thor_CD", - "type": "GMSL", - "description": "GMSL CSI-CD with external MAX96712", - "csiPort": "csi-cd", - "phyMode": "dphy", - "desTxPort": 1, - "deserInfo": { - "name": "Max96712GmslDeserializer", - "i2cAddress": "0x29" - }, - "powerControlInfo": { - "deserializerInfo": { - "name": "TegraDeserPowerDriver", - "gpioPinIndex": 3211264 - } - }, - "i2cDevice": 9, - "desI2CPort": 0, - "groupInitProg": true - } - ] - } - ], "cameraConfigs": [ { "name": "MOCK_S56C_2_SHF3L", - "type": "GMSL", - "description": "Mock rig: S56C stereo head camera on link 0, SHF3L left and right wrist cameras on links 1 and 2", "cameraModules": [ { "name": "S56C_HEAD", @@ -129,207 +18,58 @@ "linkIndex": 2, "transportIndex": 1 } - ], - "transportSettings": [ - { - "name": "transportSettings_jetson-thor_AB", - "type": "GMSL", - "csiPort": "csi-gh", - "phyMode": "dphy", - "desTxPort": 1, - "deserInfo": { - "name": "Max96712GmslDeserializer", - "i2cAddress": "0x2d" - }, - "powerControlInfo": { - "deserializerInfo": { - "name": "TegraDeserPowerDriver", - "gpioPinIndex": 9043968 - } - }, - "i2cDevice": 10, - "desI2CPort": 0, - "groupInitProg": false - }, - { - "name": "transportSettings_jetson-thor_CD", - "type": "GMSL", - "csiPort": "csi-cd", - "phyMode": "dphy", - "desTxPort": 1, - "deserInfo": { - "name": "Max96712GmslDeserializer", - "i2cAddress": "0x29" - }, - "powerControlInfo": { - "deserializerInfo": { - "name": "TegraDeserPowerDriver", - "gpioPinIndex": 3211264 - } - }, - "i2cDevice": 9, - "desI2CPort": 0, - "groupInitProg": true - } ] }, { "name": "S56C_HEAD", - "type": "GMSL", - "description": "S56C stereo head camera, left eye then right eye", - "moduleDriverName": "SHW5G", - "serInfo": { - "name": "MAX96717", - "i2cAddress": "0x40" - }, - "linkMode": "LINK_MODE_GMSL2_6GBPS", - "fsyncMode": "osc_manual", - "mipiSettings": { - "dphyRate": 2500000, - "phyMode": "dphy" - }, "sensorInfo": [ { - "name": "S56C_LEFT", - "id": 0, - "description": "S56C stereo head camera, left eye then right eye", - "i2cAddress": "0x10", - "numContext": 1, - "isTriggerModeEnabled": true, - "deviceIndex": 0, "virtualChannels": [ { - "vcIdSrc": 2, - "vcIdDst": 2, - "cfa": "gbrg", - "embeddedTopLines": 1, - "embeddedBottomLines": 0, - "inputFormat": "raw10", "width": 2560, "height": 1984, - "fps": 30.0, - "isEmbeddedDataTypeEnabled": true + "fps": 30.0 } ] }, { - "name": "S56C_RIGHT", - "id": 1, - "description": "S56C stereo head camera, left eye then right eye", - "i2cAddress": "0x10", - "numContext": 1, - "isTriggerModeEnabled": true, - "deviceIndex": 0, "virtualChannels": [ { - "vcIdSrc": 2, - "vcIdDst": 2, - "cfa": "gbrg", - "embeddedTopLines": 1, - "embeddedBottomLines": 0, - "inputFormat": "raw10", "width": 2560, "height": 1984, - "fps": 30.0, - "isEmbeddedDataTypeEnabled": true + "fps": 30.0 } ] } - ], - "eepromInfo": { - "name": "BR24H128", - "i2cAddress": "0x50" - } + ] }, { "name": "SHF3L_LEFT", - "type": "GMSL", - "description": "SHF3L left wrist camera", - "moduleDriverName": "SHW5G", - "serInfo": { - "name": "MAX96717", - "i2cAddress": "0x40" - }, - "linkMode": "LINK_MODE_GMSL2_6GBPS", - "fsyncMode": "osc_manual", - "mipiSettings": { - "dphyRate": 2500000, - "phyMode": "dphy" - }, "sensorInfo": [ { - "name": "SHF3L", - "id": 0, - "description": "SHF3L left wrist camera", - "i2cAddress": "0x10", - "numContext": 1, - "isTriggerModeEnabled": true, - "deviceIndex": 0, "virtualChannels": [ { - "vcIdSrc": 2, - "vcIdDst": 2, - "cfa": "gbrg", - "embeddedTopLines": 1, - "embeddedBottomLines": 0, - "inputFormat": "raw10", "width": 1920, "height": 1536, - "fps": 30.0, - "isEmbeddedDataTypeEnabled": true + "fps": 30.0 } ] } - ], - "eepromInfo": { - "name": "BR24H128", - "i2cAddress": "0x50" - } + ] }, { "name": "SHF3L_RIGHT", - "type": "GMSL", - "description": "SHF3L right wrist camera", - "moduleDriverName": "SHW5G", - "serInfo": { - "name": "MAX96717", - "i2cAddress": "0x40" - }, - "linkMode": "LINK_MODE_GMSL2_6GBPS", - "fsyncMode": "osc_manual", - "mipiSettings": { - "dphyRate": 2500000, - "phyMode": "dphy" - }, "sensorInfo": [ { - "name": "SHF3L", - "id": 0, - "description": "SHF3L right wrist camera", - "i2cAddress": "0x10", - "numContext": 1, - "isTriggerModeEnabled": true, - "deviceIndex": 0, "virtualChannels": [ { - "vcIdSrc": 2, - "vcIdDst": 2, - "cfa": "gbrg", - "embeddedTopLines": 1, - "embeddedBottomLines": 0, - "inputFormat": "raw10", "width": 1920, "height": 1536, - "fps": 30.0, - "isEmbeddedDataTypeEnabled": true + "fps": 30.0 } ] } - ], - "eepromInfo": { - "name": "BR24H128", - "i2cAddress": "0x50" - } + ] } ] } From 0158e834b086999aaf99e534c7000ca3fdcc9ae1 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sun, 25 Oct 2026 00:04:23 -0400 Subject: [PATCH 08/10] Drop the redundant default-picture comparison test The rig tests cover the fixed pictures, the wrist mirror and masking. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- .../plugins/sensing_mock/test_mock_gif.cpp | 57 ------------------- 1 file changed, 57 deletions(-) diff --git a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp index 946d0e58b1..1413427570 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp @@ -376,63 +376,6 @@ constexpr size_t kFrames = 12; } // namespace -TEST_CASE("by default both sensors show the waving hand, one the mirror image of the other", "[unit][gpu][sensing]") -{ - ScopedEnv scale("SENSING_MOCK_SCALE_DIV", "8"); - ScopedEnv fps("SENSING_MOCK_FPS", "20"); - SiplCamera camera(make_sipl_config()); - REQUIRE(camera.sensors().size() == 2); - const auto both = record(camera, kFrames); - - // The two sensors show different pictures, and the second is the first flipped left to right. - const auto pairs = shared(both[0], both[1]); - REQUIRE(pairs.size() >= 8); - for (const auto& [sequence, first, second] : pairs) - { - INFO("capture " << sequence); - CHECK(mean_abs_diff(*second, *first, true) < 3.0); - CHECK(mean_abs_diff(*second, *first, false) > 6.0); - } - - // The picture is the animated waving hand: transparent corners read as the background gray, and - // the picture moves with the clock. - for (const auto& [sequence, shot] : both[0]) - CHECK(near(shot.at(2, 2), MockGifSource::kBackground, MockGifSource::kBackground, MockGifSource::kBackground, 8)); - CHECK(mean_abs_diff(both[0].begin()->second, both[0].rbegin()->second) > 2.0); - - // Same time, same picture as naming the bundled hand explicitly: the first sensor is the plain - // hand, and the explicit setting does not flip the second. - Frames explicit_first; - Frames explicit_second; - { - ScopedPicture hand(bundled("noto-waving-hand.gif")); - SiplCamera named(make_sipl_config()); - auto frames = record(named, kFrames); - explicit_first = std::move(frames[0]); - explicit_second = std::move(frames[1]); - } - const auto first_vs_named = shared(both[0], explicit_first); - REQUIRE(first_vs_named.size() >= 8); - for (const auto& [sequence, shot, named] : first_vs_named) - CHECK(mean_abs_diff(*shot, *named) < 3.0); - const auto second_vs_named = shared(both[1], explicit_second); - REQUIRE(second_vs_named.size() >= 8); - for (const auto& [sequence, shot, named] : second_vs_named) - CHECK(mean_abs_diff(*shot, *named, true) < 3.0); // flipped, as the default second sensor is - - // And it is not the robot. - Frames robot; - { - ScopedPicture face(bundled("noto-robot.gif")); - SiplCamera named(make_sipl_config()); - robot = std::move(record(named, kFrames)[0]); - } - const auto first_vs_robot = shared(both[0], robot); - REQUIRE(first_vs_robot.size() >= 8); - for (const auto& [sequence, shot, face] : first_vs_robot) - CHECK(mean_abs_diff(*shot, *face) > 15.0); -} - TEST_CASE("a test picture is shown as it is, on every sensor", "[unit][gpu][sensing]") { TempFile gif(".gif"); From d51ff65e47113f601f765200094224690ba63f4f Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sun, 25 Oct 2026 00:32:57 -0400 Subject: [PATCH 09/10] Make the mock encoder produce real H.264 Feed the NV12 buffer to libavcodec's h264_nvenc, set up as the Jetson V4L2 encoder is: High profile, Level 5.2, VBR with a peak, an IDR every GOP with SPS/PPS in front of it, no B-frames, microsecond stamps. The tests decode the output. Drops the fake payload generator and its latency and depth knobs. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- src/plugins/sensing/CMakeLists.txt | 2 +- src/plugins/sensing/README.md | 15 +- src/plugins/sensing/core/CMakeLists.txt | 6 +- .../sensing/core/mock_jetson_encoder.cpp | 265 +++++++++++------- tests/cpp/plugins/sensing_mock/CMakeLists.txt | 5 +- .../sensing_mock/mock_test_support.hpp | 129 +++++++-- .../sensing_mock/test_mock_encoder.cpp | 150 +++++----- .../sensing_mock/test_mock_pipeline.cpp | 21 +- 8 files changed, 369 insertions(+), 224 deletions(-) diff --git a/src/plugins/sensing/CMakeLists.txt b/src/plugins/sensing/CMakeLists.txt index 3acff5b1bd..5f0ecf8d3f 100644 --- a/src/plugins/sensing/CMakeLists.txt +++ b/src/plugins/sensing/CMakeLists.txt @@ -6,7 +6,7 @@ # ============================================================================== # Two backends sit behind SiplCamera and JetsonEncoder: # JETSON SIPL capture and the V4L2 M2M encoder, which exist only on L4T. -# MOCK synthetic frames and a stand-in encoder, for any machine with an +# MOCK synthetic frames and an NVENC encoder (libavcodec), for any machine with an # NVIDIA GPU and the CUDA toolkit. Everything above those two classes # (SensingCamera, sinks, CUDA IPC, the app) is the same code. # SENSING_BACKEND=AUTO picks JETSON on aarch64 and MOCK elsewhere. diff --git a/src/plugins/sensing/README.md b/src/plugins/sensing/README.md index e36603a632..9560d9dee2 100644 --- a/src/plugins/sensing/README.md +++ b/src/plugins/sensing/README.md @@ -53,7 +53,7 @@ SIPL drivers. | Value | Capture | Encoder | Needs | | --- | --- | --- | --- | | `JETSON` | SIPL | V4L2 NVENC | L4T SDKs | -| `MOCK` | animated GIF frames | stand-in | NVIDIA GPU, CUDA toolkit, EGL headers (`libegl-dev`) | +| `MOCK` | animated GIF frames | NVENC through libavcodec | NVIDIA GPU, CUDA toolkit, EGL headers (`libegl-dev`), `libavcodec-dev` | | `AUTO` (default) | `JETSON` on aarch64, `MOCK` elsewhere | | | ```bash @@ -68,11 +68,12 @@ resolves `--list-sensors` and the link masks from the vendor JSON, builds a synt surface from a GIF frame, and runs the real conversion on it. Both sensors share a capture TSC, as the rig's fsync does. -The mock encoder's output is **not decodable video**: Annex-B start codes, SPS/PPS/IDR and -non-IDR NAL types, and a payload naming the frame. It reproduces what callers depend on: a -fixed pool of input buffers that makes `submit()` refuse frames when the encoder falls behind, -a warm-up delay before the first unit, and capture stamps that come back at microsecond -resolution. +The mock encoder produces real, decodable H.264 on the desktop GPU's NVENC (libavcodec's `h264_nvenc`), +configured as the Jetson encoder is: High profile, Level 5.2, VBR with a peak, an IDR every GOP with +SPS/PPS in front of it, no B-frames. It reproduces what callers depend on: a fixed pool of input +buffers that makes `submit()` refuse frames when the encoder falls behind, capture stamps that come +back at microsecond resolution, and the tail returned as one unit at end of stream. Bitstream bytes, +unit sizes and latency are NVENC's, not the Jetson engine's. Every sensor shows an animated GIF chosen by its **module name** (never its pipeline index, which the link mask renumbers), from the bundled Animated Noto Emoji in [`mock_assets/`](mock_assets/) (CC BY 4.0; @@ -97,8 +98,6 @@ Environment, read once at construction: | `SENSING_MOCK_FPS` | override the platform config's frame rate | | `SENSING_MOCK_SCALE_DIV` | divide the reported width and height (kept even); useful for tests | | `SENSING_MOCK_FAIL_AFTER` | every sensor fails after N frames, as a link drop would | -| `SENSING_MOCK_ENCODER_LATENCY_MS` | per-frame encode time | -| `SENSING_MOCK_ENCODER_DEPTH` | frames held back before the first unit (default 2) | A pass on the mock says nothing about SIPL, the ISP output format, NvSci buffer lifetime, the V4L2 encoder or real frame rates. Those are only checked on an Orin with SENSING cameras. diff --git a/src/plugins/sensing/core/CMakeLists.txt b/src/plugins/sensing/core/CMakeLists.txt index 98d06b3bc2..9ddc30adc8 100644 --- a/src/plugins/sensing/core/CMakeLists.txt +++ b/src/plugins/sensing/core/CMakeLists.txt @@ -117,7 +117,11 @@ if(SENSING_BACKEND_RESOLVED STREQUAL "JETSON") ${V4L2_LIBRARY} ) else() - target_link_libraries(sensing_core PRIVATE nlohmann_json::nlohmann_json) + # The mock encoder produces real H.264 with libavcodec's h264_nvenc (development package: + # libavcodec-dev; the desktop driver supplies NVENC). + find_package(PkgConfig REQUIRED) + pkg_check_modules(AVCODEC REQUIRED IMPORTED_TARGET libavcodec libavutil) + target_link_libraries(sensing_core PRIVATE nlohmann_json::nlohmann_json PkgConfig::AVCODEC) target_include_directories(sensing_core SYSTEM PRIVATE ${stb_image_SOURCE_DIR}) target_compile_definitions(sensing_core PRIVATE SENSING_MOCK_ASSET_DIR="${CMAKE_CURRENT_SOURCE_DIR}/../mock_assets") endif() diff --git a/src/plugins/sensing/core/mock_jetson_encoder.cpp b/src/plugins/sensing/core/mock_jetson_encoder.cpp index facd1a521e..20ff725fb4 100644 --- a/src/plugins/sensing/core/mock_jetson_encoder.cpp +++ b/src/plugins/sensing/core/mock_jetson_encoder.cpp @@ -1,33 +1,30 @@ // SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 // -// Mock backend for JetsonEncoder: no V4L2 and no NVENC. It runs the real -// RGBA-to-NV12 kernel and the same device-to-host copy as the V4L2 path, then -// a worker thread turns each frame into an Annex-B *shaped* unit. +// Mock backend for JetsonEncoder: real H.264 from the desktop GPU's NVENC, through libavcodec's +// h264_nvenc, set up as the Jetson V4L2 encoder is (High profile, Level 5.2, VBR with a peak, an IDR +// every GOP with SPS/PPS in front of it, no B-frames). It runs the real RGBA-to-NV12 kernel and the same +// device-to-host copy as the V4L2 path, then a worker thread feeds the NV12 buffer to the encoder. // -// The units are NOT decodable video. Each carries start codes, plausible NAL -// types (SPS/PPS/IDR on the first frame of every GOP, non-IDR slices between) -// and a payload that records the frame, so sinks and tests can follow a frame -// through. What it reproduces from the real encoder is the behaviour callers -// depend on: -// * a fixed pool of input buffers, so submit() returns false when the encoder -// falls behind rather than blocking; -// * warm-up: a unit appears only after more frames have been submitted; -// * the capture stamp comes back at microsecond resolution (a timeval), not -// the nanoseconds it went in with; -// * end_of_stream() returns the remaining tail as one concatenated unit. -// -// Environment (read once, at construction): -// SENSING_MOCK_ENCODER_LATENCY_MS=N per-frame encode time (default 0) -// SENSING_MOCK_ENCODER_DEPTH=N frames held back before output (default 2) +// Reproduced from the real encoder: a fixed pool of input buffers, so submit() returns false when the +// encoder falls behind rather than blocking; the capture stamp coming back at microsecond resolution (a +// V4L2 timeval); end_of_stream() returning the remaining tail as one concatenated unit. NVENC's own +// latency and unit sizes are what they are, and are not those of the Jetson engine. #include "jetson_encoder.hpp" #include "rgba_to_nv12.cuh" +extern "C" +{ +#include +#include +#include +#include +} + #include -#include #include -#include +#include #include #include #include @@ -51,24 +48,14 @@ void check_cuda(cudaError_t err, const char* what) throw std::runtime_error(std::string("mock JetsonEncoder: ") + what + ": " + cudaGetErrorString(err)); } -uint32_t env_uint(const char* name, uint32_t fallback) -{ - const char* value = std::getenv(name); - return value ? static_cast(std::strtoul(value, nullptr, 10)) : fallback; -} - -void append_start_code(std::vector& out) -{ - for (uint8_t byte : { 0, 0, 0, 1 }) - out.push_back(byte); -} - -/// One byte becomes two printable bytes, so a payload never contains the zero -/// bytes that would read as a start code. -void append_hex(std::vector& out, uint32_t value, int nibbles) +void check_av(int err, const char* what) { - for (int shift = (nibbles - 1) * 4; shift >= 0; shift -= 4) - out.push_back(static_cast('A' + ((value >> shift) & 0xF))); + if (err < 0) + { + char text[AV_ERROR_MAX_STRING_SIZE]; + av_strerror(err, text, sizeof(text)); + throw std::runtime_error(std::string("mock JetsonEncoder: ") + what + ": " + text); + } } } // namespace @@ -78,20 +65,32 @@ struct JetsonEncoder::Impl struct Job { uint32_t buffer = 0; - uint32_t frame = 0; int64_t timestamp_ns = 0; }; explicit Impl(const EncoderConfig& cfg) : config(cfg) { - gop = config.gop ? config.gop : (config.fps ? config.fps * 5 : 150); - latency = std::chrono::milliseconds(env_uint("SENSING_MOCK_ENCODER_LATENCY_MS", 0)); - depth = env_uint("SENSING_MOCK_ENCODER_DEPTH", 2); - - check_cuda(cudaMallocPitch(reinterpret_cast(&d_y), &d_y_pitch, config.width, config.height), - "cudaMallocPitch(Y)"); - check_cuda(cudaMallocPitch(reinterpret_cast(&d_uv), &d_uv_pitch, config.width, config.height / 2), - "cudaMallocPitch(UV)"); + const AVCodec* codec = avcodec_find_encoder_by_name("h264_nvenc"); + if (!codec) + throw std::runtime_error("mock JetsonEncoder: libavcodec has no h264_nvenc encoder"); + context = avcodec_alloc_context3(codec); + if (!context) + throw std::runtime_error("mock JetsonEncoder: cannot allocate the encoder context"); + + // A constructor that throws never runs the destructor, so free what was allocated here. + try + { + open(codec); + check_cuda(cudaMallocPitch(reinterpret_cast(&d_y), &d_y_pitch, config.width, config.height), + "cudaMallocPitch(Y)"); + check_cuda(cudaMallocPitch(reinterpret_cast(&d_uv), &d_uv_pitch, config.width, config.height / 2), + "cudaMallocPitch(UV)"); + } + catch (...) + { + release(); + throw; + } // Host staging, standing in for the V4L2 output plane's mapped buffers. const size_t nv12_bytes = static_cast(config.width) * config.height * 3 / 2; @@ -112,37 +111,89 @@ struct JetsonEncoder::Impl wake.notify_all(); if (worker.joinable()) worker.join(); + release(); + } + + void open(const AVCodec* codec) + { + const uint32_t fps = config.fps ? config.fps : 30; + const uint32_t gop = config.gop ? config.gop : fps * 5; + context->width = static_cast(config.width); + context->height = static_cast(config.height); + context->pix_fmt = AV_PIX_FMT_NV12; + // Microseconds, as a V4L2 timeval carries them. + context->time_base = { 1, 1'000'000 }; + context->framerate = { static_cast(fps), 1 }; + context->bit_rate = config.bitrate_bps; + context->gop_size = static_cast(gop); + context->max_b_frames = 0; + // Parameter sets stay in the stream, in front of every IDR, so a receiver can join mid-stream. + context->flags &= ~AV_CODEC_FLAG_GLOBAL_HEADER; + av_opt_set(context->priv_data, "profile", "high", 0); + av_opt_set(context->priv_data, "level", "5.2", 0); + av_opt_set(context->priv_data, "forced-idr", "1", 0); + av_opt_set(context->priv_data, "zerolatency", "1", 0); + if (config.peak_bitrate_bps > config.bitrate_bps) + { + context->rc_max_rate = config.peak_bitrate_bps; + context->rc_buffer_size = static_cast(config.peak_bitrate_bps); + av_opt_set(context->priv_data, "rc", "vbr", 0); + } + else + { + av_opt_set(context->priv_data, "rc", "cbr", 0); + } + check_av(avcodec_open2(context, codec, nullptr), "cannot open h264_nvenc"); + + frame = av_frame_alloc(); + packet = av_packet_alloc(); + if (!frame || !packet) + throw std::runtime_error("mock JetsonEncoder: cannot allocate a frame or packet"); + frame->format = AV_PIX_FMT_NV12; + frame->width = context->width; + frame->height = context->height; + check_av(av_frame_get_buffer(frame, 0), "cannot allocate the input frame"); + } + + void release() + { cudaFree(d_y); cudaFree(d_uv); + av_packet_free(&packet); + av_frame_free(&frame); + avcodec_free_context(&context); } - std::vector build_unit(const Job& job, uint32_t checksum) const + // Moves every packet the encoder has ready into `out`. Worker thread, or end_of_stream() once the + // worker is idle. + void collect(std::deque& out) { - std::vector unit; - if (job.frame % gop == 0) + for (;;) { - append_start_code(unit); - unit.push_back(0x67); // SPS - append_hex(unit, config.width, 4); - append_hex(unit, config.height, 4); - append_hex(unit, config.fps, 4); - append_start_code(unit); - unit.push_back(0x68); // PPS - append_hex(unit, 0, 2); + const int err = avcodec_receive_packet(context, packet); + if (err == AVERROR(EAGAIN) || err == AVERROR_EOF) + return; + check_av(err, "avcodec_receive_packet"); + EncodedUnit unit; + unit.data.assign(packet->data, packet->data + packet->size); + unit.timestamp_ns = packet->pts * 1000; + av_packet_unref(packet); + out.push_back(std::move(unit)); } - const bool idr = job.frame % gop == 0; - append_start_code(unit); - unit.push_back(idr ? 0x65 : 0x41); - for (char c : { 'M', 'O', 'C', 'K' }) - unit.push_back(static_cast(c)); - append_hex(unit, job.frame, 8); - append_hex(unit, checksum, 8); - - // Size follows the bitrate budget (capped), so throughput looks plausible. - const size_t budget = config.bitrate_bps / 8 / std::max(config.fps, 1); - const size_t target = std::min(std::max(budget, unit.size()), 64 * 1024); - unit.resize(target, 'U'); - return unit; + } + + void encode(const Job& job) + { + check_av(av_frame_make_writable(frame), "av_frame_make_writable"); + const uint8_t* nv12 = host_buffers[job.buffer].data(); + const size_t width = config.width; + for (uint32_t row = 0; row < config.height; ++row) + std::memcpy(frame->data[0] + static_cast(row) * frame->linesize[0], nv12 + row * width, width); + const uint8_t* uv = nv12 + width * config.height; + for (uint32_t row = 0; row < config.height / 2; ++row) + std::memcpy(frame->data[1] + static_cast(row) * frame->linesize[1], uv + row * width, width); + frame->pts = job.timestamp_ns / 1000; + check_av(avcodec_send_frame(context, frame), "avcodec_send_frame"); } void run() @@ -161,29 +212,25 @@ struct JetsonEncoder::Impl busy = true; } - if (latency.count() > 0) - std::this_thread::sleep_for(latency); - - // Cheap content fingerprint: a sparse sum over the NV12 bytes. - uint32_t checksum = 0; - const auto& nv12 = host_buffers[job.buffer]; - for (size_t i = 0; i < nv12.size(); i += 97) - checksum = checksum * 31u + nv12[i]; - - EncodedUnit unit; - unit.data = build_unit(job, checksum); - // A V4L2 timeval holds microseconds; the stamp does not survive finer. - unit.timestamp_ns = (job.timestamp_ns / 1000) * 1000; + std::deque produced; + std::string error; + try + { + encode(job); + collect(produced); + } + catch (const std::exception& e) + { + error = e.what(); + } { std::lock_guard lock(mutex); free_buffers.push_back(job.buffer); - held.push_back(std::move(unit)); - while (held.size() > depth) - { - ready.push_back(std::move(held.front())); - held.pop_front(); - } + for (auto& unit : produced) + ready.push_back(std::move(unit)); + if (!error.empty() && failure.empty()) + failure = error; busy = false; } idle.notify_all(); @@ -191,9 +238,10 @@ struct JetsonEncoder::Impl } EncoderConfig config; - uint32_t gop = 150; - uint32_t depth = 2; - std::chrono::milliseconds latency{ 0 }; + + AVCodecContext* context = nullptr; + AVFrame* frame = nullptr; + AVPacket* packet = nullptr; uint8_t* d_y = nullptr; uint8_t* d_uv = nullptr; @@ -201,17 +249,17 @@ struct JetsonEncoder::Impl size_t d_uv_pitch = 0; std::vector> host_buffers; - uint32_t submitted = 0; std::mutex mutex; std::condition_variable wake; std::condition_variable idle; std::deque free_buffers; std::deque jobs; - std::deque held; // encoded, still inside the warm-up window std::deque ready; // what poll() hands out + std::string failure; bool busy = false; bool closing = false; + bool flushed = false; std::thread worker; }; @@ -234,6 +282,8 @@ bool JetsonEncoder::submit(uintptr_t rgba_device_ptr, std::size_t row_pitch_byte uint32_t buffer = 0; { std::lock_guard lock(impl.mutex); + if (!impl.failure.empty()) + throw std::runtime_error(impl.failure); if (impl.free_buffers.empty()) return false; // encoder still busy; drop this frame rather than block buffer = impl.free_buffers.front(); @@ -256,7 +306,7 @@ bool JetsonEncoder::submit(uintptr_t rgba_device_ptr, std::size_t row_pitch_byte { std::lock_guard lock(impl.mutex); - impl.jobs.push_back({ buffer, impl.submitted++, timestamp_ns }); + impl.jobs.push_back({ buffer, timestamp_ns }); } impl.wake.notify_one(); return true; @@ -265,6 +315,8 @@ bool JetsonEncoder::submit(uintptr_t rgba_device_ptr, std::size_t row_pitch_byte EncodedUnit JetsonEncoder::poll() { std::lock_guard lock(m_impl->mutex); + if (!m_impl->failure.empty()) + throw std::runtime_error(m_impl->failure); if (m_impl->ready.empty()) return {}; auto out = std::move(m_impl->ready.front()); @@ -277,19 +329,24 @@ EncodedUnit JetsonEncoder::end_of_stream() auto& impl = *m_impl; std::unique_lock lock(impl.mutex); impl.idle.wait(lock, [&] { return impl.jobs.empty() && !impl.busy; }); + if (!impl.failure.empty()) + throw std::runtime_error(impl.failure); + + if (impl.flushed) + return {}; + impl.flushed = true; + + // Flush the encoder, then concatenate what is left, so the tail carries the stamp of the last unit. + check_av(avcodec_send_frame(impl.context, nullptr), "avcodec_send_frame(flush)"); + impl.collect(impl.ready); - // The warm-up window empties at EOS, then the tail is concatenated, so it - // carries the stamp of the last unit in it. EncodedUnit out; - for (auto* queue : { &impl.ready, &impl.held }) + for (auto& chunk : impl.ready) { - for (auto& chunk : *queue) - { - out.data.insert(out.data.end(), chunk.data.begin(), chunk.data.end()); - out.timestamp_ns = chunk.timestamp_ns; - } - queue->clear(); + out.data.insert(out.data.end(), chunk.data.begin(), chunk.data.end()); + out.timestamp_ns = chunk.timestamp_ns; } + impl.ready.clear(); return out; } diff --git a/tests/cpp/plugins/sensing_mock/CMakeLists.txt b/tests/cpp/plugins/sensing_mock/CMakeLists.txt index 96becb0f9b..1b13b47254 100644 --- a/tests/cpp/plugins/sensing_mock/CMakeLists.txt +++ b/tests/cpp/plugins/sensing_mock/CMakeLists.txt @@ -1,10 +1,12 @@ # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -# Exercises the plugin on the MOCK backend: synthetic frames and a stand-in +# Exercises the plugin on the MOCK backend: synthetic frames and an NVENC # encoder, but the real SensingCamera, sinks, CUDA kernels and CUDA IPC. Needs an # NVIDIA GPU, not a Jetson. find_package(CUDAToolkit REQUIRED) +find_package(PkgConfig REQUIRED) +pkg_check_modules(AVCODEC REQUIRED IMPORTED_TARGET libavcodec libavutil) add_executable(sensing_mock_tests gif_fixture.hpp @@ -22,6 +24,7 @@ target_link_libraries(sensing_mock_tests sensing::core CUDA::cudart CUDA::cuda_driver + PkgConfig::AVCODEC ) target_compile_definitions(sensing_mock_tests PRIVATE diff --git a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp index a7106ad011..4c1cab4224 100644 --- a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -9,6 +9,12 @@ #include #include +extern "C" +{ +#include +#include +} + #include #include #include @@ -216,54 +222,117 @@ inline Shot grab(plugins::sensing::SiplCamera& camera, uint32_t sensor) return shot; } -/// One NAL unit of the mock encoder's Annex-B output. +/// One NAL unit of an Annex-B stream. `type` is the nal_unit_type (7 SPS, 8 PPS, 5 IDR slice, 1 other +/// slice); `payload` follows the one-byte NAL header. struct Nal { uint8_t type = 0; std::vector payload; }; +constexpr uint8_t kNalSlice = 1; +constexpr uint8_t kNalIdr = 5; +constexpr uint8_t kNalSei = 6; +constexpr uint8_t kNalSps = 7; +constexpr uint8_t kNalPps = 8; + inline std::vector split_annex_b(const std::vector& data) { + // Offsets just past each 00 00 01 start code. + std::vector starts; + for (size_t i = 0; i + 3 <= data.size(); ++i) + { + if (data[i] == 0 && data[i + 1] == 0 && data[i + 2] == 1) + starts.push_back(i + 3); + } std::vector nals; - for (size_t i = 0; i + 4 < data.size();) + for (size_t n = 0; n < starts.size(); ++n) { - if (data[i] == 0 && data[i + 1] == 0 && data[i + 2] == 0 && data[i + 3] == 1) - { - Nal nal; - nal.type = data[i + 4]; - size_t end = i + 5; - while (end + 3 < data.size() && - !(data[end] == 0 && data[end + 1] == 0 && data[end + 2] == 0 && data[end + 3] == 1)) - ++end; - if (end + 3 >= data.size()) - end = data.size(); - nal.payload.assign(data.begin() + i + 5, data.begin() + end); - nals.push_back(std::move(nal)); - i = end; - } - else - { - ++i; - } + size_t end = n + 1 < starts.size() ? starts[n + 1] - 3 : data.size(); + while (end > starts[n] && data[end - 1] == 0 && n + 1 < starts.size()) + --end; // the zero of a four-byte start code + if (starts[n] >= end) + continue; + Nal nal; + nal.type = data[starts[n]] & 0x1F; + nal.payload.assign(data.begin() + starts[n] + 1, data.begin() + end); + nals.push_back(std::move(nal)); } return nals; } -/// The frame number the mock records after the "MOCK" magic, as 8 hex-like letters. -inline uint32_t frame_of(const Nal& slice) +inline bool is_slice(const Nal& nal) { - REQUIRE(slice.payload.size() >= 12); - REQUIRE(std::string(slice.payload.begin(), slice.payload.begin() + 4) == "MOCK"); - uint32_t value = 0; - for (int i = 0; i < 8; ++i) - value = (value << 4) | static_cast(slice.payload[4 + i] - 'A'); - return value; + return nal.type == kNalSlice || nal.type == kNalIdr; } -inline bool is_slice(const Nal& nal) +/// One decoded picture: its size and the mean of its luma plane. +struct Decoded +{ + int width = 0; + int height = 0; + double luma = 0; +}; + +/// Decodes an Annex-B stream with libavcodec's software H.264 decoder, in output order. +inline std::vector decode_annex_b(const std::vector& data) { - return nal.type == 0x65 || nal.type == 0x41; + const AVCodec* codec = avcodec_find_decoder(AV_CODEC_ID_H264); + REQUIRE(codec != nullptr); + AVCodecContext* context = avcodec_alloc_context3(codec); + AVCodecParserContext* parser = av_parser_init(AV_CODEC_ID_H264); + AVPacket* packet = av_packet_alloc(); + AVFrame* frame = av_frame_alloc(); + REQUIRE((context && parser && packet && frame)); + REQUIRE(avcodec_open2(context, codec, nullptr) == 0); + + std::vector out; + auto drain = [&] + { + while (avcodec_receive_frame(context, frame) == 0) + { + Decoded picture; + picture.width = frame->width; + picture.height = frame->height; + double sum = 0; + for (int y = 0; y < frame->height; ++y) + { + const uint8_t* row = frame->data[0] + static_cast(y) * frame->linesize[0]; + for (int x = 0; x < frame->width; ++x) + sum += row[x]; + } + picture.luma = sum / (static_cast(frame->width) * frame->height); + out.push_back(picture); + } + }; + size_t offset = 0; + while (offset < data.size()) + { + const int used = av_parser_parse2(parser, context, &packet->data, &packet->size, data.data() + offset, + static_cast(data.size() - offset), AV_NOPTS_VALUE, AV_NOPTS_VALUE, 0); + REQUIRE(used >= 0); + offset += static_cast(used); + if (packet->size > 0) + { + REQUIRE(avcodec_send_packet(context, packet) == 0); + drain(); + } + } + // The parser holds the last NAL until it sees what follows; flush it. + av_parser_parse2(parser, context, &packet->data, &packet->size, nullptr, 0, AV_NOPTS_VALUE, AV_NOPTS_VALUE, 0); + if (packet->size > 0) + { + REQUIRE(avcodec_send_packet(context, packet) == 0); + drain(); + } + avcodec_send_packet(context, nullptr); + drain(); + + av_frame_free(&frame); + av_packet_free(&packet); + av_parser_close(parser); + avcodec_free_context(&context); + return out; } } // namespace sensing_test diff --git a/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp b/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp index 041e7a7363..ac528986aa 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp @@ -17,21 +17,26 @@ using namespace std::chrono_literals; namespace { -constexpr uint32_t kWidth = 64; -constexpr uint32_t kHeight = 48; +constexpr uint32_t kWidth = 320; +constexpr uint32_t kHeight = 240; +/// A solid gray RGBA frame on the device. class Frame { public: Frame() { cudaMallocPitch(reinterpret_cast(&m_ptr), &m_pitch, kWidth * 4, kHeight); - cudaMemset2D(m_ptr, m_pitch, 0x80, kWidth * 4, kHeight); + set(0x80); } ~Frame() { cudaFree(m_ptr); } + void set(int gray) + { + cudaMemset2D(m_ptr, m_pitch, gray, kWidth * 4, kHeight); + } uintptr_t ptr() const { return reinterpret_cast(m_ptr); @@ -52,10 +57,26 @@ EncoderConfig make_config() config.width = kWidth; config.height = kHeight; config.fps = 30; - config.bitrate_bps = 1'000'000; + config.bitrate_bps = 4'000'000; + config.peak_bitrate_bps = 6'000'000; return config; } +// Submits `count` frames, waiting out a full input pool, then returns the whole stream. +std::vector encode_stream(JetsonEncoder& encoder, Frame& frame, int count, int64_t first_stamp_ns = 1'000'000'000) +{ + std::vector stream; + for (int i = 0; i < count; ++i) + { + REQUIRE(wait_for([&] { return encoder.submit(frame.ptr(), frame.pitch(), first_stamp_ns + i * 33'333'000); })); + for (auto unit = encoder.poll(); !unit.empty(); unit = encoder.poll()) + stream.insert(stream.end(), unit.data.begin(), unit.data.end()); + } + const auto tail = encoder.end_of_stream(); + stream.insert(stream.end(), tail.data.begin(), tail.data.end()); + return stream; +} + } // namespace TEST_CASE("encoder rejects zero or odd geometry", "[unit][gpu][sensing]") @@ -68,40 +89,52 @@ TEST_CASE("encoder rejects zero or odd geometry", "[unit][gpu][sensing]") CHECK_THROWS_AS(JetsonEncoder(config), std::runtime_error); } -TEST_CASE("encoder holds frames back during warm-up, then emits them in order", "[unit][gpu][sensing]") +TEST_CASE("the stream opens with SPS, PPS and an IDR slice, High profile at Level 5.2", "[unit][gpu][sensing]") { Frame frame; JetsonEncoder encoder(make_config()); + std::vector nals; + for (auto& nal : split_annex_b(encode_stream(encoder, frame, 3))) + { + if (nal.type != kNalSei) // the encoder may add SEI; it is not part of the shape under test + nals.push_back(std::move(nal)); + } - REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000)); - std::this_thread::sleep_for(100ms); - CHECK(encoder.poll().empty()); // nothing yet: the encoder wants more input first - - REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'033'000'000)); - REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'066'000'000)); - REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'100'000'000)); - - std::vector units; - REQUIRE(wait_for( - [&] - { - for (auto unit = encoder.poll(); !unit.empty(); unit = encoder.poll()) - units.push_back(std::move(unit)); - return units.size() >= 2; - })); - - // The first unit opens the stream: SPS, PPS, IDR slice. - const auto first = split_annex_b(units[0].data); - REQUIRE(first.size() == 3); - CHECK(first[0].type == 0x67); - CHECK(first[1].type == 0x68); - CHECK(first[2].type == 0x65); - CHECK(frame_of(first[2]) == 0); - - const auto second = split_annex_b(units[1].data); - REQUIRE(second.size() == 1); - CHECK(second[0].type == 0x41); - CHECK(frame_of(second[0]) == 1); + REQUIRE(nals.size() >= 4); + CHECK(nals[0].type == kNalSps); + CHECK(nals[1].type == kNalPps); + CHECK(nals[2].type == kNalIdr); + CHECK(nals[3].type == kNalSlice); + REQUIRE(nals[0].payload.size() >= 3); + CHECK(nals[0].payload[0] == 100); // profile_idc: High + CHECK(nals[0].payload[2] == 52); // level_idc: 5.2 +} + +TEST_CASE("the output decodes, in order, at the configured size", "[unit][gpu][sensing]") +{ + Frame frame; + JetsonEncoder encoder(make_config()); + + // Four frames of rising brightness, so order and content can both be read back. + std::vector stream; + for (int i = 0; i < 4; ++i) + { + frame.set(40 + 50 * i); + REQUIRE(wait_for([&] { return encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i * 33'333'000); })); + // The next frame is only written once the copy of this one is done: submit() has returned. + } + const auto tail = encoder.end_of_stream(); + stream.insert(stream.end(), tail.data.begin(), tail.data.end()); + + const auto pictures = decode_annex_b(stream); + REQUIRE(pictures.size() == 4); + for (size_t i = 0; i < pictures.size(); ++i) + { + CHECK(pictures[i].width == static_cast(kWidth)); + CHECK(pictures[i].height == static_cast(kHeight)); + if (i > 0) + CHECK(pictures[i].luma > pictures[i - 1].luma + 20); + } } TEST_CASE("encoder returns the capture stamp at microsecond resolution", "[unit][gpu][sensing]") @@ -118,43 +151,23 @@ TEST_CASE("encoder returns the capture stamp at microsecond resolution", "[unit] CHECK(tail.timestamp_ns != stamp); } -TEST_CASE("encoder marks an IDR at every GOP boundary", "[unit][gpu][sensing]") +TEST_CASE("encoder marks an IDR at every GOP boundary, with parameter sets in front", "[unit][gpu][sensing]") { Frame frame; auto config = make_config(); config.gop = 3; JetsonEncoder encoder(config); - // Seven frames against six input buffers: submit() may refuse one until the encoder frees a buffer. - for (int i = 0; i < 7; ++i) - REQUIRE(wait_for([&] { return encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i * 33'000'000); })); - const auto nals = split_annex_b(encoder.end_of_stream().data); - std::vector idr_frames; - for (const auto& nal : nals) + std::vector slice_types; + size_t parameter_sets = 0; + for (const auto& nal : split_annex_b(encode_stream(encoder, frame, 7))) { - if (nal.type == 0x65) - idr_frames.push_back(frame_of(nal)); + if (is_slice(nal)) + slice_types.push_back(nal.type); + parameter_sets += nal.type == kNalSps ? 1 : 0; } - CHECK(idr_frames == std::vector{ 0, 3, 6 }); -} - -TEST_CASE("encoder refuses frames when every input buffer is in flight", "[unit][gpu][sensing]") -{ - ScopedEnv latency("SENSING_MOCK_ENCODER_LATENCY_MS", "300"); - Frame frame; - JetsonEncoder encoder(make_config()); - - int accepted = 0; - int rejected = 0; - for (int i = 0; i < 10; ++i) - (encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i) ? accepted : rejected)++; - - // Six input buffers, and the encoder is too slow to return any in the loop. - CHECK(accepted == 6); - CHECK(rejected == 4); - - // Once the encoder frees one, submission works again. - CHECK(wait_for([&] { return encoder.submit(frame.ptr(), frame.pitch(), 2'000'000'000); }, 2s)); + CHECK(slice_types == std::vector{ kNalIdr, kNalSlice, kNalSlice, kNalIdr, kNalSlice, kNalSlice, kNalIdr }); + CHECK(parameter_sets == 3); } TEST_CASE("end of stream returns the tail as one unit carrying the last stamp", "[unit][gpu][sensing]") @@ -165,14 +178,11 @@ TEST_CASE("end of stream returns the tail as one unit carrying the last stamp", REQUIRE(encoder.submit(frame.ptr(), frame.pitch(), 1'000'000'000 + i * 1'000'000)); const auto tail = encoder.end_of_stream(); - std::vector frames; - for (const auto& nal : split_annex_b(tail.data)) - { - if (is_slice(nal)) - frames.push_back(frame_of(nal)); - } - CHECK(frames == std::vector{ 0, 1, 2, 3 }); CHECK(tail.timestamp_ns == 1'003'000'000); + size_t slices = 0; + for (const auto& nal : split_annex_b(tail.data)) + slices += is_slice(nal) ? 1 : 0; + CHECK(slices >= 1); CHECK(encoder.poll().empty()); CHECK(encoder.end_of_stream().empty()); diff --git a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp index df7b3e47a8..9fe7de270a 100644 --- a/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp +++ b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp @@ -129,24 +129,27 @@ TEST_CASE("an encoded stream writes Annex-B data and one metadata record per uni const auto nals = split_annex_b(bytes); REQUIRE_FALSE(nals.empty()); - CHECK(nals[0].type == 0x67); // the stream opens with its parameter sets - std::vector frames; + CHECK(nals[0].type == kNalSps); // the stream opens with its parameter sets + size_t slices = 0; for (const auto& nal : nals) + slices += is_slice(nal) ? 1 : 0; + REQUIRE(slices >= 10); + + // It is real video: every slice decodes, at the sensor's size. + const auto pictures = decode_annex_b(bytes); + CHECK(pictures.size() == slices); + for (const auto& picture : pictures) { - if (is_slice(nal)) - frames.push_back(frame_of(nal)); + CHECK(picture.width == 320); + CHECK(picture.height == 248); } - REQUIRE(frames.size() >= 10); - // The encoder preserves submission order; the mailbox may skip frames but never reorders. - for (size_t i = 1; i < frames.size(); ++i) - CHECK(frames[i] > frames[i - 1]); const auto records = records_for(*recorder, 0); REQUIRE(records.size() >= 5); for (size_t i = 0; i < records.size(); ++i) CHECK(records[i].sequence_number == i); // contiguous from zero, one per dispatched unit // The final flush folds the warm-up tail into one unit, so units can be fewer than frames. - CHECK(records.size() <= frames.size()); + CHECK(records.size() <= slices); } TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", "[unit][gpu][sensing]") From 0c8c07abbc1e4c69073528358d6ae51b6679cbc9 Mon Sep 17 00:00:00 2001 From: Gareth Morgan Date: Sun, 25 Oct 2026 00:35:39 -0400 Subject: [PATCH 10/10] Tighten the mock encoder's setup and the test decoder Cover the whole constructor with the cleanup guard, check the encoder option calls, and make the test decoder pad its input and fail on any decoder error or damaged picture. Co-Authored-By: Claude Sonnet 5.5 Signed-off-by: Gareth Morgan --- .../sensing/core/mock_jetson_encoder.cpp | 48 +++++++++++++------ .../sensing_mock/mock_test_support.hpp | 45 ++++++++++------- 2 files changed, 62 insertions(+), 31 deletions(-) diff --git a/src/plugins/sensing/core/mock_jetson_encoder.cpp b/src/plugins/sensing/core/mock_jetson_encoder.cpp index 20ff725fb4..e6a30853f9 100644 --- a/src/plugins/sensing/core/mock_jetson_encoder.cpp +++ b/src/plugins/sensing/core/mock_jetson_encoder.cpp @@ -85,21 +85,21 @@ struct JetsonEncoder::Impl "cudaMallocPitch(Y)"); check_cuda(cudaMallocPitch(reinterpret_cast(&d_uv), &d_uv_pitch, config.width, config.height / 2), "cudaMallocPitch(UV)"); + + // Host staging, standing in for the V4L2 output plane's mapped buffers. + const size_t nv12_bytes = static_cast(config.width) * config.height * 3 / 2; + for (uint32_t i = 0; i < kInputBuffers; ++i) + { + host_buffers.emplace_back(nv12_bytes); + free_buffers.push_back(i); + } + worker = std::thread([this] { run(); }); } catch (...) { release(); throw; } - - // Host staging, standing in for the V4L2 output plane's mapped buffers. - const size_t nv12_bytes = static_cast(config.width) * config.height * 3 / 2; - for (uint32_t i = 0; i < kInputBuffers; ++i) - { - host_buffers.emplace_back(nv12_bytes); - free_buffers.push_back(i); - } - worker = std::thread([this] { run(); }); } ~Impl() @@ -129,19 +129,37 @@ struct JetsonEncoder::Impl context->max_b_frames = 0; // Parameter sets stay in the stream, in front of every IDR, so a receiver can join mid-stream. context->flags &= ~AV_CODEC_FLAG_GLOBAL_HEADER; - av_opt_set(context->priv_data, "profile", "high", 0); - av_opt_set(context->priv_data, "level", "5.2", 0); - av_opt_set(context->priv_data, "forced-idr", "1", 0); - av_opt_set(context->priv_data, "zerolatency", "1", 0); + check_av(av_opt_set(context->priv_data, "profile", "high", 0), + "av_opt_set(" + "profile" + ")"); + check_av(av_opt_set(context->priv_data, "level", "5.2", 0), + "av_opt_set(" + "level" + ")"); + check_av(av_opt_set(context->priv_data, "forced-idr", "1", 0), + "av_opt_set(" + "forced-idr" + ")"); + check_av(av_opt_set(context->priv_data, "zerolatency", "1", 0), + "av_opt_set(" + "zerolatency" + ")"); if (config.peak_bitrate_bps > config.bitrate_bps) { context->rc_max_rate = config.peak_bitrate_bps; context->rc_buffer_size = static_cast(config.peak_bitrate_bps); - av_opt_set(context->priv_data, "rc", "vbr", 0); + check_av(av_opt_set(context->priv_data, "rc", "vbr", 0), + "av_opt_set(" + "rc" + ")"); } else { - av_opt_set(context->priv_data, "rc", "cbr", 0); + check_av(av_opt_set(context->priv_data, "rc", "cbr", 0), + "av_opt_set(" + "rc" + ")"); } check_av(avcodec_open2(context, codec, nullptr), "cannot open h264_nvenc"); diff --git a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp index 4c1cab4224..2bd18c035f 100644 --- a/tests/cpp/plugins/sensing_mock/mock_test_support.hpp +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -274,9 +274,15 @@ struct Decoded double luma = 0; }; -/// Decodes an Annex-B stream with libavcodec's software H.264 decoder, in output order. -inline std::vector decode_annex_b(const std::vector& data) +/// Decodes an Annex-B stream with libavcodec's software H.264 decoder, in output order. Any parser or +/// decoder error, or a picture flagged as damaged, fails the calling test. +inline std::vector decode_annex_b(const std::vector& stream) { + // The parser reads a little past the end of its input. + std::vector data(stream); + data.resize(stream.size() + AV_INPUT_BUFFER_PADDING_SIZE, 0); + const size_t size = stream.size(); + const AVCodec* codec = avcodec_find_decoder(AV_CODEC_ID_H264); REQUIRE(codec != nullptr); AVCodecContext* context = avcodec_alloc_context3(codec); @@ -289,8 +295,14 @@ inline std::vector decode_annex_b(const std::vector& data) std::vector out; auto drain = [&] { - while (avcodec_receive_frame(context, frame) == 0) + for (;;) { + const int err = avcodec_receive_frame(context, frame); + if (err == AVERROR(EAGAIN) || err == AVERROR_EOF) + return; + REQUIRE(err == 0); + CHECK((frame->flags & AV_FRAME_FLAG_CORRUPT) == 0); + CHECK((frame->decode_error_flags) == 0); Decoded picture; picture.width = frame->width; picture.height = frame->height; @@ -305,27 +317,28 @@ inline std::vector decode_annex_b(const std::vector& data) out.push_back(picture); } }; - size_t offset = 0; - while (offset < data.size()) + auto send = [&] { - const int used = av_parser_parse2(parser, context, &packet->data, &packet->size, data.data() + offset, - static_cast(data.size() - offset), AV_NOPTS_VALUE, AV_NOPTS_VALUE, 0); - REQUIRE(used >= 0); - offset += static_cast(used); if (packet->size > 0) { REQUIRE(avcodec_send_packet(context, packet) == 0); drain(); } - } - // The parser holds the last NAL until it sees what follows; flush it. - av_parser_parse2(parser, context, &packet->data, &packet->size, nullptr, 0, AV_NOPTS_VALUE, AV_NOPTS_VALUE, 0); - if (packet->size > 0) + }; + size_t offset = 0; + while (offset < size) { - REQUIRE(avcodec_send_packet(context, packet) == 0); - drain(); + const int used = av_parser_parse2(parser, context, &packet->data, &packet->size, data.data() + offset, + static_cast(size - offset), AV_NOPTS_VALUE, AV_NOPTS_VALUE, 0); + REQUIRE(used >= 0); + offset += static_cast(used); + send(); } - avcodec_send_packet(context, nullptr); + // The parser holds the last NAL until it sees what follows; flush it, then the decoder. + REQUIRE(av_parser_parse2( + parser, context, &packet->data, &packet->size, nullptr, 0, AV_NOPTS_VALUE, AV_NOPTS_VALUE, 0) >= 0); + send(); + REQUIRE(avcodec_send_packet(context, nullptr) == 0); drain(); av_frame_free(&frame);