Commit Graph
16 Commits
Author SHA1 Message Date
ClaudeandClaude Opus 5.5 19db7d362f refinement: free-memory preflight for every level before any device work
Restores the PeriodicBlockStepper safeguard dropped in 8e1feff, for all
FixedPatchStepper levels: both G buffers, halo pairs and boundary columns
of every level are compared with cuda.mem_get_info() before the stream,
any module or any buffer exists, raising MemoryError('level buffers exceed
free device memory'). Transfer tables are not counted (lower bound).

Tests: the single-level MemoryError case is back beside the param-layout
check (no stream, no allocation); host test_single_level covers depth 1
and depth 2 (no stream, build or allocation) and the wait failure after a
good step (count stays 1). Driver-mocking fixtures report ample free memory.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-27 01:10:15 +00:00
ClaudeandClaude Opus 5.5 8e1feffcad lbm: delete PeriodicBlockStepper; a single blocked level is FixedPatchStepper(refinements=[]) (S2)
The second owner of one blocked level goes. Its capability is the
zero-refinement FixedPatchStepper added in f764d68/985db14 (periodic SRT
level, stationary passive circles, shifted storage, NaN spare, single wait),
which matched it array for array on GPU at 1318fc2 (3 splits x +-circle x
shifted/unshifted) and passes the host lifecycle and fault contracts.

API change for callers:
  PeriodicBlockStepper(cfg, BlockLayout2D.build(ny, nx, by, bx), context=c, objects=o)
  -> FixedPatchStepper(cfg, context=c, refinements=[], level_blocks=[(by, bx)],
                       level_objects=[tuple(o)])
  download() returns a 1-tuple; static-circle tables are
  _level_boundaries[0].tables (was .boundary); no free-memory preflight.

- lbm/block_stepper.py deleted; MAX_GRID_DIM lives in refinement_stepper.
- tests/validation/multigrid/test_block_stepper.py: same frozen GPU gates
  (split identity, FP64 oracle, circle across seams) and pre-CUDA rejection
  cases on the single level; the old-vs-new bitwise test is retired with the
  old class (its result is the 1318fc2 record).
- tests/unit/test_block_step_sync.py folded into test_single_level.py
  (step(n) vs n x step(1) roles and one wait, step(0), mid-batch launch
  failure), which also gains bad-state rejection before upload, download
  failure and activation-failure retry on close.
- test_shifted_storage and test_raw_exchange_gpu use the single level; the
  raw-exchange reference feeds the AMR Jet column and raw ledger through the
  root level boundary instead of patching a private kernel field.
- Pinned historical runners (taylor_green/run_case.py, run_jet_circle_cv.py)
  are unchanged: they only run on trees that still ship PeriodicBlockStepper.
  No alias.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-27 01:02:08 +00:00
ClaudeandClaude Opus 5.5 1318fc205c test: single-level rejections include LES and window-exchange observation
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 22:37:40 +00:00
ClaudeandClaude Opus 5.5 985db14249 refinement: single level matches PeriodicBlockStepper lifecycle; host contract tests
Review follow-ups for S1 (f764d68):
- a single level starts its spare buffer NaN, as PeriodicBlockStepper does,
  so stale halos cannot hide (multilevel initialization unchanged);
- projected explosion is rejected for a single level;
- new tests/unit/test_single_level.py: NaN spare after initialize, per-step
  launch order halo -> BFL (circles only) -> BlockStep with a null kind and
  one wait per step(n), circle solids excluded from the state mask;
- the GPU bitwise comparison also covers shifted storage.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 22:36:05 +00:00
ClaudeandClaude Opus 5.5 f764d68c81 refinement: FixedPatchStepper(refinements=[]) owns a single blocked level (S1)
A zero-refinement FixedPatchStepper is the PeriodicBlockStepper capability:
one doubly periodic blocked SRT level (no LES, Jets or exchange
observation) whose level_objects[0] are stationary passive circles. The
circle tables come from the same checks, now blocks.static_circle_tables
(moved from block_stepper, pycuda-free); BFL, halo and BlockStep launches
are the shared helpers, with an all-ACTIVE (null) kind. Base-level bodies
stay rejected for depth >= 2. PeriodicBlockStepper and its callers are
unchanged (deleted only after S1's GPU result).

Tests: host contract rejections before CUDA for the single level; GPU
bitwise comparison against PeriodicBlockStepper for every split, with and
without the circle (the S2 gate).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 22:34:52 +00:00
ClaudeandClaude Opus 5.5 d318911221 cuda: compiler_v2 module becomes compiler
The package already exposed it as CelerisLab.cuda.compiler; the module now
has that name and every tip caller, doc and current-tree test imports it
directly (no alias).

Historical-revision tooling is not edited: the Taylor-Green harness
(run_case.py, hash-pinned check_projected_pairs.py, assess receipts) only
runs against a135252, whose package has cuda/compiler_v2.py, and the J
campaign runners and dated docs describe past trees. Its two current-tree
tests now import that package from git (historical_package) instead of the
tip, and the kernel-root test first checks the accepted case.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 19:15:50 +00:00
ClaudeandClaude Opus 5.5 18e8ef45f0 blocks: one definition of a blocked level's config and kernel launches (C1a)
PeriodicBlockStepper and FixedPatchStepper each built the blocked level's
kernel config (block window NX/NY, neutralized wall/slot/inlet fields,
omega/tau check) and each spelled out the raw BlockHaloCopy and
BlockCurvedBoundary argument lists. These now live once in lbm/blocks.py
(level_config, launch_halo_copy, launch_curved_boundary), so a kernel ABI
change has one host site. Both owners keep their own lifecycle; no new
class, no behaviour or default change. PeriodicBlockStepper now checks its
level config before any device work (same errors). Dense refinement levels
keep their own config check.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 19:12:11 +00:00
ClaudeandClaude Opus 5.5 2b3c2a7529 kernels: name compilation entries by storage route (uniform.cu, blocked.cu)
kernel_v2.cu -> uniform.cu (whole-domain NX/NY, Simulation) and
blocks.cu -> blocked.cu (one block window, every blocked level), with
InitTubeFlow_v2 -> InitTubeFlow. All live callers, include comments,
CONFIG.md, lifecycle doc and current-tree tests move together; no alias.
refinement.cu keeps its name until the dense-route decision.
Historical-revision tooling that names the old files (LES source pins,
timing/J campaign runners, TG run_case, supported_settings citations at
888c8038, development notes) is unchanged on purpose.
No arithmetic change; the PTX symbol InitTubeFlow_v2 is renamed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 14:07:34 +00:00
ClaudeandClaude Opus 5.5 ee7074a53d tests: restore the negative checks masked by stale fixtures (WP-B groups B-E)
Test-side only; no production, manifest or pin change. Group A identity
checks (test_source_pin and the jets.py pins) are unchanged and stay
failing as recorded.

B test_refined_dfg: receipt/plan fixtures load the frozen manifest JSON
  that verify_source returns for a matching tree (frozen_manifest), and
  the two manifest-selection assertions call verify_source with only its
  working-tree hash comparison neutralised (selected_manifest), so the
  malformed/collision/observation-mode/checkout receipt checks run again.
  test_source_pin remains the single identity check.
C test_channel_root: int(owner)-1 (uint64-int became float on NumPy 1);
  test_shifted_storage: FP32 NaN fill so the arrays stay FP32. Both reach
  the intended immutable/representation/FP32/density checks.
D test_block_stepper fake gets _ddf_shifting; test_fine_surface_gpu sets
  _boundary_hosts/_level_jets instead of the read-only fine_boundary view;
  test_refinement_followup fake_backend sets _layouts/_ddf_shifting/
  _level_cfgs instead of the read-only fine_layout view (as reviewed in
  32da2d73).
E test_kernel_contracts_reject_before_upload: an inconsistent pair tau is
  now first rejected by the hierarchy chain check; the test asserts that
  rejection, then bypasses _check_chain to prove the per-level omega/tau
  guard still rejects, and asserts no mem_alloc/memcpy_htod in either case.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-21 07:04:47 +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