Merge a454f023 (claude/jet-shapes-once) for integration review

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

src/CelerisLab/body/jets.py

@@ -207,19 +207,30 @@ def surface_link_velocity(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
                           jets=(), hits=None, *, _shifted_blocked=False):
     """Per-link prescribed surface velocity (n, 2) FP32 and aperture mask.
 
+    See _surface_links, which also returns the unit-amplitude jet shapes.
+    """
+    return _surface_links(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry, jets, hits,
+                          _shifted_blocked=_shifted_blocked)[:2]
+
+
+def _surface_links(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
+                   jets=(), hits=None, *, _shifted_blocked=False):
+    """(velocity, aperture, jet shape, jet id); the last two are None without jets.
+
     Host-only and storage-neutral: rows follow the given physical cut-link
-    order. Returns (None, None) when neither an object carries surface
+    order. Returns all None when neither an object carries surface
     actuation nor a jet is given. Jets need true FP64 ``hits`` (n, 2) and may
     not share a link with each other or with a legacy aperture.
     Shifted storage is privately enabled only by the all-blocked stepper.
     """
     enabled = [obj for obj in objects if obj.surface_actuation is not None]
     if not enabled and not jets:
-        return None, None
+        return None, None, None, None
     from .preprocess.flag_overlay import validate_cut_link_ownership
     validate_cut_link_ownership(flags, fluid_idx, dir, cfg.nq)
     surface_velocity = np.zeros((len(dir), 2), np.float32)
     surface_aperture = np.zeros(len(dir), bool)
+    shape = jet_id = None
     for obj in enabled:
         obj.validate_surface_actuation(cfg)
         selected = body_id == obj.obj_id
@@ -252,4 +263,4 @@ def surface_link_velocity(objects, cfg, flags, fluid_idx, dir, body_id, rx, ry,
             raise ValueError("Jet overlaps a SurfaceActuation aperture")
         surface_velocity[driven] = jet_link_velocity(jets, shape, jet_id)[driven]
         surface_aperture |= driven
-    return surface_velocity, surface_aperture
+    return surface_velocity, surface_aperture, shape, jet_id

src/CelerisLab/lbm/curved_links.py

@@ -94,12 +94,10 @@ class CurvedLinkSoA:
         """
         self.surface_velocity = self.aperture = None
         self.jets, self.jet_shape, self.jet_id = tuple(jets), None, None
-        self.surface_velocity, self.aperture = surface_link_velocity(
+        from ..body.jets import _surface_links
+        self.surface_velocity, self.aperture, self.jet_shape, self.jet_id = _surface_links(
             objects, cfg, flags, self.fluid_idx, self.dir, self.body_id, self.rx, self.ry,
             jets=self.jets, hits=hits)
-        if self.jets:
-            from ..body.jets import jet_link_shapes
-            self.jet_shape, self.jet_id = jet_link_shapes(self.jets, hits)
 
     def set_jet_amplitude(self, index, amplitude, stream=None):
         """Rewrite jet rows from cached shapes; re-upload only the velocity column.

src/CelerisLab/lbm/kernels/blocks.cu

@@ -168,21 +168,15 @@ extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int hal
         stream_pull_load(k, f, in + base, j);
         if (cell_kind == KIND_ACTIVE) {
             if (d_block_wall.nx != 0) {
-                const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
-                const unsigned int nbx = d_block_wall.nx / bx;
-                const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo;
-                const unsigned int gy = (blockIdx.z / nbx) * by + y - halo;
+                unsigned int gx, gy;
+                block_global_xy(x, y, halo, d_block_wall.nx, gx, gy);
                 if (gy == 1)
                     reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(
                         gx, d_block_wall.nx, d_block_wall.slot_width, d_block_wall.amplitude));
                 else if (gy == d_block_wall.ny - 2) reconstruct_top_reservoir(f, RHO);
             }
 #if COLLISION_MODEL == 2
-            // 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)));
+            collide_mrt_unforced(f);  // SRT path below unchanged.
 #else
             float rho, ux, uy, feq[NQ];
             compute_rho_u(f, rho, ux, uy);

src/CelerisLab/lbm/kernels/core/layout.cuh

@@ -53,6 +53,18 @@ __device__ __forceinline__ void index_from_thread(unsigned int& x,
     k = (unsigned long)y * NX + x;
 }
 
+// Global cell of window cell (x, y) in block blockIdx.z of a row-major,
+// full-domain block layout domain_nx cells wide. NX/NY are the local strides.
+__device__ __forceinline__ void block_global_xy(unsigned int x, unsigned int y,
+                                                unsigned int halo, unsigned int domain_nx,
+                                                unsigned int& gx, unsigned int& gy)
+{
+    const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
+    const unsigned int nbx = domain_nx / bx;
+    gx = (blockIdx.z % nbx) * bx + x - halo;
+    gy = (blockIdx.z / nbx) * by + y - halo;
+}
+
 // ---------------------------------------------------------------------------
 // Neighbor list (periodic BCs via modular arithmetic)
 // ---------------------------------------------------------------------------

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

@@ -114,6 +114,16 @@ __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,10 +37,7 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 
 __device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
 #if COLLISION_MODEL == 2
-    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)));
+    collide_mrt_unforced(f);
 #else
     float rho, ux, uy, feq[NQ];
     compute_rho_u(f, rho, ux, uy);

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

@@ -15,10 +15,8 @@ extern "C" __global__ void BlockChannelStep(
     unsigned int x, y; unsigned long k;
     index_from_thread(x, y, k);
     if (x < halo || y < halo || x >= NX - halo || y >= NY - halo) return;
-    const unsigned int bx = NX - 2 * halo, by = NY - 2 * halo;
-    const unsigned int nbx = d_block_channel.nx / bx;
-    const unsigned int gx = (blockIdx.z % nbx) * bx + x - halo;
-    const unsigned int gy = (blockIdx.z / nbx) * by + y - halo;
+    unsigned int gx, gy;
+    block_global_xy(x, y, halo, d_block_channel.nx, gx, gy);
     const unsigned long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
     const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
     float f[NQ];
@@ -55,11 +53,7 @@ extern "C" __global__ void BlockChannelStep(
         }
         if (!wall && !outlet) {
 #if COLLISION_MODEL == 2
-            // 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)));
+            collide_mrt_unforced(f);
 #else
             // Match uniform SRT arithmetic, including its shifted equilibrium path.
             float rho, ux, uy, feq[NQ], forcing[NQ] = {};

src/CelerisLab/lbm/refinement_tables.py

@@ -8,7 +8,7 @@ from dataclasses import dataclass, replace
 import numpy as np
 
 from .blocks import blocked_link_tables
-from ..body.jets import jet_link_shapes, surface_link_velocity
+from ..body.jets import _surface_links
 from .descriptors import D2Q9_EX, D2Q9_EY
 from .refinement import ACTIVE, OUTER_GHOST
 
@@ -142,11 +142,13 @@ def _prepare_blocked_boundary(boundary, layout, cfg, jets, *, _shifted_blocked=F
     tables.update(body_id=np.asarray(body_id, np.int32), rx=np.asarray(rx, np.float32),
                   ry=np.asarray(ry, np.float32), aperture=np.zeros(len(dirs), bool))
     actuated = bool(jets) or any(o.surface_actuation is not None for o in boundary.objects)
+    jet_cache = None
     if actuated:
         cfg = replace(cfg)
         cfg.validate()
-        tables["surface_velocity"], tables["aperture"] = surface_link_velocity(
+        tables["surface_velocity"], tables["aperture"], shape, jet_id = _surface_links(
             boundary.objects, cfg, boundary.flags, fluid_idx, dirs, body_id, rx, ry,
             jets=jets, hits=boundary.curved_hits, _shifted_blocked=_shifted_blocked)
-    return _BlockedBoundary(boundary, tables, jets, actuated,
-                            jet_link_shapes(jets, boundary.curved_hits) if jets else None)
+        if jets:
+            jet_cache = (shape, jet_id)
+    return _BlockedBoundary(boundary, tables, jets, actuated, jet_cache)

tests/unit/test_block_step_sync.py

@@ -1,6 +1,7 @@
 """Host contract of PeriodicBlockStepper.step(n): one stream wait per call (mocked driver)."""
 import ast
 from contextlib import nullcontext
+import importlib.util
 from pathlib import Path
 import types
 from unittest.mock import Mock
@@ -8,14 +9,15 @@ import unittest
 
 import numpy as np
 
-import CelerisLab.lbm.block_stepper as b
 from CelerisLab.config import LBMConfig
 from CelerisLab.lbm.blocks import BlockLayout2D
 
 
 def host_module():
     """block_stepper with pycuda/compiler replaced by mocks; no GPU context."""
-    path = Path(b.__file__)
+    # Locate the source without importing it: importing block_stepper loads pycuda.driver,
+    # which CPU-only import-boundary tests (test_build_contracts) must never see.
+    path = Path(importlib.util.find_spec("CelerisLab.lbm.block_stepper").origin)
     tree = ast.parse(path.read_text())
     tree.body = [node for node in tree.body if not (
         isinstance(node, ast.Import) and any(a.name == "pycuda.driver" for a in node.names)

tests/unit/test_refinement_mrt.py

@@ -37,8 +37,13 @@ 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_no_force("), 1)
+            self.assertEqual(source.count("collide_mrt_unforced(f);"), 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()