embedded3 S2-6: the oblique-wall instrument (tests/embedded3_wall_position_oblique.rs: z-flow / in-plane Poiseuille + Couette linear exactness; gate oblique_wall_position_is_second_order) and the closures it found, host + device, default OFF — the transverse centroid correction (RTX_E3_DIFFUSION_TRANSVERSE=1, wall_order bit 7), the fine distance floor (RTX_E3_DISTANCE_FLOOR=fine, bit 8); kernel geometry factored into cut_cv; DFG 2D-1 x-shift knob (RTX_E3_DFG_SHIFT_X); FLAG_X0 0.6 -> 0.25 in the three flag tests (the flag of every record was detached)

Co-Authored-By: Claude Fable 5.1 <[email protected]>

crates/specialized/rtx-cfd/src/kernels/cuda/e3_cut.cu

@@ -15,6 +15,7 @@
  */
 #define CUT_INERTIA_FLOOR 0.1
 #define CUT_DISTANCE_FLOOR 0.05
+#define CUT_DISTANCE_FLOOR_FINE 0.01
 
 struct E3Cut {
     const double *a_u, *a_v, *a_w;      /* apertures per face */
@@ -61,20 +62,15 @@ __device__ __forceinline__ double cut_ap_at(const E3Params& g, const E3Cut& m, i
     return fidx < 0 ? dflt : cut_ap(m, cc)[fidx];
 }
 
-/* The predicted value of the open face of component c at (i, j, k). */
-__device__ double cut_face_update(const E3Params& g, const E3Ptrs& f, const E3Cut& m, int c, int i, int j, int k, int fidx)
+/* The control volume's geometry of the face of component c at (i, j, k)
+   (cutwall.rs cv_geometry): side apertures, the closing wall vector, the
+   floored wall distance. */
+__device__ void cut_cv(const E3Params& g, const E3Cut& m, int c, int i, int j, int k, int fidx,
+                       double* alpha_out, double* apm, double* app, double* wall, double* distance_out)
 {
     double h[3] = { g.dx, g.dy, g.dz };
-    double n[3] = { (double)g.nx, (double)g.ny, (double)g.nz };
     double area[3] = { g.dy * g.dz, g.dx * g.dz, g.dx * g.dy };
-    double rho = g.rho, mu = g.nu * g.rho;
-    int sides[3][2] = { { g.bx0, g.bx1 }, { g.by0, g.by1 }, { g.bz0, g.bz1 } };
-    int p[3] = { i, j, k };
-    const double* old_c = cut_old(f, c);
-    double u0 = old_c[fidx];
-    /* the control volume's geometry (cutwall.rs cv_geometry) */
     double alpha = cut_ap(m, c)[fidx];
-    double apm[3], app[3];
     int cm[3] = { i, j, k }; cm[c] -= 1;   /* cell minus */
     int cp[3] = { i, j, k };               /* cell plus */
     for (int d = 0; d < 3; ++d) {
@@ -92,7 +88,6 @@ __device__ double cut_face_update(const E3Params& g, const E3Ptrs& f, const E3Cu
             app[d] = 0.5 * (cut_ap_at(g, m, d, q1[0], q1[1], q1[2], 1.0) + cut_ap_at(g, m, d, q2[0], q2[1], q2[2], 1.0));
         }
     }
-    double wall[3];
     for (int d = 0; d < 3; ++d) wall[d] = -(app[d] - apm[d]) * area[d];
     double h_min = fmin(fmin(g.dx, g.dy), g.dz);
     const double* dist = c == 0 ? m.d_u : (c == 1 ? m.d_v : m.d_w);
@@ -102,7 +97,64 @@ __device__ double cut_face_update(const E3Params& g, const E3Ptrs& f, const E3Cu
         double a_w0 = sqrt(wall[0] * wall[0] + wall[1] * wall[1] + wall[2] * wall[2]);
         if (a_w0 > 0.0) { double n_c = wall[c] / a_w0; n_t = sqrt(fmax(1.0 - n_c * n_c, 0.0)); }
     }
-    double distance = fmax(dist[fidx] + 0.5 * h[c] * (1.0 - alpha) * n_t, CUT_DISTANCE_FLOOR * h_min);
+    /* bit 8 of wall_order: the fine distance floor (S2-6) */
+    double floor_h = (g.wall_order & 256) ? CUT_DISTANCE_FLOOR_FINE : CUT_DISTANCE_FLOOR;
+    *alpha_out = alpha;
+    *distance_out = fmax(dist[fidx] + 0.5 * h[c] * (1.0 - alpha) * n_t, floor_h * h_min);
+}
+
+/* S2-6 (bit 7 of wall_order): the transverse part of the centroid diffusion
+   between the face (fidx, its geometry given) and its open neighbour fq at
+   q along d — grad(u) . ds_perp with the cut faces' own wall-normal
+   gradients (u - U_b)/d_f (cut_predictor.rs `transverse`). */
+__device__ double cut_transverse(const E3Params& g, const E3Cut& m, int c, int d, int fidx, int fq, const int* q,
+                                 double sign, double u0, double ub, double alpha, const double* wall, double distance,
+                                 double uq, const double* sh, const double* ubt)
+{
+    double ds[3] = { 0.0, 0.0, 0.0 };
+    int any = 0;
+    for (int e = 0; e < 3; ++e) {
+        if (e == d) continue;
+        ds[e] = sign * (sh[3 * fq + e] - sh[3 * fidx + e]);
+        if (ds[e] != 0.0) any = 1;
+    }
+    if (!any) return 0.0;
+    double alpha_q, apm_q[3], app_q[3], wall_q[3], distance_q;
+    cut_cv(g, m, c, q[0], q[1], q[2], fq, &alpha_q, apm_q, app_q, wall_q, &distance_q);
+    double gr[3] = { 0.0, 0.0, 0.0 };
+    double count = 0.0;
+    double a0 = sqrt(wall[0] * wall[0] + wall[1] * wall[1] + wall[2] * wall[2]);
+    if (a0 > 0.0 && alpha < 1.0) {
+        double slope = (u0 - ub) / (distance * a0);
+        for (int e = 0; e < 3; ++e) gr[e] -= slope * wall[e];
+        count += 1.0;
+    }
+    double aq = sqrt(wall_q[0] * wall_q[0] + wall_q[1] * wall_q[1] + wall_q[2] * wall_q[2]);
+    if (aq > 0.0 && alpha_q < 1.0) {
+        double slope = (uq - ubt[fq]) / (distance_q * aq);
+        for (int e = 0; e < 3; ++e) gr[e] -= slope * wall_q[e];
+        count += 1.0;
+    }
+    if (count == 0.0) return 0.0;
+    return (gr[0] * ds[0] + gr[1] * ds[1] + gr[2] * ds[2]) / count;
+}
+
+/* The predicted value of the open face of component c at (i, j, k). */
+__device__ double cut_face_update(const E3Params& g, const E3Ptrs& f, const E3Cut& m, int c, int i, int j, int k, int fidx)
+{
+    double h[3] = { g.dx, g.dy, g.dz };
+    double n[3] = { (double)g.nx, (double)g.ny, (double)g.nz };
+    double area[3] = { g.dy * g.dz, g.dx * g.dz, g.dx * g.dy };
+    double rho = g.rho, mu = g.nu * g.rho;
+    int sides[3][2] = { { g.bx0, g.bx1 }, { g.by0, g.by1 }, { g.bz0, g.bz1 } };
+    int p[3] = { i, j, k };
+    const double* old_c = cut_old(f, c);
+    double u0 = old_c[fidx];
+    /* the control volume's geometry (cutwall.rs cv_geometry) */
+    double alpha, apm[3], app[3], wall[3], distance;
+    cut_cv(g, m, c, i, j, k, fidx, &alpha, apm, app, wall, &distance);
+    int cm[3] = { i, j, k }; cm[c] -= 1;   /* cell minus */
+    int cp[3] = { i, j, k };               /* cell plus */
     const double* ubt = c == 0 ? m.ub_u : (c == 1 ? m.ub_v : m.ub_w);
     double ub = ubt[fidx];
     /* face position */
@@ -194,6 +246,21 @@ __device__ double cut_face_update(const E3Params& g, const E3Ptrs& f, const E3Cu
            the part beyond 1/h point-implicit (S2-5) */
         int centroid = (g.wall_order & 64) && d != c;
         const double* sh = c == 0 ? m.s_u : (c == 1 ? m.s_v : m.s_w);
+        /* bit 7 of wall_order: the transverse centroid correction, cross and own direction (S2-6) */
+        if ((g.wall_order & 128) && (g.wall_order & 64)) {
+            if (f_up1 >= 0 && !solid_up && cut_ap(m, c)[f_up1] > 0.0) {
+                int qn[3] = { i, j, k }; qn[d] += 1;
+                double dl = h[d];
+                if (d != c) { dl = h[d] + (sh[3 * f_up1 + d] - sh[3 * fidx + d]); dl = fmin(fmax(dl, 0.25 * h[d]), 2.0 * h[d]); }
+                diff -= mu * g_plus * a_d * cut_transverse(g, m, c, d, fidx, f_up1, qn, 1.0, u0, ub, alpha, wall, distance, up1, sh, ubt) / dl;
+            }
+            if (f_dn1 >= 0 && !solid_dn && cut_ap(m, c)[f_dn1] > 0.0) {
+                int qn[3] = { i, j, k }; qn[d] -= 1;
+                double dl = h[d];
+                if (d != c) { dl = h[d] - (sh[3 * f_dn1 + d] - sh[3 * fidx + d]); dl = fmin(fmax(dl, 0.25 * h[d]), 2.0 * h[d]); }
+                diff += mu * g_minus * a_d * cut_transverse(g, m, c, d, fidx, f_dn1, qn, -1.0, u0, ub, alpha, wall, distance, dn1, sh, ubt) / dl;
+            }
+        }
         if (solid_up) { double kx = mu * g_plus * a_d / delta_x; wall_implicit += kx; wall_rhs += kx * up1; }
         else if (f_up1 >= 0 && centroid) {
             double dl = h[d] + (sh[3 * f_up1 + d] - sh[3 * fidx + d]);

crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/closure.rs

@@ -17,6 +17,15 @@ impl Mask {
         [t[c][3 * f], t[c][3 * f + 1], t[c][3 * f + 2]]
     }
 
+    /// The wall-distance floor in units of the smallest spacing (S2-6).
+    pub(super) fn distance_floor(&self) -> f64 {
+        if self.distance_floor_fine {
+            super::cutwall::DISTANCE_FLOOR_FINE
+        } else {
+            super::cutwall::DISTANCE_FLOOR
+        }
+    }
+
     /// The per-face shift tables (three components interleaved).
     #[must_use]
     pub fn face_shift_tables(&self) -> Option<&[Vec<f64>; 3]> {

crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/cutwall.rs

@@ -25,6 +25,8 @@ use super::wall::{FaceKind, Mask};
 pub(super) const INERTIA_FLOOR: f64 = 0.1;
 /// The wall-distance floor of a face, in units of the smallest spacing.
 pub(super) const DISTANCE_FLOOR: f64 = 0.05;
+/// The fine floor (S2-6, `Parameters::distance_floor_fine`).
+pub(super) const DISTANCE_FLOOR_FINE: f64 = 0.01;
 /// Virtual merging: a cell whose fluid fraction (at either end of the
 /// step) stays below this shares its pressure unknown with a neighbour.
 pub(super) const MERGE_FRACTION: f64 = 0.1;
@@ -211,6 +213,8 @@ impl Mask {
             density: 1.0,
             wall_order: 1,
             wall_distance_oblique: false,
+            diffusion_transverse: false,
+            distance_floor_fine: false,
             wall_exchange_axis: false,
             grad_weights: None,
             diffusion_centroid: false,
@@ -398,7 +402,8 @@ impl Mask {
         } else {
             1.0
         };
-        let distance = (phi_face + 0.5 * h[c] * (1.0 - alpha) * n_t).max(DISTANCE_FLOOR * h_min);
+        let distance =
+            (phi_face + 0.5 * h[c] * (1.0 - alpha) * n_t).max(self.distance_floor() * h_min);
         CvGeometry {
             alpha,
             ap,

crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/cut_predictor.rs

@@ -197,6 +197,54 @@ impl Solver {
                 delta.clamp(0.25 * h[d], 2.0 * h[d])
             };
             let solid = |q: [i64; 3]| mask.wall_exchange_axis && ap(c, q) == Some(0.0);
+            // S2-6: the centroids of two neighbours with different apertures
+            // are offset ACROSS the difference's axis (in the cross and in
+            // the own direction); the two-point difference then carries
+            // `∇u · Δs_⊥`, removed with the cut faces' own wall-normal
+            // gradients `(u − U_b)/d_f` (old values: explicit).
+            let transverse = |q: [i64; 3], uq: f64, sign: f64| -> f64 {
+                let sq = mask.face_shift(c, q);
+                let mut ds = [0.0; 3];
+                for e2 in 0..3 {
+                    if e2 != d {
+                        ds[e2] = sign * (sq[e2] - shift0[e2]);
+                    }
+                }
+                if ds == [0.0; 3] {
+                    return 0.0;
+                }
+                let cvq = mask.cv_geometry(c, q);
+                let ubq = mask.surface_velocity_at(body, lat.face_position(c, q), c, t_old);
+                let mut g = [0.0; 3];
+                let mut count = 0.0;
+                for (cvk, uk, ubk) in [(&cv, u0, ub), (&cvq, uq, ubq)] {
+                    let a = (cvk.wall[0] * cvk.wall[0]
+                        + cvk.wall[1] * cvk.wall[1]
+                        + cvk.wall[2] * cvk.wall[2])
+                        .sqrt();
+                    if a > 0.0 && cvk.alpha < 1.0 {
+                        let slope = (uk - ubk) / (cvk.distance * a);
+                        for e2 in 0..3 {
+                            g[e2] -= slope * cvk.wall[e2];
+                        }
+                        count += 1.0;
+                    }
+                }
+                if count == 0.0 {
+                    return 0.0;
+                }
+                (g[0] * ds[0] + g[1] * ds[1] + g[2] * ds[2]) / count
+            };
+            if mask.diffusion_transverse {
+                for (nb, sign, gap) in [(up1, 1.0, g_plus), (dn1, -1.0, g_minus)] {
+                    let q = add(p, ed, sign as i64);
+                    let Some(uq) = nb else { continue };
+                    if solid(q) || !ap(c, q).is_some_and(|a| a > 0.0) {
+                        continue;
+                    }
+                    diff -= sign * mu * gap * a_d * transverse(q, uq, sign) / spacing(q, sign);
+                }
+            }
             diff += match up1 {
                 Some(un) if solid(add(p, ed, 1)) => {
                     let k = mu * g_plus * a_d / mask.exchange_delta(&cv, d);

crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/device.rs

@@ -292,7 +292,14 @@ impl DeviceStep {
             wall_order: i32::from(self.solver.params.wall_order)
                 + 16 * i32::from(self.solver.params.wall_distance_oblique)
                 + 32 * i32::from(self.solver.params.wall_exchange_axis)
-                + 64 * i32::from(self.solver.params.diffusion_centroid),
+                + 64 * i32::from(self.solver.params.diffusion_centroid)
+                // bit 7: the transverse centroid correction; bit 8: the fine floor (S2-6).
+                + 128
+                    * i32::from(
+                        self.solver.params.diffusion_transverse
+                            && self.solver.params.diffusion_centroid,
+                    )
+                + 256 * i32::from(self.solver.params.distance_floor_fine),
             dx: g.dx,
             dy: g.dy,
             dz: g.dz,

crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/step/mod.rs

@@ -102,6 +102,18 @@ pub struct Parameters {
     /// by `½h(1 − α)(1 − |n_t|)` at every h (S2-5). `Parameters::default()`
     /// reads `RTX_E3_WALL_DISTANCE=oblique` (default: the recorded form).
     pub wall_distance_oblique: bool,
+    /// The transverse part of the centroid diffusion (S2-6): two neighbour
+    /// faces with different apertures have open-part centroids offset
+    /// ACROSS the axis of their difference, so the two-point difference
+    /// carries `∇u · Δs_⊥` — a first-order wall position on oblique walls
+    /// (≈ 0.1 h inside the body). Removed with the faces' own wall-normal
+    /// gradients `(u − U_b)/d_f`, in the cross AND the own direction.
+    /// Needs `diffusion_centroid`. `RTX_E3_DIFFUSION_TRANSVERSE=1`.
+    pub diffusion_transverse: bool,
+    /// The wall-distance floor at 0.01 of the smallest spacing instead of
+    /// 0.05: the coarse floor doubles the distance of faces with α < 0.1
+    /// (the flat wall's θ 0.95 excess, S2-6). `RTX_E3_DISTANCE_FLOOR=fine`.
+    pub distance_floor_fine: 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
@@ -145,6 +157,9 @@ impl Default for Parameters {
                 .unwrap_or(1),
             wall_distance_oblique: std::env::var("RTX_E3_WALL_DISTANCE")
                 .is_ok_and(|v| v == "oblique"),
+            diffusion_transverse: std::env::var("RTX_E3_DIFFUSION_TRANSVERSE")
+                .is_ok_and(|v| v == "1"),
+            distance_floor_fine: std::env::var("RTX_E3_DISTANCE_FLOOR").is_ok_and(|v| v == "fine"),
             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).
@@ -281,6 +296,9 @@ impl Solver {
             m.wall_order = self.params.wall_order;
             m.wall_distance_oblique = self.params.wall_distance_oblique;
             m.wall_exchange_axis = self.params.wall_exchange_axis;
+            m.diffusion_transverse =
+                self.params.diffusion_transverse && self.params.diffusion_centroid;
+            m.distance_floor_fine = self.params.distance_floor_fine;
             m.diffusion_centroid = self.params.diffusion_centroid;
             if self.params.diffusion_centroid {
                 m.compute_face_shifts();

crates/specialized/rtx-cfd/src/solvers/incompressible/embedded3/wall.rs

@@ -101,6 +101,9 @@ pub struct Mask {
     pub(super) wall_order: u8,
     /// The oblique wall distance of the cut faces (S2-5).
     pub(super) wall_distance_oblique: bool,
+    /// The transverse centroid correction and the fine distance floor (S2-6).
+    pub(super) diffusion_transverse: bool,
+    pub(super) distance_floor_fine: 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.
@@ -524,6 +527,8 @@ impl Mask {
             density: 1.0,
             wall_order: 1,
             wall_distance_oblique: false,
+            diffusion_transverse: false,
+            distance_floor_fine: false,
             wall_exchange_axis: false,
             grad_weights: None,
             diffusion_centroid: false,