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/core/schema/golden/sensing.bfbs b/src/core/schema/golden/sensing.bfbs new file mode 100644 index 0000000000..a506fce2a5 Binary files /dev/null and b/src/core/schema/golden/sensing.bfbs differ diff --git a/src/plugins/sensing/CMakeLists.txt b/src/plugins/sensing/CMakeLists.txt index a168711c0c..5f0ecf8d3f 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 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. # # 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..9560d9dee2 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,59 @@ 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` | animated GIF frames | NVENC through libavcodec | NVIDIA GPU, CUDA toolkit, EGL headers (`libegl-dev`), `libavcodec-dev` | +| `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, 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 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; +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 (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, +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_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 | + +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..9ddc30adc8 100644 --- a/src/plugins/sensing/core/CMakeLists.txt +++ b/src/plugins/sensing/core/CMakeLists.txt @@ -1,17 +1,70 @@ # 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() + # 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) + 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() + # 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 + ) +endif() + +add_library(sensing_core STATIC + ${SENSING_COMMON_SOURCES} + ${SENSING_BACKEND_SOURCES} ) add_library(sensing::core ALIAS sensing_core) @@ -27,19 +80,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 +88,40 @@ 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() + # 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_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_jetson_encoder.cpp b/src/plugins/sensing/core/mock_jetson_encoder.cpp new file mode 100644 index 0000000000..e6a30853f9 --- /dev/null +++ b/src/plugins/sensing/core/mock_jetson_encoder.cpp @@ -0,0 +1,372 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// 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. +// +// 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 +#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)); +} + +void check_av(int err, const char* what) +{ + 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 + +struct JetsonEncoder::Impl +{ + struct Job + { + uint32_t buffer = 0; + int64_t timestamp_ns = 0; + }; + + explicit Impl(const EncoderConfig& cfg) : config(cfg) + { + 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)"); + + // 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; + } + } + + ~Impl() + { + { + std::lock_guard lock(mutex); + closing = true; + } + 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; + 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); + check_av(av_opt_set(context->priv_data, "rc", "vbr", 0), + "av_opt_set(" + "rc" + ")"); + } + else + { + 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"); + + 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); + } + + // 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) + { + for (;;) + { + 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)); + } + } + + 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() + { + for (;;) + { + Job job; + { + std::unique_lock lock(mutex); + wake.wait(lock, [&] { return closing || !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(); + busy = true; + } + + 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); + for (auto& unit : produced) + ready.push_back(std::move(unit)); + if (!error.empty() && failure.empty()) + failure = error; + busy = false; + } + idle.notify_all(); + } + } + + EncoderConfig config; + + AVCodecContext* context = nullptr; + AVFrame* frame = nullptr; + AVPacket* packet = nullptr; + + 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; + + std::mutex mutex; + std::condition_variable wake; + std::condition_variable idle; + std::deque free_buffers; + std::deque jobs; + std::deque ready; // what poll() hands out + std::string failure; + bool busy = false; + bool closing = false; + bool flushed = 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.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(); + 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, timestamp_ns }); + } + impl.wake.notify_one(); + return true; +} + +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()); + 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; }); + 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); + + EncodedUnit out; + for (auto& chunk : impl.ready) + { + out.data.insert(out.data.end(), chunk.data.begin(), chunk.data.end()); + out.timestamp_ns = chunk.timestamp_ns; + } + impl.ready.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..188f1c5113 --- /dev/null +++ b/src/plugins/sensing/core/mock_nv12_source.cu @@ -0,0 +1,166 @@ +// 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 +{ + +// 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 +{ + 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]); +} + +// 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 = 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) + { + 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); + } + + 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 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); +} + +__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); +} + +// 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, + bool mirror, + int shift_x) +{ + 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; + + 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; + // 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; + 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_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, + bool mirror, + int shift_x, + cudaStream_t stream) +{ + 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 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, mirror, shift_x); +} + + +} // 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..1a79a37a58 --- /dev/null +++ b/src/plugins/sensing/core/mock_nv12_source.cuh @@ -0,0 +1,42 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +namespace plugins +{ +namespace sensing +{ + +/** + * @brief Paint one synthetic ISP0 frame from a GIF frame already in device memory. + * + * 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. + * @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, + int src_width, + int src_height, + uint8_t* y_plane, + int y_pitch, + uint8_t* uv_plane, + int uv_pitch, + int width, + int height, + bool mirror, + int shift_x, + 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..bf9cfd24d8 --- /dev/null +++ b/src/plugins/sensing/core/mock_platform_config.cpp @@ -0,0 +1,109 @@ +// 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()) + { + // 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; + } + + 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..491b2c1911 --- /dev/null +++ b/src/plugins/sensing/core/mock_sipl_camera.cpp @@ -0,0 +1,417 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// 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. +// +// 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 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 +// 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 "mock_test_hooks.hpp" +#include "sipl_camera.hpp" +#include "yuv_to_rgba.cuh" + +#include +#include +#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; + +// 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() +{ + 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)); +} + +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 + { + 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::shared_ptr gif; + bool mirror = false; + int shift_x = 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) + { + 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); + choose_picture(info, *sensor); + sensors.push_back(std::move(sensor)); + } + } + + ~Impl() + { + stop(); + } + + 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"); + } + + /// 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; + + std::string path = g_test_picture; + if (path.empty()) + { + if (!look) + 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) + { + // 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); + sensor.gif = slot; + } + + 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, 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_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. + 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 * period_ns / 1e9, 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 }; + std::map> gif_cache; + 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/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/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/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/mock_rig.json b/src/plugins/sensing/mock_assets/mock_rig.json new file mode 100644 index 0000000000..ce0753682b --- /dev/null +++ b/src/plugins/sensing/mock_assets/mock_rig.json @@ -0,0 +1,75 @@ +{ + "cameraConfigs": [ + { + "name": "MOCK_S56C_2_SHF3L", + "cameraModules": [ + { + "name": "S56C_HEAD", + "linkIndex": 0, + "transportIndex": 1 + }, + { + "name": "SHF3L_LEFT", + "linkIndex": 1, + "transportIndex": 1 + }, + { + "name": "SHF3L_RIGHT", + "linkIndex": 2, + "transportIndex": 1 + } + ] + }, + { + "name": "S56C_HEAD", + "sensorInfo": [ + { + "virtualChannels": [ + { + "width": 2560, + "height": 1984, + "fps": 30.0 + } + ] + }, + { + "virtualChannels": [ + { + "width": 2560, + "height": 1984, + "fps": 30.0 + } + ] + } + ] + }, + { + "name": "SHF3L_LEFT", + "sensorInfo": [ + { + "virtualChannels": [ + { + "width": 1920, + "height": 1536, + "fps": 30.0 + } + ] + } + ] + }, + { + "name": "SHF3L_RIGHT", + "sensorInfo": [ + { + "virtualChannels": [ + { + "width": 1920, + "height": 1536, + "fps": 30.0 + } + ] + } + ] + } + ] +} 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/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/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..1b13b47254 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/CMakeLists.txt @@ -0,0 +1,36 @@ +# 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 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 + 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 +) + +target_link_libraries(sensing_mock_tests + PRIVATE + Catch2::Catch2WithMain + sensing::core + CUDA::cudart + CUDA::cuda_driver + PkgConfig::AVCODEC +) + +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" +) + +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 new file mode 100644 index 0000000000..2bd18c035f --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/mock_test_support.hpp @@ -0,0 +1,351 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "gif_fixture.hpp" + +#include +#include +#include + +extern "C" +{ +#include +#include +} + +#include +#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) + { + 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() + { + ::unsetenv(m_name.c_str()); + } + ScopedEnv(const ScopedEnv&) = delete; + ScopedEnv& operator=(const ScopedEnv&) = delete; + +private: + 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)) +{ + 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; +} + +/// 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; +}; + +/// 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") + { + 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 = std::make_unique(m_gif.path()); + } + +private: + TempFile m_gif; + ScopedEnv m_scale; + std::unique_ptr 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) +{ + 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; +} + +/// Waits for one frame from `sensor` and copies it to the host. +inline Shot grab(plugins::sensing::SiplCamera& camera, uint32_t sensor) +{ + 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 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 n = 0; n < starts.size(); ++n) + { + 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; +} + +inline bool is_slice(const Nal& nal) +{ + return nal.type == kNalSlice || nal.type == kNalIdr; +} + +/// 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. 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); + 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 = [&] + { + 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; + 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); + } + }; + auto send = [&] + { + if (packet->size > 0) + { + REQUIRE(avcodec_send_packet(context, packet) == 0); + drain(); + } + }; + size_t offset = 0; + while (offset < size) + { + 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(); + } + // 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); + 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_cuda_ipc_consumer.cpp b/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp new file mode 100644 index 0000000000..95cf117abc --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_cuda_ipc_consumer.cpp @@ -0,0 +1,293 @@ +// 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]") +{ + FastMock fast; + 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 whole, in order, and the picture moves", "[unit][gpu][sensing]") +{ + FastMock fast; + 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); + REQUIRE(consumer.hello().pitch == static_cast(consumer.hello().width) * 4); + + uint64_t last_sequence = 0; + uint64_t last_stamp = 0; + std::vector colours; + for (int i = 0; i < 20; ++i) + { + ipc::FrameReady frame{}; + REQUIRE(consumer.next_frame(frame, 2000ms)); + 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); + CHECK(is_whole(shot)); // not torn between two captures + last_sequence = frame.sequence; + last_stamp = frame.timestamp_ns; + 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]") +{ + FastMock fast; + 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..9fc23779c1 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_camera.cpp @@ -0,0 +1,229 @@ +// 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 whole, with rising sequence and stamps, and the picture moves", "[unit][gpu][sensing]") +{ + FastMock fast; + 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; + std::vector shots; + 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; + + 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]") +{ + FastMock fast; + 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]") +{ + FastMock fast; + SiplCamera camera(make_sipl_config()); + camera.start(); + + std::optional held; + REQUIRE(wait_for([&] { return (held = camera.latest(0)).has_value(); })); + 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); + } + + // 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]") +{ + FastMock fast; + 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]") +{ + FastMock fast; + 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..ac528986aa --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_encoder.cpp @@ -0,0 +1,189 @@ +// 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 = 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); + 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); + } + 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 = 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]") +{ + 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("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(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]") +{ + 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, with parameter sets in front", "[unit][gpu][sensing]") +{ + Frame frame; + auto config = make_config(); + config.gop = 3; + JetsonEncoder encoder(config); + + std::vector slice_types; + size_t parameter_sets = 0; + for (const auto& nal : split_annex_b(encode_stream(encoder, frame, 7))) + { + if (is_slice(nal)) + slice_types.push_back(nal.type); + parameter_sets += nal.type == kNalSps ? 1 : 0; + } + 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]") +{ + 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(); + 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_gif.cpp b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp new file mode 100644 index 0000000000..1413427570 --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_gif.cpp @@ -0,0 +1,523 @@ +// 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 +#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"); + ScopedPicture 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"); + ScopedPicture source(gif.path()); + 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"); + ScopedPicture source(gif.path()); + + 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"); + ScopedPicture source(gif.path()); + + 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"); + ScopedPicture source(gif.path()); + + 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"); + ScopedPicture source(gif.path()); + + 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) }); + ScopedPicture source(gif.path()); + 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); +} + +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); + } +} + +namespace +{ + +// 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( + [&] + { + bool done = true; + for (size_t s = 0; s < frames.size(); ++s) + { + 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); + frames[s][view->sequence] = std::move(shot); + } + done = done && frames[s].size() >= count; + } + return done; + }, + 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("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"); + ScopedPicture source(gif.path()); + + SiplCamera camera(make_sipl_config()); + camera.start(); + const Shot first = grab(camera, 0); + const Shot second = grab(camera, 1); + camera.stop(); + + // 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)); + CHECK(near(shot->at(shot->width * 3 / 4, shot->height / 2), 0, 0, 255)); + } +} + +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 is an error", "[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("no mock picture for module 'UNKNOWN_CAM'")); +} 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..9fe7de270a --- /dev/null +++ b/tests/cpp/plugins/sensing_mock/test_mock_pipeline.cpp @@ -0,0 +1,215 @@ +// 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 + +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]") +{ + FastMock fast; + 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]") +{ + FastMock fast; + TempFile output_file(".h264"); + const std::string& output = output_file.path(); + + 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); + REQUIRE(bytes.size() > 4); + + const auto nals = split_annex_b(bytes); + REQUIRE_FALSE(nals.empty()); + 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) + { + CHECK(picture.width == 320); + CHECK(picture.height == 248); + } + + 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() <= slices); +} + +TEST_CASE("an ipc-only stream records metadata with the sensor's capture TSC", "[unit][gpu][sensing]") +{ + FastMock fast; + 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); +} + +TEST_CASE("an encoded stream keeps its capture TSC through the encoder", "[unit][gpu][sensing]") +{ + FastMock fast; + TempFile output_file(".h264"); + const std::string& output = output_file.path(); + + 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); + } + + 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); +}