From 60f6c9e56742f3226cd5d956c0843f73f3cb4332 Mon Sep 17 00:00:00 2001 From: awe Date: Fri, 25 Sep 2026 17:01:22 +0300 Subject: [PATCH] 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 --- .../gpr_cpp/src/kernel_src.hpp | 19 +++++---- data_acq_and_processing/gpr_cpp/src/sar3d.cpp | 42 +++++++++++-------- 2 files changed, 36 insertions(+), 25 deletions(-) diff --git a/data_acq_and_processing/gpr_cpp/src/kernel_src.hpp b/data_acq_and_processing/gpr_cpp/src/kernel_src.hpp index ff144d7..6ca1da7 100644 --- a/data_acq_and_processing/gpr_cpp/src/kernel_src.hpp +++ b/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 = blockIdx.x * blockDim.x + threadIdx.x; - gid < total; gid += (long)gridDim.x * blockDim.x) { + for (long gid = gid_begin + blockIdx.x * blockDim.x + threadIdx.x; + gid < gid_end; 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].y = (float)ay; - cnt[gid] = c; + acc[gid - gid_begin].x = (float)ax; + acc[gid - gid_begin].y = (float)ay; + cnt[gid - gid_begin] = c; } } )CUDA"; diff --git a/data_acq_and_processing/gpr_cpp/src/sar3d.cpp b/data_acq_and_processing/gpr_cpp/src/sar3d.cpp index ff9555e..97fb721 100644 --- a/data_acq_and_processing/gpr_cpp/src/sar3d.cpp +++ b/data_acq_and_processing/gpr_cpp/src/sar3d.cpp @@ -84,32 +84,40 @@ VolumeRaw SarEngine::aperture_volume(const std::vector& frames, c digr.upload(igr.data(), n_contrib); dwgt.upload(wgt.data(), n_contrib); DeviceBuffer 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 dacc(nvox); - DeviceBuffer dcnt(nvox); + + // Объём считается чанками с переиспользуемыми acc/cnt-буферами: полные + // буферы (~115 МБ на сетку 9.6 млн вокселей) могут не влезать в остаток + // VRAM, когда карту делят десктоп-процессы. + const size_t chunk_vox = std::min(nvox, 2'000'000); + DeviceBuffer dacc(chunk_vox); + DeviceBuffer 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; - 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, - &comp_pow, &w_max, &use_comp, dacc.addr(), dcnt.addr()}; - - int block = 256; - int gridDim = (int)std::min(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); + 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, &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(65535, ((long)(end - begin) + 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); + dcnt.download(out.cnt.data() + begin, end - begin); + } return out; }