Add continuous TensorRT camera inference and ready-to-run Orin package

2026-09-23 11:28:29 +08:00
commit 71723ef3b4
53 changed files with 27755 additions and 0 deletions
@@ -0,0 +1,5 @@
+*.sh text eol=lf
+*.engine -text
+*.ko -text
+/orin/** -text
+/config/door_mask.json -text
@@ -0,0 +1,4 @@
+/build/
+/engines/
+/logs/
+*.tmp
@@ -0,0 +1,23 @@
+cmake_minimum_required(VERSION 3.24)
+if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
+  set(CMAKE_CUDA_ARCHITECTURES native)
+endif()
+project(camera_risk LANGUAGES CXX CUDA)
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CUDA_STANDARD 17)
+find_package(CUDAToolkit REQUIRED)
+find_package(OpenCV REQUIRED COMPONENTS core imgproc imgcodecs highgui)
+find_package(Threads REQUIRED)
+find_path(TENSORRT_INCLUDE_DIR NvInfer.h PATH_SUFFIXES aarch64-linux-gnu x86_64-linux-gnu REQUIRED)
+find_library(TENSORRT_LIBRARY nvinfer REQUIRED)
+add_library(camera_risk_core
+  src/trt_engine.cpp src/pose_engine.cu src/gpu_preprocess.cu src/uyvy_converter.cu
+  src/detector.cpp src/door_risk.cpp src/camera.cpp src/camera_pipeline.cpp src/config.cpp)
+target_include_directories(camera_risk_core PUBLIC include third_party ${TENSORRT_INCLUDE_DIR} ${OpenCV_INCLUDE_DIRS})
+target_link_libraries(camera_risk_core PUBLIC ${TENSORRT_LIBRARY} ${OpenCV_LIBS} Threads::Threads CUDA::cudart CUDA::nppc CUDA::nppig CUDA::nppidei CUDA::nppicc)
+target_compile_options(camera_risk_core PRIVATE $<$<COMPILE_LANGUAGE:CXX>:-Wall;-Wextra;-Wpedantic>)
+add_library(camera_risk_visualizer src/visualizer.cpp)
+target_link_libraries(camera_risk_visualizer PUBLIC camera_risk_core)
+add_executable(camera_risk src/main.cpp)
+target_link_libraries(camera_risk PRIVATE camera_risk_visualizer)
@@ -0,0 +1,76 @@
+# Vision Anti-Pinch
+
+C++ / TensorRT 实时摄像头风险推理，算法来自最新 run_risk.py。
+持续取图、推理并更新内存中的 FrameResult；可视化可选。
+
+## Orin 直接运行
+
+适用环境：Jetson AGX Orin、JetPack 5.1.2 / L4T R35.4.1、CUDA 11.4、TensorRT 8.5.2、OpenCV 4.5.4。
+仓库已包含该板端编译的程序和重新构建的 TensorRT 引擎，无需先编译或转换。
+
+```bash
+git clone https://git.kendeji.fun/zmv/vision-anti-pinch.git
+cd vision-anti-pinch/orin
+./scripts/initialize_camera.sh  # 开机后初始化 /dev/video5 摄像头链路
+./run.sh                      # 持续推理，结果只保存在内存
+```
+
+可选板端本地可视化：
+
+```bash
+./run.sh --visualize
+```
+
+Ctrl+C / SIGTERM 停止程序。没有有限帧选项；正常运行和退出不输出统计。
+关闭预览窗口或按 q/Esc 只关闭预览，推理继续。没有图像网络服务、端口转发、JSONL 或视频输出。
+
+摄像头为 /dev/video5，1920×1536 UYVY，30 FPS。
+当前 door mask 是指定的 constant_door_mask.json 副本，标定状态为 false；装车后应更换与实际画面匹配的门区并更新标定状态。
+即使不开启可视化，风险计算仍需要门区。
+
+## 下游获取结果
+
+main.cpp 注释中也包含以下示例。下游通过自己的调用位置或线程读取：
+
+```cpp
+auto result = pipeline.latest();
+if (result) {
+    const auto& risk = result->risk;
+    const auto& people = result->persons;
+    // result->frame_index / track_ids / boxes
+    // result->image / door_mask_pixels / capture / timing
+    // people[i].selected / smoothed / temporal：各 133 个关键点
+}
+```
+
+latest() 返回 shared_ptr<const FrameResult>；首帧尚未完成时为空。
+读取不删除结果，也不产生消费状态。推理不关心是否有下游读取。
+下游持有的结果保持有效，后续更新不会改变旧快照；历史保存由下游自行负责。
+公共定义位于 include/camera_pipeline.hpp 和 include/frame_result.hpp。
+
+## 从源码构建
+
+```bash
+./scripts/build.sh
+./scripts/initialize_camera.sh
+./run_camera.sh
+```
+
+源码构建要求 CMake 3.24+、CUDA/NPP、TensorRT 和 OpenCV 开发依赖。
+构建脚本会优先使用 Orin 上已有的 ~/miniconda3/envs/yolo/bin/cmake。
+源码配置直接使用 orin/engines/ 中的板端引擎。该引擎仅适用于上述 Orin 环境；仓库不包含 ONNX 或模型转换脚本。
+独立库目标为 camera_risk_core；可视化在 camera_risk_visualizer。
+
+## 目录
+
+- src/、include/：采集、推理、结果变量及可选绘制。
+- config/：源码构建程序的配置和固定门区。
+- scripts/：构建和摄像头初始化脚本。
+- orin/：已编译程序、TensorRT 引擎、配置、当前 GMSL 驱动及启动脚本。
+- third_party/：nlohmann JSON 及其许可证。
+
+NPP 用于 UYVY 转换、YOLO 缩放填充、RTMW 仿射裁剪；CUDA 用于通道转换与归一化。
+风险几何计算仍包含 CPU/OpenCV 操作；算法内部的临时 JSON 对象只在内存中使用。
+
+当前摄像头驱动启动时偶发首帧 ERROR，应用会丢弃并计数；持续错误、采集超时或推理错误会明确报错并退出。
+现有 GMSL 驱动停止采集时可产生 videobuf2 缓冲区告警，v4l2-ctl 同样复现；需在驱动侧完成量产修复和验证。
@@ -0,0 +1,22 @@
+{
+  "camera": {
+    "device": "/dev/video5",
+    "width": 1920,
+    "height": 1536,
+    "fps": 30,
+    "buffers": 4,
+    "timeout_ms": 2000
+  },
+  "inference": {
+    "engine_directory": "../orin/engines",
+    "detector_confidence": 0.3,
+    "max_people": 3,
+    "queue_capacity": 2
+  },
+  "door": {
+    "mask": "door_mask.json",
+    "warning_relative": 0.3,
+    "release_relative": 0.45,
+    "calibration_verified": false
+  }
+}
@@ -0,0 +1,77 @@
+{
+  "version": "x-anylabeling",
+  "flags": {},
+  "shapes": [
+    {
+      "label": "door",
+      "points": [
+        [
+          212.86,
+          145.919
+        ],
+        [
+          378.891,
+          851.857
+        ],
+        [
+          468.293,
+          996.446
+        ],
+        [
+          536.408,
+          1294.131
+        ],
+        [
+          553.437,
+          1404.7
+        ],
+        [
+          561.951,
+          1532.279
+        ],
+        [
+          51.086,
+          1532.279
+        ],
+        [
+          8.514,
+          1515.268
+        ],
+        [
+          0.0,
+          1468.489
+        ],
+        [
+          8.514,
+          26.845
+        ],
+        [
+          25.543,
+          5.582
+        ],
+        [
+          131.973,
+          -2.924
+        ],
+        [
+          191.574,
+          1.329
+        ],
+        [
+          208.603,
+          145.919
+        ]
+      ],
+      "group_id": null,
+      "shape_type": "polygon",
+      "flags": {}
+    }
+  ],
+  "imagePath": "constant_door_mask_reference.jpg",
+  "imageData": null,
+  "imageHeight": 1536,
+  "imageWidth": 1920,
+  "sourceVideo": "D:\\huang.xuewei\\CMS\\video\\20260921104502\\segment_009.mp4",
+  "sourceFrameIndex": 244,
+  "constantDoorMask": true
+}
@@ -0,0 +1,36 @@
+#pragma once
+#include <condition_variable>
+#include <deque>
+#include <mutex>
+#include <optional>
+#include <stdexcept>
+
+namespace cms {
+// FIFO with bounded ownership. Normal close drains; cancellation may discard.
+template<class T> class BoundedQueue {
+public:
+    explicit BoundedQueue(std::size_t capacity):capacity_(capacity) {
+        if(!capacity)throw std::invalid_argument("Queue capacity must be positive");
+    }
+    bool push(T item) {
+        std::unique_lock<std::mutex> lock(mutex_);
+        changed_.wait(lock,[&]{return closed_||items_.size()<capacity_;});
+        if(closed_)return false;
+        items_.push_back(std::move(item));changed_.notify_all();return true;
+    }
+    std::optional<T> pop() {
+        std::unique_lock<std::mutex> lock(mutex_);
+        changed_.wait(lock,[&]{return closed_||!items_.empty();});
+        if(items_.empty())return {};
+        T item=std::move(items_.front());items_.pop_front();changed_.notify_all();return item;
+    }
+    void close(bool discard=false) {
+        std::lock_guard<std::mutex> lock(mutex_);
+        closed_=true;if(discard)items_.clear();changed_.notify_all();
+    }
+private:
+    const std::size_t capacity_;
+    std::mutex mutex_;std::condition_variable changed_;
+    std::deque<T> items_;bool closed_{};
+};
+}
@@ -0,0 +1,31 @@
+#pragma once
+#include "image.hpp"
+#include <atomic>
+#include <memory>
+#include <optional>
+#include <string>
+
+namespace cms {
+struct CameraConfig {
+    std::string device{"/dev/video5"};
+    int width{1920},height{1536},fps{30},buffers{4},timeout_ms{2000};
+};
+struct CapturedFrame {
+    CaptureMetadata metadata;
+    std::shared_ptr<const BgrImage> image;
+};
+// The capture thread exclusively owns read(); requestStop() is thread safe.
+class Camera {
+public:
+    explicit Camera(CameraConfig config);
+    ~Camera();
+    Camera(Camera const&)=delete;
+    Camera& operator=(Camera const&)=delete;
+    std::optional<CapturedFrame> read();
+    void requestStop() noexcept;
+private:
+    struct Impl;
+    std::unique_ptr<Impl> impl_;
+};
+std::uint64_t monotonicNs();
+}
@@ -0,0 +1,43 @@
+#pragma once
+#include "camera.hpp"
+#include "frame_result.hpp"
+#include <filesystem>
+#include <memory>
+
+namespace cms {
+struct PipelineConfig {
+    CameraConfig camera;
+    std::filesystem::path engines,door_mask;
+    double detector_confidence{0.30};
+    int max_people{3};
+    double warning_relative{0.30},release_relative{0.45};
+    std::size_t queue_capacity{2};
+    bool door_mask_calibration_verified{false};
+};
+struct PipelineStats {
+    std::uint64_t captured{},published{},driver_sequence_gaps{},discarded_error_frames{};
+    double mean_capture_to_result_ms{};
+};
+// One lifecycle per instance. start() warms engines before starting capture.
+// Runs continuously until requestStop() or a worker failure.
+// latest() reads the immutable result variable without removing or acknowledging it.
+// No result queue, consumer registration or frame-count termination.
+// Stop is safe from any thread;
+// destruction closes queues, interrupts capture and joins every worker.
+class CameraPipeline {
+public:
+    explicit CameraPipeline(PipelineConfig config);
+    ~CameraPipeline();
+    CameraPipeline(CameraPipeline const&)=delete;
+    CameraPipeline& operator=(CameraPipeline const&)=delete;
+    void start();
+    FrameResultPtr latest() const;
+    bool running() const noexcept;
+    void requestStop() noexcept;
+    void join();
+    PipelineStats stats() const;
+private:
+    struct Impl;
+    std::unique_ptr<Impl> impl_;
+};
+}
@@ -0,0 +1,122 @@
+#pragma once
+
+#include <NvInfer.h>
+#include <NvInferVersion.h>
+#include <cuda_runtime.h>
+
+#include <cstring>
+#include <cstdint>
+#include <filesystem>
+#include <fstream>
+#include <iostream>
+#include <memory>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+namespace fangjia {
+
+inline void checkCuda(cudaError_t status, char const* expression, char const* file, int line) {
+    if (status != cudaSuccess) {
+        throw std::runtime_error(
+            std::string("CUDA failure: ") + expression + " at " + file + ":" +
+            std::to_string(line) + ": " + cudaGetErrorString(status));
+    }
+}
+
+#define CUDA_CHECK(expr) ::fangjia::checkCuda((expr), #expr, __FILE__, __LINE__)
+
+class TrtLogger final : public nvinfer1::ILogger {
+public:
+    explicit TrtLogger(Severity minimum = Severity::kWARNING) : minimum_(minimum) {}
+
+    void log(Severity severity, char const* message) noexcept override {
+        if (severity <= minimum_) {
+            std::cerr << "[TensorRT] " << message << '\n';
+        }
+    }
+
+private:
+    Severity minimum_;
+};
+
+template <typename T>
+struct TrtDeleter {
+    void operator()(T* object) const noexcept { delete object; }
+};
+
+template <typename T>
+using TrtUniquePtr = std::unique_ptr<T, TrtDeleter<T>>;
+
+inline std::vector<std::uint8_t> readBinary(std::filesystem::path const& path) {
+    std::ifstream stream(path, std::ios::binary | std::ios::ate);
+    if (!stream) {
+        throw std::runtime_error("Cannot open " + path.string());
+    }
+    auto const size = stream.tellg();
+    if(size<=0)throw std::runtime_error("Empty or unreadable engine: "+path.string());
+    stream.seekg(0);
+    std::vector<std::uint8_t> bytes(static_cast<std::size_t>(size));
+    if(!stream.read(reinterpret_cast<char*>(bytes.data()), size))throw std::runtime_error("Truncated engine: "+path.string());
+    return bytes;
+}
+
+inline void writeBinary(std::filesystem::path const& path, void const* data, std::size_t size) {
+    std::filesystem::create_directories(path.parent_path());
+    std::ofstream stream(path, std::ios::binary);
+    if (!stream) {
+        throw std::runtime_error("Cannot write " + path.string());
+    }
+    stream.write(static_cast<char const*>(data), static_cast<std::streamsize>(size));
+}
+
+inline std::size_t volume(nvinfer1::Dims const& dims) {
+    std::size_t value = 1;
+    for (int index = 0; index < dims.nbDims; ++index) {
+        if (dims.d[index] < 0) {
+            throw std::runtime_error("Unresolved dynamic tensor dimension");
+        }
+        value *= static_cast<std::size_t>(dims.d[index]);
+    }
+    return value;
+}
+
+inline std::size_t dataTypeSize(nvinfer1::DataType type) {
+    switch (type) {
+    case nvinfer1::DataType::kFLOAT: return 4;
+    case nvinfer1::DataType::kHALF: return 2;
+    case nvinfer1::DataType::kINT8: return 1;
+    case nvinfer1::DataType::kINT32: return 4;
+#if NV_TENSORRT_MAJOR >= 10
+    case nvinfer1::DataType::kINT64: return 8;
+#endif
+    case nvinfer1::DataType::kBOOL: return 1;
+    default: throw std::runtime_error("Unsupported TensorRT data type");
+    }
+}
+
+class CudaEventTimer {
+public:
+    CudaEventTimer() {
+        CUDA_CHECK(cudaEventCreate(&start_));
+        CUDA_CHECK(cudaEventCreate(&end_));
+    }
+    ~CudaEventTimer() {
+        cudaEventDestroy(start_);
+        cudaEventDestroy(end_);
+    }
+    void start(cudaStream_t stream) { CUDA_CHECK(cudaEventRecord(start_, stream)); }
+    float stop(cudaStream_t stream) {
+        CUDA_CHECK(cudaEventRecord(end_, stream));
+        CUDA_CHECK(cudaEventSynchronize(end_));
+        float milliseconds = 0.0F;
+        CUDA_CHECK(cudaEventElapsedTime(&milliseconds, start_, end_));
+        return milliseconds;
+    }
+
+private:
+    cudaEvent_t start_{};
+    cudaEvent_t end_{};
+};
+
+} // namespace fangjia
@@ -0,0 +1,5 @@
+#pragma once
+#include "camera_pipeline.hpp"
+namespace cms {
+PipelineConfig loadConfig(std::filesystem::path const& path);
+}
@@ -0,0 +1,18 @@
+#pragma once
+#include "pose_engine.hpp"
+#include "gpu_preprocess.hpp"
+namespace cms {
+class Detector {
+public:
+    Detector(std::filesystem::path const& path, fangjia::TrtLogger& logger, float confidence, int maxPeople);
+    std::vector<fangjia::Box> infer(cv::Mat const& frame,cudaStream_t stream,int filterHeight=0);
+private:
+    fangjia::TrtEngine engine_;
+    float confidence_;
+    int maxPeople_, width_, height_;
+    std::string input_,output_;
+    std::unique_ptr<YoloPreprocess> preprocess_;
+    std::vector<float> outputHost_;
+};
+cv::Mat maskOtherPeople(cv::Mat const& frame,std::vector<fangjia::Box> const& boxes,size_t target);
+}
@@ -0,0 +1,35 @@
+#pragma once
+#include <nlohmann/json.hpp>
+#include <opencv2/core.hpp>
+#include <filesystem>
+#include <map>
+#include <set>
+using Json = nlohmann::ordered_json;
+namespace cms {
+struct DoorOptions {
+    std::filesystem::path constant, directory;
+    std::string format{"auto"};
+    std::set<std::string> labels{"door", "车门"};
+    int offset{};
+    bool require{};
+    double warning{0.30}, release{0.45};
+};
+class DoorRisk {
+public:
+    DoorRisk(DoorOptions options, cv::Size size);
+    void initialize() { load(0); }
+    Json evaluate(int frame, Json& record);
+    cv::Mat const& mask() const { return mask_; }
+private:
+    void load(int frame);
+    cv::Mat rasterize(std::filesystem::path const& path) const;
+    DoorOptions options_;
+    cv::Size size_;
+    std::string format_, status_;
+    std::filesystem::path loadedPath_;
+    std::map<int, std::filesystem::path> paths_;
+    cv::Mat mask_, distance_;
+    bool cached_{};
+};
+double rounded(double value, int digits);
+}
@@ -0,0 +1,80 @@
+#pragma once
+// Public downstream contract. Owns its data; no GPU, JSON or network handles.
+#include "image.hpp"
+#include <array>
+#include <cstdint>
+#include <memory>
+#include <optional>
+#include <string>
+#include <vector>
+
+namespace cms {
+struct DoorMaskResult {
+    std::string status;
+    int annotation_index{};
+    std::optional<std::string> path;
+    std::string format;
+    int area_px{};
+};
+struct RegionRisk {
+    std::string status;
+    std::optional<double> risk_score;
+    std::string risk_level;
+    std::string decision;
+    std::string source;
+    int track_id{};
+    std::string side;
+    std::optional<std::string> state;
+    // Unavailable evidence has no geometric measurements, matching the reference.
+    std::optional<int> region_area_px, door_overlap_px;
+    std::optional<double> region_overlap_ratio, min_distance_px;
+    std::optional<double> region_scale_px, distance_relative;
+};
+struct RiskResult {
+    std::string status;
+    std::optional<double> risk_score;
+    std::string risk_level;
+    std::string decision;
+    std::optional<std::string> source;
+    std::optional<int> track_id;
+    std::optional<std::string> side;
+    DoorMaskResult door_mask;
+    double warning_distance_relative{}, release_distance_relative{};
+    std::vector<RegionRisk> regions;
+};
+struct PosePoint { double x{},y{},score{}; };
+struct HandState {
+    std::string state, source;
+    bool limb_consensus{}, distance_eligible{}, alarm_eligible{};
+};
+struct PersonResult {
+    int track_id{};
+    double confidence{};
+    std::array<double,4> box{};
+    std::array<PosePoint,133> selected{}, smoothed{}, temporal{};
+    std::array<HandState,2> hands; // left, right
+};
+struct HandRegion {
+    int person_id{},track_id{};
+    std::string side;
+    std::array<PosePoint,21> keypoints{};
+};
+struct StageTiming {
+    double detector_ms{}, pose_ms{}, fusion_risk_ms{}, capture_to_result_ms{};
+};
+struct FrameResult {
+    int frame_index{};
+    std::vector<int> track_ids;
+    std::vector<std::array<double,4>> boxes;
+    RiskResult risk;
+    CaptureMetadata capture;
+    StageTiming timing;
+    std::vector<PersonResult> persons;
+    std::vector<HandRegion> hand_regions;
+    // Both images belong to this result. The constant mask is shared across frames.
+    std::shared_ptr<const BgrImage> image;
+    std::shared_ptr<const BinaryMask> door_mask_pixels;
+    bool door_mask_calibration_verified{};
+};
+using FrameResultPtr = std::shared_ptr<const FrameResult>;
+}
@@ -0,0 +1,19 @@
+#pragma once
+#include "common.hpp"
+#include <npp.h>
+#include <opencv2/core.hpp>
+namespace cms {
+inline void nppCheck(NppStatus status) {if(status<0)throw std::runtime_error("NPP error "+std::to_string(int(status)));}
+inline NppStreamContext nppContext(cudaStream_t stream) {
+    NppStreamContext c{};c.hStream=stream;CUDA_CHECK(cudaGetDevice(&c.nCudaDeviceId));cudaDeviceProp p{};CUDA_CHECK(cudaGetDeviceProperties(&p,c.nCudaDeviceId));
+    c.nMultiProcessorCount=p.multiProcessorCount;c.nMaxThreadsPerMultiProcessor=p.maxThreadsPerMultiProcessor;c.nMaxThreadsPerBlock=p.maxThreadsPerBlock;c.nSharedMemPerBlock=p.sharedMemPerBlock;c.nCudaDevAttrComputeCapabilityMajor=p.major;c.nCudaDevAttrComputeCapabilityMinor=p.minor;CUDA_CHECK(cudaStreamGetFlags(stream,&c.nStreamFlags));return c;
+}
+class YoloPreprocess {
+public:
+    YoloPreprocess(int width,int height):width_(width),height_(height) {CUDA_CHECK(cudaMalloc(&letter_,size_t(width)*height*3));}
+    ~YoloPreprocess(){cudaFree(source_);cudaFree(letter_);}
+    void enqueue(cv::Mat const& frame,float* destination,cudaStream_t stream);
+private:
+    int width_,height_;size_t capacity_{};unsigned char *source_{},*letter_{};
+};
+}
@@ -0,0 +1,25 @@
+#pragma once
+#include <cstdint>
+#include <vector>
+
+namespace cms {
+// Immutable when published. Stride is explicit; no borrowed camera/GPU memory.
+struct BgrImage {
+    int width{}, height{}, stride{};
+    std::vector<std::uint8_t> pixels;
+};
+struct BinaryMask {
+    int width{}, height{}, stride{};
+    std::vector<std::uint8_t> pixels;
+};
+struct CaptureMetadata {
+    std::uint64_t frame_index{};
+    std::uint32_t driver_sequence{};
+    std::uint64_t driver_timestamp_ns{};
+    std::uint64_t received_monotonic_ns{};
+    std::uint64_t capture_monotonic_ns{};
+    bool driver_timestamp_monotonic{};
+    std::uint32_t skipped_driver_frames{};
+    std::uint32_t discarded_error_frames{};
+};
+}
@@ -0,0 +1,20 @@
+#pragma once
+#include "frame_result.hpp"
+#include <atomic>
+#include <stdexcept>
+
+namespace cms {
+// One immutable result variable. Readers neither remove nor acknowledge results.
+class LatestResult {
+public:
+    void publish(FrameResultPtr result) {
+        if(!result)throw std::invalid_argument("Cannot publish a null result");
+        std::atomic_store_explicit(&latest_,std::move(result),std::memory_order_release);
+    }
+    FrameResultPtr load() const {
+        return std::atomic_load_explicit(&latest_,std::memory_order_acquire);
+    }
+private:
+    FrameResultPtr latest_;
+};
+}
@@ -0,0 +1,100 @@
+#pragma once
+
+#include <array>
+#include <cstdint>
+#include <string>
+#include <utility>
+#include <vector>
+
+namespace fangjia {
+
+struct Box {
+    float x1{};
+    float y1{};
+    float x2{};
+    float y2{};
+    float score{};
+    int trackId{};
+};
+
+struct Keypoint {
+    float x{};
+    float y{};
+    float score{};
+};
+
+struct PersonPose {
+    Box box;
+    std::array<Keypoint, 133> keypoints{};
+};
+
+enum class LimbState {
+    Observed,
+    Provisional,
+    Recovered,
+    Reassigned,
+    Interpolated,
+    PredictedShort,
+    ReachableMedium,
+    ReachableLong,
+};
+
+inline char const* limbStateName(LimbState state) {
+    switch (state) {
+    case LimbState::Observed: return "observed";
+    case LimbState::Provisional: return "provisional";
+    case LimbState::Recovered: return "recovered";
+    case LimbState::Reassigned: return "reassigned";
+    case LimbState::Interpolated: return "interpolated";
+    case LimbState::PredictedShort: return "predicted_short";
+    case LimbState::ReachableMedium: return "reachable_medium";
+    case LimbState::ReachableLong: return "reachable_long";
+    }
+    return "unknown";
+}
+
+struct LimbObservation {
+    int trackId{};
+    bool left{};
+    LimbState state{LimbState::ReachableLong};
+    Keypoint elbow{};
+    Keypoint wrist{};
+    std::array<Keypoint, 21> hand{};
+    Box personBox{};
+    bool handValid{};
+    bool alarmEligible{};
+    float handQuality{};
+};
+
+// Exact pre-V2 hand evidence consumed by the Python risk stage.  These points
+// come from raw/isolated fusion before temporal reconstruction.
+struct GlobalHand {
+    int personId{};
+    int trackId{};
+    bool left{};
+    std::string source;
+    std::string state;
+    bool distanceEligible{};
+    bool alarmEligible{};
+    float score{};
+    float spanPx{};
+    std::array<Keypoint, 21> keypoints{};
+    std::vector<std::pair<int, bool>> aliases;
+};
+
+struct FusionMeta {
+    bool limbConsensus{}, handValid{}, distanceEligible{}, alarmEligible{};
+    double handScore{};
+    std::string handSource, handState;
+};
+
+struct PoseFrameResult {
+    std::vector<std::array<FusionMeta, 2>> metadata;
+    std::vector<PersonPose> selectedPeople;
+    std::vector<PersonPose> smoothedPeople;
+    std::vector<PersonPose> v2People;
+    std::vector<LimbObservation> limbs;
+    std::vector<GlobalHand> globalHands;
+};
+
+} // namespace fangjia
@@ -0,0 +1,108 @@
+#pragma once
+
+#include "pipeline_types.hpp"
+#include "trt_engine.hpp"
+
+#include <opencv2/core.hpp>
+
+#include <deque>
+#include <array>
+#include <optional>
+#include <string>
+#include <unordered_map>
+
+namespace fangjia {
+
+class BoxTracker {
+public:
+    std::vector<int> update(std::vector<Box>& boxes, int frameIndex, int width, int height);
+
+private:
+    struct Track { Box box; int lastFrame{}; };
+    int nextId_{1};
+    std::unordered_map<int, Track> tracks_;
+};
+
+class RtmwEngine {
+public:
+    RtmwEngine(
+        std::filesystem::path const& enginePath,
+        TrtLogger& logger,
+        int maxSourceWidth,
+        int maxSourceHeight,
+        int maxPeople = 3);
+    ~RtmwEngine();
+
+    RtmwEngine(RtmwEngine const&) = delete;
+    RtmwEngine& operator=(RtmwEngine const&) = delete;
+
+    std::vector<PersonPose> infer(
+        cv::Mat const& frame,
+        std::vector<Box> const& boxes,
+        cudaStream_t stream,
+        float* preprocessMs = nullptr,
+        float* inferenceMs = nullptr,
+        float* postprocessMs = nullptr);
+
+    void warmup(cudaStream_t stream);
+
+private:
+    TrtEngine engine_;
+    int maxSourceWidth_{};
+    int maxSourceHeight_{};
+    int maxPeople_{};
+    std::uint8_t* sourceDevice_{};
+    std::uint8_t* cropsDevice_{};
+    float* networkDevice_{};
+    std::vector<double> centersHost_;
+    std::vector<double> scalesHost_;
+    std::vector<float> simccXHost_;
+    std::vector<float> simccYHost_;
+};
+
+class PoseFusionTracker {
+public:
+    PoseFrameResult update(
+        std::vector<PersonPose> const& raw,
+        std::vector<PersonPose> const& isolated,
+        int frameIndex);
+
+private:
+    struct Vec2 { double x{}; double y{}; };
+    struct SideFusion {
+        std::array<Vec2, 3> limbPrevious{};
+        std::array<Vec2, 3> limbVelocity{};
+        std::array<Vec2, 21> handPrevious{};
+        std::array<Vec2, 21> handVelocity{};
+        bool hasLimbPrevious{};
+        bool hasLimbVelocity{};
+        bool hasHandPrevious{};
+        bool hasHandVelocity{};
+        std::string limbSource;
+        std::string handSource;
+    };
+    struct Endpoint {
+        bool hasWrist{};
+        Vec2 wrist{};
+        Vec2 velocity{};
+        std::array<Vec2, 21> handShape{};
+        bool hasHandShape{};
+        int lastObservedFrame{-1};
+        int missedFrames{};
+        bool hasLastElbow{};
+        Vec2 lastElbow{};
+        bool hasLastDirection{};
+        Vec2 lastDirection{};
+    };
+    struct PersonState {
+        std::array<SideFusion, 2> fusion;
+        std::array<Endpoint, 2> endpoints;
+        std::array<std::deque<std::array<double, 2>>, 2> boneSamples;
+        std::array<Vec2, 133> smoothed{};
+        bool hasSmoothed{};
+        int lastFrame{-1};
+    };
+    std::unordered_map<int, PersonState> states_;
+};
+
+} // namespace fangjia
@@ -0,0 +1,38 @@
+#pragma once
+#include "door_risk.hpp"
+#include "pipeline_types.hpp"
+namespace cms {
+inline Json xy(std::vector<fangjia::PersonPose> const& people, int digits=3) {
+    Json a=Json::array(); for(auto const& p:people) { Json k=Json::array();for(auto const& v:p.keypoints) k.push_back({rounded(v.x,digits),rounded(v.y,digits)});a.push_back(k); } return a;
+}
+inline Json scores(std::vector<fangjia::PersonPose> const& people) {
+    Json a=Json::array();for(auto const& p:people) {Json k=Json::array();for(auto const& v:p.keypoints) k.push_back(v.score);a.push_back(k);}return a;
+}
+inline Json makeRecord(int frame,std::vector<fangjia::PersonPose> const& raw,std::vector<fangjia::PersonPose> const& isolated,fangjia::PoseFrameResult const& fused, bool diagnostics=true, bool visual=true) {
+    Json r={{"frame_index",frame},{"track_ids",Json::array()},{"boxes",Json::array()},{"det_confidences",Json::array()}};
+    for(auto const&p:raw) {auto const&b=p.box;r["track_ids"].push_back(b.trackId);r["boxes"].push_back({rounded(b.x1,3),rounded(b.y1,3),rounded(b.x2,3),rounded(b.y2,3)});r["det_confidences"].push_back(rounded(b.score,4));}
+    if(diagnostics) {
+    r["raw_keypoints"]=xy(raw,9);r["raw_scores"]=scores(raw);r["isolated_keypoints"]=xy(isolated,9);r["isolated_scores"]=scores(isolated);
+    r["inference_boxes"]=Json::array();for(auto const&p:raw)r["inference_boxes"].push_back({p.box.x1,p.box.y1,p.box.x2,p.box.y2});
+    }
+    if(visual||diagnostics) {
+    r["temporal_selected_keypoints"]=xy(fused.selectedPeople);r["temporal_keypoints"]=xy(fused.smoothedPeople);r["v2_keypoints"]=xy(fused.v2People);
+    r["temporal_metadata"]=Json::array();r["v2_metadata"]=Json::array();
+    for(size_t p=0;p<fused.metadata.size();++p) {
+        Json m=Json::object(),v=Json::object();
+        for(int s=0;s<2;++s) {auto side=s==0?"left":"right";auto const&t=fused.metadata[p][s];
+            m[side]={{"limb_state",t.limbConsensus?"consensus":"inferred"},{"hand_source",t.handSource},{"hand_state",t.handState},{"distance_eligible",t.distanceEligible},{"alarm_eligible",t.alarmEligible},{"selected_hand_valid",t.handValid},{"selected_hand_score",t.handScore}};
+            v[side]={{"state",fangjia::limbStateName(fused.limbs[p*2+s].state)}};
+        } r["temporal_metadata"].push_back(m);r["v2_metadata"].push_back(v);
+    }
+    }
+    r["global_hands"]=Json::array();
+    for(auto const& h:fused.globalHands) {
+        Json points=Json::array(),aliases=Json::array();double cx=0,cy=0;
+        for(auto const& k:h.keypoints) {points.push_back({rounded(k.x,3),rounded(k.y,3)});cx+=k.x;cy+=k.y;}
+        for(auto const&a:h.aliases) aliases.push_back({{"person_id",a.first},{"side",a.second?"left":"right"}});
+        r["global_hands"].push_back({{"person_id",h.personId},{"side",h.left?"left":"right"},{"source",h.source},{"state",h.state},{"distance_eligible",h.distanceEligible},{"alarm_eligible",h.alarmEligible},{"score",rounded(h.score,4)},{"center",{rounded(cx/21,3),rounded(cy/21,3)}},{"span_px",rounded(h.spanPx,3)},{"keypoints",points},{"aliases",aliases}});
+    }
+    return r;
+}
+}
@@ -0,0 +1,34 @@
+#pragma once
+#include "frame_result.hpp"
+#include "door_risk.hpp"
+namespace cms {
+template<class T> inline std::optional<T> optionalField(Json const& value,char const* key) {
+    auto it=value.find(key);return it==value.end()||it->is_null()?std::nullopt:std::optional<T>(it->get<T>());
+}
+inline FrameResult makeFrameResult(Json const& record,Json const& risk,
+                                  CaptureMetadata capture) {
+    FrameResult out;
+    out.frame_index=record.at("frame_index").get<int>();
+    out.track_ids=record.at("track_ids").get<std::vector<int>>();
+    out.boxes=record.at("boxes").get<std::vector<std::array<double,4>>>();
+    out.capture=capture;
+    auto& r=out.risk;
+    r.status=risk.at("status");r.risk_score=optionalField<double>(risk,"risk_score");
+    r.risk_level=risk.at("risk_level");r.decision=risk.at("decision");
+    r.source=optionalField<std::string>(risk,"source");r.track_id=optionalField<int>(risk,"track_id");r.side=optionalField<std::string>(risk,"side");
+    auto const& d=risk.at("door_mask");
+    r.door_mask={d.at("status"),d.at("annotation_index"),optionalField<std::string>(d,"path"),d.at("format"),d.at("area_px")};
+    r.warning_distance_relative=risk.at("warning_distance_relative");r.release_distance_relative=risk.at("release_distance_relative");
+    for(auto const& j:risk.at("regions")) {
+        RegionRisk v;
+        v.status=j.at("status");v.risk_score=optionalField<double>(j,"risk_score");
+        v.risk_level=j.at("risk_level");v.decision=j.at("decision");v.source=j.at("source");
+        v.track_id=j.at("track_id");v.side=j.at("side");v.state=optionalField<std::string>(j,"state");
+        v.region_area_px=optionalField<int>(j,"region_area_px");v.door_overlap_px=optionalField<int>(j,"door_overlap_px");
+        v.region_overlap_ratio=optionalField<double>(j,"region_overlap_ratio");v.min_distance_px=optionalField<double>(j,"min_distance_px");
+        v.region_scale_px=optionalField<double>(j,"region_scale_px");v.distance_relative=optionalField<double>(j,"distance_relative");
+        r.regions.push_back(std::move(v));
+    }
+    return out;
+}
+}
@@ -0,0 +1,36 @@
+#pragma once
+#include "camera_pipeline.hpp"
+#include <atomic>
+#include <csignal>
+#include <pthread.h>
+#include <stdexcept>
+#include <thread>
+
+namespace cms {
+// Construct before starting workers so all threads inherit the signal mask.
+// Stop requests run in an ordinary thread, never an asynchronous signal handler.
+class StopSignals {
+public:
+    explicit StopSignals(CameraPipeline& pipeline):pipeline_(pipeline) {
+        sigemptyset(&signals_);sigaddset(&signals_,SIGINT);sigaddset(&signals_,SIGTERM);
+        if(pthread_sigmask(SIG_BLOCK,&signals_,&old_)!=0)throw std::runtime_error("Cannot block termination signals");
+        try {
+            thread_=std::thread([this]{
+                while(!done_) {
+                    timespec timeout{0,100000000};
+                    int signal=sigtimedwait(&signals_,nullptr,&timeout);
+                    if(signal==SIGINT||signal==SIGTERM){pipeline_.requestStop();break;}
+                }
+            });
+        }catch(...){pthread_sigmask(SIG_SETMASK,&old_,nullptr);throw;}
+    }
+    ~StopSignals(){done_=true;if(thread_.joinable())thread_.join();pthread_sigmask(SIG_SETMASK,&old_,nullptr);}
+    StopSignals(StopSignals const&)=delete;
+    StopSignals& operator=(StopSignals const&)=delete;
+private:
+    CameraPipeline& pipeline_;
+    sigset_t signals_{},old_{};
+    std::atomic<bool> done_{false};
+    std::thread thread_;
+};
+}
@@ -0,0 +1,49 @@
+#pragma once
+
+#include "common.hpp"
+
+#include <unordered_map>
+
+namespace fangjia {
+
+class TrtEngine {
+public:
+    TrtEngine(std::filesystem::path const& enginePath, TrtLogger& logger);
+    ~TrtEngine();
+
+    TrtEngine(TrtEngine const&) = delete;
+    TrtEngine& operator=(TrtEngine const&) = delete;
+
+    void setInputShape(std::string const& name, nvinfer1::Dims const& shape);
+    void prepare();
+    void bindExternal(std::string const& name, void* pointer, std::size_t bytes);
+    void useOwned(std::string const& name);
+    bool enqueue(cudaStream_t stream);
+
+    void* device(std::string const& name);
+    std::size_t bytes(std::string const& name) const;
+    nvinfer1::Dims shape(std::string const& name) const;
+    nvinfer1::DataType dataType(std::string const& name) const;
+    std::vector<std::string> inputs() const;
+    std::vector<std::string> outputs() const;
+
+private:
+    struct Buffer {
+        void* owned{};
+        void* active{};
+        std::size_t capacity{};
+        std::size_t required{};
+        bool external{};
+    };
+
+    Buffer& buffer(std::string const& name);
+    Buffer const& buffer(std::string const& name) const;
+
+    TrtUniquePtr<nvinfer1::IRuntime> runtime_;
+    TrtUniquePtr<nvinfer1::ICudaEngine> engine_;
+    TrtUniquePtr<nvinfer1::IExecutionContext> context_;
+    std::filesystem::path enginePath_;
+    std::unordered_map<std::string, Buffer> buffers_;
+};
+
+} // namespace fangjia
@@ -0,0 +1,18 @@
+#pragma once
+#include "image.hpp"
+#include <cstddef>
+#include <memory>
+namespace cms {
+// GPU NPP BT.601 UYVY -> RGB followed by a CUDA BGR channel swap.
+class UyvyConverter {
+public:
+    UyvyConverter(int width,int height);
+    ~UyvyConverter();
+    UyvyConverter(UyvyConverter const&)=delete;
+    UyvyConverter& operator=(UyvyConverter const&)=delete;
+    std::shared_ptr<const BgrImage> convert(void const* data,std::size_t stride);
+private:
+    struct Impl;
+    std::unique_ptr<Impl> impl_;
+};
+}
@@ -0,0 +1,11 @@
+#pragma once
+#include "frame_result.hpp"
+#include <opencv2/core.hpp>
+
+namespace cms {
+class CameraPipeline;
+// Pure drawing from the immutable result. Never re-runs inference.
+cv::Mat renderResult(FrameResult const& result);
+// Optional observer: closing the window returns without stopping the pipeline.
+void runLocalVisualization(CameraPipeline const& pipeline);
+}
@@ -0,0 +1,22 @@
+# Orin 可直接运行目录
+
+本目录包含实际在 Jetson AGX Orin 上编译的 aarch64 程序和 TensorRT FP16 引擎。
+目标环境：L4T R35.4.1 / JetPack 5.1.2，CUDA 11.4，TensorRT 8.5.2，OpenCV 4.5.4。
+使用系统安装的 CUDA、TensorRT、OpenCV 共享库；不适用于任意 JetPack 版本或 x86 主机。
+
+```bash
+cd orin
+sha256sum -c SHA256SUMS
+./scripts/initialize_camera.sh
+./run.sh
+```
+
+可视化可选：`./run.sh --visualize`，需要板端本地桌面 DISPLAY。
+默认持续推理直到 Ctrl+C / SIGTERM；正常退出没有统计输出。
+预览窗口关闭后推理继续。结果通过 C++ CameraPipeline::latest() 获取，不通过端口或结果文件传递。
+
+程序为 build/camera_risk，配置为 config/camera.json，引擎为 engines/*.engine。
+./run.sh 从所在目录解析路径，也可以在其他工作目录通过绝对路径调用。
+
+bin/drivers 中 GMSL 模块对应内核 5.10.120-tegra。初始化脚本检查内核版本；加载模块需要 root 或现有 sudo 权限。
+门区保留当前室内联调设置，door.calibration_verified=false。
@@ -0,0 +1,11 @@
+926a5165a4ebffe9568cf694018afb097cfb69b6e696a19939e5d364a9e3dfa5  README.md
+6d5e874fd27fcfd75d303d7e35d9672028af71041f52dcc82a1e83676faefa65  bin/drivers/max9295.ko
+8a78d04ef8be44e9b62a5e4ee315d3d6d679e34aedff02a234bbf4d40e222933  bin/drivers/max9296.ko
+3da628587477cc35d5e3f8aee86e718e2ddf16808e893e00ba6e164b2307e843  bin/drivers/sgx-yuv-gmsl2.ko
+b5102fe0a529349bb8e9700ae1d8e8ddeaa7c4279cb28e7826c5cc32dce2bb1d  build/camera_risk
+f2334c43e6a48480260c770d46a3ab0edb855f6f52bc8b91d1559e69edfad043  config/camera.json
+9ba768d559fd3b64e833b50c62c70f2adfd6c169620aea5fba32cfb1b4c5b0e9  config/door_mask.json
+ebcd09e8c2e35008adbb0c1d222e580b17e0df1c8093df3481ba3374cd9a8174  engines/rtmw_l_384.engine
+f6a65b926ac1144484f4aa93cd26d7997d32d62310ab93fe867f2bd4412222e9  engines/yolo11s.engine
+7f42b0e7f9396e0d82313a892fc62c153261dab1206a206a466bce6376f58a36  run.sh
+54e1d63e3aa72ad9b056a224d55968d4947ca6703a8d1c405570de5707a7e7ef  scripts/initialize_camera.sh
@@ -0,0 +1,22 @@
+{
+  "camera": {
+    "device": "/dev/video5",
+    "width": 1920,
+    "height": 1536,
+    "fps": 30,
+    "buffers": 4,
+    "timeout_ms": 2000
+  },
+  "inference": {
+    "engine_directory": "../engines",
+    "detector_confidence": 0.30,
+    "max_people": 3,
+    "queue_capacity": 2
+  },
+  "door": {
+    "mask": "door_mask.json",
+    "warning_relative": 0.30,
+    "release_relative": 0.45,
+    "calibration_verified": false
+  }
+}
@@ -0,0 +1,77 @@
+{
+  "version": "x-anylabeling",
+  "flags": {},
+  "shapes": [
+    {
+      "label": "door",
+      "points": [
+        [
+          212.86,
+          145.919
+        ],
+        [
+          378.891,
+          851.857
+        ],
+        [
+          468.293,
+          996.446
+        ],
+        [
+          536.408,
+          1294.131
+        ],
+        [
+          553.437,
+          1404.7
+        ],
+        [
+          561.951,
+          1532.279
+        ],
+        [
+          51.086,
+          1532.279
+        ],
+        [
+          8.514,
+          1515.268
+        ],
+        [
+          0.0,
+          1468.489
+        ],
+        [
+          8.514,
+          26.845
+        ],
+        [
+          25.543,
+          5.582
+        ],
+        [
+          131.973,
+          -2.924
+        ],
+        [
+          191.574,
+          1.329
+        ],
+        [
+          208.603,
+          145.919
+        ]
+      ],
+      "group_id": null,
+      "shape_type": "polygon",
+      "flags": {}
+    }
+  ],
+  "imagePath": "constant_door_mask_reference.jpg",
+  "imageData": null,
+  "imageHeight": 1536,
+  "imageWidth": 1920,
+  "sourceVideo": "D:\\huang.xuewei\\CMS\\video\\20260921104502\\segment_009.mp4",
+  "sourceFrameIndex": 244,
+  "constantDoorMask": true
+}
@@ -0,0 +1,4 @@
+#!/usr/bin/env bash
+set -euo pipefail
+root="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
+exec "$root/build/camera_risk" --config "$root/config/camera.json" "$@"
@@ -0,0 +1,21 @@
+#!/usr/bin/env bash
+set -euo pipefail
+root="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")/.." && pwd)"
+device=/dev/video5
+loaded() { test -d "/sys/module/$1"; }
+root_command() { if [[ $EUID -eq 0 ]]; then "$@"; else sudo -n "$@"; fi; }
+if [[ -c "$device" ]] && loaded max9295 && loaded max9296 && loaded sgx_yuv_gmsl2; then
+  echo "Camera already initialized: $device"
+  exit 0
+fi
+for driver in max9295 max9296 sgx-yuv-gmsl2; do
+  file="$root/bin/drivers/$driver.ko"
+  [[ -f "$file" ]] || { echo "Missing driver: $file" >&2; exit 1; }
+  [[ "$(modinfo -F vermagic "$file")" == "$(uname -r) "* ]] || { echo "Kernel mismatch: $file" >&2; exit 1; }
+done
+loaded max9295 || root_command insmod "$root/bin/drivers/max9295.ko"
+loaded max9296 || root_command insmod "$root/bin/drivers/max9296.ko"
+loaded sgx_yuv_gmsl2 || root_command insmod "$root/bin/drivers/sgx-yuv-gmsl2.ko" gmsl_mode=2,2,2,2
+for _ in {1..20}; do [[ -c "$device" ]] && break; sleep 0.25; done
+[[ -c "$device" ]] || { echo "$device was not created; verify camera/DTB" >&2; exit 1; }
+v4l2-ctl -d "$device" -c sensor_mode=1,trig_pin=0xffff0007
@@ -0,0 +1,4 @@
+#!/usr/bin/env bash
+set -euo pipefail
+root="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd)"
+exec "$root/build/camera_risk" --config "$root/config/camera.json" "$@"
@@ -0,0 +1,7 @@
+#!/usr/bin/env bash
+set -euo pipefail
+root="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")/.." && pwd)"
+cmake_bin="${CMAKE_COMMAND:-cmake}"
+if [[ -x "$HOME/miniconda3/envs/yolo/bin/cmake" ]]; then cmake_bin="${CMAKE_COMMAND:-$HOME/miniconda3/envs/yolo/bin/cmake}"; fi
+"$cmake_bin" -S "$root" -B "$root/build" -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_COMPILER="${CUDACXX:-/usr/local/cuda/bin/nvcc}" "$@"
+"$cmake_bin" --build "$root/build" -j4
@@ -0,0 +1,4 @@
+#!/usr/bin/env bash
+set -euo pipefail
+root="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")/.." && pwd)"
+exec "$root/orin/scripts/initialize_camera.sh" "$@"
@@ -0,0 +1,146 @@
+#include "camera.hpp"
+#include "uyvy_converter.hpp"
+#include <linux/videodev2.h>
+#include <sys/eventfd.h>
+#include <sys/ioctl.h>
+#include <sys/mman.h>
+#include <poll.h>
+#include <fcntl.h>
+#include <unistd.h>
+#include <time.h>
+#include <cerrno>
+#include <cstring>
+#include <limits>
+#include <stdexcept>
+
+namespace cms {
+std::uint64_t monotonicNs() {
+    timespec time{};
+    if(clock_gettime(CLOCK_MONOTONIC,&time)!=0)throw std::runtime_error("clock_gettime failed");
+    return std::uint64_t(time.tv_sec)*1000000000ULL+time.tv_nsec;
+}
+namespace {
+void checkedIoctl(int fd,unsigned long request,void* argument,char const* operation) {
+    int status;
+    do {status=ioctl(fd,request,argument);} while(status<0&&errno==EINTR);
+    if(status<0)throw std::runtime_error(std::string(operation)+": "+std::strerror(errno));
+}
+}
+struct Camera::Impl {
+    struct Buffer {void* data{};std::size_t size{};};
+    CameraConfig config;
+    int fd{-1},stopFd{-1};bool streaming{};
+    std::uint32_t stride{},lastSequence{};bool haveSequence{};
+    std::uint64_t index{};
+    std::uint32_t consecutiveErrors{};
+    std::atomic<bool> stopped{false};
+    std::vector<Buffer> buffers;
+    std::unique_ptr<UyvyConverter> converter;
+    explicit Impl(CameraConfig c):config(std::move(c)) {
+        if(config.width<=0||config.width>8192||config.height<=0||config.height>8192||config.width%2||config.fps<=0||config.fps>240||config.buffers<2||config.buffers>32||config.timeout_ms<1)
+            throw std::invalid_argument("Invalid camera configuration");
+        try {
+            stopFd=eventfd(0,EFD_CLOEXEC|EFD_NONBLOCK);
+            if(stopFd<0)throw std::runtime_error("Cannot create camera stop event");
+            fd=open(config.device.c_str(),O_RDWR|O_NONBLOCK|O_CLOEXEC);
+            if(fd<0)throw std::runtime_error("Cannot open "+config.device+": "+std::strerror(errno));
+            v4l2_capability capability{};checkedIoctl(fd,VIDIOC_QUERYCAP,&capability,"VIDIOC_QUERYCAP");
+            auto caps=(capability.capabilities&V4L2_CAP_DEVICE_CAPS)?capability.device_caps:capability.capabilities;
+            if(!(caps&V4L2_CAP_VIDEO_CAPTURE)||!(caps&V4L2_CAP_STREAMING))throw std::runtime_error("Camera does not support single-plane streaming capture");
+            v4l2_format format{};format.type=V4L2_BUF_TYPE_VIDEO_CAPTURE;
+            format.fmt.pix.width=config.width;format.fmt.pix.height=config.height;
+            format.fmt.pix.pixelformat=V4L2_PIX_FMT_UYVY;format.fmt.pix.field=V4L2_FIELD_NONE;
+            checkedIoctl(fd,VIDIOC_S_FMT,&format,"VIDIOC_S_FMT");
+            if(format.fmt.pix.pixelformat!=V4L2_PIX_FMT_UYVY||format.fmt.pix.width!=static_cast<unsigned>(config.width)||format.fmt.pix.height!=static_cast<unsigned>(config.height))
+                throw std::runtime_error("Camera rejected the configured UYVY resolution; refusing mismatched mask coordinates");
+            stride=format.fmt.pix.bytesperline;
+            if(stride<static_cast<unsigned>(config.width*2))throw std::runtime_error("Invalid camera row stride");
+            v4l2_streamparm params{};params.type=V4L2_BUF_TYPE_VIDEO_CAPTURE;
+            if(ioctl(fd,VIDIOC_G_PARM,&params)==0&&(params.parm.capture.capability&V4L2_CAP_TIMEPERFRAME)) {
+                params.parm.capture.timeperframe.numerator=1;params.parm.capture.timeperframe.denominator=config.fps;
+                checkedIoctl(fd,VIDIOC_S_PARM,&params,"VIDIOC_S_PARM");
+                auto const& interval=params.parm.capture.timeperframe;
+                if(!interval.numerator||interval.denominator!=static_cast<unsigned>(config.fps)*interval.numerator)
+                    throw std::runtime_error("Camera rejected requested frame rate");
+            }
+            v4l2_requestbuffers request{};request.count=config.buffers;
+            request.type=V4L2_BUF_TYPE_VIDEO_CAPTURE;request.memory=V4L2_MEMORY_MMAP;
+            checkedIoctl(fd,VIDIOC_REQBUFS,&request,"VIDIOC_REQBUFS");
+            if(request.count<2)throw std::runtime_error("Insufficient V4L2 buffers");
+            buffers.resize(request.count);
+            for(unsigned i=0;i<request.count;++i) {
+                v4l2_buffer buffer{};buffer.type=request.type;buffer.memory=request.memory;buffer.index=i;
+                checkedIoctl(fd,VIDIOC_QUERYBUF,&buffer,"VIDIOC_QUERYBUF");
+                if(buffer.length<std::size_t(stride)*config.height)throw std::runtime_error("Camera buffer smaller than negotiated frame");
+                buffers[i].size=buffer.length;
+                auto mapping=mmap(nullptr,buffer.length,PROT_READ|PROT_WRITE,MAP_SHARED,fd,buffer.m.offset);
+                if(mapping==MAP_FAILED)throw std::runtime_error("Camera mmap failed");
+                buffers[i].data=mapping;checkedIoctl(fd,VIDIOC_QBUF,&buffer,"VIDIOC_QBUF");
+            }
+            converter=std::make_unique<UyvyConverter>(config.width,config.height);
+            auto type=V4L2_BUF_TYPE_VIDEO_CAPTURE;checkedIoctl(fd,VIDIOC_STREAMON,&type,"VIDIOC_STREAMON");streaming=true;
+        } catch(...) {release();throw;}
+    }
+    ~Impl(){release();}
+    void release() noexcept {
+        if(streaming){auto type=V4L2_BUF_TYPE_VIDEO_CAPTURE;ioctl(fd,VIDIOC_STREAMOFF,&type);streaming=false;}
+        for(auto const& b:buffers)if(b.data)munmap(b.data,b.size);
+        buffers.clear();
+        if(fd>=0){close(fd);fd=-1;}
+        if(stopFd>=0){close(stopFd);stopFd=-1;}
+    }
+    void stop() noexcept {
+        stopped=true;std::uint64_t one=1;
+        if(stopFd>=0){auto ignored=write(stopFd,&one,sizeof(one));(void)ignored;}
+    }
+    std::optional<CapturedFrame> readFrame() {
+        while(!stopped.load()) {
+            pollfd descriptors[2]={{fd,POLLIN,0},{stopFd,POLLIN,0}};
+            int ready=poll(descriptors,2,config.timeout_ms);
+            if(ready<0&&errno==EINTR)continue;
+            if(stopped.load()||(descriptors[1].revents&POLLIN))return {};
+            if(ready==0)throw std::runtime_error("Camera capture timeout: "+config.device);
+            if(ready<0||(descriptors[0].revents&(POLLERR|POLLHUP|POLLNVAL)))throw std::runtime_error("Camera disconnected or poll failed");
+            v4l2_buffer buffer{};buffer.type=V4L2_BUF_TYPE_VIDEO_CAPTURE;buffer.memory=V4L2_MEMORY_MMAP;
+            int status=ioctl(fd,VIDIOC_DQBUF,&buffer);
+            if(status<0&&(errno==EAGAIN||errno==EINTR))continue;
+            if(status<0)throw std::runtime_error("VIDIOC_DQBUF: "+std::string(std::strerror(errno)));
+            if(buffer.index>=buffers.size())throw std::runtime_error("Invalid V4L2 buffer index");
+            // The GMSL driver can mark its startup buffer as corrupt. Requeue
+            // it without conversion or inference; persistent corruption is fatal.
+            if(buffer.flags&V4L2_BUF_FLAG_ERROR) {
+                checkedIoctl(fd,VIDIOC_QBUF,&buffer,"VIDIOC_QBUF(error frame)");
+                if(++consecutiveErrors>=static_cast<unsigned>(config.fps))
+                    throw std::runtime_error("Camera produced one second of consecutive error frames");
+                continue;
+            }
+            CapturedFrame frame;
+            try {
+                if(buffer.bytesused<std::size_t(stride)*config.height||buffer.bytesused>buffers[buffer.index].size)
+                    throw std::runtime_error("Invalid camera frame: flags="+std::to_string(buffer.flags)+", used="+std::to_string(buffer.bytesused)+", required="+std::to_string(std::size_t(stride)*config.height)+", allocated="+std::to_string(buffers[buffer.index].size));
+                auto& m=frame.metadata;
+                m.discarded_error_frames=consecutiveErrors;consecutiveErrors=0;
+                m.driver_sequence=buffer.sequence;m.received_monotonic_ns=monotonicNs();
+                m.driver_timestamp_ns=std::uint64_t(buffer.timestamp.tv_sec)*1000000000ULL+std::uint64_t(buffer.timestamp.tv_usec)*1000;
+                m.driver_timestamp_monotonic=(buffer.flags&V4L2_BUF_FLAG_TIMESTAMP_MASK)==V4L2_BUF_FLAG_TIMESTAMP_MONOTONIC;
+                m.capture_monotonic_ns=m.driver_timestamp_monotonic&&m.driver_timestamp_ns>0&&m.driver_timestamp_ns<=m.received_monotonic_ns?m.driver_timestamp_ns:m.received_monotonic_ns;
+                if(haveSequence) {
+                    auto delta=static_cast<std::uint32_t>(buffer.sequence-lastSequence);
+                    if(delta==0||delta>0x7fffffffU)throw std::runtime_error("Camera sequence reset or repeated frame");
+                    m.skipped_driver_frames=delta-1;index+=delta;
+                }
+                m.frame_index=index;
+                lastSequence=buffer.sequence;haveSequence=true;
+                // Convert before requeuing. No published result references MMAP memory.
+                frame.image=converter->convert(buffers[buffer.index].data,stride);
+            } catch(...) {ioctl(fd,VIDIOC_QBUF,&buffer);throw;}
+            checkedIoctl(fd,VIDIOC_QBUF,&buffer,"VIDIOC_QBUF");return frame;
+        }
+        return {};
+    }
+};
+Camera::Camera(CameraConfig config):impl_(std::make_unique<Impl>(std::move(config))){}
+Camera::~Camera()=default;
+std::optional<CapturedFrame> Camera::read(){return impl_->readFrame();}
+void Camera::requestStop() noexcept {impl_->stop();}
+}
@@ -0,0 +1,178 @@
+#include "camera_pipeline.hpp"
+#include "bounded_queue.hpp"
+#include "latest_result.hpp"
+#include "detector.hpp"
+#include "door_risk.hpp"
+#include "records.hpp"
+#include "result_conversion.hpp"
+#include <opencv2/core.hpp>
+#include <atomic>
+#include <cmath>
+#include <cstring>
+#include <exception>
+#include <mutex>
+#include <limits>
+#include <thread>
+
+namespace cms {
+namespace {
+struct Stream {
+    cudaStream_t value{};
+    Stream(){CUDA_CHECK(cudaStreamCreateWithFlags(&value,cudaStreamNonBlocking));}
+    ~Stream(){cudaStreamDestroy(value);}
+};
+cv::Mat imageView(BgrImage const& image) {
+    return cv::Mat(image.height,image.width,CV_8UC3,const_cast<std::uint8_t*>(image.pixels.data()),image.stride);
+}
+double elapsedMs(std::uint64_t start){return (monotonicNs()-start)*1e-6;}
+void attachPoses(FrameResult& result,fangjia::PoseFrameResult const& fused) {
+    for(std::size_t i=0;i<fused.selectedPeople.size();++i) {
+        auto const& source=fused.selectedPeople[i];
+        PersonResult person;person.track_id=source.box.trackId;person.confidence=source.box.score;
+        person.box={source.box.x1,source.box.y1,source.box.x2,source.box.y2};
+        for(std::size_t k=0;k<133;++k) {
+            auto const&a=source.keypoints[k];auto const&b=fused.smoothedPeople[i].keypoints[k];auto const&c=fused.v2People[i].keypoints[k];
+            person.selected[k]={a.x,a.y,a.score};person.smoothed[k]={b.x,b.y,b.score};person.temporal[k]={c.x,c.y,c.score};
+        }
+        for(int side=0;side<2;++side){auto const& m=fused.metadata[i][side];person.hands[side]={m.handState,m.handSource,m.limbConsensus,m.distanceEligible,m.alarmEligible};}
+        result.persons.push_back(std::move(person));
+    }
+    for(auto const& source:fused.globalHands) {
+        HandRegion hand;hand.person_id=source.personId;hand.track_id=source.trackId;hand.side=source.left?"left":"right";
+        for(std::size_t k=0;k<21;++k)hand.keypoints[k]={source.keypoints[k].x,source.keypoints[k].y,source.keypoints[k].score};
+        result.hand_regions.push_back(std::move(hand));
+    }
+}
+}
+struct CameraPipeline::Impl {
+    struct Detected {CapturedFrame frame;std::vector<fangjia::Box> boxes;double elapsed_ms{};};
+    PipelineConfig config;
+    BoundedQueue<CapturedFrame> captures;
+    BoundedQueue<Detected> detections;
+    LatestResult results;
+    fangjia::TrtLogger logger;
+    std::unique_ptr<Detector> detector;
+    std::unique_ptr<fangjia::RtmwEngine> pose;
+    std::unique_ptr<DoorRisk> risk;
+    std::shared_ptr<const BinaryMask> doorPixels;
+    std::unique_ptr<Camera> camera;
+    std::thread captureThread,detectorThread,inferenceThread;
+    mutable std::mutex stateMutex;
+    std::exception_ptr failure;
+    bool started{};
+    std::atomic<bool> stopped{false},active{false};
+    std::atomic<std::uint64_t> captured{0},published{0},gaps{0},errorFrames{0};
+    double totalLatency{};
+
+    explicit Impl(PipelineConfig c):config(std::move(c)),captures(config.queue_capacity),detections(config.queue_capacity) {
+        if(config.queue_capacity>8)throw std::invalid_argument("Queue capacity must be between 1 and 8");
+        if(config.max_people<1||config.max_people>16||!std::isfinite(config.detector_confidence)||config.detector_confidence<=0||config.detector_confidence>=1)
+            throw std::invalid_argument("Invalid detector configuration");
+        if(!std::filesystem::is_regular_file(config.door_mask))throw std::invalid_argument("Missing door mask: "+config.door_mask.string());
+    }
+    ~Impl(){stop();joinWorkers();}
+    void joinWorkers() noexcept {
+        if(captureThread.joinable())captureThread.join();
+        if(detectorThread.joinable())detectorThread.join();
+        if(inferenceThread.joinable())inferenceThread.join();
+    }
+    void stop() noexcept {
+        stopped=true;active=false;
+        // Camera stays alive until all threads have joined. start()/destruction
+        // are owner-thread operations; requestStop() only interrupts existing work.
+        {std::lock_guard<std::mutex> lock(stateMutex);if(camera)camera->requestStop();}
+        captures.close(true);detections.close(true);
+    }
+    void fail() noexcept {
+        {std::lock_guard<std::mutex> lock(stateMutex);if(!failure)failure=std::current_exception();}
+        stop();
+    }
+    void rethrowFailure() const {
+        std::lock_guard<std::mutex> lock(stateMutex);if(failure)std::rethrow_exception(failure);
+    }
+    void start() {
+        {std::lock_guard<std::mutex> lock(stateMutex);if(started||stopped)throw std::logic_error("CameraPipeline can only be started once");started=true;}
+        try {
+            detector=std::make_unique<Detector>(config.engines/"yolo11s.engine",logger,config.detector_confidence,config.max_people);
+            pose=std::make_unique<fangjia::RtmwEngine>(config.engines/"rtmw_l_384.engine",logger,config.camera.width,config.camera.height,config.max_people);
+            DoorOptions options;options.constant=config.door_mask;options.warning=config.warning_relative;options.release=config.release_relative;
+            risk=std::make_unique<DoorRisk>(options,cv::Size(config.camera.width,config.camera.height));risk->initialize();
+            if(risk->mask().empty()||cv::countNonZero(risk->mask())==0)throw std::runtime_error("Door mask has no usable region");
+            auto pixels=std::make_shared<BinaryMask>();pixels->width=config.camera.width;pixels->height=config.camera.height;pixels->stride=config.camera.width;
+            auto const& mask=risk->mask();pixels->pixels.resize(std::size_t(mask.cols)*mask.rows);
+            for(int y=0;y<mask.rows;++y)std::memcpy(pixels->pixels.data()+std::size_t(y)*mask.cols,mask.ptr(y),mask.cols);
+            doorPixels=std::move(pixels);
+            Stream warmup;
+            pose->warmup(warmup.value);
+            detector->infer(cv::Mat(config.camera.height,config.camera.width,CV_8UC3,cv::Scalar()),warmup.value);
+            {std::lock_guard<std::mutex> lock(stateMutex);camera=std::make_unique<Camera>(config.camera);if(stopped)camera->requestStop();}
+            active=!stopped.load();
+            captureThread=std::thread([this]{try{captureLoop();}catch(...){fail();}});
+            detectorThread=std::thread([this]{try{detectorLoop();}catch(...){fail();}});
+            inferenceThread=std::thread([this]{try{inferenceLoop();}catch(...){fail();}});
+        }catch(...){stop();joinWorkers();throw;}
+    }
+    void captureLoop() {
+        while(!stopped) {
+            auto frame=camera->read();if(!frame)break;
+            gaps+=frame->metadata.skipped_driver_frames;
+            errorFrames+=frame->metadata.discarded_error_frames;
+            if(!captures.push(std::move(*frame)))break;
+            ++captured;
+        }
+        captures.close();
+    }
+    void detectorLoop() {
+        Stream stream;
+        while(!stopped) {
+            auto frame=captures.pop();if(!frame)break;
+            auto startedAt=monotonicNs();auto boxes=detector->infer(imageView(*frame->image),stream.value);
+            if(!detections.push({std::move(*frame),std::move(boxes),elapsedMs(startedAt)}))break;
+        }
+        detections.close();
+    }
+    void inferenceLoop() {
+        Stream stream;fangjia::BoxTracker boxesTracker;fangjia::PoseFusionTracker fusion;
+        while(!stopped) {
+            auto detected=detections.pop();if(!detected)break;
+            if(detected->frame.metadata.frame_index>static_cast<std::uint64_t>(std::numeric_limits<int>::max()))throw std::runtime_error("Frame index limit reached; restart pipeline");
+            auto index=static_cast<int>(detected->frame.metadata.frame_index);
+            auto& boxes=detected->boxes;
+            auto frame=imageView(*detected->frame.image);
+            boxesTracker.update(boxes,index,frame.cols,frame.rows);
+            auto poseStart=monotonicNs();
+            auto raw=pose->infer(frame,boxes,stream.value);auto isolated=raw;
+            for(std::size_t p=0;p<boxes.size();++p) {
+                auto const& box=boxes[p];double area=std::max(double(box.x2-box.x1)*(box.y2-box.y1),1.0);
+                bool needsIsolation=false;
+                for(std::size_t q=0;q<boxes.size();++q)if(q!=p&&std::max(double(boxes[q].x2-boxes[q].x1)*(boxes[q].y2-boxes[q].y1),1.0)<0.5*area)needsIsolation=true;
+                if(needsIsolation){auto masked=maskOtherPeople(frame,boxes,p);isolated[p]=pose->infer(masked,{box},stream.value).at(0);}
+            }
+            auto poseElapsed=elapsedMs(poseStart);auto fusionStart=monotonicNs();
+            auto fused=fusion.update(raw,isolated,index);
+            auto record=makeRecord(index,raw,isolated,fused,false,false);
+            auto riskData=risk->evaluate(index,record);
+            auto result=std::make_shared<FrameResult>(makeFrameResult(record,riskData,detected->frame.metadata));
+            attachPoses(*result,fused);
+            result->image=std::move(detected->frame.image);result->door_mask_pixels=doorPixels;
+            result->door_mask_calibration_verified=config.door_mask_calibration_verified;
+            result->timing={detected->elapsed_ms,poseElapsed,elapsedMs(fusionStart),elapsedMs(result->capture.capture_monotonic_ns)};
+            FrameResultPtr immutable=result;
+            results.publish(std::move(immutable));
+            {std::lock_guard<std::mutex> lock(stateMutex);totalLatency+=result->timing.capture_to_result_ms;++published;}
+        }
+        active=false;
+    }
+};
+CameraPipeline::CameraPipeline(PipelineConfig config):impl_(std::make_unique<Impl>(std::move(config))){}
+CameraPipeline::~CameraPipeline()=default;
+void CameraPipeline::start(){impl_->start();}
+FrameResultPtr CameraPipeline::latest()const{return impl_->results.load();}
+bool CameraPipeline::running()const noexcept{return impl_->active.load();}
+void CameraPipeline::requestStop()noexcept{impl_->stop();}
+void CameraPipeline::join(){impl_->joinWorkers();{std::lock_guard<std::mutex> lock(impl_->stateMutex);impl_->camera.reset();}impl_->rethrowFailure();}
+PipelineStats CameraPipeline::stats()const {
+    std::lock_guard<std::mutex> lock(impl_->stateMutex);auto count=impl_->published.load();
+    return {impl_->captured.load(),count,impl_->gaps.load(),impl_->errorFrames.load(),count?impl_->totalLatency/count:0};
+}
+}
@@ -0,0 +1,24 @@
+#include "config.hpp"
+#include "door_risk.hpp"
+#include <fstream>
+#include <stdexcept>
+
+namespace cms {
+PipelineConfig loadConfig(std::filesystem::path const& path) {
+    std::ifstream file(path);
+    if(!file)throw std::runtime_error("Cannot read configuration: "+path.string());
+    auto j=Json::parse(file);PipelineConfig config;
+    auto const& camera=j.at("camera");
+    config.camera.device=camera.at("device");config.camera.width=camera.at("width");
+    config.camera.height=camera.at("height");config.camera.fps=camera.at("fps");
+    config.camera.buffers=camera.at("buffers");config.camera.timeout_ms=camera.at("timeout_ms");
+    auto resolve=[&](std::string const& value){auto p=std::filesystem::path(value);return std::filesystem::absolute(p.is_absolute()?p:path.parent_path()/p);};
+    auto const& inference=j.at("inference");
+    config.engines=resolve(inference.at("engine_directory"));config.max_people=inference.at("max_people");
+    config.detector_confidence=inference.at("detector_confidence");config.queue_capacity=inference.at("queue_capacity");
+    auto const& door=j.at("door");config.door_mask=resolve(door.at("mask"));
+    config.warning_relative=door.at("warning_relative");config.release_relative=door.at("release_relative");
+    config.door_mask_calibration_verified=door.at("calibration_verified");
+    return config;
+}
+}
@@ -0,0 +1,70 @@
+#include "detector.hpp"
+#include <opencv2/imgproc.hpp>
+#include <algorithm>
+#include <cmath>
+namespace cms {
+using fangjia::Box;
+Detector::Detector(std::filesystem::path const& path,fangjia::TrtLogger& logger,float confidence,int maxPeople):engine_(path,logger),confidence_(confidence),maxPeople_(maxPeople) {
+    if(engine_.inputs().size()!=1||engine_.outputs().size()!=1) throw std::runtime_error("Expected YOLO raw ONNX output (one input/output)");
+    input_=engine_.inputs()[0];output_=engine_.outputs()[0];
+    auto d=engine_.shape(input_);
+    if(d.nbDims!=4||d.d[0]!=1||d.d[1]!=3||d.d[2]<=0||d.d[3]<=0) throw std::runtime_error("Expected fixed YOLO NCHW input");
+    height_=d.d[2];width_=d.d[3];
+    engine_.prepare();
+    d=engine_.shape(output_);
+    if(d.nbDims!=3||d.d[0]!=1||d.d[1]!=84) throw std::runtime_error("Expected YOLO output [1,84,N]");
+    if(engine_.dataType(input_)!=nvinfer1::DataType::kFLOAT || engine_.dataType(output_)!=nvinfer1::DataType::kFLOAT) throw std::runtime_error("Expected float32 engine I/O");
+    preprocess_=std::make_unique<YoloPreprocess>(width_,height_);outputHost_.resize(engine_.bytes(output_)/sizeof(float));
+}
+std::vector<Box> Detector::infer(cv::Mat const& frame,cudaStream_t stream,int filterHeight) {
+    double gain=std::min(double(width_)/frame.cols,double(height_)/frame.rows);
+    int rw=int(std::nearbyint(frame.cols*gain)),rh=int(std::nearbyint(frame.rows*gain));
+    int left=int(std::nearbyint((width_-rw)/2.0-0.1)),top=int(std::nearbyint((height_-rh)/2.0-0.1));
+    preprocess_->enqueue(frame,static_cast<float*>(engine_.device(input_)),stream);
+    if(!engine_.enqueue(stream)) throw std::runtime_error("YOLO inference failed");
+    CUDA_CHECK(cudaMemcpyAsync(outputHost_.data(),engine_.device(output_),engine_.bytes(output_),cudaMemcpyDeviceToHost,stream));
+    CUDA_CHECK(cudaStreamSynchronize(stream));
+    int n=engine_.shape(output_).d[2];std::vector<Box> candidates;
+    for(int i=0;i<n;++i) {
+        int cls=0; float score=outputHost_[4*n+i];
+        for(int c=1;c<80;++c) if(outputHost_[(4+c)*n+i]>score) {cls=c;score=outputHost_[(4+c)*n+i];}
+        if(cls!=0||score<=confidence_||!std::isfinite(score)) continue;
+        float x=outputHost_[i],y=outputHost_[n+i],w=outputHost_[2*n+i],h=outputHost_[3*n+i];
+        if(!std::isfinite(x)||!std::isfinite(y)||!std::isfinite(w)||!std::isfinite(h)||w<=0||h<=0) continue;
+        candidates.push_back({x-w/2,y-h/2,x+w/2,y+h/2,score,0});
+    }
+    std::stable_sort(candidates.begin(),candidates.end(),[](auto const&a,auto const&b){return a.score>b.score;});
+    std::vector<Box> kept;
+    for(auto b:candidates) {
+        bool suppressed=false;
+        for(auto const& a:kept) {
+            float inter=std::max(0.F,std::min(a.x2,b.x2)-std::max(a.x1,b.x1))*std::max(0.F,std::min(a.y2,b.y2)-std::max(a.y1,b.y1));
+            float uni=(a.x2-a.x1)*(a.y2-a.y1)+(b.x2-b.x1)*(b.y2-b.y1)-inter;
+            if(inter/uni>0.30F) {suppressed=true;break;}
+        }
+        if(!suppressed) kept.push_back(b);
+        if(kept.size()==300) break;
+    }
+    std::vector<Box> filtered;
+    for(auto b:kept) {
+        b.x1=std::clamp(float((b.x1-left)/gain),0.F,float(frame.cols));b.x2=std::clamp(float((b.x2-left)/gain),0.F,float(frame.cols));
+        b.y1=std::clamp(float((b.y1-top)/gain),0.F,float(frame.rows));b.y2=std::clamp(float((b.y2-top)/gain),0.F,float(frame.rows));
+        if((b.x1+b.x2)*0.5F>0.80F*frame.cols||b.y2-b.y1<0.20F*(filterHeight>0?filterHeight:frame.rows)) continue;
+        filtered.push_back(b);
+    }
+    if(int(filtered.size())>maxPeople_) filtered.resize(maxPeople_);
+    std::stable_sort(filtered.begin(),filtered.end(),[](auto const&a,auto const&b){return a.x1<b.x1;});
+    return filtered;
+}
+cv::Mat maskOtherPeople(cv::Mat const& frame,std::vector<Box> const& boxes,size_t target) {
+    auto masked=frame.clone(); auto const&t=boxes.at(target);
+    double targetArea=std::max(double(t.x2-t.x1)*(t.y2-t.y1),1.0);
+    for(size_t i=0;i<boxes.size();++i) {
+        auto const& b=boxes[i]; if(i==target||std::max(double(b.x2-b.x1)*(b.y2-b.y1),1.0)>=0.5*targetArea) continue;
+        int x1=std::clamp(int(std::nearbyint(b.x1)),0,frame.cols-1),x2=std::clamp(int(std::nearbyint(b.x2)),0,frame.cols-1);
+        int y1=std::clamp(int(std::nearbyint(b.y1)),0,frame.rows-1),y2=std::clamp(int(std::nearbyint(b.y2)),0,frame.rows-1);
+        if(x2>=x1&&y2>=y1) cv::rectangle(masked,{x1,y1},{x2,y2},{124,116,104},-1);
+    }
+    return masked;
+}
+}
@@ -0,0 +1,137 @@
+#include "door_risk.hpp"
+#include <opencv2/imgproc.hpp>
+#include <opencv2/imgcodecs.hpp>
+#include <algorithm>
+#include <cmath>
+#include <fstream>
+#include <regex>
+#include <stdexcept>
+namespace cms {
+double rounded(double v, int d) { double s=std::pow(10.0,d); return std::nearbyint(v*s)/s; }
+static std::string lower(std::string s) { for(auto& c:s) if(static_cast<unsigned char>(c)<128) c=static_cast<char>(std::tolower(c)); return s; }
+static bool raster(std::filesystem::path const& p) { auto e=lower(p.extension().string()); return e==".png"||e==".bmp"||e==".tif"||e==".tiff"; }
+DoorRisk::DoorRisk(DoorOptions o, cv::Size size):options_(std::move(o)),size_(size) {
+    if(options_.constant.empty()==options_.directory.empty()) throw std::runtime_error("Provide exactly one constant mask or door-mask directory");
+    if(options_.warning<0 || options_.release<=options_.warning) throw std::runtime_error("Risk thresholds require 0 <= warning < release");
+    format_=options_.format;
+    if(!options_.constant.empty()) {
+        if(!std::filesystem::is_regular_file(options_.constant)) throw std::runtime_error("Missing mask: "+options_.constant.string());
+        if(format_=="auto") format_=lower(options_.constant.extension().string())==".json"?"xanylabeling":"raster";
+    } else {
+        std::vector<std::filesystem::path> files;
+        for(auto const& e:std::filesystem::recursive_directory_iterator(options_.directory)) if(e.is_regular_file()) files.push_back(e.path());
+        std::sort(files.begin(),files.end());
+        if(format_=="auto") format_=std::any_of(files.begin(),files.end(),[](auto const& p){return p.extension()==".json";})?"xanylabeling":"raster";
+        std::regex digits("[0-9]+");
+        for(auto const& p:files) {
+            if(format_=="xanylabeling"?p.extension()!=".json":!raster(p)) continue;
+            int index=-1; std::string stem=p.stem().string();
+            for(std::sregex_iterator i(stem.begin(),stem.end(),digits),end;i!=end;++i) index=std::stoi(i->str());
+            if(index<0 && format_=="xanylabeling") { std::ifstream in(p); auto j=Json::parse(in); if(j.contains("frame_index")) index=j["frame_index"].get<int>(); else if(j.contains("frameIndex")) index=j["frameIndex"].get<int>(); }
+            if(index<0) continue;
+            if(!paths_.emplace(index,p).second) throw std::runtime_error("Duplicate mask frame: "+std::to_string(index));
+        }
+        if(paths_.empty()) throw std::runtime_error("No indexed door masks in "+options_.directory.string());
+    }
+    if(format_!="raster" && format_!="xanylabeling") throw std::runtime_error("Unknown door mask format");
+}
+cv::Mat DoorRisk::rasterize(std::filesystem::path const& path) const {
+    cv::Mat mask(size_,CV_8U,cv::Scalar(0));
+    if(format_=="raster") {
+        auto img=cv::imread(path.string(),cv::IMREAD_UNCHANGED);
+        if(img.empty()) throw std::runtime_error("Cannot read mask: "+path.string());
+        if(img.channels()>1) { std::vector<cv::Mat> ch; cv::split(img,ch); img=ch[0]; for(size_t i=1;i<std::min(ch.size(),size_t(3));++i) cv::max(img,ch[i],img); }
+        if(img.size()!=size_) cv::resize(img,img,size_,0,0,cv::INTER_NEAREST);
+        return img>0;
+    }
+    std::ifstream in(path); auto j=Json::parse(in);
+    double w=j.contains("imageWidth")&&!j["imageWidth"].is_null()?j["imageWidth"].get<double>():size_.width;
+    double h=j.contains("imageHeight")&&!j["imageHeight"].is_null()?j["imageHeight"].get<double>():size_.height;
+    if(w==0) w=size_.width;
+    if(h==0) h=size_.height;
+    float sx=static_cast<float>(size_.width/std::max(w,1.0)),sy=static_cast<float>(size_.height/std::max(h,1.0));
+    for(auto const& s:j.value("shapes",Json::array())) {
+        if(!options_.labels.count(lower(s.value("label",std::string())))) continue;
+        std::vector<cv::Point> points;
+        for(auto const& p:s.value("points",Json::array())) points.emplace_back(static_cast<int>(std::nearbyint(p[0].get<float>()*sx)),static_cast<int>(std::nearbyint(p[1].get<float>()*sy)));
+        auto type=lower(s.value("shape_type",std::string("polygon")));
+        if(points.size()<2) continue;
+        if(type=="rectangle") cv::rectangle(mask,points[0],points[1],255,-1);
+        else if(type=="circle") cv::circle(mask,points[0],static_cast<int>(std::nearbyint(cv::norm(points[1]-points[0]))),255,-1);
+        else if(points.size()>=3) cv::fillPoly(mask,std::vector<std::vector<cv::Point>>{points},255);
+    }
+    return mask;
+}
+void DoorRisk::load(int frame) {
+    if(cached_) return;
+    loadedPath_.clear(); mask_.release(); distance_.release();
+    if(!options_.constant.empty()) loadedPath_=options_.constant;
+    else { auto it=paths_.find(frame+options_.offset); if(it!=paths_.end()) loadedPath_=it->second; }
+    if(loadedPath_.empty()) {
+        status_="missing";
+        if(options_.require) throw std::runtime_error("No door mask for frame "+std::to_string(frame));
+        return;
+    }
+    mask_=rasterize(loadedPath_);
+    status_=cv::countNonZero(mask_)?"loaded":"empty";
+    if(status_=="loaded") cv::distanceTransform(mask_==0,distance_,cv::DIST_L2,cv::DIST_MASK_PRECISE);
+    cached_=!options_.constant.empty();
+}
+Json DoorRisk::evaluate(int frame, Json& record) {
+    load(frame);
+    Json regions=Json::array(),fallbacks=Json::array(),nearest=nullptr;
+    std::set<int> withHands;
+    auto measure=[&](cv::Mat const& region, std::string const& source,int track,std::string const& side,std::string const& state) {
+        Json m={{"status","door_mask_"+status_},{"risk_score",nullptr},{"risk_level","unavailable"},{"decision","EVIDENCE_UNAVAILABLE"},{"source",source},{"track_id",track},{"side",side}};
+        if(status_=="loaded") {
+            int area=cv::countNonZero(region); cv::Mat intersection; cv::bitwise_and(region,mask_,intersection); int overlap=cv::countNonZero(intersection);
+            double distance=0; cv::minMaxLoc(distance_,&distance,nullptr,nullptr,nullptr,region);
+            double scale=std::max(std::sqrt(double(area)),1.0),relative=distance/scale;
+            double score=overlap?1:std::clamp(1-relative/options_.release,0.0,1.0);
+            std::string level=overlap?"danger":relative<=options_.warning?"warning":relative<=options_.release?"watch":"clear";
+            std::string decision=overlap?"DANGER_INTERSECTION":relative<=options_.warning?"WARNING_NEAR":relative<=options_.release?"WATCH":"CLEAR";
+            m={{"status","calculated"},{"risk_score",rounded(score,5)},{"risk_level",level},{"decision",decision},{"source",source},{"track_id",track},{"side",side},{"state",state},{"region_area_px",area},{"door_overlap_px",overlap},{"region_overlap_ratio",rounded(double(overlap)/std::max(area,1),6)},{"min_distance_px",rounded(distance,3)},{"region_scale_px",rounded(scale,3)},{"distance_relative",rounded(relative,6)}};
+            if(nearest.is_null() || m["risk_score"].get<double>()>nearest["risk_score"].get<double>()) nearest=m;
+        }
+        regions.push_back(m); return m;
+    };
+    auto const ids=record.value("track_ids",Json::array());
+    auto trackFor=[&](int pos){return pos>=0&&pos<int(ids.size())?ids[pos].get<int>():pos+1;};
+    if(!record.contains("global_hands")) record["global_hands"]=Json::array();
+    for(auto& hand:record["global_hands"]) {
+        int pos=hand.value("person_id",0)-1; std::vector<cv::Point> points;
+        bool finite=true;
+        for(auto const& p:hand.value("keypoints",Json::array())) {
+            float x=p[0].get<float>(),y=p[1].get<float>();
+            if(!std::isfinite(x)||!std::isfinite(y)) {finite=false;break;}
+            points.emplace_back(static_cast<int>(std::nearbyint(x)),static_cast<int>(std::nearbyint(y)));
+        }
+        if(!finite||points.size()<3) continue;
+        std::vector<cv::Point> hull; cv::convexHull(points,hull);
+        cv::Mat region(size_,CV_8U,cv::Scalar(0)); cv::fillConvexPoly(region,hull,255);
+        if(!cv::countNonZero(region)) continue;
+        withHands.insert(pos);
+        hand["risk"]=measure(region,"RTMW_HAND_KEYPOINTS",trackFor(pos),hand.value("side",std::string("unknown")),hand.value("state",std::string("unknown")));
+    }
+    auto const boxes=record.value("boxes",Json::array());
+    for(int pos=0;pos<int(boxes.size());++pos) {
+        if(withHands.count(pos)) continue;
+        auto const& b=boxes[pos]; if(b.size()<4) continue;
+        int x1=std::clamp(int(std::nearbyint(b[0].get<double>())),0,size_.width-1),y1=std::clamp(int(std::nearbyint(b[1].get<double>())),0,size_.height-1);
+        int x2=std::clamp(int(std::nearbyint(b[2].get<double>())),0,size_.width-1),y2=std::clamp(int(std::nearbyint(b[3].get<double>())),0,size_.height-1);
+        if(x1>x2) std::swap(x1,x2);
+        if(y1>y2) std::swap(y1,y2);
+        if(x2<=x1||y2<=y1) continue;
+        cv::Mat region(size_,CV_8U,cv::Scalar(0)); cv::rectangle(region,{x1,y1},{x2,y2},255,-1);
+        Json fallback={{"track_id",trackFor(pos)},{"side","body"},{"state","rtmw_hand_unavailable"},{"body_zone_box",b}};
+        fallback["risk"]=measure(region,"YOLO_BODY_BOX_FALLBACK",trackFor(pos),"body","rtmw_hand_unavailable"); fallbacks.push_back(fallback);
+    }
+    record["body_fallbacks"]=fallbacks;
+    Json result={{"status",status_!="loaded"?"door_mask_"+status_:nearest.is_null()?"person_region_missing":"calculated"}};
+    for(auto const* key:{"risk_score","risk_level","decision","source","track_id","side"}) result[key]=nearest.is_null()?Json(nullptr):nearest[key];
+    if(nearest.is_null()) {result["risk_level"]="unavailable";result["decision"]="EVIDENCE_UNAVAILABLE";}
+    result["door_mask"]={{"status",status_},{"annotation_index",options_.constant.empty()?frame+options_.offset:frame},{"path",loadedPath_.empty()?Json(nullptr):Json(std::filesystem::absolute(loadedPath_).string())},{"format",format_},{"area_px",mask_.empty()?0:cv::countNonZero(mask_)}};
+    result["warning_distance_relative"]=options_.warning; result["release_distance_relative"]=options_.release; result["regions"]=regions;
+    return result;
+}
+}
@@ -0,0 +1,19 @@
+#include "gpu_preprocess.hpp"
+#include <algorithm>
+#include <cmath>
+namespace cms {
+__global__ void yoloNormalize(float* out,unsigned char const* in,int pixels) {
+    int i=blockDim.x*blockIdx.x+threadIdx.x;if(i>=pixels)return;
+    out[i]=__fdiv_rn(float(in[3*i+2]),255.F);out[pixels+i]=__fdiv_rn(float(in[3*i+1]),255.F);out[2*pixels+i]=__fdiv_rn(float(in[3*i]),255.F);
+}
+void YoloPreprocess::enqueue(cv::Mat const& frame,float* destination,cudaStream_t stream) {
+    double scale=std::min(double(width_)/frame.cols,double(height_)/frame.rows);
+    int rw=int(std::nearbyint(frame.cols*scale)),rh=int(std::nearbyint(frame.rows*scale));
+    int left=int(std::nearbyint((width_-rw)/2.0-.1)),top=int(std::nearbyint((height_-rh)/2.0-.1));
+        size_t bytes=size_t(frame.cols)*frame.rows*3;if(capacity_<bytes){if(source_)CUDA_CHECK(cudaFree(source_));CUDA_CHECK(cudaMalloc(&source_,bytes));capacity_=bytes;}
+        CUDA_CHECK(cudaMemcpy2DAsync(source_,frame.cols*3,frame.data,frame.step,frame.cols*3,frame.rows,cudaMemcpyHostToDevice,stream));
+        auto c=nppContext(stream);Npp8u fill[3]={114,114,114};nppCheck(nppiSet_8u_C3R_Ctx(fill,letter_,width_*3,{width_,height_},c));
+        nppCheck(nppiResizeSqrPixel_8u_C3R_Ctx(source_,{frame.cols,frame.rows},frame.cols*3,{0,0,frame.cols,frame.rows},letter_,width_*3,{left,top,rw,rh},double(rw)/frame.cols,double(rh)/frame.rows,left,top,NPPI_INTER_LINEAR,c));
+    int pixels=width_*height_;yoloNormalize<<<(pixels+255)/256,256,0,stream>>>(destination,letter_,pixels);CUDA_CHECK(cudaGetLastError());
+}
+}
@@ -0,0 +1,43 @@
+#include "camera_pipeline.hpp"
+#include "config.hpp"
+#include "stop_signals.hpp"
+#include "visualizer.hpp"
+#include <functional>
+#include <iostream>
+#include <map>
+
+int main(int argc,char** argv)try {
+    auto root=std::filesystem::canonical("/proc/self/exe").parent_path().parent_path();
+    auto configPath=root/"config/camera.json";
+    bool visualize=false;
+    std::map<std::string,std::function<void(int&)>> arguments{
+        {"--visualize",[&](int&){visualize=true;}},
+        {"--config",[&](int&i){if(++i>=argc)throw std::invalid_argument("--config needs a file");configPath=argv[i];}},
+        {"--help",[](int&){std::cout<<"camera_risk [--config FILE] [--visualize]\nContinuous camera inference; latest FrameResult stays in memory. Stop with Ctrl+C or SIGTERM.\n";std::exit(0);}}
+    };
+    for(int i=1;i<argc;++i){auto handler=arguments.find(argv[i]);if(handler==arguments.end())throw std::invalid_argument("Unknown option: "+std::string(argv[i]));handler->second(i);}
+
+    cms::CameraPipeline pipeline(cms::loadConfig(configPath));
+    cms::StopSignals signals(pipeline);
+    pipeline.start();
+    /*
+     * 下游使用示例：在下游自己的线程/调用位置读取最新结果。
+     * 首帧尚未完成时 result 为空；读取不会影响取图和推理。
+     *
+     * auto result = pipeline.latest();
+     * if (result) {
+     *     const auto& risk = result->risk;
+     *     const auto& people = result->persons;
+     *     // result->frame_index / track_ids / boxes
+     *     // result->image / door_mask_pixels / capture / timing
+     *     // people[i].selected / smoothed / temporal：各 133 个关键点
+     * }
+     */
+    // Optional observer. Closing/failing the window never requests inference stop.
+    if(visualize) {
+        try {cms::runLocalVisualization(pipeline);}
+        catch(std::exception const& e){std::cerr<<"Visualization unavailable; inference continues: "<<e.what()<<std::endl;}
+    }
+    pipeline.join();
+    return 0;
+}catch(std::exception const& e){std::cerr<<"Camera pipeline failed: "<<e.what()<<std::endl;return 1;}
@@ -0,0 +1,875 @@
+#include "pose_engine.hpp"
+#include "common.hpp"
+#include "gpu_preprocess.hpp"
+
+#include <opencv2/imgproc.hpp>
+
+#include <algorithm>
+#include <chrono>
+#include <cmath>
+#include <numeric>
+#include <set>
+#include <tuple>
+
+namespace fangjia {
+namespace {
+
+constexpr int kInputWidth = 288;
+constexpr int kInputHeight = 384;
+constexpr int kKeypoints = 133;
+constexpr int kSimccX = 576;
+constexpr int kSimccY = 768;
+
+__global__ void poseNormalizeKernel(
+    float* output, std::uint8_t const* input, int people, int inputStep) {
+    int index = blockIdx.x * blockDim.x + threadIdx.x;
+    int pixels = kInputWidth * kInputHeight;
+    int total = people * pixels;
+    if (index >= total) return;
+    int person = index / pixels;
+    int pixel = index - person * pixels;
+    int y = pixel / kInputWidth;
+    int x = pixel - y * kInputWidth;
+    auto const* value = input + person * inputStep * kInputHeight + y * inputStep + x * 3;
+    float* destination = output + static_cast<std::size_t>(person) * 3 * pixels;
+    destination[pixel] = static_cast<float>((static_cast<double>(value[0]) - 123.675) / 58.395);
+    destination[pixels + pixel] = static_cast<float>((static_cast<double>(value[1]) - 116.28) / 57.12);
+    destination[2 * pixels + pixel] = static_cast<float>((static_cast<double>(value[2]) - 103.53) / 57.375);
+}
+
+float iou(Box const& a, Box const& b) {
+    float x1 = std::max(a.x1, b.x1);
+    float y1 = std::max(a.y1, b.y1);
+    float x2 = std::min(a.x2, b.x2);
+    float y2 = std::min(a.y2, b.y2);
+    float intersection = std::max(x2 - x1, 0.0F) * std::max(y2 - y1, 0.0F);
+    float areaA = std::max(a.x2 - a.x1, 0.0F) * std::max(a.y2 - a.y1, 0.0F);
+    float areaB = std::max(b.x2 - b.x1, 0.0F) * std::max(b.y2 - b.y1, 0.0F);
+    return intersection / std::max(areaA + areaB - intersection, 1.0F);
+}
+
+float distance(float x1, float y1, float x2, float y2) {
+    return std::hypot(x1 - x2, y1 - y2);
+}
+
+} // namespace
+
+std::vector<int> BoxTracker::update(
+    std::vector<Box>& boxes, int frameIndex, int width, int height) {
+    // Tracks older than Python's max_age can never be matched again.
+    // Reclaim them to bound memory during continuous camera operation.
+    for(auto it=tracks_.begin();it!=tracks_.end();) {
+        if(frameIndex-it->second.lastFrame>30)it=tracks_.erase(it);
+        else ++it;
+    }
+    struct Pair { float score; int detection; int track; };
+    std::vector<Pair> pairs;
+    float imageDiagonal = std::hypot(static_cast<float>(width), static_cast<float>(height));
+    for (int d = 0; d < static_cast<int>(boxes.size()); ++d) {
+        float cx = 0.5F * (boxes[d].x1 + boxes[d].x2);
+        float cy = 0.5F * (boxes[d].y1 + boxes[d].y2);
+        for (auto const& [id, track] : tracks_) {
+            if (frameIndex - track.lastFrame > 30) continue;
+            float pcx = 0.5F * (track.box.x1 + track.box.x2);
+            float pcy = 0.5F * (track.box.y1 + track.box.y2);
+            float overlap = iou(boxes[d], track.box);
+            float centerDistance = distance(cx, cy, pcx, pcy) / imageDiagonal;
+            if (overlap >= 0.05F || centerDistance <= 0.18F) {
+                pairs.push_back({2.0F * overlap + std::max(0.0F, 1.0F - centerDistance / 0.18F), d, id});
+            }
+        }
+    }
+    std::sort(pairs.begin(), pairs.end(), [](Pair const& a, Pair const& b) {
+        if (a.score != b.score) return a.score > b.score;
+        if (a.detection != b.detection) return a.detection > b.detection;
+        return a.track > b.track;
+    });
+    std::vector<int> result(boxes.size(), 0);
+    std::unordered_map<int, bool> assigned;
+    for (auto const& pair : pairs) {
+        if (result[pair.detection] || assigned[pair.track]) continue;
+        result[pair.detection] = pair.track;
+        assigned[pair.track] = true;
+    }
+    for (int d = 0; d < static_cast<int>(boxes.size()); ++d) {
+        if (!result[d]) result[d] = nextId_++;
+        boxes[d].trackId = result[d];
+        tracks_[result[d]] = {boxes[d], frameIndex};
+    }
+    return result;
+}
+
+RtmwEngine::RtmwEngine(
+    std::filesystem::path const& enginePath, TrtLogger& logger,
+    int maxSourceWidth, int maxSourceHeight, int maxPeople)
+    : engine_(enginePath, logger), maxSourceWidth_(maxSourceWidth),
+      maxSourceHeight_(maxSourceHeight), maxPeople_(maxPeople),
+      centersHost_(maxPeople * 2), scalesHost_(maxPeople * 2),
+      simccXHost_(static_cast<std::size_t>(maxPeople) * kKeypoints * kSimccX),
+      simccYHost_(static_cast<std::size_t>(maxPeople) * kKeypoints * kSimccY) {
+    engine_.setInputShape("input", nvinfer1::Dims4{1,3,kInputHeight,kInputWidth});
+    engine_.prepare();
+    auto xs=engine_.shape("simcc_x"),ys=engine_.shape("simcc_y");
+    if(xs.nbDims!=3||ys.nbDims!=3||xs.d[0]!=1||ys.d[0]!=1||xs.d[1]!=kKeypoints||ys.d[1]!=kKeypoints||xs.d[2]!=kSimccX||ys.d[2]!=kSimccY)
+        throw std::runtime_error("RTMW engine I/O shape does not match the 133-point model");
+    for(auto name:{"input","simcc_x","simcc_y"})if(engine_.dataType(name)!=nvinfer1::DataType::kFLOAT)
+        throw std::runtime_error("RTMW requires float32 I/O bindings");
+    std::size_t const values = static_cast<std::size_t>(maxPeople) * 3 * kInputWidth * kInputHeight;
+    CUDA_CHECK(cudaMalloc(
+        &cropsDevice_, static_cast<std::size_t>(maxPeople_) * kInputWidth * kInputHeight * 3));
+    CUDA_CHECK(cudaMalloc(&networkDevice_, values * sizeof(float)));
+    CUDA_CHECK(cudaMalloc(&sourceDevice_, static_cast<size_t>(maxSourceWidth_)*maxSourceHeight_*3));
+}
+
+RtmwEngine::~RtmwEngine() {
+    if (sourceDevice_) cudaFree(sourceDevice_);
+    if (cropsDevice_) cudaFree(cropsDevice_);
+    if (networkDevice_) cudaFree(networkDevice_);
+}
+
+void RtmwEngine::warmup(cudaStream_t stream) {
+    std::vector<Box> boxes{{0, 0, static_cast<float>(maxSourceWidth_), static_cast<float>(maxSourceHeight_), 1, 1}};
+    cv::Mat black(maxSourceHeight_, maxSourceWidth_, CV_8UC3, cv::Scalar(0, 0, 0));
+    for (int i = 0; i < 2; ++i) infer(black, boxes, stream);
+}
+
+std::vector<PersonPose> RtmwEngine::infer(
+    cv::Mat const& frame, std::vector<Box> const& boxes, cudaStream_t stream,
+    float* preprocessMs, float* inferenceMs, float* postprocessMs) {
+    int batch = std::min(static_cast<int>(boxes.size()), maxPeople_);
+    if (batch <= 0) return {};
+    auto const preprocessStarted = std::chrono::steady_clock::now();
+    int const cropStep = kInputWidth * 3;
+    int const pixels = kInputWidth * kInputHeight;
+    if (frame.cols>maxSourceWidth_ || frame.rows>maxSourceHeight_) throw std::runtime_error("Pose input exceeds allocation");
+    auto npp = cms::nppContext(stream);
+    CUDA_CHECK(cudaMemcpy2DAsync(sourceDevice_, frame.cols*3, frame.data, frame.step, frame.cols*3, frame.rows, cudaMemcpyHostToDevice, stream));
+    for (int i = 0; i < batch; ++i) {
+        auto const& box = boxes[i];
+        double width = (static_cast<double>(box.x2) - static_cast<double>(box.x1)) * 1.25;
+        double height = (static_cast<double>(box.y2) - static_cast<double>(box.y1)) * 1.25;
+        constexpr double aspect = static_cast<double>(kInputWidth) / kInputHeight;
+        if (width > height * aspect) height = width / aspect;
+        else width = height * aspect;
+        centersHost_[i * 2] = 0.5 * (static_cast<double>(box.x1) + static_cast<double>(box.x2));
+        centersHost_[i * 2 + 1] = 0.5 * (static_cast<double>(box.y1) + static_cast<double>(box.y2));
+        scalesHost_[i * 2] = width;
+        scalesHost_[i * 2 + 1] = height;
+
+        cv::Point2f source[3]{
+            {static_cast<float>(centersHost_[i * 2]), static_cast<float>(centersHost_[i * 2 + 1])},
+            {static_cast<float>(centersHost_[i * 2]), static_cast<float>(centersHost_[i * 2 + 1] - width * 0.5)},
+            {static_cast<float>(centersHost_[i * 2] - width * 0.5), static_cast<float>(centersHost_[i * 2 + 1] - width * 0.5)}};
+        cv::Point2f destination[3]{
+            {kInputWidth * 0.5F, kInputHeight * 0.5F},
+            {kInputWidth * 0.5F, kInputHeight * 0.5F - kInputWidth * 0.5F},
+            {0.0F, kInputHeight * 0.5F - kInputWidth * 0.5F}};
+        cv::Mat affine = cv::getAffineTransform(source, destination);
+            auto* cropDevice=cropsDevice_+static_cast<size_t>(i)*cropStep*kInputHeight;
+            Npp8u fill[3]={0,0,0};
+            cms::nppCheck(nppiSet_8u_C3R_Ctx(fill,cropDevice,cropStep,{kInputWidth,kInputHeight},npp));
+            double coeff[2][3];
+            for(int row=0;row<2;++row) for(int col=0;col<3;++col) coeff[row][col]=affine.at<double>(row,col);
+            cms::nppCheck(nppiWarpAffine_8u_C3R_Ctx(sourceDevice_,{frame.cols,frame.rows},frame.cols*3,{0,0,frame.cols,frame.rows},cropDevice,cropStep,{0,0,kInputWidth,kInputHeight},coeff,NPPI_INTER_LINEAR,npp));
+    }
+    int const totalPixels = batch * pixels;
+    poseNormalizeKernel<<<(totalPixels + 255) / 256, 256, 0, stream>>>(
+        networkDevice_, cropsDevice_, batch, cropStep);
+    CUDA_CHECK(cudaGetLastError());
+    CUDA_CHECK(cudaStreamSynchronize(stream));
+    if (preprocessMs) {
+        *preprocessMs = static_cast<float>(std::chrono::duration<double, std::milli>(
+            std::chrono::steady_clock::now() - preprocessStarted).count());
+    }
+
+    // RTMPose's Python API invokes the ONNX/TensorRT session once per person.
+    // Preserve that execution contract instead of batching people together.
+    engine_.setInputShape("input", nvinfer1::Dims4{1, 3, kInputHeight, kInputWidth});
+    engine_.prepare();
+    CudaEventTimer timer;
+    float inferTime = 0.0F;
+    std::size_t const inputValues = static_cast<std::size_t>(3) * pixels;
+    std::size_t const xValues = static_cast<std::size_t>(kKeypoints) * kSimccX;
+    std::size_t const yValues = static_cast<std::size_t>(kKeypoints) * kSimccY;
+    for (int person = 0; person < batch; ++person) {
+        CUDA_CHECK(cudaMemcpyAsync(
+            engine_.device("input"), networkDevice_ + static_cast<std::size_t>(person) * inputValues,
+            inputValues * sizeof(float), cudaMemcpyDeviceToDevice, stream));
+        timer.start(stream);
+        if (!engine_.enqueue(stream)) throw std::runtime_error("RTMW enqueue failed");
+        inferTime += timer.stop(stream);
+        CUDA_CHECK(cudaMemcpyAsync(
+            simccXHost_.data() + static_cast<std::size_t>(person) * xValues,
+            engine_.device("simcc_x"), xValues * sizeof(float), cudaMemcpyDeviceToHost, stream));
+        CUDA_CHECK(cudaMemcpyAsync(
+            simccYHost_.data() + static_cast<std::size_t>(person) * yValues,
+            engine_.device("simcc_y"), yValues * sizeof(float), cudaMemcpyDeviceToHost, stream));
+    }
+    if (inferenceMs) *inferenceMs = inferTime;
+
+    auto const postprocessStarted = std::chrono::steady_clock::now();
+    CUDA_CHECK(cudaStreamSynchronize(stream));
+    std::vector<PersonPose> result(batch);
+    for (int person = 0; person < batch; ++person) {
+        result[person].box = boxes[person];
+        for (int keypoint = 0; keypoint < kKeypoints; ++keypoint) {
+            auto xBegin = simccXHost_.begin() + (person * kKeypoints + keypoint) * kSimccX;
+            auto yBegin = simccYHost_.begin() + (person * kKeypoints + keypoint) * kSimccY;
+            auto xIt = std::max_element(xBegin, xBegin + kSimccX);
+            auto yIt = std::max_element(yBegin, yBegin + kSimccY);
+            float score = 0.5F * (*xIt + *yIt);
+            if(!std::isfinite(score))throw std::runtime_error("Non-finite RTMW output");
+            float localX = score > 0.0F ? static_cast<float>(xIt - xBegin) * 0.5F : -0.5F;
+            float localY = score > 0.0F ? static_cast<float>(yIt - yBegin) * 0.5F : -0.5F;
+            double scaleX = scalesHost_[person * 2];
+            double scaleY = scalesHost_[person * 2 + 1];
+            result[person].keypoints[keypoint] = {
+                static_cast<float>(static_cast<double>(localX) / kInputWidth * scaleX + centersHost_[person * 2] - scaleX * 0.5),
+                static_cast<float>(static_cast<double>(localY) / kInputHeight * scaleY + centersHost_[person * 2 + 1] - scaleY * 0.5),
+                score};
+        }
+    }
+    float post = static_cast<float>(std::chrono::duration<double, std::milli>(
+        std::chrono::steady_clock::now() - postprocessStarted).count());
+    if (postprocessMs) *postprocessMs = post;
+    return result;
+}
+
+PoseFrameResult PoseFusionTracker::update(
+    std::vector<PersonPose> const& rawPeople,
+    std::vector<PersonPose> const& isolatedPeople,
+    int frameIndex) {
+    for(auto it=states_.begin();it!=states_.end();) {
+        if(it->second.lastFrame>=0&&frameIndex-it->second.lastFrame>30)it=states_.erase(it);
+        else ++it;
+    }
+    struct SideMeta {
+        bool limbConsensus{};
+        bool handValid{};
+        bool distanceEligible{};
+        bool alarmEligible{};
+        double handScore{};
+        double linkToSelectedLimb{};
+        std::string handSource;
+        std::string handState;
+    };
+    struct FusedPerson {
+        PersonPose pose;
+        std::array<SideMeta, 2> meta;
+    };
+    struct Measurement {
+        std::array<Vec2, 21> hand{};
+        bool precise{};
+        bool provisional{};
+        double quality{};
+        std::string source{"model"};
+    };
+    auto vec = [](Keypoint const& point) { return Vec2{point.x, point.y}; };
+    auto add = [](Vec2 a, Vec2 b) { return Vec2{a.x + b.x, a.y + b.y}; };
+    auto sub = [](Vec2 a, Vec2 b) { return Vec2{a.x - b.x, a.y - b.y}; };
+    auto mul = [](Vec2 a, double value) { return Vec2{a.x * value, a.y * value}; };
+    auto norm = [](Vec2 a) { return std::hypot(a.x, a.y); };
+    auto median = [](std::vector<double> values) {
+        if (values.empty()) return 0.0;
+        std::sort(values.begin(), values.end());
+        auto middle = values.size() / 2;
+        return values.size() % 2 ? values[middle] : 0.5 * (values[middle - 1] + values[middle]);
+    };
+    auto percentile90 = [](std::vector<double> values) {
+        if (values.empty()) return 0.0;
+        std::sort(values.begin(), values.end());
+        double position = 0.9 * (values.size() - 1);
+        auto lower = static_cast<std::size_t>(position);
+        auto upper = std::min(lower + 1, values.size() - 1);
+        return values[lower] * (upper - position) + values[upper] * (position - lower);
+    };
+    auto boxDiagonal = [](Box const& box) {
+        return std::max(std::hypot(box.x2 - box.x1, box.y2 - box.y1), 1.0F);
+    };
+    auto normalized = [](Keypoint const& point, Box const& box) {
+        return Vec2{
+            (point.x - box.x1) / std::max(box.x2 - box.x1, 1.0F),
+            (point.y - box.y1) / std::max(box.y2 - box.y1, 1.0F)};
+    };
+    auto denormalized = [](Vec2 point, Box const& box) {
+        return Vec2{
+            point.x * std::max(box.x2 - box.x1, 1.0F) + box.x1,
+            point.y * std::max(box.y2 - box.y1, 1.0F) + box.y1};
+    };
+    auto solveTwoBoneElbow = [&](Vec2 shoulder, Vec2 wrist, double upperLength,
+                                 double forearmLength, Vec2 reference) {
+        Vec2 direction = sub(wrist, shoulder);
+        double length = norm(direction);
+        if (length < 1e-6) return reference;
+        Vec2 unit = mul(direction, 1.0 / length);
+        if (length >= upperLength + forearmLength) return add(shoulder, mul(unit, upperLength));
+        if (length <= std::abs(upperLength - forearmLength)) {
+            double sign = upperLength >= forearmLength ? 1.0 : -1.0;
+            return add(shoulder, mul(unit, upperLength * sign));
+        }
+        double along = (upperLength * upperLength - forearmLength * forearmLength + length * length) /
+            (2.0 * length);
+        double height = std::sqrt(std::max(upperLength * upperLength - along * along, 0.0));
+        Vec2 base = add(shoulder, mul(unit, along));
+        Vec2 perpendicular{-unit.y, unit.x};
+        Vec2 first = add(base, mul(perpendicular, height));
+        Vec2 second = sub(base, mul(perpendicular, height));
+        return norm(sub(first, reference)) <= norm(sub(second, reference)) ? first : second;
+    };
+    struct Candidate {
+        PersonPose const* pose{};
+        std::string source;
+        std::array<Vec2, 3> limb{};
+        std::array<Vec2, 21> hand{};
+        double limbScore{};
+        double handScore{};
+        double handSpan{};
+        double wristLink{};
+        bool limbValid{};
+        bool handValid{};
+    };
+    auto buildCandidate = [&](PersonPose const& pose, int side, std::string source) {
+        int shoulder = side == 0 ? 5 : 6;
+        int elbow = side == 0 ? 7 : 8;
+        int wrist = side == 0 ? 9 : 10;
+        int handStart = side == 0 ? 91 : 112;
+        Candidate candidate;
+        candidate.pose = &pose;
+        candidate.source = std::move(source);
+        candidate.limb = {normalized(pose.keypoints[shoulder], pose.box),
+                          normalized(pose.keypoints[elbow], pose.box),
+                          normalized(pose.keypoints[wrist], pose.box)};
+        double minX = pose.keypoints[handStart].x, maxX = minX;
+        double minY = pose.keypoints[handStart].y, maxY = minY;
+        for (int index = 0; index < 21; ++index) {
+            candidate.hand[index] = normalized(pose.keypoints[handStart + index], pose.box);
+            minX = std::min<double>(minX, pose.keypoints[handStart + index].x);
+            maxX = std::max<double>(maxX, pose.keypoints[handStart + index].x);
+            minY = std::min<double>(minY, pose.keypoints[handStart + index].y);
+            maxY = std::max<double>(maxY, pose.keypoints[handStart + index].y);
+            candidate.handScore += pose.keypoints[handStart + index].score;
+        }
+        candidate.handScore /= 21.0;
+        candidate.limbScore = (pose.keypoints[shoulder].score + pose.keypoints[elbow].score + pose.keypoints[wrist].score) / 3.0;
+        double upper = norm(sub(candidate.limb[0], candidate.limb[1]));
+        double forearm = norm(sub(candidate.limb[1], candidate.limb[2]));
+        candidate.handSpan = std::hypot(maxX - minX, maxY - minY) / boxDiagonal(pose.box);
+        candidate.wristLink = distance(
+            pose.keypoints[wrist].x, pose.keypoints[wrist].y,
+            pose.keypoints[handStart].x, pose.keypoints[handStart].y) / boxDiagonal(pose.box);
+        candidate.limbValid = upper >= 0.025 && upper <= 0.8 && forearm >= 0.025 && forearm <= 0.8;
+        candidate.handValid = candidate.handSpan >= 0.008 && candidate.handSpan <= 0.35 && candidate.wristLink <= 0.12;
+        return candidate;
+    };
+    auto robustDistance = [&](auto const& points, auto const& target, bool available) {
+        if (!available) return 0.0;
+        std::vector<double> distances;
+        for (std::size_t index = 0; index < points.size(); ++index) {
+            distances.push_back(norm(sub(points[index], target[index])));
+        }
+        return median(std::move(distances));
+    };
+
+    std::vector<FusedPerson> fused;
+    fused.reserve(rawPeople.size());
+    for (std::size_t personIndex = 0; personIndex < rawPeople.size(); ++personIndex) {
+        auto const& raw = rawPeople[personIndex];
+        auto const& isolated = personIndex < isolatedPeople.size() ? isolatedPeople[personIndex] : raw;
+        auto& state = states_[raw.box.trackId];
+        FusedPerson output;
+        output.pose = raw;
+        for (int side = 0; side < 2; ++side) {
+            int shoulder = side == 0 ? 5 : 6;
+            int elbow = side == 0 ? 7 : 8;
+            int wrist = side == 0 ? 9 : 10;
+            int handStart = side == 0 ? 91 : 112;
+            auto& track = state.fusion[side];
+            std::array<Candidate, 2> candidates{
+                buildCandidate(raw, side, "raw"), buildCandidate(isolated, side, "isolated")};
+            int limbChoice = 0;
+            double bestLimbScore = -1e30;
+            for (int choice = 0; choice < 2; ++choice) {
+                std::array<Vec2, 3> target{};
+                if (track.hasLimbPrevious) {
+                    for (int i = 0; i < 3; ++i) target[i] = add(track.limbPrevious[i], track.hasLimbVelocity ? mul(track.limbVelocity[i], 0.65) : Vec2{});
+                }
+                double temporal = robustDistance(candidates[choice].limb, target, track.hasLimbPrevious);
+                double score = candidates[choice].limbScore + (candidates[choice].limbValid ? 1.2 : -5.0)
+                    + (track.limbSource == candidates[choice].source ? 0.35 : 0.0) - 7.0 * temporal;
+                if (score > bestLimbScore) { bestLimbScore = score; limbChoice = choice; }
+            }
+            int handChoice = 0;
+            double bestHandScore = -1e30;
+            double chosenTemporal = 0.0;
+            double chosenLink = 1.0;
+            for (int choice = 0; choice < 2; ++choice) {
+                std::array<Vec2, 21> target{};
+                if (track.hasHandPrevious) {
+                    for (int i = 0; i < 21; ++i) target[i] = add(track.handPrevious[i], track.hasHandVelocity ? mul(track.handVelocity[i], 0.65) : Vec2{});
+                }
+                double temporal = robustDistance(candidates[choice].hand, target, track.hasHandPrevious);
+                auto const& limbPose = *candidates[limbChoice].pose;
+                double link = distance(
+                    candidates[choice].pose->keypoints[handStart].x,
+                    candidates[choice].pose->keypoints[handStart].y,
+                    limbPose.keypoints[wrist].x, limbPose.keypoints[wrist].y) / boxDiagonal(raw.box);
+                double linkPenalty = 1.5 * std::min(std::max(link - 0.12, 0.0), 0.2);
+                double score = candidates[choice].handScore + (candidates[choice].handValid ? 3.0 : -7.0)
+                    + (track.handSource == candidates[choice].source ? 0.45 : 0.0)
+                    + (candidates[choice].source == candidates[limbChoice].source ? 0.15 : 0.0)
+                    - linkPenalty - 8.0 * std::min(temporal, 0.15);
+                if (score > bestHandScore) {
+                    bestHandScore = score;
+                    handChoice = choice;
+                    chosenTemporal = temporal;
+                    chosenLink = link;
+                }
+            }
+            double endpointAdvantage = candidates[handChoice].handScore - candidates[1 - handChoice].handScore;
+            if (handChoice != limbChoice && candidates[handChoice].handValid && endpointAdvantage >= 0.6) {
+                limbChoice = handChoice;
+                // fuse_person() rewrites this link after its visible-endpoint
+                // override.  Temporal recovery consumes the rewritten value.
+                chosenLink = candidates[handChoice].wristLink;
+            }
+            auto const& limb = candidates[limbChoice];
+            auto const& hand = candidates[handChoice];
+            for (int index : {shoulder, elbow, wrist}) output.pose.keypoints[index] = limb.pose->keypoints[index];
+            for (int index = 0; index < 21; ++index) output.pose.keypoints[handStart + index] = hand.pose->keypoints[handStart + index];
+            std::vector<double> limbDistances, handDistances;
+            for (int index = 0; index < 3; ++index) limbDistances.push_back(norm(sub(candidates[0].limb[index], candidates[1].limb[index])));
+            for (int index = 0; index < 21; ++index) handDistances.push_back(norm(sub(candidates[0].hand[index], candidates[1].hand[index])));
+            double handDisagreement = percentile90(handDistances);
+            double scoreAdvantage = candidates[0].handScore - candidates[1].handScore;
+            auto& meta = output.meta[side];
+            meta.limbConsensus = median(limbDistances) <= 0.05;
+            meta.handValid = hand.handValid;
+            meta.distanceEligible = hand.handValid && (
+                handDisagreement <= 0.06 ||
+                (hand.source == "raw" && scoreAdvantage >= 0.6 && chosenTemporal <= 0.08));
+            meta.alarmEligible = meta.distanceEligible || (hand.handValid && hand.source == "raw" && scoreAdvantage >= 0.6);
+            meta.handScore = std::nearbyint(hand.handScore * 1e4) / 1e4;
+            meta.linkToSelectedLimb = std::nearbyint(chosenLink * 1e5) / 1e5;
+            meta.handSource = hand.source;
+            if (!hand.handValid) meta.handState = "rejected";
+            else if (handDisagreement <= 0.06) meta.handState = "consensus";
+            else if (hand.source == "raw" && scoreAdvantage >= 0.6 && chosenTemporal <= 0.08) meta.handState = "visible_endpoint";
+            else if (hand.source == "isolated") meta.handState = "inferred";
+            else meta.handState = "conflict";
+            if (track.hasLimbPrevious) {
+                for (int i = 0; i < 3; ++i) {
+                    auto delta = sub(limb.limb[i], track.limbPrevious[i]);
+                    track.limbVelocity[i] = track.hasLimbVelocity ? add(mul(track.limbVelocity[i], 0.65), mul(delta, 0.35)) : delta;
+                }
+                track.hasLimbVelocity = true;
+            }
+            if (hand.handValid) {
+                if (track.hasHandPrevious) {
+                    for (int i = 0; i < 21; ++i) {
+                        auto delta = sub(hand.hand[i], track.handPrevious[i]);
+                        track.handVelocity[i] = track.hasHandVelocity ? add(mul(track.handVelocity[i], 0.65), mul(delta, 0.35)) : delta;
+                    }
+                    track.hasHandVelocity = true;
+                }
+                track.handPrevious = hand.hand;
+                track.hasHandPrevious = true;
+                track.handSource = hand.source;
+            }
+            track.limbPrevious = limb.limb;
+            track.hasLimbPrevious = true;
+            track.limbSource = limb.source;
+        }
+        // Same-person duplicate suppression from the reference fusion.
+        auto handCenter = [&](int side) {
+            int start = side == 0 ? 91 : 112;
+            Vec2 center{};
+            for (int i = 0; i < 21; ++i) center = add(center, vec(output.pose.keypoints[start + i]));
+            return mul(center, 1.0 / 21.0);
+        };
+        if (norm(sub(handCenter(0), handCenter(1))) / boxDiagonal(raw.box) < 0.08) {
+            int weaker = output.meta[0].handScore <= output.meta[1].handScore ? 0 : 1;
+            output.meta[weaker].handState = "duplicate";
+            output.meta[weaker].distanceEligible = false;
+            output.meta[weaker].alarmEligible = false;
+        }
+        // Motion-adaptive EMA in person-box coordinates.
+        std::array<Vec2, 133> now{};
+        for (int i = 0; i < 133; ++i) now[i] = normalized(output.pose.keypoints[i], raw.box);
+        if (!state.hasSmoothed) { state.smoothed = now; state.hasSmoothed = true; }
+        else {
+            for (int i = 0; i < 133; ++i) {
+                double displacement = norm(sub(now[i], state.smoothed[i]));
+                double alpha = 0.32 + 0.58 * std::clamp(displacement / 0.08, 0.0, 1.0);
+                state.smoothed[i] = add(mul(now[i], alpha), mul(state.smoothed[i], 1.0 - alpha));
+            }
+        }
+        fused.push_back(std::move(output));
+    }
+
+    // Exact build_global_hands() evidence: selected fusion output, stable
+    // score ordering, then the same cross-person duplicate collapse.
+    std::vector<GlobalHand> globalCandidates;
+    for (std::size_t person = 0; person < fused.size(); ++person) {
+        for (int side = 0; side < 2; ++side) {
+            auto const& meta = fused[person].meta[side];
+            if (!meta.alarmEligible) continue;
+            int const start = side == 0 ? 91 : 112;
+            GlobalHand hand;
+            hand.personId = static_cast<int>(person) + 1;
+            hand.trackId = fused[person].pose.box.trackId;
+            hand.left = side == 0;
+            hand.source = meta.handSource;
+            hand.state = meta.handState;
+            hand.distanceEligible = meta.distanceEligible;
+            hand.alarmEligible = meta.alarmEligible;
+            hand.score = static_cast<float>(meta.handScore);
+            float minX = fused[person].pose.keypoints[start].x;
+            float maxX = minX;
+            float minY = fused[person].pose.keypoints[start].y;
+            float maxY = minY;
+            for (int point = 0; point < 21; ++point) {
+                hand.keypoints[point] = fused[person].pose.keypoints[start + point];
+                minX = std::min(minX, hand.keypoints[point].x);
+                maxX = std::max(maxX, hand.keypoints[point].x);
+                minY = std::min(minY, hand.keypoints[point].y);
+                maxY = std::max(maxY, hand.keypoints[point].y);
+            }
+            hand.spanPx = std::hypot(maxX - minX, maxY - minY);
+            globalCandidates.push_back(std::move(hand));
+        }
+    }
+    std::stable_sort(globalCandidates.begin(), globalCandidates.end(), [](auto const& first, auto const& second) {
+        return first.score > second.score;
+    });
+    std::vector<GlobalHand> globalHands;
+    for (auto& candidate : globalCandidates) {
+        Vec2 candidateCenter{};
+        for (auto const& point : candidate.keypoints) candidateCenter = add(candidateCenter, vec(point));
+        candidateCenter = mul(candidateCenter, 1.0 / 21.0);
+        auto duplicate = globalHands.end();
+        for (auto iterator = globalHands.begin(); iterator != globalHands.end(); ++iterator) {
+            Vec2 existingCenter{};
+            for (auto const& point : iterator->keypoints) existingCenter = add(existingCenter, vec(point));
+            existingCenter = mul(existingCenter, 1.0 / 21.0);
+            double const threshold = std::clamp(
+                0.45 * std::max<double>(candidate.spanPx, iterator->spanPx), 18.0, 70.0);
+            if (norm(sub(candidateCenter, existingCenter)) <= threshold) {
+                duplicate = iterator;
+                break;
+            }
+        }
+        if (duplicate == globalHands.end()) globalHands.push_back(std::move(candidate));
+        else duplicate->aliases.emplace_back(candidate.personId, candidate.left);
+    }
+
+    // Build the model measurements, then collapse cross-person duplicates.
+    std::vector<std::array<std::optional<Measurement>, 2>> measurements(fused.size());
+    for (std::size_t person = 0; person < fused.size(); ++person) {
+        for (int side = 0; side < 2; ++side) {
+            if (!fused[person].meta[side].alarmEligible) continue;
+            Measurement measurement;
+            int start = side == 0 ? 91 : 112;
+            for (int i = 0; i < 21; ++i) measurement.hand[i] = vec(fused[person].pose.keypoints[start + i]);
+            measurement.precise = fused[person].meta[side].distanceEligible;
+            measurement.provisional = !measurement.precise;
+            measurement.quality = fused[person].meta[side].handScore;
+            measurements[person][side] = measurement;
+        }
+    }
+    auto handSpan = [&](Measurement const& measurement) {
+        double minX = measurement.hand[0].x, maxX = minX, minY = measurement.hand[0].y, maxY = minY;
+        for (auto point : measurement.hand) { minX = std::min(minX, point.x); maxX = std::max(maxX, point.x); minY = std::min(minY, point.y); maxY = std::max(maxY, point.y); }
+        return std::hypot(maxX - minX, maxY - minY);
+    };
+    auto handCenter = [&](Measurement const& measurement) {
+        Vec2 center{}; for (auto point : measurement.hand) center = add(center, point); return mul(center, 1.0 / 21.0);
+    };
+    for (std::size_t firstPerson = 0; firstPerson < fused.size(); ++firstPerson) {
+        for (int firstSide = 0; firstSide < 2; ++firstSide) {
+            if (!measurements[firstPerson][firstSide]) continue;
+            for (std::size_t secondPerson = firstPerson + 1; secondPerson < fused.size(); ++secondPerson) {
+                for (int secondSide = 0; secondSide < 2; ++secondSide) {
+                    if (!measurements[firstPerson][firstSide] || !measurements[secondPerson][secondSide]) continue;
+                    auto const& first = *measurements[firstPerson][firstSide];
+                    auto const& second = *measurements[secondPerson][secondSide];
+                    double firstSpan = handSpan(first), secondSpan = handSpan(second);
+                    if (norm(sub(handCenter(first), handCenter(second))) > std::max(70.0, 0.5 * std::max(firstSpan, secondSpan))) continue;
+                    double firstArea = std::max((fused[firstPerson].pose.box.x2 - fused[firstPerson].pose.box.x1) * (fused[firstPerson].pose.box.y2 - fused[firstPerson].pose.box.y1), 1.0F);
+                    double secondArea = std::max((fused[secondPerson].pose.box.x2 - fused[secondPerson].pose.box.x1) * (fused[secondPerson].pose.box.y2 - fused[secondPerson].pose.box.y1), 1.0F);
+                    if (firstArea >= 2.5 * secondArea &&
+                        (firstSpan >= 1.45 * secondSpan || first.quality >= second.quality + 0.6)) {
+                        measurements[secondPerson][secondSide].reset();
+                        states_[fused[secondPerson].pose.box.trackId].endpoints[secondSide] = {};
+                    } else if (secondArea >= 2.5 * firstArea &&
+                               (secondSpan >= 1.45 * firstSpan || second.quality >= first.quality + 0.6)) {
+                        measurements[firstPerson][firstSide].reset();
+                        states_[fused[firstPerson].pose.box.trackId].endpoints[firstSide] = {};
+                    }
+                }
+            }
+        }
+    }
+
+    // Match resolve_global_hand_assignments(): a large foreground track may reclaim
+    // a nearby hand that RTMW emitted only from a smaller background crop.
+    struct Reassignment {
+        double normalizedDistance{};
+        std::size_t targetPerson{};
+        int targetSide{};
+        std::size_t sourcePerson{};
+        int sourceSide{};
+        Vec2 sourceCenter{};
+    };
+    std::vector<Reassignment> reassignments;
+    for (std::size_t targetPerson = 0; targetPerson < fused.size(); ++targetPerson) {
+        for (int targetSide = 0; targetSide < 2; ++targetSide) {
+            if (measurements[targetPerson][targetSide]) continue;
+            auto& targetEndpoint = states_[fused[targetPerson].pose.box.trackId].endpoints[targetSide];
+            if (!targetEndpoint.hasWrist || !targetEndpoint.hasHandShape || targetEndpoint.missedFrames > 30) continue;
+            double targetArea = std::max(
+                (fused[targetPerson].pose.box.x2 - fused[targetPerson].pose.box.x1) *
+                (fused[targetPerson].pose.box.y2 - fused[targetPerson].pose.box.y1), 1.0F);
+            Vec2 shapeCenter{};
+            for (auto point : targetEndpoint.handShape) shapeCenter = add(shapeCenter, point);
+            shapeCenter = mul(shapeCenter, 1.0 / 21.0);
+            Vec2 predictedCenter = add(targetEndpoint.wrist, shapeCenter);
+            Measurement targetShape;
+            for (int point = 0; point < 21; ++point) targetShape.hand[point] = targetEndpoint.handShape[point];
+            double targetSpan = handSpan(targetShape);
+            double gate = 0.20 * boxDiagonal(fused[targetPerson].pose.box);
+            for (std::size_t sourcePerson = 0; sourcePerson < fused.size(); ++sourcePerson) {
+                if (sourcePerson == targetPerson) continue;
+                for (int sourceSide = 0; sourceSide < 2; ++sourceSide) {
+                    auto const& source = measurements[sourcePerson][sourceSide];
+                    if (!source || source->source != "model") continue;
+                    double sourceArea = std::max(
+                        (fused[sourcePerson].pose.box.x2 - fused[sourcePerson].pose.box.x1) *
+                        (fused[sourcePerson].pose.box.y2 - fused[sourcePerson].pose.box.y1), 1.0F);
+                    if (targetArea < 2.5 * sourceArea) continue;
+                    Vec2 sourceCenter = handCenter(*source);
+                    double sourceSpan = handSpan(*source);
+                    double candidateDistance = norm(sub(sourceCenter, predictedCenter));
+                    if (candidateDistance > gate || targetSpan < 1.45 * std::max(sourceSpan, 1.0)) continue;
+                    reassignments.push_back({
+                        candidateDistance / std::max(gate, 1e-6), targetPerson, targetSide,
+                        sourcePerson, sourceSide, sourceCenter});
+                }
+            }
+        }
+    }
+    std::sort(reassignments.begin(), reassignments.end(), [](auto const& first, auto const& second) {
+        return std::tie(first.normalizedDistance, first.targetPerson, first.targetSide,
+                        first.sourcePerson, first.sourceSide) <
+               std::tie(second.normalizedDistance, second.targetPerson, second.targetSide,
+                        second.sourcePerson, second.sourceSide);
+    });
+    std::set<std::pair<std::size_t, int>> usedTargets, usedSources;
+    for (auto const& item : reassignments) {
+        auto targetKey = std::make_pair(item.targetPerson, item.targetSide);
+        auto sourceKey = std::make_pair(item.sourcePerson, item.sourceSide);
+        if (usedTargets.count(targetKey) || usedSources.count(sourceKey) ||
+            measurements[item.targetPerson][item.targetSide] ||
+            !measurements[item.sourcePerson][item.sourceSide]) continue;
+        auto& targetEndpoint = states_[fused[item.targetPerson].pose.box.trackId].endpoints[item.targetSide];
+        Vec2 shapeCenter{};
+        for (auto point : targetEndpoint.handShape) shapeCenter = add(shapeCenter, point);
+        shapeCenter = mul(shapeCenter, 1.0 / 21.0);
+        Vec2 reassignedWrist = sub(item.sourceCenter, shapeCenter);
+        Measurement reassigned;
+        for (int point = 0; point < 21; ++point) {
+            reassigned.hand[point] = add(reassignedWrist, targetEndpoint.handShape[point]);
+        }
+        reassigned.provisional = true;
+        reassigned.source = "global_reassignment";
+        reassigned.quality = measurements[item.sourcePerson][item.sourceSide]->quality;
+        measurements[item.targetPerson][item.targetSide] = reassigned;
+        measurements[item.sourcePerson][item.sourceSide].reset();
+        auto& sourceEndpoint = states_[fused[item.sourcePerson].pose.box.trackId].endpoints[item.sourceSide];
+        sourceEndpoint = {};
+        usedTargets.insert(targetKey);
+        usedSources.insert(sourceKey);
+    }
+
+    PoseFrameResult result;
+    result.globalHands = std::move(globalHands);
+    result.selectedPeople.reserve(fused.size());
+    result.smoothedPeople.reserve(fused.size());
+    result.v2People.reserve(fused.size());
+    for (auto const& person : fused) {
+        std::array<FusionMeta,2> meta{};
+        for (int side=0;side<2;++side) {
+            auto const& m=person.meta[side];
+            meta[side]={m.limbConsensus,m.handValid,m.distanceEligible,m.alarmEligible,m.handScore,m.handSource,m.handState};
+        }
+        result.metadata.push_back(meta);
+        result.selectedPeople.push_back(person.pose);
+        PersonPose smoothed = person.pose;
+        auto const& state = states_[person.pose.box.trackId];
+        for (int keypoint = 0; keypoint < 133; ++keypoint) {
+            Vec2 point = denormalized(state.smoothed[keypoint], person.pose.box);
+            smoothed.keypoints[keypoint].x = static_cast<float>(point.x);
+            smoothed.keypoints[keypoint].y = static_cast<float>(point.y);
+        }
+        result.smoothedPeople.push_back(smoothed);
+        result.v2People.push_back(std::move(smoothed));
+    }
+    for (std::size_t person = 0; person < fused.size(); ++person) {
+        auto& state = states_[fused[person].pose.box.trackId];
+        int frameGap = state.lastFrame < 0 ? 0 : std::max(frameIndex - state.lastFrame - 1, 0);
+        for (auto& endpoint : state.endpoints) endpoint.missedFrames += frameGap;
+        state.lastFrame = frameIndex;
+        for (int side = 0; side < 2; ++side) {
+            if (!fused[person].meta[side].limbConsensus || !fused[person].meta[side].distanceEligible) continue;
+            int shoulderIndex = side == 0 ? 5 : 6;
+            int elbowIndex = side == 0 ? 7 : 8;
+            int wristIndex = side == 0 ? 9 : 10;
+            Vec2 shoulder = vec(fused[person].pose.keypoints[shoulderIndex]);
+            Vec2 elbow = vec(fused[person].pose.keypoints[elbowIndex]);
+            Vec2 wrist = vec(fused[person].pose.keypoints[wristIndex]);
+            double diagonal = boxDiagonal(fused[person].pose.box);
+            double upper = norm(sub(shoulder, elbow)) / diagonal;
+            double forearm = norm(sub(elbow, wrist)) / diagonal;
+            if (upper >= 0.04 && upper <= 0.65 && forearm >= 0.04 && forearm <= 0.65) {
+                auto& samples = state.boneSamples[side];
+                samples.push_back({upper, forearm});
+                while (samples.size() > 120) samples.pop_front();
+            }
+        }
+        for (int side = 0; side < 2; ++side) {
+            int shoulderIndex = side == 0 ? 5 : 6;
+            int elbowIndex = side == 0 ? 7 : 8;
+            int wristIndex = side == 0 ? 9 : 10;
+            int handStart = side == 0 ? 91 : 112;
+            auto& endpoint = state.endpoints[side];
+            auto measurement = measurements[person][side];
+            // Reference temporal recovery gate.
+            bool const recoverableHandState = fused[person].meta[side].handState == "conflict" ||
+                fused[person].meta[side].handState == "inferred";
+            if (!measurement && endpoint.hasWrist && endpoint.hasHandShape && endpoint.missedFrames <= 30 &&
+                recoverableHandState && fused[person].meta[side].handValid &&
+                fused[person].meta[side].handScore >= 2.0 &&
+                fused[person].meta[side].linkToSelectedLimb <= 0.08) {
+                Vec2 measured = vec(fused[person].pose.keypoints[handStart]);
+                Vec2 prediction = add(endpoint.wrist, endpoint.velocity);
+                double gate = std::min(0.20, 0.05 + 0.008 * std::max(endpoint.missedFrames, 1)) * boxDiagonal(fused[person].pose.box);
+                if (norm(sub(measured, prediction)) <= gate) {
+                    Measurement recovered;
+                    for (int i = 0; i < 21; ++i) recovered.hand[i] = vec(fused[person].pose.keypoints[handStart + i]);
+                    recovered.provisional = true; recovered.quality = fused[person].meta[side].handScore; recovered.source = "temporal_recovery";
+                    measurement = recovered;
+                }
+            }
+            LimbState outputState;
+            std::array<Vec2, 21> outputHand{};
+            bool hasOutput = false;
+            bool returningAfterLong = measurement && endpoint.missedFrames > 12;
+            if (measurement) {
+                Vec2 measuredWrist = measurement->hand[0];
+                std::array<Vec2, 21> measuredShape{};
+                for (int i = 0; i < 21; ++i) measuredShape[i] = sub(measurement->hand[i], measuredWrist);
+                if (!endpoint.hasWrist || endpoint.missedFrames > 12) {
+                    endpoint.wrist = measuredWrist; endpoint.velocity = {}; endpoint.handShape = measuredShape;
+                } else {
+                    Vec2 prediction = add(endpoint.wrist, endpoint.velocity);
+                    Vec2 residual = sub(measuredWrist, prediction);
+                    double alpha = measurement->precise ? 0.78 : 0.88;
+                    double beta = measurement->precise ? 0.18 : 0.24;
+                    endpoint.wrist = add(prediction, mul(residual, alpha));
+                    endpoint.velocity = add(mul(endpoint.velocity, 0.82), mul(residual, beta));
+                    double shapeAlpha = measurement->precise ? 0.55 : 0.70;
+                    for (int i = 0; i < 21; ++i) endpoint.handShape[i] = add(mul(measuredShape[i], shapeAlpha), mul(endpoint.handShape[i], 1.0 - shapeAlpha));
+                }
+                endpoint.hasWrist = endpoint.hasHandShape = true;
+                endpoint.lastObservedFrame = frameIndex;
+                endpoint.missedFrames = 0;
+                outputState = measurement->source == "global_reassignment" ? LimbState::Reassigned : measurement->source == "temporal_recovery" ? LimbState::Recovered : measurement->precise ? LimbState::Observed : measurement->provisional ? LimbState::Provisional : LimbState::Interpolated;
+                for (int i = 0; i < 21; ++i) outputHand[i] = add(endpoint.wrist, endpoint.handShape[i]);
+                hasOutput = true;
+            } else {
+                ++endpoint.missedFrames;
+                if (!endpoint.hasWrist || endpoint.missedFrames > 30) outputState = LimbState::ReachableLong;
+                else if (endpoint.missedFrames <= 12) {
+                    endpoint.wrist = add(endpoint.wrist, endpoint.velocity);
+                    endpoint.velocity = mul(endpoint.velocity, 0.86);
+                    for (int i = 0; i < 21; ++i) outputHand[i] = add(endpoint.wrist, endpoint.handShape[i]);
+                    hasOutput = true; outputState = LimbState::PredictedShort;
+                } else outputState = LimbState::ReachableMedium;
+            }
+            Vec2 smoothedShoulder = denormalized(state.smoothed[shoulderIndex], fused[person].pose.box);
+            Vec2 smoothedElbow = denormalized(state.smoothed[elbowIndex], fused[person].pose.box);
+            Keypoint shoulder = fused[person].pose.keypoints[shoulderIndex];
+            shoulder.x = static_cast<float>(smoothedShoulder.x);
+            shoulder.y = static_cast<float>(smoothedShoulder.y);
+            Keypoint elbow = fused[person].pose.keypoints[elbowIndex];
+            elbow.x = static_cast<float>(smoothedElbow.x);
+            elbow.y = static_cast<float>(smoothedElbow.y);
+            Keypoint wrist = result.smoothedPeople[person].keypoints[wristIndex];
+            std::array<Keypoint, 21> hand{};
+            for (int i = 0; i < 21; ++i) {
+                hand[i] = result.smoothedPeople[person].keypoints[handStart + i];
+            }
+            if (hasOutput) {
+                auto& ownSamples = state.boneSamples[side];
+                auto& oppositeSamples = state.boneSamples[1 - side];
+                auto const& samples = ownSamples.size() < 10 && oppositeSamples.size() >= 10
+                    ? oppositeSamples : ownSamples;
+                double upperRatio = 0.24, forearmRatio = 0.22;
+                if (!samples.empty()) {
+                    std::vector<double> upperValues, forearmValues;
+                    upperValues.reserve(samples.size());
+                    forearmValues.reserve(samples.size());
+                    for (auto const& sample : samples) {
+                        upperValues.push_back(sample[0]);
+                        forearmValues.push_back(sample[1]);
+                    }
+                    upperRatio = median(std::move(upperValues));
+                    forearmRatio = median(std::move(forearmValues));
+                }
+                double diagonal = boxDiagonal(fused[person].pose.box);
+                double upperLength = upperRatio * diagonal;
+                double forearmLength = forearmRatio * diagonal;
+                double totalLength = upperLength + forearmLength;
+                if (totalLength > 0.95 * diagonal) {
+                    double lengthScale = 0.95 * diagonal / totalLength;
+                    upperLength *= lengthScale;
+                    forearmLength *= lengthScale;
+                }
+                Vec2 referenceElbow = returningAfterLong || !endpoint.hasLastElbow
+                    ? vec(fused[person].pose.keypoints[elbowIndex]) : endpoint.lastElbow;
+                if (!fused[person].meta[side].limbConsensus || outputState != LimbState::Observed) {
+                    Vec2 solved = solveTwoBoneElbow(
+                        smoothedShoulder, outputHand[0], upperLength, forearmLength, referenceElbow);
+                    elbow.x = static_cast<float>(solved.x);
+                    elbow.y = static_cast<float>(solved.y);
+                }
+                endpoint.lastElbow = vec(elbow);
+                endpoint.hasLastElbow = true;
+                Vec2 armDirection = sub(outputHand[0], smoothedShoulder);
+                if (norm(armDirection) > 1e-6) {
+                    endpoint.lastDirection = mul(armDirection, 1.0 / norm(armDirection));
+                    endpoint.hasLastDirection = true;
+                }
+                wrist.x = outputHand[0].x; wrist.y = outputHand[0].y;
+                for (int i = 0; i < 21; ++i) {
+                    hand[i].x = outputHand[i].x; hand[i].y = outputHand[i].y;
+                }
+            }
+            auto& v2Pose = result.v2People[person];
+            v2Pose.keypoints[shoulderIndex] = shoulder;
+            v2Pose.keypoints[elbowIndex] = elbow;
+            v2Pose.keypoints[wristIndex] = wrist;
+            if (hasOutput) {
+                for (int point = 0; point < 21; ++point) v2Pose.keypoints[handStart + point] = hand[point];
+            }
+            result.limbs.push_back({
+                fused[person].pose.box.trackId, side == 0, outputState,
+                elbow, wrist, hand, fused[person].pose.box,
+                fused[person].meta[side].distanceEligible,
+                fused[person].meta[side].alarmEligible,
+                static_cast<float>(fused[person].meta[side].handScore)});
+        }
+    }
+    return result;
+}
+
+} // namespace fangjia
@@ -0,0 +1,125 @@
+#include "trt_engine.hpp"
+
+#include <algorithm>
+
+namespace fangjia {
+
+TrtEngine::TrtEngine(std::filesystem::path const& enginePath, TrtLogger& logger)
+    : enginePath_(enginePath) {
+    auto bytes = readBinary(enginePath);
+    runtime_.reset(nvinfer1::createInferRuntime(logger));
+    if (!runtime_) throw std::runtime_error("createInferRuntime failed");
+    engine_.reset(runtime_->deserializeCudaEngine(bytes.data(), bytes.size()));
+    if (!engine_) throw std::runtime_error("Cannot deserialize " + enginePath.string());
+    context_.reset(engine_->createExecutionContext());
+    if (!context_) throw std::runtime_error("createExecutionContext failed");
+    for (int index = 0; index < engine_->getNbIOTensors(); ++index) {
+        buffers_.emplace(engine_->getIOTensorName(index), Buffer{});
+    }
+}
+
+TrtEngine::~TrtEngine() {
+    for (auto& [name, item] : buffers_) {
+        (void) name;
+        if (item.owned) cudaFree(item.owned);
+    }
+}
+
+void TrtEngine::setInputShape(std::string const& name, nvinfer1::Dims const& shape) {
+    if (!context_->setInputShape(name.c_str(), shape)) {
+        throw std::runtime_error("TensorRT rejected input shape for " + name);
+    }
+}
+
+void TrtEngine::prepare() {
+    for (int index = 0; index < engine_->getNbIOTensors(); ++index) {
+        std::string const name = engine_->getIOTensorName(index);
+        auto const dims = context_->getTensorShape(name.c_str());
+        for (int axis = 0; axis < dims.nbDims; ++axis) {
+            if (dims.d[axis] < 0) {
+                throw std::runtime_error(
+                    "TensorRT dimensions are incomplete for " + enginePath_.string() +
+                    ":" + name);
+            }
+        }
+        auto& item = buffer(name);
+        item.required = volume(dims) * dataTypeSize(engine_->getTensorDataType(name.c_str()));
+        if (item.external) {
+            if (item.capacity < item.required) {
+                throw std::runtime_error("External buffer is too small for " + name);
+            }
+            continue;
+        }
+        if (item.capacity < item.required) {
+            if (item.owned) CUDA_CHECK(cudaFree(item.owned));
+            CUDA_CHECK(cudaMalloc(&item.owned, item.required));
+            item.capacity = item.required;
+        }
+        item.active = item.owned;
+    }
+}
+
+void TrtEngine::bindExternal(std::string const& name, void* pointer, std::size_t bytesValue) {
+    auto& item = buffer(name);
+    item.active = pointer;
+    item.capacity = bytesValue;
+    item.external = true;
+}
+
+void TrtEngine::useOwned(std::string const& name) {
+    auto& item = buffer(name);
+    item.external = false;
+    item.active = item.owned;
+}
+
+bool TrtEngine::enqueue(cudaStream_t stream) {
+    for (int index = 0; index < engine_->getNbIOTensors(); ++index) {
+        std::string const name = engine_->getIOTensorName(index);
+        auto& item = buffer(name);
+        if (!item.active) throw std::runtime_error("No device buffer for " + name);
+        if (!context_->setTensorAddress(name.c_str(), item.active)) {
+            throw std::runtime_error("setTensorAddress failed for " + name);
+        }
+    }
+    return context_->enqueueV3(stream);
+}
+
+void* TrtEngine::device(std::string const& name) { return buffer(name).active; }
+std::size_t TrtEngine::bytes(std::string const& name) const { return buffer(name).required; }
+nvinfer1::Dims TrtEngine::shape(std::string const& name) const {
+    return context_->getTensorShape(name.c_str());
+}
+nvinfer1::DataType TrtEngine::dataType(std::string const& name) const {
+    return engine_->getTensorDataType(name.c_str());
+}
+
+std::vector<std::string> TrtEngine::inputs() const {
+    std::vector<std::string> result;
+    for (int index = 0; index < engine_->getNbIOTensors(); ++index) {
+        char const* name = engine_->getIOTensorName(index);
+        if (engine_->getTensorIOMode(name) == nvinfer1::TensorIOMode::kINPUT) result.emplace_back(name);
+    }
+    return result;
+}
+
+std::vector<std::string> TrtEngine::outputs() const {
+    std::vector<std::string> result;
+    for (int index = 0; index < engine_->getNbIOTensors(); ++index) {
+        char const* name = engine_->getIOTensorName(index);
+        if (engine_->getTensorIOMode(name) == nvinfer1::TensorIOMode::kOUTPUT) result.emplace_back(name);
+    }
+    return result;
+}
+
+TrtEngine::Buffer& TrtEngine::buffer(std::string const& name) {
+    auto iterator = buffers_.find(name);
+    if (iterator == buffers_.end()) throw std::runtime_error("Unknown tensor " + name);
+    return iterator->second;
+}
+TrtEngine::Buffer const& TrtEngine::buffer(std::string const& name) const {
+    auto iterator = buffers_.find(name);
+    if (iterator == buffers_.end()) throw std::runtime_error("Unknown tensor " + name);
+    return iterator->second;
+}
+
+} // namespace fangjia
@@ -0,0 +1,48 @@
+#include "uyvy_converter.hpp"
+#include "gpu_preprocess.hpp"
+#include <cstring>
+
+namespace cms {
+__global__ void swapRedBlue(unsigned char* image,int count) {
+    int i=blockIdx.x*blockDim.x+threadIdx.x;
+    if(i>=count)return;
+    unsigned char r=image[3*i];image[3*i]=image[3*i+2];image[3*i+2]=r;
+}
+struct UyvyConverter::Impl {
+    int width,height;
+    cudaStream_t stream{};
+    unsigned char *uyvy{},*bgr{},*host{};
+    Impl(int w,int h):width(w),height(h) {
+        try {
+            CUDA_CHECK(cudaStreamCreateWithFlags(&stream,cudaStreamNonBlocking));
+            CUDA_CHECK(cudaMalloc(&uyvy,std::size_t(w)*h*2));
+            CUDA_CHECK(cudaMalloc(&bgr,std::size_t(w)*h*3));
+            CUDA_CHECK(cudaMallocHost(&host,std::size_t(w)*h*3));
+        } catch(...) {release();throw;}
+    }
+    ~Impl(){release();}
+    void release() noexcept {
+        if(stream)cudaStreamSynchronize(stream);
+        if(uyvy)cudaFree(uyvy);
+        if(bgr)cudaFree(bgr);
+        if(host)cudaFreeHost(host);
+        if(stream)cudaStreamDestroy(stream);
+    }
+};
+UyvyConverter::UyvyConverter(int w,int h):impl_(std::make_unique<Impl>(w,h)){}
+UyvyConverter::~UyvyConverter()=default;
+std::shared_ptr<const BgrImage> UyvyConverter::convert(void const* data,std::size_t stride) {
+    auto& s=*impl_;
+    CUDA_CHECK(cudaMemcpy2DAsync(s.uyvy,s.width*2,data,stride,s.width*2,s.height,cudaMemcpyHostToDevice,s.stream));
+    nppCheck(nppiCbYCr422ToRGB_8u_C2C3R_Ctx(s.uyvy,s.width*2,s.bgr,s.width*3,{s.width,s.height},nppContext(s.stream)));
+    int pixels=s.width*s.height;
+    swapRedBlue<<<(pixels+255)/256,256,0,s.stream>>>(s.bgr,pixels);
+    CUDA_CHECK(cudaGetLastError());
+    CUDA_CHECK(cudaMemcpyAsync(s.host,s.bgr,std::size_t(pixels)*3,cudaMemcpyDeviceToHost,s.stream));
+    CUDA_CHECK(cudaStreamSynchronize(s.stream));
+    auto image=std::make_shared<BgrImage>();
+    image->width=s.width;image->height=s.height;image->stride=s.width*3;
+    image->pixels.assign(s.host,s.host+std::size_t(pixels)*3);
+    return image;
+}
+}
@@ -0,0 +1,97 @@
+#include "visualizer.hpp"
+#include "camera_pipeline.hpp"
+#include <opencv2/highgui.hpp>
+#include <opencv2/highgui/highgui_c.h>
+#include <opencv2/imgproc.hpp>
+#include <algorithm>
+#include <cmath>
+#include <iomanip>
+#include <sstream>
+#include <stdexcept>
+
+namespace cms {
+namespace {
+constexpr char windowName[]="CMS Camera Risk";
+cv::Point point(PosePoint const& p){return {int(std::nearbyint(p.x)),int(std::nearbyint(p.y))};}
+cv::Scalar color(std::string const& level) {
+    if(level=="danger")return {40,40,255};
+    if(level=="warning")return {0,160,255};
+    if(level=="watch")return {0,240,240};
+    if(level=="clear")return {80,220,90};
+    return {190,190,190};
+}
+void segment(cv::Mat& image,cv::Point a,cv::Point b,cv::Scalar c,int thickness,bool dashed) {
+    if(!dashed){cv::line(image,a,b,c,thickness,cv::LINE_AA);return;}
+    cv::Point2d d=b-a;double length=cv::norm(d);if(length<1)return;d*=1.0/length;
+    for(double x=0;x<length;x+=14)cv::line(image,cv::Point2d(a)+d*x,cv::Point2d(a)+d*std::min(x+8,length),c,thickness,cv::LINE_AA);
+}
+}
+cv::Mat renderResult(FrameResult const& r) {
+    if(!r.image)throw std::invalid_argument("Result has no frame image");
+    auto const& image=*r.image;
+    if(image.width<=0||image.height<=0||image.stride<image.width*3||image.pixels.size()<std::size_t(image.stride)*image.height)
+        throw std::invalid_argument("Malformed result image");
+    cv::Mat view(image.height,image.width,CV_8UC3,const_cast<std::uint8_t*>(image.pixels.data()),image.stride);
+    auto canvas=view.clone();
+    if(r.door_mask_pixels) {
+        auto const& m=*r.door_mask_pixels;
+        if(m.width!=image.width||m.height!=image.height||m.stride<m.width||m.pixels.size()<std::size_t(m.stride)*m.height)
+            throw std::invalid_argument("Result mask does not match image");
+        cv::Mat mask(m.height,m.width,CV_8U,const_cast<std::uint8_t*>(m.pixels.data()),m.stride);
+        auto tint=canvas.clone();tint.setTo(cv::Scalar(255,180,0),mask);
+        cv::addWeighted(canvas,0.75,tint,0.25,0,canvas);
+        std::vector<std::vector<cv::Point>> contours;cv::findContours(mask.clone(),contours,cv::RETR_EXTERNAL,cv::CHAIN_APPROX_SIMPLE);
+        cv::drawContours(canvas,contours,-1,{255,180,0},2);
+    }
+    for(auto const& person:r.persons) {
+        auto const& b=person.box;
+        cv::rectangle(canvas,{int(b[0]),int(b[1])},{int(b[2]),int(b[3])},{0,220,255},2);
+        std::ostringstream text;text<<"ID "<<person.track_id<<" "<<std::fixed<<std::setprecision(2)<<person.confidence;
+        cv::putText(canvas,text.str(),{int(b[0])+4,std::max(180,int(b[1])+25)},cv::FONT_HERSHEY_SIMPLEX,.8,{0,220,255},2,cv::LINE_AA);
+        for(int s=0;s<2;++s) {
+            auto const& state=person.hands[s];auto const& k=person.selected;
+            cv::Scalar arm=s==0?cv::Scalar(40,40,255):cv::Scalar(30,230,30);
+            for(auto e:std::vector<std::pair<int,int>>{{5+s,7+s},{7+s,9+s}}) {
+                segment(canvas,point(k[e.first]),point(k[e.second]),state.limb_consensus?arm:cv::Scalar(0,140,255),3,!state.limb_consensus);
+                cv::circle(canvas,point(k[e.second]),4,arm,-1,cv::LINE_AA);
+            }
+            if(state.state=="rejected")continue;
+            int start=s==0?91:112;
+            auto hand=state.distance_eligible?(s==0?cv::Scalar(255,80,220):cv::Scalar(255,190,40)):cv::Scalar(0,140,255);
+            for(int f=0;f<5;++f)for(int j=0;j<4;++j){int a=j==0?start:start+f*4+j;int b=start+1+f*4+j;segment(canvas,point(k[a]),point(k[b]),hand,2,!state.distance_eligible);}
+            for(int j=0;j<21;++j)cv::circle(canvas,point(k[start+j]),3,hand,-1,cv::LINE_AA);
+        }
+    }
+    for(auto const& h:r.hand_regions) {
+        std::vector<cv::Point> points,hull;for(auto const& p:h.keypoints)points.push_back(point(p));cv::convexHull(points,hull);
+        auto risk=std::find_if(r.risk.regions.begin(),r.risk.regions.end(),[&](auto const& item){return item.track_id==h.track_id&&item.side==h.side;});
+        auto c=color(risk==r.risk.regions.end()?"unavailable":risk->risk_level);
+        cv::polylines(canvas,std::vector<std::vector<cv::Point>>{hull},true,c,2,cv::LINE_AA);
+    }
+    for(auto const& region:r.risk.regions)if(region.source=="YOLO_BODY_BOX_FALLBACK") {
+        auto it=std::find(r.track_ids.begin(),r.track_ids.end(),region.track_id);
+        if(it!=r.track_ids.end()){auto const& b=r.boxes.at(it-r.track_ids.begin());cv::putText(canvas,"BODY FALLBACK",{int(b[0])+4,int(b[1])+50},cv::FONT_HERSHEY_SIMPLEX,.65,{0,160,255},2,cv::LINE_AA);}
+    }
+    cv::rectangle(canvas,{0,0},{canvas.cols,148},{20,24,30},-1);
+    std::ostringstream headline;headline<<"RISK ";if(r.risk.risk_score)headline<<std::fixed<<std::setprecision(3)<<*r.risk.risk_score;else headline<<"N/A";headline<<"  "<<r.risk.risk_level;
+    cv::putText(canvas,headline.str(),{24,40},cv::FONT_HERSHEY_SIMPLEX,1.0,color(r.risk.risk_level),2,cv::LINE_AA);
+    std::ostringstream status;status<<"camera frame "<<r.frame_index<<" | seq "<<r.capture.driver_sequence<<" | latency "<<std::fixed<<std::setprecision(1)<<r.timing.capture_to_result_ms<<" ms";
+    cv::putText(canvas,status.str(),{24,82},cv::FONT_HERSHEY_SIMPLEX,.72,{235,235,235},2,cv::LINE_AA);
+    std::string calibration=r.door_mask_calibration_verified?"door calibration verified":"door calibration NOT verified (bring-up)";
+    cv::putText(canvas,calibration+" | q/Esc: close preview",{24,123},cv::FONT_HERSHEY_SIMPLEX,.70,{0,190,255},2,cv::LINE_AA);
+    return canvas;
+}
+void runLocalVisualization(CameraPipeline const& pipeline) {
+    struct Window {
+        Window(){cv::namedWindow(windowName,cv::WINDOW_NORMAL);cv::resizeWindow(windowName,1000,800);}
+        ~Window(){try{if(cvGetWindowHandle(windowName))cv::destroyWindow(windowName);}catch(...) {}}
+    } window;
+    FrameResultPtr shown;
+    while(pipeline.running()) {
+        auto result=pipeline.latest();
+        if(result&&result!=shown){cv::imshow(windowName,renderResult(*result));shown=std::move(result);}
+        int key=cv::waitKey(16)&0xff;
+        if(key==27||key=='q'||key=='Q'||cvGetWindowHandle(windowName)==nullptr)return;
+    }
+}
+}
@@ -0,0 +1,21 @@
+MIT License 
+
+Copyright (c) 2013-2022 Niels Lohmann
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.