diff --git a/crates/specialized/rtx-cfd/src/kernels/cuda/e3_cut.cu b/crates/specialized/rtx-cfd/src/kernels/cuda/e3_cut.cu index bb71d10..be125a7 100644 --- a/crates/specialized/rtx-cfd/src/kernels/cuda/e3_cut.cu +++ b/crates/specialized/rtx-cfd/src/kernels/cuda/e3_cut.cu @@ -235,14 +235,20 @@ __device__ double cut_face_update(const E3Params& g, const E3Ptrs& f, const E3Cu (void)n; (void)x; /* convection through the plus face */ double u_plus, delta_plus; - if (f_up1 >= 0) { + /* bit 10 of wall_order (A2): a prescribed neighbour face exchanges no convective momentum */ + const int* opn = c == 0 ? m.open_u : (c == 1 ? m.open_v : m.open_w); + int presc_up = (g.wall_order & 1024) && f_up1 >= 0 && !opn[f_up1]; + int presc_dn = (g.wall_order & 1024) && f_dn1 >= 0 && !opn[f_dn1]; + if (presc_up) { u_plus = u0; delta_plus = 0.0; } + else if (f_up1 >= 0) { if (g.scheme == SCHEME_UPWIND) delta_plus = 0.0; else if (m_plus >= 0.0) delta_plus = face_corr3(g.scheme, f_dn1 >= 0, dn1, u0, up1); else delta_plus = face_corr3(g.scheme, f_up2 >= 0, up2, up1, u0); u_plus = upwind3(m_plus, u0, up1); } else { u_plus = upwind3(m_plus, u0, beyond_p); delta_plus = 0.0; } double u_minus, delta_minus; - if (f_dn1 >= 0) { + if (presc_dn) { u_minus = u0; delta_minus = 0.0; } + else if (f_dn1 >= 0) { if (g.scheme == SCHEME_UPWIND) delta_minus = 0.0; else if (m_minus >= 0.0) delta_minus = face_corr3(g.scheme, f_dn2 >= 0, dn2, dn1, u0); else delta_minus = face_corr3(g.scheme, f_up1 >= 0, up1, u0, dn1); diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/cutwall.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/cutwall.rs index 8d4f60a..544687c 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/cutwall.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/cutwall.rs @@ -217,6 +217,7 @@ impl Mask { diffusion_transverse: false, distance_floor_fine: false, wall_advancing: false, + exchange_convection_off: false, wall_exchange_axis: false, grad_weights: None, diffusion_centroid: false, diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/exchange.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/exchange.rs index 8a06e1d..0fcc1a2 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/exchange.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/exchange.rs @@ -191,7 +191,9 @@ impl Mask { scheme.face_correction(up2, un, u0) }; let u_face = upwind(m_plus, u0, un) + delta; - convective[c] -= -rho * m_plus * (u_face - u0); + if !self.exchange_convection_off { + convective[c] -= -rho * m_plus * (u_face - u0); + } force[c] -= mu * cv.ap[d][1] * a_d * (un - u0) / solid_spacing(1.0); } @@ -208,7 +210,9 @@ impl Mask { scheme.face_correction(up1, u0, ud) }; let u_face = upwind(m_minus, ud, u0) + delta; - convective[c] -= rho * m_minus * (u_face - u0); + if !self.exchange_convection_off { + convective[c] -= rho * m_minus * (u_face - u0); + } force[c] -= mu * cv.ap[d][0] * a_d * (ud - u0) / solid_spacing(-1.0); } diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/poisson/mod.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/poisson/mod.rs index e91763b..5e440e7 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/poisson/mod.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/poisson/mod.rs @@ -13,6 +13,7 @@ mod hierarchy; mod pcg; mod problem; +#[cfg(feature = "cuda")] pub use device_cg::GuessBasis; pub use export::{LevelExport, export_hierarchy, vcycle_f32_reference}; pub use hierarchy::Hierarchy; diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/cut_predictor.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/cut_predictor.rs index 1209abd..d3aaaa0 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/cut_predictor.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/cut_predictor.rs @@ -163,7 +163,19 @@ impl Solver { } }; // Convection through the plus face. + // A2: a prescribed neighbour face exchanges no convective momentum. + let prescribed = |q: [i64; 3]| -> bool { + mask.exchange_convection_off + && lat.face(c, q).is_some_and(|f| { + (match c { + 0 => mask.u_kind(f), + 1 => mask.v_kind(f), + _ => mask.w_kind(f), + }) != super::super::wall::FaceKind::Fluid + }) + }; let (u_plus, delta_plus) = match up1 { + Some(_) if prescribed(add(p, ed, 1)) => (u0, 0.0), Some(un) => { let delta = if scheme == ConvectionScheme::Upwind { 0.0 @@ -177,6 +189,7 @@ impl Solver { None => (Self::upwind(m_plus, u0, beyond(true)), 0.0), }; let (u_minus, delta_minus) = match dn1 { + Some(_) if prescribed(add(p, ed, -1)) => (u0, 0.0), Some(ud) => { let delta = if scheme == ConvectionScheme::Upwind { 0.0 diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/device.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/device.rs index cd1ba8e..e28165e 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/device.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/device.rs @@ -304,7 +304,9 @@ impl DeviceStep { ) + 256 * i32::from(self.solver.params.distance_floor_fine) // bit 9: the advancing-wall friction closure (A3-i). - + 512 * i32::from(self.solver.params.wall_advancing), + + 512 * i32::from(self.solver.params.wall_advancing) + // bit 10: no convective exchange with prescribed faces (A2). + + 1024 * i32::from(self.solver.params.exchange_convection_off), dx: g.dx, dy: g.dy, dz: g.dz, diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/mod.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/mod.rs index 53caba5..42f152c 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/mod.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/mod.rs @@ -120,6 +120,12 @@ pub struct Parameters { /// normal to itself sets up (thinner than a cell on the fast flag). f(0) = 1 /// exactly: static cases unchanged. `RTX_E3_WALL_ADVANCING=1`. pub wall_advancing: bool, + /// A2: no convective momentum exchange with a PRESCRIBED neighbour face — + /// the flux through a control-volume side that adjoins one advects the + /// volume's own value (the upwinded exchange was a first-order momentum + /// sink the wall collected: 4 % of the flag's drag at ny 62). + /// `RTX_E3_EXCHANGE_CONVECTION=off`. + pub exchange_convection_off: bool, /// The diffusive exchange of a fluid face with a SOLID neighbour face /// over the axis distance to the wall, `δ = min(h, d_f/|n_d|)`, and /// implicit — instead of the full `h`, which places the no-slip value @@ -171,6 +177,8 @@ impl Default for Parameters { distance_floor_fine: std::env::var("RTX_E3_DISTANCE_FLOOR") .map_or(true, |v| v != "coarse"), wall_advancing: std::env::var("RTX_E3_WALL_ADVANCING").is_ok_and(|v| v == "1"), + exchange_convection_off: std::env::var("RTX_E3_EXCHANGE_CONVECTION") + .is_ok_and(|v| v == "off"), wall_exchange_axis: std::env::var("RTX_E3_WALL_EXCHANGE").is_ok_and(|v| v == "axis"), pressure_centroid: std::env::var("RTX_E3_PRESSURE_CENTROID").is_ok_and(|v| v == "1"), // ON by default since S2-5 (`=0` reproduces the records before it). @@ -311,6 +319,7 @@ impl Solver { self.params.diffusion_transverse && self.params.diffusion_centroid; m.distance_floor_fine = self.params.distance_floor_fine; m.wall_advancing = self.params.wall_advancing; + m.exchange_convection_off = self.params.exchange_convection_off; m.diffusion_centroid = self.params.diffusion_centroid; if self.params.diffusion_centroid { m.compute_face_shifts(); diff --git a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/wall.rs b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/wall.rs index 6ef8c11..bd7a30a 100644 --- a/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/wall.rs +++ b/crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/wall.rs @@ -105,6 +105,7 @@ pub struct Mask { pub(super) diffusion_transverse: bool, pub(super) distance_floor_fine: bool, pub(super) wall_advancing: bool, + pub(super) exchange_convection_off: bool, /// The axis-distance implicit wall exchange (S2-5). pub(super) wall_exchange_axis: bool, /// The centroid prototype's pressure-gradient weights per u / v / w face. @@ -531,6 +532,7 @@ impl Mask { diffusion_transverse: false, distance_floor_fine: false, wall_advancing: false, + exchange_convection_off: false, wall_exchange_axis: false, grad_weights: None, diffusion_centroid: false,