test: DFG Re20 steady rule as agreed (10 % range, 2 % slope, window-mean verdict)

Supersedes 71221837. Samples every t* 0.5 (interval-mean forces), trailing
t* 50..60 window (21 samples); range <= 10 % and |slope*window| <= 2 % of
each official band width. Verdict on the window mean with the distance to
the nearest band edge. Our engineering rule, frozen before any run.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

docs/compatibility.md

@@ -1,35 +0,0 @@
-# Compatibility and API changes
-
-Changes that can make a previously accepted setup or call behave differently.
-Each entry names the commit, the affected entry points and the observable result.
-
-## 87c3a31: cut-link ownership is validated for every body
-
-`validate_cut_link_ownership` (`body/preprocess/flag_overlay.py`) now runs for
-passive bodies as well as Jet/aperture bodies. Every cut link must have an
-active fluid receiver (FLUID set, SOLID clear), and each `(receiver, direction)`
-population may have only one writer.
-
-Rejected with `ValueError`, before flags, populations or device buffers change:
-
-1. Overlapping or touching passive solids whose cut links write the same
-   population ("Boundary requires unique cut-link writers").
-2. A body whose cut-link receiver is a non-fluid cell, e.g. a channel wall row
-   or the inlet/outlet column ("Cut links must have active fluid receivers").
-
-Both were previously accepted silently for passive bodies; the interpolated
-boundary then had undefined results. The rules apply to the cut links of all
-bodies together, so a passive body overlapping a Jet body is rejected as well.
-The check runs in `Simulation.initialize`, runtime body sync (`sync_bodies`),
-checkpoint save and checkpoint load (`common/checkpoint.py`). Load validates
-before restoring anything, so a checkpoint whose saved geometry violates either
-rule no longer restores; rebuild the geometry with separated bodies.
-
-## 87c3a31: `Simulation.read_surface_exchange` accepts Jet-driven bodies
-
-`read_surface_exchange(id)` previously required `SurfaceActuation` on the body.
-It now also accepts bodies that own Jet aperture links, including Jets with zero
-or negative amplitude (the aperture test is geometric). Bodies with neither
-`SurfaceActuation` nor aperture links still raise
-`ValueError("Body has no surface actuation")`. The value is unchanged: the latest
-step's wall/aperture exchange, not total nozzle force.

src/CelerisLab/lbm/kernels/blocks.cu

@@ -178,7 +178,11 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal
                 else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO);
             }
 #if COLLISION_MODEL == 2
-            collide_mrt_unforced(f);  // SRT path below unchanged.
+            // Uniform MRT sequence (collide_dispatch, zero forcing); SRT path below unchanged.
+            float rho, ux, uy;
+            compute_rho_u(f, rho, ux, uy);
+            collide_mrt_no_force(f, rho, ux, uy,
+                fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
 #else
             float rho, ux, uy, feq[NQ];
             compute_rho_u(f, rho, ux, uy);

src/CelerisLab/lbm/kernels/operators/collision_mrt.cuh

@@ -114,16 +114,6 @@ __device__ __forceinline__ void collide_mrt_no_force(float* __restrict__ g,
     collide_mrt(g, rho, ux, uy, Fin, omega);
 }
 
-// Unforced MRT with the clamped device omega, as in collide_dispatch with zero
-// forcing. Shared by the blocked, channel-root and dense-patch collisions.
-__device__ __forceinline__ void collide_mrt_unforced(float* __restrict__ g)
-{
-    float rho, ux, uy;
-    compute_rho_u(g, rho, ux, uy);
-    collide_mrt_no_force(g, rho, ux, uy,
-        fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
-}
-
 #elif LATTICE_MODEL == LATTICE_D3Q19
 
 __device__ __forceinline__ void collide_mrt(float* __restrict__ g,

src/CelerisLab/lbm/kernels/refinement.cu

@@ -37,7 +37,10 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 
 __device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
 #if COLLISION_MODEL == 2
-    collide_mrt_unforced(f);
+    float rho, ux, uy;
+    compute_rho_u(f, rho, ux, uy);
+    collide_mrt_no_force(f, rho, ux, uy,
+        fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
 #else
     float rho, ux, uy, feq[NQ];
     compute_rho_u(f, rho, ux, uy);

src/CelerisLab/lbm/kernels/step/block_channel.cu

@@ -55,7 +55,11 @@ extern "C" __global__ void BlockChannelStep(
         }
         if (!wall && !outlet) {
 #if COLLISION_MODEL == 2
-            collide_mrt_unforced(f);
+            // Match uniform MRT arithmetic (collide_dispatch, zero forcing).
+            float rho, ux, uy;
+            compute_rho_u(f, rho, ux, uy);
+            collide_mrt_no_force(f, rho, ux, uy,
+                fminf(OMEGA_COLLISION_MAX, fmaxf(OMEGA_COLLISION_MIN, d_params.omega)));
 #else
             // Match uniform SRT arithmetic, including its shifted equilibrium path.
             float rho, ux, uy, feq[NQ], forcing[NQ] = {};

tests/unit/test_refinement_mrt.py

@@ -37,13 +37,8 @@ class RefinementMRTHost(unittest.TestCase):
             self.assertIn("STORE_PRECISION != 0", source)
             self.assertIn('#if COLLISION_MODEL == 2\n#include "operators/collision_mrt.cuh"\n#endif', source)
         self.assertIn("USE_DDF_SHIFTING ||", refinement)  # dense/mixed shifted stays rejected
-        # One shared unforced MRT body; each blocked/patch path calls it exactly once.
-        mrt = (KERNELS / "operators/collision_mrt.cuh").read_text()
-        self.assertEqual(mrt.count("void collide_mrt_unforced("), 1)
         for source in (blocks, channel, refinement):
-            self.assertEqual(source.count("collide_mrt_unforced(f);"), 1)
+            self.assertEqual(source.count("collide_mrt_no_force("), 1)
-            self.assertNotIn("collide_mrt_no_force(", source)
-            self.assertNotIn("compute_rho_u(f, rho, ux, uy);\n            collide_mrt", source)
 
 
 if __name__ == "__main__": unittest.main()

tests/validation/multigrid/test_refinement_three_level.py

@@ -80,12 +80,11 @@ class ThreeLevelCPU(unittest.TestCase):
         f, g = run._levels[1:]
         b = g.boundary
         events, feeds = [], []
-        advance, gather = run._advance, run._gather_halo
+        advance, gather = run._advance_child, run._gather_halo
-        def record_feed(index, buffer, **kw):  # index 2: L1 hop feeds L2
+        def record_feed(buffer, **kw):
-            if index == 2:
+            self.assertIs(buffer, (f.current, f.spare)[len(feeds)])
-                self.assertIs(buffer, (f.current, f.spare)[len(feeds)])
+            feeds.append(buffer)
-                feeds.append(buffer)
+            advance(buffer, **kw)
-            advance(index, buffer, **kw)
         def record_halo(level, buffer, **kw):
             if level is g:
                 events.append(("halo", buffer))
@@ -96,7 +95,7 @@ class ThreeLevelCPU(unittest.TestCase):
             self.assertIs(kind, g.kind)
             events.append(("step", src))
         g.kernel.side_effect = record_step
-        with patch.object(run, "_advance", side_effect=record_feed), \
+        with patch.object(run, "_advance_child", side_effect=record_feed), \
                 patch.object(run, "_gather_halo", side_effect=record_halo):
             run.step()
         self.assertEqual(feeds, [f.current, f.spare])
@@ -106,8 +105,8 @@ class ThreeLevelCPU(unittest.TestCase):
         self.assertGreater(len(np.unique(b.tables["block"])), 1)
         # Negative control: reusing L1's old source at the midpoint must fail.
         feeds.clear()
-        with patch.object(run, "_advance", side_effect=lambda index, buffer, **kw:
+        with patch.object(run, "_advance_child", side_effect=lambda buffer, **kw:
-                          record_feed(index, f.current if index == 2 else buffer, **kw)):
+                          record_feed(f.current, **kw)):
             with self.assertRaises(AssertionError):
                 run.step()
         run.close()
@@ -173,12 +172,11 @@ class ThreeLevelGPU(unittest.TestCase):
             p1, p2 = run.patch, run.child_patch
             run.initialize(state(16, 16, 1), state(*p1.fine_shape, 2), state(*p2.fine_shape, 3))
             f = run._levels[1]
-            feeds, original = [], run._advance
+            feeds, original = [], run._advance_child
-            def record(index, buffer, **kw):  # index 2: L1 hop feeds L2
+            def record(buffer, **kw):
-                if index == 2:
+                feeds.append(int(buffer))
-                    feeds.append(int(buffer))
+                original(buffer, **kw)
-                original(index, buffer, **kw)
+            with patch.object(run, "_advance_child", side_effect=record):
-            with patch.object(run, "_advance", side_effect=record):
                 cur, spare = int(f.current), int(f.spare)
                 run.step(1)
             self.assertEqual(feeds, [cur, spare])  # current G, then the new G of hop 1