simulation: keep the owner context active through moving diagnostics readback

moving_diagnostics is context_bound, so stream synchronization and the
DTOH copy run inside one owner activation. Tests: owner stays active
through sync and download; launch order moving -> static curved -> OneStep
with a real static link count and the buffer swap between steps.

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

docs/moving_bodies.md

@@ -338,6 +338,31 @@ recorded as not runnable rather than replaced by another test.
   run with a moving body raises before upload/launch; stationary unaffected). No kernels, no GPU runs, no
   y-only VIV oscillator (S4), no force RHS.
 
+## 10a. S2a — fixed-centre moving links (implemented on `claude/moving-bodies-s2a`)
+
+- `MovingCircleLinks` (`kernels/step/moving_circle.cu`, dense double buffer)
+  runs before the static curved kernel when moving tables are attached: one
+  thread per (band cell, direction) slot, FP64 `{cx, cy, r}` at the applied
+  offset, q/hit/clamp/fallback in the order of `build_cut_links` with explicit
+  `_rn` operations (host/device bit parity is a test target, not a claim),
+  then the unchanged `apply_bouzidi_link` with spin wall velocity only.
+- `obs` keeps its contract (CME and torque from `apply_bouzidi_link`). A
+  separate per-body diagnostics buffer receives `cme`, the Wen term
+  `u_w * dm`, `link_dm` and the link count from the same link outputs; it is
+  zeroed with obs and only exists when moving circles do. Optional
+  `slot_debug` (q, fallback, rx, ry per slot) is null in production.
+- Displacement is disabled: `set_body(offset=..., vx=..., vy=...)` on a
+  translating circle raises before mutation, before and after initialize, so
+  no pending target can exist and `run`/`step`/direct `LBMStepper.step` all
+  hold at the applied geometry. `initialize()` never reads or consumes
+  `pending`. `stepper.blocked` is cleared only when the moving tables are
+  attached (initialize, stepper recreation).
+- S3 rule (not implemented): once targets are re-enabled, any pending target
+  must raise in `run`/`step` and in `LBMStepper.step` alike before any launch,
+  swap or step-count change; zero-step calls never consume it.
+- Not validation: a pass of the zero-motion gates says nothing about moving
+  geometry. GPU parity (G0a/G0b) is unverified until run on a free device.
+
 ## 11. Remaining decisions
 
 Structural RHS (section 6); motion-record layout; refill details; coupling

src/CelerisLab/body/motion.py

@@ -166,3 +166,56 @@ def enumerate_donor_coverage(radius: float, positions: int = 8, directions: int
                 c_new = (c_old[0] + MAX_STEP * math.cos(a), c_old[1] + MAX_STEP * math.sin(a))
                 count += len(empty_donor_cells(c_old, c_new, radius))
     return count
+
+
+class MovingTables:
+    """Device tables of the translating circles for ``MovingCircleLinks`` (fixed centre).
+
+    Immutable per run: band table (x0, y0, width, height, first_slot, body_id) and
+    FP64 geometry (cx, cy, r) at the applied offset; plus the separate per-body
+    diagnostics (cme_x, cme_y, wen_x, wen_y, link_dm, n_links), never part of obs.
+    """
+    BAND_INTS, DIAG_FLOATS = 6, 6
+
+    def __init__(self, band, geom):
+        import numpy as np
+        import pycuda.driver as cuda
+        self.n_bodies = len(geom) // 3
+        self.n_slots = int(band[-2] + band[-4] * band[-3] * 8)
+        self.band_gpu = cuda.to_device(np.asarray(band, np.int32))
+        self.geom_gpu = cuda.to_device(np.asarray(geom, np.float64))
+        self.diag_gpu = cuda.to_device(np.zeros(self.n_bodies * self.DIAG_FLOATS, np.float32))
+        self.slot_debug_gpu = None  # optional (q, fallback, rx, ry) per slot for parity tests
+
+    @classmethod
+    def upload(cls, circles):
+        """``circles``: [(MovingCircle, body_id)] in body order."""
+        return cls(*cls.layout(circles))
+
+    @staticmethod
+    def layout(circles):
+        """Host band and geometry tables (no device work)."""
+        band, geom, first = [], [], 0
+        for motion, body_id in circles:
+            x0, x1, y0, y1 = motion.band()
+            width, height = x1 - x0 + 1, y1 - y0 + 1
+            band += [x0, y0, width, height, first, int(body_id)]
+            first += width * height * 8
+            geom += [*motion.centre(motion.applied), motion.radius]
+        return band, geom
+
+    def enable_slot_debug(self):
+        import numpy as np
+        import pycuda.driver as cuda
+        self.slot_debug_gpu = cuda.to_device(np.zeros(self.n_slots * 4, np.float32))
+
+    def zero_diagnostics(self, stream):
+        import pycuda.driver as cuda
+        cuda.memset_d32_async(self.diag_gpu, 0, self.n_bodies * self.DIAG_FLOATS, stream)
+
+    def download_diagnostics(self):
+        import numpy as np
+        import pycuda.driver as cuda
+        out = np.empty((self.n_bodies, self.DIAG_FLOATS), np.float32)
+        cuda.memcpy_dtoh(out, self.diag_gpu)
+        return out

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

@@ -0,0 +1,104 @@
+// Translating-circle cut links, fixed centre (S2a; docs/moving_bodies.md, M2).
+// One thread per (swept-band cell, direction) slot. Geometry comes from an FP64
+// {cx, cy, r} table per body; the operation order mirrors
+// CircleGeometry.build_cut_links / find_circle_intersection with explicit _rn
+// operations (bit parity with the host builder is a test target, not a claim).
+// Moving circles are never in the packed static links, so this kernel is the
+// only writer of their solid slots.
+#ifndef CELERIS_MOVING_CIRCLE_CU
+#define CELERIS_MOVING_CIRCLE_CU
+
+#if DIM == 2
+constexpr int MOVING_BAND_INTS = 6;   // x0, y0, width, height, first_slot, body_id
+constexpr int MOVING_DIAG_FLOATS = 6; // cme_x, cme_y, wen_x, wen_y, link_dm, n_links
+
+__device__ __forceinline__ double mc_sq(double v) { return __dmul_rn(v, v); }
+
+__device__ __forceinline__ bool mc_inside(int x, int y, double cx, double cy, double r)
+{
+    return __dadd_rn(mc_sq(__dsub_rn((double)x, cx)), mc_sq(__dsub_rn((double)y, cy))) < mc_sq(r);
+}
+
+// Smallest root t in [0, 1] of |p + t c - centre| = r (find_circle_intersection order).
+__device__ __forceinline__ bool mc_hit(int x, int y, int ex, int ey,
+                                       double cx, double cy, double r, double& t)
+{
+    const double dx = (double)ex, dy = (double)ey;
+    const double px = __dsub_rn((double)x, cx), py = __dsub_rn((double)y, cy);
+    const double a = __dadd_rn(mc_sq(dx), mc_sq(dy));
+    const double b = __dmul_rn(2.0, __dadd_rn(__dmul_rn(dx, px), __dmul_rn(dy, py)));
+    const double c = __dsub_rn(__dadd_rn(mc_sq(px), mc_sq(py)), mc_sq(r));
+    const double det = __dsub_rn(mc_sq(b), __dmul_rn(__dmul_rn(4.0, a), c));
+    if (det < 0.0) return false;
+    const double s = __dsqrt_rn(det), a2 = __dmul_rn(2.0, a);
+    const double t1 = __ddiv_rn(__dsub_rn(-b, s), a2);
+    const double t2 = __ddiv_rn(__dadd_rn(-b, s), a2);
+    if (0.0 <= t1 && t1 <= 1.0) { t = t1; return true; }
+    if (0.0 <= t2 && t2 <= 1.0) { t = t2; return true; }
+    return false;
+}
+
+__global__ void MovingCircleLinks(
+    fpxx* fi, const uint16_t* flag,
+    const int* band, const double* geom,
+    unsigned int n_bodies, unsigned int n_slots,
+    const float* action, float* obs, float* diag, float* slot_debug)
+{
+    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
+    if (tid >= n_slots) return;
+    unsigned int m = 0;
+    while (m + 1 < n_bodies && tid >= (unsigned int)band[(m + 1) * MOVING_BAND_INTS + 4]) ++m;
+    const int* b = band + m * MOVING_BAND_INTS;
+    const unsigned int local = tid - (unsigned int)b[4];
+    const unsigned int dir = local % 8u + 1u;
+    const int x = b[0] + (int)((local / 8u) % (unsigned int)b[2]);
+    const int y = b[1] + (int)((local / 8u) / (unsigned int)b[2]);
+    const double cx = geom[3 * m], cy = geom[3 * m + 1], r = geom[3 * m + 2];
+    const int ex = d_cx[dir], ey = d_cy[dir];
+    float* debug = slot_debug == nullptr ? nullptr : slot_debug + 4 * tid;
+    if (debug != nullptr) debug[0] = -1.0f;
+
+    double t;
+    if (mc_inside(x, y, cx, cy, r) || !mc_inside(x + ex, y + ey, cx, cy, r)
+            || !mc_hit(x, y, ex, ey, cx, cy, r, t))
+        return;
+    const double hx = __dadd_rn((double)x, __dmul_rn(t, (double)ex));
+    const double hy = __dadd_rn((double)y, __dmul_rn(t, (double)ey));
+    const double hdx = __dsub_rn(hx, (double)x), hdy = __dsub_rn(hy, (double)y);
+    const double norm = __dsqrt_rn((double)(ex * ex + ey * ey));
+    double qd = __ddiv_rn(__dsqrt_rn(__dadd_rn(mc_sq(hdx), mc_sq(hdy))), norm);
+    qd = fmax(0.01, fmin(qd, 1.0));
+    const float q = (float)qd;
+    const float rx = (float)__dsub_rn(hx, cx), ry = (float)__dsub_rn(hy, cy);
+
+    // Same two-stage order as build_cut_links + pack_cut_links_to_soa: the fallback
+    // class is chosen from the clamped FP64 q (q < 1/2 needs a fluid donor at x - c);
+    // only then are q, rx, ry rounded to FP32 for the BFL branch and spin velocity.
+    unsigned char fallback = CURVED_FALLBACK_BOUZIDI;
+    if (qd < 0.5) {
+        const int xd = x - ex, yd = y - ey;
+        if (xd < 0 || xd >= (int)NX || yd < 0 || yd >= (int)NY || mc_inside(xd, yd, cx, cy, r)
+                || (flag[linear_index((unsigned int)xd, (unsigned int)yd)] & FLAG_FLUID) == 0)
+            fallback = CURVED_FALLBACK_HALFWAY;
+    }
+    if (debug != nullptr) { debug[0] = q; debug[1] = (float)fallback; debug[2] = rx; debug[3] = ry; }
+
+    const int bid = b[5];
+    const float omega = action_omega(action, bid);
+    const float Uw = -omega * ry, Vw = omega * rx;  // fixed centre: no translation term
+    float exchange[3], dm;
+    apply_bouzidi_link(dir, q, linear_index((unsigned int)x, (unsigned int)y), fi,
+                       Uw, Vw, rx, ry, fallback, obs, bid, exchange, &dm);
+    if (diag != nullptr) {
+        float* d = diag + m * MOVING_DIAG_FLOATS;
+        atomicAdd(&d[0], exchange[0]);
+        atomicAdd(&d[1], exchange[1]);
+        atomicAdd(&d[2], Uw * dm);   // Wen 2014 correction -u_w (f_t - f_r)
+        atomicAdd(&d[3], Vw * dm);
+        atomicAdd(&d[4], dm);
+        atomicAdd(&d[5], 1.0f);
+    }
+}
+#endif  // DIM == 2
+
+#endif  // CELERIS_MOVING_CIRCLE_CU

src/CelerisLab/lbm/kernels/uniform.cu

@@ -71,6 +71,7 @@ extern "C"
 #include "step/one_step_double.cu"
 #include "step/one_step_esopull.cu"
 #include "step/aux_kernels.cu"
+#include "step/moving_circle.cu"
 #include "step/esopull_macro.cu"
 #include "step/esopull_sync.cu"
 

src/CelerisLab/lbm/stepper.py

@@ -23,6 +23,7 @@ class LBMStepper:
         # Set by the owner when a step would be wrong (translating circles
         # excluded from the packed links before the moving kernels exist).
         self.blocked: str | None = None
+        self.moving = None  # MovingTables attached by the owner (translating circles)
 
         if self._esopull:
             self.step_fn = module.get_function("EsoPullStep")
@@ -38,6 +39,7 @@ class LBMStepper:
         self.apply_esopull_curved_fn = module.get_function("ApplyEsoPullCurvedKernel")
         self.sensor_fn = module.get_function("SensorKernel")
         self.force_region_fn = module.get_function("ForceRegionKernel")
+        self.moving_fn = None if self._esopull else module.get_function("MovingCircleLinks")
 
         tpb = cfg.threads_per_block
         self.block = (tpb, 1, 1)
@@ -86,6 +88,8 @@ class LBMStepper:
                     f.ddf_gpu, self.boundary_values_gpu, np.uint64(self._step_count),
                     block=self.block, grid=self.boundary_grid, **launch_kw,
                 )
+            if self.moving is not None:
+                self._launch_moving(action_gpu, obs_gpu, **launch_kw)
             self._launch_curved(action_gpu, obs_gpu, **launch_kw)
 
             if self._esopull:
@@ -111,6 +115,15 @@ class LBMStepper:
             # Any completed GPU step invalidates host DDF mirror.
             f.invalidate_host_ddf_cache()
 
+    def _launch_moving(self, action_gpu, obs_gpu, **launch_kw):
+        m, tpb = self.moving, self.cfg.threads_per_block
+        self.moving_fn(
+            self.field.ddf_gpu, self.field.flag_gpu, m.band_gpu, m.geom_gpu,
+            np.uint32(m.n_bodies), np.uint32(m.n_slots), action_gpu, obs_gpu,
+            m.diag_gpu, m.slot_debug_gpu if m.slot_debug_gpu is not None else np.uintp(0),
+            block=(tpb, 1, 1), grid=((m.n_slots + tpb - 1) // tpb, 1, 1), **launch_kw,
+        )
+
     def _launch_curved(self, action_gpu, obs_gpu, **launch_kw):
         f = self.field
         if f.n_curved == 0:

src/CelerisLab/simulation.py

@@ -28,8 +28,9 @@ from .cuda.context import CudaContext, context_bound
 from .cuda import compiler
 from .lbm.descriptors import FLUID
 
-# S1: translating circles are declared and validated, but no moving kernel exists yet.
-MOVING_BLOCKED = ("translating circles need the moving kernels (not implemented); "
+# Translating circles step only through their moving tables (fixed centre in S2a).
+MOVING_BLOCKED = ("translating circles need the moving kernels and their tables (attached "
+                  "at initialize); "
                   "positive-step runs are rejected")
 from .lbm.field import LBMField
 from .lbm.stepper import LBMStepper
@@ -212,6 +213,30 @@ class Simulation:
         # Device geometry is unknown after a failed moving step: same lifecycle as Jets.
         self._failed = f"moving circle step failed: {exc!r}"
 
+    def _attach_moving(self) -> None:
+        # Direct stepper calls stay blocked until the moving tables exist.
+        tables = getattr(self, "_moving_tables", None)
+        self.stepper.moving = tables
+        self.stepper.blocked = MOVING_BLOCKED if self._moving_bodies() and tables is None else None
+
+    def _build_moving_tables(self):
+        moving = self._moving_bodies()
+        if not moving:
+            return None
+        from .body.motion import MovingTables
+        return MovingTables.upload([(obj.motion, obj.obj_id) for obj in moving])
+
+    @context_bound
+    def moving_diagnostics(self) -> np.ndarray:
+        """Per translating circle (in body order): cme_x, cme_y, wen_x, wen_y, link_dm,
+        n_links accumulated since obs were last zeroed (diagnostics only)."""
+        self._require_usable()
+        tables = self.stepper.moving
+        if tables is None:
+            return np.zeros((0, 6), np.float32)
+        self.stream.synchronize()
+        return tables.download_diagnostics()
+
     def _moving_bodies(self):
         return [obj for obj in self.bodies.objects if obj.motion is not None]
 
@@ -329,18 +354,17 @@ class Simulation:
             omega: rotation speed (lattice units).  ``None`` = no change.
             vx: translation velocity x (future).  ``None`` = no change.
             vy: translation velocity y (future).  ``None`` = no change.
-            offset: target centre offset of a translating circle; with vx/vy
-                (default 0) it replaces the pending target after validation
-                against the applied offset (docs/moving_bodies.md, 4a).
+            offset: target centre offset of a translating circle. Displacement
+                is disabled until the moving cell update exists (S3): any offset
+                raises before mutation; translating circles hold at the applied
+                offset and accept omega only.
         """
         obj = self.bodies.get(id)
         motion = obj.motion
         if motion is not None:
-            if offset is None and (vx is not None or vy is not None):
-                raise ValueError("a translating circle takes vx/vy only with an offset")
-            target = None if offset is None else motion.check_target(
-                offset, (0.0 if vx is None else vx, 0.0 if vy is None else vy))
-            vx = vy = None  # the motion target, not ObjectState, owns translation
+            if offset is not None or vx is not None or vy is not None:
+                raise ValueError("translating-circle displacement is disabled until S3; "
+                                 "only omega may be set")
         elif offset is not None:
             raise ValueError(f"body {id} was not declared with a translation")
         proposed = dict(vars(obj.state))
@@ -358,8 +382,6 @@ class Simulation:
         if vy is not None:
             obj.state.vy = float(vy)
             changed = True
-        if motion is not None and target is not None:
-            motion.pending = target
         if changed:
             self.bodies._refresh_action_from_objects()
 
@@ -516,7 +538,7 @@ class Simulation:
             self.field, self._module, self.lbm_cfg,
         )
         self.stepper._step_count = _prev_step_count
-        self.stepper.blocked = MOVING_BLOCKED if self._moving_bodies() else None
+        self._attach_moving()
 
         self._assert_object_count_contract(expected_count=expected_n_objects)
 
@@ -543,7 +565,7 @@ class Simulation:
         # change. Compact-list construction publishes no outputs on failure.
         next_flags = self.bodies.build_flags(self.field.build_channel_flags())
         self._validate_moving_bodies()
-        self.stepper.blocked = MOVING_BLOCKED if self._moving_bodies() else None
+        self._attach_moving()
         if self.bodies.jets and self.bodies.count == 0:
             raise ValueError("Jets need bodies with cut links")
         self.bodies.build_compact_lists_for(self.bodies.objects, next_flags)
@@ -560,6 +582,8 @@ class Simulation:
             self.bodies.sync_to_gpu(self.field, rebuild_flags=False)
         else:
             self.bodies.ensure_packed_buffers()
+        self._moving_tables = self._build_moving_tables()
+        self._attach_moving()
         self._assert_runtime_contracts()
         self._initialized = True
 
@@ -676,7 +700,7 @@ class Simulation:
         checkpoint_interval = operator.index(checkpoint_interval)
         if steps < 0 or checkpoint_interval < 0:
             raise ValueError("steps and checkpoint_interval must be nonnegative")
-        if steps > 0 and self._moving_bodies():
+        if steps > 0 and self._moving_bodies() and self.stepper.moving is None:
             raise RuntimeError(MOVING_BLOCKED)
         self._assert_runtime_contracts()
         # Discard any uncommitted body edits before stepping.
@@ -694,6 +718,8 @@ class Simulation:
         # Zero obs segments before step group
         if zero_obs:
             self.bodies.zero_obs_async(stream)
+            if self.stepper.moving is not None:
+                self.stepper.moving.zero_diagnostics(stream)
 
         if checkpoint_interval > 0:
             done = 0
@@ -807,6 +833,7 @@ class Simulation:
                 cuda.Context.synchronize()
                 # Release owned GPU resources while their context is current.
                 self.stepper = None
+                self._moving_tables = None
                 self.bodies = None
                 self.field = None
                 self._module = None

tests/unit/test_moving_bodies.py

@@ -117,7 +117,8 @@ def moving_simulation(monkeypatch, *, moving=True):
                                  zero_obs_async=lambda s: writes.append("obs"),
                                  increment_obs_steps=lambda n: writes.append("count"),
                                  download_obs_full_async=lambda s: writes.append("download"))
-    sim.stepper = SimpleNamespace(step=lambda n, **k: writes.append(("step", n)), blocked=None)
+    sim.stepper = SimpleNamespace(step=lambda n, **k: writes.append(("step", n)), blocked=None,
+                                  moving=None)
     sim._assert_runtime_contracts = lambda: None
     return sim, writes
 
@@ -217,17 +218,19 @@ def test_add_body_and_set_body_route_targets(monkeypatch):
     still = sim.add_body("circle", (100., 40.), 5.)
     motion = sim.bodies.get(bid).motion
     assert sim.stepper.blocked is not None  # declaring a moving circle blocks direct steps
-    sim.set_body(bid, offset=(0., .3), vy=.05)
-    sim.set_body(bid, offset=(0., .4), vy=.05)  # last wins
-    for kwargs in (dict(offset=(0., .6)), dict(vy=.05), dict(vx=.05),
-                   dict(offset=(0., .1), vy=.2), dict(offset=(0., 0.), vy=.01)):
-        with pytest.raises(ValueError):
-            sim.set_body(bid, **kwargs)
-        assert motion.pending == ((0., .4), (0., .05))  # invalid input keeps the target
-    sim.set_body(bid, omega=.01)  # omega-only keeps the pending target
+    for initialized in (False, True):  # S2a: displacement disabled before and after initialize
+        sim._initialized = initialized
+        for kwargs in (dict(offset=(0., .3), vy=.05), dict(offset=(0., 0.)), dict(vy=.05),
+                       dict(vx=.05), dict(offset=(0., .1), omega=.01)):
+            with pytest.raises(ValueError, match="disabled until S3"):
+                sim.set_body(bid, **kwargs)
+            assert motion.pending is None and motion.applied == (0., 0.)
+    sim._initialized = False
     obj = sim.bodies.get(bid)
+    assert obj.state.omega == 0.
+    sim.set_body(bid, omega=.01)  # spin on hold stays allowed
     assert obj.state.omega == .01 and (obj.state.vx, obj.state.vy) == (0., 0.)
-    assert motion.pending == ((0., .4), (0., .05)) and motion.applied == (0., 0.)
+    assert motion.pending is None
     with pytest.raises(ValueError, match="not declared"):
         sim.set_body(still, offset=(0., .1))
     sim.set_body(still, vx=.01)  # stationary API unchanged
@@ -240,6 +243,84 @@ def test_add_body_and_set_body_route_targets(monkeypatch):
                          **kwargs)
 
 
+def test_moving_tables_layout_and_attachment(monkeypatch):
+    from CelerisLab.body.motion import MovingTables
+    a = MovingCircle((50., 40.), 10., Translation((-1., 1., -3., 3.), .1))
+    b = MovingCircle((90.25, 30.5), 5., Translation((0., 0., 0., 0.), .1))
+    band, geom = MovingTables.layout([(a, 2), (b, 5)])
+    wa, ha = 63 - 37 + 1, 55 - 25 + 1
+    assert band == [37, 25, wa, ha, 0, 2, 83, 23, 16, 16, wa * ha * 8, 5]
+    assert geom == [50., 40., 10., 90.25, 30.5, 5.]
+    sim = routed_simulation(monkeypatch)
+    sim.add_body("circle", (50., 40.), 10., translation=Translation((0., 0., 0., 0.), .1))
+    sim._moving_tables = "tables"
+    sim._attach_moving()
+    assert sim.stepper.moving == "tables" and sim.stepper.blocked is None
+    sim._moving_tables = None
+    sim._attach_moving()
+    assert sim.stepper.moving is None and sim.stepper.blocked is not None
+
+
+class CountingOwner(Owner):
+    """Owner double that counts nested activations like CudaContext."""
+    def __init__(self):
+        super().__init__()
+        self.active = 0
+
+    def activate(self):
+        from contextlib import contextmanager
+
+        @contextmanager
+        def scope():
+            self.active += 1
+            try:
+                yield self
+            finally:
+                self.active -= 1
+        return scope()
+
+
+def test_public_diagnostics_keep_owner_active_through_sync_and_download(monkeypatch):
+    sim = make_simulation(monkeypatch)
+    sim.ctx, sim._failed = CountingOwner(), None
+    seen = []
+    sim._stream = SimpleNamespace(synchronize=lambda: seen.append(("sync", sim.ctx.active)))
+    expected = np.arange(6, dtype=np.float32).reshape(1, 6)
+
+    def download():
+        seen.append(("download", sim.ctx.active))
+        return expected
+    sim.stepper = SimpleNamespace(moving=SimpleNamespace(download_diagnostics=download))
+    assert sim.moving_diagnostics() is expected
+    assert [name for name, _ in seen] == ["sync", "download"]
+    assert all(active > 0 for _, active in seen) and sim.ctx.active == 0
+
+
+def test_stepper_launches_moving_links_before_static_curved():
+    from CelerisLab.lbm.stepper import LBMStepper
+    launches = []
+    module = SimpleNamespace(get_function=lambda name: lambda *a, **k: launches.append((name, a, k)))
+    cfg = SimpleNamespace(streaming="double_buffer", threads_per_block=128, is_d2q9=True,
+                          nx=64, ny=32, nz=1)
+    curved = SimpleNamespace(**{f"{name}_gpu": name for name in (
+        "fluid_idx", "dir", "q", "rx", "ry", "rz", "fallback_class", "body_id")},
+        surface_velocity=None, surface_velocity_gpu=None, link_exchange_gpu=None)
+    field = SimpleNamespace(ddf_gpu="ddf", temp_gpu="temp", flag_gpu="flag", n_curved=5,
+                            curved=curved, invalidate_host_ddf_cache=lambda: None)
+    stepper = LBMStepper(field, module, cfg)
+    stepper._launch_force_region = stepper._launch_sensor = lambda *a, **k: None
+    stepper.moving = SimpleNamespace(band_gpu="band", geom_gpu="geom", diag_gpu="diag",
+                                     slot_debug_gpu=None, n_bodies=1, n_slots=300)
+    stepper.step(2, action_gpu="act", obs_gpu="obs")
+    names = [name for name, _, _ in launches]
+    assert names == ["MovingCircleLinks", "CurvedBoundaryKernel", "OneStep"] * 2
+    assert stepper.step_count == 2
+    assert launches[1][1][0] == "ddf" and launches[3][1][0] == "temp"  # swap between steps
+    args, kw = launches[0][1], launches[0][2]
+    assert args[:4] == ("ddf", "flag", "band", "geom") and int(args[5]) == 300
+    assert kw["grid"] == (3, 1, 1) and kw["block"] == (128, 1, 1)
+
+
 def test_recompile_installs_a_guarded_stepper(monkeypatch):
     from CelerisLab.lbm.stepper import LBMStepper
     sim = routed_simulation(monkeypatch)

tests/unit/test_state_regressions.py

@@ -297,7 +297,7 @@ def test_passive_overlap_failure_keeps_live_state(device, entry, second_x):
                           stream=SimpleNamespace(synchronize=device.synchronize),
                           stepper=SimpleNamespace(blocked=None))
     for name in ("_require_usable", "_moving_bodies", "_reject_moving_bodies",
-                 "_validate_moving_bodies"):
+                 "_validate_moving_bodies", "_attach_moving"):
         setattr(sim, name, getattr(Simulation, name).__get__(sim))
     with pytest.raises(ValueError, match="active fluid receivers|unique cut-link writers"):
         if entry == "sync":