sar3d: port the Multi_antenna_3D_clean_v1 notebook algorithm

The old sar3d core was a port of the superseded Horns_3D notebook and
produced wrong volumes. The compute core now follows the reference
notebook exactly:

- segmented_interleaved frame timing rebuilt from per-trace capture
  stamps (uniform tx_per_segment split as the no-stamps fallback), with
  the coarse inter-segment dtau and the intra-sweep dtau applied as one
  phase ramp before the IFFT;
- ground scene coordinates with the ground_depth_m lift (scene Y=0 is
  the ground), full-region grid instead of a window around a center
  depth, geo_ref compensation anchored on the ground plane, power 0.3;
- notebook finalization in C++: |acc/cnt|, depth mask, normalization
  without the Y-edge planes, separable Gaussian over (Z, X), level;
- CUDA kernel computes geometry in double: float32 loses ~3e-4 of an
  A-scan sample in tau and drifts ~1e-4 from the float64 reference.

Live-config keys follow: the center_depth/half_width/grid_n window knobs
are replaced by ground/voxel/y-range/margin/sigma/power/tx_per_segment.

Parity against the notebook npz exports (frames 0034-0036 of the
20260922_mine_fast_true dataset): max|diff| <= 1.6e-6 off the depth-mask
knife edge, identical argmax, level rel diff <= 7e-7. Verified via
make sar3d_parity --dump-dir + scripts/sar3d_volume_compare.py; the
executable spec lives in scripts/sar3d_reference_check.py (bit-tight
against the npz, including the animation shifted-grid arithmetic).

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

data_acq_and_processing/gpr_cpp/src/ascan.cpp

@@ -9,19 +9,69 @@
 
 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;
+TxWindows uniform_tx_windows(const std::vector<int>& tx_order, double acq_s, int tx_per_segment) {
+    TxWindows windows;
+    if (tx_order.empty()) return windows;
+    int per_seg = std::max(1, tx_per_segment);
+    int n_seg = std::max(1, (int)((tx_order.size() + per_seg - 1) / per_seg));
+    double t_seg = acq_s / n_seg;
+    for (size_t k = 0; k < tx_order.size(); ++k) {
+        int seg = (int)k / per_seg;
+        windows[tx_order[k]] = {seg * t_seg, seg * t_seg + t_seg};
+    }
+    return windows;
+}
+
+PairTimestamps compute_pair_timestamps(const TxWindows& tx_windows_s, double acq_s,
+                                       const std::vector<int>& tx_order,
+                                       const GeometryConfig& geom, const MotionConfig& mot) {
+    // Аппаратный порядок: сперва tx_order (run.combos), затем прочие ключи окон.
+    std::vector<int> hw_order;
+    for (int tx : tx_order)
+        if (tx_windows_s.count(tx) && std::find(hw_order.begin(), hw_order.end(), tx) == hw_order.end())
+            hw_order.push_back(tx);
+    for (auto& [tx, _] : tx_windows_s)
+        if (std::find(hw_order.begin(), hw_order.end(), tx) == hw_order.end())
+            hw_order.push_back(tx);
+
+    // Tx с одинаковым окном — один сегмент; сегменты по возрастанию окна,
+    // слот — порядок появления в hw_order (порт группировки ноутбука).
+    std::vector<std::pair<std::pair<double, double>, std::vector<int>>> groups;
+    for (int tx : hw_order) {
+        auto win = tx_windows_s.at(tx);
+        auto found = std::find_if(groups.begin(), groups.end(),
+                                  [&](const auto& g) { return g.first == win; });
+        if (found == groups.end())
+            groups.push_back({win, {tx}});
+        else
+            found->second.push_back(tx);
+    }
+    std::sort(groups.begin(), groups.end(),
+              [](const auto& a, const auto& b) { return a.first < b.first; });
+
+    struct SegSlot { double t0, t1; int slot, n_slots; };
+    std::map<int, SegSlot> seg_of;
+    for (auto& [win, txs] : groups)
+        for (size_t slot = 0; slot < txs.size(); ++slot)
+            seg_of[txs[slot]] = {win.first, win.second, (int)slot, (int)txs.size()};
+
+    double cos_theta = std::cos(mot.look_angle_deg * M_PI / 180.0);
+    double t_ref = 0.5 * acq_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};
-        }
+    for (int i = 0; i < geom.n_tx; ++i) {
+        auto it = seg_of.find(geom.output_idx[i]);
+        if (it == seg_of.end())
+            throw std::runtime_error("pair timestamps: нет окна для Tx");
+        const SegSlot& s = it->second;
+        PairTiming timing;
+        timing.tx_interleave_slot = s.slot;
+        timing.tx_interleave_slots = s.n_slots;
+        timing.segment_start_s = s.t0;
+        timing.segment_time_s = s.t1 - s.t0;
+        timing.t_center_s = 0.5 * (s.t0 + s.t1);
+        double dz = mot.direction_sign * mot.speed_m_s * (timing.t_center_s - t_ref) * cos_theta;
+        timing.dtau_pair_s = 2.0 * dz / geom.v;
+        for (int j = 0; j < geom.n_rx; ++j) ts[{i, j}] = timing;
     }
     return ts;
 }
@@ -37,7 +87,7 @@ static double median_of_diff(const std::vector<double>& f) {
 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) {
+    double f_start, double f_stop) {
     std::vector<c128> out = s21;
     int n_full = (int)freq_full.size();
     // индексы в полосе
@@ -47,21 +97,23 @@ std::vector<c128> apply_intra_sweep_phase_correction(
     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;
+    // Сегмент проходит все n_full частот за своё окно; на каждой частоте по
+    // очереди работают его n_slots передатчиков, поэтому событие частоты i для
+    // слота slot имеет номер (n_slots*i + slot) из (n_slots*n_full).
+    int n_slots = std::max(1, info.tx_interleave_slots);
     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;
+    double dt_event = info.segment_time_s / ((double)n_slots * n_full);
+    double cos_theta = std::cos(cfg.look_angle_deg * M_PI / 180.0);
+    double pair_dtau = cfg.apply_pair_motion_correction ? info.dtau_pair_s : 0.0;
+    if (!cfg.apply_freq_phase_correction && pair_dtau == 0.0) 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 dtau_intra = 0.0;
+        if (cfg.apply_freq_phase_correction) {
+            double t_abs = info.segment_start_s + (double)(n_slots * i + slot) * dt_event;
+            double dt_intra = t_abs - info.t_center_s;
+            dtau_intra = 2.0 * cfg.direction_sign * cfg.speed_m_s * cos_theta * dt_intra / wave_speed;
+        }
+        double phi = 2.0 * M_PI * freq_full[i] * (dtau_intra + pair_dtau);
         double ang = cfg.motion_phase_sign * phi;
         out[i] *= c128(std::cos(ang), std::sin(ang));
     }
@@ -163,7 +215,7 @@ AScanSet preprocess_snapshot(
         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);
+                                                   bp.f_start, bp.f_stop);
         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);

data_acq_and_processing/gpr_cpp/src/geometry.cpp

@@ -1,5 +1,6 @@
 #include "gpr/geometry.hpp"
 
+#include <algorithm>
 #include <cmath>
 
 namespace gpr {
@@ -48,20 +49,53 @@ std::array<double, 3> point_to_scene(double x, double y, double z, const MotionC
                                      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 ys = y * std::cos(p) - z * std::sin(p) + sar.ground_depth_m;
     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) {
+double ground_plane_y(double depth, double ground_depth_m, double pitch_rad) {
+    return (-ground_depth_m + depth * std::sin(pitch_rad)) / std::cos(pitch_rad);
+}
+
+Grid build_full_region_grid(const GeometryConfig& geom, const BPConfig& bp, const SARConfig& sar,
+                            double pitch_deg) {
+    // numpy: linspace(lo, hi, max(round(span/step)+1, 2)); nearbyint = банковское
+    // округление python round()
+    auto axis = [](double lo, double hi, double step, std::vector<double>& out) {
+        int n = std::max((int)std::nearbyint((hi - lo) / step) + 1, 2);
         out.resize(n);
-        for (int i = 0; i < n; ++i) out[i] = a + (b - a) * i / (double)(n - 1);
+        for (int i = 0; i < n; ++i) out[i] = lo + (hi - lo) * 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);
+    double x_lo = geom.x_tx[0], x_hi = geom.x_tx[0];
+    for (double x : geom.x_tx) { x_lo = std::min(x_lo, x); x_hi = std::max(x_hi, x); }
+    for (double x : geom.x_rx) { x_lo = std::min(x_lo, x); x_hi = std::max(x_hi, x); }
+    x_lo -= sar.x_margin_m;
+    x_hi += sar.x_margin_m;
+
+    Grid g;
+    axis(x_lo, x_hi, sar.voxel_m[0], g.x);
+    axis(sar.y_min_m, sar.y_max_m, sar.voxel_m[1], g.y);
+    if (sar.coord_mode_ground) {
+        // Радарная z линейна по Z сцены при фиксированной высоте:
+        // z_rad(Y, Z=0) = -(Y - ground_depth)*sin(p), z = (depth - z_rad)/cos(p).
+        double p = pitch_deg * M_PI / 180.0;
+        double cos_p = std::cos(p), sin_p = std::sin(p);
+        double z_lo = 0.0, z_hi = 0.0;
+        bool first = true;
+        for (double y : {sar.y_min_m, sar.y_max_m}) {
+            for (double depth : {bp.min_depth, bp.max_depth}) {
+                double z_rad0 = -(y - sar.ground_depth_m) * sin_p;
+                double z = (depth - z_rad0) / cos_p;
+                if (first) { z_lo = z_hi = z; first = false; }
+                z_lo = std::min(z_lo, z);
+                z_hi = std::max(z_hi, z);
+            }
+        }
+        axis(z_lo, z_hi, sar.voxel_m[2], g.z);
+    } else {
+        axis(bp.min_depth, bp.max_depth, sar.voxel_m[2], g.z);
+    }
     return g;
 }
 

data_acq_and_processing/gpr_cpp/src/kernel_src.hpp

@@ -1,28 +1,30 @@
-// Исходник CUDA-ядра для NVRTC (порт bp_aperture_contribution из gpr_core/kernels.py).
-// Self-contained complex<float> (без cupy-заголовков).
+// Исходник CUDA-ядра для NVRTC (порт sar_frame_contribution из
+// Multi_antenna_3D_clean_v1.ipynb). Self-contained complex<float>.
+//
+// Геометрия (координаты, дальности, tau, вес компенсации) считается в double:
+// эталон ноутбука — float64, а в float32 позиция выборки A-скана (тысячи
+// отсчётов) теряет ~3e-4 отсчёта и расходится с эталоном на ~1e-4. A-сканы
+// остаются float32 — их значения O(1) и линейная интерполяция не теряет.
 #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 double t0, const double inv_dt, const int n_fft, const double inv_v,
+    const double* __restrict__ xt, const double* __restrict__ yt, const double* __restrict__ zt,
+    const double* __restrict__ xr, const double* __restrict__ yr, const double* __restrict__ zr,
+    const double* __restrict__ inv_geo_ref, const float* __restrict__ wgt,
+    const double* __restrict__ gx, const double* __restrict__ gy, const double* __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,
+    const double cos_p, const double sin_p, const double ground_depth,
+    const double min_depth, const double max_depth,
+    const double comp_pow, const double w_max, const int use_comp,
     cf* __restrict__ acc, float* __restrict__ cnt)
 {
     const long total = (long)NX * NY * NZ;
@@ -31,39 +33,45 @@ extern "C" __global__ void bp_aperture_contribution(
         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;
+        double x = gx[ix];
+        double zs = gz[iz];
+        // Сцена → оси радара: Y=0 сцены — земля, ground_depth ниже нижнего ряда
+        // антенн (порт sar_scene_to_radar_yz).
+        double y0 = gy[iy] - ground_depth;
+        double y_rad = y0 * cos_p + zs * sin_p;
+        double z_rad = -y0 * sin_p + zs * cos_p;
         bool depth_ok = (z_rad >= min_depth) && (z_rad <= max_depth);
 
-        cf a = {0.f, 0.f};
+        double ax = 0.0, ay = 0.0;
         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);
+            double dxt = x - xt[i], dyt = y_rad - yt[i], dzt = z_rad - zt[i];
+            double dxr = x - xr[i], dyr = y_rad - yr[i], dzr = z_rad - zr[i];
+            double Rt = sqrt(dxt*dxt + dyt*dyt + dzt*dzt);
+            double Rr = sqrt(dxr*dxr + dyr*dyr + dzr*dzr);
+            double pos = ((Rt + Rr) * inv_v - t0) * inv_dt;
+            bool inside = (pos >= 0.0) && (pos <= (double)(n_fft - 1));
+            if (!inside || !depth_ok) continue;
+            int k = (int)floor(pos);
             if (k < 0) k = 0;
             if (k > n_fft - 2) k = n_fft - 2;
-            float f = pos - (float)k;
+            float f = (float)(pos - (double)k);
             const cf* h = H + (long)i * n_fft;
-            cf hv = h[k] + f * (h[k+1] - h[k]);
+            float hx = h[k].x + f * (h[k+1].x - h[k].x);
+            float hy = h[k].y + f * (h[k+1].y - h[k].y);
+            double w = (double)wgt[i];
             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;
+                double ratio = (1.0/(Rt*Rr + 1e-12) + 1e-30) * inv_geo_ref[i];
+                double ww = 1.0/(pow(ratio, comp_pow) + 1e-12);
+                ww = fmin(fmax(ww, 0.0), w_max);
+                w *= ww;
             }
-            if (depth_ok) { a = a + (w * hv); c += w; }
+            ax += w * (double)hx;
+            ay += w * (double)hy;
+            c += wgt[i];
         }
-        acc[gid] = a;
+        acc[gid].x = (float)ax;
+        acc[gid].y = (float)ay;
         cnt[gid] = c;
     }
 }

data_acq_and_processing/gpr_cpp/src/sar3d.cpp

@@ -31,11 +31,18 @@ VolumeRaw SarEngine::aperture_volume(const std::vector<ApertureFrame>& frames, c
     int NX = grid.NX(), NY = grid.NY(), NZ = grid.NZ();
     size_t nvox = (size_t)NX * NY * NZ;
 
+    double pitch = mot_.pitch_deg * M_PI / 180.0;
+    // Опорное расхождение пары — на плоскости земли на глубине comp_ref_depth.
+    double geo_ref_y = sar_.coord_mode_ground
+                           ? ground_plane_y(bp_.comp_ref_depth, sar_.ground_depth_m, pitch)
+                           : 0.0;
+
     // Плоские массивы по (кадр, пара)
     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);
+    std::vector<double> xt(n_contrib), yt(n_contrib), zt(n_contrib);
+    std::vector<double> xr(n_contrib), yr(n_contrib), zr(n_contrib);
+    std::vector<double> igr(n_contrib);
+    std::vector<float> wgt(n_contrib);
     int idx = 0;
     for (int f = 0; f < n_frames; ++f) {
         const AScanSet& a = *frames[f].ascans;
@@ -49,49 +56,47 @@ VolumeRaw SarEngine::aperture_volume(const std::vector<ApertureFrame>& frames, c
             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));
+            xt[idx] = geom_.x_tx[i] + dx; yt[idx] = geom_.y_tx[i] + dy;
+            zt[idx] = geom_.z_tx[i] + dz;
+            xr[idx] = geom_.x_rx[j] + dx; yr[idx] = geom_.y_rx[j] + dy;
+            zr[idx] = geom_.z_rx[j] + dz;
+            igr[idx] = 1.0 / pair_geo_ref(geom_, bp_, i, j, geo_ref_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;
+    double cos_p = 1.0, sin_p = 0.0, ground_depth = 0.0;
     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);
+        cos_p = std::cos(pitch);
+        sin_p = std::sin(pitch);
+        ground_depth = sar_.ground_depth_m;
     }
 
     // Заливка на устройство
     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);
+    DeviceBuffer<double> dxt(n_contrib), dyt(n_contrib), dzt(n_contrib);
+    DeviceBuffer<double> dxr(n_contrib), dyr(n_contrib), dzr(n_contrib);
+    DeviceBuffer<double> digr(n_contrib);
+    DeviceBuffer<float> 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<double> dgx(NX), dgy(NY), dgz(NZ);
+    dgx.upload(grid.x.data(), NX); dgy.upload(grid.y.data(), NY); dgz.upload(grid.z.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;
+    double inv_dt = 1.0 / dt, inv_v = 1.0 / geom_.v;
+    double min_d = bp_.min_depth, max_d = bp_.max_depth;
+    double comp_pow = bp_.comp_range_power, w_max = 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,
+        dH.addr(), &t0, &inv_dt, &n_fft, &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,
+        &NX, &NY, &NZ, &n_contrib, &cos_p, &sin_p, &ground_depth, &min_d, &max_d,
         &comp_pow, &w_max, &use_comp, dacc.addr(), dcnt.addr()};
 
     int block = 256;
@@ -108,4 +113,113 @@ VolumeRaw SarEngine::aperture_volume(const std::vector<ApertureFrame>& frames, c
     return out;
 }
 
+namespace {
+
+// scipy.ndimage.gaussian_filter: ядро радиуса int(truncate*sigma+0.5), reflect.
+std::vector<double> gaussian_kernel_1d(double sigma, int radius) {
+    std::vector<double> w(2 * radius + 1);
+    double sum = 0.0;
+    for (int i = -radius; i <= radius; ++i) {
+        w[i + radius] = std::exp(-0.5 * i * i / (sigma * sigma));
+        sum += w[i + radius];
+    }
+    for (auto& v : w) v /= sum;
+    return w;
+}
+
+// Корреляция вдоль оси с reflect-границами ('d c b a | a b c d').
+inline int reflect_index(int i, int n) {
+    while (i < 0 || i >= n) {
+        if (i < 0) i = -i - 1;
+        if (i >= n) i = 2 * n - i - 1;
+    }
+    return i;
+}
+
+}  // namespace
+
+FinishedVolume finish_volume(const VolumeRaw& raw, const Grid& grid, double pitch_deg,
+                             const BPConfig& bp, const SARConfig& sar) {
+    int NX = raw.NX, NY = raw.NY, NZ = raw.NZ;
+    size_t nvox = (size_t)NX * NY * NZ;
+    FinishedVolume out;
+    out.NX = NX; out.NY = NY; out.NZ = NZ;
+
+    std::vector<double> vol(nvox);
+    for (size_t i = 0; i < nvox; ++i)
+        vol[i] = std::abs(c128(raw.acc[i].real(), raw.acc[i].imag())) / ((double)raw.cnt[i] + 1e-12);
+
+    // Маска глубины в радарных осях (в ядре невалидные вкладов не получали,
+    // но |0/(0+eps)| уже 0; обнуление здесь повторяет ноутбук буквально).
+    if (sar.coord_mode_ground) {
+        double p = pitch_deg * M_PI / 180.0;
+        double cos_p = std::cos(p), sin_p = std::sin(p);
+        for (int iy = 0; iy < NY; ++iy) {
+            double y0 = grid.y[iy] - sar.ground_depth_m;
+            for (int iz = 0; iz < NZ; ++iz) {
+                double z_rad = -y0 * sin_p + grid.z[iz] * cos_p;
+                if (z_rad >= bp.min_depth && z_rad <= bp.max_depth) continue;
+                double* row = vol.data() + ((size_t)iy * NZ + iz) * NX;
+                std::fill(row, row + NX, 0.0);
+            }
+        }
+    }
+
+    // Плоскости Y для нормировки: без верхних/нижних norm_edge_margin_m.
+    std::vector<bool> ref_plane(NY, true);
+    double margin = sar.norm_edge_margin_m;
+    double y_span = grid.y.back() - grid.y.front();
+    if (margin > 0.0 && y_span > 2.5 * margin) {
+        for (int iy = 0; iy < NY; ++iy)
+            ref_plane[iy] = grid.y[iy] >= grid.y.front() + margin &&
+                            grid.y[iy] <= grid.y.back() - margin;
+    }
+    auto ref_max = [&](const std::vector<double>& v) {
+        double m = 0.0;
+        for (int iy = 0; iy < NY; ++iy) {
+            if (!ref_plane[iy]) continue;
+            const double* plane = v.data() + (size_t)iy * NZ * NX;
+            for (size_t i = 0; i < (size_t)NZ * NX; ++i) m = std::max(m, plane[i]);
+        }
+        return m + 1e-12;
+    };
+
+    double level_raw = ref_max(vol);
+    for (auto& v : vol) v /= level_raw;
+
+    // Гаусс sigma=(0, s, s): сглаживание по Z и X внутри каждого Y-плана
+    // (порядок осей как у scipy: сначала ось 1, затем ось 2).
+    if (sar.smooth_sigma > 0.0) {
+        int radius = (int)(4.0 * sar.smooth_sigma + 0.5);
+        const auto kernel = gaussian_kernel_1d(sar.smooth_sigma, radius);
+        std::vector<double> tmp((size_t)NZ * NX);
+        for (int iy = 0; iy < NY; ++iy) {
+            double* plane = vol.data() + (size_t)iy * NZ * NX;
+            // вдоль Z (ось 1)
+            for (int iz = 0; iz < NZ; ++iz)
+                for (int ix = 0; ix < NX; ++ix) {
+                    double s = 0.0;
+                    for (int k = -radius; k <= radius; ++k)
+                        s += kernel[k + radius] * plane[(size_t)reflect_index(iz + k, NZ) * NX + ix];
+                    tmp[(size_t)iz * NX + ix] = s;
+                }
+            // вдоль X (ось 2)
+            for (int iz = 0; iz < NZ; ++iz)
+                for (int ix = 0; ix < NX; ++ix) {
+                    double s = 0.0;
+                    for (int k = -radius; k <= radius; ++k)
+                        s += kernel[k + radius] * tmp[(size_t)iz * NX + reflect_index(ix + k, NX)];
+                    plane[(size_t)iz * NX + ix] = s;
+                }
+        }
+    }
+
+    double level_smooth = ref_max(vol);
+    out.volume.resize(nvox);
+    for (size_t i = 0; i < nvox; ++i)
+        out.volume[i] = (float)std::max(vol[i] / level_smooth, 0.0);
+    out.level = level_raw * level_smooth;
+    return out;
+}
+
 }  // namespace gpr