Commit Graph
13 Commits
Author SHA1 Message Date
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