Commit Graph
17 Commits
Author SHA1 Message Date
ClaudeandClaude Opus 5.5 aebef52345 Wire Jets into FixedPatchStepper fine level (dense and blocked)
FixedPatchStepper(..., jets=()) passes jets (coarse cell-edge frame, as
fine_objects) to prepare_fine_boundary, which runs the shared assignment
checks before CUDA. Dense fine path: assign_surface with jets and the
boundary's true hits. Blocked path: surface_link_velocity with jets/hits;
unit-amplitude shapes cached. set_jet_amplitude (new/ready only) validates
first, then rewrites the dense SoA column or the blocked jet rows and
re-uploads only the velocity buffer after a stream sync. Scheduling,
kernels and launch ABI unchanged.

tests/validation/multigrid/test_refinement_jets.py: mocked-driver CPU
checks and frozen GPU gates G1-G3 (layout identity, jet vs legacy <=1e-6,
amplitude switch).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 08:42:51 +00:00
Frank14f 906e0d9a0e Support blocked parent storage in fixed patch stepping 2026-10-15 16:25:26 +08:00
ClaudeandClaude Opus 5.5 5c37d75dc4 Validate physical fine_cfg before surface_link_velocity
The blocked actuated path called the helper with an unvalidated fine
config (nq=0), so receiver*nq+dir collapsed distinct writers and the
unique-writer guard fired. Validate after wall-aware sanitization, as the
dense path does before assign_surface; the test's fine_cfg helper too.
Unique-writer guard unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 04:49:34 +00:00
ClaudeandClaude Opus 5.5 1a274817cd Test: Astra's actuated seam fixture (angle 0.7, half-width 0.85)
Asserts 21 aperture links, all with nonzero velocity, and 3 halo-writer
plus 2 halo-donor aperture rows in each decomposed layout (actuation
coverage across seams, not only geometric cut-link coverage). Velocity
is given in coarse units; prepare_fine_boundary scales it once. The GPU
gate remains storage equivalence; no zero-mass-change claim.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 04:45:57 +00:00
ClaudeandClaude Opus 5.5 7bf6101460 Blocked fine storage: velocity-only surface actuation on fixed circles
- curved_links.surface_link_velocity: pure host helper extracted from
  CurvedLinkSoA.assign_surface (which now calls it; dense path unchanged).
  Rows follow physical cut-link order.
- blocks.cu BlockCurvedBoundary: one nullable trailing surface_velocity
  argument (2 floats/link), same convention as CurvedBoundaryKernel; no
  exchange output. Both call sites pass it explicitly (PeriodicBlockStepper
  passes null; single-level behaviour unchanged).
- FixedPatchStepper blocked mode accepts actuated stationary circles,
  validated against the physical fine_cfg; uploads the velocity table only
  when any actuation spec exists and fails before BFL if that buffer is
  missing (even at zero amplitude). Blocked exchange diagnostics are
  unavailable; dense exchange behaviour unchanged. No force aggregation,
  no motion, no new numerical formula.
Tests: 1x1 seam (0,0), receivers in 4 blocks, all q<.5 links fallback 0;
dense SoA never built in blocked mode; launch order incl. null vs velocity
pointer; missing-buffer guard at 0 and 0.002; helper == assign_surface;
aperture crosses a seam; GPU: periodic and slot_reservoir, +/-0.002,
8 coarse steps, dense vs 1x1/2x2/4x2 exact values + NaN masks.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 04:45:17 +00:00
ClaudeandClaude Opus 5.5 2bbf82121e Blocked fine storage: passive stationary circles via receiver-local BFL
FixedPatchStepper(fine_block=..., fine_objects=[circle]) now accepts
passive stationary circles; any surface actuation is rejected before CUDA
(dense fine path unchanged). prepare_fine_boundary still decides legality
on the dense fine grid; its SoA is mapped by the new array-based storage
adapter blocks.blocked_link_tables (map_cut_links: unique receiver owner
and local writer, row order kept; blocked flags sized nblocks*L), which
PeriodicBlockStepper._static_circles now also uses (single-level legality
builder unchanged). Per fine hop on the same buffer: halo -> existing
BlockCurvedBoundary -> BlockStep(flags, kind). No CUDA changes.
Tests: fixture 76 links / 68 q<.5 and 8 seam writes + 8 seam donors per
decomposed layout (both modes); adapter flags sizing; mocked launch-order
contract; GPU: periodic and slot_reservoir, 8 coarse steps, dense fine
body vs 1x1/2x2/4x2 fine blocks, exact values + NaN masks on masked
downloads and raw final arrays, ACTIVE fluid finite, solid/OUTER NaN.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 04:25:09 +00:00
ClaudeandClaude Opus 5.5 59921a0181 Blocked fine: fix INNER assertion, skip empty halo uploads, explicit fixture
- GPU test: after hop 2 some INNER_GHOST populations are NaN (pulled from
  OUTER poisoned on hop 1; 236 slots in pinned CPUREF); assert ACTIVE
  finite and OUTER NaN, keep the full dense-vs-blocked raw value/NaN-mask
  comparison.
- A 1x1 non-periodic fine layout has zero halo pairs: skip the halo table
  uploads (CUDA rejects mem_alloc(0)); _fine_halo already skips launch.
  CPU guard test with driver doubles.
- Drop the class-level fine_layout default; the followup fake backend now
  sets fine_layout = None explicitly.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 03:52:42 +00:00
ClaudeandClaude Opus 5.5 dc5850c198 Add opt-in blocked fine storage to FixedPatchStepper
fine_block=(by, bx) partitions the ENTIRE fine_shape (incl. the four-cell
patch ghost band) with a non-periodic one-cell-halo BlockLayout2D; body-
free only (fine_objects rejected in blocked mode; dense fine-body and
actuation path unchanged). Per coarse step on the same stream:
BlockExplode into ghost owner cores -> gather -> BlockStep(kind) ->
gather -> BlockStep(kind) -> unchanged PatchCoarseStep reading remapped
owner-core child slots in the original row/child order. Public arrays
keep the dense fine shape; only the private allocation is blocked.
blocks.cu: explicit BlockExplode; BlockStep gains nullable kind
(OUTER skip pull + NaN, INNER pull only, ACTIVE pull + SRT).
PeriodicBlockStepper passes an explicit null kind (raw launch ABI).
refinement.cu/refinement.py and coarse numerics unchanged.
Tests: CPU rejections and table remap (dense (q,cell) and child order
preserved, owner cores, kinds); GPU gate: periodic and slot_reservoir,
8 coarse steps, dense vs 1x1/2x2/4x2 fine blocks bit-identical in
masked downloads and raw coarse/unpacked fine incl. ghost NaN masks.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 03:48:54 +00:00
ClaudeandClaude Opus 5.5 3e44e90ca5 Test: negative-radius regression at integer center (16, 12)
At the half-integer center the -0.1/-0.5 circles were already rejected
as unresolved, so the case did not exercise the finite/positive-radius
guard; an integer center makes the squared radius mark one solid cell.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 03:46:14 +00:00
ClaudeandClaude Opus 5.5 eeda651626 Reject nonfinite or nonpositive circle geometry; tuple-normalize objects
_static_circles now checks finite center/radius and radius > 0 before
build_flag_mask (a negative radius squared to a one-cell solid), matching
prepare_fine_boundary. objects is converted to a tuple before the
emptiness test so an empty iterable takes the body-free path.
CPU tests: negative/NaN radius and nonfinite center rejected before CUDA;
empty iterator skips circle preparation.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 03:29:46 +00:00
ClaudeandClaude Opus 5.5 1a2d58312d Add stationary passive circles across same-level block seams
blocks.cu: BlockCurvedBoundary passes the receiver-local core cell and the
block-offset fi to the unchanged apply_bouzidi_link (Uw=Vw=0, obs=nullptr),
so k_s and the q<1/2 donor follow the window image the next pull reads.
BlockStep takes nullable block-cell flags; solid cores store NaN each step.
Adds core/flags.cuh, core/obs.cuh and boundary/curved_boundary.cuh.
block_stepper.py: objects=() accepts stationary passive circles only; host
preparation rejects before CUDA: non-circle/moving/actuated/sensor roles,
overlap, domain-edge contact, and any fluid core (all rows/columns, periodic
neighbors) whose pull from a solid lacks exactly one cut link. Uses
map_cut_links for unique receiver owner/local writer. Order per step:
halo copy -> BFL -> pull/collide, no copy in between. Fluid-only
initialize validation; solids NaN.
Tests: CPU rejections/tables; GPU 32x24, nu .1, 8 steps, circle
(15.5,11.5) R4.3 with halo-write and halo-donor presence asserted,
1x1/2x2/4x2 bit-identical, finite fluid/NaN solid, FP64 BFL+pull/SRT
<= 3e-6 on fluid; drift reported. Body-free gate unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 03:25:36 +00:00
ClaudeandClaude Opus 5.5 7071ed32a6 Test fixture: set _step_count for a complete ready state
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 02:59:40 +00:00
ClaudeandClaude Opus 5.5 7281d91fd6 Release the compiled module on close
close() now clears _module (initialized to None) inside the active
context together with the function references; lifecycle test asserts
it is retained on failed activation and cleared after close.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 02:59:08 +00:00
ClaudeandClaude Opus 5.5 4645e48a2c Test fixture: give fake backend mock halo table pointers
step() reads _dst/_src before the launch; without them the failed-launch
test raised AttributeError instead of the intended RuntimeError.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 02:55:39 +00:00
ClaudeandClaude Opus 5.5 4029a0b27e Add body-free periodic GPU consumer for BlockLayout2D
blocks.cu: separate compilation entry with NX/NY = block window (halo
included), blockIdx.z selects the block; BlockHaloCopy refreshes halos
from owner cores, BlockStep pulls/collides core cells with the existing
compute_rho_u/compute_feq/collide_srt_no_force. Existing kernels and
refinement numerics unchanged.
block_stepper.py: PeriodicBlockStepper (D2Q9/SRT/FP32 double buffers,
doubly periodic layout); launch-limit and free-memory checks before any
allocation; uploads only halo pairs; spare buffer starts NaN; lifecycle
mirrors FixedPatchStepper.
Tests: CPU contract doubles (rejections before CUDA/allocation, initialize
guards, failure/close/retry) and one GPU case: 32x24, 8 steps, 1x1/2x2/4x2
splits bit-identical, FP64 oracle max abs <= 3e-6, drift reported.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 02:49:09 +00:00
Frank14f 263f3b0c7a Preserve bounded validation evidence and historical failures 2026-10-15 10:13:51 +08:00
Frank14f cdbb9c42f6 Add jet, surface and fixed-patch tests and runnable examples 2026-10-15 10:13:49 +08:00