Commit Graph
34 Commits
Author SHA1 Message Date
TRON 60ad6b8b0d Generalize synchronized remap to multiple pinned body rectangles 2026-10-16 11:32:42 -07:00
TRON 864a220dbe Add research measured-strain fluid refinement example 2026-10-16 11:09:49 -07:00
TRON 6a7e9226db Add bounded pinned-circle manual L2 remapping 2026-10-16 10:41:56 -07:00
TRON 766658185f Add explicit synchronized L2 runtime continuation 2026-10-16 10:12:40 -07:00
TRON a24e6d2742 Support bounded three-level separated patch unions 2026-10-16 09:18:07 -07:00
TRON 20fb16ded5 Support stationary bodies and geometry Jets on nested L2 2026-10-16 08:49:08 -07:00
ClaudeandClaude Opus 5.5 64665319a4 Multipatch: fine_objects and jets on several patches (level-wide projection)
prepare_fine_boundary accepts a FixedPatchSet2D: objects are projected once
onto the level-wide fine grid (FixedPatchSet2D.fine_origin = (0, 0), same
coarse cell-edge -> fine cell transform), each circle must lie inside one
patch, coarse transfer cells are checked against the covered union, and body
ids stay object positions, so the existing jet one-body-per-jet, overlap and
(receiver, dir) writer checks run once over the whole level (caller obj_id
cannot bypass them). The stepper no longer rejects fine_objects/jets with
several patches; blocked link tables, BFL and the jet amplitude cache are
reused unchanged. No kernel change. README/docstring scope updated.

Tests: CPU translation/ids/rejections (outside one patch, one jet over two
bodies at zero amplitude); GPU one body via the new path == single-patch path
(both modes), two bodies (+jets, amplitude switch) == tiled single (bitwise).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-16 14:59:48 +00:00
ClaudeandClaude Opus 5.5 76555de4b8 Three levels: nested body-free child patch on a single chain
FixedPatchStepper(..., child=(x0,y0,x1,y1), child_block=(by,bx)) adds an L2
patch in L1 window cells (requires fine_block/base_block, one L1 patch, no
bodies/jets). nest_patch() builds the L1->L2 FixedPatch2D (2*nu on L1) and
the composed L1 kind (inherited INNER/OUTER kept, child-covered ACTIVE made
OUTER), rejecting from the actual tables: coverage outside ACTIVE, donors or
receivers not uncovered ACTIVE, surviving ACTIVE pulls of covered/OUTER cells
without a child-supplied arrival, and overlap with L1's restriction footprint.

Step order per L1 hop: explode L1(src)->L2, two L2 hops (halo, step), L1 halo,
receiver-local L2->L1 staging into src, L1 step. The second hop thus feeds L2
from the new L1 G; L2 is never refilled from L0; L0 stays once per step (1:2:4).
Existing kernels only; a small _Pair holds the child transfer bindings (explode
from the child module, parent window stride). Two-level API unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-16 14:35:57 +00:00
ClaudeandClaude Opus 5.5 6301887093 Multipatch: validate fine block before modulo, raw-core GPU checks, README usage
FixedPatchSet2D.fine_layout rejects non-positive, float, bool or wrong-length
block sizes with ValueError before any modulo (was ZeroDivisionError/TypeError).
The face gap-4 GPU test now checks raw owner cores directly by kind/covered
for every patch (ACTIVE finite, OUTER/holes/covered NaN; INNER not asserted),
not only the masked download and reversed-order parity. README documents the
body-free multipatch call, scope and level-wide fine shape.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-16 14:22:00 +00:00
ClaudeandClaude Opus 5.5 0397b6ebb7 Multipatch: CPU guards and frozen GPU gates through the new interface
One rectangle via the multipatch path is bit-identical to the old
single-patch blocked path (both modes, raw masks); S1 two-patch 16x32 equals
the 16x16 one-patch run tiled (bitwise), in both rectangle orders; face
gap-4 pair is order-invariant with exact covered/hole masks. Mass printed
as a diagnostic only.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-16 14:01:41 +00:00
ClaudeandClaude Opus 5.5 99ddb5fc27 Test: fine jet ownership uses remapped physical owners
Caller obj_id 99 spanning 2x2 fine blocks: driven rows carry the single
remapped owner 0 across several block ids, accepted; source id unchanged.
Caller ids (99, 7) under one tube are rejected at amplitude 0 and .002
on the 32x24 fixture; source ids unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-16 02:42:04 +00:00
ClaudeandClaude Opus 5.5 64d54ab23e Jets: one jet may drive links of only one body
surface_link_velocity (the shared host check used by CurvedLinkSoA,
ObjectManager sync, prepare_fine_boundary and FixedPatchStepper) now
rejects a jet whose selected cut links carry more than one body id, at
any amplitude, before uploads. Identity is the scene-wide body id of the
physical cut-link list, so one body spanning storage blocks stays valid.
No body selector; empty-selection semantics unchanged.

Tests: two bodies under one jet rejected (amplitude 2e-3 and 0) on the
Simulation/manager path and on the fine path; separate jets on separate
bodies accepted on both; one fine body crossing 2x2 block seams accepted.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-16 02:39:58 +00:00
ClaudeandClaude Opus 5.5 063508d1a3 Jet amplitude setter: failed state on device errors, commit after upload
FixedPatchStepper.set_jet_amplitude validates index/amplitude and builds
the new rows before touching the device; invalid input leaves the
instance usable. A copy/sync error marks it failed (as initialize/step/
download do) and re-raises; host state (_jets, blocked velocity table)
commits only after a successful upload. CurvedLinkSoA.set_jet_amplitude
likewise uploads a new array first and commits jets/surface_velocity
afterwards. Focused mocked-driver tests for dense and blocked copy
failure, invalid input and IndexError.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-15 08:49:35 +00:00
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
Frank14fandCursor 6e3756c587 fix(esopull): correct init layout and pre-streaming semantics (v0.5.1)
EsoPull curved boundaries and wall BCs now use consistent backing-layout
reads; InitEsoPull writes equilibrium in t=0 EsoPull layout. Cache N_OBJS
after compile and atomic config header writes to avoid parallel races.
Adds config screening tools, flume configs, and FP16S/EsoPull diagnosis doc.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-06-27 22:32:01 +08:00
Frank14fandCursor 04c2bc75ea feat(obs): unified zero_obs control and time-normalised readback
- Replace split zero_force_segment / zero_sensor_segment with unified
  zero_obs_async() — a single memset covers all three obs segments
  (force, torque, sensor), resetting the step accumulator.
- Add _obs_accum_steps counter so read_*(normalize=True) returns the
  physically meaningful per-step average for all telemetry fields.
- Sensor now always applies area-normalisation internally; the normalize
  parameter only controls the additional time-normalisation step.
- run() gains zero_obs=True parameter (default) to control reset-on-step.
- 7 new integration tests covering accumulation, zeroing, and normalise.
- Fix bug in test_sensor_accuracy.py (undefined loop variable i).
- Bump version to 0.4.0 for the API change.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-06-21 00:50:20 +08:00
Frank14fandCursor d5b7e98750 feat(compat): FP16S and ddf_shifting compatibility, EsoPull curved closure
Phase A: FP16S store precision verification
- Kan99b K2 FP16S: quantization sensitivity documented (St 0.170 -> 0.142)
- Sah04 S2 FP16S: PASS (St error 1.53% within 5% gate)
- Sah04 S4 FP16S: diverges at high blockage (known limitation)

Phase B: ddf_shifting code fixes
- Fix inlet west_velocity_rho_closure for shifted DDF (common.cuh)
- Fix curved force/torque accumulation for shifted DDF (curved_boundary.cuh, aux_kernels.cu)
- Fix host upload_ddf() asymmetry (field.py)
- Add checkpoint streaming/ddf_shifting match check (checkpoint.py)
- MRT shifting fix: MRT is NOT shift-invariant; unshift/reshift around collision
- Generalize inlet knowns repair from Zou-He to all west inlet schemes

Phase C: EsoPull curved boundary semantic closure (from round 2)
- streaming/esopull_semantic_helpers.cuh: single truth for physical-value semantics
- step/esopull_macro.cu: MacroscopicEsoPullKernel for correct GPU diagnostics
- SensorKernel, ForceRegionKernel share semantic helpers
- Kan99b K2: bit-identical to double-buffer
- Code-level comments document compatibility boundaries
- README updated with compatibility matrix
- output/round3_compatibility_summary.md: full round documentation

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-06-03 10:48:42 +08:00
Frank14f 2e052480c2 重构body api,性能分析,项目整理 2026-05-31 01:42:58 +08:00