From 0f8a428f074347ddfc7b65a3a83e8fde9e5356cf Mon Sep 17 00:00:00 2001 From: Claude Date: Fri, 30 Oct 2026 07:35:02 +0000 Subject: [PATCH] lbm: fixed-centre moving-circle cut links (S2a) MovingCircleLinks computes the cut links of each translating circle from an FP64 {cx, cy, r} table over its swept band (same q/hit/clamp/fallback order as build_cut_links, explicit _rn operations) and applies the unchanged apply_bouzidi_link with spin wall velocity. obs keeps its contract; a separate per-body diagnostics buffer records CME, the Wen term u_w*dm, link_dm and the link count. Moving tables are built at initialize and attached to the stepper, which then launches the kernel before the static curved kernel; until then direct and run steps stay blocked. Displacement is disabled: set_body offset/vx/vy on a translating circle raises before mutation. Contract section 10a documents S2a and the S3 pending rule. Co-Authored-By: Claude Opus 5.5 --- docs/moving_bodies.md | 25 +++++ src/CelerisLab/body/motion.py | 53 +++++++++ .../lbm/kernels/step/moving_circle.cu | 102 ++++++++++++++++++ src/CelerisLab/lbm/kernels/uniform.cu | 1 + src/CelerisLab/lbm/stepper.py | 13 +++ src/CelerisLab/simulation.py | 56 +++++++--- tests/unit/test_moving_bodies.py | 61 +++++++++-- tests/unit/test_state_regressions.py | 2 +- 8 files changed, 287 insertions(+), 26 deletions(-) create mode 100644 src/CelerisLab/lbm/kernels/step/moving_circle.cu diff --git a/docs/moving_bodies.md b/docs/moving_bodies.md index aa64a63..c45cb9c 100644 --- a/docs/moving_bodies.md +++ b/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 diff --git a/src/CelerisLab/body/motion.py b/src/CelerisLab/body/motion.py index 9746207..134e4fb 100644 --- a/src/CelerisLab/body/motion.py +++ b/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 diff --git a/src/CelerisLab/lbm/kernels/step/moving_circle.cu b/src/CelerisLab/lbm/kernels/step/moving_circle.cu new file mode 100644 index 0000000..c58fe2d --- /dev/null +++ b/src/CelerisLab/lbm/kernels/step/moving_circle.cu @@ -0,0 +1,102 @@ +// 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 fallback rule as build_cut_links (q < 1/2 needs a fluid donor at x - c). + unsigned char fallback = CURVED_FALLBACK_BOUZIDI; + if (q < 0.5f) { + 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 diff --git a/src/CelerisLab/lbm/kernels/uniform.cu b/src/CelerisLab/lbm/kernels/uniform.cu index 051dd18..6304732 100644 --- a/src/CelerisLab/lbm/kernels/uniform.cu +++ b/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" diff --git a/src/CelerisLab/lbm/stepper.py b/src/CelerisLab/lbm/stepper.py index facfe54..15cc16a 100644 --- a/src/CelerisLab/lbm/stepper.py +++ b/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: diff --git a/src/CelerisLab/simulation.py b/src/CelerisLab/simulation.py index fc8f7d1..1bb90d4 100644 --- a/src/CelerisLab/simulation.py +++ b/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,29 @@ 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]) + + 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 +353,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 +381,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 +537,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 +564,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 +581,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 +699,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 +717,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 +832,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 diff --git a/tests/unit/test_moving_bodies.py b/tests/unit/test_moving_bodies.py index f5b2814..c757707 100644 --- a/tests/unit/test_moving_bodies.py +++ b/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,44 @@ 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 + + +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) + field = SimpleNamespace(ddf_gpu="ddf", temp_gpu="temp", flag_gpu="flag", n_curved=0, + 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", "OneStep"] * 2 and stepper.step_count == 2 + 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) diff --git a/tests/unit/test_state_regressions.py b/tests/unit/test_state_regressions.py index f7b1ff7..4c1df44 100644 --- a/tests/unit/test_state_regressions.py +++ b/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":