revert gpr heatmap threshold, keep it for sar3d only

The display threshold from 'add treshhold' (e113ef6) was meant for the
sar3d projections; the GPR heatmap goes back to full-range levels and
the constant moves into sar3d_plot_mixin.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

.gitignore

@@ -3,6 +3,8 @@ __pycache__/
 *.py[codz]
 *$py.class
 build*
+# build* иначе прячет скрипт сборки вендоренного gpr_cpp (его build/ остаётся игнорируемым)
+!data_acq_and_processing/gpr_cpp/build_cmake.sh
 .codex
 .DS_Store
 # C extensions
@@ -86,7 +88,7 @@ target/
 
 # Jupyter Notebook
 .ipynb_checkpoints
-
+202609144
 # IPython
 profile_default/
 ipython_config.py

Makefile

@@ -8,7 +8,43 @@ VISA_LDFLAGS := $(LDFLAGS) -lvisa
 BUILD_DIR := build
 BIN_DIR := $(BUILD_DIR)/bin
 
+# --- gpr_cpp: vendored C++/CUDA 3D SAR library (originally from VNA_AlGo).
+# Built automatically by the $(GPR_CPP_LIB) rule below via its own CMake
+# (build_cmake.sh). CUDA headers and libnvrtc come from the pip nvidia-*
+# packages inside this repo's venv (same detection as build_cmake.sh); when
+# those are absent (e.g. Jetson, where CUDA comes with JetPack) the system
+# toolkit at /usr/local/cuda is used. libcuda is always the system driver.
+GPR_CPP_DIR ?= data_acq_and_processing/gpr_cpp
+GPR_CPP_LIB ?= $(GPR_CPP_DIR)/build/libgpr_core_cpp.a
+# Absolute so the derived -rpath (libnvrtc dir) survives running the binary
+# from outside the repo root.
+GPR_VENV ?= $(CURDIR)/.venv
+GPR_CUDA_INC_DIR ?= $(dir $(firstword $(wildcard $(GPR_VENV)/lib/python*/site-packages/nvidia/*/include/nvrtc.h)))
+# cuda.h can live in a DIFFERENT pip package than nvrtc.h (Jetson: cuda_runtime
+# vs cuda_nvrtc); the unified cu13 package on dev machines has both in one dir.
+GPR_CUDA_RT_INC_DIR ?= $(dir $(firstword $(wildcard $(GPR_VENV)/lib/python*/site-packages/nvidia/*/include/cuda.h)))
+GPR_NVRTC_LIB ?= $(firstword $(wildcard $(GPR_VENV)/lib/python*/site-packages/nvidia/*/lib/libnvrtc.so*))
+# System-CUDA fallback. An empty GPR_CUDA_INC_DIR would also corrupt the compile
+# command: a bare `-I` swallows the next argument.
+ifeq ($(strip $(GPR_CUDA_INC_DIR)),)
+GPR_CUDA_INC_DIR := /usr/local/cuda/include
+endif
+ifeq ($(strip $(GPR_CUDA_RT_INC_DIR)),)
+GPR_CUDA_RT_INC_DIR := /usr/local/cuda/include
+endif
+ifeq ($(strip $(GPR_NVRTC_LIB)),)
+GPR_NVRTC_LIB := $(firstword $(wildcard /usr/local/cuda/lib64/libnvrtc.so*))
+endif
+# Deduplicated CUDA include flags (a single dir on dev machines).
+GPR_CUDA_INC_FLAGS := $(addprefix -I,$(sort $(GPR_CUDA_INC_DIR) $(GPR_CUDA_RT_INC_DIR)))
+GPR_CUDA_LIB_DIR := $(dir $(GPR_NVRTC_LIB))
+# libnvrtc from pip ships only as libnvrtc.so.N (no .so symlink), so link the
+# full path; the rpath makes the binary find it without LD_LIBRARY_PATH.
+DATA_PROCESSOR_LDFLAGS := $(LDFLAGS) $(GPR_CPP_LIB) $(GPR_NVRTC_LIB) -lcuda -Wl,-rpath,$(GPR_CUDA_LIB_DIR)
+
 INCLUDES := \
+	-I$(GPR_CPP_DIR)/include \
+	$(GPR_CUDA_INC_FLAGS) \
 	-Idata_acq_and_processing/common_cpp/ipc/include \
 	-Idata_acq_and_processing/common_cpp/config/include \
 	-Idata_acq_and_processing/third_party \
@@ -55,6 +91,7 @@ PROCESSOR_SOURCES := \
 	data_acq_and_processing/processing/processors/src/bscan_processor.cpp \
 	data_acq_and_processing/processing/processors/src/gpr_processor.cpp \
 	data_acq_and_processing/processing/processors/src/passthrough_processor.cpp \
+	data_acq_and_processing/processing/processors/src/sar3d_processor.cpp \
 	data_acq_and_processing/processing/data_processor/src/processing_live_config.cpp \
 	data_acq_and_processing/processing/data_processor/src/data_processor.cpp \
 	data_acq_and_processing/processing/data_processor/src/main.cpp \
@@ -102,12 +139,24 @@ TARGETS := \
 	$(BIN_DIR)/data_preprocessor \
 	$(BIN_DIR)/data_processor
 
-.PHONY: all clean kamil_adc_collector
+.PHONY: all clean kamil_adc_collector sar3d_parity
 
 all: $(TARGETS)
 
 kamil_adc_collector: $(BIN_DIR)/kamil_adc_collector
 
+# On-demand parity check of the sar3d processor against the gpr_cpp
+# standalone_bp example (same .npy dataset). Not part of `all`.
+SAR3D_PARITY_OBJS := $(addprefix $(BUILD_DIR)/, \
+	data_acq_and_processing/processing/tools/sar3d_parity_main.o \
+	data_acq_and_processing/processing/processors/src/sar3d_processor.o)
+
+sar3d_parity: $(BIN_DIR)/sar3d_parity
+
+$(BIN_DIR)/sar3d_parity: $(SAR3D_PARITY_OBJS) $(GPR_CPP_LIB)
+	@mkdir -p $(BIN_DIR)
+	$(CXX) $(SAR3D_PARITY_OBJS) -o $@ $(DATA_PROCESSOR_LDFLAGS)
+
 $(BIN_DIR)/sweep_orchestrator: $(SWEEP_ORCH_OBJS)
 	@mkdir -p $(BIN_DIR)
 	$(CXX) $(SWEEP_ORCH_OBJS) -o $@ $(ORCH_LDFLAGS)
@@ -116,9 +165,17 @@ $(BIN_DIR)/data_preprocessor: $(PREPROCESSOR_OBJS)
 	@mkdir -p $(BIN_DIR)
 	$(CXX) $(PREPROCESSOR_OBJS) -o $@ $(LDFLAGS)
 
-$(BIN_DIR)/data_processor: $(DATA_PROCESSOR_OBJS)
+$(BIN_DIR)/data_processor: $(DATA_PROCESSOR_OBJS) $(GPR_CPP_LIB)
 	@mkdir -p $(BIN_DIR)
-	$(CXX) $(DATA_PROCESSOR_OBJS) -o $@ $(LDFLAGS)
+	$(CXX) $(DATA_PROCESSOR_OBJS) -o $@ $(DATA_PROCESSOR_LDFLAGS)
+
+# Vendored gpr_cpp static lib: delegate to its CMake build (also produces the
+# standalone_bp example used by sar3d_parity cross-checks).
+GPR_CPP_SRC_FILES := $(wildcard $(GPR_CPP_DIR)/src/*) $(wildcard $(GPR_CPP_DIR)/include/gpr/*) \
+	$(wildcard $(GPR_CPP_DIR)/examples/*) $(GPR_CPP_DIR)/CMakeLists.txt $(GPR_CPP_DIR)/build_cmake.sh
+
+$(GPR_CPP_LIB): $(GPR_CPP_SRC_FILES)
+	GPR_BUILD_PYTHON=OFF VENV=$(abspath $(GPR_VENV)) bash $(GPR_CPP_DIR)/build_cmake.sh
 
 # The collector compiles into its own build/kamil/ tree with collector-only
 # includes (no VISA), so the shared driver/config sources it reuses never collide
@@ -141,4 +198,4 @@ $(BUILD_DIR)/%.o: %.cpp
 -include $(DEPFILES)
 
 clean:
-	rm -rf $(BUILD_DIR)
+	rm -rf $(BUILD_DIR) $(GPR_CPP_DIR)/build

README.md

@@ -60,7 +60,7 @@ controls plus a pixel-identical live plot through an embedded web UI.
 
 | Path | Contents |
 | --- | --- |
-| `data_acq_and_processing/` | C++ pipeline: `sweep_orchestrator/`, `preprocessing/`, `processing/` (including `locator/`), shared `common_cpp/` (config, IPC rings), and vendored `third_party/` (Eigen, nlohmann/json). |
+| `data_acq_and_processing/` | C++ pipeline: `sweep_orchestrator/`, `preprocessing/`, `processing/` (including `locator/`), shared `common_cpp/` (config, IPC rings), the vendored `gpr_cpp/` C++/CUDA 3D SAR library, and vendored `third_party/` (Eigen, nlohmann/json). |
 | `python_app/` | Python application: `gui/`, `orchestration/` (process supervisor, shared-memory readers, config), `hardware_full/` (device drivers and services), `webui/`, `storage/`, `models/`, `workflows/`, `scripts/` (producers and tools), and `tests/`. |
 | `run_config_examples/` | Ready-made configuration profiles, one per hardware setup. |
 | `deploy/` | systemd daemon installer for Raspberry Pi appliances. |
@@ -74,6 +74,13 @@ controls plus a pixel-identical live plot through an embedded web UI.
   specific (GPIO12/13/18/19, `dtoverlay=pwm-2chan`).
 - A C++20 compiler and `make` — Eigen and nlohmann/json are vendored, so no extra C++ dependencies are needed.
 - **Python 3.10 or newer**, in a virtual environment.
+- For the `sar3d` processor (3D SAR): an NVIDIA GPU with the driver installed
+  (`libcuda`) and `cmake` (the vendored `gpr_cpp` library at
+  `data_acq_and_processing/gpr_cpp/` builds via its own CMake, invoked
+  automatically by `make`). CUDA headers/libnvrtc are taken from the pip
+  `nvidia-*` packages in this repo's venv (installed by `requirements.txt` on
+  x86_64), otherwise from the system toolkit at `/usr/local/cuda`
+  (Jetson/JetPack, apt `cuda-toolkit`).
 
 ## Build and run
 
@@ -86,6 +93,21 @@ python3 -m venv .venv
 make                     # build the C++ binaries into build/bin/
 ```
 
+`make` also builds the vendored 3D SAR library
+(`data_acq_and_processing/gpr_cpp/`, CMake + NVRTC — needs `cmake`; CUDA
+headers come from the pip `nvidia-*` packages in `.venv`, or `/usr/local/cuda`
+as a fallback).
+
+To verify the sar3d processor against the gpr_cpp reference on the recorded
+dataset, build and run the on-demand parity tool:
+
+```bash
+make sar3d_parity
+./build/bin/sar3d_parity <preprocessed_dataset_dir> <center_snapshot_name>
+# compare the printed peak/argmax with
+# ./data_acq_and_processing/gpr_cpp/build/standalone_bp on the same arguments
+```
+
 Run the desktop GUI:
 
 ```bash
@@ -99,6 +121,69 @@ Run the desktop GUI:
 `./start.sh --help`. The acquisition device (LibreVNA, Kamil ADC, …) is detected
 automatically from the active config's `radar.model` — there is no per-device flag.
 
+## Jetson (aarch64) bring-up
+
+Ubuntu 22.04 on Jetson has glibc 2.35, and PyQt6 6.8+ ships only
+`manylinux_2_39` aarch64 wheels — `requirements.txt` therefore pins PyQt6 to
+`<6.8` on aarch64 (6.7.1 is the last release with a compatible wheel). With
+that pin the stock `./start.sh` provisioning works; the extra steps are only
+for the CUDA-backed `sar3d` processor:
+
+```bash
+# 1. copy the replay dataset (for dataset-replay profiles) and point
+#    radar.replay.dataset_dir in the profile at the local copy — the path in
+#    run_config_examples/*.json is machine-specific.
+
+# 2. normal launch; the vendored gpr_cpp builds against the JetPack system CUDA
+#    (requirements.txt installs no pip nvidia-* packages on aarch64, and
+#    build_cmake.sh falls back to /usr/local/cuda automatically; override with
+#    CUDA_INC=... CUDA_LIB=... if the toolkit lives elsewhere):
+./start.sh --profile run_config_examples/run_config_replay_vna.example.json
+```
+
+If linking fails with `cannot find -lcuda`, the driver stub lives at
+`/usr/local/cuda/lib64/stubs/libcuda.so`; on JetPack the real driver is
+`/usr/lib/aarch64-linux-gnu/nvidia/libcuda.so.1` (present after boot).
+
+## Running without hardware (emulation)
+
+Two ways to run the full pipeline with no radar attached:
+
+**Mock producer (synthetic data).** The default `run_config.json` already sets
+`radar.driver_mode: "mock"`, so the producer generates synthetic S-parameters.
+The preprocessor still requires S21 calibration/reference sets; create them once:
+
+```bash
+.venv/bin/python -m python_app.scripts.manual_smoke_run --duration 1.0  # creates smoke_cal/smoke_ref
+./start.sh   # then in the GUI: Preprocessing Panel -> pick the sets -> Start
+```
+
+(Alternatively use the GUI's "Create neutral sets" button, which makes
+identity calibration + zero reference so preprocessing is a pass-through.)
+
+**Dataset replay (recorded scene, meaningful sar3d picture).** When
+`radar.replay.dataset_dir` is set in the profile, the supervisor launches
+`python_app.scripts.replay_raw_producer` instead of the normal producer: it
+replays recorded snapshots into the raw ring at the recorded cadence, looping
+forever with a monotonic timeline. Pair it with NEUTRAL preprocess sets so the
+already-preprocessed recording passes through the preprocessor bit-exactly.
+A ready profile for the VNA_AlGo tank005 dataset is provided:
+
+```bash
+# one-time: neutral sets for this profile (or use the GUI button)
+.venv/bin/python -m python_app.scripts.create_neutral_preprocess_sets \
+    --config run_config_examples/run_config_replay_vna.example.json
+# interactive GUI (mode sar3d is pre-selected by the profile's gui section):
+./start.sh --profile run_config_examples/run_config_replay_vna.example.json
+# or a fully headless end-to-end check:
+.venv/bin/python -m python_app.scripts.replay_smoke_run --duration 30
+```
+
+Note: the sar3d processor allocates GPU memory; if the desktop (e.g. a browser)
+has filled the GPU, sar3d logs `out of memory` and publishes empty results until
+memory frees. Reduce `sar3d_grid_nx/ny/nz` in the processing panel to shrink
+the footprint.
+
 The web interface is served in-process at `http://<host>:8080`. It mirrors the live
 plot and offers Start / Stop / Capture controls and the active processor settings.
 

data_acq_and_processing/gpr_cpp/CMakeLists.txt

@@ -0,0 +1,78 @@
+# 3.22 = apt cmake на Ubuntu 22.04 (Jetson/JetPack); фич новее здесь нет.
+cmake_minimum_required(VERSION 3.22)
+project(gpr_cpp LANGUAGES CXX)   # CXX-only: CUDA компилируется в рантайме через NVRTC
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+if(NOT CMAKE_BUILD_TYPE)
+  set(CMAKE_BUILD_TYPE Release)
+endif()
+set(CMAKE_CXX_FLAGS_RELEASE "-O3")
+
+# bindings/pymodule.cpp не вендорен в radar_system — включение потребует вернуть его из VNA_AlGo
+option(GPR_BUILD_PYTHON "Build pybind11 module" OFF)
+option(GPR_BUILD_TESTS  "Build C++ parity tests" OFF)
+option(GPR_BUILD_EXAMPLES "Build standalone C++ example" ON)
+
+# Python окружение (venv) — источник pybind11 и pip CUDA-библиотек (nvrtc, cudart).
+if(NOT DEFINED PYTHON_EXECUTABLE)
+  set(PYTHON_EXECUTABLE python3)
+endif()
+
+# CUDA include/lib из pip-пакетов nvidia-*; задаются -DGPR_CUDA_INC / -DGPR_CUDA_LIB
+# (детект делает build_cmake.sh). Заголовки: nvrtc.h, cuda.h (Driver API).
+# GPR_CUDA_INC может быть СПИСКОМ каталогов (через ';'): в раздельных pip-пакетах
+# (Jetson: cuda_nvrtc + cuda_runtime) nvrtc.h и cuda.h лежат в разных include/.
+if(NOT DEFINED GPR_CUDA_INC OR NOT DEFINED GPR_CUDA_LIB)
+  message(FATAL_ERROR "Задайте -DGPR_CUDA_INC=<dir(ы) с nvrtc.h,cuda.h; ';'-список> и -DGPR_CUDA_LIB=<dir с libnvrtc.so*>")
+endif()
+
+# libcuda (драйвер) — системный
+find_library(CUDA_DRIVER_LIB cuda PATHS /usr/lib/x86_64-linux-gnu /usr/lib/aarch64-linux-gnu
+             /usr/lib REQUIRED)
+# libnvrtc из pip
+find_library(NVRTC_LIB NAMES nvrtc libnvrtc.so.13 libnvrtc.so.12 nvrtc.so
+             PATHS ${GPR_CUDA_LIB} REQUIRED NO_DEFAULT_PATH)
+
+# PocketFFT (header-only, vendored)
+add_library(pocketfft INTERFACE)
+target_include_directories(pocketfft INTERFACE ${CMAKE_CURRENT_SOURCE_DIR}/third_party/pocketfft)
+
+# --- ядро библиотеки (usable без Python) ---
+add_library(gpr_core_cpp STATIC
+  src/geometry.cpp
+  src/ascan.cpp
+  src/device.cpp
+  src/nvrtc_kernel.cpp
+  src/sar3d.cpp)
+target_include_directories(gpr_core_cpp PUBLIC
+  ${CMAKE_CURRENT_SOURCE_DIR}/include
+  ${GPR_CUDA_INC})
+target_link_libraries(gpr_core_cpp PUBLIC pocketfft ${NVRTC_LIB} ${CUDA_DRIVER_LIB})
+# rpath на каталог pip libnvrtc — чтобы бинарник/модуль грузились без LD_LIBRARY_PATH
+set_target_properties(gpr_core_cpp PROPERTIES POSITION_INDEPENDENT_CODE ON
+  BUILD_RPATH "${GPR_CUDA_LIB}" INSTALL_RPATH "${GPR_CUDA_LIB}")
+
+if(GPR_BUILD_PYTHON)
+  execute_process(COMMAND ${PYTHON_EXECUTABLE} -m pybind11 --cmakedir
+                  OUTPUT_VARIABLE PB11_DIR OUTPUT_STRIP_TRAILING_WHITESPACE RESULT_VARIABLE PB11_RC)
+  if(PB11_RC EQUAL 0)
+    find_package(pybind11 CONFIG REQUIRED HINTS ${PB11_DIR})
+  else()
+    find_package(pybind11 CONFIG REQUIRED)
+  endif()
+  pybind11_add_module(_gpr_cpp bindings/pymodule.cpp)
+  target_link_libraries(_gpr_cpp PRIVATE gpr_core_cpp)
+  set_target_properties(_gpr_cpp PROPERTIES BUILD_RPATH "${GPR_CUDA_LIB}" INSTALL_RPATH "${GPR_CUDA_LIB}")
+endif()
+
+if(GPR_BUILD_TESTS)
+  add_executable(gpr_test_parity tests/cpp/test_parity.cpp)
+  target_link_libraries(gpr_test_parity PRIVATE gpr_core_cpp)
+endif()
+
+if(GPR_BUILD_EXAMPLES)
+  add_executable(standalone_bp examples/standalone_bp.cpp)
+  target_link_libraries(standalone_bp PRIVATE gpr_core_cpp)
+  set_target_properties(standalone_bp PROPERTIES BUILD_RPATH "${GPR_CUDA_LIB}" INSTALL_RPATH "${GPR_CUDA_LIB}")
+endif()

data_acq_and_processing/gpr_cpp/README.md

@@ -0,0 +1,87 @@
+# gpr_cpp — C++/CUDA счёт-ядро MIMO GPR 3D SAR
+
+> Вендорено из репозитория VNA_AlGo (`gpr_cpp/`, коммит `544dc59`); python-биндинги
+> (`bindings/`, pybind11) при переносе не взяты — сборка здесь только C++
+> (`GPR_BUILD_PYTHON=OFF`). Собирается из корневого Makefile radar_system
+> (цель `$(GPR_CPP_LIB)` вызывает `build_cmake.sh`).
+
+Порт вычислительного ядра из Python/CuPy (`gpr_core`) в самостоятельную
+C++/CUDA библиотеку для интеграции в C++ проект.
+Алгоритм — см. `ALGORITHM.md` в репозитории VNA_AlGo.
+
+## Что внутри (граница C++ ↔ Python)
+
+**В C++ (`libgpr_core_cpp`):** геометрия, препроцессинг A-сканов (bg-subtract,
+фазовая коррекция движения, k0-вставка + Ханн + IFFT через PocketFFT, нормировка),
+3D SAR апертурный объём (одно CUDA-ядро на всю апертуру). Возвращает **сырые**
+`acc`(complex64)/`cnt`(float32).
+
+**Остаётся в Python:** загрузка `.npy`, финализация объёма (abs+gate+gaussian+
+нормировка), CLEAN-детекция, трек, оценка скорости, plotly.
+
+## Ключевое решение: CUDA без nvcc (NVRTC + Driver API)
+
+На целевых машинах (x86+T600, Jetson) нет полного CUDA toolkit — только runtime +
+**NVRTC** (как у cupy). Поэтому ядро **не компилируется nvcc'ом**: его исходник
+встроен строкой ([src/kernel_src.hpp](src/kernel_src.hpp)), компилируется в рантайме
+через NVRTC в **cubin под точную арх устройства** (запрос compute capability →
+`sm_XY`, без JIT PTX), грузится и запускается через **CUDA Driver API**. Хост-часть
+собирается обычным g++. Плюс: без toolkit и без привязки к SM — один код для
+x86 (sm_75) и Jetson (sm_87).
+
+## Сборка
+
+```bash
+bash build_cmake.sh            # авто-детект pip CUDA-путей; собирает всё
+# артефакты в build/: libgpr_core_cpp.a, _gpr_cpp*.so, standalone_bp
+```
+Требуется: g++ ≥13, cmake ≥3.24, pybind11 (`pip install pybind11`), pip-пакеты
+`nvidia-*` (nvrtc+cudart, идут с cupy), системный `libcuda`. rpath на каталог
+libnvrtc прописан — бинарники/модуль грузятся без `LD_LIBRARY_PATH`.
+
+## Публичный C++ API ([include/gpr/](include/gpr/))
+
+```cpp
+#include "gpr/config.hpp"    // GeometryConfig/BPConfig/MotionConfig/SARConfig (дефолты = config.py)
+#include "gpr/ascan.hpp"     // preprocess_snapshot(...) -> AScanSet
+#include "gpr/sar3d.hpp"     // SarEngine::aperture_volume(frames, grid) -> VolumeRaw{acc,cnt}
+#include "gpr/geometry.hpp"  // frame_offsets, frame_weights, build_grid_scene, point_to_scene
+
+gpr::SarEngine engine(geom, bp, mot, sar);          // компилит ядро под устройство
+auto vol = engine.aperture_volume(frames, grid);    // один запуск ядра на апертуру
+// vol.acc / vol.cnt — сырые; финализацию (gaussian/normalize) делает вызывающий
+```
+Пример без Python: [examples/standalone_bp.cpp](examples/standalone_bp.cpp) — грузит
+S21 из `.npy`, строит объём, печатает максимум.
+```bash
+./build/standalone_bp <preprocessed_dir> 0038_id39_ns11361438137486
+```
+
+## Python-модуль
+
+```python
+import _gpr_cpp                     # из gpr_cpp/build
+eng = _gpr_cpp.SarEngine(_gpr_cpp.GeometryConfig(), _gpr_cpp.BPConfig(),
+                         _gpr_cpp.MotionConfig(), _gpr_cpp.SARConfig())
+a = _gpr_cpp.preprocess_snapshot(s21, freq, background, mot, geom, bp, True)  # -> AScanSet
+acc, cnt = eng.aperture_volume([(a, dx, dy, dz, w), ...], gx, gy, gz)         # numpy
+```
+
+## Паритет (проверено на x86 T600)
+
+| Этап | Порог | Факт |
+|---|---|---|
+| config дефолты vs dataclass | точно | OK |
+| A-сканы raw H vs numpy (float64) | <1e-9 | 1.7e-12 |
+| A-сканы норм. H (FFT-ULP через медиану) | <1e-4 | 2.5e-6 |
+| **3D объём C++ ядро vs cupy** (float32) | <5e-3 | **3e-7, argmax идентичен** |
+
+Тесты: [../tests/test_parity_cpp.py](../tests/test_parity_cpp.py) (C++ vs cupy).
+
+## Статус
+
+- [x] x86 (T600, sm_75, CUDA 13): собрано, паритет, standalone работает.
+- [ ] Jetson (Orin NX, sm_87, CUDA 12.6): порт тривиален (NVRTC сам берёт sm_87);
+      собрать `build_cmake.sh` на Jetson с его venv.
+- Опц.: ветка `impl='cpp'` в `gpr_core/sar3d.py` (адаптер `gpr_cpp_backend.py`),
+      кэш device-A-сканов (сейчас каждый вызов грузит стек на GPU).

data_acq_and_processing/gpr_cpp/build_cmake.sh

@@ -0,0 +1,87 @@
+#!/usr/bin/env bash
+# CMake-сборка libgpr_core_cpp + standalone (без nvcc; NVRTC/driver).
+#
+# Источник CUDA (nvrtc.h / libnvrtc), по убыванию приоритета:
+#   1. env CUDA_INC / CUDA_LIB — заданы явно;
+#   2. pip-пакеты nvidia-* в venv ($VENV, по умолчанию ../../.venv — venv
+#      radar_system) — дев-машина;
+#   3. системный toolkit /usr/local/cuda — Jetson/JetPack, apt cuda-toolkit.
+#
+# GPR_BUILD_PYTHON=ON вернул бы pybind-модуль, но bindings/ не вендорен в
+# radar_system (см. README) — по умолчанию OFF.
+set -eu
+HERE="$(cd "$(dirname "$0")" && pwd)"
+VENV="${VENV:-$HERE/../../.venv}"
+PY="$VENV/bin/python"
+GPR_BUILD_PYTHON="${GPR_BUILD_PYTHON:-OFF}"
+
+CUDA_INC="${CUDA_INC:-}"
+CUDA_LIB="${CUDA_LIB:-}"
+
+if [ -z "$CUDA_INC" ] || [ -z "$CUDA_LIB" ]; then
+    if [ -x "$PY" ]; then
+        # Печатается только когда найдены nvrtc.h, cuda.h И библиотека — частичная
+        # детекция не может подсунуть мусорный путь. nvrtc.h и cuda.h могут лежать
+        # в РАЗНЫХ каталогах (раздельные pip-пакеты cuda_nvrtc/cuda_runtime на
+        # Jetson) — тогда CUDA_INC становится ';'-списком для CMake.
+        PIP_CUDA=$("$PY" - <<'P' 2>/dev/null || true
+import glob, os, site
+sp = site.getsitepackages()[0]
+nvrtc_hdr = glob.glob(os.path.join(sp, 'nvidia', '**', 'include', 'nvrtc.h'), recursive=True)
+cuda_hdr = glob.glob(os.path.join(sp, 'nvidia', '**', 'include', 'cuda.h'), recursive=True)
+lib = sorted(glob.glob(os.path.join(sp, 'nvidia', '**', 'lib', 'libnvrtc.so*'), recursive=True))
+if nvrtc_hdr and cuda_hdr and lib:
+    inc_dirs = []
+    for hdr in (nvrtc_hdr[0], cuda_hdr[0]):
+        d = os.path.dirname(hdr)
+        if d not in inc_dirs:
+            inc_dirs.append(d)
+    print(';'.join(inc_dirs))
+    print(os.path.dirname(lib[0]))
+P
+)
+        if [ -n "$PIP_CUDA" ]; then
+            CUDA_INC=$(printf '%s\n' "$PIP_CUDA" | sed -n 1p)
+            CUDA_LIB=$(printf '%s\n' "$PIP_CUDA" | sed -n 2p)
+        fi
+    fi
+fi
+
+if [ -z "$CUDA_INC" ] || [ -z "$CUDA_LIB" ]; then
+    if [ -e /usr/local/cuda/include/nvrtc.h ]; then
+        CUDA_INC=/usr/local/cuda/include
+        CUDA_LIB=/usr/local/cuda/lib64
+        echo "[build_cmake] pip nvidia-* не найдены; использую системный CUDA: $CUDA_INC"
+    else
+        echo "[build_cmake] CUDA не найден: нет ни pip nvidia-* в $VENV, ни /usr/local/cuda." >&2
+        echo "[build_cmake] Задайте CUDA_INC=<dir с nvrtc.h,cuda.h> и CUDA_LIB=<dir с libnvrtc.so*>." >&2
+        exit 1
+    fi
+fi
+
+PYTHON_ARG=()
+if [ "$GPR_BUILD_PYTHON" = "ON" ] && [ -x "$PY" ]; then
+    PYTHON_ARG=(-DPYTHON_EXECUTABLE="$PY")
+fi
+
+# cmake: сначала из venv (pip install cmake — свежий без sudo), потом системный.
+CMAKE="$VENV/bin/cmake"
+if [ ! -x "$CMAKE" ]; then
+    CMAKE=$(command -v cmake || true)
+fi
+if [ -z "$CMAKE" ]; then
+    echo "[build_cmake] cmake не найден. Поставьте его в venv: \"$VENV/bin/pip\" install cmake" >&2
+    echo "[build_cmake] (или системно: sudo apt install cmake)" >&2
+    exit 1
+fi
+
+"$CMAKE" -S "$HERE" -B "$HERE/build" \
+  ${PYTHON_ARG[@]+"${PYTHON_ARG[@]}"} \
+  -DGPR_BUILD_PYTHON="$GPR_BUILD_PYTHON" \
+  -DGPR_CUDA_INC="$CUDA_INC" \
+  -DGPR_CUDA_LIB="$CUDA_LIB" \
+  -DCMAKE_BUILD_TYPE=Release "$@" 2>&1 | tail -8
+"$CMAKE" --build "$HERE/build" -j4 2>&1 | tail -15
+echo "--- артефакты ---"
+ls -la "$HERE/build/libgpr_core_cpp.a" "$HERE/build/standalone_bp" "$HERE"/build/_gpr_cpp*.so 2>/dev/null | awk '{print $5,$9}'
+echo "CUDA_LIB=$CUDA_LIB  (для LD_LIBRARY_PATH при запуске)"

data_acq_and_processing/gpr_cpp/examples/standalone_bp.cpp

@@ -0,0 +1,145 @@
+// Автономный пример: C++ без Python. Грузит S21 снапшота из .npy,
+// строит A-сканы и 3D SAR-объём апертуры, печатает максимум |acc/cnt|.
+//
+// Использование:
+//   standalone_bp <preprocessed_dir> <center_snapshot_name>
+// где preprocessed_dir содержит папки снапшотов с i{rx}_o{tx}_s21.npy и _freq.npy.
+#include <algorithm>
+#include <cmath>
+#include <complex>
+#include <cstdint>
+#include <cstdio>
+#include <dirent.h>
+#include <fstream>
+#include <map>
+#include <string>
+#include <vector>
+
+#include "gpr/ascan.hpp"
+#include "gpr/config.hpp"
+#include "gpr/geometry.hpp"
+#include "gpr/sar3d.hpp"
+
+using namespace gpr;
+
+// --- минимальный загрузчик numpy .npy (v1.0, 1-D complex128 или float64) ---
+struct NpyArray { std::vector<double> f64; std::vector<c128> c128v; bool is_complex = false; };
+
+static NpyArray load_npy(const std::string& path) {
+    std::ifstream in(path, std::ios::binary);
+    if (!in) throw std::runtime_error("не открыть " + path);
+    char magic[6]; in.read(magic, 6);
+    uint8_t ver[2]; in.read((char*)ver, 2);
+    uint16_t hlen; in.read((char*)&hlen, 2);
+    std::string hdr(hlen, '\0'); in.read(&hdr[0], hlen);
+    auto sp = hdr.find("'shape':");
+    auto lp = hdr.find('(', sp), rp = hdr.find(')', sp);
+    long n = std::stol(hdr.substr(lp + 1, rp - lp - 1));
+    NpyArray a;
+    auto has = [&](const char* s) { return hdr.find(s) != std::string::npos; };
+    if (has("c8")) {  // complex64 → c128
+        a.is_complex = true; a.c128v.resize(n);
+        std::vector<std::complex<float>> tmp(n);
+        in.read((char*)tmp.data(), n * sizeof(std::complex<float>));
+        for (long i = 0; i < n; ++i) a.c128v[i] = c128(tmp[i].real(), tmp[i].imag());
+    } else if (has("c16")) {  // complex128
+        a.is_complex = true; a.c128v.resize(n);
+        in.read((char*)a.c128v.data(), n * sizeof(c128));
+    } else if (has("f4")) {  // float32 → double
+        a.f64.resize(n);
+        std::vector<float> tmp(n); in.read((char*)tmp.data(), n * sizeof(float));
+        for (long i = 0; i < n; ++i) a.f64[i] = tmp[i];
+    } else {  // float64
+        a.f64.resize(n); in.read((char*)a.f64.data(), n * sizeof(double));
+    }
+    return a;
+}
+
+int main(int argc, char** argv) {
+    if (argc < 3) { printf("usage: %s <preprocessed_dir> <center_snapshot>\n", argv[0]); return 1; }
+    std::string root = argv[1], center = argv[2];
+
+    GeometryConfig geom; geom.finalize();
+    BPConfig bp; MotionConfig mot; SARConfig sar;
+    auto pairs = geom.pair_keys();
+    auto tx = geom.output_idx, rx = geom.input_idx;
+    std::sort(tx.begin(), tx.end()); std::sort(rx.begin(), rx.end());
+
+    // список снапшотов, выбор апертуры вокруг center
+    std::vector<std::string> snaps;
+    if (DIR* d = opendir(root.c_str())) {
+        for (dirent* e; (e = readdir(d));) {
+            std::string n = e->d_name;
+            if (!n.empty() && isdigit(n[0])) snaps.push_back(n);
+        }
+        closedir(d);
+    }
+    std::sort(snaps.begin(), snaps.end());
+    int ci = -1; for (int i = 0; i < (int)snaps.size(); ++i) if (snaps[i] == center) ci = i;
+    if (ci < 0) { printf("центр %s не найден\n", center.c_str()); return 1; }
+    int lo = std::max(0, ci - sar.frame_radius), hi = std::min((int)snaps.size(), ci + sar.frame_radius + 1);
+    std::vector<std::string> aperture(snaps.begin() + lo, snaps.begin() + hi);
+    printf("апертура: %d кадров вокруг %s\n", (int)aperture.size(), center.c_str());
+
+    auto load_snapshot = [&](const std::string& snap,
+                             std::map<PairKey, std::vector<c128>>& s21,
+                             std::map<PairKey, std::vector<double>>& freq) {
+        for (int p = 0; p < (int)pairs.size(); ++p) {
+            int i_tx = pairs[p].first, i_rx = pairs[p].second;
+            int phys_tx = tx[i_tx], phys_rx = rx[i_rx];
+            char fn[512];
+            snprintf(fn, sizeof(fn), "%s/%s/i%d_o%d_s21.npy", root.c_str(), snap.c_str(), phys_rx, phys_tx);
+            s21[{i_tx, i_rx}] = load_npy(fn).c128v;
+            snprintf(fn, sizeof(fn), "%s/%s/i%d_o%d_freq.npy", root.c_str(), snap.c_str(), phys_rx, phys_tx);
+            freq[{i_tx, i_rx}] = load_npy(fn).f64;
+        }
+    };
+
+    auto ts = compute_pair_timestamps(mot, geom);
+    SarEngine engine(geom, bp, mot, sar);
+
+    // держим AScanSet живыми
+    std::vector<std::shared_ptr<AScanSet>> keep;
+    std::vector<ApertureFrame> frames;
+    std::vector<double> frame_times;
+    for (auto& snap : aperture) {
+        // время из суффикса _nsNNNN
+        auto pos = snap.rfind("_ns");
+        double t = (pos != std::string::npos) ? std::stod(snap.substr(pos + 3)) * 1e-9 : 0.0;
+        frame_times.push_back(t);
+    }
+    double t_ref = frame_times[ci - lo];
+    auto offs = frame_offsets(frame_times, t_ref, mot);
+    auto wts = frame_weights(offs, sar.edge_frame_weight);
+
+    for (int f = 0; f < (int)aperture.size(); ++f) {
+        std::map<PairKey, std::vector<c128>> s21;
+        std::map<PairKey, std::vector<double>> freq;
+        load_snapshot(aperture[f], s21, freq);
+        auto a = std::make_shared<AScanSet>(
+            preprocess_snapshot(s21, freq, nullptr, mot, geom, bp, ts, true));
+        keep.push_back(a);
+        frames.push_back(ApertureFrame{a.get(), offs[f], wts[f]});
+    }
+
+    // центр сетки: под центральным снапшотом на глубине 5.6 м (сцена)
+    auto center_scene = point_to_scene(0.0, 0.0, 5.6, mot, sar);
+    Grid grid = build_grid_scene(center_scene, sar);
+    printf("сетка %dx%dx%d, счёт объёма...\n", grid.NX(), grid.NY(), grid.NZ());
+
+    VolumeRaw vol = engine.aperture_volume(frames, grid);
+
+    // финализация: |acc/cnt|, максимум
+    double vmax = 0; size_t amax = 0;
+    for (size_t i = 0; i < vol.acc.size(); ++i) {
+        double v = std::abs(vol.acc[i]) / (vol.cnt[i] + 1e-12);
+        if (v > vmax) { vmax = v; amax = i; }
+    }
+    int NX = vol.NX, NZ = vol.NZ;
+    int ix = amax % NX, iz = (amax / NX) % NZ, iy = amax / ((size_t)NX * NZ);
+    printf("МАКСИМУМ объёма: |acc/cnt|=%.4e в вокселе (ix=%d,iy=%d,iz=%d)\n", vmax, ix, iy, iz);
+    printf("  X=%.1f см, Y=%.1f см, Z=%.1f см\n",
+           grid.x[ix] * 100, grid.y[iy] * 100, grid.z[iz] * 100);
+    printf("OK: автономный C++ путь отработал без Python.\n");
+    return 0;
+}

data_acq_and_processing/gpr_cpp/include/gpr/ascan.hpp

@@ -0,0 +1,58 @@
+// Препроцессинг S21 → комплексные A-сканы (host, float64). Порт preprocess.py.
+#pragma once
+#include <map>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "gpr/config.hpp"
+#include "gpr/types.hpp"
+
+namespace gpr {
+
+using PairKey = std::pair<int, int>;
+
+// Тайминги пары (нужны slot/n_slots/t_center для фазовой коррекции).
+struct PairTiming {
+    int tx_interleave_slot = 0;
+    int tx_interleave_slots = 1;
+    double t_center_s = 0.0;
+};
+using PairTimestamps = std::map<PairKey, PairTiming>;
+
+PairTimestamps compute_pair_timestamps(const MotionConfig&, const GeometryConfig&);
+
+// A-сканы одного снапшота, общая ось времени у всех пар (assert).
+struct AScanSet {
+    int n_pairs = 0;
+    int n_fft = 0;
+    double t0 = 0.0, dt = 0.0;               // общая ось: t_k = t0 + k*dt
+    std::vector<PairKey> pairs;              // порядок как в geom.pair_keys()
+    std::vector<c128> H;                     // [n_pairs * n_fft] float64
+    const c128* pair_ptr(int p) const { return H.data() + (size_t)p * n_fft; }
+};
+
+// Порт compute_pair_ascans: bg-subtract → phase correction → k0-IFFT → норма.
+// s21/freq/background — по ключам (i_tx,i_rx); background может быть nullptr.
+AScanSet preprocess_snapshot(
+    const std::map<PairKey, std::vector<c128>>& s21,
+    const std::map<PairKey, std::vector<double>>& freq,
+    const std::map<PairKey, std::vector<c128>>* background,
+    const MotionConfig&, const GeometryConfig&, const BPConfig&,
+    const PairTimestamps&, bool normalize = true);
+
+// Отдельные шаги (для тестов/переиспользования):
+std::vector<c128> apply_intra_sweep_phase_correction(
+    const std::vector<c128>& s21, const std::vector<double>& freq_full,
+    const PairTiming&, const MotionConfig&, double wave_speed,
+    double f_start, double f_stop, int n_slots_default);
+
+// Возвращает t (ось), h (комплексный A-скан). n_fft через out-параметр.
+void compute_ascan_bp(const std::vector<c128>& s21, const std::vector<double>& freq,
+                      double f_start, double f_stop, bool window, int oversample,
+                      std::vector<double>& t_out, std::vector<c128>& h_out);
+
+double pair_norm_scale(const std::vector<c128>& h, const std::vector<double>& t,
+                       double wave_speed, const BPConfig&);
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/include/gpr/config.hpp

@@ -0,0 +1,102 @@
+// Конфигурация счёт-ядра. Дефолты 1:1 повторяют gpr_core/config.py.
+#pragma once
+#include <algorithm>
+#include <array>
+#include <cmath>
+#include <cstdint>
+#include <utility>
+#include <vector>
+
+namespace gpr {
+
+struct GeometryConfig {
+    std::vector<int> input_idx{0, 1, 2, 3};
+    std::vector<int> output_idx{0, 1};
+    // (x,y,z) метры; порядок по физическим индексам как в config.py
+    std::vector<std::array<double, 3>> tx_positions{{0.785, 0.0, 0.0}, {-0.785, 0.0, 0.0}};
+    std::vector<std::array<double, 3>> rx_positions{
+        {-0.42, 0.0, 0.0}, {0.48, 0.0, 0.0}, {0.18, 0.0, 0.0}, {-0.12, 0.0, 0.0}};
+    double eps_r = 1.0;
+    double bp_plane_y = 0.0;
+
+    // --- производные (заполняются finalize) ---
+    std::vector<double> x_tx, y_tx, z_tx;   // отсортированы по phys-индексу
+    std::vector<double> x_rx, y_rx, z_rx;
+    int n_tx = 0, n_rx = 0;
+    double v = 0.0;
+
+    void finalize() {
+        auto sort_xyz = [](std::vector<int> idx, const std::vector<std::array<double, 3>>& pos,
+                           std::vector<double>& X, std::vector<double>& Y, std::vector<double>& Z) {
+            std::vector<int> order(idx.size());
+            for (size_t i = 0; i < idx.size(); ++i) order[i] = (int)i;
+            std::sort(order.begin(), order.end(), [&](int a, int b) { return idx[a] < idx[b]; });
+            X.clear(); Y.clear(); Z.clear();
+            for (int o : order) { X.push_back(pos[o][0]); Y.push_back(pos[o][1]); Z.push_back(pos[o][2]); }
+        };
+        sort_xyz(output_idx, tx_positions, x_tx, y_tx, z_tx);
+        sort_xyz(input_idx, rx_positions, x_rx, y_rx, z_rx);
+        n_tx = (int)output_idx.size();
+        n_rx = (int)input_idx.size();
+        v = 3e8 / std::sqrt(eps_r);
+    }
+
+    std::vector<std::pair<int, int>> pair_keys() const {
+        std::vector<std::pair<int, int>> p;
+        for (int i = 0; i < n_tx; ++i)
+            for (int j = 0; j < n_rx; ++j) p.emplace_back(i, j);
+        return p;
+    }
+};
+
+struct BPConfig {
+    double f_start = 20 * 1e8;   // 2 ГГц
+    double f_stop = 6.0e9;
+    double min_depth = 3.2;
+    double max_depth = 10.0;
+    double grid_x_pad = 2.0;
+    double z_min = 0.1;
+    int nx = 300;
+    int nz = 300;
+    int oversample = 8;
+    bool window = true;
+    bool pair_normalize = true;
+    double pair_norm_percentile = 50.0;
+    double pair_norm_depth_min = 3.2;
+    double pair_norm_depth_max = 10.0;
+    double pair_norm_eps = 1e-15;
+    double comp_range_power = 0.25;
+    double comp_weight_max = 5.0;
+    double comp_ref_depth = 5.0;
+    double smooth_sigma = 1.5;
+    bool bg_subtract = true;
+    double cf_eps = 1e-12;
+};
+
+struct MotionConfig {
+    double speed_m_s = 0.854;
+    double look_angle_deg = 6.0;
+    double pitch_deg = 6.0;
+    double tx_sweep_time_s = 0.079;
+    double tx_switch_time_s = 1e-5;
+    // frequency_timing_mode всегда interleaved_tx_by_frequency
+    double direction_sign = +1.0;
+    bool apply_freq_phase_correction = true;
+    double motion_phase_sign = -1.0;
+
+    double frame_time_s(int n_tx) const {
+        return n_tx * tx_sweep_time_s + std::max(0, n_tx - 1) * tx_switch_time_s;
+    }
+};
+
+struct SARConfig {
+    int frame_radius = 6;
+    double edge_frame_weight = 0.75;
+    std::array<double, 3> half_width_m{1.8, 1.0, 1.8};
+    std::array<int, 3> grid{150, 120, 160};   // NX, NY, NZ
+    bool coord_mode_ground = true;             // 'ground' vs 'radar_local'
+    bool use_compensation = true;
+    bool normalize_each_frame = true;
+};
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/include/gpr/device.hpp

@@ -0,0 +1,52 @@
+// CUDA Driver API RAII: контекст, поток, буфер устройства.
+#pragma once
+#include <cuda.h>
+
+#include <cstddef>
+#include <stdexcept>
+#include <string>
+
+namespace gpr {
+
+void cu_check(CUresult r, const char* what);
+
+// Инициализация драйвера + primary context на device 0. Возвращает (major,minor) CC.
+struct CudaContext {
+    CUdevice dev = 0;
+    CUcontext ctx = nullptr;
+    int cc_major = 0, cc_minor = 0;
+    CudaContext();
+    ~CudaContext();
+    std::string arch_flag() const;  // "--gpu-architecture=sm_XY"
+};
+
+template <class T>
+class DeviceBuffer {
+    CUdeviceptr ptr_ = 0;
+    size_t n_ = 0;
+
+   public:
+    DeviceBuffer() = default;
+    explicit DeviceBuffer(size_t n) { alloc(n); }
+    ~DeviceBuffer() { free(); }
+    DeviceBuffer(const DeviceBuffer&) = delete;
+    DeviceBuffer& operator=(const DeviceBuffer&) = delete;
+    DeviceBuffer(DeviceBuffer&& o) noexcept : ptr_(o.ptr_), n_(o.n_) { o.ptr_ = 0; o.n_ = 0; }
+
+    void alloc(size_t n) {
+        free();
+        if (n) cu_check(cuMemAlloc(&ptr_, n * sizeof(T)), "cuMemAlloc");
+        n_ = n;
+    }
+    void free() {
+        if (ptr_) cuMemFree(ptr_);
+        ptr_ = 0; n_ = 0;
+    }
+    void upload(const T* host, size_t n) { cu_check(cuMemcpyHtoD(ptr_, host, n * sizeof(T)), "HtoD"); }
+    void download(T* host, size_t n) const { cu_check(cuMemcpyDtoH(host, ptr_, n * sizeof(T)), "DtoH"); }
+    CUdeviceptr ptr() const { return ptr_; }
+    CUdeviceptr* addr() { return &ptr_; }
+    size_t size() const { return n_; }
+};
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/include/gpr/geometry.hpp

@@ -0,0 +1,28 @@
+// Геометрия: опорное расхождение пары, смещения/веса кадров, построение сетки.
+#pragma once
+#include <vector>
+
+#include "gpr/config.hpp"
+#include "gpr/types.hpp"
+
+namespace gpr {
+
+// 1/geo_ref пары (порт pair_geo_ref, скаляр float64).
+double pair_geo_ref(const GeometryConfig&, const BPConfig&, int i_tx, int i_rx, double yy);
+
+// Смещения кадров относительно опорного (порт sar_frame_offsets).
+// frame_times[k] — время кадра (сек, от ns-таймстемпа); t_ref — время опорного.
+std::vector<FrameOffset> frame_offsets(const std::vector<double>& frame_times, double t_ref,
+                                       const MotionConfig&);
+
+// Косинусные веса апертуры (порт sar_frame_weights); dt берётся относительно центра.
+std::vector<double> frame_weights(const std::vector<FrameOffset>& offsets, double edge_weight);
+
+// Сетка вокруг точки сцены (порт sar_build_grid_scene).
+Grid build_grid_scene(const std::array<double, 3>& center_scene, const SARConfig&);
+
+// point_to_scene: оси радара → сцена (поворот на pitch, если ground).
+std::array<double, 3> point_to_scene(double x, double y, double z, const MotionConfig&,
+                                     const SARConfig&);
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/include/gpr/nvrtc_kernel.hpp

@@ -0,0 +1,25 @@
+// Компиляция строки CUDA-ядра через NVRTC → cubin под арх устройства → CUfunction.
+#pragma once
+#include <cuda.h>
+
+#include <string>
+
+namespace gpr {
+
+class NvrtcKernel {
+    CUmodule module_ = nullptr;
+    CUfunction func_ = nullptr;
+
+   public:
+    NvrtcKernel() = default;
+    ~NvrtcKernel();
+    NvrtcKernel(const NvrtcKernel&) = delete;
+    NvrtcKernel& operator=(const NvrtcKernel&) = delete;
+
+    // Компилирует src в cubin под arch_flag ("--gpu-architecture=sm_XY"),
+    // загружает модуль и берёт функцию func_name.
+    void compile(const char* src, const std::string& arch_flag, const char* func_name);
+    CUfunction func() const { return func_; }
+};
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/include/gpr/sar3d.hpp

@@ -0,0 +1,44 @@
+// 3D SAR: сборка взвешенного объёма апертуры за один запуск ядра.
+#pragma once
+#include <memory>
+#include <vector>
+
+#include "gpr/ascan.hpp"
+#include "gpr/config.hpp"
+#include "gpr/device.hpp"
+#include "gpr/geometry.hpp"
+#include "gpr/nvrtc_kernel.hpp"
+#include "gpr/types.hpp"
+
+namespace gpr {
+
+// Один кадр апертуры: A-сканы + смещение + вес.
+struct ApertureFrame {
+    const AScanSet* ascans;
+    FrameOffset offset;
+    double weight;
+};
+
+// Сырой результат объёма (Python делает finish: abs+gate+gaussian+normalize).
+struct VolumeRaw {
+    std::vector<c64> acc;    // [NY*NZ*NX]
+    std::vector<float> cnt;  // [NY*NZ*NX]
+    int NX = 0, NY = 0, NZ = 0;
+};
+
+class SarEngine {
+   public:
+    SarEngine(const GeometryConfig&, const BPConfig&, const MotionConfig&, const SARConfig&);
+    // Весь объём апертуры за один запуск ядра; возвращает сырые acc/cnt (host).
+    VolumeRaw aperture_volume(const std::vector<ApertureFrame>& frames, const Grid&);
+
+   private:
+    GeometryConfig geom_;
+    BPConfig bp_;
+    MotionConfig mot_;
+    SARConfig sar_;
+    std::unique_ptr<CudaContext> ctx_;
+    std::unique_ptr<NvrtcKernel> kernel_;
+};
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/include/gpr/types.hpp

@@ -0,0 +1,25 @@
+// Базовые типы счёт-ядра.
+#pragma once
+#include <complex>
+#include <cstdint>
+#include <vector>
+
+namespace gpr {
+
+using c64 = std::complex<float>;    // complex64, layout == cuComplex (2×float)
+using c128 = std::complex<double>;  // float64 для препроцессинга/IFFT
+
+// Ground-scene сетка (порт dict из geometry.py: x,y,z оси в метрах).
+struct Grid {
+    std::vector<double> x, y, z;
+    int NX() const { return (int)x.size(); }
+    int NY() const { return (int)y.size(); }
+    int NZ() const { return (int)z.size(); }
+};
+
+// Смещение антенн кадра апертуры (порт sar_frame_offsets).
+struct FrameOffset {
+    double dx_m = 0.0, dy_m = 0.0, dz_m = 0.0, dt_s = 0.0, horizontal_m = 0.0;
+};
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/src/ascan.cpp

@@ -0,0 +1,198 @@
+#include "gpr/ascan.hpp"
+
+#include <algorithm>
+#include <cmath>
+#include <numeric>
+#include <stdexcept>
+
+#include "pocketfft_hdronly.h"
+
+namespace gpr {
+
+PairTimestamps compute_pair_timestamps(const MotionConfig& mot, const GeometryConfig& geom) {
+    // interleaved_tx_by_frequency: слот Tx = порядок в отсортированном output_idx.
+    std::vector<int> tx = geom.output_idx, rx = geom.input_idx;
+    std::sort(tx.begin(), tx.end());
+    std::sort(rx.begin(), rx.end());
+    int n_slots = (int)tx.size();
+    double frame_center = 0.5 * n_slots * mot.tx_sweep_time_s;
+    PairTimestamps ts;
+    for (int slot = 0; slot < (int)tx.size(); ++slot) {
+        int i_tx = slot;  // индекс в отсортированном == логический
+        for (int j = 0; j < (int)rx.size(); ++j) {
+            ts[{i_tx, j}] = PairTiming{slot, n_slots, frame_center};
+        }
+    }
+    return ts;
+}
+
+static double median_of_diff(const std::vector<double>& f) {
+    std::vector<double> d(f.size() - 1);
+    for (size_t i = 0; i + 1 < f.size(); ++i) d[i] = f[i + 1] - f[i];
+    std::sort(d.begin(), d.end());
+    size_t n = d.size();
+    return (n % 2) ? d[n / 2] : 0.5 * (d[n / 2 - 1] + d[n / 2]);
+}
+
+std::vector<c128> apply_intra_sweep_phase_correction(
+    const std::vector<c128>& s21, const std::vector<double>& freq_full,
+    const PairTiming& info, const MotionConfig& cfg, double wave_speed,
+    double f_start, double f_stop, int n_slots_default) {
+    std::vector<c128> out = s21;
+    int n_full = (int)freq_full.size();
+    // индексы в полосе
+    std::vector<int> idx;
+    for (int i = 0; i < n_full; ++i)
+        if (freq_full[i] >= f_start && freq_full[i] <= f_stop) idx.push_back(i);
+    if (idx.empty()) throw std::runtime_error("phase corr: пустая полоса");
+    if (n_full < 2) return out;
+
+    int n_slots = info.tx_interleave_slots ? info.tx_interleave_slots : n_slots_default;
+    int slot = info.tx_interleave_slot;
+    double dt_base = cfg.tx_sweep_time_s / (n_full - 1);
+    // t_all по всему свипу; t_center — середина ПОЛНОГО свипа (как в preprocess.py)
+    double t_first = (double)(n_slots * 0 + slot) * dt_base;
+    double t_last = (double)(n_slots * (n_full - 1) + slot) * dt_base;
+    double t_center = 0.5 * (t_first + t_last);
+    double theta = cfg.look_angle_deg * M_PI / 180.0;
+    if (!cfg.apply_freq_phase_correction) return out;
+    for (int i : idx) {
+        double t_abs = (double)(n_slots * i + slot) * dt_base;
+        double dt_intra = t_abs - t_center;
+        double delta_range = cfg.direction_sign * cfg.speed_m_s * std::cos(theta) * dt_intra;
+        double dtau = 2.0 * delta_range / wave_speed;
+        double phi = 2.0 * M_PI * freq_full[i] * dtau;
+        double ang = cfg.motion_phase_sign * phi;
+        out[i] *= c128(std::cos(ang), std::sin(ang));
+    }
+    return out;
+}
+
+static int next_pow2(int x) {
+    int p = 1;
+    while (p < x) p <<= 1;
+    return p;
+}
+
+void compute_ascan_bp(const std::vector<c128>& s21, const std::vector<double>& freq,
+                      double f_start, double f_stop, bool window, int oversample,
+                      std::vector<double>& t_out, std::vector<c128>& h_out) {
+    std::vector<double> freq_cut;
+    std::vector<c128> s_cut;
+    for (size_t i = 0; i < freq.size(); ++i)
+        if (freq[i] >= f_start && freq[i] <= f_stop) {
+            freq_cut.push_back(freq[i]);
+            s_cut.push_back(s21[i]);
+        }
+    if (freq_cut.size() < 2) throw std::runtime_error("ascan: <2 точек в полосе");
+
+    double df = median_of_diff(freq_cut);
+    int n = (int)freq_cut.size();
+    int k0 = (int)std::nearbyint(freq_cut[0] / df);  // round-half-to-even как python round
+    int min_len = 2 * (k0 + n - 1);
+    int n_fft_base = next_pow2(min_len);
+    int n_fft = n_fft_base * oversample;
+    double dt = 1.0 / ((double)n_fft * df);
+
+    // окно Ханна (numpy np.hanning: 0.5 - 0.5*cos(2*pi*i/(N-1)))
+    std::vector<c128> H(n_fft, c128(0.0, 0.0));
+    for (int i = 0; i < n; ++i) {
+        c128 v = s_cut[i];
+        if (window) {
+            double w = 0.5 - 0.5 * std::cos(2.0 * M_PI * i / (n - 1));
+            v *= w;
+        }
+        H[k0 + i] = v;
+    }
+
+    // IFFT через pocketfft (backward + 1/n_fft == numpy ifft)
+    h_out.assign(n_fft, c128(0.0, 0.0));
+    pocketfft::shape_t shape{(size_t)n_fft};
+    pocketfft::stride_t stride{(ptrdiff_t)sizeof(c128)};
+    pocketfft::c2c(shape, stride, stride, {0}, pocketfft::BACKWARD,
+                   H.data(), h_out.data(), (double)(1.0 / n_fft));
+
+    t_out.resize(n_fft);
+    for (int i = 0; i < n_fft; ++i) t_out[i] = i * dt;
+}
+
+double pair_norm_scale(const std::vector<c128>& h, const std::vector<double>& t,
+                       double wave_speed, const BPConfig& bp) {
+    std::vector<double> vals;
+    vals.reserve(h.size());
+    for (size_t i = 0; i < h.size(); ++i) {
+        double z = t[i] * wave_speed / 2.0;
+        if (z >= bp.pair_norm_depth_min && z <= bp.pair_norm_depth_max)
+            vals.push_back(std::abs(h[i]));
+    }
+    if (vals.empty())
+        for (auto& c : h) vals.push_back(std::abs(c));
+    std::sort(vals.begin(), vals.end());
+    // numpy.percentile linear interpolation
+    double q = bp.pair_norm_percentile / 100.0;
+    double pos = q * (vals.size() - 1);
+    int lo = (int)std::floor(pos);
+    int hi = std::min(lo + 1, (int)vals.size() - 1);
+    double scale = vals[lo] + (pos - lo) * (vals[hi] - vals[lo]);
+    if (!std::isfinite(scale) || scale <= bp.pair_norm_eps) scale = 1.0;
+    return scale;
+}
+
+AScanSet preprocess_snapshot(
+    const std::map<PairKey, std::vector<c128>>& s21,
+    const std::map<PairKey, std::vector<double>>& freq,
+    const std::map<PairKey, std::vector<c128>>* background,
+    const MotionConfig& mot, const GeometryConfig& geom, const BPConfig& bp,
+    const PairTimestamps& ts, bool normalize) {
+    AScanSet out;
+    out.pairs = geom.pair_keys();
+    out.n_pairs = (int)out.pairs.size();
+
+    std::vector<std::vector<c128>> Hs;
+    std::vector<std::vector<double>> Ts;
+    for (auto& key : out.pairs) {
+        auto its = s21.find(key);
+        if (its == s21.end()) throw std::runtime_error("preprocess: нет пары в s21");
+        std::vector<c128> s = its->second;  // копия
+        if (background) {
+            auto ib = background->find(key);
+            if (ib != background->end())
+                for (size_t i = 0; i < s.size(); ++i) s[i] -= ib->second[i];
+        }
+        const std::vector<double>& fr = freq.at(key);
+        auto itt = ts.find(key);
+        if (itt != ts.end())
+            s = apply_intra_sweep_phase_correction(s, fr, itt->second, mot, geom.v,
+                                                   bp.f_start, bp.f_stop, geom.n_tx);
+        std::vector<double> t;
+        std::vector<c128> h;
+        compute_ascan_bp(s, fr, bp.f_start, bp.f_stop, bp.window, bp.oversample, t, h);
+        Hs.push_back(std::move(h));
+        Ts.push_back(std::move(t));
+    }
+
+    // общая ось времени — проверка
+    int n_fft = (int)Ts[0].size();
+    double t0 = Ts[0][0];
+    double dt = (Ts[0].back() - t0) / (n_fft - 1);
+    for (auto& t : Ts)
+        if ((int)t.size() != n_fft || std::abs(t[0] - t0) > 1e-18)
+            throw std::runtime_error("A-сканы: разная ось времени у пар");
+
+    if (normalize) {
+        for (int p = 0; p < out.n_pairs; ++p) {
+            double scale = pair_norm_scale(Hs[p], Ts[p], geom.v, bp);
+            for (auto& c : Hs[p]) c /= (scale + bp.pair_norm_eps);
+        }
+    }
+
+    out.n_fft = n_fft;
+    out.t0 = t0;
+    out.dt = dt;
+    out.H.resize((size_t)out.n_pairs * n_fft);
+    for (int p = 0; p < out.n_pairs; ++p)
+        std::copy(Hs[p].begin(), Hs[p].end(), out.H.begin() + (size_t)p * n_fft);
+    return out;
+}
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/src/device.cpp

@@ -0,0 +1,34 @@
+#include "gpr/device.hpp"
+
+#include <cstdio>
+
+namespace gpr {
+
+void cu_check(CUresult r, const char* what) {
+    if (r != CUDA_SUCCESS) {
+        const char* s = nullptr;
+        cuGetErrorString(r, &s);
+        throw std::runtime_error(std::string("CUDA error in ") + what + ": " + (s ? s : "?"));
+    }
+}
+
+CudaContext::CudaContext() {
+    cu_check(cuInit(0), "cuInit");
+    cu_check(cuDeviceGet(&dev, 0), "cuDeviceGet");
+    cu_check(cuDeviceGetAttribute(&cc_major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, dev),
+             "cc_major");
+    cu_check(cuDeviceGetAttribute(&cc_minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev),
+             "cc_minor");
+    cu_check(cuDevicePrimaryCtxRetain(&ctx, dev), "primaryCtxRetain");
+    cu_check(cuCtxSetCurrent(ctx), "ctxSetCurrent");
+}
+
+CudaContext::~CudaContext() {
+    if (ctx) cuDevicePrimaryCtxRelease(dev);
+}
+
+std::string CudaContext::arch_flag() const {
+    return "--gpu-architecture=sm_" + std::to_string(cc_major) + std::to_string(cc_minor);
+}
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/src/geometry.cpp

@@ -0,0 +1,68 @@
+#include "gpr/geometry.hpp"
+
+#include <cmath>
+
+namespace gpr {
+
+double pair_geo_ref(const GeometryConfig& g, const BPConfig& bp, int i_tx, int i_rx, double yy) {
+    double xc = (g.x_tx[i_tx] + g.x_rx[i_rx]) / 2.0;
+    double Rtx = std::sqrt(std::pow(xc - g.x_tx[i_tx], 2) + std::pow(yy - g.y_tx[i_tx], 2) +
+                           std::pow(bp.comp_ref_depth - g.z_tx[i_tx], 2));
+    double Rrx = std::sqrt(std::pow(xc - g.x_rx[i_rx], 2) + std::pow(yy - g.y_rx[i_rx], 2) +
+                           std::pow(bp.comp_ref_depth - g.z_rx[i_rx], 2));
+    return 1.0 / (Rtx * Rrx + 1e-12) + 1e-30;
+}
+
+std::vector<FrameOffset> frame_offsets(const std::vector<double>& frame_times, double t_ref,
+                                       const MotionConfig& cfg) {
+    std::vector<FrameOffset> out(frame_times.size());
+    double pitch = cfg.pitch_deg * M_PI / 180.0;
+    for (size_t k = 0; k < frame_times.size(); ++k) {
+        double t = frame_times[k] - t_ref;
+        double h = cfg.direction_sign * cfg.speed_m_s * t;
+        out[k].dx_m = 0.0;
+        out[k].dy_m = h * std::sin(pitch);
+        out[k].dz_m = h * std::cos(pitch);
+        out[k].dt_s = t;
+        out[k].horizontal_m = h;
+    }
+    return out;
+}
+
+std::vector<double> frame_weights(const std::vector<FrameOffset>& offs, double edge_weight) {
+    size_t n = offs.size();
+    std::vector<double> w(n, 1.0);
+    double edge = std::min(std::max(edge_weight, 0.0), 1.0);
+    if (n == 0 || edge >= 1.0) return w;
+    double dmax = 0.0;
+    for (auto& o : offs) dmax = std::max(dmax, std::abs(o.dt_s));
+    if (dmax <= 1e-12) return w;
+    for (size_t i = 0; i < n; ++i) {
+        double r = std::min(std::abs(offs[i].dt_s) / dmax, 1.0);
+        w[i] = edge + (1.0 - edge) * 0.5 * (1.0 + std::cos(M_PI * r));
+    }
+    return w;
+}
+
+std::array<double, 3> point_to_scene(double x, double y, double z, const MotionConfig& mot,
+                                     const SARConfig& sar) {
+    if (!sar.coord_mode_ground) return {x, y, z};
+    double p = mot.pitch_deg * M_PI / 180.0;
+    double ys = y * std::cos(p) - z * std::sin(p);
+    double zs = y * std::sin(p) + z * std::cos(p);
+    return {x, ys, zs};
+}
+
+Grid build_grid_scene(const std::array<double, 3>& c, const SARConfig& sar) {
+    Grid g;
+    auto lin = [](double a, double b, int n, std::vector<double>& out) {
+        out.resize(n);
+        for (int i = 0; i < n; ++i) out[i] = a + (b - a) * i / (double)(n - 1);
+    };
+    lin(c[0] - sar.half_width_m[0], c[0] + sar.half_width_m[0], sar.grid[0], g.x);
+    lin(c[1] - sar.half_width_m[1], c[1] + sar.half_width_m[1], sar.grid[1], g.y);
+    lin(c[2] - sar.half_width_m[2], c[2] + sar.half_width_m[2], sar.grid[2], g.z);
+    return g;
+}
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/src/kernel_src.hpp

@@ -0,0 +1,72 @@
+// Исходник CUDA-ядра для NVRTC (порт bp_aperture_contribution из gpr_core/kernels.py).
+// Self-contained complex<float> (без cupy-заголовков).
+#pragma once
+
+namespace gpr {
+
+inline const char* APERTURE_KERNEL_SRC = R"CUDA(
+struct cf { float x, y; };
+__device__ __forceinline__ cf operator+(cf a, cf b){ return {a.x+b.x, a.y+b.y}; }
+__device__ __forceinline__ cf operator-(cf a, cf b){ return {a.x-b.x, a.y-b.y}; }
+__device__ __forceinline__ cf operator*(float s, cf a){ return {s*a.x, s*a.y}; }
+
+// Весь взвешенный объём апертуры за один запуск: один поток на воксель,
+// внутренний цикл по n_contrib = n_frames*n_pairs. acc/cnt пишутся один раз.
+extern "C" __global__ void bp_aperture_contribution(
+    const cf* __restrict__ H,
+    const float t0, const float inv_dt, const int n_fft, const float inv_v,
+    const float* __restrict__ xt, const float* __restrict__ yt, const float* __restrict__ zt,
+    const float* __restrict__ xr, const float* __restrict__ yr, const float* __restrict__ zr,
+    const float* __restrict__ inv_geo_ref, const float* __restrict__ wgt,
+    const float* __restrict__ gx, const float* __restrict__ gy, const float* __restrict__ gz,
+    const int NX, const int NY, const int NZ, const int n_contrib,
+    const float cos_p, const float sin_p,
+    const float min_depth, const float max_depth,
+    const float comp_pow, const float w_max, const int use_comp,
+    cf* __restrict__ acc, float* __restrict__ cnt)
+{
+    const long total = (long)NX * NY * NZ;
+    for (long gid = blockIdx.x * blockDim.x + threadIdx.x;
+         gid < total; gid += (long)gridDim.x * blockDim.x) {
+        int ix = gid % NX;
+        int iz = (gid / NX) % NZ;
+        int iy = gid / ((long)NX * NZ);
+        float x = gx[ix];
+        float ys = gy[iy], zs = gz[iz];
+        float y_rad = ys * cos_p + zs * sin_p;
+        float z_rad = -ys * sin_p + zs * cos_p;
+        bool depth_ok = (z_rad >= min_depth) && (z_rad <= max_depth);
+
+        cf a = {0.f, 0.f};
+        float c = 0.f;
+        for (int i = 0; i < n_contrib; ++i) {
+            float w = wgt[i];
+            float dxt = x - xt[i], dyt = y_rad - yt[i], dzt = z_rad - zt[i];
+            float dxr = x - xr[i], dyr = y_rad - yr[i], dzr = z_rad - zr[i];
+            float Rt = sqrtf(dxt*dxt + dyt*dyt + dzt*dzt);
+            float Rr = sqrtf(dxr*dxr + dyr*dyr + dzr*dzr);
+            float pos = ((Rt + Rr) * inv_v - t0) * inv_dt;
+            bool inside = (pos >= 0.f) && (pos <= (float)(n_fft - 1));
+            if (!inside) continue;
+            int k = (int)floorf(pos);
+            if (k < 0) k = 0;
+            if (k > n_fft - 2) k = n_fft - 2;
+            float f = pos - (float)k;
+            const cf* h = H + (long)i * n_fft;
+            cf hv = h[k] + f * (h[k+1] - h[k]);
+            if (use_comp) {
+                float ratio = (1.f/(Rt*Rr + 1e-12f) + 1e-30f) * inv_geo_ref[i];
+                float denom = (comp_pow == 0.25f) ? sqrtf(sqrtf(ratio)) : powf(ratio, comp_pow);
+                float ww = 1.f/(denom + 1e-12f);
+                ww = fminf(fmaxf(ww, 0.f), w_max);
+                hv = ww * hv;
+            }
+            if (depth_ok) { a = a + (w * hv); c += w; }
+        }
+        acc[gid] = a;
+        cnt[gid] = c;
+    }
+}
+)CUDA";
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/src/nvrtc_kernel.cpp

@@ -0,0 +1,48 @@
+#include "gpr/nvrtc_kernel.hpp"
+
+#include <nvrtc.h>
+
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+#include "gpr/device.hpp"
+
+namespace gpr {
+
+static void nvrtc_check(nvrtcResult r, const char* what) {
+    if (r != NVRTC_SUCCESS)
+        throw std::runtime_error(std::string("NVRTC error in ") + what + ": " +
+                                 nvrtcGetErrorString(r));
+}
+
+NvrtcKernel::~NvrtcKernel() {
+    if (module_) cuModuleUnload(module_);
+}
+
+void NvrtcKernel::compile(const char* src, const std::string& arch_flag, const char* func_name) {
+    nvrtcProgram prog;
+    nvrtc_check(nvrtcCreateProgram(&prog, src, "gpr_kernel.cu", 0, nullptr, nullptr),
+                "createProgram");
+    std::vector<const char*> opts{arch_flag.c_str()};
+    nvrtcResult cr = nvrtcCompileProgram(prog, (int)opts.size(), opts.data());
+    if (cr != NVRTC_SUCCESS) {
+        size_t logsz = 0;
+        nvrtcGetProgramLogSize(prog, &logsz);
+        std::string log(logsz, '\0');
+        if (logsz) nvrtcGetProgramLog(prog, &log[0]);
+        nvrtcDestroyProgram(&prog);
+        throw std::runtime_error("NVRTC compile failed:\n" + log);
+    }
+    // cubin под точную арх (без JIT PTX — драйвер грузит SASS напрямую)
+    size_t cubsz = 0;
+    nvrtc_check(nvrtcGetCUBINSize(prog, &cubsz), "getCUBINSize");
+    std::string cubin(cubsz, '\0');
+    nvrtc_check(nvrtcGetCUBIN(prog, &cubin[0]), "getCUBIN");
+    nvrtcDestroyProgram(&prog);
+
+    cu_check(cuModuleLoadDataEx(&module_, cubin.data(), 0, nullptr, nullptr), "moduleLoad");
+    cu_check(cuModuleGetFunction(&func_, module_, func_name), "getFunction");
+}
+
+}  // namespace gpr

data_acq_and_processing/gpr_cpp/src/sar3d.cpp

@@ -0,0 +1,111 @@
+#include "gpr/sar3d.hpp"
+
+#include <algorithm>
+#include <cmath>
+#include <stdexcept>
+
+#include "kernel_src.hpp"
+
+namespace gpr {
+
+SarEngine::SarEngine(const GeometryConfig& g, const BPConfig& bp, const MotionConfig& mot,
+                     const SARConfig& sar)
+    : geom_(g), bp_(bp), mot_(mot), sar_(sar) {
+    geom_.finalize();
+    ctx_ = std::make_unique<CudaContext>();
+    kernel_ = std::make_unique<NvrtcKernel>();
+    kernel_->compile(APERTURE_KERNEL_SRC, ctx_->arch_flag(), "bp_aperture_contribution");
+}
+
+VolumeRaw SarEngine::aperture_volume(const std::vector<ApertureFrame>& frames, const Grid& grid) {
+    if (frames.empty()) throw std::runtime_error("aperture_volume: нет кадров");
+    auto pairs = geom_.pair_keys();
+    int n_pairs = (int)pairs.size();
+    int n_frames = (int)frames.size();
+    int n_contrib = n_frames * n_pairs;
+
+    const AScanSet& a0 = *frames[0].ascans;
+    int n_fft = a0.n_fft;
+    double t0 = a0.t0, dt = a0.dt;
+
+    int NX = grid.NX(), NY = grid.NY(), NZ = grid.NZ();
+    size_t nvox = (size_t)NX * NY * NZ;
+
+    // Плоские массивы по (кадр, пара)
+    std::vector<c64> H_stack((size_t)n_contrib * n_fft);
+    std::vector<float> xt(n_contrib), yt(n_contrib), zt(n_contrib);
+    std::vector<float> xr(n_contrib), yr(n_contrib), zr(n_contrib);
+    std::vector<float> igr(n_contrib), wgt(n_contrib);
+    int idx = 0;
+    for (int f = 0; f < n_frames; ++f) {
+        const AScanSet& a = *frames[f].ascans;
+        if (a.n_fft != n_fft || std::abs(a.t0 - t0) > 1e-18)
+            throw std::runtime_error("кадры апертуры: разная ось времени");
+        double dx = frames[f].offset.dx_m, dy = frames[f].offset.dy_m, dz = frames[f].offset.dz_m;
+        float w = (float)frames[f].weight;
+        for (int p = 0; p < n_pairs; ++p) {
+            int i = pairs[p].first, j = pairs[p].second;
+            // A-скан пары p → в стек (каст c128→c64)
+            const c128* src = a.pair_ptr(p);
+            c64* dst = H_stack.data() + (size_t)idx * n_fft;
+            for (int k = 0; k < n_fft; ++k) dst[k] = c64((float)src[k].real(), (float)src[k].imag());
+            xt[idx] = (float)(geom_.x_tx[i] + dx); yt[idx] = (float)(geom_.y_tx[i] + dy);
+            zt[idx] = (float)(geom_.z_tx[i] + dz);
+            xr[idx] = (float)(geom_.x_rx[j] + dx); yr[idx] = (float)(geom_.y_rx[j] + dy);
+            zr[idx] = (float)(geom_.z_rx[j] + dz);
+            igr[idx] = (float)(1.0 / pair_geo_ref(geom_, bp_, i, j, geom_.bp_plane_y));
+            wgt[idx] = w;
+            ++idx;
+        }
+    }
+
+    std::vector<float> gx(grid.x.begin(), grid.x.end());
+    std::vector<float> gy(grid.y.begin(), grid.y.end());
+    std::vector<float> gz(grid.z.begin(), grid.z.end());
+
+    float cos_p = 1.f, sin_p = 0.f;
+    if (sar_.coord_mode_ground) {
+        double p = mot_.pitch_deg * M_PI / 180.0;
+        cos_p = (float)std::cos(p); sin_p = (float)std::sin(p);
+    }
+
+    // Заливка на устройство
+    DeviceBuffer<c64> dH(H_stack.size()); dH.upload(H_stack.data(), H_stack.size());
+    DeviceBuffer<float> dxt(n_contrib), dyt(n_contrib), dzt(n_contrib);
+    DeviceBuffer<float> dxr(n_contrib), dyr(n_contrib), dzr(n_contrib);
+    DeviceBuffer<float> digr(n_contrib), dwgt(n_contrib);
+    dxt.upload(xt.data(), n_contrib); dyt.upload(yt.data(), n_contrib); dzt.upload(zt.data(), n_contrib);
+    dxr.upload(xr.data(), n_contrib); dyr.upload(yr.data(), n_contrib); dzr.upload(zr.data(), n_contrib);
+    digr.upload(igr.data(), n_contrib); dwgt.upload(wgt.data(), n_contrib);
+    DeviceBuffer<float> dgx(NX), dgy(NY), dgz(NZ);
+    dgx.upload(gx.data(), NX); dgy.upload(gy.data(), NY); dgz.upload(gz.data(), NZ);
+    DeviceBuffer<c64> dacc(nvox);
+    DeviceBuffer<float> dcnt(nvox);
+
+    float f_t0 = (float)t0, f_inv_dt = (float)(1.0 / dt), f_inv_v = (float)(1.0 / geom_.v);
+    float min_d = (float)bp_.min_depth, max_d = (float)bp_.max_depth;
+    float comp_pow = (float)bp_.comp_range_power, w_max = (float)bp_.comp_weight_max;
+    int use_comp = sar_.use_compensation ? 1 : 0;
+
+    void* args[] = {
+        dH.addr(), &f_t0, &f_inv_dt, &n_fft, &f_inv_v,
+        dxt.addr(), dyt.addr(), dzt.addr(), dxr.addr(), dyr.addr(), dzr.addr(),
+        digr.addr(), dwgt.addr(), dgx.addr(), dgy.addr(), dgz.addr(),
+        &NX, &NY, &NZ, &n_contrib, &cos_p, &sin_p, &min_d, &max_d,
+        &comp_pow, &w_max, &use_comp, dacc.addr(), dcnt.addr()};
+
+    int block = 256;
+    int gridDim = (int)std::min<long>(65535, ((long)nvox + block - 1) / block);
+    cu_check(cuLaunchKernel(kernel_->func(), gridDim, 1, 1, block, 1, 1, 0, 0, args, nullptr),
+             "launch aperture");
+    cu_check(cuCtxSynchronize(), "sync");
+
+    VolumeRaw out;
+    out.NX = NX; out.NY = NY; out.NZ = NZ;
+    out.acc.resize(nvox); out.cnt.resize(nvox);
+    dacc.download(out.acc.data(), nvox);
+    dcnt.download(out.cnt.data(), nvox);
+    return out;
+}
+
+}  // namespace gpr

data_acq_and_processing/processing/data_processor/include/processing_live_config.hpp

@@ -90,6 +90,25 @@ struct ProcessingLiveConfig {
     float gpr_visible_x_max_m = 2.0F;
     float gpr_visible_z_min_m = 0.0F;
     float gpr_visible_z_max_m = 14.0F;
+    // 3D SAR (sar3d processor) parameters; defaults mirror VNA_AlGo gpr_core.
+    // Frequency band, depth gate, antenna selection and motion (speed, look
+    // angle, direction, phase correction) are shared with the gpr_* fields.
+    std::uint32_t sar3d_frame_radius = 6U;
+    float sar3d_edge_frame_weight = 0.75F;
+    float sar3d_pitch_deg = 6.0F;
+    float sar3d_center_depth_m = 5.6F;
+    float sar3d_half_width_x_m = 1.8F;
+    float sar3d_half_width_y_m = 1.0F;
+    float sar3d_half_width_z_m = 1.8F;
+    std::uint32_t sar3d_grid_nx = 150U;
+    std::uint32_t sar3d_grid_ny = 120U;
+    std::uint32_t sar3d_grid_nz = 160U;
+    bool sar3d_use_compensation = true;
+    bool sar3d_normalize_pairs = true;
+    // Bottom-view projection angle around the vertical Y axis, degrees:
+    // 0 = front view (X horizontal), 90 = side view (Z horizontal),
+    // intermediate values project along the rotated horizontal direction.
+    float sar3d_bottom_view_deg = 0.0F;
     // When true, the data_processor ignores socket-supplied `vlc` updates and
     // keeps using `gpr_speed_m_s` from this file. Mirrored from the GUI's
     // "ignore socket speed" checkbox.

data_acq_and_processing/processing/data_processor/src/data_processor.cpp

@@ -13,6 +13,7 @@
 #include "bscan_processor.hpp"
 #include "gpr_processor.hpp"
 #include "passthrough_processor.hpp"
+#include "sar3d_processor.hpp"
 
 namespace radar::processing {
 namespace {
@@ -280,6 +281,10 @@ auto create_default_processors() -> ProcessorRegistry {
         auto processor = std::make_unique<LegacyGprProcessor>();
         processors.emplace(processor->name(), std::move(processor));
     }
+    {
+        auto processor = std::make_unique<Sar3dProcessor>();
+        processors.emplace(processor->name(), std::move(processor));
+    }
     return processors;
 }
 

data_acq_and_processing/processing/data_processor/src/processing_live_config.cpp

@@ -386,6 +386,46 @@ void apply_legacy_gpr_algorithm_alias(ProcessingLiveConfig& config, const std::s
         config.gpr_object_approach_min_frames =
             parse_u32_number(*found, "processing.gpr_object_approach_min_frames");
     }
+    if (const auto found = root.find("sar3d_frame_radius"); found != root.end()) {
+        config.sar3d_frame_radius = parse_u32_number(*found, "processing.sar3d_frame_radius");
+    }
+    for (const auto& [key, target] : {
+             std::pair{"sar3d_edge_frame_weight", &config.sar3d_edge_frame_weight},
+             std::pair{"sar3d_pitch_deg", &config.sar3d_pitch_deg},
+             std::pair{"sar3d_center_depth_m", &config.sar3d_center_depth_m},
+             std::pair{"sar3d_half_width_x_m", &config.sar3d_half_width_x_m},
+             std::pair{"sar3d_half_width_y_m", &config.sar3d_half_width_y_m},
+             std::pair{"sar3d_half_width_z_m", &config.sar3d_half_width_z_m},
+             std::pair{"sar3d_bottom_view_deg", &config.sar3d_bottom_view_deg},
+         }) {
+        if (const auto found = root.find(key); found != root.end()) {
+            if (!found->is_number()) {
+                throw std::runtime_error(std::string("processing.") + key + " must be number");
+            }
+            *target = static_cast<float>(found->get<double>());
+        }
+    }
+    for (const auto& [key, target] : {
+             std::pair{"sar3d_grid_nx", &config.sar3d_grid_nx},
+             std::pair{"sar3d_grid_ny", &config.sar3d_grid_ny},
+             std::pair{"sar3d_grid_nz", &config.sar3d_grid_nz},
+         }) {
+        if (const auto found = root.find(key); found != root.end()) {
+            *target = parse_u32_number(*found, std::string("processing.") + key);
+        }
+    }
+    if (const auto found = root.find("sar3d_use_compensation"); found != root.end()) {
+        if (!found->is_boolean()) {
+            throw std::runtime_error("processing.sar3d_use_compensation must be bool");
+        }
+        config.sar3d_use_compensation = found->get<bool>();
+    }
+    if (const auto found = root.find("sar3d_normalize_pairs"); found != root.end()) {
+        if (!found->is_boolean()) {
+            throw std::runtime_error("processing.sar3d_normalize_pairs must be bool");
+        }
+        config.sar3d_normalize_pairs = found->get<bool>();
+    }
     if (const auto found = root.find("ignore_socket_speed"); found != root.end()) {
         if (!found->is_boolean()) {
             throw std::runtime_error("processing.ignore_socket_speed must be bool");

data_acq_and_processing/processing/processors/include/sar3d_processor.hpp

@@ -0,0 +1,34 @@
+#pragma once
+
+#include <memory>
+#include <span>
+#include <string>
+
+#include "processor_interface.hpp"
+
+namespace radar::processing {
+
+// 3D SAR over a synthesized aperture, computed by the external gpr_cpp
+// CUDA/NVRTC library (VNA_AlGo). Emits 2D max-projections of the volume
+// (side XZ and top XY views) as ImageF32 collection payloads.
+class Sar3dProcessor final : public ProcessorInterface {
+  public:
+    Sar3dProcessor();
+    ~Sar3dProcessor() override;
+
+    [[nodiscard]] auto name() const -> std::string override;
+    [[nodiscard]] auto process_collection(
+        const config::RunConfig& run_config,
+        const ipc::PreprocessedCollection& collection,
+        std::span<const ipc::PreprocessedCollection> previous_collections,
+        const ProcessingLiveConfig& live_config
+    ) -> ipc::ResultCollection override;
+
+  private:
+    // Pimpl keeps gpr/CUDA headers (cuda.h, nvrtc.h) out of the shared
+    // processor include path; only sar3d_processor.cpp needs them.
+    struct State;
+    std::unique_ptr<State> state_;
+};
+
+}  // namespace radar::processing

data_acq_and_processing/processing/processors/src/sar3d_processor.cpp

@@ -0,0 +1,568 @@
+#include "sar3d_processor.hpp"
+
+#include <algorithm>
+#include <bit>
+#include <cmath>
+#include <cstdint>
+#include <iostream>
+#include <map>
+#include <optional>
+#include <stdexcept>
+#include <vector>
+
+#include "gpr/ascan.hpp"
+#include "gpr/config.hpp"
+#include "gpr/geometry.hpp"
+#include "gpr/sar3d.hpp"
+
+namespace radar::processing {
+namespace {
+
+constexpr std::size_t kAscanCacheLimit = 64U;
+constexpr double kVolumeEps = 1e-12;
+
+void hash_mix(std::size_t& seed, std::uint64_t value) {
+    seed ^= value + 0x9e3779b97f4a7c15ULL + (seed << 6U) + (seed >> 2U);
+}
+
+void hash_mix(std::size_t& seed, double value) {
+    hash_mix(seed, std::bit_cast<std::uint64_t>(value));
+}
+
+[[nodiscard]] auto contains_u32(const std::vector<std::uint32_t>& values, std::uint32_t value) -> bool {
+    return std::find(values.begin(), values.end(), value) != values.end();
+}
+
+// Requested positions filtered to the available set; empty request selects all.
+// Mirrors the coherent-BP processor's selection semantics.
+[[nodiscard]] auto selected_positions(
+    const std::vector<std::uint32_t>& requested,
+    std::vector<std::uint32_t> available
+) -> std::vector<std::uint32_t> {
+    std::sort(available.begin(), available.end());
+    available.erase(std::unique(available.begin(), available.end()), available.end());
+    if (requested.empty()) {
+        return available;
+    }
+    std::vector<std::uint32_t> result{};
+    for (const auto value : requested) {
+        if (contains_u32(available, value) && !contains_u32(result, value)) {
+            result.push_back(value);
+        }
+    }
+    std::sort(result.begin(), result.end());
+    return result;
+}
+
+struct Selection {
+    // Sorted ascending switch positions; the logical pair key (i_tx, i_rx) is
+    // the index of a position within these lists.
+    std::vector<std::uint32_t> inputs{};
+    std::vector<std::uint32_t> outputs{};
+};
+
+[[nodiscard]] auto build_selection(
+    const config::RunConfig& run_config,
+    const ProcessingLiveConfig& live_config
+) -> Selection {
+    std::vector<std::uint32_t> available_inputs{};
+    available_inputs.reserve(run_config.gpr.rx_geometry.size());
+    for (const auto& entry : run_config.gpr.rx_geometry) {
+        available_inputs.push_back(entry.input_pos);
+    }
+    std::vector<std::uint32_t> available_outputs{};
+    available_outputs.reserve(run_config.gpr.tx_geometry.size());
+    for (const auto& entry : run_config.gpr.tx_geometry) {
+        available_outputs.push_back(entry.output_pos);
+    }
+    Selection selection{};
+    selection.inputs = selected_positions(live_config.gpr_input_positions, std::move(available_inputs));
+    selection.outputs = selected_positions(live_config.gpr_output_positions, std::move(available_outputs));
+    return selection;
+}
+
+[[nodiscard]] auto build_geometry(
+    const config::RunConfig& run_config,
+    const Selection& selection
+) -> gpr::GeometryConfig {
+    gpr::GeometryConfig geom{};
+    geom.input_idx.clear();
+    geom.output_idx.clear();
+    geom.rx_positions.clear();
+    geom.tx_positions.clear();
+    for (const auto position : selection.inputs) {
+        for (const auto& entry : run_config.gpr.rx_geometry) {
+            if (entry.input_pos == position) {
+                geom.input_idx.push_back(static_cast<int>(position));
+                geom.rx_positions.push_back({
+                    static_cast<double>(entry.x_m),
+                    static_cast<double>(entry.y_m),
+                    static_cast<double>(entry.z_m),
+                });
+                break;
+            }
+        }
+    }
+    for (const auto position : selection.outputs) {
+        for (const auto& entry : run_config.gpr.tx_geometry) {
+            if (entry.output_pos == position) {
+                geom.output_idx.push_back(static_cast<int>(position));
+                geom.tx_positions.push_back({
+                    static_cast<double>(entry.x_m),
+                    static_cast<double>(entry.y_m),
+                    static_cast<double>(entry.z_m),
+                });
+                break;
+            }
+        }
+    }
+    geom.eps_r = static_cast<double>(run_config.gpr.relative_permittivity);
+    geom.finalize();
+    return geom;
+}
+
+[[nodiscard]] auto build_bp_config(const ProcessingLiveConfig& live_config) -> gpr::BPConfig {
+    gpr::BPConfig bp{};
+    bp.f_start = static_cast<double>(live_config.gpr_start_freq_mhz) * 1e6;
+    bp.f_stop = static_cast<double>(live_config.gpr_stop_freq_mhz) * 1e6;
+    bp.min_depth = static_cast<double>(live_config.gpr_min_depth_m);
+    bp.max_depth = static_cast<double>(live_config.gpr_max_depth_m);
+    bp.pair_normalize = live_config.sar3d_normalize_pairs;
+    bp.pair_norm_depth_min = bp.min_depth;
+    bp.pair_norm_depth_max = bp.max_depth;
+    return bp;
+}
+
+[[nodiscard]] auto build_motion_config(const ProcessingLiveConfig& live_config) -> gpr::MotionConfig {
+    gpr::MotionConfig motion{};
+    motion.speed_m_s = static_cast<double>(live_config.gpr_speed_m_s);
+    motion.look_angle_deg = static_cast<double>(live_config.gpr_look_angle_deg);
+    motion.pitch_deg = static_cast<double>(live_config.sar3d_pitch_deg);
+    motion.direction_sign = static_cast<double>(live_config.gpr_direction_sign);
+    motion.apply_freq_phase_correction = live_config.gpr_apply_freq_phase_correction;
+    return motion;
+}
+
+[[nodiscard]] auto build_sar_config(const ProcessingLiveConfig& live_config) -> gpr::SARConfig {
+    gpr::SARConfig sar{};
+    sar.frame_radius = static_cast<int>(std::max<std::uint32_t>(0U, live_config.sar3d_frame_radius));
+    sar.edge_frame_weight = static_cast<double>(live_config.sar3d_edge_frame_weight);
+    sar.half_width_m = {
+        static_cast<double>(live_config.sar3d_half_width_x_m),
+        static_cast<double>(live_config.sar3d_half_width_y_m),
+        static_cast<double>(live_config.sar3d_half_width_z_m),
+    };
+    sar.grid = {
+        static_cast<int>(std::max<std::uint32_t>(2U, live_config.sar3d_grid_nx)),
+        static_cast<int>(std::max<std::uint32_t>(2U, live_config.sar3d_grid_ny)),
+        static_cast<int>(std::max<std::uint32_t>(2U, live_config.sar3d_grid_nz)),
+    };
+    sar.use_compensation = live_config.sar3d_use_compensation;
+    return sar;
+}
+
+[[nodiscard]] auto ascan_fingerprint(
+    const Selection& selection,
+    const gpr::GeometryConfig& geom,
+    const gpr::BPConfig& bp,
+    const gpr::MotionConfig& motion
+) -> std::size_t {
+    std::size_t seed = 0U;
+    for (const auto position : selection.inputs) {
+        hash_mix(seed, static_cast<std::uint64_t>(position));
+    }
+    for (const auto position : selection.outputs) {
+        hash_mix(seed, static_cast<std::uint64_t>(static_cast<std::uint64_t>(position) | (1ULL << 40U)));
+    }
+    hash_mix(seed, geom.eps_r);
+    hash_mix(seed, bp.f_start);
+    hash_mix(seed, bp.f_stop);
+    hash_mix(seed, bp.pair_norm_depth_min);
+    hash_mix(seed, bp.pair_norm_depth_max);
+    hash_mix(seed, static_cast<std::uint64_t>(bp.pair_normalize ? 1U : 0U));
+    hash_mix(seed, motion.speed_m_s);
+    hash_mix(seed, motion.look_angle_deg);
+    hash_mix(seed, motion.direction_sign);
+    hash_mix(seed, static_cast<std::uint64_t>(motion.apply_freq_phase_correction ? 1U : 0U));
+    return seed;
+}
+
+[[nodiscard]] auto engine_fingerprint(
+    const gpr::GeometryConfig& geom,
+    const gpr::BPConfig& bp,
+    const gpr::MotionConfig& motion,
+    const gpr::SARConfig& sar
+) -> std::size_t {
+    std::size_t seed = 1U;
+    for (std::size_t index = 0U; index < geom.tx_positions.size(); ++index) {
+        hash_mix(seed, static_cast<std::uint64_t>(geom.output_idx[index]));
+        hash_mix(seed, geom.tx_positions[index][0]);
+        hash_mix(seed, geom.tx_positions[index][1]);
+        hash_mix(seed, geom.tx_positions[index][2]);
+    }
+    for (std::size_t index = 0U; index < geom.rx_positions.size(); ++index) {
+        hash_mix(seed, static_cast<std::uint64_t>(geom.input_idx[index]));
+        hash_mix(seed, geom.rx_positions[index][0]);
+        hash_mix(seed, geom.rx_positions[index][1]);
+        hash_mix(seed, geom.rx_positions[index][2]);
+    }
+    hash_mix(seed, geom.eps_r);
+    hash_mix(seed, bp.f_start);
+    hash_mix(seed, bp.f_stop);
+    hash_mix(seed, bp.min_depth);
+    hash_mix(seed, bp.max_depth);
+    hash_mix(seed, bp.comp_range_power);
+    hash_mix(seed, bp.comp_weight_max);
+    hash_mix(seed, bp.comp_ref_depth);
+    hash_mix(seed, motion.pitch_deg);
+    hash_mix(seed, static_cast<std::uint64_t>(sar.use_compensation ? 1U : 0U));
+    return seed;
+}
+
+// Midpoint of the collection's capture window; monotonic_ns is the fallback for
+// producers that never stamped capture times.
+[[nodiscard]] auto collection_time_s(const ipc::PreprocessedCollection& collection) -> double {
+    if (collection.capture_end_ns > collection.capture_start_ns) {
+        return 0.5 * (static_cast<double>(collection.capture_start_ns) + static_cast<double>(collection.capture_end_ns)) * 1e-9;
+    }
+    return static_cast<double>(collection.monotonic_ns) * 1e-9;
+}
+
+// One Tx's interleaved sweep spans the whole capture window, so the per-Tx
+// sweep time is capture_span / N_tx (same derivation as the coherent BP path).
+[[nodiscard]] auto tx_sweep_time_s(const ipc::PreprocessedCollection& collection, int n_tx) -> double {
+    if (collection.capture_end_ns <= collection.capture_start_ns || n_tx <= 0) {
+        return 0.0;
+    }
+    const double span_s = static_cast<double>(collection.capture_end_ns - collection.capture_start_ns) * 1e-9;
+    return span_s / static_cast<double>(n_tx);
+}
+
+[[nodiscard]] auto build_image_payload(
+    const std::string& processing_name,
+    const std::vector<double>& x_axis,
+    const std::vector<double>& y_axis,
+    std::vector<float> values
+) -> ipc::ResultPayload {
+    ipc::ResultPayload payload{};
+    payload.processing_name = processing_name;
+    payload.kind = ipc::ResultKind::ImageF32;
+    payload.image_x_axis.reserve(x_axis.size());
+    for (const auto value : x_axis) {
+        payload.image_x_axis.push_back(static_cast<float>(value));
+    }
+    payload.image_y_axis.reserve(y_axis.size());
+    for (const auto value : y_axis) {
+        payload.image_y_axis.push_back(static_cast<float>(value));
+    }
+    payload.image_values = std::move(values);
+    return payload;
+}
+
+}  // namespace
+
+struct Sar3dProcessor::State {
+    struct CacheEntry {
+        double time_s = 0.0;
+        // Null when the collection could not be preprocessed (e.g. missing
+        // pairs); cached so a bad collection is not retried every frame.
+        std::shared_ptr<gpr::AScanSet> ascans{};
+    };
+
+    // Keyed by collection_id (monotonically increasing), so evicting from the
+    // front drops the oldest collections.
+    std::map<std::uint64_t, CacheEntry> ascan_cache{};
+    std::size_t ascan_fp = 0U;
+
+    std::unique_ptr<gpr::SarEngine> engine{};
+    std::size_t engine_fp = 0U;
+
+    std::string last_error{};
+    std::uint64_t suppressed_errors = 0U;
+
+    auto get_or_compute_ascans(
+        const ipc::PreprocessedCollection& collection,
+        const Selection& selection,
+        const gpr::GeometryConfig& geom,
+        const gpr::BPConfig& bp,
+        gpr::MotionConfig motion
+    ) -> const CacheEntry& {
+        if (const auto found = ascan_cache.find(collection.collection_id); found != ascan_cache.end()) {
+            return found->second;
+        }
+
+        CacheEntry entry{};
+        entry.time_s = collection_time_s(collection);
+
+        std::map<gpr::PairKey, std::vector<gpr::c128>> s21{};
+        std::map<gpr::PairKey, std::vector<double>> freq{};
+        for (const auto& trace : collection.traces) {
+            const auto input_it = std::find(selection.inputs.begin(), selection.inputs.end(), trace.combo.input_pos);
+            const auto output_it = std::find(selection.outputs.begin(), selection.outputs.end(), trace.combo.output_pos);
+            if (input_it == selection.inputs.end() || output_it == selection.outputs.end()) {
+                continue;
+            }
+            const gpr::PairKey key{
+                static_cast<int>(output_it - selection.outputs.begin()),
+                static_cast<int>(input_it - selection.inputs.begin()),
+            };
+            auto& samples = s21[key];
+            samples.clear();
+            samples.reserve(trace.s21.size());
+            for (const auto& sample : trace.s21) {
+                samples.emplace_back(static_cast<double>(sample.re), static_cast<double>(sample.im));
+            }
+            auto& axis = freq[key];
+            axis.assign(trace.frequency_hz.begin(), trace.frequency_hz.end());
+        }
+
+        if (s21.size() == selection.inputs.size() * selection.outputs.size()) {
+            motion.tx_sweep_time_s = tx_sweep_time_s(collection, geom.n_tx);
+            const auto timestamps = gpr::compute_pair_timestamps(motion, geom);
+            try {
+                entry.ascans = std::make_shared<gpr::AScanSet>(gpr::preprocess_snapshot(
+                    s21, freq, /*background=*/nullptr, motion, geom, bp, timestamps, bp.pair_normalize
+                ));
+            } catch (const std::exception& exc) {
+                // A malformed collection (e.g. empty band after a sweep change)
+                // is cached as a skipped frame instead of failing the aperture.
+                report_error(std::string{"preprocess_snapshot: "} + exc.what());
+                entry.ascans = nullptr;
+            }
+        }
+
+        auto [inserted, _] = ascan_cache.emplace(collection.collection_id, std::move(entry));
+        while (ascan_cache.size() > kAscanCacheLimit) {
+            ascan_cache.erase(ascan_cache.begin());
+        }
+        return inserted->second;
+    }
+
+    void report_error(const std::string& message) {
+        if (message == last_error) {
+            ++suppressed_errors;
+            return;
+        }
+        if (suppressed_errors > 0U) {
+            std::cerr << "sar3d: previous error repeated " << suppressed_errors << " times\n";
+        }
+        std::cerr << "sar3d: " << message << '\n';
+        last_error = message;
+        suppressed_errors = 0U;
+    }
+};
+
+Sar3dProcessor::Sar3dProcessor() : state_(std::make_unique<State>()) {}
+
+Sar3dProcessor::~Sar3dProcessor() = default;
+
+auto Sar3dProcessor::name() const -> std::string {
+    return "sar3d";
+}
+
+auto Sar3dProcessor::process_collection(
+    const config::RunConfig& run_config,
+    const ipc::PreprocessedCollection& collection,
+    std::span<const ipc::PreprocessedCollection> previous_collections,
+    const ProcessingLiveConfig& live_config
+) -> ipc::ResultCollection {
+    ipc::ResultCollection results{};
+    results.collection_id = collection.collection_id;
+    results.monotonic_ns = collection.monotonic_ns;
+
+    try {
+        const auto selection = build_selection(run_config, live_config);
+        if (selection.inputs.empty() || selection.outputs.empty()) {
+            return results;
+        }
+        const auto geom = build_geometry(run_config, selection);
+        const auto bp = build_bp_config(live_config);
+        const auto motion = build_motion_config(live_config);
+        const auto sar = build_sar_config(live_config);
+        if (!(bp.max_depth > bp.min_depth) || !(bp.f_stop > bp.f_start)) {
+            return results;
+        }
+
+        const auto ascan_fp = ascan_fingerprint(selection, geom, bp, motion);
+        if (ascan_fp != state_->ascan_fp) {
+            state_->ascan_cache.clear();
+            state_->ascan_fp = ascan_fp;
+        }
+
+        // Aperture window: the newest 2r+1 collections, current one last.
+        const std::size_t window_size = 2U * static_cast<std::size_t>(sar.frame_radius) + 1U;
+        std::vector<const ipc::PreprocessedCollection*> window{};
+        window.reserve(window_size);
+        const std::size_t history_take = std::min(previous_collections.size(), window_size - 1U);
+        for (std::size_t index = previous_collections.size() - history_take; index < previous_collections.size(); ++index) {
+            window.push_back(&previous_collections[index]);
+        }
+        window.push_back(&collection);
+
+        std::vector<std::shared_ptr<gpr::AScanSet>> ascan_sets{};
+        std::vector<double> frame_times{};
+        ascan_sets.reserve(window.size());
+        frame_times.reserve(window.size());
+        for (const auto* entry_collection : window) {
+            const auto& entry = state_->get_or_compute_ascans(*entry_collection, selection, geom, bp, motion);
+            if (entry.ascans == nullptr) {
+                continue;
+            }
+            ascan_sets.push_back(entry.ascans);
+            frame_times.push_back(entry.time_s);
+        }
+        if (ascan_sets.empty()) {
+            return results;
+        }
+
+        // aperture_volume requires one shared time axis; drop frames whose axis
+        // differs from the current collection's (e.g. after a sweep-grid change).
+        const auto& reference = *ascan_sets.back();
+        for (std::size_t index = ascan_sets.size(); index-- > 0U;) {
+            const auto& candidate = *ascan_sets[index];
+            if (candidate.n_fft != reference.n_fft || std::abs(candidate.t0 - reference.t0) > 1e-18 ||
+                candidate.n_pairs != reference.n_pairs) {
+                ascan_sets.erase(ascan_sets.begin() + static_cast<std::ptrdiff_t>(index));
+                frame_times.erase(frame_times.begin() + static_cast<std::ptrdiff_t>(index));
+            }
+        }
+        if (ascan_sets.empty()) {
+            return results;
+        }
+
+        // Symmetric aperture: the reference frame sits `radius` frames before the
+        // newest one, so the published volume lags the input by `radius` frames.
+        const std::size_t center_index =
+            ascan_sets.size() > static_cast<std::size_t>(sar.frame_radius)
+                ? ascan_sets.size() - 1U - static_cast<std::size_t>(sar.frame_radius)
+                : 0U;
+        const auto offsets = gpr::frame_offsets(frame_times, frame_times[center_index], motion);
+        const auto weights = gpr::frame_weights(offsets, sar.edge_frame_weight);
+        std::vector<gpr::ApertureFrame> frames{};
+        frames.reserve(ascan_sets.size());
+        for (std::size_t index = 0U; index < ascan_sets.size(); ++index) {
+            frames.push_back(gpr::ApertureFrame{ascan_sets[index].get(), offsets[index], weights[index]});
+        }
+
+        const auto engine_fp = engine_fingerprint(geom, bp, motion, sar);
+        if (state_->engine == nullptr || state_->engine_fp != engine_fp) {
+            state_->engine = std::make_unique<gpr::SarEngine>(geom, bp, motion, sar);
+            state_->engine_fp = engine_fp;
+        }
+
+        const auto center_scene =
+            gpr::point_to_scene(0.0, 0.0, static_cast<double>(live_config.sar3d_center_depth_m), motion, sar);
+        const gpr::Grid grid = gpr::build_grid_scene(center_scene, sar);
+        const gpr::VolumeRaw volume = state_->engine->aperture_volume(frames, grid);
+
+        // Finalization (port of gpr_core sar_finish_volume, minus smoothing):
+        // |acc/(cnt+eps)|, ground-coordinate depth gate, normalization to peak.
+        //
+        // Two projections match the legacy render conventions (X across,
+        // Z range/"вдаль", Y up):
+        // - top view `sar3d_proj_xz`: max over Y, axes X (horizontal) / Z;
+        // - bottom view `sar3d_proj_bottom`: vertical axis is always Y, the
+        //   horizontal axis u is the ground direction rotated by
+        //   `sar3d_bottom_view_deg` around Y: 0° = front view (u = X, max
+        //   along Z), 90° = side view (u = Z, max along X), intermediate
+        //   angles project along the rotated perpendicular.
+        const int nx = volume.NX;
+        const int ny = volume.NY;
+        const int nz = volume.NZ;
+        const double pitch_rad = motion.pitch_deg * M_PI / 180.0;
+        const double cos_p = std::cos(pitch_rad);
+        const double sin_p = std::sin(pitch_rad);
+
+        const double theta_deg = std::clamp(static_cast<double>(live_config.sar3d_bottom_view_deg), 0.0, 90.0);
+        const double theta_rad = theta_deg * M_PI / 180.0;
+        const double cos_t = std::cos(theta_rad);
+        const double sin_t = std::sin(theta_rad);
+        // Exact-axis fast paths keep the front/side views on the native grid.
+        const bool front_view = theta_deg <= 1e-9;
+        const bool side_view = theta_deg >= 90.0 - 1e-9;
+        std::vector<double> u_axis{};
+        if (front_view) {
+            u_axis = grid.x;
+        } else if (side_view) {
+            u_axis = grid.z;
+        } else {
+            const double u_a = grid.x.front() * cos_t + grid.z.front() * sin_t;
+            const double u_b = grid.x.front() * cos_t + grid.z.back() * sin_t;
+            const double u_c = grid.x.back() * cos_t + grid.z.front() * sin_t;
+            const double u_d = grid.x.back() * cos_t + grid.z.back() * sin_t;
+            const double u_min = std::min(std::min(u_a, u_b), std::min(u_c, u_d));
+            const double u_max = std::max(std::max(u_a, u_b), std::max(u_c, u_d));
+            u_axis.resize(static_cast<std::size_t>(nx));
+            for (int iu = 0; iu < nx; ++iu) {
+                u_axis[static_cast<std::size_t>(iu)] = u_min + (u_max - u_min) * iu / static_cast<double>(nx - 1);
+            }
+        }
+        const int nu = static_cast<int>(u_axis.size());
+        const double u_start = u_axis.front();
+        const double inv_du = nu > 1 ? static_cast<double>(nu - 1) / (u_axis.back() - u_start) : 0.0;
+
+        std::vector<float> projection_xz(static_cast<std::size_t>(nz) * nx, 0.0F);
+        std::vector<float> projection_bottom(static_cast<std::size_t>(ny) * nu, 0.0F);
+        float peak = 0.0F;
+        for (int iy = 0; iy < ny; ++iy) {
+            for (int iz = 0; iz < nz; ++iz) {
+                const double z_rad = -grid.y[iy] * sin_p + grid.z[iz] * cos_p;
+                const bool depth_ok = z_rad >= bp.min_depth && z_rad <= bp.max_depth;
+                if (!depth_ok) {
+                    continue;
+                }
+                const std::size_t row = (static_cast<std::size_t>(iy) * nz + iz) * nx;
+                const std::size_t bottom_row = static_cast<std::size_t>(iy) * nu;
+                const double z_sin_t = grid.z[iz] * sin_t;
+                for (int ix = 0; ix < nx; ++ix) {
+                    const std::size_t voxel = row + static_cast<std::size_t>(ix);
+                    const float value = std::abs(volume.acc[voxel]) /
+                        (volume.cnt[voxel] + static_cast<float>(kVolumeEps));
+                    peak = std::max(peak, value);
+                    auto& xz_cell = projection_xz[static_cast<std::size_t>(iz) * nx + ix];
+                    xz_cell = std::max(xz_cell, value);
+                    int bin = ix;
+                    if (side_view) {
+                        bin = iz;
+                    } else if (!front_view) {
+                        const double u = grid.x[ix] * cos_t + z_sin_t;
+                        bin = static_cast<int>(std::lround((u - u_start) * inv_du));
+                        bin = std::clamp(bin, 0, nu - 1);
+                    }
+                    auto& bottom_cell = projection_bottom[bottom_row + static_cast<std::size_t>(bin)];
+                    bottom_cell = std::max(bottom_cell, value);
+                }
+            }
+        }
+        if (peak > 0.0F) {
+            for (auto& value : projection_xz) {
+                value /= peak;
+            }
+            for (auto& value : projection_bottom) {
+                value /= peak;
+            }
+        }
+
+        results.collection_payloads.push_back(
+            build_image_payload("sar3d_proj_xz", grid.x, grid.z, std::move(projection_xz))
+        );
+        results.collection_payloads.push_back(
+            build_image_payload("sar3d_proj_bottom", u_axis, grid.y, std::move(projection_bottom))
+        );
+        ipc::ResultPayload peak_payload{};
+        peak_payload.processing_name = "sar3d_peak";
+        peak_payload.kind = ipc::ResultKind::ScalarF32;
+        peak_payload.scalar_value = peak;
+        results.collection_payloads.push_back(std::move(peak_payload));
+    } catch (const std::exception& exc) {
+        // A CUDA/NVRTC or data error must not kill the data_processor loop:
+        // publish an empty result and keep going. Repeated identical errors are
+        // logged once with a suppression count.
+        state_->report_error(exc.what());
+        results.collection_payloads.clear();
+    }
+
+    return results;
+}
+
+}  // namespace radar::processing

data_acq_and_processing/processing/tools/sar3d_parity_main.cpp

@@ -0,0 +1,284 @@
+// Parity check for the sar3d processor against the VNA_AlGo reference path.
+//
+// Loads the same .npy snapshot dataset the gpr_cpp `standalone_bp` example
+// uses, feeds it through Sar3dProcessor (combo mapping, timestamps, A-scan
+// cache, aperture, finalization) and prints the volume peak and the argmax of
+// both projections. Compare against `standalone_bp <dir> <center>` output:
+// the peak |acc/cnt| and the (X, Y, Z) of the maximum voxel must match.
+//
+// Usage: sar3d_parity <preprocessed_dir> <center_snapshot_name>
+#include <algorithm>
+#include <cctype>
+#include <complex>
+#include <cstdint>
+#include <cstdio>
+#include <dirent.h>
+#include <fstream>
+#include <span>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+#include "processing_live_config.hpp"
+#include "run_config.hpp"
+#include "sar3d_processor.hpp"
+#include "shared_types.hpp"
+
+namespace {
+
+// Minimal numpy .npy loader (v1.0, 1-D complex128/complex64/float64/float32),
+// mirroring gpr_cpp/examples/standalone_bp.cpp.
+struct NpyArray {
+    std::vector<double> f64{};
+    std::vector<std::complex<double>> c128{};
+};
+
+NpyArray load_npy(const std::string& path) {
+    std::ifstream in(path, std::ios::binary);
+    if (!in) {
+        throw std::runtime_error("cannot open " + path);
+    }
+    char magic[6];
+    in.read(magic, 6);
+    std::uint8_t version[2];
+    in.read(reinterpret_cast<char*>(version), 2);
+    std::uint16_t header_len = 0;
+    in.read(reinterpret_cast<char*>(&header_len), 2);
+    std::string header(header_len, '\0');
+    in.read(header.data(), header_len);
+    const auto shape_pos = header.find("'shape':");
+    const auto left = header.find('(', shape_pos);
+    const auto right = header.find(')', shape_pos);
+    const long count = std::stol(header.substr(left + 1, right - left - 1));
+
+    NpyArray array{};
+    const auto has = [&header](const char* token) { return header.find(token) != std::string::npos; };
+    if (has("c8")) {
+        array.c128.resize(count);
+        std::vector<std::complex<float>> tmp(count);
+        in.read(reinterpret_cast<char*>(tmp.data()), count * sizeof(std::complex<float>));
+        for (long i = 0; i < count; ++i) {
+            array.c128[i] = std::complex<double>(tmp[i].real(), tmp[i].imag());
+        }
+    } else if (has("c16")) {
+        array.c128.resize(count);
+        in.read(reinterpret_cast<char*>(array.c128.data()), count * sizeof(std::complex<double>));
+    } else if (has("f4")) {
+        array.f64.resize(count);
+        std::vector<float> tmp(count);
+        in.read(reinterpret_cast<char*>(tmp.data()), count * sizeof(float));
+        for (long i = 0; i < count; ++i) {
+            array.f64[i] = tmp[i];
+        }
+    } else {
+        array.f64.resize(count);
+        in.read(reinterpret_cast<char*>(array.f64.data()), count * sizeof(double));
+    }
+    return array;
+}
+
+// VNA_AlGo reference geometry (gpr_core/config.py) keyed by switch positions.
+auto build_run_config() -> radar::config::RunConfig {
+    radar::config::RunConfig config{};
+    config.gpr.relative_permittivity = 1.0F;
+    config.gpr.tx_geometry = {
+        {.output_pos = 0U, .x_m = 0.785F, .y_m = 0.0F, .z_m = 0.0F},
+        {.output_pos = 1U, .x_m = -0.785F, .y_m = 0.0F, .z_m = 0.0F},
+    };
+    config.gpr.rx_geometry = {
+        {.input_pos = 0U, .x_m = -0.42F, .y_m = 0.0F, .z_m = 0.0F},
+        {.input_pos = 1U, .x_m = 0.48F, .y_m = 0.0F, .z_m = 0.0F},
+        {.input_pos = 2U, .x_m = 0.18F, .y_m = 0.0F, .z_m = 0.0F},
+        {.input_pos = 3U, .x_m = -0.12F, .y_m = 0.0F, .z_m = 0.0F},
+    };
+    return config;
+}
+
+// Live config matching the VNA_AlGo defaults (BPConfig/MotionConfig/SARConfig).
+auto build_live_config() -> radar::processing::ProcessingLiveConfig {
+    radar::processing::ProcessingLiveConfig live{};
+    live.processor_mode = "sar3d";
+    live.gpr_start_freq_mhz = 2000.0F;
+    live.gpr_stop_freq_mhz = 6000.0F;
+    live.gpr_min_depth_m = 3.2F;
+    live.gpr_max_depth_m = 10.0F;
+    live.gpr_speed_m_s = 0.854F;
+    live.gpr_look_angle_deg = 6.0F;
+    live.gpr_direction_sign = 1.0F;
+    live.gpr_apply_freq_phase_correction = true;
+    return live;
+}
+
+constexpr double kTxSweepTimeS = 0.079;
+constexpr int kTxCount = 2;
+constexpr int kRxCount = 4;
+
+auto load_collection(
+    const std::string& root,
+    const std::string& snapshot,
+    std::uint64_t collection_id
+) -> radar::ipc::PreprocessedCollection {
+    radar::ipc::PreprocessedCollection collection{};
+    collection.collection_id = collection_id;
+
+    // Frame time from the `_ns<...>` folder suffix, identical to standalone_bp.
+    const auto ns_pos = snapshot.rfind("_ns");
+    if (ns_pos == std::string::npos) {
+        throw std::runtime_error("snapshot has no _ns suffix: " + snapshot);
+    }
+    const auto mid_ns = static_cast<std::uint64_t>(std::stoull(snapshot.substr(ns_pos + 3)));
+    // The processor derives tx_sweep_time from capture_span / n_tx and the frame
+    // time from the capture midpoint, so set the window symmetric around mid_ns.
+    const auto span_ns = static_cast<std::uint64_t>(kTxCount * kTxSweepTimeS * 1e9);
+    collection.capture_start_ns = mid_ns - span_ns / 2U;
+    collection.capture_end_ns = mid_ns + span_ns / 2U;
+    collection.monotonic_ns = mid_ns;
+
+    for (int rx = 0; rx < kRxCount; ++rx) {
+        for (int tx = 0; tx < kTxCount; ++tx) {
+            char path[512];
+            std::snprintf(path, sizeof(path), "%s/%s/i%d_o%d_s21.npy", root.c_str(), snapshot.c_str(), rx, tx);
+            const auto s21 = load_npy(path);
+            std::snprintf(path, sizeof(path), "%s/%s/i%d_o%d_freq.npy", root.c_str(), snapshot.c_str(), rx, tx);
+            const auto freq = load_npy(path);
+
+            radar::ipc::SweepTraceBlock trace{};
+            trace.combo = radar::ipc::ComboKey{
+                .input_pos = static_cast<std::uint32_t>(rx),
+                .output_pos = static_cast<std::uint32_t>(tx),
+            };
+            trace.frequency_hz.assign(freq.f64.begin(), freq.f64.end());
+            trace.s21.reserve(s21.c128.size());
+            for (const auto& sample : s21.c128) {
+                trace.s21.push_back(radar::ipc::Complex32{
+                    .re = static_cast<float>(sample.real()),
+                    .im = static_cast<float>(sample.imag()),
+                });
+            }
+            trace.s11.resize(trace.s21.size());
+            trace.capture_start_ns = collection.capture_start_ns;
+            trace.capture_end_ns = collection.capture_end_ns;
+            collection.traces.push_back(std::move(trace));
+        }
+    }
+    return collection;
+}
+
+void print_projection_argmax(
+    const radar::ipc::ResultPayload& payload,
+    const char* col_axis_name,
+    const char* row_axis_name
+) {
+    std::size_t argmax = 0U;
+    float best = -1.0F;
+    for (std::size_t index = 0U; index < payload.image_values.size(); ++index) {
+        if (payload.image_values[index] > best) {
+            best = payload.image_values[index];
+            argmax = index;
+        }
+    }
+    const auto nx = payload.image_x_axis.size();
+    const auto col = argmax % nx;
+    const auto row = argmax / nx;
+    std::printf(
+        "%s: max=%.6f at %s=%.1f cm, %s=%.1f cm\n",
+        payload.processing_name.c_str(),
+        static_cast<double>(best),
+        col_axis_name,
+        static_cast<double>(payload.image_x_axis[col]) * 100.0,
+        row_axis_name,
+        static_cast<double>(payload.image_y_axis[row]) * 100.0
+    );
+}
+
+}  // namespace
+
+int main(int argc, char** argv) {
+    if (argc < 3) {
+        std::printf("usage: %s <preprocessed_dir> <center_snapshot>\n", argv[0]);
+        return 1;
+    }
+    const std::string root = argv[1];
+    const std::string center = argv[2];
+
+    std::vector<std::string> snapshots{};
+    if (DIR* dir = opendir(root.c_str())) {
+        for (dirent* entry = nullptr; (entry = readdir(dir)) != nullptr;) {
+            const std::string name = entry->d_name;
+            if (!name.empty() && (std::isdigit(name[0]) != 0)) {
+                snapshots.push_back(name);
+            }
+        }
+        closedir(dir);
+    }
+    std::sort(snapshots.begin(), snapshots.end());
+    const auto center_it = std::find(snapshots.begin(), snapshots.end(), center);
+    if (center_it == snapshots.end()) {
+        std::printf("center %s not found\n", center.c_str());
+        return 1;
+    }
+    const auto center_index = static_cast<std::size_t>(center_it - snapshots.begin());
+
+    const auto live = build_live_config();
+    const auto radius = static_cast<std::size_t>(live.sar3d_frame_radius);
+    if (center_index < radius || center_index + radius >= snapshots.size()) {
+        std::printf("not enough snapshots around center for radius %zu\n", radius);
+        return 1;
+    }
+
+    // The processor centers the aperture `radius` frames before the newest
+    // collection, so feeding [center-r .. center+r] in order reproduces the
+    // symmetric aperture of standalone_bp around `center`.
+    std::vector<radar::ipc::PreprocessedCollection> history{};
+    for (std::size_t index = center_index - radius; index <= center_index + radius; ++index) {
+        history.push_back(load_collection(root, snapshots[index], index));
+    }
+    std::printf("aperture: %zu frames around %s\n", history.size(), center.c_str());
+
+    const auto run_config = build_run_config();
+    radar::processing::Sar3dProcessor processor{};
+    radar::ipc::ResultCollection result{};
+    for (std::size_t index = 0U; index < history.size(); ++index) {
+        result = processor.process_collection(
+            run_config,
+            history[index],
+            std::span<const radar::ipc::PreprocessedCollection>(history.data(), index),
+            live
+        );
+    }
+
+    if (result.collection_payloads.empty()) {
+        std::printf("FAIL: sar3d produced no payloads\n");
+        return 1;
+    }
+    for (const auto& payload : result.collection_payloads) {
+        if (payload.kind == radar::ipc::ResultKind::ScalarF32) {
+            std::printf("%s: %.6e (compare with standalone_bp max |acc/cnt|)\n",
+                        payload.processing_name.c_str(), static_cast<double>(payload.scalar_value));
+        }
+    }
+    for (const auto& payload : result.collection_payloads) {
+        if (payload.kind == radar::ipc::ResultKind::ImageF32 && payload.processing_name == "sar3d_proj_xz") {
+            print_projection_argmax(payload, "X", "Z");
+        }
+    }
+
+    // Bottom view across the selectable angles: the target's u coordinate must
+    // move smoothly from its X (front, 0 deg) to its Z (side, 90 deg).
+    const auto last_previous =
+        std::span<const radar::ipc::PreprocessedCollection>(history.data(), history.size() - 1U);
+    for (const float angle_deg : {0.0F, 30.0F, 60.0F, 90.0F}) {
+        auto live_angle = live;
+        live_angle.sar3d_bottom_view_deg = angle_deg;
+        const auto angle_result =
+            processor.process_collection(run_config, history.back(), last_previous, live_angle);
+        for (const auto& payload : angle_result.collection_payloads) {
+            if (payload.kind == radar::ipc::ResultKind::ImageF32 && payload.processing_name == "sar3d_proj_bottom") {
+                std::printf("bottom @ %2.0f deg -> ", static_cast<double>(angle_deg));
+                print_projection_argmax(payload, "u", "Y");
+            }
+        }
+    }
+    std::printf("OK: sar3d processor path completed\n");
+    return 0;
+}

python_app/gui/app_window.py

@@ -371,6 +371,9 @@ class AppWindow(
         self._gpr_region_contours = []
         self._gpr_geometry_signature = None
         self._gpr_selected_geometry = None
+        self._sar3d_lookup_table = None
+        self._sar3d_xz_image_item = None
+        self._sar3d_xy_image_item = None
         self._phase_viewbox = None
         self._history_run_signature = None
         self._processor_run_signature = None

python_app/gui/controllers/app_window_config/live_processing_mixin.py

@@ -13,13 +13,13 @@ from PyQt6.QtWidgets import (
 )
 
 from python_app.gui.runtime.constraints import validate_processing_mode_constraints
+from python_app.gui.runtime.sar3d_views import sar3d_bottom_view_deg_from_label
 from python_app.orchestration.live_processing_config import ProcessingLiveConfig
 
 _GPR_MODES = ("gpr", "legacy_gpr")
-
-
-def _is_legacy_gpr(window) -> bool:
-    return window._processing_mode.currentText() == "legacy_gpr"
+# Modes that consume the shared gpr_* live-config fields (band, depth gate, antenna
+# selection, motion): both GPR imagers plus the 3D SAR processor.
+_GPR_LIKE_MODES = ("gpr", "legacy_gpr", "sar3d")
 
 
 def _attr(name: str):
@@ -27,9 +27,16 @@ def _attr(name: str):
     return lambda window: getattr(window, name)
 
 
-def _dual(gpr_attr: str, legacy_attr: str):
-    """Widget getter for a gpr_* field with separate gpr / legacy_gpr widgets."""
-    return lambda window: getattr(window, legacy_attr if _is_legacy_gpr(window) else gpr_attr)
+def _dual(default_attr: str, mode_attrs: dict[str, str] | None = None):
+    """Widget getter for a field with per-mode widgets (gpr / legacy_gpr / sar3d).
+
+    `mode_attrs` maps a processor mode to its dedicated widget attribute; any
+    other mode (including inactive ones) falls back to `default_attr`.
+    """
+    attrs = dict(mode_attrs or {})
+    return lambda window: getattr(
+        window, attrs.get(window._processing_mode.currentText(), default_attr)
+    )
 
 
 # The ONE authoritative field <-> widget map (inverse of `_live_processing_config`),
@@ -50,17 +57,17 @@ _WEB_LIVE_SCHEMA = [
     ("bscan_start_freq_mhz", "B-scan", ("bscan",), _attr("_bscan_start_freq_mhz")),
     ("bscan_stop_freq_mhz", "B-scan", ("bscan",), _attr("_bscan_stop_freq_mhz")),
     ("legacy_gpr_mode", "Mode", ("legacy_gpr",), _attr("_legacy_gpr_config_mode")),
-    ("gpr_input_positions", "Geometry & depth", _GPR_MODES, _dual("_gpr_input_positions_input", "_legacy_gpr_input_positions_input")),
-    ("gpr_output_positions", "Geometry & depth", _GPR_MODES, _dual("_gpr_output_positions_input", "_legacy_gpr_output_positions_input")),
-    ("gpr_min_depth_m", "Geometry & depth", _GPR_MODES, _dual("_gpr_min_depth_m", "_legacy_gpr_min_depth_m")),
-    ("gpr_max_depth_m", "Geometry & depth", _GPR_MODES, _dual("_gpr_max_depth_m", "_legacy_gpr_max_depth_m")),
-    ("gpr_start_freq_mhz", "Geometry & depth", _GPR_MODES, _dual("_gpr_start_freq_mhz", "_legacy_gpr_start_freq_mhz")),
-    ("gpr_stop_freq_mhz", "Geometry & depth", _GPR_MODES, _dual("_gpr_stop_freq_mhz", "_legacy_gpr_stop_freq_mhz")),
+    ("gpr_input_positions", "Geometry & depth", _GPR_LIKE_MODES, _dual("_gpr_input_positions_input", {"legacy_gpr": "_legacy_gpr_input_positions_input", "sar3d": "_sar3d_input_positions_input"})),
+    ("gpr_output_positions", "Geometry & depth", _GPR_LIKE_MODES, _dual("_gpr_output_positions_input", {"legacy_gpr": "_legacy_gpr_output_positions_input", "sar3d": "_sar3d_output_positions_input"})),
+    ("gpr_min_depth_m", "Geometry & depth", _GPR_LIKE_MODES, _dual("_gpr_min_depth_m", {"legacy_gpr": "_legacy_gpr_min_depth_m", "sar3d": "_sar3d_min_depth_m"})),
+    ("gpr_max_depth_m", "Geometry & depth", _GPR_LIKE_MODES, _dual("_gpr_max_depth_m", {"legacy_gpr": "_legacy_gpr_max_depth_m", "sar3d": "_sar3d_max_depth_m"})),
+    ("gpr_start_freq_mhz", "Geometry & depth", _GPR_LIKE_MODES, _dual("_gpr_start_freq_mhz", {"legacy_gpr": "_legacy_gpr_start_freq_mhz", "sar3d": "_sar3d_start_freq_mhz"})),
+    ("gpr_stop_freq_mhz", "Geometry & depth", _GPR_LIKE_MODES, _dual("_gpr_stop_freq_mhz", {"legacy_gpr": "_legacy_gpr_stop_freq_mhz", "sar3d": "_sar3d_stop_freq_mhz"})),
     ("gpr_imaging_plane_y_m", "Geometry & depth", ("gpr",), _attr("_gpr_imaging_plane_y_m")),
     ("gpr_range_comp_power", "Imaging", ("gpr",), _attr("_gpr_range_comp_power")),
     ("gpr_score_mode", "Imaging", ("gpr",), _attr("_gpr_score_mode")),
-    ("gpr_background_subtract_enabled", "Imaging", _GPR_MODES, _dual("_gpr_background_subtract_enabled", "_legacy_gpr_background_subtract_enabled")),
-    ("gpr_background_mean_count", "Imaging", _GPR_MODES, _dual("_gpr_background_mean_count", "_legacy_gpr_background_mean_count")),
+    ("gpr_background_subtract_enabled", "Imaging", _GPR_MODES, _dual("_gpr_background_subtract_enabled", {"legacy_gpr": "_legacy_gpr_background_subtract_enabled"})),
+    ("gpr_background_mean_count", "Imaging", _GPR_MODES, _dual("_gpr_background_mean_count", {"legacy_gpr": "_legacy_gpr_background_mean_count"})),
     ("gpr_remove_sidelobe_objects_enabled", "Imaging", ("gpr",), _attr("_gpr_remove_sidelobe_objects_enabled")),
     ("gpr_object_min_frac", "Detection", ("gpr",), _attr("_gpr_object_min_frac")),
     ("gpr_max_detected_objects_to_draw", "Detection", ("gpr",), _attr("_gpr_max_detected_objects_to_draw")),
@@ -70,17 +77,30 @@ _WEB_LIVE_SCHEMA = [
     ("gpr_snr_thresh", "Detection", ("legacy_gpr",), _attr("_legacy_gpr_snr_thresh")),
     ("gpr_snr_comp_max", "Detection", ("legacy_gpr",), _attr("_legacy_gpr_snr_comp_max")),
     ("legacy_gpr_min_visible_pair_count", "Detection", ("legacy_gpr",), _attr("_legacy_gpr_min_visible_pair_count")),
-    ("gpr_visible_x_min_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_x_min_m", "_legacy_gpr_visible_x_min_m")),
-    ("gpr_visible_x_max_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_x_max_m", "_legacy_gpr_visible_x_max_m")),
-    ("gpr_visible_z_min_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_z_min_m", "_legacy_gpr_visible_z_min_m")),
-    ("gpr_visible_z_max_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_z_max_m", "_legacy_gpr_visible_z_max_m")),
+    ("gpr_visible_x_min_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_x_min_m", {"legacy_gpr": "_legacy_gpr_visible_x_min_m"})),
+    ("gpr_visible_x_max_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_x_max_m", {"legacy_gpr": "_legacy_gpr_visible_x_max_m"})),
+    ("gpr_visible_z_min_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_z_min_m", {"legacy_gpr": "_legacy_gpr_visible_z_min_m"})),
+    ("gpr_visible_z_max_m", "Visible window", _GPR_MODES, _dual("_gpr_visible_z_max_m", {"legacy_gpr": "_legacy_gpr_visible_z_max_m"})),
     ("gpr_motion_mode", "Motion", ("gpr",), _attr("_gpr_motion_mode")),
-    ("gpr_look_angle_deg", "Motion", _GPR_MODES, _dual("_gpr_look_angle_deg", "_legacy_gpr_look_angle_deg")),
-    ("gpr_direction_sign", "Motion", ("gpr",), _attr("_gpr_direction_sign")),
-    ("gpr_apply_freq_phase_correction", "Motion", ("legacy_gpr",), _attr("_legacy_gpr_apply_freq_phase_correction")),
+    ("gpr_look_angle_deg", "Motion", _GPR_LIKE_MODES, _dual("_gpr_look_angle_deg", {"legacy_gpr": "_legacy_gpr_look_angle_deg", "sar3d": "_sar3d_look_angle_deg"})),
+    ("gpr_direction_sign", "Motion", ("gpr", "sar3d"), _dual("_gpr_direction_sign", {"sar3d": "_sar3d_direction_sign"})),
+    ("gpr_apply_freq_phase_correction", "Motion", ("legacy_gpr", "sar3d"), _dual("_legacy_gpr_apply_freq_phase_correction", {"sar3d": "_sar3d_apply_freq_phase_correction"})),
     ("gpr_reference_mode", "Motion", ("legacy_gpr",), _attr("_legacy_gpr_reference_mode")),
-    ("ignore_socket_speed", "Motion", _GPR_MODES, _dual("_gpr_ignore_socket_speed_enabled", "_legacy_gpr_ignore_socket_speed_enabled")),
-    ("gpr_speed_m_s", "Motion", _GPR_MODES, _dual("_gpr_speed_m_s", "_legacy_gpr_speed_m_s")),
+    ("ignore_socket_speed", "Motion", _GPR_LIKE_MODES, _dual("_gpr_ignore_socket_speed_enabled", {"legacy_gpr": "_legacy_gpr_ignore_socket_speed_enabled", "sar3d": "_sar3d_ignore_socket_speed_enabled"})),
+    ("gpr_speed_m_s", "Motion", _GPR_LIKE_MODES, _dual("_gpr_speed_m_s", {"legacy_gpr": "_legacy_gpr_speed_m_s", "sar3d": "_sar3d_speed_m_s"})),
+    ("sar3d_frame_radius", "3D SAR", ("sar3d",), _attr("_sar3d_frame_radius")),
+    ("sar3d_edge_frame_weight", "3D SAR", ("sar3d",), _attr("_sar3d_edge_frame_weight")),
+    ("sar3d_pitch_deg", "3D SAR", ("sar3d",), _attr("_sar3d_pitch_deg")),
+    ("sar3d_center_depth_m", "3D SAR", ("sar3d",), _attr("_sar3d_center_depth_m")),
+    ("sar3d_half_width_x_m", "3D SAR", ("sar3d",), _attr("_sar3d_half_width_x_m")),
+    ("sar3d_half_width_y_m", "3D SAR", ("sar3d",), _attr("_sar3d_half_width_y_m")),
+    ("sar3d_half_width_z_m", "3D SAR", ("sar3d",), _attr("_sar3d_half_width_z_m")),
+    ("sar3d_grid_nx", "3D SAR", ("sar3d",), _attr("_sar3d_grid_nx")),
+    ("sar3d_grid_ny", "3D SAR", ("sar3d",), _attr("_sar3d_grid_ny")),
+    ("sar3d_grid_nz", "3D SAR", ("sar3d",), _attr("_sar3d_grid_nz")),
+    ("sar3d_use_compensation", "3D SAR", ("sar3d",), _attr("_sar3d_use_compensation")),
+    ("sar3d_normalize_pairs", "3D SAR", ("sar3d",), _attr("_sar3d_normalize_pairs")),
+    ("sar3d_bottom_view_deg", "3D SAR", ("sar3d",), _attr("_sar3d_bottom_view")),
 ]
 _WEB_LIVE_GETTERS = {field: getter for field, _group, _modes, getter in _WEB_LIVE_SCHEMA}
 
@@ -101,7 +121,7 @@ _NON_WIDGET_LIVE_FIELDS = frozenset({
 # config). They apply immediately through the same redraw as a live edit but write
 # nothing to processing_live.json.
 _WEB_DISPLAY_SCHEMA = [
-    ("render_mode", "Display", _GPR_MODES, _dual("_gpr_render_mode", "_legacy_gpr_render_mode")),
+    ("render_mode", "Display", _GPR_MODES, _dual("_gpr_render_mode", {"legacy_gpr": "_legacy_gpr_render_mode"})),
     ("subtract_mean_ascan", "Display", ("bscan",), _attr("_bscan_subtract_mean_ascan")),
     ("combo_filter", "Display", ("pass_through",), _attr("_pass_through_combo_filter_input")),
 ]
@@ -113,8 +133,8 @@ _WEB_STABLE_SCHEMA = [
     # Coherent BP fixes the medium to eps_r = 1 (Horns_motion_3libre.py), so relative
     # permittivity is a legacy-GPR-only knob; BP ignores it.
     ("relative_permittivity", "Geometry & medium", ("legacy_gpr",), _attr("_gpr_relative_permittivity")),
-    ("tx_geometry", "Geometry & medium", _GPR_MODES, _attr("_gpr_tx_geometry_input")),
-    ("rx_geometry", "Geometry & medium", _GPR_MODES, _attr("_gpr_rx_geometry_input")),
+    ("tx_geometry", "Geometry & medium", _GPR_LIKE_MODES, _attr("_gpr_tx_geometry_input")),
+    ("rx_geometry", "Geometry & medium", _GPR_LIKE_MODES, _attr("_gpr_rx_geometry_input")),
 ]
 
 # Every field the web form renders and may write back, across all three categories.
@@ -222,6 +242,17 @@ class AppWindowLiveProcessingMixin:
         """Return whether a mode uses GPR result payloads and plot surface."""
         return mode in {"gpr", "legacy_gpr"}
 
+    @staticmethod
+    def _is_gpr_like_processing_mode(mode: str) -> bool:
+        """Return whether a mode consumes the shared gpr_* fields and Tx/Rx geometry.
+
+        Covers both GPR imagers and the 3D SAR processor. The latter renders on its
+        own plot surface, so GPR-object/locator branches must keep using
+        :meth:`_is_gpr_processing_mode`; this predicate gates only the shared
+        behaviors (result drain on settings changes, geometry page visibility).
+        """
+        return mode in set(_GPR_LIKE_MODES)
+
     def _validate_processing_mode_selection(self, mode: str) -> None:
         """Validate requested processing mode against current stable/live GUI state."""
         validate_processing_mode_constraints(
@@ -256,6 +287,10 @@ class AppWindowLiveProcessingMixin:
         # Switch-position CSV text -> integer lists.
         values["gpr_input_positions"] = self._parse_csv_int_list(values["gpr_input_positions"])
         values["gpr_output_positions"] = self._parse_csv_int_list(values["gpr_output_positions"])
+        # Bottom-view combo label -> projection angle in degrees.
+        values["sar3d_bottom_view_deg"] = sar3d_bottom_view_deg_from_label(
+            str(values["sar3d_bottom_view_deg"])
+        )
         # Order-normalize the pass-through Y range so a reversed entry still works.
         y_min_db, y_max_db = values["pass_through_y_min_db"], values["pass_through_y_max_db"]
         values["pass_through_y_min_db"] = min(y_min_db, y_max_db)
@@ -352,7 +387,7 @@ class AppWindowLiveProcessingMixin:
             return
         try:
             current_mode = self._processing_mode.currentText()
-            if self._is_gpr_processing_mode(current_mode):
+            if self._is_gpr_like_processing_mode(current_mode):
                 self._drain_results_until_quiet(timeout_s=0.05, poll_s=0.005)
 
             self._write_live_processing_config()
@@ -364,16 +399,16 @@ class AppWindowLiveProcessingMixin:
                 # recompute of everything on screen. Debounced, so dragging a spin box
                 # sends one burst instead of one per step.
                 self._schedule_bscan_history_reprocess()
-            elif self._is_gpr_processing_mode(current_mode):
+            elif self._is_gpr_like_processing_mode(current_mode):
                 latest = self._drain_results_until_quiet(timeout_s=0.8, poll_s=0.02)
                 collection = latest
                 if collection is None and self._result_history:
                     collection = self._result_history[-1]
-                if collection is not None:
-                    if not self._draw_results(collection):
+                if collection is None or not self._draw_results(collection):
+                    if current_mode == "sar3d":
+                        self._clear_sar3d_plot()
+                    else:
                         self._clear_gpr_plot()
-                else:
-                    self._clear_gpr_plot()
             elif self._result_history:
                 self._draw_results(self._result_history[-1])
             else:
@@ -427,9 +462,10 @@ class AppWindowLiveProcessingMixin:
             "bscan": 1,
             "gpr": 2,
             "legacy_gpr": 3,
+            "sar3d": 4,
         }
         self._processing_mode_pages.setCurrentIndex(mode_to_page.get(mode, 0))
-        self._gpr_common_page.setVisible(self._is_gpr_processing_mode(mode))
+        self._gpr_common_page.setVisible(self._is_gpr_like_processing_mode(mode))
         current_page = self._processing_mode_pages.currentWidget()
         if current_page is not None:
             self._processing_mode_pages.setFixedHeight(current_page.sizeHint().height())
@@ -509,9 +545,21 @@ class AppWindowLiveProcessingMixin:
                 f"render_mode={self._legacy_gpr_render_mode.currentText()}, "
                 f"min_pairs={self._legacy_gpr_min_visible_pair_count.value()})"
             )
+        elif mode == "sar3d":
+            self._log(
+                "Processing mode selected: sar3d "
+                f"(inputs={self._sar3d_input_positions_input.text().strip() or '<all>'}, "
+                f"outputs={self._sar3d_output_positions_input.text().strip() or '<all>'}, "
+                f"depth={self._sar3d_min_depth_m.value():g}..{self._sar3d_max_depth_m.value():g} m, "
+                f"freq={self._sar3d_start_freq_mhz.value():g}..{self._sar3d_stop_freq_mhz.value():g} MHz, "
+                f"frame_radius={self._sar3d_frame_radius.value()}, "
+                f"grid={self._sar3d_grid_nx.value()}x{self._sar3d_grid_ny.value()}x{self._sar3d_grid_nz.value()}, "
+                f"center_depth={self._sar3d_center_depth_m.value():g} m, "
+                f"pitch={self._sar3d_pitch_deg.value():g} deg)"
+            )
 
     def _clear_history_mode_caches(self) -> None:
-        """Drop cached render state for pass-through, B-scan, and GPR views."""
+        """Drop cached render state for pass-through, B-scan, GPR, and 3D SAR views."""
         self._clear_bscan_plot_history()
         if hasattr(self, "_bscan_plot"):
             self._bscan_plot.clear()
@@ -519,6 +567,8 @@ class AppWindowLiveProcessingMixin:
         self._clear_trace_plots()
         if hasattr(self, "_gpr_plot"):
             self._clear_gpr_plot()
+        if hasattr(self, "_sar3d_xz_plot"):
+            self._clear_sar3d_plot()
 
     def _redraw_after_history_deletion(self) -> None:
         """Refresh plot immediately after destructive history deletion."""
@@ -534,6 +584,11 @@ class AppWindowLiveProcessingMixin:
                 return
             self._clear_gpr_plot()
             return
+        if self._processing_mode.currentText() == "sar3d":
+            if self._result_history and self._draw_results(self._result_history[-1]):
+                return
+            self._clear_sar3d_plot()
+            return
         if self._result_history:
             self._draw_results(self._result_history[-1])
             return

python_app/gui/controllers/app_window_config/profile_io_mixin.py

@@ -10,6 +10,7 @@ from pathlib import Path
 from PyQt6.QtCore import QSignalBlocker
 from PyQt6.QtWidgets import QFileDialog
 
+from python_app.gui.runtime.sar3d_views import sar3d_bottom_view_label_from_deg
 from python_app.models.gui_profile_model import GuiProfileModel
 from python_app.models.run_config_model import RunConfigModel
 from python_app.orchestration.preprocess_assets import (
@@ -336,6 +337,29 @@ class AppWindowConfigProfileIOMixin:
             self._legacy_gpr_visible_x_max_m,
             self._legacy_gpr_visible_z_min_m,
             self._legacy_gpr_visible_z_max_m,
+            self._sar3d_input_positions_input,
+            self._sar3d_output_positions_input,
+            self._sar3d_min_depth_m,
+            self._sar3d_max_depth_m,
+            self._sar3d_start_freq_mhz,
+            self._sar3d_stop_freq_mhz,
+            self._sar3d_speed_m_s,
+            self._sar3d_ignore_socket_speed_enabled,
+            self._sar3d_look_angle_deg,
+            self._sar3d_direction_sign,
+            self._sar3d_apply_freq_phase_correction,
+            self._sar3d_frame_radius,
+            self._sar3d_edge_frame_weight,
+            self._sar3d_pitch_deg,
+            self._sar3d_center_depth_m,
+            self._sar3d_half_width_x_m,
+            self._sar3d_half_width_y_m,
+            self._sar3d_half_width_z_m,
+            self._sar3d_grid_nx,
+            self._sar3d_grid_ny,
+            self._sar3d_grid_nz,
+            self._sar3d_use_compensation,
+            self._sar3d_normalize_pairs,
             self._save_count,
             self._record_count,
             self._save_path_input,
@@ -524,6 +548,35 @@ class AppWindowConfigProfileIOMixin:
             self._legacy_gpr_visible_z_min_m.setValue(float(gui_state.processing.legacy_gpr.visible_z_min_m))
             self._legacy_gpr_visible_z_max_m.setValue(float(gui_state.processing.legacy_gpr.visible_z_max_m))
 
+            sar3d_state = gui_state.processing.sar3d
+            self._sar3d_input_positions_input.setText(str(sar3d_state.input_positions))
+            self._sar3d_output_positions_input.setText(str(sar3d_state.output_positions))
+            self._sar3d_min_depth_m.setValue(float(sar3d_state.min_depth_m))
+            self._sar3d_max_depth_m.setValue(float(sar3d_state.max_depth_m))
+            self._sar3d_start_freq_mhz.setValue(float(sar3d_state.start_freq_mhz))
+            self._sar3d_stop_freq_mhz.setValue(float(sar3d_state.stop_freq_mhz))
+            self._sar3d_speed_m_s.setValue(float(sar3d_state.speed_m_s))
+            self._sar3d_ignore_socket_speed_enabled.setChecked(bool(sar3d_state.ignore_socket_speed_enabled))
+            self._sar3d_look_angle_deg.setValue(float(sar3d_state.look_angle_deg))
+            self._sar3d_direction_sign.setValue(float(sar3d_state.direction_sign))
+            self._sar3d_apply_freq_phase_correction.setChecked(bool(sar3d_state.apply_freq_phase_correction))
+            self._sar3d_frame_radius.setValue(int(sar3d_state.frame_radius))
+            self._sar3d_edge_frame_weight.setValue(float(sar3d_state.edge_frame_weight))
+            self._sar3d_pitch_deg.setValue(float(sar3d_state.pitch_deg))
+            self._sar3d_center_depth_m.setValue(float(sar3d_state.center_depth_m))
+            self._sar3d_half_width_x_m.setValue(float(sar3d_state.half_width_x_m))
+            self._sar3d_half_width_y_m.setValue(float(sar3d_state.half_width_y_m))
+            self._sar3d_half_width_z_m.setValue(float(sar3d_state.half_width_z_m))
+            self._sar3d_grid_nx.setValue(int(sar3d_state.grid_nx))
+            self._sar3d_grid_ny.setValue(int(sar3d_state.grid_ny))
+            self._sar3d_grid_nz.setValue(int(sar3d_state.grid_nz))
+            self._sar3d_use_compensation.setChecked(bool(sar3d_state.use_compensation))
+            self._sar3d_normalize_pairs.setChecked(bool(sar3d_state.normalize_pairs))
+            self._set_combo_current_text(
+                self._sar3d_bottom_view,
+                sar3d_bottom_view_label_from_deg(float(sar3d_state.bottom_view_deg)),
+            )
+
             self._save_count.setValue(int(gui_state.data_actions.save_count))
             self._record_count.setValue(int(gui_state.data_actions.record_count))
             self._save_path_input.setText(str(gui_state.data_actions.save_path))

python_app/gui/controllers/app_window_config/radar_limits_mixin.py

@@ -182,6 +182,8 @@ class AppWindowRadarLimitsMixin:
             "_gpr_stop_freq_mhz": "GPR Stop MHz",
             "_legacy_gpr_start_freq_mhz": "Legacy GPR Start MHz",
             "_legacy_gpr_stop_freq_mhz": "Legacy GPR Stop MHz",
+            "_sar3d_start_freq_mhz": "3D SAR Start MHz",
+            "_sar3d_stop_freq_mhz": "3D SAR Stop MHz",
         }
         widgets = [getattr(self, widget_name) for widget_name in widget_names]
         previous_values = {widget_name: getattr(self, widget_name).value() for widget_name in widget_names}
@@ -232,5 +234,7 @@ class AppWindowRadarLimitsMixin:
                 "_gpr_stop_freq_mhz",
                 "_legacy_gpr_start_freq_mhz",
                 "_legacy_gpr_stop_freq_mhz",
+                "_sar3d_start_freq_mhz",
+                "_sar3d_stop_freq_mhz",
             )
         )

python_app/gui/controllers/app_window_config/state_builders.py

@@ -13,6 +13,7 @@ from python_app.models.gui_profile_model import (
     GuiPreprocessDialogStateModel,
     GuiProcessingStateModel,
     GuiProfileModel,
+    GuiSar3dStateModel,
     GuiStateModel,
     GuiSwitchStateModel,
 )
@@ -24,6 +25,7 @@ from python_app.orchestration.preprocess_assets import (
     VISIBLE_PREPROCESS_ASSET_KEYS,
     preprocess_asset_model,
 )
+from python_app.gui.runtime.sar3d_views import sar3d_bottom_view_deg_from_label
 from python_app.storage.npz_store import radar_key_from_config
 
 
@@ -285,6 +287,10 @@ class AppWindowConfigStateBuildersMixin:
                     visible_z_min_m=0.0,
                     visible_z_max_m=14.0,
                 ),
+                sar3d=GuiSar3dStateModel(
+                    input_positions=self._default_gpr_input_positions_from_config(config),
+                    output_positions=self._default_gpr_output_positions_from_config(config),
+                ),
             ),
             data_actions=GuiDataActionsStateModel(
                 save_count=10,
@@ -418,6 +424,34 @@ class AppWindowConfigStateBuildersMixin:
                     visible_z_min_m=float(self._legacy_gpr_visible_z_min_m.value()),
                     visible_z_max_m=float(self._legacy_gpr_visible_z_max_m.value()),
                 ),
+                sar3d=GuiSar3dStateModel(
+                    input_positions=self._sar3d_input_positions_input.text().strip(),
+                    output_positions=self._sar3d_output_positions_input.text().strip(),
+                    min_depth_m=float(self._sar3d_min_depth_m.value()),
+                    max_depth_m=float(self._sar3d_max_depth_m.value()),
+                    start_freq_mhz=float(self._sar3d_start_freq_mhz.value()),
+                    stop_freq_mhz=float(self._sar3d_stop_freq_mhz.value()),
+                    speed_m_s=float(self._sar3d_speed_m_s.value()),
+                    ignore_socket_speed_enabled=bool(self._sar3d_ignore_socket_speed_enabled.isChecked()),
+                    look_angle_deg=float(self._sar3d_look_angle_deg.value()),
+                    direction_sign=float(self._sar3d_direction_sign.value()),
+                    apply_freq_phase_correction=bool(self._sar3d_apply_freq_phase_correction.isChecked()),
+                    frame_radius=int(self._sar3d_frame_radius.value()),
+                    edge_frame_weight=float(self._sar3d_edge_frame_weight.value()),
+                    pitch_deg=float(self._sar3d_pitch_deg.value()),
+                    center_depth_m=float(self._sar3d_center_depth_m.value()),
+                    half_width_x_m=float(self._sar3d_half_width_x_m.value()),
+                    half_width_y_m=float(self._sar3d_half_width_y_m.value()),
+                    half_width_z_m=float(self._sar3d_half_width_z_m.value()),
+                    grid_nx=int(self._sar3d_grid_nx.value()),
+                    grid_ny=int(self._sar3d_grid_ny.value()),
+                    grid_nz=int(self._sar3d_grid_nz.value()),
+                    use_compensation=bool(self._sar3d_use_compensation.isChecked()),
+                    normalize_pairs=bool(self._sar3d_normalize_pairs.isChecked()),
+                    bottom_view_deg=sar3d_bottom_view_deg_from_label(
+                        self._sar3d_bottom_view.currentText()
+                    ),
+                ),
             ),
             data_actions=GuiDataActionsStateModel(
                 save_count=int(self._save_count.value()),
@@ -472,7 +506,7 @@ class AppWindowConfigStateBuildersMixin:
                 combo_text = self._combos_text.text()
                 config.combos = parse_combos_from_text(combo_text)
                 config.ensure_combos()
-            if self._processing_mode.currentText() not in {"gpr", "legacy_gpr"} and self._switches_are_effectively_static(config):
+            if self._processing_mode.currentText() not in {"gpr", "legacy_gpr", "sar3d"} and self._switches_are_effectively_static(config):
                 config.combos = [ComboModel(input=0, output=0)]
 
         for key in PREPROCESS_ASSET_KEYS:

python_app/gui/controllers/app_window_plot/__init__.py

@@ -3,11 +3,13 @@
 from python_app.gui.controllers.app_window_plot.bscan_plot_mixin import AppWindowBscanPlotMixin
 from python_app.gui.controllers.app_window_plot.bscan_replay_mixin import AppWindowBscanReplayMixin
 from python_app.gui.controllers.app_window_plot.gpr_plot_mixin import AppWindowGprPlotMixin
+from python_app.gui.controllers.app_window_plot.sar3d_plot_mixin import AppWindowSar3dPlotMixin
 from python_app.gui.controllers.app_window_plot.trace_plot_mixin import AppWindowTracePlotMixin
 
 __all__ = [
     "AppWindowBscanPlotMixin",
     "AppWindowBscanReplayMixin",
     "AppWindowGprPlotMixin",
+    "AppWindowSar3dPlotMixin",
     "AppWindowTracePlotMixin",
 ]

python_app/gui/controllers/app_window_plot/sar3d_plot_mixin.py

@@ -0,0 +1,162 @@
+"""3D SAR projection rendering helpers.
+
+The sar3d processor publishes two normalized max-projections of the imaged
+volume as collection-level ImageF32 payloads, matching the legacy render
+conventions (X across, Z range/"вдаль", Y up):
+
+- ``sar3d_proj_xz`` — top view: X (m) horizontal, Z/range (m) vertical
+  (max over the height Y), like the classic GPR heatmap;
+- ``sar3d_proj_bottom`` — switchable bottom view: height Y (m) vertical, the
+  horizontal axis is the ground direction rotated by ``sar3d_bottom_view_deg``
+  around Y (0° = front view X, 90° = side view Z, 30°/60° in between);
+
+plus ``sar3d_peak`` (ScalarF32), the volume peak before normalization. Both
+projections are drawn on their own PlotWidgets stacked in one plot-stack page
+(top view on top), mirroring how the pass-through page stacks its two plots.
+"""
+
+from __future__ import annotations
+
+from PyQt6.QtCore import QRectF
+import numpy as np
+import pyqtgraph as pg
+
+from python_app.gui.runtime.sar3d_views import sar3d_bottom_view_deg_from_label
+from python_app.models.dataset_model import ResultCollection
+
+# Fraction of the projection maximum below which the heatmap is not highlighted
+# (same as cloud-threshold=0.75 in the VNA_AlGo 3D SAR animation): the color
+# scale is stretched over [frac, 1], everything below falls into the darkest
+# LUT color.
+HEATMAP_DISPLAY_THRESHOLD_FRAC = 0.75
+
+
+class AppWindowSar3dPlotMixin:
+    """Renders 3D SAR volume projections published by the sar3d processor."""
+
+    def _sar3d_bottom_view_labels(self) -> tuple[str, str]:
+        """Return (plot title, horizontal axis label) for the active bottom view.
+
+        The plot surfaces are created before the processing section builds the
+        bottom-view combo, so fall back to the front view until it exists.
+        """
+        combo = getattr(self, "_sar3d_bottom_view", None)
+        degrees = sar3d_bottom_view_deg_from_label(combo.currentText()) if combo is not None else 0.0
+        if degrees <= 1e-9:
+            return "3D SAR front view (X-Y)", "X"
+        if degrees >= 90.0 - 1e-9:
+            return "3D SAR side view (Z-Y)", "Range Z"
+        return f"3D SAR view @ {degrees:g}\N{DEGREE SIGN}", f"u({degrees:g}\N{DEGREE SIGN})"
+
+    def _configure_sar3d_plot_axes(self) -> None:
+        """Apply persistent axis labels and view settings to both projection plots."""
+        _bottom_title, bottom_x_label = self._sar3d_bottom_view_labels()
+        for plot, x_label, y_label in (
+            (self._sar3d_xz_plot, "X", "Range Z"),
+            (self._sar3d_xy_plot, bottom_x_label, "Height Y"),
+        ):
+            plot_item = plot.getPlotItem()
+            plot_item.showAxis("left", show=True)
+            plot_item.showAxis("bottom", show=True)
+            plot_item.setClipToView(True)
+            plot.setLabel("bottom", x_label, units="m")
+            plot.setLabel("left", y_label, units="m")
+            view_box = plot.getViewBox()
+            view_box.invertY(False)
+            view_box.enableAutoRange(x=False, y=False)
+            self._disable_gpr_plot_interaction(plot, view_box)
+
+    def _ensure_sar3d_plot_items(self) -> None:
+        """Create persistent projection image items once and reuse them on redraw."""
+        if self._sar3d_xz_image_item is not None:
+            return
+
+        self._configure_sar3d_plot_axes()
+        if self._sar3d_lookup_table is None:
+            # Same palette as the GPR accumulator so both imagers read alike.
+            self._sar3d_lookup_table = self._build_lut(["#081c15", "#1b4332", "#ffd166", "#f94144"])
+
+        self._sar3d_xz_image_item = pg.ImageItem(axisOrder="row-major")
+        self._sar3d_xz_image_item.setZValue(0)
+        self._sar3d_xz_image_item.hide()
+        self._sar3d_xz_plot.addItem(self._sar3d_xz_image_item)
+
+        self._sar3d_xy_image_item = pg.ImageItem(axisOrder="row-major")
+        self._sar3d_xy_image_item.setZValue(0)
+        self._sar3d_xy_image_item.hide()
+        self._sar3d_xy_plot.addItem(self._sar3d_xy_image_item)
+
+    def _clear_sar3d_plot(self) -> None:
+        """Clear both 3D SAR projection surfaces."""
+        if not hasattr(self, "_sar3d_xz_plot"):
+            return
+        self._configure_sar3d_plot_axes()
+        if self._sar3d_xz_image_item is not None:
+            self._sar3d_xz_image_item.hide()
+        if self._sar3d_xy_image_item is not None:
+            self._sar3d_xy_image_item.hide()
+        bottom_title, _bottom_x_label = self._sar3d_bottom_view_labels()
+        self._sar3d_xz_plot.setTitle("3D SAR top view (X-Z)")
+        self._sar3d_xy_plot.setTitle(bottom_title)
+
+    def _draw_sar3d_projection(self, plot: pg.PlotWidget, image_item: pg.ImageItem, payload) -> bool:
+        """Draw one normalized projection payload onto its plot; return success."""
+        image = np.asarray(payload.image, dtype=np.float32)
+        x_axis = np.asarray(payload.image_x_axis, dtype=np.float32)
+        y_axis = np.asarray(payload.image_y_axis, dtype=np.float32)
+        if image.ndim != 2 or image.size == 0 or x_axis.size == 0 or y_axis.size == 0:
+            image_item.hide()
+            return False
+
+        x_min = float(x_axis[0])
+        x_max = float(x_axis[-1])
+        y_min = float(y_axis[0])
+        y_max = float(y_axis[-1])
+        rect = QRectF(x_min, y_min, max(x_max - x_min, 1e-6), max(y_max - y_min, 1e-6))
+
+        image_item.setImage(image, autoLevels=False)
+        image_item.setRect(rect)
+        image_item.setLookupTable(self._sar3d_lookup_table)
+        # Projections are normalized to [0, 1] by the processor; stretch the
+        # colors over [threshold, 1] like the 3D animation cloud cutoff.
+        image_item.setLevels((HEATMAP_DISPLAY_THRESHOLD_FRAC, 1.0))
+        image_item.show()
+        plot.setXRange(x_min, x_max, padding=0.0)
+        plot.setYRange(y_min, y_max, padding=0.0)
+        return True
+
+    def _draw_sar3d_projections(self, collection: ResultCollection) -> bool:
+        """Draw the latest collection-level 3D SAR top and bottom projections."""
+        xz_payload = self._collection_payload_by_name(collection, "sar3d_proj_xz", kind=3)
+        bottom_payload = self._collection_payload_by_name(collection, "sar3d_proj_bottom", kind=3)
+        if xz_payload is None and bottom_payload is None:
+            self._clear_sar3d_plot()
+            return False
+
+        for plot in (self._sar3d_xz_plot, self._sar3d_xy_plot):
+            plot.setUpdatesEnabled(False)
+        try:
+            self._ensure_sar3d_plot_items()
+            # Axis labels track the bottom-view selector (payload axes change with it).
+            self._configure_sar3d_plot_axes()
+            xz_drawn = xz_payload is not None and self._draw_sar3d_projection(
+                self._sar3d_xz_plot, self._sar3d_xz_image_item, xz_payload
+            )
+            bottom_drawn = bottom_payload is not None and self._draw_sar3d_projection(
+                self._sar3d_xy_plot, self._sar3d_xy_image_item, bottom_payload
+            )
+            if not xz_drawn and not bottom_drawn:
+                self._clear_sar3d_plot()
+                return False
+
+            peak_payload = self._collection_payload_by_name(collection, "sar3d_peak", kind=2)
+            peak_suffix = ""
+            if peak_payload is not None:
+                peak_suffix = f" | peak={float(peak_payload.scalar_value):.3g}"
+            bottom_title, _bottom_x_label = self._sar3d_bottom_view_labels()
+            self._sar3d_xz_plot.setTitle(f"3D SAR top view (X-Z){peak_suffix}")
+            self._sar3d_xy_plot.setTitle(bottom_title)
+        finally:
+            for plot in (self._sar3d_xz_plot, self._sar3d_xy_plot):
+                plot.setUpdatesEnabled(True)
+        return True

python_app/gui/controllers/app_window_plot/trace_plot_mixin.py

@@ -78,7 +78,7 @@ class AppWindowTracePlotMixin:
 
     def _on_trace_visibility_changed(self, *_args) -> None:
         """Redraw pass-through traces when magnitude/phase toggles changed."""
-        if self._processing_mode.currentText() in {"bscan", "gpr", "legacy_gpr"}:
+        if self._processing_mode.currentText() in {"bscan", "gpr", "legacy_gpr", "sar3d"}:
             return
         if self._result_history:
             self._draw_results(self._result_history[-1])

python_app/gui/controllers/app_window_plot_mixin.py

@@ -6,6 +6,7 @@ from python_app.gui.controllers.app_window_plot import (
     AppWindowBscanPlotMixin,
     AppWindowBscanReplayMixin,
     AppWindowGprPlotMixin,
+    AppWindowSar3dPlotMixin,
     AppWindowTracePlotMixin,
 )
 from python_app.models.dataset_model import ResultCollection
@@ -16,8 +17,9 @@ class AppWindowPlotMixin(
     AppWindowBscanPlotMixin,
     AppWindowBscanReplayMixin,
     AppWindowGprPlotMixin,
+    AppWindowSar3dPlotMixin,
 ):
-    """Routes plotting to trace, B-scan, or GPR-specific mixins."""
+    """Routes plotting to trace, B-scan, GPR, or 3D SAR-specific mixins."""
 
     def _draw_preferred_collection(
         self,
@@ -35,6 +37,8 @@ class AppWindowPlotMixin(
             return self._draw_bscan_heatmap(collection)
         if self._processing_mode.currentText() in {"gpr", "legacy_gpr"}:
             return self._draw_gpr_map(collection)
+        if self._processing_mode.currentText() == "sar3d":
+            return self._draw_sar3d_projections(collection)
         return self._draw_trace_lines(collection)
 
     def _result_collection_has_trace(self, collection: ResultCollection) -> bool:

python_app/gui/controllers/app_window_ui_mixin.py

@@ -74,6 +74,7 @@ class AppWindowUiMixin:
         self._plot_stack = QStackedWidget(root)
         self._build_bscan_plot_page()
         self._build_gpr_plot_page()
+        self._build_sar3d_plot_page()
         self._build_trace_plot_page()
 
         # Default view on startup is pass-through traces.
@@ -139,6 +140,27 @@ class AppWindowUiMixin:
         self._configure_gpr_plot_axes()
         self._plot_stack.addWidget(self._gpr_plot)
 
+    def _build_sar3d_plot_page(self) -> None:
+        """Create 3D SAR page with stacked side (X-Z) and top (X-Y) projection plots."""
+        # 3D SAR surface: container with two projection heatmaps, mirroring the
+        # pass-through page layout (two stacked PlotWidgets on one stack page).
+        self._sar3d_plots_container = QWidget()
+        sar3d_layout = QVBoxLayout(self._sar3d_plots_container)
+        sar3d_layout.setContentsMargins(0, 0, 0, 0)
+        sar3d_layout.setSpacing(6)
+
+        self._sar3d_xz_plot = self._create_plot_widget(background="#0f141c")
+        self._sar3d_xz_plot.showGrid(x=True, y=True, alpha=0.2)
+        sar3d_layout.addWidget(self._sar3d_xz_plot, stretch=1)
+
+        self._sar3d_xy_plot = self._create_plot_widget(background="#0f141c")
+        self._sar3d_xy_plot.showGrid(x=True, y=True, alpha=0.2)
+        sar3d_layout.addWidget(self._sar3d_xy_plot, stretch=1)
+
+        self._configure_sar3d_plot_axes()
+        self._clear_sar3d_plot()
+        self._plot_stack.addWidget(self._sar3d_plots_container)
+
     def _build_settings_toggle(self, root_layout: QHBoxLayout) -> None:
         """Create narrow button used to collapse or show settings panel."""
         self._settings_toggle_button = QPushButton("<")
@@ -286,6 +308,7 @@ class AppWindowUiMixin:
 
         `bscan`   -> show `self._bscan_plot`
         `gpr` and `legacy_gpr` -> show `self._gpr_plot`
+        `sar3d`   -> show `self._sar3d_plots_container`
         otherwise -> show `self._trace_plots_container`
         """
         if mode == "bscan":
@@ -294,6 +317,9 @@ class AppWindowUiMixin:
         if mode in {"gpr", "legacy_gpr"}:
             self._plot_stack.setCurrentWidget(self._gpr_plot)
             return
+        if mode == "sar3d":
+            self._plot_stack.setCurrentWidget(self._sar3d_plots_container)
+            return
         self._plot_stack.setCurrentWidget(self._trace_plots_container)
 
     @staticmethod

python_app/gui/controllers/sections/processing_section.py

@@ -20,6 +20,10 @@ from PyQt6.QtWidgets import (
 
 from python_app.gui.controllers.app_window_config.state_builders import GUI_SAVE_HISTORY_LIMIT
 from python_app.gui.controllers.sections.layout_helpers import FormRow, build_two_column_form_widget
+from python_app.gui.runtime.sar3d_views import (
+    SAR3D_BOTTOM_VIEW_CHOICES,
+    sar3d_bottom_view_label_from_deg,
+)
 
 
 def _format_geometry_row(position: int, x_m: float, y_m: float, z_m: float) -> str:
@@ -59,7 +63,7 @@ def _build_processing_mode_page(parent: QStackedWidget, rows: list[FormRow], *,
 
 
 def build_processing_group(owner) -> QGroupBox:
-    """Create processing mode section with pass-through, B-scan, and GPR pages."""
+    """Create processing mode section with pass-through, B-scan, GPR, and 3D SAR pages."""
     group = QGroupBox("Processing")
     form = QFormLayout(group)
     form.setFieldGrowthPolicy(QFormLayout.FieldGrowthPolicy.AllNonFixedFieldsGrow)
@@ -68,9 +72,10 @@ def build_processing_group(owner) -> QGroupBox:
     bscan_defaults = processing_defaults.bscan
     gpr_live_defaults = processing_defaults.gpr
     legacy_gpr_defaults = processing_defaults.legacy_gpr
+    sar3d_defaults = processing_defaults.sar3d
 
     owner._processing_mode = QComboBox()
-    owner._processing_mode.addItems(["pass_through", "bscan", "gpr", "legacy_gpr"])
+    owner._processing_mode.addItems(["pass_through", "bscan", "gpr", "legacy_gpr", "sar3d"])
     owner._set_combo_current_text(owner._processing_mode, processing_defaults.selected_mode)
 
     owner._processing_mode_pages = QStackedWidget(group)
@@ -570,6 +575,185 @@ def build_processing_group(owner) -> QGroupBox:
     )
     owner._processing_mode_pages.addWidget(legacy_gpr_page)
 
+    # 3D SAR page: shared gpr_* knobs get their own widget set (like legacy_gpr)
+    # plus the sar3d-specific aperture/volume knobs consumed by the C++ processor.
+    owner._sar3d_input_positions_input = QLineEdit(str(sar3d_defaults.input_positions))
+    owner._sar3d_input_positions_input.setPlaceholderText("0,1,2")
+
+    owner._sar3d_output_positions_input = QLineEdit(str(sar3d_defaults.output_positions))
+    owner._sar3d_output_positions_input.setPlaceholderText("0,1")
+
+    owner._sar3d_min_depth_m = QDoubleSpinBox()
+    owner._sar3d_min_depth_m.setDecimals(2)
+    owner._sar3d_min_depth_m.setRange(0.0, 50.0)
+    owner._sar3d_min_depth_m.setSingleStep(0.1)
+    owner._sar3d_min_depth_m.setValue(float(sar3d_defaults.min_depth_m))
+
+    owner._sar3d_max_depth_m = QDoubleSpinBox()
+    owner._sar3d_max_depth_m.setDecimals(2)
+    owner._sar3d_max_depth_m.setRange(0.1, 50.0)
+    owner._sar3d_max_depth_m.setSingleStep(0.1)
+    owner._sar3d_max_depth_m.setValue(float(sar3d_defaults.max_depth_m))
+
+    owner._sar3d_start_freq_mhz = QDoubleSpinBox()
+    owner._sar3d_start_freq_mhz.setDecimals(1)
+    owner._sar3d_start_freq_mhz.setRange(100.0, 8800.0)
+    owner._sar3d_start_freq_mhz.setSingleStep(10.0)
+    owner._sar3d_start_freq_mhz.setValue(float(sar3d_defaults.start_freq_mhz))
+
+    owner._sar3d_stop_freq_mhz = QDoubleSpinBox()
+    owner._sar3d_stop_freq_mhz.setDecimals(1)
+    owner._sar3d_stop_freq_mhz.setRange(100.0, 8800.0)
+    owner._sar3d_stop_freq_mhz.setSingleStep(10.0)
+    owner._sar3d_stop_freq_mhz.setValue(float(sar3d_defaults.stop_freq_mhz))
+
+    owner._sar3d_look_angle_deg = QDoubleSpinBox()
+    owner._sar3d_look_angle_deg.setDecimals(2)
+    owner._sar3d_look_angle_deg.setRange(-90.0, 90.0)
+    owner._sar3d_look_angle_deg.setSingleStep(0.5)
+    owner._sar3d_look_angle_deg.setValue(float(sar3d_defaults.look_angle_deg))
+    owner._sar3d_look_angle_deg.setToolTip(
+        "Radar look angle vs the range axis (deg). Used by intra-sweep motion correction."
+    )
+
+    owner._sar3d_direction_sign = QDoubleSpinBox()
+    owner._sar3d_direction_sign.setDecimals(0)
+    owner._sar3d_direction_sign.setRange(-1.0, 1.0)
+    owner._sar3d_direction_sign.setSingleStep(2.0)
+    owner._sar3d_direction_sign.setValue(float(sar3d_defaults.direction_sign))
+    owner._sar3d_direction_sign.setToolTip(
+        "Motion direction along range: +1 (later frequencies deeper) or -1 (shallower)."
+    )
+
+    owner._sar3d_ignore_socket_speed_enabled = QCheckBox("Ignore socket speed")
+    owner._sar3d_ignore_socket_speed_enabled.setChecked(bool(sar3d_defaults.ignore_socket_speed_enabled))
+    owner._sar3d_ignore_socket_speed_enabled.setToolTip(
+        "When checked, use the speed below for motion correction instead of the "
+        "value streamed over the locator socket."
+    )
+
+    owner._sar3d_speed_m_s = QDoubleSpinBox()
+    owner._sar3d_speed_m_s.setDecimals(3)
+    owner._sar3d_speed_m_s.setRange(-100.0, 100.0)
+    owner._sar3d_speed_m_s.setSingleStep(0.05)
+    owner._sar3d_speed_m_s.setValue(float(sar3d_defaults.speed_m_s))
+    owner._sar3d_speed_m_s.setToolTip(
+        "Radar speed (m/s) used by intra-sweep motion correction when 'Ignore socket speed' is on."
+    )
+
+    owner._sar3d_apply_freq_phase_correction = QCheckBox("Apply intra-sweep phase correction")
+    owner._sar3d_apply_freq_phase_correction.setChecked(bool(sar3d_defaults.apply_freq_phase_correction))
+
+    owner._sar3d_frame_radius = QSpinBox()
+    owner._sar3d_frame_radius.setRange(0, 25)
+    owner._sar3d_frame_radius.setValue(int(sar3d_defaults.frame_radius))
+    owner._sar3d_frame_radius.setToolTip(
+        "Synthetic-aperture half width in frames: 2*radius+1 consecutive frames form one aperture."
+    )
+
+    owner._sar3d_edge_frame_weight = QDoubleSpinBox()
+    owner._sar3d_edge_frame_weight.setDecimals(2)
+    owner._sar3d_edge_frame_weight.setRange(0.0, 1.0)
+    owner._sar3d_edge_frame_weight.setSingleStep(0.05)
+    owner._sar3d_edge_frame_weight.setValue(float(sar3d_defaults.edge_frame_weight))
+    owner._sar3d_edge_frame_weight.setToolTip(
+        "Weight of the outermost aperture frames (1 = flat window, 0 = strong taper)."
+    )
+
+    owner._sar3d_pitch_deg = QDoubleSpinBox()
+    owner._sar3d_pitch_deg.setDecimals(2)
+    owner._sar3d_pitch_deg.setRange(-90.0, 90.0)
+    owner._sar3d_pitch_deg.setSingleStep(0.5)
+    owner._sar3d_pitch_deg.setValue(float(sar3d_defaults.pitch_deg))
+    owner._sar3d_pitch_deg.setToolTip("Antenna pitch above the horizon (deg); tilts the imaged volume.")
+
+    owner._sar3d_center_depth_m = QDoubleSpinBox()
+    owner._sar3d_center_depth_m.setDecimals(2)
+    owner._sar3d_center_depth_m.setRange(0.0, 50.0)
+    owner._sar3d_center_depth_m.setSingleStep(0.1)
+    owner._sar3d_center_depth_m.setValue(float(sar3d_defaults.center_depth_m))
+    owner._sar3d_center_depth_m.setToolTip("Range distance to the center of the imaged volume (m).")
+
+    owner._sar3d_half_width_x_m = QDoubleSpinBox()
+    owner._sar3d_half_width_x_m.setDecimals(2)
+    owner._sar3d_half_width_x_m.setRange(0.01, 50.0)
+    owner._sar3d_half_width_x_m.setSingleStep(0.1)
+    owner._sar3d_half_width_x_m.setValue(float(sar3d_defaults.half_width_x_m))
+
+    owner._sar3d_half_width_y_m = QDoubleSpinBox()
+    owner._sar3d_half_width_y_m.setDecimals(2)
+    owner._sar3d_half_width_y_m.setRange(0.01, 50.0)
+    owner._sar3d_half_width_y_m.setSingleStep(0.1)
+    owner._sar3d_half_width_y_m.setValue(float(sar3d_defaults.half_width_y_m))
+
+    owner._sar3d_half_width_z_m = QDoubleSpinBox()
+    owner._sar3d_half_width_z_m.setDecimals(2)
+    owner._sar3d_half_width_z_m.setRange(0.01, 50.0)
+    owner._sar3d_half_width_z_m.setSingleStep(0.1)
+    owner._sar3d_half_width_z_m.setValue(float(sar3d_defaults.half_width_z_m))
+
+    owner._sar3d_grid_nx = QSpinBox()
+    owner._sar3d_grid_nx.setRange(2, 1024)
+    owner._sar3d_grid_nx.setValue(int(sar3d_defaults.grid_nx))
+
+    owner._sar3d_grid_ny = QSpinBox()
+    owner._sar3d_grid_ny.setRange(2, 1024)
+    owner._sar3d_grid_ny.setValue(int(sar3d_defaults.grid_ny))
+
+    owner._sar3d_grid_nz = QSpinBox()
+    owner._sar3d_grid_nz.setRange(2, 1024)
+    owner._sar3d_grid_nz.setValue(int(sar3d_defaults.grid_nz))
+
+    owner._sar3d_use_compensation = QCheckBox("Use per-voxel compensation")
+    owner._sar3d_use_compensation.setChecked(bool(sar3d_defaults.use_compensation))
+
+    owner._sar3d_normalize_pairs = QCheckBox("Normalize Tx/Rx pairs")
+    owner._sar3d_normalize_pairs.setChecked(bool(sar3d_defaults.normalize_pairs))
+
+    owner._sar3d_bottom_view = QComboBox()
+    owner._sar3d_bottom_view.addItems([label for label, _deg in SAR3D_BOTTOM_VIEW_CHOICES])
+    owner._set_combo_current_text(
+        owner._sar3d_bottom_view,
+        sar3d_bottom_view_label_from_deg(float(sar3d_defaults.bottom_view_deg)),
+    )
+    owner._sar3d_bottom_view.setToolTip(
+        "Bottom plot projection: front view (X-Y), side view (Z-Y) or an "
+        "intermediate rotation around the vertical axis. The top plot always "
+        "shows the X-Z top view."
+    )
+
+    sar3d_page = _build_processing_mode_page(
+        owner._processing_mode_pages,
+        [
+            ("Input positions", owner._sar3d_input_positions_input),
+            ("Output positions", owner._sar3d_output_positions_input),
+            ("Min depth m", owner._sar3d_min_depth_m),
+            ("Max depth m", owner._sar3d_max_depth_m),
+            ("Start MHz", owner._sar3d_start_freq_mhz),
+            ("Stop MHz", owner._sar3d_stop_freq_mhz),
+            ("Look angle deg", owner._sar3d_look_angle_deg),
+            ("Direction sign", owner._sar3d_direction_sign),
+            owner._sar3d_ignore_socket_speed_enabled,
+            ("Speed m/s", owner._sar3d_speed_m_s),
+            owner._sar3d_apply_freq_phase_correction,
+            ("Frame radius", owner._sar3d_frame_radius),
+            ("Edge frame weight", owner._sar3d_edge_frame_weight),
+            ("Pitch deg", owner._sar3d_pitch_deg),
+            ("Center depth m", owner._sar3d_center_depth_m),
+            ("Half width X m", owner._sar3d_half_width_x_m),
+            ("Half width Y m", owner._sar3d_half_width_y_m),
+            ("Half width Z m", owner._sar3d_half_width_z_m),
+            ("Grid Nx", owner._sar3d_grid_nx),
+            ("Grid Ny", owner._sar3d_grid_ny),
+            ("Grid Nz", owner._sar3d_grid_nz),
+            ("Bottom view", owner._sar3d_bottom_view),
+            owner._sar3d_use_compensation,
+            owner._sar3d_normalize_pairs,
+        ],
+        split_index=11,
+    )
+    owner._processing_mode_pages.addWidget(sar3d_page)
+
     owner._processing_mode.currentTextChanged.connect(owner._on_processing_mode_changed)
     owner._show_magnitude_checkbox.toggled.connect(owner._on_trace_visibility_changed)
     owner._show_phase_checkbox.toggled.connect(owner._on_trace_visibility_changed)
@@ -642,6 +826,30 @@ def build_processing_group(owner) -> QGroupBox:
     owner._legacy_gpr_visible_x_max_m.valueChanged.connect(owner._on_gpr_locator_window_changed)
     owner._legacy_gpr_visible_z_min_m.valueChanged.connect(owner._on_gpr_locator_window_changed)
     owner._legacy_gpr_visible_z_max_m.valueChanged.connect(owner._on_gpr_locator_window_changed)
+    owner._sar3d_input_positions_input.editingFinished.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_output_positions_input.editingFinished.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_min_depth_m.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_max_depth_m.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_start_freq_mhz.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_stop_freq_mhz.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_look_angle_deg.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_direction_sign.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_ignore_socket_speed_enabled.toggled.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_speed_m_s.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_apply_freq_phase_correction.toggled.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_frame_radius.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_edge_frame_weight.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_pitch_deg.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_center_depth_m.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_half_width_x_m.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_half_width_y_m.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_half_width_z_m.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_grid_nx.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_grid_ny.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_grid_nz.valueChanged.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_use_compensation.toggled.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_normalize_pairs.toggled.connect(owner._on_processing_live_settings_changed)
+    owner._sar3d_bottom_view.currentTextChanged.connect(owner._on_processing_live_settings_changed)
 
     owner._set_processing_mode_page(owner._processing_mode.currentText())
     return group

python_app/gui/runtime/constraints.py

@@ -24,7 +24,9 @@ def validate_processing_mode_constraints(
             )
         return
 
-    if processing_mode not in {"gpr", "legacy_gpr"}:
+    # sar3d shares the GPR combo/geometry requirements: it images from the same
+    # multi-Tx/Rx pair set selected through the gpr_* live-config fields.
+    if processing_mode not in {"gpr", "legacy_gpr", "sar3d"}:
         return
 
     configured_combos = {(int(combo.input), int(combo.output)) for combo in config.combos}

python_app/gui/runtime/sar3d_views.py

@@ -0,0 +1,30 @@
+"""Shared mapping for the sar3d bottom-view projection choices.
+
+The bottom plot of the sar3d mode projects the imaged volume along a
+horizontal direction rotated around the vertical Y axis: 0° = front view
+(X horizontal), 90° = side view (Z horizontal). One shared mapping keeps the
+combo box, the live-config build, the profile persistence and the plot labels
+consistent.
+"""
+
+from __future__ import annotations
+
+SAR3D_BOTTOM_VIEW_CHOICES: tuple[tuple[str, float], ...] = (
+    ("front (0°)", 0.0),
+    ("30°", 30.0),
+    ("60°", 60.0),
+    ("side (90°)", 90.0),
+)
+
+
+def sar3d_bottom_view_deg_from_label(label: str) -> float:
+    """Map a bottom-view combo label to its projection angle in degrees."""
+    for choice_label, degrees in SAR3D_BOTTOM_VIEW_CHOICES:
+        if choice_label == label:
+            return degrees
+    return 0.0
+
+
+def sar3d_bottom_view_label_from_deg(degrees: float) -> str:
+    """Map a projection angle to the nearest bottom-view combo label."""
+    return min(SAR3D_BOTTOM_VIEW_CHOICES, key=lambda choice: abs(choice[1] - float(degrees)))[0]

python_app/hardware_full/laser_control/exceptions.py

@@ -68,6 +68,24 @@ class PortNotFoundError(CommunicationError):
         super().__init__(message)
 
 
+class PortBusyError(CommunicationError):
+    """Raised when the serial port is already held by another program.
+
+    A laser controller is a single half-duplex request/response bus: only one
+    process may own it at a time. When the port is opened exclusively
+    (``exclusive=True``), a second exclusive open fails with this error instead
+    of silently sharing the bus and corrupting both sides' traffic.
+    """
+
+    def __init__(self, port: str = None):
+        target = f"'{port}'" if port else "the laser controller port"
+        super().__init__(
+            f"Serial port {target} is busy — another program is already using "
+            f"the laser controller. Close it first (only one program may own the "
+            f"port at a time)."
+        )
+
+
 class DeviceNotRespondingError(CommunicationError):
     """Exception raised when device doesn't respond to commands."""
 

python_app/hardware_full/laser_control/protocol.py

@@ -33,6 +33,7 @@ from .models import Measurements, VariationType
 from .exceptions import (
     CommunicationError,
     PortNotFoundError,
+    PortBusyError,
     CRCError,
     ProtocolError,
 )
@@ -184,20 +185,36 @@ class Protocol:
     # ---- Connection management
 
     def connect(self) -> None:
-        """Open the serial port. Auto-detects the device path when port is None."""
+        """Open the serial port. Auto-detects the device path when port is None.
+
+        The port is opened with ``exclusive=True`` so only one program can own the
+        laser controller at a time. This is required because the board is a single
+        half-duplex request/response bus: if two processes opened it, their
+        commands and responses would interleave and corrupt each other. A second
+        opener fails fast with :class:`PortBusyError` instead of silently sharing
+        the bus. Note the lock is advisory (``flock``): every program that opens
+        this port must also use ``exclusive=True`` for the guarantee to hold.
+        """
         port = self._port_name or self._detect_port()
         try:
             self._serial = serial.Serial(
                 port=port,
                 baudrate=BAUDRATE,
                 timeout=SERIAL_TIMEOUT_SEC,
+                exclusive=True,
             )
+        except serial.SerialException as exc:
+            # pyserial raises SerialException with errno EAGAIN/EWOULDBLOCK (11)
+            # when another process already holds the exclusive flock.
+            if getattr(exc, "errno", None) in (11,) or "exclusively lock" in str(exc).lower():
+                logger.error("Laser serial port '%s' is busy (held by another program)", port)
+                raise PortBusyError(port) from exc
+            logger.error("Cannot open laser serial port '%s': %s", port, exc)
+            raise CommunicationError(f"Cannot connect to port '{port}': {exc}") from exc
         except Exception as exc:
             logger.error("Cannot open laser serial port '%s': %s", port, exc)
-            raise CommunicationError(
-                f"Cannot connect to port '{port}': {exc}"
-            ) from exc
-        logger.debug("Laser serial port opened: %s @ %d baud", port, BAUDRATE)
+            raise CommunicationError(f"Cannot connect to port '{port}': {exc}") from exc
+        logger.debug("Laser serial port opened (exclusive): %s @ %d baud", port, BAUDRATE)
 
     def disconnect(self) -> None:
         """Close the serial port if open."""

python_app/models/gui_profile_codec.py

@@ -14,6 +14,7 @@ from python_app.models.gui_profile_schema import (
     GuiPreprocessDialogStateModel,
     GuiProcessingStateModel,
     GuiProfileModel,
+    GuiSar3dStateModel,
     GuiStateModel,
     GuiSwitchStateModel,
 )
@@ -106,6 +107,7 @@ def gui_profile_from_dict(payload: dict[str, Any]) -> GuiProfileModel:
     bscan_object = _as_dict(processing_object.get("bscan"), "gui.processing.bscan")
     gpr_object = _as_dict(processing_object.get("gpr"), "gui.processing.gpr")
     legacy_gpr_object = _as_dict(processing_object.get("legacy_gpr"), "gui.processing.legacy_gpr")
+    sar3d_object = _as_dict(processing_object.get("sar3d"), "gui.processing.sar3d")
     root_gpr_object = payload.get("gpr")
 
     selected_mode = _optional_string(
@@ -452,9 +454,142 @@ def gui_profile_from_dict(payload: dict[str, Any]) -> GuiProfileModel:
             visible_z_min_m=legacy_float("visible_z_min_m", gui.processing.legacy_gpr.visible_z_min_m),
             visible_z_max_m=legacy_float("visible_z_max_m", gui.processing.legacy_gpr.visible_z_max_m),
         ),
+        sar3d=GuiSar3dStateModel(
+            input_positions=_optional_string(
+                sar3d_object,
+                "input_positions",
+                gui.processing.sar3d.input_positions,
+                "gui.processing.sar3d",
+            ),
+            output_positions=_optional_string(
+                sar3d_object,
+                "output_positions",
+                gui.processing.sar3d.output_positions,
+                "gui.processing.sar3d",
+            ),
+            min_depth_m=_optional_float(
+                sar3d_object,
+                "min_depth_m",
+                gui.processing.sar3d.min_depth_m,
+                "gui.processing.sar3d",
+            ),
+            max_depth_m=_optional_float(
+                sar3d_object,
+                "max_depth_m",
+                gui.processing.sar3d.max_depth_m,
+                "gui.processing.sar3d",
+            ),
+            start_freq_mhz=_optional_float(
+                sar3d_object,
+                "start_freq_mhz",
+                gui.processing.sar3d.start_freq_mhz,
+                "gui.processing.sar3d",
+            ),
+            stop_freq_mhz=_optional_float(
+                sar3d_object,
+                "stop_freq_mhz",
+                gui.processing.sar3d.stop_freq_mhz,
+                "gui.processing.sar3d",
+            ),
+            speed_m_s=_optional_float(
+                sar3d_object,
+                "speed_m_s",
+                gui.processing.sar3d.speed_m_s,
+                "gui.processing.sar3d",
+            ),
+            ignore_socket_speed_enabled=_optional_bool(
+                sar3d_object,
+                "ignore_socket_speed_enabled",
+                gui.processing.sar3d.ignore_socket_speed_enabled,
+                "gui.processing.sar3d",
+            ),
+            look_angle_deg=_optional_float(
+                sar3d_object,
+                "look_angle_deg",
+                gui.processing.sar3d.look_angle_deg,
+                "gui.processing.sar3d",
+            ),
+            direction_sign=_optional_float(
+                sar3d_object,
+                "direction_sign",
+                gui.processing.sar3d.direction_sign,
+                "gui.processing.sar3d",
+            ),
+            apply_freq_phase_correction=_optional_bool(
+                sar3d_object,
+                "apply_freq_phase_correction",
+                gui.processing.sar3d.apply_freq_phase_correction,
+                "gui.processing.sar3d",
+            ),
+            frame_radius=_optional_int(
+                sar3d_object,
+                "frame_radius",
+                gui.processing.sar3d.frame_radius,
+                "gui.processing.sar3d",
+            ),
+            edge_frame_weight=_optional_float(
+                sar3d_object,
+                "edge_frame_weight",
+                gui.processing.sar3d.edge_frame_weight,
+                "gui.processing.sar3d",
+            ),
+            pitch_deg=_optional_float(
+                sar3d_object,
+                "pitch_deg",
+                gui.processing.sar3d.pitch_deg,
+                "gui.processing.sar3d",
+            ),
+            center_depth_m=_optional_float(
+                sar3d_object,
+                "center_depth_m",
+                gui.processing.sar3d.center_depth_m,
+                "gui.processing.sar3d",
+            ),
+            half_width_x_m=_optional_float(
+                sar3d_object,
+                "half_width_x_m",
+                gui.processing.sar3d.half_width_x_m,
+                "gui.processing.sar3d",
+            ),
+            half_width_y_m=_optional_float(
+                sar3d_object,
+                "half_width_y_m",
+                gui.processing.sar3d.half_width_y_m,
+                "gui.processing.sar3d",
+            ),
+            half_width_z_m=_optional_float(
+                sar3d_object,
+                "half_width_z_m",
+                gui.processing.sar3d.half_width_z_m,
+                "gui.processing.sar3d",
+            ),
+            grid_nx=_optional_int(sar3d_object, "grid_nx", gui.processing.sar3d.grid_nx, "gui.processing.sar3d"),
+            grid_ny=_optional_int(sar3d_object, "grid_ny", gui.processing.sar3d.grid_ny, "gui.processing.sar3d"),
+            grid_nz=_optional_int(sar3d_object, "grid_nz", gui.processing.sar3d.grid_nz, "gui.processing.sar3d"),
+            use_compensation=_optional_bool(
+                sar3d_object,
+                "use_compensation",
+                gui.processing.sar3d.use_compensation,
+                "gui.processing.sar3d",
+            ),
+            normalize_pairs=_optional_bool(
+                sar3d_object,
+                "normalize_pairs",
+                gui.processing.sar3d.normalize_pairs,
+                "gui.processing.sar3d",
+            ),
+            bottom_view_deg=_optional_float(
+                sar3d_object,
+                "bottom_view_deg",
+                gui.processing.sar3d.bottom_view_deg,
+                "gui.processing.sar3d",
+            ),
+        ),
     )
-    if gui.processing.selected_mode not in {"pass_through", "bscan", "gpr", "legacy_gpr"}:
-        raise ValueError("gui.processing.selected_mode must be one of: pass_through, bscan, gpr, legacy_gpr")
+    if gui.processing.selected_mode not in {"pass_through", "bscan", "gpr", "legacy_gpr", "sar3d"}:
+        raise ValueError(
+            "gui.processing.selected_mode must be one of: pass_through, bscan, gpr, legacy_gpr, sar3d"
+        )
     if gui.processing.bscan.axis not in {"abs", "real", "phase"}:
         raise ValueError("gui.processing.bscan.axis must be one of: abs, real, phase")
     if gui.processing.gpr.render_mode not in {"heatmap", "objects_only"}:
@@ -489,6 +624,12 @@ def gui_profile_from_dict(payload: dict[str, Any]) -> GuiProfileModel:
         raise ValueError("gui.processing.legacy_gpr.snr_comp_max must be >= 0")
     if gui.processing.legacy_gpr.min_visible_pair_count < 1:
         raise ValueError("gui.processing.legacy_gpr.min_visible_pair_count must be >= 1")
+    if gui.processing.sar3d.frame_radius < 0:
+        raise ValueError("gui.processing.sar3d.frame_radius must be >= 0")
+    if not 0.0 <= gui.processing.sar3d.edge_frame_weight <= 1.0:
+        raise ValueError("gui.processing.sar3d.edge_frame_weight must be within [0, 1]")
+    if min(gui.processing.sar3d.grid_nx, gui.processing.sar3d.grid_ny, gui.processing.sar3d.grid_nz) < 2:
+        raise ValueError("gui.processing.sar3d.grid_nx/ny/nz must each be >= 2")
 
     data_actions_object = _as_dict(gui_object.get("data_actions"), "gui.data_actions")
     gui.data_actions = GuiDataActionsStateModel(
@@ -640,6 +781,32 @@ def gui_profile_to_dict(model: GuiProfileModel) -> dict[str, Any]:
                 "visible_z_min_m": gui.processing.legacy_gpr.visible_z_min_m,
                 "visible_z_max_m": gui.processing.legacy_gpr.visible_z_max_m,
             },
+            "sar3d": {
+                "input_positions": gui.processing.sar3d.input_positions,
+                "output_positions": gui.processing.sar3d.output_positions,
+                "min_depth_m": gui.processing.sar3d.min_depth_m,
+                "max_depth_m": gui.processing.sar3d.max_depth_m,
+                "start_freq_mhz": gui.processing.sar3d.start_freq_mhz,
+                "stop_freq_mhz": gui.processing.sar3d.stop_freq_mhz,
+                "speed_m_s": gui.processing.sar3d.speed_m_s,
+                "ignore_socket_speed_enabled": gui.processing.sar3d.ignore_socket_speed_enabled,
+                "look_angle_deg": gui.processing.sar3d.look_angle_deg,
+                "direction_sign": gui.processing.sar3d.direction_sign,
+                "apply_freq_phase_correction": gui.processing.sar3d.apply_freq_phase_correction,
+                "frame_radius": gui.processing.sar3d.frame_radius,
+                "edge_frame_weight": gui.processing.sar3d.edge_frame_weight,
+                "pitch_deg": gui.processing.sar3d.pitch_deg,
+                "center_depth_m": gui.processing.sar3d.center_depth_m,
+                "half_width_x_m": gui.processing.sar3d.half_width_x_m,
+                "half_width_y_m": gui.processing.sar3d.half_width_y_m,
+                "half_width_z_m": gui.processing.sar3d.half_width_z_m,
+                "grid_nx": gui.processing.sar3d.grid_nx,
+                "grid_ny": gui.processing.sar3d.grid_ny,
+                "grid_nz": gui.processing.sar3d.grid_nz,
+                "use_compensation": gui.processing.sar3d.use_compensation,
+                "normalize_pairs": gui.processing.sar3d.normalize_pairs,
+                "bottom_view_deg": gui.processing.sar3d.bottom_view_deg,
+            },
         },
         "data_actions": {
             "save_count": gui.data_actions.save_count,

python_app/models/gui_profile_model.py

@@ -10,6 +10,7 @@ from python_app.models.gui_profile_schema import (
     GuiPreprocessDialogStateModel,
     GuiProcessingStateModel,
     GuiProfileModel,
+    GuiSar3dStateModel,
     GuiStateModel,
     GuiSwitchStateModel,
 )
@@ -23,6 +24,7 @@ __all__ = [
     "GuiPreprocessDialogStateModel",
     "GuiProcessingStateModel",
     "GuiProfileModel",
+    "GuiSar3dStateModel",
     "GuiStateModel",
     "GuiSwitchStateModel",
     "gui_profile_from_dict",

python_app/models/gui_profile_schema.py

@@ -123,6 +123,43 @@ class GuiLegacyGprStateModel:
     visible_z_max_m: float = 14.0
 
 
+@dataclass(slots=True)
+class GuiSar3dStateModel:
+    """UI-only defaults for the 3D SAR (sar3d) live settings.
+
+    The band, depth gate, antenna selection and motion knobs mirror the shared
+    gpr_* live-config fields; the sar3d_* knobs mirror the VNA_AlGo gpr_core
+    defaults used by the C++ sar3d processor.
+    """
+
+    input_positions: str = ""
+    output_positions: str = ""
+    min_depth_m: float = 2.0
+    max_depth_m: float = 14.0
+    start_freq_mhz: float = 3000.0
+    stop_freq_mhz: float = 6000.0
+    speed_m_s: float = 0.0
+    ignore_socket_speed_enabled: bool = False
+    look_angle_deg: float = 0.0
+    direction_sign: float = 1.0
+    apply_freq_phase_correction: bool = True
+    frame_radius: int = 6
+    edge_frame_weight: float = 0.75
+    pitch_deg: float = 6.0
+    center_depth_m: float = 5.6
+    half_width_x_m: float = 1.8
+    half_width_y_m: float = 1.0
+    half_width_z_m: float = 1.8
+    grid_nx: int = 150
+    grid_ny: int = 120
+    grid_nz: int = 160
+    use_compensation: bool = True
+    normalize_pairs: bool = True
+    # Bottom-view projection angle around the vertical Y axis, degrees:
+    # 0 = front view (X-Y), 90 = side view (Z-Y).
+    bottom_view_deg: float = 0.0
+
+
 @dataclass(slots=True)
 class GuiProcessingStateModel:
     """UI-only processing-section defaults."""
@@ -132,6 +169,7 @@ class GuiProcessingStateModel:
     bscan: GuiBscanStateModel = field(default_factory=GuiBscanStateModel)
     gpr: GuiGprStateModel = field(default_factory=GuiGprStateModel)
     legacy_gpr: GuiLegacyGprStateModel = field(default_factory=GuiLegacyGprStateModel)
+    sar3d: GuiSar3dStateModel = field(default_factory=GuiSar3dStateModel)
 
 
 @dataclass(slots=True)

python_app/models/run_config_codec.py

@@ -175,6 +175,7 @@ def run_config_from_dict(payload: dict[str, Any]) -> RunConfigModel:
         "run.locator_server",
     )
     multi_device_payload = _as_dict(radar_payload.get("multi_device"), "radar.multi_device")
+    replay_payload = _as_dict(radar_payload.get("replay"), "radar.replay")
     kamil_adc_payload = _as_dict(radar_payload.get("kamil_adc"), "radar.kamil_adc")
     laser_control_payload = _as_dict(radar_payload.get("laser_control"), "radar.laser_control")
 
@@ -235,6 +236,13 @@ def run_config_from_dict(payload: dict[str, Any]) -> RunConfigModel:
         "stimulus_switch_positions",
         model.radar.multi_device.stimulus_switch_positions,
     )
+    model.radar.replay.dataset_dir = _read_str(
+        replay_payload, "dataset_dir", model.radar.replay.dataset_dir
+    )
+    model.radar.replay.loop = _read_bool(replay_payload, "loop", model.radar.replay.loop)
+    model.radar.replay.rate_scale = _read_float(
+        replay_payload, "rate_scale", model.radar.replay.rate_scale
+    )
     model.radar.kamil_adc.project_dir = _read_str(
         kamil_adc_payload, "project_dir", model.radar.kamil_adc.project_dir
     )
@@ -519,6 +527,11 @@ def run_config_to_dict(model: RunConfigModel) -> dict[str, Any]:
                 "input_switch_positions": model.radar.multi_device.input_switch_positions,
                 "stimulus_switch_positions": model.radar.multi_device.stimulus_switch_positions
             },
+            "replay": {
+                "dataset_dir": model.radar.replay.dataset_dir,
+                "loop": model.radar.replay.loop,
+                "rate_scale": model.radar.replay.rate_scale,
+            },
             "kamil_adc": {
                 "project_dir": model.radar.kamil_adc.project_dir,
                 "executable_path": model.radar.kamil_adc.executable_path,

python_app/models/run_config_schema.py

@@ -52,6 +52,25 @@ class RadarMultiDeviceModel:
     stimulus_switch_positions: int = 1
 
 
+@dataclass(slots=True)
+class RadarReplayModel:
+    """Recorded-dataset replay settings (no hardware).
+
+    When ``dataset_dir`` is non-empty the supervisor launches
+    ``python_app.scripts.replay_raw_producer`` instead of the model's normal
+    acquisition producer. The dataset is a directory of snapshot folders named
+    ``NNNN_idK_ns<timestamp>``, each holding ``i{input}_o{output}_s21.npy`` /
+    ``_freq.npy`` files (the VNA_AlGo preprocessed-capture layout).
+    """
+
+    dataset_dir: str = ""
+    # Restart from the first snapshot when the dataset ends; the replay
+    # timeline keeps advancing monotonically across the seam.
+    loop: bool = True
+    # Playback speed multiplier: 1.0 replays at the recorded frame cadence.
+    rate_scale: float = 1.0
+
+
 @dataclass(slots=True)
 class KamilAdcPhaseCalibrationModel:
     """Affine law mapping the reference signal's unwrapped phase to frequency.
@@ -159,6 +178,7 @@ class RadarModel:
     visa_library: str = ""
     sweep: RadarSweepModel = field(default_factory=RadarSweepModel)
     multi_device: RadarMultiDeviceModel = field(default_factory=RadarMultiDeviceModel)
+    replay: RadarReplayModel = field(default_factory=RadarReplayModel)
     kamil_adc: KamilAdcModel = field(default_factory=KamilAdcModel)
     laser_control: LaserControlModel = field(default_factory=LaserControlModel)
 

python_app/orchestration/live_processing_config.py

@@ -72,6 +72,24 @@ class ProcessingLiveConfig:
     gpr_visible_x_max_m: float = 2.0
     gpr_visible_z_min_m: float = 0.0
     gpr_visible_z_max_m: float = 14.0
+    # 3D SAR (sar3d processor) parameters; defaults mirror VNA_AlGo gpr_core.
+    # The frequency band, depth gate, antenna selection and motion parameters
+    # are shared with the gpr_* fields above.
+    sar3d_frame_radius: int = 6
+    sar3d_edge_frame_weight: float = 0.75
+    sar3d_pitch_deg: float = 6.0
+    sar3d_center_depth_m: float = 5.6
+    sar3d_half_width_x_m: float = 1.8
+    sar3d_half_width_y_m: float = 1.0
+    sar3d_half_width_z_m: float = 1.8
+    sar3d_grid_nx: int = 150
+    sar3d_grid_ny: int = 120
+    sar3d_grid_nz: int = 160
+    sar3d_use_compensation: bool = True
+    sar3d_normalize_pairs: bool = True
+    # Bottom-view projection angle around the vertical Y axis, degrees:
+    # 0 = front view (X horizontal), 90 = side view (Z horizontal).
+    sar3d_bottom_view_deg: float = 0.0
     # When true, the C++ data_processor ignores socket-supplied vlc updates
     # and keeps using `gpr_speed_m_s` from this file.
     ignore_socket_speed: bool = False
@@ -136,6 +154,19 @@ class ProcessingLiveConfig:
             "gpr_visible_x_max_m": float(self.gpr_visible_x_max_m),
             "gpr_visible_z_min_m": float(self.gpr_visible_z_min_m),
             "gpr_visible_z_max_m": float(self.gpr_visible_z_max_m),
+            "sar3d_frame_radius": int(self.sar3d_frame_radius),
+            "sar3d_edge_frame_weight": float(self.sar3d_edge_frame_weight),
+            "sar3d_pitch_deg": float(self.sar3d_pitch_deg),
+            "sar3d_center_depth_m": float(self.sar3d_center_depth_m),
+            "sar3d_half_width_x_m": float(self.sar3d_half_width_x_m),
+            "sar3d_half_width_y_m": float(self.sar3d_half_width_y_m),
+            "sar3d_half_width_z_m": float(self.sar3d_half_width_z_m),
+            "sar3d_grid_nx": int(self.sar3d_grid_nx),
+            "sar3d_grid_ny": int(self.sar3d_grid_ny),
+            "sar3d_grid_nz": int(self.sar3d_grid_nz),
+            "sar3d_use_compensation": bool(self.sar3d_use_compensation),
+            "sar3d_normalize_pairs": bool(self.sar3d_normalize_pairs),
+            "sar3d_bottom_view_deg": float(self.sar3d_bottom_view_deg),
             "ignore_socket_speed": bool(self.ignore_socket_speed),
             "reprocess_current_result": bool(self.reprocess_current_result),
             "history_command_seq": int(self.history_command_seq),

python_app/orchestration/process_supervisor.py

@@ -283,6 +283,19 @@ class ProcessSupervisor:
 
     def _acquisition_command(self, config_path: Path) -> list[str]:
         """Return acquisition producer command selected by radar.model."""
+        # A configured replay dataset overrides the model's normal producer:
+        # recorded snapshots are replayed into the raw ring instead of
+        # acquiring from (mock) hardware.
+        replay_dir = self._read_replay_dataset_dir(config_path)
+        if replay_dir:
+            logger.debug("Selecting replay producer for dataset_dir=%s", replay_dir)
+            return [
+                sys.executable,
+                "-m",
+                "python_app.scripts.replay_raw_producer",
+                "--config",
+                str(config_path),
+            ]
         radar_model = self._read_radar_model(config_path)
         logger.debug("Selecting acquisition producer for radar.model=%s", radar_model)
         if radar_model in {"librevna_multi", "sn9000"}:
@@ -318,6 +331,20 @@ class ProcessSupervisor:
             return "librevna"
         return str(radar_payload.get("model", "librevna"))
 
+    @staticmethod
+    def _read_replay_dataset_dir(config_path: Path) -> str:
+        """Read `radar.replay.dataset_dir` cheaply, '' when replay is off."""
+        payload = json.loads(config_path.read_text(encoding="utf-8"))
+        if not isinstance(payload, dict):
+            return ""
+        radar_payload = payload.get("radar")
+        if not isinstance(radar_payload, dict):
+            return ""
+        replay_payload = radar_payload.get("replay")
+        if not isinstance(replay_payload, dict):
+            return ""
+        return str(replay_payload.get("dataset_dir", "")).strip()
+
     def _stop_processes(self, names: Iterable[str]) -> None:
         """Gracefully terminate processes, then force-kill on per-process timeout."""
         ordered_names = list(names)
@@ -591,6 +618,7 @@ class ProcessSupervisor:
             "build/bin/sweep_orchestrator",
             "python_app.scripts.matrix_raw_producer",
             "python_app.scripts.kamil_adc_raw_producer",
+            "python_app.scripts.replay_raw_producer",
         )
         try:
             raw = (Path("/proc") / str(pid) / "cmdline").read_bytes()

python_app/scripts/create_neutral_preprocess_sets.py

@@ -0,0 +1,90 @@
+"""CLI to create neutral S21 preprocess sets for a profile (headless).
+
+Mirrors the GUI's "Create neutral sets" button: builds calibration = 1+0j and
+reference = 0 collections for every combo of the profile and saves them under
+the profile's radar key, so the data_preprocessor becomes a bit-exact
+pass-through. Needed for replay profiles whose recorded data is already
+preprocessed (see run_config_examples/run_config_replay_vna.example.json).
+
+Usage:
+    python -m python_app.scripts.create_neutral_preprocess_sets \
+        --config run_config_examples/run_config_replay_vna.example.json \
+        [--set-name neutral_replay] [--force]
+"""
+
+from __future__ import annotations
+
+import argparse
+import logging
+from pathlib import Path
+
+from python_app.models.run_config_model import RunConfigModel
+from python_app.storage.npz_store import NpzStore, radar_key_from_config
+from python_app.workflows.kamil_adc_neutral_preprocess import (
+    build_neutral_s21_sets,
+    supports_neutral_preprocess_sets,
+)
+
+logger = logging.getLogger(__name__)
+
+PROJECT_ROOT = Path(__file__).resolve().parents[2]
+
+
+def main() -> int:
+    """Create and save the neutral S21 calibration/reference pair."""
+    parser = argparse.ArgumentParser(description="Create neutral S21 preprocess sets")
+    parser.add_argument("--config", required=True, type=Path, help="Path to a run_config profile")
+    parser.add_argument("--set-name", default="neutral_replay", help="Name for both saved sets")
+    parser.add_argument(
+        "--force", action="store_true", help="Overwrite the sets if they already exist"
+    )
+    args = parser.parse_args()
+
+    logging.basicConfig(level=logging.INFO, format="%(levelname)s %(name)s: %(message)s")
+
+    config = RunConfigModel.load_from_path(args.config)
+    config.ensure_combos()
+    if not supports_neutral_preprocess_sets(config):
+        raise RuntimeError(
+            "Neutral S21 sets are available only for kamil_adc and librevna_multi"
+        )
+
+    radar_key = radar_key_from_config(
+        model_name=config.radar.model,
+        serial=config.radar.serial,
+        sweep_start_hz=config.radar.sweep.start_hz,
+        sweep_stop_hz=config.radar.sweep.stop_hz,
+        sweep_points=config.radar.sweep.points,
+        ifbw_hz=config.radar.sweep.if_bandwidth_hz,
+        power_dbm=config.radar.sweep.power_dbm,
+        extra_serials=config.radar_key_extra_parts() or None,
+    )
+    store = NpzStore(PROJECT_ROOT / "python_app/data")
+
+    existing = [
+        kind
+        for kind in ("s21_calibration", "s21_reference")
+        if args.set_name in store.list_sets(kind, radar_key)
+    ]
+    if existing and not args.force:
+        logger.info(
+            "Sets '%s' already exist for %s (radar_key=%s); nothing to do (use --force to overwrite).",
+            args.set_name, ", ".join(existing), radar_key,
+        )
+        return 0
+
+    calibration, reference = build_neutral_s21_sets(config)
+    store.save_set("s21_calibration", radar_key, args.set_name, calibration)
+    store.save_set("s21_reference", radar_key, args.set_name, reference)
+    logger.info(
+        "Neutral S21 sets saved: set=%s, radar_key=%s, combos=%d, points=%d",
+        args.set_name,
+        radar_key,
+        len(calibration.traces),
+        int(calibration.traces[0].frequency_hz.size),
+    )
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

python_app/scripts/replay_raw_producer.py

@@ -0,0 +1,249 @@
+"""Raw producer that replays a recorded snapshot dataset (no hardware).
+
+Selected by the supervisor instead of the model's normal acquisition producer
+when ``radar.replay.dataset_dir`` is non-empty. The dataset is a directory of
+snapshot folders named ``NNNN_idK_ns<timestamp>`` (the VNA_AlGo
+preprocessed-capture layout), each holding ``i{input}_o{output}_s21.npy`` and
+``i{input}_o{output}_freq.npy`` files for every run combo.
+
+The recorded data is already background-subtracted, so the profile must pair
+this producer with NEUTRAL preprocess sets (calibration = 1, reference = 0,
+no notch) — then the C++ preprocessor passes the traces through bit-exactly
+and the processors (e.g. sar3d) see the recorded scene.
+
+Timeline: snapshot timestamps are re-anchored to this process's monotonic
+clock, preserving the recorded inter-frame gaps. On loop, the seam advances by
+the median inter-frame gap, so motion-model consumers (sar3d aperture offsets)
+never see time jump backwards. ``collection_id`` increases forever.
+"""
+
+from __future__ import annotations
+
+import argparse
+import logging
+from pathlib import Path
+import re
+import signal
+import statistics
+import threading
+import time
+
+import numpy as np
+
+from python_app.logging_setup import coerce_level
+from python_app.models.dataset_model import ComboKey, SweepCollection, TraceData
+from python_app.models.run_config_model import RunConfigModel
+from python_app.orchestration.shm import ShmRingWriter
+from python_app.storage.npz.serialize import RAW_MAGIC, serialize_trace_collection
+
+logger = logging.getLogger(__name__)
+
+_SNAPSHOT_NAME = re.compile(r"^\d+_id\d+_ns(\d+)$")
+
+
+class ReplayDataset:
+    """Recorded snapshots preloaded into memory, with their recorded timeline."""
+
+    def __init__(self, snapshots: list[list[TraceData]], timestamps_ns: list[int]) -> None:
+        self.snapshots = snapshots
+        self.timestamps_ns = timestamps_ns
+        gaps = [b - a for a, b in zip(timestamps_ns, timestamps_ns[1:])]
+        # Median gap doubles as the loop-seam step and the capture-window span;
+        # a single-snapshot dataset falls back to an arbitrary sane cadence.
+        self.median_gap_ns = int(statistics.median(gaps)) if gaps else 160_000_000
+
+    def __len__(self) -> int:
+        return len(self.snapshots)
+
+
+def _load_dataset(dataset_dir: Path, combos: list[ComboKey], expected_points: int) -> ReplayDataset:
+    """Preload every snapshot's traces; fail loudly on a malformed dataset."""
+    if not dataset_dir.is_dir():
+        raise RuntimeError(f"radar.replay.dataset_dir is not a directory: {dataset_dir}")
+
+    named: list[tuple[str, int]] = []
+    for entry in sorted(dataset_dir.iterdir()):
+        match = _SNAPSHOT_NAME.match(entry.name)
+        if entry.is_dir() and match:
+            named.append((entry.name, int(match.group(1))))
+    if not named:
+        raise RuntimeError(
+            f"No snapshot folders (NNNN_idK_ns<ts>) found in {dataset_dir}"
+        )
+
+    snapshots: list[list[TraceData]] = []
+    timestamps_ns: list[int] = []
+    for name, timestamp_ns in named:
+        folder = dataset_dir / name
+        traces: list[TraceData] = []
+        for combo in combos:
+            s21_path = folder / f"i{combo.input}_o{combo.output}_s21.npy"
+            freq_path = folder / f"i{combo.input}_o{combo.output}_freq.npy"
+            if not s21_path.is_file() or not freq_path.is_file():
+                raise RuntimeError(
+                    f"Snapshot {name} is missing files for combo "
+                    f"input={combo.input} output={combo.output}: expected "
+                    f"{s21_path.name} and {freq_path.name}. Check run.combos vs the dataset."
+                )
+            s21 = np.load(s21_path).astype(np.complex64, copy=False)
+            freq = np.load(freq_path).astype(np.float32, copy=False)
+            if s21.size != freq.size:
+                raise RuntimeError(f"{s21_path}: s21 size {s21.size} != freq size {freq.size}")
+            if expected_points and s21.size != expected_points:
+                raise RuntimeError(
+                    f"{s21_path}: {s21.size} points, but radar.sweep.points = "
+                    f"{expected_points}. Align the profile's sweep with the dataset."
+                )
+            traces.append(
+                TraceData(
+                    combo=combo,
+                    frequency_hz=freq,
+                    s11=np.zeros_like(s21),
+                    s21=s21,
+                )
+            )
+        snapshots.append(traces)
+        timestamps_ns.append(timestamp_ns)
+
+    return ReplayDataset(snapshots, timestamps_ns)
+
+
+def _build_collection(
+    dataset: ReplayDataset,
+    index: int,
+    collection_id: int,
+    frame_ns: int,
+) -> SweepCollection:
+    """Wrap one preloaded snapshot in a SweepCollection at the given frame time."""
+    span_ns = dataset.median_gap_ns
+    start_ns = frame_ns - span_ns // 2
+    end_ns = frame_ns + span_ns // 2
+    traces = []
+    for source in dataset.snapshots[index]:
+        traces.append(
+            TraceData(
+                combo=source.combo,
+                frequency_hz=source.frequency_hz,
+                s11=source.s11,
+                s21=source.s21,
+                capture_start_ns=start_ns,
+                capture_end_ns=end_ns,
+            )
+        )
+    return SweepCollection(
+        collection_id=collection_id,
+        monotonic_ns=frame_ns,
+        traces=traces,
+        capture_start_ns=start_ns,
+        capture_end_ns=end_ns,
+    )
+
+
+def main() -> int:
+    """Replay the configured dataset into the raw SHM rings until stopped."""
+    parser = argparse.ArgumentParser(description="Replay a recorded dataset to raw SHM rings")
+    parser.add_argument("--config", required=True, type=Path, help="Path to run_config.json")
+    args = parser.parse_args()
+
+    logging.basicConfig(level=logging.INFO, format="%(levelname)s %(name)s: %(message)s")
+    stop_requested = threading.Event()
+
+    def request_stop(_signum: int, _frame: object) -> None:
+        stop_requested.set()
+
+    signal.signal(signal.SIGINT, request_stop)
+    signal.signal(signal.SIGTERM, request_stop)
+
+    config = RunConfigModel.load_from_path(args.config)
+    config.apply_device_model_constraints()
+    logging.getLogger("python_app").setLevel(coerce_level(config.logging.level))
+
+    replay = config.radar.replay
+    if not replay.dataset_dir:
+        raise RuntimeError("replay_raw_producer requires a non-empty radar.replay.dataset_dir")
+    rate_scale = float(replay.rate_scale) if replay.rate_scale > 0 else 1.0
+    combos = [ComboKey(input=combo.input, output=combo.output) for combo in config.combos]
+
+    # expanduser: profiles are copied between machines, so "~/..." dataset paths
+    # must resolve against the local home instead of failing as a literal "~".
+    dataset = _load_dataset(
+        Path(replay.dataset_dir).expanduser(), combos, int(config.radar.sweep.points)
+    )
+    logger.info(
+        "Replay producer starting: dataset=%s, snapshots=%d, combos=%d, "
+        "median_gap=%.3f s, loop=%s, rate_scale=%.2f",
+        replay.dataset_dir, len(dataset), len(combos),
+        dataset.median_gap_ns / 1e9, replay.loop, rate_scale,
+    )
+
+    raw_writer = ShmRingWriter(
+        config.rings.raw.name,
+        config.rings.raw.capacity,
+        config.rings.raw.slot_size_bytes,
+    )
+    raw_tap_writer = ShmRingWriter(
+        config.rings.raw_tap.name,
+        config.rings.raw_tap.capacity,
+        config.rings.raw_tap.slot_size_bytes,
+    )
+
+    try:
+        base_ns = time.monotonic_ns()
+        wall_start_s = time.monotonic()
+        first_ts = dataset.timestamps_ns[0]
+        cycle_span_ns = dataset.timestamps_ns[-1] - first_ts + dataset.median_gap_ns
+
+        collection_id = 1
+        publish_failures = 0
+        cycle = 0
+        while not stop_requested.is_set():
+            for index in range(len(dataset)):
+                if stop_requested.is_set():
+                    break
+                # Monotonic replay timeline: recorded offsets re-anchored at
+                # base_ns, each loop shifted by one full cycle span.
+                frame_ns = (
+                    base_ns
+                    + (dataset.timestamps_ns[index] - first_ts)
+                    + cycle * cycle_span_ns
+                )
+                # Publish on the recorded cadence (scaled); waiting via the stop
+                # event keeps shutdown prompt.
+                deadline_s = wall_start_s + (frame_ns - base_ns) / 1e9 / rate_scale
+                delay_s = deadline_s - time.monotonic()
+                if delay_s > 0 and stop_requested.wait(delay_s):
+                    break
+
+                collection = _build_collection(dataset, index, collection_id, frame_ns)
+                payload = serialize_trace_collection(collection, RAW_MAGIC)
+                if raw_writer.push(payload):
+                    raw_tap_writer.push(payload)  # best-effort GUI tap; never fatal
+                else:
+                    publish_failures += 1
+                    if publish_failures == 1 or publish_failures % 100 == 0:
+                        logger.error(
+                            "Replay payload %d B exceeds ring slot %d B; dropping collection %d (drops=%d)",
+                            len(payload), raw_writer.slot_size_bytes, collection_id, publish_failures,
+                        )
+                if collection_id == 1 or collection_id % 50 == 0:
+                    logger.info(
+                        "replay collection %d published (snapshot %d/%d, cycle %d)",
+                        collection_id, index + 1, len(dataset), cycle,
+                    )
+                collection_id += 1
+                if not config.runtime.continuous:
+                    stop_requested.set()
+                    break
+            if not replay.loop:
+                break
+            cycle += 1
+    finally:
+        raw_tap_writer.close()
+        raw_writer.close()
+
+    logger.info("replay raw producer stopped")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

python_app/scripts/replay_smoke_run.py

@@ -0,0 +1,248 @@
+"""Headless smoke: replay the recorded VNA_AlGo dataset through the full pipeline.
+
+Loads the replay profile, creates neutral preprocess sets, launches the whole
+pipeline (replay_raw_producer -> data_preprocessor -> data_processor) via the
+ProcessSupervisor with unique ring names, forces the sar3d processor through
+processing_live.json, and asserts that:
+
+- sar3d projection payloads arrive with a positive peak;
+- collection ids increase monotonically (including across a dataset loop);
+- the preprocessor is transparent: a preprocessed_tap trace equals the
+  recorded snapshot's s21 within float32 rounding.
+
+Usage:
+    python -m python_app.scripts.replay_smoke_run \
+        [--config run_config_examples/run_config_replay_vna.example.json] \
+        [--duration 20.0]
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+import logging
+import os
+from pathlib import Path
+import time
+
+import numpy as np
+
+from python_app.models.run_config_model import RunConfigModel
+from python_app.orchestration.config_writer import ConfigWriter
+from python_app.orchestration.live_processing_config import (
+    ProcessingLiveConfig,
+    ProcessingLiveConfigWriter,
+)
+from python_app.orchestration.process_supervisor import ProcessSupervisor
+from python_app.orchestration.shm import ShmRingReader
+from python_app.orchestration.shm.decoder import PREPROC_MAGIC, decode_trace_collection
+from python_app.storage.npz_store import NpzStore, radar_key_from_config
+from python_app.workflows.kamil_adc_neutral_preprocess import build_neutral_s21_sets
+
+logger = logging.getLogger(__name__)
+
+PROJECT_ROOT = Path(__file__).resolve().parents[2]
+
+
+def _make_unique_ring_name(prefix: str) -> str:
+    """Build unique POSIX SHM ring name."""
+    return f"/{prefix}_{os.getpid()}_{time.monotonic_ns()}"
+
+
+def _sar3d_live_config(profile_path: Path) -> ProcessingLiveConfig:
+    """Build a sar3d live config from the profile's gui.processing.sar3d section.
+
+    RunConfigModel does not carry the GUI section (it belongs to the GUI profile
+    layer), so the sar3d defaults are read straight from the profile JSON.
+    """
+    payload = json.loads(profile_path.read_text(encoding="utf-8"))
+    sar3d = (
+        payload.get("gui", {}).get("processing", {}).get("sar3d", {})
+        if isinstance(payload, dict)
+        else {}
+    )
+
+    def _positions(text: str) -> list[int]:
+        return [int(value) for value in str(text).split(",") if value.strip()]
+
+    live = ProcessingLiveConfig(processor_mode="sar3d", ignore_socket_speed=True)
+    if not isinstance(sar3d, dict) or not sar3d:
+        return live
+    live.gpr_input_positions = _positions(sar3d.get("input_positions", ""))
+    live.gpr_output_positions = _positions(sar3d.get("output_positions", ""))
+    live.gpr_min_depth_m = float(sar3d.get("min_depth_m", live.gpr_min_depth_m))
+    live.gpr_max_depth_m = float(sar3d.get("max_depth_m", live.gpr_max_depth_m))
+    live.gpr_start_freq_mhz = float(sar3d.get("start_freq_mhz", live.gpr_start_freq_mhz))
+    live.gpr_stop_freq_mhz = float(sar3d.get("stop_freq_mhz", live.gpr_stop_freq_mhz))
+    live.gpr_speed_m_s = float(sar3d.get("speed_m_s", live.gpr_speed_m_s))
+    live.gpr_look_angle_deg = float(sar3d.get("look_angle_deg", live.gpr_look_angle_deg))
+    live.gpr_direction_sign = float(sar3d.get("direction_sign", live.gpr_direction_sign))
+    live.gpr_apply_freq_phase_correction = bool(
+        sar3d.get("apply_freq_phase_correction", live.gpr_apply_freq_phase_correction)
+    )
+    live.sar3d_frame_radius = int(sar3d.get("frame_radius", live.sar3d_frame_radius))
+    live.sar3d_edge_frame_weight = float(
+        sar3d.get("edge_frame_weight", live.sar3d_edge_frame_weight)
+    )
+    live.sar3d_pitch_deg = float(sar3d.get("pitch_deg", live.sar3d_pitch_deg))
+    live.sar3d_center_depth_m = float(sar3d.get("center_depth_m", live.sar3d_center_depth_m))
+    live.sar3d_half_width_x_m = float(sar3d.get("half_width_x_m", live.sar3d_half_width_x_m))
+    live.sar3d_half_width_y_m = float(sar3d.get("half_width_y_m", live.sar3d_half_width_y_m))
+    live.sar3d_half_width_z_m = float(sar3d.get("half_width_z_m", live.sar3d_half_width_z_m))
+    live.sar3d_grid_nx = int(sar3d.get("grid_nx", live.sar3d_grid_nx))
+    live.sar3d_grid_ny = int(sar3d.get("grid_ny", live.sar3d_grid_ny))
+    live.sar3d_grid_nz = int(sar3d.get("grid_nz", live.sar3d_grid_nz))
+    live.sar3d_use_compensation = bool(sar3d.get("use_compensation", live.sar3d_use_compensation))
+    live.sar3d_normalize_pairs = bool(sar3d.get("normalize_pairs", live.sar3d_normalize_pairs))
+    live.sar3d_bottom_view_deg = float(sar3d.get("bottom_view_deg", live.sar3d_bottom_view_deg))
+    return live
+
+
+def _check_preprocessor_transparency(
+    tap_reader: ShmRingReader, dataset_dir: Path
+) -> bool:
+    """Compare one preprocessed trace against the recorded snapshot's s21."""
+    payload = tap_reader.peek_latest_payload()
+    if payload is None:
+        logger.warning("No preprocessed_tap payload available for transparency check")
+        return False
+    preprocessed = decode_trace_collection(payload, PREPROC_MAGIC)
+
+    # Any snapshot works: neutral preprocessing is time-invariant, so match the
+    # trace against every recorded snapshot and accept the best one.
+    trace = preprocessed.traces[0]
+    combo_file = f"i{trace.combo.input}_o{trace.combo.output}_s21.npy"
+    best_delta = np.inf
+    for folder in sorted(p for p in dataset_dir.iterdir() if p.name[:1].isdigit()):
+        recorded = np.load(folder / combo_file).astype(np.complex64)
+        if recorded.size != np.asarray(trace.s21).size:
+            continue
+        delta = float(np.max(np.abs(np.asarray(trace.s21) - recorded)))
+        best_delta = min(best_delta, delta)
+        if best_delta == 0.0:
+            break
+    logger.info("Preprocessor transparency: max|delta| vs best snapshot = %.3e", best_delta)
+    return best_delta < 1e-6
+
+
+def main() -> int:
+    """Run the replay pipeline smoke and report pass/fail."""
+    parser = argparse.ArgumentParser(description="Headless replay pipeline smoke")
+    parser.add_argument(
+        "--config",
+        type=Path,
+        default=PROJECT_ROOT / "run_config_examples/run_config_replay_vna.example.json",
+    )
+    parser.add_argument("--duration", type=float, default=20.0)
+    args = parser.parse_args()
+
+    logging.basicConfig(level=logging.INFO, format="%(levelname)s %(name)s: %(message)s")
+
+    config = RunConfigModel.load_from_path(args.config)
+    config.ensure_combos()
+    # expanduser: "~/..." dataset paths from machine-portable profiles.
+    dataset_dir = Path(config.radar.replay.dataset_dir).expanduser()
+    if not dataset_dir.is_dir():
+        raise RuntimeError(f"Replay dataset not found: {dataset_dir}")
+
+    config.rings.raw.name = _make_unique_ring_name("radar_raw_replay_smoke")
+    config.rings.raw_tap.name = _make_unique_ring_name("radar_raw_tap_replay_smoke")
+    config.rings.preprocessed.name = _make_unique_ring_name("radar_preprocessed_replay_smoke")
+    config.rings.preprocessed_tap.name = _make_unique_ring_name("radar_preprocessed_tap_replay_smoke")
+    config.rings.results.name = _make_unique_ring_name("radar_results_replay_smoke")
+    # Private live-config path so the smoke never clobbers the GUI's file.
+    live_path = PROJECT_ROOT / "python_app/runtime/processing_live_replay_smoke.json"
+    config.runtime.processing_live_config_path = str(
+        live_path.relative_to(PROJECT_ROOT)
+    )
+
+    radar_key = radar_key_from_config(
+        model_name=config.radar.model,
+        serial=config.radar.serial,
+        sweep_start_hz=config.radar.sweep.start_hz,
+        sweep_stop_hz=config.radar.sweep.stop_hz,
+        sweep_points=config.radar.sweep.points,
+        ifbw_hz=config.radar.sweep.if_bandwidth_hz,
+        power_dbm=config.radar.sweep.power_dbm,
+        extra_serials=config.radar_key_extra_parts() or None,
+    )
+    store = NpzStore(PROJECT_ROOT / "python_app/data")
+    set_name = config.preprocess.s21.calibration.set_name or "neutral_replay"
+    if set_name not in store.list_sets("s21_calibration", radar_key):
+        calibration, reference = build_neutral_s21_sets(config)
+        store.save_set("s21_calibration", radar_key, set_name, calibration)
+        store.save_set("s21_reference", radar_key, set_name, reference)
+        logger.info("Created neutral sets '%s' (radar_key=%s)", set_name, radar_key)
+    config.preprocess.s21.calibration.set_name = set_name
+    config.preprocess.s21.reference.set_name = set_name
+
+    ProcessingLiveConfigWriter(live_path).write(_sar3d_live_config(args.config))
+
+    config_writer = ConfigWriter(PROJECT_ROOT / "python_app/runtime")
+    config_writer.prepare_preprocess_bundles(store, radar_key, config)
+    config_path = config_writer.write(
+        config, PROJECT_ROOT / "python_app/runtime/run_config_replay_smoke.json"
+    )
+
+    supervisor = ProcessSupervisor(PROJECT_ROOT)
+    result_reader: ShmRingReader | None = None
+    tap_reader: ShmRingReader | None = None
+    failures: list[str] = []
+    try:
+        supervisor.start(config_path)
+        result_reader = ShmRingReader(config.rings.results.name, open_timeout_s=15.0)
+        tap_reader = ShmRingReader(config.rings.preprocessed_tap.name, open_timeout_s=15.0)
+
+        deadline = time.monotonic() + args.duration
+        received = 0
+        sar3d_results = 0
+        last_peak = 0.0
+        last_id = 0
+        ids_monotonic = True
+        while time.monotonic() < deadline:
+            result = result_reader.pop_result_collection()
+            if result is None:
+                time.sleep(0.02)
+                continue
+            received += 1
+            if last_id and result.collection_id <= last_id:
+                ids_monotonic = False
+            last_id = result.collection_id
+            names = {p.processing_name: p for p in result.collection_payloads}
+            if "sar3d_proj_xz" in names and "sar3d_proj_bottom" in names:
+                sar3d_results += 1
+                peak_payload = names.get("sar3d_peak")
+                if peak_payload is not None:
+                    last_peak = float(peak_payload.scalar_value)
+
+        logger.info(
+            "Results: total=%d, with sar3d projections=%d, last collection_id=%d, last peak=%.4e",
+            received, sar3d_results, last_id, last_peak,
+        )
+        if received == 0:
+            failures.append("no result collections received")
+        if sar3d_results == 0:
+            failures.append("no sar3d projection payloads received")
+        if not (last_peak > 0.0):
+            failures.append(f"sar3d peak is not positive: {last_peak}")
+        if not ids_monotonic:
+            failures.append("collection ids were not monotonically increasing")
+        if not _check_preprocessor_transparency(tap_reader, dataset_dir):
+            failures.append("preprocessor is not transparent vs recorded snapshots")
+    finally:
+        supervisor.stop_all()
+        if result_reader is not None:
+            result_reader.close()
+        if tap_reader is not None:
+            tap_reader.close()
+
+    if failures:
+        for failure in failures:
+            logger.error("FAIL: %s", failure)
+        return 1
+    logger.info("REPLAY SMOKE OK")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())

requirements.txt

@@ -3,8 +3,17 @@ libusb1>=3.1
 pyserial>=3.5,<4
 pyvisa>=1.14
 PyVISA-py>=0.7  # pure-Python VISA backend ("@py"), required by the SN9000 (and any visa_library="@py") path
-PyQt6>=6.6
+# PyQt6 6.8+ ships only manylinux_2_39 aarch64 wheels (glibc >= 2.39); Jetson /
+# Ubuntu 22.04 has glibc 2.35, so pip would fall back to a source build that
+# dies on qmake. 6.7.1 is the last release with a manylinux_2_28 aarch64 wheel.
+PyQt6>=6.6,<6.8; platform_machine == "aarch64"
+PyQt6>=6.6; platform_machine != "aarch64"
 pyqtgraph>=0.13.7
 rpi-hardware-pwm>=0.2.2,<1
 fastapi>=0.110           # embedded web UI (browser view + controls), served in-process
 uvicorn[standard]>=0.29  # ASGI server for the web UI, run on a background thread
+# CUDA headers + libnvrtc for building the vendored gpr_cpp library (3D SAR).
+# x86_64 dev machines only: on Jetson (aarch64) CUDA comes with JetPack at
+# /usr/local/cuda and build_cmake.sh falls back to it automatically.
+nvidia-cuda-nvrtc==13.3.33; platform_machine == "x86_64"   # nvrtc.h + libnvrtc
+nvidia-cuda-runtime==13.3.29; platform_machine == "x86_64" # cuda.h (Driver API)

run_config_examples/run_config_replay_vna.example.json

@@ -0,0 +1,205 @@
+{
+  "_comment": "Replay of the recorded VNA_AlGo tank005 dataset through the full pipeline (no hardware). The replay producer feeds the raw ring; pair it with NEUTRAL preprocess sets (calibration=1, reference=0) so the already-preprocessed recording passes through unchanged. Create them via GUI 'Create neutral sets' (name: neutral_replay) or python -m python_app.scripts.create_neutral_preprocess_sets --config <this file>.",
+  "radar": {
+    "model": "librevna_multi",
+    "serial": "206830985532",
+    "remote_host": "127.0.0.1",
+    "remote_port": 50209,
+    "driver_mode": "mock",
+    "mock_signal_hz": 5000000.0,
+    "visa_library": "",
+    "multi_device": {
+      "slave_serials": [
+        "206930965532",
+        "206930A15532"
+      ],
+      "force_external_reference": true,
+      "recovery_attempts": 3,
+      "output_switch_positions": 1,
+      "input_switch_positions": 1,
+      "stimulus_switch_positions": 1
+    },
+    "replay": {
+      "dataset_dir": "/home/awe/Documents/VNA_AlGo/librevna_multi_2-6GHz_751pts_-3dBm_tank005/librevna_multi_2-6GHz_751pts_-3dBm_tank005/preprocessed",
+      "loop": true,
+      "rate_scale": 1.0
+    },
+    "sweep": {
+      "start_hz": 2000000000.0,
+      "stop_hz": 6000000000.0,
+      "if_bandwidth_hz": 50000.0,
+      "stimulus_power_dbm": -3.0,
+      "points": 751
+    }
+  },
+  "switches": {
+    "port1": {
+      "name": "port1",
+      "driver_mode": "mock",
+      "driver": "h7992",
+      "radar_port": 1,
+      "positions": 2,
+      "default_position": 0,
+      "gpio_chip": "/dev/gpiochip0",
+      "pin_a": 17,
+      "pin_b": 27,
+      "invert_logic": false
+    },
+    "port2": {
+      "name": "port2",
+      "driver_mode": "mock",
+      "driver": "h7992",
+      "radar_port": 2,
+      "positions": 4,
+      "default_position": 0,
+      "gpio_chip": "/dev/gpiochip0",
+      "pin_a": 22,
+      "pin_b": 23,
+      "invert_logic": false
+    }
+  },
+  "control_button": {
+    "enabled": false,
+    "gpio_chip": "/dev/gpiochip0",
+    "pin": -1,
+    "active_low": true,
+    "bias": "",
+    "debounce_ms": 50,
+    "action": "capture_tmp_reference"
+  },
+  "run": {
+    "settling_ms": 0,
+    "idle_sleep_ms": 2,
+    "continuous": true,
+    "processing_live_config_path": "python_app/runtime/processing_live.json",
+    "locator_server": {
+      "device_id": 3,
+      "protocol_version": 1,
+      "host": "0.0.0.0",
+      "port": 8888,
+      "max_payload_bytes": 65536,
+      "client_queue_size": 32,
+      "logger_name": "locator_runtime"
+    },
+    "combos": [
+      { "input": 0, "output": 0 },
+      { "input": 1, "output": 0 },
+      { "input": 2, "output": 0 },
+      { "input": 3, "output": 0 },
+      { "input": 0, "output": 1 },
+      { "input": 1, "output": 1 },
+      { "input": 2, "output": 1 },
+      { "input": 3, "output": 1 }
+    ]
+  },
+  "preprocess": {
+    "s21": {
+      "calibration": {
+        "set_name": "neutral_replay",
+        "bundle_path": ""
+      },
+      "reference": {
+        "set_name": "neutral_replay",
+        "bundle_path": ""
+      }
+    },
+    "s11": {
+      "calibration": {
+        "open": { "set_name": "", "bundle_path": "" },
+        "short": { "set_name": "", "bundle_path": "" },
+        "load": { "set_name": "", "bundle_path": "" }
+      },
+      "reference": { "set_name": "", "bundle_path": "" }
+    },
+    "notch": {
+      "enabled": false,
+      "bands_hz": [],
+      "taper_width_hz": 40000000.0,
+      "taper_type": "cosine"
+    }
+  },
+  "gpr": {
+    "relative_permittivity": 1.0,
+    "tx_geometry": [
+      { "output_pos": 0, "x_m": 0.785, "y_m": 0.0, "z_m": 0.0 },
+      { "output_pos": 1, "x_m": -0.785, "y_m": 0.0, "z_m": 0.0 }
+    ],
+    "rx_geometry": [
+      { "input_pos": 0, "x_m": -0.42, "y_m": 0.0, "z_m": 0.0 },
+      { "input_pos": 1, "x_m": 0.48, "y_m": 0.0, "z_m": 0.0 },
+      { "input_pos": 2, "x_m": 0.18, "y_m": 0.0, "z_m": 0.0 },
+      { "input_pos": 3, "x_m": -0.12, "y_m": 0.0, "z_m": 0.0 }
+    ]
+  },
+  "rings": {
+    "raw": {
+      "name": "/radar_raw_simulator",
+      "capacity": 50,
+      "slot_size_bytes": 2097152
+    },
+    "raw_tap": {
+      "name": "/radar_raw_tap_simulator",
+      "capacity": 50,
+      "slot_size_bytes": 2097152
+    },
+    "preprocessed": {
+      "name": "/radar_preprocessed_simulator",
+      "capacity": 50,
+      "slot_size_bytes": 2097152
+    },
+    "preprocessed_tap": {
+      "name": "/radar_preprocessed_tap_simulator",
+      "capacity": 50,
+      "slot_size_bytes": 2097152
+    },
+    "results": {
+      "name": "/radar_results_simulator",
+      "capacity": 50,
+      "slot_size_bytes": 2097152
+    }
+  },
+  "gui": {
+    "version": 1,
+    "switches": {
+      "combo_mode": "text",
+      "combos_text": "0:0,1:0,2:0,3:0,0:1,1:1,2:1,3:1",
+      "single_input": "0",
+      "single_output": "0"
+    },
+    "processing": {
+      "selected_mode": "sar3d",
+      "sar3d": {
+        "input_positions": "0,1,2,3",
+        "output_positions": "0,1",
+        "min_depth_m": 3.2,
+        "max_depth_m": 10.0,
+        "start_freq_mhz": 2000.0,
+        "stop_freq_mhz": 6000.0,
+        "speed_m_s": 0.854,
+        "ignore_socket_speed_enabled": true,
+        "look_angle_deg": 6.0,
+        "direction_sign": 1.0,
+        "apply_freq_phase_correction": true,
+        "frame_radius": 6,
+        "edge_frame_weight": 0.75,
+        "pitch_deg": 6.0,
+        "center_depth_m": 5.6,
+        "half_width_x_m": 1.8,
+        "half_width_y_m": 1.0,
+        "half_width_z_m": 1.8,
+        "grid_nx": 150,
+        "grid_ny": 120,
+        "grid_nz": 160,
+        "use_compensation": true,
+        "normalize_pairs": true,
+        "bottom_view_deg": 0.0
+      }
+    },
+    "preprocess_dialog": {
+      "set_name": "neutral_replay",
+      "radar_config_dir": "",
+      "use_all_radar_configs": false,
+      "median_sweep_count": 5
+    }
+  }
+}

start.sh

@@ -225,14 +225,29 @@ run_privileged() {
 }
 
 ensure_system_dependencies() {
-    if command -v pkg-config >/dev/null 2>&1 && pkg-config --exists libusb-1.0; then
+    # Qt >= 6.5 xcb platform plugin needs libxcb-cursor0; aarch64/Jetson images
+    # (and some minimal desktops) ship without it, and the GUI then dies at
+    # startup with "could not load the Qt platform plugin xcb". apt-only: other
+    # distros/managers name these differently and usually have them preinstalled.
+    local qt_gui_packages_missing=0
+    if command -v apt-get >/dev/null 2>&1 && command -v dpkg-query >/dev/null 2>&1; then
+        local package
+        for package in libxcb-cursor0 libxkbcommon-x11-0; do
+            if ! dpkg-query -W -f='${Status}' "${package}" 2>/dev/null | grep -q "install ok installed"; then
+                qt_gui_packages_missing=1
+            fi
+        done
+    fi
+
+    if command -v pkg-config >/dev/null 2>&1 && pkg-config --exists libusb-1.0 \
+        && ((qt_gui_packages_missing == 0)); then
         return
     fi
 
     echo "[start.sh] Installing system dependency: libusb-1.0 dev headers..."
     if command -v apt-get >/dev/null 2>&1; then
         run_privileged apt-get update
-        run_privileged apt-get install -y pkg-config libusb-1.0-0-dev
+        run_privileged apt-get install -y pkg-config libusb-1.0-0-dev libxcb-cursor0 libxkbcommon-x11-0
     elif command -v dnf >/dev/null 2>&1; then
         run_privileged dnf install -y pkgconf-pkg-config libusb1-devel
     elif command -v yum >/dev/null 2>&1; then