sar3d: chunk the aperture kernel over voxels to survive low free VRAM

Full-volume acc/cnt device buffers (~115 MB for the 9.6M-voxel grid) plus
the CUDA context did not fit next to a desktop session that had taken the
T600's 4 GB down to ~190 MB free, so every frame died with cuMemAlloc OOM
and the GUI showed no sar3d picture. The kernel now processes
[gid_begin, gid_end) ranges and the host launches it in 2M-voxel chunks
reusing one ~24 MB buffer pair. Parity vs the reference volumes is
unchanged (max diff 1.3e-6).

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

data_acq_and_processing/gpr_cpp/src/kernel_src.hpp

@@ -12,8 +12,11 @@ namespace gpr {
 inline const char* APERTURE_KERNEL_SRC = R"CUDA(
 struct cf { float x, y; };
 
-// Весь взвешенный объём апертуры за один запуск: один поток на воксель,
+// Взвешенный объём апертуры: один поток на воксель, внутренний цикл по
-// внутренний цикл по n_contrib = n_frames*n_pairs. acc/cnt пишутся один раз.
+// n_contrib = n_frames*n_pairs. Обрабатывается диапазон вокселей
+// [gid_begin, gid_end) — хост запускает ядро чанками, переиспользуя небольшие
+// acc/cnt-буферы (полный объём в VRAM не обязан помещаться: на десктопе с
+// занятой другими процессами картой свободных может быть <200 МБ).
 extern "C" __global__ void bp_aperture_contribution(
     const cf* __restrict__ H,
     const double t0, const double inv_dt, const int n_fft, const double inv_v,
@@ -22,14 +25,14 @@ extern "C" __global__ void bp_aperture_contribution(
     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 long gid_begin, const long gid_end,
     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;
+    for (long gid = gid_begin + blockIdx.x * blockDim.x + threadIdx.x;
-    for (long gid = blockIdx.x * blockDim.x + threadIdx.x;
+         gid < gid_end; gid += (long)gridDim.x * blockDim.x) {
-         gid < total; gid += (long)gridDim.x * blockDim.x) {
         int ix = gid % NX;
         int iz = (gid / NX) % NZ;
         int iy = gid / ((long)NX * NZ);
@@ -70,9 +73,9 @@ extern "C" __global__ void bp_aperture_contribution(
             ay += w * (double)hy;
             c += wgt[i];
         }
-        acc[gid].x = (float)ax;
+        acc[gid - gid_begin].x = (float)ax;
-        acc[gid].y = (float)ay;
+        acc[gid - gid_begin].y = (float)ay;
-        cnt[gid] = c;
+        cnt[gid - gid_begin] = c;
     }
 }
 )CUDA";

data_acq_and_processing/gpr_cpp/src/sar3d.cpp

@@ -84,32 +84,40 @@ VolumeRaw SarEngine::aperture_volume(const std::vector<ApertureFrame>& frames, c
     digr.upload(igr.data(), n_contrib); dwgt.upload(wgt.data(), n_contrib);
     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);
+    // Объём считается чанками с переиспользуемыми acc/cnt-буферами: полные
+    // буферы (~115 МБ на сетку 9.6 млн вокселей) могут не влезать в остаток
+    // VRAM, когда карту делят десктоп-процессы.
+    const size_t chunk_vox = std::min<size_t>(nvox, 2'000'000);
+    DeviceBuffer<c64> dacc(chunk_vox);
+    DeviceBuffer<float> dcnt(chunk_vox);
 
     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;
 
+    VolumeRaw out;
+    out.NX = NX; out.NY = NY; out.NZ = NZ;
+    out.acc.resize(nvox); out.cnt.resize(nvox);
+    int block = 256;
+    for (size_t begin = 0; begin < nvox; begin += chunk_vox) {
+        const size_t end = std::min(nvox, begin + chunk_vox);
+        long gid_begin = (long)begin, gid_end = (long)end;
         void* args[] = {
             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, &ground_depth, &min_d, &max_d,
+            &NX, &NY, &NZ, &n_contrib, &gid_begin, &gid_end,
+            &cos_p, &sin_p, &ground_depth, &min_d, &max_d,
             &comp_pow, &w_max, &use_comp, dacc.addr(), dcnt.addr()};
-
+        int gridDim = (int)std::min<long>(65535, ((long)(end - begin) + block - 1) / block);
-    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");
-
+        dacc.download(out.acc.data() + begin, end - begin);
-    VolumeRaw out;
+        dcnt.download(out.cnt.data() + begin, end - begin);
-    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;
 }