From 12605a02642a5efdd184f0de8613fdb4e53fe295 Mon Sep 17 00:00:00 2001 From: Weiqun Zhang Date: Mon, 10 Aug 2026 09:31:26 -0700 Subject: [PATCH] Hoist the eb_flow tracer device vector out of the MFIter loop. It was constructed per box and destroyed right after the async ParallelFor that captures its data() pointer, so a kernel could read freed device memory. At function scope it outlives the loop, whose MFIter destructor synchronizes every stream used. The values are box-independent, so this also drops a per-box host-to-device copy. Fixes #170. --- src/boundary_conditions/incflo_set_bcs.cpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/src/boundary_conditions/incflo_set_bcs.cpp b/src/boundary_conditions/incflo_set_bcs.cpp index dde7a5c6..c07d70ce 100644 --- a/src/boundary_conditions/incflo_set_bcs.cpp +++ b/src/boundary_conditions/incflo_set_bcs.cpp @@ -362,6 +362,15 @@ incflo::set_eb_tracer (int lev, Real /*time*/, MultiFab& eb_tracer, int nghost) const auto& factory = dynamic_cast(eb_tracer.Factory()); + // These are box-independent, and must outlive every kernel launched in the + // MFIter loop below, which captures eb_flow_tracer. + int num_trac = m_ntrac; + Gpu::DeviceVector eb_flow_tracer_dv; + for(int n(0); n < num_trac; ++n) { + eb_flow_tracer_dv.push_back(m_eb_flow.tracer[n]); + } + Real* eb_flow_tracer = eb_flow_tracer_dv.data(); + #ifdef _OPENMP #pragma omp parallel if (Gpu::notInLaunchRegion()) #endif @@ -387,14 +396,6 @@ incflo::set_eb_tracer (int lev, Real /*time*/, MultiFab& eb_tracer, int nghost) Real norm_tol_lo = Real(-1.) - (normal_tol + pad); Real norm_tol_hi = Real(-1.) + (normal_tol + pad); - int num_trac = m_ntrac; - // Create a device vector - Gpu::DeviceVector eb_flow_tracer_dv; - for(int n(0); n < num_trac; ++n) { - eb_flow_tracer_dv.push_back(m_eb_flow.tracer[n]); - } - Real* eb_flow_tracer = eb_flow_tracer_dv.data(); - ParallelFor(bx, [flags_arr,eb_tracer_arr,norm_arr,has_normal,normal, norm_tol_lo, norm_tol_hi, num_trac, eb_flow_tracer] AMREX_GPU_DEVICE (int i, int j, int k) noexcept