make gpr update

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