test: dense/blocked parity rows for MRT bodies and two fine bodies

Extends DenseBlockedRouteParity (test-only, on 1573942) with the residual dense-supported combinations named before any dense-route deletion: - test_mrt_bodies: MRT on a periodic root with a passive circle, an actuated circle (+/-) and a velocity Jet with an amplitude switch (MRT slot/reservoir is rejected on every route: on_site_velocity is SRT-only); - test_two_bodies: two circles on a wider patch (4, 4, 28, 20), periodic and slot/reservoir, Jet on one + passive other, and two actuated (+/-); routes D/D, D/B [None, (8, 8)], B/B [(12, 8), (8, 8)]. Same gate: exact FP32 bits and NaN masks at 0, 1, 2, 8 and after re-initialize; D/B vs B/B raw-exchange ledgers equal for every body. Docstring table updated; no tolerance, no src change. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
test: ledger comparator compares labels exactly and numbers with NaN-equality
2026-10-27 10:28:11 +00:00 · 2026-10-27 10:15:10 +00:00 · 2026-10-27 10:06:38 +00:00 · 2026-10-27 01:18:41 +00:00 · 2026-10-27 01:10:56 +00:00 · 2026-10-27 01:10:15 +00:00
77 changed files with 1973 additions and 1121 deletions
@@ -25,7 +25,7 @@ under the alternate compile branch.
 - No LES, no external forcing (`d_params` force fields stay zero).
 - Shifted Jets stay SRT-only: `body/jets.py` keeps its existing guard. This extension does not
  provide shifted MRT Jets.
- `PeriodicBlockStepper` (uniform blocked stepper) remains SRT-only.
+- `PeriodicBlockStepper` (uniform blocked stepper) remained SRT-only; it was removed in 8e1feff and its single-level capability, still SRT-only, is `FixedPatchStepper(refinements=[])`.

 ## Evidence status
 CPU: `tests/unit/test_refinement_mrt.py` (host guards, source contract). GPU:
@@ -0,0 +1,393 @@
+# Evidence index (draft, R0)
+
+Active-tree evidence at candidate 0f8fdbc30ea640c764b38ee598d8f4d135c81c85, classified for the release cleanup
+(thread C0C78UFKGRM 1793021833.593279). Draft for review: nothing is removed by this file.
+Recovery of any listed file: `git show 0f8fdbc3:<path>` (blob id given; an immutable tag is requested as anchor).
+
+Classes: K = conclusion kept in the active tree; K-code = tested research code kept; D-snapshot = pre-Gitea
+patch/diff transport (code landed in git from ed1a6ee on; not yet verified hunk by hunk); P-pin = generated hash inventories, git-status dumps and manifests;
+R-raw = run outputs, logs, signals, frames, one-shot run scripts and handoff notes. D, P, R leave the active tree in R1.
+No file under src/ or a test reads any of these paths (git grep at 0f8fdbc3); runners under tests/validation
+that read research/j* are handled in R2.
+
+| Class | Files | Bytes |
+|---|---:|---:|
+| K | 31 | 196,506 |
+| K-code | 3 | 24,324 |
+| D-snapshot | 18 | 338,948 |
+| P-pin | 50 | 979,177 |
+| R-raw | 142 | 3,790,554 |
+
+## (root)
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| BASELINE_JET_AMR.txt | R-raw | 5,244 | 62aa2a66dc05 |
+
+## output
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| round3_compatibility_summary.md | R-raw | 4,366 | f243318f5591 |
+
+## research/fine_surface_20261014
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 54,507 | 71a4543ead37 |
+| cpu_device_guard_cmd.txt | R-raw | 241 | 06e0f35d92b3 |
+| cpu_device_guard_exit.txt | R-raw | 2 | 573541ac9702 |
+| cpu_device_guard_stdout.txt | R-raw | 98 | 424d584a031c |
+| cpu_existing_runs.json | R-raw | 650 | 3179536407b7 |
+| cpu_final_run.json | R-raw | 386 | 523fadb3fec9 |
+| cpu_final_stdout.txt | R-raw | 99 | 9cfe73da3858 |
+| cpu_lifecycle_run.json | R-raw | 387 | 5b403b9e2f30 |
+| cpu_lifecycle_stdout.txt | R-raw | 112 | 3c5c49b285fa |
+| cpu_maps_stdout.txt | R-raw | 95 | 71dd87f9faf9 |
+| cpu_preallocation_run.json | R-raw | 337 | fa6e96d98cdd |
+| cpu_preallocation_stdout.txt | R-raw | 98 | 5d643576f251 |
+| cpu_shared_kernel_stdout.txt | R-raw | 98 | 8a0c046bb9b1 |
+| cpu_surface_command.txt | R-raw | 113 | 2a5ff31e8d06 |
+| cpu_surface_exit.txt | R-raw | 2 | 573541ac9702 |
+| cpu_surface_projection_command.txt | R-raw | 171 | 70b7f5c62c1a |
+| cpu_surface_projection_exit.txt | R-raw | 2 | 573541ac9702 |
+| cpu_surface_projection_stdout.txt | R-raw | 113 | efd37a149cc1 |
+| cpu_surface_stdout.txt | R-raw | 99 | 9cfe73da3858 |
+| gpu_preparation_run1.py | R-raw | 14,043 | ee586d035b7a |
+| gpu_run.json | R-raw | 2,943 | 15b8ba70d528 |
+| gpu_stdout.txt | R-raw | 5,424 | c5216e0267f3 |
+| gpu_test_run1.py | R-raw | 15,479 | f30376abc97f |
+| incremental.diff | D-snapshot | 53,853 | 318d92e7d177 |
+| manifest.json | P-pin | 8,411 | 8281338416af |
+| report.txt | K | 8,373 | 6b9e5123e414 |
+
+## research/fine_surface_joint_20261014
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline_00.json | P-pin | 21,302 | 4ae101ebd02b |
+| baseline_01.json | P-pin | 26,191 | 548cf7b3b494 |
+| baseline_02.json | P-pin | 22,440 | e0421e7a4052 |
+| baseline_03.json | P-pin | 8,584 | da99b45f32a4 |
+| blow_boundary.csv | R-raw | 12,536 | e13363b69457 |
+| blow_collisions.csv | R-raw | 19,145 | 5d5b395486bb |
+| blow_ledger.csv | R-raw | 15,977 | b43bbfd0b8d2 |
+| blow_propagation.csv | R-raw | 9,714 | f06d80d296ca |
+| blow_signals.npz | R-raw | 221,114 | 528adf75baa2 |
+| blow_summary.json | R-raw | 4,113 | 32285ac15be3 |
+| command.txt | R-raw | 232 | 510ecb2925dd |
+| criteria.txt | K | 7,480 | dd36e76a3be1 |
+| frozen.json | R-raw | 11,972 | b5a3324e8686 |
+| gpu_run.json | R-raw | 1,955 | 67ac4785f7e4 |
+| incremental.diff | D-snapshot | 625 | e03c132d22df |
+| integrity.json | P-pin | 826 | a8b6e4be88d2 |
+| manifest.json | P-pin | 4,388 | d2af50260c39 |
+| regions.json | R-raw | 27,525 | 7bba26904221 |
+| report.txt | K | 8,207 | d7571e068ad3 |
+| result.json | R-raw | 11,442 | 9580c2686e70 |
+| review.txt | K | 2,117 | bbe1d8a83d47 |
+| runner.py | R-raw | 5,847 | 812973acc000 |
+| stdout.txt | R-raw | 19,940 | c9218e2e48fc |
+| suction_boundary.csv | R-raw | 12,505 | b870be05e6c0 |
+| suction_collisions.csv | R-raw | 19,173 | 2de48c0a84ab |
+| suction_ledger.csv | R-raw | 16,208 | 92cf1f44911e |
+| suction_propagation.csv | R-raw | 9,720 | f017bcc3ca6f |
+| suction_signals.npz | R-raw | 220,058 | 66d06cf2d909 |
+| suction_summary.json | R-raw | 4,126 | d70a1d371950 |
+| validate.py | R-raw | 19,576 | 37441febd168 |
+| zero_boundary.csv | R-raw | 11,491 | 65c92e48c974 |
+| zero_collisions.csv | R-raw | 18,915 | 027cb8ff6b23 |
+| zero_ledger.csv | R-raw | 16,051 | 34a99bb89dd7 |
+| zero_signals.npz | R-raw | 127,958 | a966f093a9cd |
+| zero_summary.json | R-raw | 1,508 | 6e033f991438 |
+
+## research/fixed_patch_jet_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 69,721 | 91ad9c4d1e88 |
+| before.json | R-raw | 690 | 06e1b6278a29 |
+| fixed_patch_jet.py | R-raw | 9,646 | 6fbdcd9a076b |
+| incremental.patch | D-snapshot | 9,891 | b43129d82c69 |
+| manifest.json | P-pin | 1,967 | 08daa0248f60 |
+| notes.txt | K | 2,214 | 50070fe54efb |
+| results.json | R-raw | 1,100 | d959a6a1cc2c |
+| stdout.txt | R-raw | 1,272 | eccdc0f6b16a |
+
+## research/fixed_patch_jet_closure_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 72,692 | 8a157af82ecf |
+| cpu_results.json | R-raw | 628 | 8e4a1c090c93 |
+| cpu_stdout.txt | R-raw | 590 | d800b783fe47 |
+| incremental.patch | D-snapshot | 3,490 | a9a486c68afd |
+| layout_uncropped.png | R-raw | 51,785 | 0cf1cfa5e99d |
+| manifest.json | P-pin | 1,849 | 31cf0c739241 |
+| notes.txt | K | 3,023 | 0f7e7bdfd0e6 |
+
+## research/fixed_patch_jet_records_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 71,228 | a8e967f76e08 |
+| before.json | R-raw | 690 | ebd6f4251fc2 |
+| fixed_patch_jet.py | R-raw | 10,650 | f6b299885da0 |
+| incremental.patch | D-snapshot | 5,539 | c7f4d0c02ccc |
+| manifest.json | P-pin | 2,057 | 1d616c2541df |
+| notes.txt | K | 2,449 | b58cec14d359 |
+| results.json | R-raw | 1,136 | 788af64a6ebd |
+| stdout.txt | R-raw | 20,750 | a765b06deefa |
+
+## research/j1-2026-10-12
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| REVIEW.md | K | 7,667 | 8f469daef054 |
+| baseline-status.txt | P-pin | 2,601 | 64ca05ab952e |
+| baseline.json | P-pin | 23,052 | ab5852bcb0ef |
+| incremental.diff | D-snapshot | 29,912 | dd9c4e688888 |
+| results.json | R-raw | 68,294 | 56547aa8a678 |
+
+## research/j2-slot-2026-10-12
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| REVIEW.md | K | 12,930 | 69fa4d3bccd3 |
+| baseline-status.txt | P-pin | 2,825 | e711afb0016c |
+| baseline.json | P-pin | 24,660 | 286e023da94f |
+| incremental.diff | D-snapshot | 25,231 | 649f3b5d168d |
+| results.json | R-raw | 477,249 | b4272e25179d |
+
+## research/j3-coverage-2026-10-12
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| REVIEW.md | K | 11,409 | 5887cf73dce2 |
+| baseline-status.txt | P-pin | 3,055 | 78e88356caa3 |
+| baseline.json | P-pin | 30,407 | 095972547a7f |
+| incremental.diff | D-snapshot | 20,114 | f0d502dc5c5b |
+| negative-budget-20744.h5 | R-raw | 88,017 | 693b133be84d |
+| negative-end-20760.h5 | R-raw | 88,174 | 429270d263d4 |
+| restart-21280.h5 | R-raw | 88,062 | 74900affae38 |
+| restart-reference-21296.h5 | R-raw | 88,154 | 7510ee97885b |
+| results.json | R-raw | 296,314 | 7806b10e65cb |
+| verification.json | K | 1,468 | 78a07c47aea4 |
+
+## research/j3-lag-2026-10-12
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| REVIEW.md | K | 11,732 | 34fe5dbb2cae |
+| baseline-status.txt | P-pin | 2,968 | 6ae6094ea9d7 |
+| baseline.json | P-pin | 26,505 | 5c289d14ff30 |
+| end-21280.h5 | R-raw | 92,366 | f30b209179ab |
+| incremental.diff | D-snapshot | 15,790 | 5c6c485e1f75 |
+| results.json | R-raw | 21,718 | 2bf58939129c |
+| start-20760.h5 | R-raw | 92,247 | ae0abd05190c |
+
+## research/j3-offline-2026-10-12
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| REVIEW.md | K | 6,271 | a76e1b486bd2 |
+| lags.csv | R-raw | 1,104 | dd3136dd2dc2 |
+| localization.png | R-raw | 136,013 | 74b56f3e814c |
+| regions.csv | R-raw | 2,126 | 5d361e6872af |
+| results.json | R-raw | 8,852 | bf77b0f534d4 |
+
+## research/j3-reservoir-2026-10-12
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| REVIEW.md | K | 13,563 | a74b67a64e0d |
+| baseline-status.txt | P-pin | 2,896 | a318e077dfae |
+| baseline.json | P-pin | 25,571 | 060502ea7dd8 |
+| incremental.diff | D-snapshot | 22,898 | 8e2c79a86c98 |
+| results.json | R-raw | 182,747 | 7fd5e51a1785 |
+
+## research/refinement_example_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 8,286 | 591478100da8 |
+| before.json | R-raw | 818 | 7ab0962775b1 |
+| incremental.patch | D-snapshot | 8,682 | 7ad8719d0167 |
+| manifest.json | P-pin | 1,798 | bf1eee8a7cf8 |
+| notes.txt | K | 1,185 | 7ff019ad023d |
+| results.json | R-raw | 1,118 | 94697b0988b9 |
+| stdout.txt | R-raw | 325 | 501446f557f4 |
+
+## research/refinement_followup_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.txt | K | 4,887 | af65e889f20e |
+| baseline.json | P-pin | 5,343 | 7c60ad8ad692 |
+| cleanup.json | R-raw | 1,253 | 15cf342fba2e |
+| cpu_results.json | R-raw | 1,056 | 1166dffd2a09 |
+| cpu_results_attempt1.json | R-raw | 642 | d139d5508105 |
+| cpu_stdout.txt | R-raw | 499 | d1b33ea20563 |
+| cpu_stdout_attempt1.txt | R-raw | 879 | 6a9fe7fce39a |
+| gpu_before.json | R-raw | 589 | 46762a0ab207 |
+| gpu_before_no_launch.json | R-raw | 589 | 074982d1b2d4 |
+| gpu_rectangle_before.json | R-raw | 589 | 831e3793339f |
+| gpu_rectangle_before_no_launch_1791886104981638567.json | R-raw | 589 | c33885d50bfd |
+| gpu_rectangle_results.json | R-raw | 1,003 | 9c6da219ac6c |
+| gpu_rectangle_stdout.txt | R-raw | 332 | 9185bf303954 |
+| gpu_results.json | R-raw | 1,189 | 53e4f1c44de1 |
+| gpu_seed1_before.json | R-raw | 589 | c77ea3485f84 |
+| gpu_seed1_results.json | R-raw | 996 | b718e8af84e4 |
+| gpu_seed1_stdout.txt | R-raw | 306 | 4d7de0897137 |
+| gpu_stdout.txt | R-raw | 68 | b0956fe3cdb0 |
+| incremental.patch | D-snapshot | 16,836 | 9df25d8c177a |
+| manifest.json | P-pin | 4,410 | 177704d6f1ae |
+| run_gpu.py | R-raw | 5,237 | b5af3f0d4f79 |
+
+## research/refinement_prepare_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.txt | K | 8,939 | 8e210c28c7c5 |
+| baseline.json | P-pin | 46,348 | d5f257216b8d |
+| incremental.patch | D-snapshot | 15,751 | 0ddb32550584 |
+| manifest.json | P-pin | 2,113 | 9fd94ccf2c3e |
+| validation.json | R-raw | 3,567 | f057e7cbcd35 |
+| validation.txt | R-raw | 873 | 594daa45c58c |
+
+## research/refinement_slot_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 67,717 | d498081e7c50 |
+| before.json | R-raw | 690 | 430e76b3aa8a |
+| cpu_results.json | R-raw | 1,421 | 4f6bc4d9428f |
+| incremental.patch | D-snapshot | 18,676 | 67a3e8cc9e68 |
+| manifest.json | P-pin | 2,270 | 6515458290f4 |
+| notes.txt | K | 2,612 | 08a713ba4b37 |
+| results.json | R-raw | 1,160 | 03e683027ba8 |
+| stdout.txt | R-raw | 643 | d57b91769826 |
+
+## research/refinement_slot_followup_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 68,536 | 89a26ce79438 |
+| before.json | R-raw | 690 | e707df192ff2 |
+| incremental.patch | D-snapshot | 495 | c6f6fa11237d |
+| manifest.json | P-pin | 1,920 | 4840243a33aa |
+| notes.txt | K | 1,549 | 4b77583ca116 |
+| planar_velocity.cuh | R-raw | 1,911 | 87db7466c843 |
+| planar_velocity.sha256 | P-pin | 251 | 6694c4a1c1f0 |
+| results.json | R-raw | 1,050 | 6ef6eda854c6 |
+| stdout.txt | R-raw | 1,160 | 461860351c0c |
+
+## research/refinement_step_20261013
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.txt | K | 10,671 | 88d6f24743f5 |
+| baseline.json | P-pin | 40,620 | 20f61559dd38 |
+| cpu_validation.json | R-raw | 1,231 | 7999fd30018c |
+| gpu_failure_validation.json | R-raw | 868 | 5f6db65ee4f9 |
+| gpu_failure_validation.txt | R-raw | 248 | d9a570a1e9e9 |
+| gpu_validation.json | R-raw | 672 | 2028f59e1055 |
+| gpu_validation.txt | R-raw | 632 | 9699918ee127 |
+| incremental.patch | D-snapshot | 26,516 | d95f9347633d |
+| manifest.json | P-pin | 2,075 | 32a7f1a7309e |
+
+## research/repair-2026-10-10
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| TRON_HANDOFF.md | R-raw | 25,208 | a339786c7944 |
+| baseline-hashes.json | P-pin | 16,020 | 9b4e24cb1d10 |
+| baseline-status.txt | P-pin | 429 | 515be249b638 |
+| checkpoint-domain-2026-10-12.diff | D-snapshot | 14,443 | 035ed630dda1 |
+| classic-results.json | R-raw | 244,729 | 7297aa0d5bca |
+| classic-signals.npz | R-raw | 630,332 | 8cf6d1218aa0 |
+| final-gpu.log | R-raw | 191 | a0eec1a8ea31 |
+| final-unit.log | R-raw | 188 | 94efba2f1d42 |
+| provenance.json | P-pin | 1,020 | 25e22b4f0696 |
+| repair-architecture.md | K | 9,498 | 12051c3e8571 |
+| repair-numerical.md | K | 10,769 | 0c5757fe5f36 |
+| repair-state.md | K | 8,481 | 3c7f2d5b687a |
+| simplification-checks.log | R-raw | 1,045 | fa9d464be740 |
+
+## research/shear_refinement
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.md | K-code | 4,142 | 67bdd8ed622b |
+| boundary_shear_tagging.py | K-code | 13,859 | b5415a1c4ea1 |
+| example.py | K-code | 6,323 | 66295d3f0273 |
+
+## research/surface_actuation_20261014
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| baseline.json | P-pin | 54,139 | 0ad06888dbeb |
+| cpu_checkpoint_run.json | R-raw | 400 | 39681b60ff15 |
+| cpu_checkpoint_stdout.txt | R-raw | 114 | 929e4b8d4f46 |
+| cpu_initial.txt | R-raw | 620 | 29213c58a679 |
+| cpu_surface_command.txt | R-raw | 118 | 6d8c615a6e03 |
+| cpu_surface_exit.txt | R-raw | 2 | 573541ac9702 |
+| cpu_surface_stdout.txt | R-raw | 99 | ee9237fc26e5 |
+| gpu_run.json | R-raw | 1,698 | dc8f87eafaa9 |
+| gpu_stdout.txt | R-raw | 159 | b262891e52c3 |
+| incremental.diff | D-snapshot | 50,206 | 43090a1a8b11 |
+| manifest.json | P-pin | 7,527 | 35daa01271cb |
+| report.txt | K | 5,078 | 1162a6479f9f |
+
+## research/wall_seam_pulse_20261014
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| analyze.py | R-raw | 10,343 | ed7cf6999a49 |
+| interpretation.txt | K | 3,514 | 90935db2aa99 |
+| results.json | R-raw | 15,137 | 678a0759871a |
+| run.json | R-raw | 1,170 | 53e303fd0d5d |
+| sources.csv | R-raw | 2,313 | 40d3736b8ef1 |
+| stdout.txt | R-raw | 1,179 | 756d8c38e213 |
+
+## tests/validation/multigrid/cpuref_0_3_evidence
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.txt | K | 4,133 | 62897bfd9812 |
+| README_grok.md | K | 1,751 | 2d36e05661a8 |
+| baseline.json | P-pin | 39,291 | 188491096547 |
+| manifest.json | P-pin | 3,495 | 10c19f5fe5ca |
+| results_horsea.json | R-raw | 2,281 | 3b163384e376 |
+| stdout_grok.txt | R-raw | 7,143 | 5a0c270ea4ae |
+| stdout_horsea.txt | R-raw | 9,013 | adff8bb8d295 |
+
+## tests/validation/multigrid/local_stress_v1_evidence
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.txt | K | 8,239 | 3513fd294223 |
+| TABLE_local_stress_observable_v1.md | K | 2,967 | 6115dcde0f46 |
+| baseline.json | P-pin | 43,991 | ab631ea4f2ad |
+| comparison.json | R-raw | 13,248 | d7d1b36eb247 |
+| manifest.json | P-pin | 4,132 | c6c44d8125cd |
+| results_horsea.json | R-raw | 2,671 | 79bdf001f948 |
+| run_once.py | R-raw | 4,313 | fa9c5239e993 |
+| source_manifest.json | P-pin | 622 | 0516f1860407 |
+| stdout_horsea.txt | R-raw | 17,091 | a20c554693df |
+| stdout_local_stress_observable_v1.txt | R-raw | 12,238 | 692ffe1a05da |
+
+## tests/validation/multigrid/shear_v0_2_evidence
+
+| File | Class | Bytes | Blob |
+|---|---|---:|---|
+| README.txt | K | 7,220 | 772f74ad40ff |
+| baseline.json | P-pin | 41,751 | 70d9c9fa71c0 |
+| comparison.json | R-raw | 7,679 | 0af8736f9a01 |
+| joint_report.md | K | 6,110 | eed871221298 |
+| manifest.json | P-pin | 3,796 | b22d03a1ba9e |
+| results_horsea.json | R-raw | 1,957 | b2cd48a07a08 |
+| run_once.py | R-raw | 3,245 | 0b8d5e21f326 |
+| source_manifest.json | P-pin | 574 | 4d1c742690c7 |
+| stdout_grok.txt | R-raw | 7,950 | 0e80b1aba7f6 |
+| stdout_horsea.txt | R-raw | 11,421 | 3c867cdac197 |
@@ -0,0 +1,25 @@
+# Device state lifecycle (uniform Simulation, FixedPatchStepper)
+
+State is grouped by what a change invalidates. This file describes the code as
+it is; it adds no framework. Line references are to candidate 0f8fdbc3.
+
+| Class | Contents | How it changes | Owner rule |
+|---|---|---|---|
+| 1. Module and its globals | Compile-time (visible to the compiler, fixed per module): `config.h` macros incl. `NX`/`NY`, `COLLISION_MODEL`, `BLOCK_ROOT_BC`, `WALL_*`. Runtime per-module globals (not visible to the compiler; written by host code): `__constant__` `d_params`, `d_w`, `d_cx`, `d_cy`, `d_block_channel`, `d_block_wall`; plus bound kernel handles | Globals are written once by the owner right after `load_module`. Uniform grid: `Simulation._recompile` (construction, body-topology sync) loads a new module and re-uploads; `update_runtime_params` rewrites `d_params` (omega, force, n_objects) between runs. Steppers: never rewritten after construction, so they are immutable for the instance's lifetime although they are runtime values. The same quantity can be a macro in one route and a global in another (dense slot root: `WALL_SLOT_WIDTH`/`WALL_NORMAL_VELOCITY` macros; blocked slot root: `d_block_wall`), which can change generated arithmetic. | Every owner loads its own module (`compiler.load_module` = `module_from_file`, no cache). A loaded module is never shared between live owners: globals are per module, so sharing would share mutable `d_params`. |
+| 2. Geometry, layout, tables | flags, curved-link tables, block layouts and halo pairs, transfer (explode/coalesce) tables, block kinds | Immutable for the life of an instance. Uniform grid: rebuilt by body sync. `FixedPatchStepper`: a geometry change is a new instance (`remap_child`), committed only after a successful transfer. | No in-place table edits. |
+| 3. Controls and dependent boundary columns | Jet amplitude -> per-link `surface_velocity`; body actions (uniform grid) | Between steps, without a table or module rebuild (`set_jet_amplitude`). | Host copy commits only after the device upload succeeded; a device error fails the instance. |
+| 4. Per-step state | population double buffers and their parity, raw-exchange ledgers, observations | Device-only during `step(n)` / `run(n)`; buffer swaps are host-handle swaps. | Valid on the host only after the single wait at the end of the call. |
+
+Storage routes are separate concerns:
+
+- Original uniform grid (`uniform.cu`, `Simulation`): whole-domain `NX`/`NY`.
+- Blocked levels (`blocked.cu`; every `FixedPatchStepper` level, including the
+  single level of `refinements=[]`): `NX`/`NY` are one block window including halo.
+- Dense refinement levels (`refinement.cu`, `level_blocks[l] is None`).
+
+A single-rectangle refinement (`FixedPatch2D`) is stored in its own fine
+window frame in both dense and blocked form (`refinement_stepper.py:222-224`,
+`BlockLayout2D.build(*pair.fine_shape, ...)`). Only a rectangle set
+(`FixedPatchSet2D`) uses the level-wide frame `(2*ny, 2*nx)`, whose
+`fine_layout` requires that full grid to be divisible by the tile
+(`refinement.py:240-252`).
@@ -1,6 +1,7 @@
-# Supported settings: 2D Jet + fixed multilevel refinement (candidate c95a7ba)
+# Supported settings: 2D Jet + fixed multilevel refinement (candidate 888c8038)

-Every entry names the guard that enforces it (file:line at c95a7ba) and, where one
+Every entry names the guard that enforces it (file and function or kernel guard,
+at source revision 888c80381c006e976be2af8a452288729fc17c1b) and, where one
 exists, an existing test that exercises it. A setting not listed here is not
 claimed. Observations at the end are tied to the exact case in which they were
 measured; they are not general bounds.
@@ -9,12 +10,12 @@ measured; they are not general bounds.

 | Setting | Supported | Guard |
 |---|---|---|
-| Dimension / lattice | 2D D2Q9 only (3D refinement is future work) | `lbm/refinement.py:65`, `lbm/kernels/blocks.cu:25`, `lbm/kernels/refinement.cu:26` |
+| Dimension / lattice | 2D D2Q9 only (3D refinement is future work) | `lbm/refinement.py` `FixedPatch2D.from_config`; `#error` guards in `lbm/kernels/blocks.cu` and `lbm/kernels/refinement.cu` |
 | Precision / streaming | FP32 storage, double buffer | same guards |
-| Collision on refinement levels | SRT or MRT; TRT and LES rejected | `lbm/refinement.py:65` (re-checked for fine boundaries at `lbm/refinement.py:457`) |
-| Shifted (ddf_shifting) storage | only when every level is blocked; dense or mixed rejected | `lbm/refinement_stepper.py:198`, `:469`; `lbm/kernels/refinement.cu:26` (dense patch kernel unshifted only) |
-| Fine relaxation | 1/tau_fine must lie inside the configured omega guards | `lbm/refinement.py:80` |
-| Uniform blocked backend `PeriodicBlockStepper` | doubly periodic, SRT only, stationary passive circles | `lbm/block_stepper.py:84`, `:91`, `:45`-`:67` |
+| Collision on refinement levels | SRT or MRT; TRT and LES rejected | `lbm/refinement.py` `FixedPatch2D.from_config` (re-checked for fine boundaries in `prepare_fine_boundary`) |
+| Shifted (ddf_shifting) storage | only when every level is blocked; dense or mixed rejected | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__`, `FixedPatchStepper._check_storage`; `#error` guard in `lbm/kernels/refinement.cu` (dense patch kernel unshifted only) |
+| Fine relaxation | 1/tau_fine must lie inside the configured omega guards | `lbm/refinement.py` `FixedPatch2D.from_config` |
+| Single blocked level `FixedPatchStepper(refinements=[], level_blocks=[(by, bx)], level_objects=[circles])` (replaces `PeriodicBlockStepper`, removed in 8e1feff) | doubly periodic, SRT only, no LES/Jets/projection/exchange observation, stationary passive circles; `download()` returns a 1-tuple | `lbm/refinement_stepper.py` `_check_levels`, `FixedPatchStepper.__init__`; `lbm/blocks.py` `static_circle_tables` |

 Tests: `tests/unit/test_refinement_mrt.py`, `tests/validation/multigrid/test_refinement_mrt_gpu.py`,
 `tests/unit/test_shifted_storage.py`, `tests/unit/test_shifted_block_operators.py`.
@@ -23,17 +24,17 @@ Tests: `tests/unit/test_refinement_mrt.py`, `tests/validation/multigrid/test_ref

 | Setting | Supported | Guard |
 |---|---|---|
-| Depth | at least one refinement level; one entry per level in `level_blocks`/`level_objects`/`level_jets` | `lbm/refinement_stepper.py:72`, `:75` |
-| Single L1 rectangle | two-coarse-cell margin to the domain edge | `lbm/refinement.py:73` |
-| Several L1 rectangles (union) | require `level_blocks[0]` and `level_blocks[1]`; transfers must close (dependency checks, no fixed gap rule) | `lbm/refinement_stepper.py:84`; `lbm/refinement.py:283`-`:295` |
-| Nested levels | depth >= 3 must be all blocked; a blocked level requires blocked levels below; a scalar child needs a scalar parent | `lbm/refinement_stepper.py:82`, `:87`; `lbm/hierarchy.py:54` |
-| Nested transfer support | child only over inherited ACTIVE cells, donors/receivers uncovered ACTIVE, no child over the parent restriction footprint | `lbm/refinement.py:391`-`:407` |
-| Tile shape | positive `(by, bx)`; the level allocation must be divisible by the tile (no partial tiles yet) | `lbm/refinement_stepper.py:78`, `lbm/refinement.py:246` (shape); `lbm/blocks.py:134` (divisibility) |
-| Coarse boundary | `periodic`, `slot_reservoir` or `channel` | `lbm/refinement_stepper.py:177` |
-| `slot_reservoir` | on-site velocity bottom slot, reservoir top, velocity 0, 2 <= slot width < nx, patch rows 4..ny-4 | `lbm/refinement_stepper.py:205`, `:212`; `lbm/refinement.py:546` |
-| `channel` | all levels blocked; Zou-He inlet, pressure-NEQ outlet, stationary bounce-back walls; transfers away from the closure | `lbm/channel.py:22`, `:26`, `:43`, `:54` |
-| Explosion | `constant` (default); `projected_linear` only periodic, all-blocked SRT, no bodies/Jets, no LES (experimental) | `lbm/refinement_stepper.py:188`, `:193` |
-| L2 runtime remap (`remap_child`) | periodic three-level L2 sets only, no L0/L1 bodies, identical pinned boundary | `lbm/refinement_stepper.py:557`, `:567`-`:581` |
+| Depth | at least one refinement level; one entry per level in `level_blocks`/`level_objects`/`level_jets` | `lbm/refinement_stepper.py` `_check_levels` |
+| Single L1 rectangle | two-coarse-cell margin to the domain edge | `lbm/refinement.py` `FixedPatch2D.from_config` |
+| Several L1 rectangles (union) | require `level_blocks[0]` and `level_blocks[1]`; transfers must close (dependency checks, no fixed gap rule) | `lbm/refinement_stepper.py` `_check_levels`; `lbm/refinement.py` `_check_parent_pull_closure` |
+| Nested levels | depth >= 3 must be all blocked; a blocked level requires blocked levels below; a scalar child needs a scalar parent | `lbm/refinement_stepper.py` `_check_levels`; `lbm/hierarchy.py` `build_hierarchy` |
+| Nested transfer support | child only over inherited ACTIVE cells, donors/receivers uncovered ACTIVE, no child over the parent restriction footprint | `lbm/refinement.py` `_nested_kind` |
+| Tile shape | positive `(by, bx)`; the level allocation must be divisible by the tile (no partial tiles yet) | `lbm/refinement_stepper.py` `_check_levels`, `lbm/refinement.py` `FixedPatchSet2D.fine_layout` (shape); `lbm/blocks.py` `BlockLayout2D.build` (divisibility) |
+| Coarse boundary | `periodic`, `slot_reservoir` or `channel` | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| `slot_reservoir` | on-site velocity bottom slot, reservoir top, velocity 0, 2 <= slot width < nx, patch rows 4..ny-4 | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__`; `lbm/refinement.py` `prepare_fine_boundary` |
+| `channel` | all levels blocked; Zou-He inlet, pressure-NEQ outlet, stationary bounce-back walls; transfers away from the closure | `lbm/channel.py` `validate_channel_config`, `validate_channel_transfers` |
+| Explosion | `constant` (default); `projected_linear` only periodic, all-blocked SRT, no bodies/Jets, no LES (experimental) | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| L2 runtime remap (`remap_child`) | periodic three-level L2 sets only, no L0/L1 bodies, identical pinned boundary | `lbm/refinement_stepper.py` `FixedPatchStepper.remap_child` |

 Tests: `tests/unit/test_refinement_three_level_multipatch.py`, `tests/validation/multigrid/test_refinement_multipatch.py`,
 `test_refinement_three_level.py`, `test_refinement_three_level_multipatch_gpu.py`, `test_refinement_blocked_base.py`,
@@ -44,14 +45,14 @@ Tests: `tests/unit/test_refinement_three_level_multipatch.py`, `tests/validation

 | Setting | Supported | Guard |
 |---|---|---|
-| Bodies | stationary circles without sensor/force-region roles, entirely inside the patch union, fully resolved, disjoint, stencil away from the two-hop transfer footprint | `lbm/refinement.py:474`-`:542` |
-| Level | not on L0; only on the finest level of a nested hierarchy | `lbm/refinement_stepper.py:80`, `:89` |
-| Jet condition | velocity only; `pressure` is reserved and raises NotImplementedError | `body/jets.py:148`, `:150` |
-| Jet geometry | each Jet selects resolved links of exactly one stationary body; no Jet/Jet or Jet/SurfaceActuation overlap; abs(amplitude) < 1 lattice | `body/jets.py:155`, `:175`, `:177`, `:255`, `:261`, `:263` |
-| Jets with shifted storage | only in the all-blocked stepper with SRT and no LES | `body/jets.py:249` |
-| SurfaceActuation | unshifted storage only | `body/objects.py:155` |
-| Raw boundary exchange (force) | blocked body levels only; the window ledger requires `observe_raw_exchange` | `lbm/refinement_stepper.py:264`, `:269` |
-| Checkpoint with Jets | refused | `simulation.py:687` |
+| Bodies | stationary circles without sensor/force-region roles, entirely inside the patch union, fully resolved, disjoint, stencil away from the two-hop transfer footprint | `lbm/refinement.py` `prepare_fine_boundary` |
+| Level | not on L0; only on the finest level of a nested hierarchy | `lbm/refinement_stepper.py` `_check_levels` |
+| Jet condition | velocity only; `pressure` is reserved and raises NotImplementedError | `body/jets.py` `Jet.__post_init__` |
+| Jet geometry | each Jet selects resolved links of exactly one stationary body; no Jet/Jet or Jet/SurfaceActuation overlap; abs(amplitude) < 1 lattice | `body/jets.py` `Jet.__post_init__`, `jet_link_shapes`, `_surface_links` |
+| Jets with shifted storage | only in the all-blocked stepper with SRT and no LES | `body/jets.py` `_surface_links` |
+| SurfaceActuation | unshifted storage only | `body/objects.py` `SimObject.validate_surface_actuation` |
+| Raw boundary exchange (force) | blocked body levels only; the window ledger requires `observe_raw_exchange` | `lbm/refinement_stepper.py` `FixedPatchStepper.__init__` |
+| Checkpoint with Jets | refused | `simulation.py` `Simulation.save_checkpoint` |

 Raw exchange is the cut-link momentum reaction (wall + aperture), not a validated
 nozzle thrust or stress load. Jet amplitude is a lattice velocity, not a flow rate.
@@ -11,7 +11,7 @@ Responsibilities:

 Design::
  Packed ``obs`` layout matches ``generate_config(..., n_objects=count)`` and
-  ``config_obs.h`` macros. Host sizing uses :func:`~CelerisLab.cuda.compiler_v2.obs_layout`.
+  ``config_obs.h`` macros. Host sizing uses :func:`~CelerisLab.cuda.compiler.obs_layout`.

  Pending edits are lightweight: only the "intent" (add / remove) is stored.
  All derived quantities (flags, compact lists, GPU buffers) are built at
@@ -28,7 +28,7 @@ import pycuda.driver as cuda
 from .objects import SimObject
 from .registry import BodyRegistry
 from .preprocess.flag_overlay import merge_flag_masks, validate_cut_link_ownership
-from ..cuda.compiler_v2 import obs_layout
+from ..cuda.compiler import obs_layout
 from ..cuda.obs_layout import ObsLayout
 from ..lbm.descriptors import (
    FLUID, SOLID,
@@ -7,7 +7,7 @@ Two JSON files:
  config_body.json – object list

 LBMConfig / BodyConfig are plain dataclasses; they translate to
-macro dicts consumed by `cuda/compiler_v2.py` for header generation.
+macro dicts consumed by `cuda/compiler.py` for header generation.
 """

 import json
@@ -60,14 +60,14 @@ Python `config.py` 只负责读取和校验，不是配置位置。显式配置

 **`lbm/kernels/config/*.h` 由编译器根据 `LBMConfig` 自动生成，请勿手改。**

-仓库中检入的 `config/*.h` 可能是历史版本。`compiler_v2.generate_config` 会复制 kernel 源码到独立临时构建目录，并在副本中生成配置；不会改写包内头文件或 PTX。返回目录须显式传给 `compile_kernel(build_dir=...)`，编译结果与配置摘要均保存在该目录。`Simulation` 负责目录清理；直接使用编译器的调用者须调用 `cleanup_build_dir()`。跳过生成配置直接编译仓库中的 `kernel_v2.cu` 属不受支持用法。
+仓库中检入的 `config/*.h` 可能是历史版本。`compiler.generate_config` 会复制 kernel 源码到独立临时构建目录，并在副本中生成配置；不会改写包内头文件或 PTX。返回目录须显式传给 `compile_kernel(build_dir=...)`，编译结果与配置摘要均保存在该目录。`Simulation` 负责目录清理；直接使用编译器的调用者须调用 `cleanup_build_dir()`。跳过生成配置直接编译仓库中的 `uniform.cu` 属不受支持用法。

 ```
 config_lbm.json
  ↓  load_lbm_config()
 config.py:LBMConfig       (读取 + 校验)
  ↓  to_macros()
-cuda/compiler_v2.py       (独立构建目录中生成 config/*.h)
+cuda/compiler.py          (独立构建目录中生成 config/*.h)
  ↓
 config_grid.h             → NX, NY, NZ, NT, LATTICE_MODEL；
                            DIM / NQ 由 LATTICE_MODEL 推导写入（与 grid.lattice_model 一致，非 JSON 顶层字段）
@@ -1,4 +1,4 @@
 """CUDA build helpers (compilation itself does not need a live context)."""
-from . import compiler_v2 as compiler
+from . import compiler

 __all__ = ["compiler"]
@@ -1,11 +1,14 @@
-# CelerisLab/cuda/compiler_v2.py
+# CelerisLab/cuda/compiler.py
 """
 Kernel configuration, compilation, and module loading.

-Workflow:
+Owners call build_module(cfg, entry, arch=...), then set_params(module, omega, ...).
+build_module runs, in a private directory it removes afterwards:
  1. build_dir = generate_config(lbm_cfg, n_objects) — isolated source snapshot
-  2. compile_kernel(arch, build_dir=build_dir)       — nvcc -ptx kernel_v2.cu
+  2. compile_kernel(arch, source=entry, build_dir=build_dir) — nvcc -ptx
  3. load_module(ptx_path)              — PyCUDA module from PTX
+Entries: uniform.cu (uniform grid), blocked.cu (blocked level), refinement.cu
+(dense refinement level).
 """

 from __future__ import annotations
@@ -16,7 +19,6 @@ import shutil
 import tempfile
 import hashlib
 import json
-from contextvars import ContextVar
 from dataclasses import asdict
 from pathlib import Path
 from typing import Optional
@@ -33,9 +35,6 @@ _KERNEL_DIR = os.path.join(
 )


-_last_build_dir = ContextVar("celeris_last_build_dir", default=None)
-_last_ptx_path = ContextVar("celeris_last_ptx_path", default=None)
-

 def kernel_path(filename: str) -> str:
    return os.path.join(_KERNEL_DIR, filename)
@@ -54,14 +53,16 @@ def _write(path: str, content: str):


 def generate_config(cfg: LBMConfig, n_objects: int = 0,
-                    *, build_root: Optional[str] = None) -> str:
+                    *, build_root: Optional[str] = None, block_root_bc: int = 0) -> str:
    """Snapshot kernels and render headers in a fresh private directory.

-    Returns its path. Pass it explicitly to compile_kernel for interleaved
-    builds. Legacy generate_config(); compile_kernel() remains context-local.
-    The caller owns cleanup_build_dir() once the loaded module is no longer
+    Returns its path; pass it to compile_kernel(build_dir=...). The caller owns cleanup_build_dir() once the loaded module is no longer
    needed for recompilation/debugging. No package files are changed.
+    block_root_bc specializes a blocked root module; fine/dense builds use 0.
    """
+    if (isinstance(block_root_bc, bool) or not isinstance(block_root_bc, int)
+            or block_root_bc not in (0, 1, 2)):
+        raise ValueError("block_root_bc must be an integer mode 0, 1 or 2")
    macros = cfg.to_macros()
    if isinstance(n_objects, bool) or not isinstance(n_objects, int) or n_objects < 0:
        raise ValueError("n_objects must be a nonnegative integer")
@@ -71,20 +72,19 @@ def generate_config(cfg: LBMConfig, n_objects: int = 0,
    try:
        shutil.copytree(_KERNEL_DIR, build_dir, dirs_exist_ok=True,
                        ignore=shutil.ignore_patterns("*.ptx", "*.cubin", "*.tmp"))
-        _render_config(macros, n_objects, os.path.join(build_dir, "config"))
+        _render_config(macros, n_objects, os.path.join(build_dir, "config"),
+                       block_root_bc=block_root_bc)
        digest = hashlib.sha256()
        for path in sorted(Path(build_dir).rglob("*")):
            if path.is_file():
                digest.update(str(path.relative_to(build_dir)).encode())
                digest.update(path.read_bytes())
        metadata = {"source_sha256": digest.hexdigest(), "config": asdict(cfg),
-                    "n_objects": n_objects}
+                    "n_objects": n_objects, "block_root_bc": block_root_bc}
        _write(os.path.join(build_dir, ".celeris-build"), json.dumps(metadata, sort_keys=True))
    except BaseException:
        shutil.rmtree(build_dir)
        raise
-    _last_build_dir.set(build_dir)
-    _last_ptx_path.set(None)
    return build_dir


@@ -96,12 +96,9 @@ def cleanup_build_dir(build_dir: Optional[str]) -> None:
    if path == Path(_KERNEL_DIR).resolve() or not (path / ".celeris-build").is_file():
        raise ValueError("Refusing to remove a directory not owned by the compiler")
    shutil.rmtree(path)
-    if _last_build_dir.get() == str(build_dir):
-        _last_build_dir.set(None)
-        _last_ptx_path.set(None)


-def _render_config(m: dict, n_objects: int, cdir: str):
+def _render_config(m: dict, n_objects: int, cdir: str, *, block_root_bc: int = 0):
    _write(os.path.join(cdir, "config_grid.h"), f"""{_HEADER}\
 // Layer: Global / Grid
 #ifndef CELERIS_CONFIG_GRID_H
@@ -165,6 +162,9 @@ def _render_config(m: dict, n_objects: int, cdir: str):
 #define STORE_PRECISION  {m['STORE_PRECISION']}
 #define USE_DDF_SHIFTING {m['USE_DDF_SHIFTING']}

+// Internal blocked root mode: 0=periodic/fine, 1=slot/reservoir, 2=channel.
+#define BLOCK_ROOT_BC {block_root_bc}
+
 #define USE_LES {m['USE_LES']}
 #define LES_CS  {m['LES_CS']}
 #define LES_CLOSED_FORM {m['LES_CLOSED_FORM']}
@@ -241,13 +241,9 @@ def _render_config(m: dict, n_objects: int, cdir: str):
 # ---------------------------------------------------------------------------
 # Compilation
 # ---------------------------------------------------------------------------
-def compile_kernel(arch: str = "sm_70", source: str = "kernel_v2.cu",
-                   output: Optional[str] = None, *,
-                   build_dir: Optional[str] = None) -> str:
-    """Compile a generated private source snapshot; never use package headers."""
-    # A failed attempt must not make implicit load_module() reuse older code.
-    _last_ptx_path.set(None)
-    build_dir = build_dir or _last_build_dir.get()
+def compile_kernel(arch: str = "sm_70", source: str = "uniform.cu",
+                   output: Optional[str] = None, *, build_dir: str) -> str:
+    """Compile the snapshot in build_dir (from generate_config); never package headers."""
    if not build_dir or not (Path(build_dir) / ".celeris-build").is_file():
        raise RuntimeError("Call generate_config() and pass its build_dir before compiling")
    root = Path(build_dir).resolve()
@@ -272,23 +268,36 @@ def compile_kernel(arch: str = "sm_70", source: str = "kernel_v2.cu",
    finally:
        if os.path.exists(tmp):
            os.unlink(tmp)
-    _last_ptx_path.set(str(dst))
    return str(dst)


-def load_module(ptx_path: Optional[str] = None):
-    """Load explicit PTX or the last PTX compiled in this Python context."""
-    ptx_path = ptx_path or _last_ptx_path.get()
-    if not ptx_path:
-        raise RuntimeError("No compiled PTX; call compile_kernel() first")
+def load_module(ptx_path: str):
+    """Load one compiled PTX file as a new PyCUDA module."""
    import pycuda.driver as cuda
    return cuda.module_from_file(ptx_path)


-# ---------------------------------------------------------------------------
-# Back-compat helpers (used by test scripts)
-# ---------------------------------------------------------------------------
+def build_module(cfg: LBMConfig, source: str, *, arch: str, n_objects: int = 0,
+                 block_root_bc: int = 0):
+    """Generate, compile and load one private module; the build directory is removed.

-def compile_kernel_v2(arch: str = "sm_70") -> str:
-    """Alias for compile_kernel() kept for test-script compatibility."""
-    return compile_kernel(arch=arch)
+    The single build path of every owner (Simulation, FixedPatchStepper levels).
+    Each call returns a NEW module: callers must not share it, because its
+    globals (d_params, ...) are mutable per module.
+    """
+    build_dir = generate_config(cfg, n_objects=n_objects, block_root_bc=block_root_bc)
+    try:
+        return load_module(compile_kernel(arch, source=source, build_dir=build_dir))
+    finally:
+        cleanup_build_dir(build_dir)
+
+
+def set_params(module, omega: float, fx: float = 0.0, fy: float = 0.0, fz: float = 0.0,
+               n_objects: int = 0) -> None:
+    """Write the module's __constant__ LBMParams (core/params.cuh layout)."""
+    import struct
+    import pycuda.driver as cuda
+    ptr, size = module.get_global("d_params")
+    if size != struct.calcsize("=ffffI"):
+        raise RuntimeError("LBMParams host/device layout mismatch")
+    cuda.memcpy_htod(ptr, struct.pack("=ffffI", omega, fx, fy, fz, n_objects))
@@ -2,7 +2,7 @@
 """Packed ``obs`` buffer layout for force / torque / sensor telemetry.

 Usage::
-    Host code (e.g. ``ObjectManager``) and ``compiler_v2.generate_config`` must use
+    Host code (e.g. ``ObjectManager``) and ``compiler.generate_config`` must use
    the same ``dim`` and ``n_objects`` so allocation sizes match device macros.

 Design::
@@ -1,296 +0,0 @@
-# CelerisLab/lbm/block_stepper.py
-"""Opt-in periodic backend for one BlockLayout2D level, optional static circles.
-
-D2Q9/SRT/FP32 unforced double buffers. Each step refreshes every halo from
-its owner core, then pulls/collides core cells into the spare buffer, then
-swaps. Only halo pair tables are uploaded; dense ownership stays on the host.
-The caller owns the CudaContext and must close this backend before it.
-Optional ``objects`` are stationary passive circles (static linear BFL via
-receiver-local staging: halo refresh -> cut links -> pull, no copy between).
-No force/actuation readout, motion, walls, refinement, checkpoint or runtime
-parameter mutation; wall/slot/inlet fields in cfg are ignored.
-"""
-from dataclasses import replace
-from numbers import Integral
-import struct
-from threading import RLock
-
-import numpy as np
-import pycuda.driver as cuda
-
-from ..cuda import compiler_v2 as compiler
-from .blocks import NQ, BlockLayout2D, blocked_link_tables
-from .descriptors import D2Q9_EX, D2Q9_EY, FLUID, SOLID
-from .population_storage import _StoredPopulations, _encode_physical, _decode_stored
-
-MAX_GRID_DIM = 65535  # CUDA gridDim.y/z limit (rows, blocks)
-
-
-def _static_circles(layout, objects):
-    """Host tables for stationary circles; rejects before any CUDA work.
-
-    Every periodic neighbor of every fluid core is inspected, including domain
-    edges that CircleGeometry does not enumerate, so clipped or incomplete
-    geometry cannot pass. Each physical link has one receiver-owning block.
-    """
-    from ..body.geometry.circle import CircleGeometry
-
-    ny, nx = layout.shape
-    flags = np.full(ny * nx, FLUID, np.uint16)
-    for obj in objects:
-        if (type(obj.geometry) is not CircleGeometry or obj.is_sensor or obj.is_force_region
-                or obj.surface_actuation is not None
-                or any(getattr(obj.state, name) != 0 for name in ("vx", "vy", "omega"))
-                or any(vars(obj.control).values())):
-            raise ValueError("Block bodies must be stationary passive circles")
-        cx, cy = obj.geometry.center
-        radius = obj.geometry._radius
-        if not np.isfinite([cx, cy, radius]).all() or radius <= 0:
-            raise ValueError("Block circle center/radius must be finite with radius > 0")
-        mask = obj.geometry.build_flag_mask(nx, ny)
-        solid = (mask & SOLID) != 0
-        if not solid.any() or np.any(flags[solid] & SOLID):
-            raise ValueError("Block circles must be resolved and disjoint")
-        flags[solid] = mask[solid]
-    solid = ((flags & SOLID) != 0).reshape(ny, nx)
-    if solid[[0, -1]].any() or solid[:, [0, -1]].any():
-        raise ValueError("Block circles must not touch the domain edge")
-    links = [link for obj in objects
-             for link in obj.geometry.build_cut_links(nx, ny, domain_flags=flags)]
-    def source(r, q):
-        y, x = divmod(r, nx)
-        return (y + D2Q9_EY[q]) % ny, (x + D2Q9_EX[q]) % nx
-    receivers = {(int(link.fluid_idx), int(link.dir)) for link in links}
-    required = {(int(r), q) for r in np.flatnonzero(~solid.ravel())
-                for q in range(1, NQ) if solid[source(r, q)]}
-    if len(receivers) != len(links) or receivers != required:
-        raise ValueError("Cut links must cover every fluid pull from a solid exactly once")
-    tables, _ = blocked_link_tables(
-        layout, flags, np.array([link.fluid_idx for link in links], np.uint32),
-        np.array([link.dir for link in links], np.uint8),
-        np.array([link.q for link in links], np.float32),
-        np.array([link.fallback for link in links], np.uint8))
-    return dict(solid=solid, **tables)
-
-
-class PeriodicBlockStepper:
-    """G->G stepping of Q-major (9, ny, nx) postcollision populations.
-
-    On a device/launch failure the instance becomes unusable: close and
-    recreate it rather than assuming a partially advanced buffer rolled back.
-    """
-    def __init__(self, cfg, layout, *, context, objects=()):
-        if not isinstance(layout, BlockLayout2D) or layout.periodic != (True, True):
-            raise ValueError("PeriodicBlockStepper requires a doubly periodic BlockLayout2D")
-        checked = replace(cfg)
-        checked.validate()
-        if (not checked.is_d2q9 or checked.collision != "SRT"
-                or checked.streaming != "double_buffer" or checked.store_precision != "FP32"
-                or checked.les_enabled
-                or (checked.ny, checked.nx) != layout.shape):
-            raise ValueError("PeriodicBlockStepper requires matching D2Q9/SRT/FP32/"
-                             "double_buffer config without LES")
-        ly, lx = layout.local_shape
-        pairs = len(layout.halo_dst)
-        tpb = checked.threads_per_block
-        if layout.nblocks > MAX_GRID_DIM or ly > MAX_GRID_DIM or not 0 < pairs < 2**31:
-            raise ValueError("block count, block rows or halo pairs exceed launch limits")
-        objects = tuple(objects)  # an empty iterable takes the body-free path
-        self.boundary = _static_circles(layout, objects) if objects else None
-        links = 0 if self.boundary is None else len(self.boundary["cell"])
-        if links >= 2**31:
-            raise ValueError("cut links exceed launch limits")
-        self._ddf_shifting = bool(checked.ddf_shifting)
-        self.layout = layout
-        self.ctx = context
-        self._links = links
-        self._link_grid = ((links + tpb - 1) // tpb, 1)
-        # Wall/slot fields are not applied; neutralize them like the fine patch level.
-        self._local_cfg = replace(checked, nx=lx, ny=ly, y_wall_bc="bounce_back", top_wall_bc="velocity",
-                                  wall_slot_width=0, wall_normal_velocity=0.0, velocity=0.0)
-        self._block = (tpb, 1, 1)
-        self._grid = ((lx + tpb - 1) // tpb, ly, layout.nblocks)
-        self._copy_grid = ((pairs + tpb - 1) // tpb, 1)
-        self._bytes = int(np.prod(layout.blocked_shape)) * 4
-        self._lock = RLock()
-        self._status = "new"
-        self._step_count = 0
-        self._allocations = []
-        self._stream = self._module = self._copy = self._step = self._curved = None
-        self._tables = {}
-        with self.ctx.activate():
-            try:
-                free, _ = cuda.mem_get_info()
-                tables = sum(a.nbytes for k, a in (self.boundary or {}).items() if k != "solid")
-                if 2 * self._bytes + layout.halo_dst.nbytes + layout.halo_src.nbytes + tables > free:
-                    raise MemoryError("blocked buffers exceed free device memory")
-                self._stream = cuda.Stream()
-                config = self._local_cfg
-                config.validate()
-                if not np.isclose(config.omega, 1 / (3 * checked.viscosity + .5), rtol=1e-14, atol=0):
-                    raise RuntimeError("block omega/tau mismatch")
-                build = compiler.generate_config(config)
-                try:
-                    ptx = compiler.compile_kernel(self.ctx.sm_arch, source="blocks.cu", build_dir=build)
-                    module = compiler.load_module(ptx)
-                finally:
-                    compiler.cleanup_build_dir(build)
-                ptr, size = module.get_global("d_params")
-                if size != struct.calcsize("=ffffI"):
-                    raise RuntimeError("LBMParams host/device layout mismatch")
-                cuda.memcpy_htod(ptr, struct.pack("=ffffI", config.omega, 0, 0, 0, 0))
-                self._module = module
-                self._current = self._allocate(self._bytes)
-                self._spare = self._allocate(self._bytes)
-                self._dst = self._upload(layout.halo_dst)
-                self._src = self._upload(layout.halo_src)
-                self._copy = module.get_function("BlockHaloCopy")
-                self._step = module.get_function("BlockStep")
-                if self.boundary is not None:
-                    self._tables = {k: self._upload(a) for k, a in self.boundary.items() if k != "solid"}
-                    self._curved = module.get_function("BlockCurvedBoundary")
-            except BaseException:
-                try:
-                    self.close()
-                except Exception:
-                    pass  # Preserve the original construction failure.
-                raise
-
-    def _allocate(self, size):
-        pointer = cuda.mem_alloc(size)
-        self._allocations.append(pointer)
-        return pointer
-
-    def _upload(self, array):
-        array = np.ascontiguousarray(array)
-        pointer = self._allocate(array.nbytes)
-        cuda.memcpy_htod(pointer, array)
-        return pointer
-
-    def _require(self, *states):
-        if self._status not in states:
-            raise RuntimeError(f"Block stepper is {self._status}; expected {' or '.join(states)}")
-
-    @property
-    def step_count(self):
-        return self._step_count
-
-    def initialize(self, g):
-        """Upload dense G_0; the spare buffer starts NaN so stale halos cannot hide."""
-        with self._lock, self.ctx.activate():
-            self._require("new", "ready")
-            dense = np.array(g, dtype=np.float64 if self._ddf_shifting else np.float32,
-                             order="C", copy=True)
-            if dense.shape != (NQ, *self.layout.shape):
-                raise ValueError("Expected a Q-major (9, ny, nx) array")
-            fluid = np.ones(self.layout.shape, bool) if self.boundary is None else ~self.boundary["solid"]
-            if not np.isfinite(dense[:, fluid]).all() or (dense[:, fluid].sum(axis=0) <= 0).any():
-                raise ValueError("Fluid populations must be finite with positive density")
-            dense[:, ~fluid] = np.nan  # Solid cores are deliberately invalid.
-            encoded = _encode_physical(dense, self._ddf_shifting)
-            values = encoded[:, fluid]
-            if (not np.isfinite(values).all()
-                    or (values.sum(axis=0, dtype=np.float64) + int(self._ddf_shifting) <= 0).any()):
-                raise ValueError("Stored fluid populations must be finite with positive density")
-            blocked = self.layout.to_blocked(encoded)
-            spare = np.full_like(blocked, np.nan)
-            try:
-                cuda.memcpy_htod(self._current, blocked)
-                cuda.memcpy_htod(self._spare, spare)
-                self._step_count = 0
-                self._status = "ready"
-            except BaseException:
-                self._status = "failed"
-                raise
-
-    def step(self, n=1):
-        """Advance n steps: halo refresh, then core pull/collide, per step.
-
-        All n steps are enqueued in order on the one stream and the host waits
-        once before returning (same launches and bits as n step(1) calls).
-        Asynchronous device errors surface at that single wait, so they are not
-        attributed to the step that caused them; call step(1) repeatedly to
-        localize one. On any failure the instance is failed and step_count keeps
-        its entry value.
-        """
-        if isinstance(n, bool) or not isinstance(n, Integral) or n < 0:
-            raise ValueError("n must be a nonnegative integer")
-        with self._lock, self.ctx.activate():
-            self._require("ready")
-            if n == 0:
-                return
-            kw = dict(block=self._block, stream=self._stream)
-            try:
-                self._status = "advancing"
-                for _ in range(n):
-                    self._copy(self._current, self._dst, self._src,
-                               np.uint32(len(self.layout.halo_dst)), grid=self._copy_grid, **kw)
-                    if self._curved is not None:
-                        t = self._tables
-                        self._curved(self._current, t["block"], t["cell"], t["dir"], t["q"],
-                                     t["fallback"], np.uint32(self._links), np.uintp(0),  # no actuation
-                                     np.uintp(0), np.uintp(0), np.uintp(0),  # no observation
-                                     grid=self._link_grid, **kw)
-                    self._step(self._current, self._spare, np.uint32(self.layout.halo),
-                               self._tables.get("flags", np.uintp(0)), np.uintp(0),  # no kind
-                               grid=self._grid, **kw)
-                    self._current, self._spare = self._spare, self._current  # host handles only
-                self._stream.synchronize()
-                self._step_count += n
-                self._status = "ready"
-            except BaseException:
-                self._status = "failed"
-                raise
-
-    def download(self):
-        """Return physical core populations (FP64 when shifted); no host cache."""
-        snapshot = self._download_raw()
-        return _decode_stored(snapshot.values[0], snapshot.shifted)
-
-    def _download_raw(self):
-        """Return one tagged dense FP32 population array without conversion."""
-        with self._lock, self.ctx.activate():
-            self._require("ready")
-            blocked = np.empty(self.layout.blocked_shape, np.float32)
-            try:
-                cuda.memcpy_dtoh(blocked, self._current)
-            except BaseException:
-                self._status = "failed"
-                raise
-            return _StoredPopulations((self.layout.to_dense(blocked),), self._ddf_shifting)
-
-    def close(self):
-        """Release owned resources; retain them for retry if context activation fails."""
-        with self._lock:
-            if self._status == "closed":
-                return
-            self._status = "failed"
-            with self.ctx.activate():
-                self._status = "closed"
-                error = None
-                try:
-                    if self._stream is not None:
-                        self._stream.synchronize()
-                except Exception as exc:
-                    error = exc
-                self._copy = self._step = self._curved = self._module = None
-                self._tables = {}
-                for pointer in self._allocations:
-                    try:
-                        pointer.free()
-                    except Exception as exc:
-                        error = error or exc
-                self._allocations.clear()
-                self._stream = None
-                if error is not None:
-                    raise error
-
-    def __enter__(self):
-        return self
-
-    def __exit__(self, *exc):
-        try:
-            self.close()
-        except Exception:
-            if exc[0] is None:
-                raise
@@ -31,7 +31,7 @@ receiver-local boundary values: the solid's owner-core copy of that slot may
 be stale, which is distinct from the unique core ownership of stored state.
 No moving geometry is implied.
 """
-from dataclasses import dataclass
+from dataclasses import dataclass, replace
 from numbers import Integral
 from types import SimpleNamespace

@@ -373,3 +373,84 @@ def blocked_link_tables(layout, flags, fluid_idx, dir, q, fallback):
                  dir=columns[1].astype(np.uint8), q=columns[2].astype(np.float32),
                  fallback=columns[3].astype(np.uint8))
    return tables, mapped
+
+
+# One blocked level as FixedPatchStepper builds and launches it (any depth). Raw argument orders of the blocked.cu kernels live only here.
+
+def level_config(cfg, layout, tau):
+    """Kernel config of one blocked level: NX/NY = block window incl. halo.
+
+    Physical wall/slot/inlet fields are neutralized; a blocked root gets them
+    from its module globals (d_block_wall/d_block_channel) instead.
+    """
+    ly, lx = layout.local_shape
+    local = replace(cfg, nx=lx, ny=ly, y_wall_bc="bounce_back", top_wall_bc="velocity",
+                    wall_slot_width=0, wall_normal_velocity=0.0, velocity=0.0)
+    local.validate()
+    if not np.isclose(local.omega, 1 / tau, rtol=1e-14, atol=0):
+        raise RuntimeError("block level omega/tau mismatch")
+    return local
+
+
+def launch_halo_copy(kernel, buffer, dst, src, count, *, block, stream):
+    """BlockHaloCopy over ``count`` uploaded (dst, src) block-cell pairs."""
+    kernel(buffer, dst, src, np.uint32(count),
+           grid=((count + block[0] - 1) // block[0], 1), block=block, stream=stream)
+
+
+def launch_curved_boundary(kernel, buffer, tables, count, *, block, stream,
+                           velocity=None, rx=None, ry=None, raw=None):
+    """BlockCurvedBoundary on uploaded link tables; None = no actuation/observation."""
+    none = np.uintp(0)
+    kernel(buffer, tables["block"], tables["cell"], tables["dir"], tables["q"], tables["fallback"],
+           np.uint32(count), none if velocity is None else velocity,
+           none if rx is None else rx, none if ry is None else ry, none if raw is None else raw,
+           grid=((count + block[0] - 1) // block[0], 1), block=block, stream=stream)
+
+
+def static_circle_tables(layout, objects):
+    """Host tables for stationary circles; rejects before any CUDA work.
+
+    Every periodic neighbor of every fluid core is inspected, including domain
+    edges that CircleGeometry does not enumerate, so clipped or incomplete
+    geometry cannot pass. Each physical link has one receiver-owning block.
+    """
+    from ..body.geometry.circle import CircleGeometry
+    from .descriptors import FLUID, SOLID
+
+    ny, nx = layout.shape
+    flags = np.full(ny * nx, FLUID, np.uint16)
+    for obj in objects:
+        if (type(obj.geometry) is not CircleGeometry or obj.is_sensor or obj.is_force_region
+                or obj.surface_actuation is not None
+                or any(getattr(obj.state, name) != 0 for name in ("vx", "vy", "omega"))
+                or any(vars(obj.control).values())):
+            raise ValueError("Block bodies must be stationary passive circles")
+        cx, cy = obj.geometry.center
+        radius = obj.geometry._radius
+        if not np.isfinite([cx, cy, radius]).all() or radius <= 0:
+            raise ValueError("Block circle center/radius must be finite with radius > 0")
+        mask = obj.geometry.build_flag_mask(nx, ny)
+        solid = (mask & SOLID) != 0
+        if not solid.any() or np.any(flags[solid] & SOLID):
+            raise ValueError("Block circles must be resolved and disjoint")
+        flags[solid] = mask[solid]
+    solid = ((flags & SOLID) != 0).reshape(ny, nx)
+    if solid[[0, -1]].any() or solid[:, [0, -1]].any():
+        raise ValueError("Block circles must not touch the domain edge")
+    links = [link for obj in objects
+             for link in obj.geometry.build_cut_links(nx, ny, domain_flags=flags)]
+    def source(r, q):
+        y, x = divmod(r, nx)
+        return (y + D2Q9_EY[q]) % ny, (x + D2Q9_EX[q]) % nx
+    receivers = {(int(link.fluid_idx), int(link.dir)) for link in links}
+    required = {(int(r), q) for r in np.flatnonzero(~solid.ravel())
+                for q in range(1, NQ) if solid[source(r, q)]}
+    if len(receivers) != len(links) or receivers != required:
+        raise ValueError("Cut links must cover every fluid pull from a solid exactly once")
+    tables, _ = blocked_link_tables(
+        layout, flags, np.array([link.fluid_idx for link in links], np.uint32),
+        np.array([link.dir for link in links], np.uint8),
+        np.array([link.q for link in links], np.float32),
+        np.array([link.fallback for link in links], np.uint8))
+    return dict(solid=solid, **tables)
@@ -0,0 +1,49 @@
+// CelerisLab/lbm/kernels/blocked.cu
+// One fixed-size blocked level (lbm/blocks.py layout), specialized at compile
+// time for periodic/fine, slot/reservoir or channel stepping. Separate
+// compilation entry: NX, NY are the block window including halo, so index_f
+// and linear_index address one block; blockIdx.z selects the block, whose
+// populations start at b*NQ*TOTAL_CELLS (block-contiguous, Q-major inside).
+#include <stdint.h>
+#include <math.h>
+#include "config.h"
+#include "core/precision.cuh"
+#include "core/descriptors.cuh"
+#include "core/layout.cuh"
+#include "core/params.cuh"
+#include "core/flags.cuh"
+#include "core/obs.cuh"
+#include "core/cell_kind.cuh"
+#include "operators/macro.cuh"
+#include "operators/collision_srt.cuh"
+#if COLLISION_MODEL == 2
+#include "operators/collision_mrt.cuh"
+#endif
+#include "streaming/pull_double_buffer.cuh"
+#include "boundary/curved_boundary.cuh"
+
+#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
+#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
+#endif
+#ifndef BLOCK_ROOT_BC
+#error "Blocks require a generated BLOCK_ROOT_BC mode"
+#elif BLOCK_ROOT_BC < 0 || BLOCK_ROOT_BC > 2
+#error "BLOCK_ROOT_BC must be 0, 1 or 2"
+#endif
+
+#if BLOCK_ROOT_BC == 1
+#include "boundary/block_wall.cuh"
+#endif
+#if BLOCK_ROOT_BC != 2
+#include "operators/block_collision.cuh"
+#endif
+
+// Kernels, in definition order: transfers, boundary staging, then the one full
+// step selected for this module (BlockStep, or BlockChannelStep for a channel root).
+#include "step/block_transfer.cu"
+#include "step/block_boundary.cu"
+#if BLOCK_ROOT_BC != 2
+#include "step/block_step.cu"
+#else
+#include "step/block_channel.cu"
+#endif
@@ -1,200 +0,0 @@
-// CelerisLab/lbm/kernels/blocks.cu
-// One periodic level of fixed-size blocks (lbm/blocks.py layout). Separate
-// compilation entry: NX, NY are the block window including halo, so index_f
-// and linear_index address one block; blockIdx.z selects the block, whose
-// populations start at b*NQ*TOTAL_CELLS (block-contiguous, Q-major inside).
-#include <stdint.h>
-#include <math.h>
-#include "config.h"
-#include "core/precision.cuh"
-#include "core/descriptors.cuh"
-#include "core/layout.cuh"
-#include "core/params.cuh"
-#include "core/flags.cuh"
-#include "core/obs.cuh"
-#include "operators/macro.cuh"
-#include "operators/collision_srt.cuh"
-#if COLLISION_MODEL == 2
-#include "operators/collision_mrt.cuh"
-#endif
-#include "streaming/pull_double_buffer.cuh"
-#include "boundary/curved_boundary.cuh"
-#include "boundary/planar_velocity.cuh"
-
-#if DIM != 2 || NQ != 9 || (COLLISION_MODEL != 0 && COLLISION_MODEL != 2) || STORE_PRECISION != 0 || USE_LES || STREAMING_MODEL != 0
-#error "Blocks require FP32 D2Q9 SRT or MRT double buffers without LES"
-#endif
-
-// Fixed-patch cell kinds, matching refinement.py ACTIVE/INNER_GHOST/OUTER_GHOST.
-constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
-
-// Zero-initialized for ordinary blocks/fine levels. Set only for parent walls;
-// NX/NY remain local storage strides, never physical wall coordinates.
-struct BlockWall { unsigned int nx, ny, slot_width; float amplitude; };
-extern "C" { __constant__ BlockWall d_block_wall; }
-
-// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
-extern "C" __global__ void BlockHaloCopy(
-    fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t d = dst[row], s = src[row];
-    fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    for (unsigned int q = 0; q < NQ; ++q)
-        to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
-}
-
-// Fixed-patch explosion into blocked fine storage: target = owner block cell
-// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
-// storage is the one-window case with coarse_cells = full domain size.
-extern "C" __global__ void BlockExplode(
-    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
-    unsigned int count, unsigned long coarse_cells)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t t = target[row];
-    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    const uint64_t d = donor[row];
-    const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
-    for (unsigned int q = 0; q < NQ; ++q)
-        to[index_f(t % TOTAL_CELLS, q)] = from[q * coarse_cells + d % coarse_cells];
-}
-
-// Experimental Schornbaum--Ruede (arXiv:1508.07982, Algorithm 2)
-// and Chen et al. (doi:10.1016/j.physa.2005.09.036) population interpolation.
-// Independent implementation of the mathematical stencil; no upstream source copied.
-// Stencil cells are owner cores; a disabled axis repeats its center twice.
-extern "C" __global__ void BlockExplodeProjected(
-    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
-    unsigned int count, unsigned long coarse_cells,
-    const uint64_t* neighbors, const float* offsets)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t t = target[row], d = donor[row];
-    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
-    for (unsigned int q = 0; q < NQ; ++q) {
-        float v[4];
-        for (unsigned int a = 0; a < 4; ++a) {
-            const uint64_t n = neighbors[4ul * row + a];
-            v[a] = coarse[(n / coarse_cells) * NQ * coarse_cells
-                          + q * coarse_cells + n % coarse_cells];
-        }
-        float gx = .5f * (v[1] - v[0]), gy = .5f * (v[3] - v[2]);
-        const float ex = d_cx[q], ey = d_cy[q];
-        if (q != 0) {
-            const float along = (ex * gx + ey * gy) / (ex * ex + ey * ey);
-            gx -= ex * along;
-            gy -= ey * along;
-        }
-        to[index_f(t % TOTAL_CELLS, q)] =
-            coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
-            + offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
-    }
-}
-
-// Targets are receiver-local SOURCE slots for the immediately following pull.
-// Keep child summation order identical to PatchCoarseStep; no collision here.
-extern "C" __global__ void BlockCoalesce(
-    const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
-    unsigned int count)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const uint64_t* s = children + 4ul * row;
-    coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
-}
-
-// Static cut links, receiver-local staging: cell is the receiver's local core
-// cell, so the existing BFL writes the solid-source window image that this
-// block's next pull reads (no halo copy between this kernel and BlockStep).
-// surface_velocity (nullable, 2 floats per link): prescribed fluid velocity on
-// fixed geometry, same convention as CurvedBoundaryKernel. Optional raw rows
-// (jx,jy,torque,dm) are unique physical links, never body-observation slots.
-extern "C" __global__ void BlockCurvedBoundary(
-    fpxx* f, const uint32_t* block, const uint32_t* cell, const unsigned char* dir,
-    const float* q, const unsigned char* fallback, unsigned int count,
-    const float* surface_velocity, const float* rx, const float* ry, float* raw)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    float Uw = 0.0f, Vw = 0.0f;
-    if (surface_velocity != nullptr) {
-        const float u = surface_velocity[2 * row], v = surface_velocity[2 * row + 1];
-        if (u != 0.0f || v != 0.0f) { Uw = u; Vw = v; }
-    }
-    apply_bouzidi_link(dir[row], q[row], cell[row], f + (unsigned long)block[row] * NQ * TOTAL_CELLS,
-                       Uw, Vw, rx ? rx[row] : 0.0f, ry ? ry[row] : 0.0f,
-                       fallback[row], nullptr, -1, raw ? raw + 4ul * row : nullptr,
-                       raw ? raw + 4ul * row + 3 : nullptr);
-}
-
-// Same-stream BFL -> this row-owned accumulation -> pull step. No atomics,
-// no cross-row reduction and no history allocation proportional to substeps.
-extern "C" __global__ void BlockAccumulateExchange(
-    const float* raw, double* integral, unsigned int count, double h, double dt)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    const unsigned long k = 4ul * row;
-    const double impulse = h * dt;
-    integral[k] += impulse * (double)raw[k];
-    integral[k + 1] += impulse * (double)raw[k + 1];
-    integral[k + 2] += h * impulse * (double)raw[k + 2];
-    integral[k + 3] += h * h * (double)raw[k + 3];
-}
-
-// Pull/collide/store core cells only; halo cells of out are left untouched.
-// flags (block cell ids, nullable): solid cores store NaN and never collide.
-// kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
-// stores NaN, INNER pulls without collision, ACTIVE pulls and collides.
-extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int halo,
-                                     const uint16_t* flags, const unsigned char* kind) {
-    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 long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
-    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
-    const unsigned char cell_kind = kind != nullptr ? kind[cell] : KIND_ACTIVE;
-    float f[NQ];
-    if (cell_kind == KIND_OUTER || (flags != nullptr && is_solid(flags[cell]))) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
-    } else {
-        unsigned long j[NQ];
-        compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
-        stream_pull_load(k, f, in + base, j);
-        if (cell_kind == KIND_ACTIVE) {
-            if (d_block_wall.nx != 0) {
-                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
-            collide_mrt_unforced(f);  // SRT path below unchanged.
-#else
-            float rho, ux, uy, feq[NQ];
-            compute_rho_u(f, rho, ux, uy);
-#if USE_DDF_SHIFTING
-            // Preserve the density perturbation before adding the unit reference.
-            const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
-            const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
-            const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
-            compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
-#else
-            compute_feq(rho, ux, uy, feq);
-#endif
-            collide_srt_no_force(f, feq, d_params.omega);
-#endif
-        }
-    }
-    stream_pull_store(k, f, out + base);
-}
-
-// Root channel closure; fine and periodic levels keep BlockStep.
-#include "step/block_channel.cu"
@@ -0,0 +1,20 @@
+// Physical parent walls; NX/NY remain local storage strides.
+#ifndef CELERIS_BOUNDARY_BLOCK_WALL_CUH
+#define CELERIS_BOUNDARY_BLOCK_WALL_CUH
+
+#include "planar_velocity.cuh"
+
+struct BlockWall { unsigned int nx, ny, slot_width; float amplitude; };
+extern "C" { __constant__ BlockWall d_block_wall; }
+
+__device__ __forceinline__ void block_wall(float* f, unsigned int x, unsigned int y,
+                                         unsigned int 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);
+}
+
+#endif
@@ -15,10 +15,13 @@
 // D2Q9 y-wall half-way bounce-back for pull double-buffer streaming.
 // ---------------------------------------------------------------------------
 #if LATTICE_MODEL == LATTICE_D2Q9
-__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
-                                                  float* __restrict__ f,
-                                                  const fpxx* __restrict__ fi_in,
-                                                  unsigned long k)
+// Rows 1 and top_row (= physical ny - 2) next to walls at rows 0 and ny - 1.
+// y is the physical row; k addresses the node in fi_in (a block window may
+// have NX/NY strides different from the physical extent).
+__device__ inline void apply_wall_bb_d2q9_rows(unsigned int y, unsigned int top_row,
+                                               float* __restrict__ f,
+                                               const fpxx* __restrict__ fi_in,
+                                               unsigned long k)
 {
    if (y == 1) {
        // Directions sourced from y=0 wall in pull step: +y, +x+y, -x+y.
@@ -26,14 +29,22 @@ __device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
        f[5] = load_ddf(fi_in, index_f(k, 6u));
        f[8] = load_ddf(fi_in, index_f(k, 7u));
    }
-    else if (y == (unsigned int)(NY - 2)) {
-        // Directions sourced from y=NY-1 wall in pull step: -y, -x-y, +x-y.
+    else if (y == top_row) {
+        // Directions sourced from the y = ny - 1 wall in pull step: -y, -x-y, +x-y.
        f[4] = load_ddf(fi_in, index_f(k, 3u));
        f[6] = load_ddf(fi_in, index_f(k, 5u));
        f[7] = load_ddf(fi_in, index_f(k, 8u));
    }
 }

+__device__ inline void apply_wall_bb_d2q9_y_pull(unsigned int y,
+                                                  float* __restrict__ f,
+                                                  const fpxx* __restrict__ fi_in,
+                                                  unsigned long k)
+{
+    apply_wall_bb_d2q9_rows(y, (unsigned int)(NY - 2), f, fi_in, k);
+}
+
 // ---------------------------------------------------------------------------
 // D2Q9 y-wall free-slip (specular reflection) for pull double-buffer streaming.
 //
@@ -8,6 +8,9 @@
 #define STORE_PRECISION  0
 #define USE_DDF_SHIFTING 0

+// Internal blocked root mode: 0=periodic/fine, 1=slot/reservoir, 2=channel.
+#define BLOCK_ROOT_BC 0
+
 #define USE_LES 0
 #define LES_CS  0.160000f
 #define LES_CLOSED_FORM 1
@@ -0,0 +1,16 @@
+// CelerisLab – core/cell_kind.cuh
+// Fixed-patch cell kinds stored per block cell, matching lbm/refinement.py
+// ACTIVE/INNER_GHOST/OUTER_GHOST: OUTER is invalid (stores NaN, no pull),
+// INNER streams without collision, ACTIVE pulls and collides.
+#ifndef CELERIS_CORE_CELL_KIND_CUH
+#define CELERIS_CORE_CELL_KIND_CUH
+
+constexpr unsigned char KIND_ACTIVE = 0u, KIND_INNER = 1u, KIND_OUTER = 2u;
+
+// NaN populations mark a cell without valid state (OUTER ghost or solid core).
+__device__ __forceinline__ void fill_invalid(float* f)
+{
+    for (int q = 0; q < NQ; ++q) f[q] = NAN;
+}
+
+#endif
@@ -0,0 +1,31 @@
+// Collision used by periodic/fine and slot/reservoir BlockStep.
+#ifndef CELERIS_OPERATORS_BLOCK_COLLISION_CUH
+#define CELERIS_OPERATORS_BLOCK_COLLISION_CUH
+
+#include "../core/params.cuh"
+#include "macro.cuh"
+#include "collision_srt.cuh"
+#if COLLISION_MODEL == 2
+#include "collision_mrt.cuh"
+#endif
+
+__device__ __forceinline__ void block_collide(float* f) {
+#if COLLISION_MODEL == 2
+    collide_mrt_unforced(f);  // SRT path below unchanged.
+#else
+    float rho, ux, uy, feq[NQ];
+    compute_rho_u(f, rho, ux, uy);
+#if USE_DDF_SHIFTING
+    // Preserve the density perturbation before adding the unit reference.
+    const float drho = f[0]+f[1]+f[2]+f[3]+f[4]+f[5]+f[6]+f[7]+f[8];
+    const float jx = f[1]-f[2]+f[5]-f[6]+f[7]-f[8];
+    const float jy = f[3]-f[4]+f[5]-f[6]-f[7]+f[8];
+    compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);
+#else
+    compute_feq(rho, ux, uy, feq);
+#endif
+    collide_srt_no_force(f, feq, d_params.omega);
+#endif
+}
+
+#endif
@@ -0,0 +1,14 @@
+// CelerisLab – operators/bounce_back_swap.cuh
+// Full-way bounce-back at a wall node: swap every opposite pair (paired
+// order q, q+1). Needs only NQ; shared by the uniform and blocked steps.
+#ifndef CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+#define CELERIS_OPERATORS_BOUNCE_BACK_SWAP_CUH
+
+__device__ __forceinline__ void bounce_back_swap(float* f) {
+    #pragma unroll
+    for (int i = 1; i < NQ; i += 2) {
+        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
+    }
+}
+
+#endif
@@ -5,12 +5,7 @@
 #ifndef CELERIS_OPERATORS_HELPERS_CUH
 #define CELERIS_OPERATORS_HELPERS_CUH

-__device__ __forceinline__ void bounce_back_swap(float* f) {
-    #pragma unroll
-    for (int i = 1; i < NQ; i += 2) {
-        float t = f[i]; f[i] = f[i+1]; f[i+1] = t;
-    }
-}
+#include "bounce_back_swap.cuh"

 __device__ __forceinline__ void collide_dispatch(
    float* f, float rho_n,
@@ -9,6 +9,7 @@
 #include "core/params.cuh"
 #include "core/flags.cuh"
 #include "core/obs.cuh"
+#include "core/cell_kind.cuh"
 #include "operators/macro.cuh"
 #include "operators/collision_srt.cuh"
 #if COLLISION_MODEL == 2
@@ -35,93 +36,5 @@ __device__ unsigned int incoming_counts[NQ * TOTAL_CELLS];
 }
 #endif

-__device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
-#if COLLISION_MODEL == 2
-    collide_mrt_unforced(f);
-#else
-    float rho, ux, uy, feq[NQ];
-    compute_rho_u(f, rho, ux, uy);
-    compute_feq(rho, ux, uy, feq);
-    collide_srt_no_force(f, feq, d_params.omega);
-#endif
-#ifdef CELERIS_REFINEMENT_AUDIT
-    ++collision_counts[k];
-#endif
-}
-
-extern "C" __global__ void PatchExplode(
-    const float* coarse, float* fine, const uint32_t* targets,
-    const uint32_t* donors, unsigned int count, unsigned long coarse_cells)
-{
-    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
-    if (row >= count) return;
-    for (unsigned int q = 0; q < NQ; ++q)
-        fine[index_f(targets[row], q)] = coarse[q * coarse_cells + donors[row]];
-#ifdef CELERIS_REFINEMENT_AUDIT
-    ++explosion_counts[targets[row]];
-#endif
-}
-
-extern "C" __global__ void PatchFineStep(const float* in, float* out, const uint16_t* flags) {
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x >= NX || y >= NY) return;
-    float f[NQ];
-    // Outer halo is invalid after EACH hop, including when reusing the buffer
-    // that held explosion data. Inner halo only streams; it never collides.
-    if (x < 2 || y < 2 || x >= NX - 2 || y >= NY - 2
-            || (flags != nullptr && is_solid(flags[k]))) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
-    } else {
-        unsigned long j[NQ];
-        compute_neighbors(x, y, j);  // These interior targets never wrap.
-        stream_pull_load(k, f, in, j);
-        if (x >= 4 && y >= 4 && x < NX - 4 && y < NY - 4)
-            patch_collide(f, k);
-    }
-    stream_pull_store(k, f, out);
-}
-
-extern "C" __global__ void PatchCoarseStep(
-    const float* in, float* out, const float* fine,
-    const uint64_t* targets, const uint64_t* children, unsigned int count,
-    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
-    unsigned int wall_aware)
-{
-    unsigned int x, y; unsigned long k;
-    index_from_thread(x, y, k);
-    if (x >= NX || y >= NY) return;
-    // Both buffers were poisoned at initialization; covered cells stay invalid.
-    if (x >= x0 && x < x1 && y >= y0 && y < y1) return;
-    if (wall_aware && (y == 0 || y == NY - 1)) return;
-    unsigned long j[NQ];
-    compute_neighbors(x, y, j);
-    float f[NQ];
-    stream_pull_load(k, f, in, j);
-    if (x >= x0 - 1 && x <= x1 && y >= y0 - 1 && y <= y1) {
-        for (unsigned int q = 1; q < NQ; ++q) {
-            uint64_t slot = index_f(k, q);
-            unsigned int lo = 0, hi = count;
-            while (lo < hi) {
-                unsigned int mid = lo + (hi - lo) / 2;
-                if (targets[mid] < slot) lo = mid + 1;
-                else hi = mid;
-            }
-            if (lo < count && targets[lo] == slot) {
-                const uint64_t* s = children + 4ul * lo;
-                f[q] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
-#ifdef CELERIS_REFINEMENT_AUDIT
-                ++incoming_counts[slot];
-#endif
-            }
-        }
-    }
-    // Includes the eight grazing-corner links with UNCOVERED coarse donors.
-    // All replacement arrivals must be in place before computing rho/u.
-    if (wall_aware) {
-        if (y == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
-        else if (y == NY - 2) reconstruct_top_reservoir(f, RHO);
-    }
-    patch_collide(f, k);
-    stream_pull_store(k, f, out);
-}
+// Kernels: patch collision helper, PatchExplode, PatchFineStep, PatchCoarseStep.
+#include "step/patch_step.cu"
@@ -3,7 +3,7 @@
 // ApplyEsoPullCurvedKernel, SensorKernel, ForceRegionKernel.
 //
 // Launched on compact lists (n_curved / n_sensor / n_cells threads).
-// Included by kernel_v2.cu.  No standalone compilation.
+// Included by uniform.cu.  No standalone compilation.
 //
 // Double-buffer mode:
 //   CurvedBoundaryKernel runs BEFORE OneStep.
@@ -0,0 +1,42 @@
+// CelerisLab/lbm/kernels/step/block_boundary.cu
+// Receiver-local curved-boundary staging and its raw-exchange accumulation.
+// Part of the blocked.cu compilation entry; relies on its includes and guards.
+
+// Static cut links, receiver-local staging: cell is the receiver's local core
+// cell, so the existing BFL writes the solid-source window image that this
+// block's next pull reads (no halo copy between this kernel and BlockStep).
+// surface_velocity (nullable, 2 floats per link): prescribed fluid velocity on
+// fixed geometry, same convention as CurvedBoundaryKernel. Optional raw rows
+// (jx,jy,torque,dm) are unique physical links, never body-observation slots.
+extern "C" __global__ void BlockCurvedBoundary(
+    fpxx* f, const uint32_t* block, const uint32_t* cell, const unsigned char* dir,
+    const float* q, const unsigned char* fallback, unsigned int count,
+    const float* surface_velocity, const float* rx, const float* ry, float* raw)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    float Uw = 0.0f, Vw = 0.0f;
+    if (surface_velocity != nullptr) {
+        const float u = surface_velocity[2 * row], v = surface_velocity[2 * row + 1];
+        if (u != 0.0f || v != 0.0f) { Uw = u; Vw = v; }
+    }
+    apply_bouzidi_link(dir[row], q[row], cell[row], f + (unsigned long)block[row] * NQ * TOTAL_CELLS,
+                       Uw, Vw, rx ? rx[row] : 0.0f, ry ? ry[row] : 0.0f,
+                       fallback[row], nullptr, -1, raw ? raw + 4ul * row : nullptr,
+                       raw ? raw + 4ul * row + 3 : nullptr);
+}
+
+// Same-stream BFL -> this row-owned accumulation -> pull step. No atomics,
+// no cross-row reduction and no history allocation proportional to substeps.
+extern "C" __global__ void BlockAccumulateExchange(
+    const float* raw, double* integral, unsigned int count, double h, double dt)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const unsigned long k = 4ul * row;
+    const double impulse = h * dt;
+    integral[k] += impulse * (double)raw[k];
+    integral[k + 1] += impulse * (double)raw[k + 1];
+    integral[k + 2] += h * impulse * (double)raw[k + 2];
+    integral[k + 3] += h * h * (double)raw[k + 3];
+}
@@ -1,5 +1,12 @@
+#if BLOCK_ROOT_BC != 2
+#error "BlockChannelStep requires BLOCK_ROOT_BC == 2"
+#endif
+
 // Root-only blocked counterpart of OneStep's D2Q9/SRT channel closure.
 // NX/NY are LOCAL strides; all physical boundary decisions use global coordinates.
+#include "../core/cell_kind.cuh"
+#include "../operators/bounce_back_swap.cuh"
+#include "../boundary/bounce_back.cuh"
 #include "../boundary/inlet/common.cuh"
 #include "../boundary/inlet/zou_he_local.cuh"
 #include "../boundary/outlet/pressure_neq.cuh"
@@ -21,7 +28,7 @@ extern "C" __global__ void BlockChannelStep(
    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
    float f[NQ];
    if (kind != nullptr && kind[cell] == KIND_OUTER) {
-        for (int q = 0; q < NQ; ++q) f[q] = NAN;
+        fill_invalid(f);
    } else {
        unsigned long j[NQ];
        compute_neighbors(x, y, j);
@@ -29,27 +36,17 @@ extern "C" __global__ void BlockChannelStep(
        const bool wall = gy == 0 || gy == d_block_channel.ny - 1;
        const bool outlet = gx == d_block_channel.nx - 1;
        if (wall) {
-            // Wall owns all four corners. Ghosts pull and swap, never collide.
-            for (int q = 1; q < NQ; q += 2) {
-                const float t = f[q]; f[q] = f[q + 1]; f[q + 1] = t;
-            }
+            bounce_back_swap(f);  // Wall owns all four corners; ghosts never collide.
        } else if (gx == 0) {
            apply_zou_he_left_velocity_inlet_d2q9(
                f, inlet_target_u((float)gy, d_block_channel.ny, d_block_channel.velocity), 0.0f);
        } else if (outlet) {
            // Immutable input donor, receiver-local (also valid across a block seam).
            float donor[NQ];
-            for (int q = 0; q < NQ; ++q)
-                donor[q] = load_ddf(in + base, index_f(k - 1, q));
+            load_node(in + base, k - 1, donor);
            apply_pressure_outlet_d2q9(f, donor);
-        } else if (gy == 1) {
-            f[3] = load_ddf(in + base, index_f(k, 4));
-            f[5] = load_ddf(in + base, index_f(k, 6));
-            f[8] = load_ddf(in + base, index_f(k, 7));
-        } else if (gy == d_block_channel.ny - 2) {
-            f[4] = load_ddf(in + base, index_f(k, 3));
-            f[6] = load_ddf(in + base, index_f(k, 5));
-            f[7] = load_ddf(in + base, index_f(k, 8));
+        } else {
+            apply_wall_bb_d2q9_rows(gy, d_block_channel.ny - 2, f, in + base, k);
        }
        if (!wall && !outlet) {
 #if COLLISION_MODEL == 2
@@ -0,0 +1,35 @@
+// CelerisLab/lbm/kernels/step/block_step.cu
+// Generic blocked pull/collide/store for periodic/fine and slot/reservoir levels.
+// Part of the blocked.cu compilation entry; relies on its includes and guards.
+#if BLOCK_ROOT_BC == 2
+#error "BlockStep is not built for the channel root (BLOCK_ROOT_BC == 2)"
+#endif
+
+// Pull/collide/store core cells only; halo cells of out are left untouched.
+// flags (block cell ids, nullable): solid cores store NaN and never collide.
+// kind (block cell ids, nullable = all ACTIVE): OUTER skips the pull and
+// stores NaN, INNER pulls without collision, ACTIVE pulls and collides.
+extern "C" __global__ void BlockStep(const fpxx* in, fpxx* out, unsigned int halo,
+                                     const uint16_t* flags, const unsigned char* kind) {
+    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 long base = (unsigned long)blockIdx.z * NQ * TOTAL_CELLS;
+    const unsigned long cell = (unsigned long)blockIdx.z * TOTAL_CELLS + k;
+    const unsigned char cell_kind = kind != nullptr ? kind[cell] : KIND_ACTIVE;
+    float f[NQ];
+    if (cell_kind == KIND_OUTER || (flags != nullptr && is_solid(flags[cell]))) {
+        fill_invalid(f);
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);  // Core neighbors stay inside the window: no wrap.
+        stream_pull_load(k, f, in + base, j);
+        if (cell_kind == KIND_ACTIVE) {
+#if BLOCK_ROOT_BC == 1
+            block_wall(f, x, y, halo);
+#endif
+            block_collide(f);
+        }
+    }
+    stream_pull_store(k, f, out + base);
+}
@@ -0,0 +1,78 @@
+// CelerisLab/lbm/kernels/step/block_transfer.cu
+// Same-level halo refresh and 2:1 explosion/coalescence between blocked levels.
+// Part of the blocked.cu compilation entry; relies on its includes and guards.
+
+// Halo refresh: block cell ids b*TOTAL_CELLS+c, sources are always core cells.
+extern "C" __global__ void BlockHaloCopy(
+    fpxx* f, const uint64_t* dst, const uint64_t* src, unsigned int count)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t d = dst[row], s = src[row];
+    fpxx* to = f + (d / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    const fpxx* from = f + (s / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    for (unsigned int q = 0; q < NQ; ++q)
+        to[index_f(d % TOTAL_CELLS, q)] = from[index_f(s % TOTAL_CELLS, q)];
+}
+
+// Fixed-patch explosion into blocked fine storage: target = owner block cell
+// id of a patch-ghost cell; donors are parent block cell ids. Dense parent
+// storage is the one-window case with coarse_cells = full domain size.
+extern "C" __global__ void BlockExplode(
+    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
+    unsigned int count, unsigned long coarse_cells)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t t = target[row];
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    const uint64_t d = donor[row];
+    const fpxx* from = coarse + (d / coarse_cells) * NQ * coarse_cells;
+    for (unsigned int q = 0; q < NQ; ++q)
+        to[index_f(t % TOTAL_CELLS, q)] = from[q * coarse_cells + d % coarse_cells];
+}
+
+// Experimental Schornbaum--Ruede (arXiv:1508.07982, Algorithm 2)
+// and Chen et al. (doi:10.1016/j.physa.2005.09.036) population interpolation.
+// Independent implementation of the mathematical stencil; no upstream source copied.
+// Stencil cells are owner cores; a disabled axis repeats its center twice.
+extern "C" __global__ void BlockExplodeProjected(
+    const fpxx* coarse, fpxx* fine, const uint64_t* target, const uint64_t* donor,
+    unsigned int count, unsigned long coarse_cells,
+    const uint64_t* neighbors, const float* offsets)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t t = target[row], d = donor[row];
+    fpxx* to = fine + (t / TOTAL_CELLS) * NQ * TOTAL_CELLS;
+    for (unsigned int q = 0; q < NQ; ++q) {
+        float v[4];
+        for (unsigned int a = 0; a < 4; ++a) {
+            const uint64_t n = neighbors[4ul * row + a];
+            v[a] = coarse[(n / coarse_cells) * NQ * coarse_cells
+                          + q * coarse_cells + n % coarse_cells];
+        }
+        float gx = .5f * (v[1] - v[0]), gy = .5f * (v[3] - v[2]);
+        const float ex = d_cx[q], ey = d_cy[q];
+        if (q != 0) {
+            const float along = (ex * gx + ey * gy) / (ex * ex + ey * ey);
+            gx -= ex * along;
+            gy -= ey * along;
+        }
+        to[index_f(t % TOTAL_CELLS, q)] =
+            coarse[(d / coarse_cells) * NQ * coarse_cells + q * coarse_cells + d % coarse_cells]
+            + offsets[2ul * row] * gx + offsets[2ul * row + 1] * gy;
+    }
+}
+
+// Targets are receiver-local SOURCE slots for the immediately following pull.
+// Keep child summation order identical to PatchCoarseStep; no collision here.
+extern "C" __global__ void BlockCoalesce(
+    const fpxx* fine, fpxx* coarse, const uint64_t* target, const uint64_t* children,
+    unsigned int count)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    const uint64_t* s = children + 4ul * row;
+    coarse[target[row]] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
+}
@@ -3,7 +3,7 @@
 //
 // Decodes the EsoPull backing layout into physical distributions, then
 // computes rho/ux/uy.  Output is SoA: rho_out[N], ux_out[N], uy_out[N].
-// Included by kernel_v2.cu.  No standalone compilation.
+// Included by uniform.cu.  No standalone compilation.
 //
 // This kernel is the sole path for macroscopic diagnostics in EsoPull mode.
 // It uses the shared semantic helper layer (esopull_semantic_helpers.cuh).
@@ -5,7 +5,7 @@
 // or encodes physical DDF back to EsoPull backing layout.
 // Each thread handles one cell from a compact index list.
 //
-// Included by kernel_v2.cu.  No standalone compilation.
+// Included by uniform.cu.  No standalone compilation.
 // ============================================================================

 #ifndef CELERIS_STEP_ESOPULL_SYNC_CU
@@ -1,6 +1,6 @@
 // CelerisLab – step/init_flow.cu
 // Channel-flow initialization kernels (double-buffer and EsoPull).
-// Included by kernel_v2.cu after all headers.  No standalone compilation.
+// Included by uniform.cu after all headers.  No standalone compilation.
 // ============================================================================

 #ifndef CELERIS_STEP_INIT_FLOW_CU
@@ -73,9 +73,9 @@ __device__ __forceinline__ uint16_t finalize_domain_flag(
 }

 // ============================================================================
-// InitTubeFlow_v2 — double-buffer channel initialization
+// InitTubeFlow — double-buffer channel initialization
 // ============================================================================
-__global__ void InitTubeFlow_v2(uint16_t* flag, fpxx* fi)
+__global__ void InitTubeFlow(uint16_t* flag, fpxx* fi)
 {
 #if DIM == 2
    unsigned int x, y; unsigned long k;
@@ -1,6 +1,6 @@
 // CelerisLab – step/one_step_double.cu
 // Double-buffer PULL streaming step kernel.
-// Included by kernel_v2.cu.  No standalone compilation.
+// Included by uniform.cu.  No standalone compilation.
 // ============================================================================

 #ifndef CELERIS_STEP_ONE_STEP_DOUBLE_CU
@@ -1,6 +1,6 @@
 // CelerisLab – step/one_step_esopull.cu
 // Esoteric-Pull single-buffer step kernel.
-// Included by kernel_v2.cu.  No standalone compilation.
+// Included by uniform.cu.  No standalone compilation.
 //
 // CURVED BC: supported via PrepareEsoPullCurvedKernel + ApplyEsoPullCurvedKernel
 // in aux_kernels.cu, launched by stepper.py before EsoPullStep.
@@ -0,0 +1,95 @@
+// CelerisLab/lbm/kernels/step/patch_step.cu
+// Dense fixed-patch collision, explosion and the fine/coarse steps.
+// Part of the refinement.cu compilation entry; relies on its includes, guards
+// and the optional CELERIS_REFINEMENT_AUDIT counters declared there.
+
+__device__ __forceinline__ void patch_collide(float* f, unsigned long k) {
+#if COLLISION_MODEL == 2
+    collide_mrt_unforced(f);
+#else
+    float rho, ux, uy, feq[NQ];
+    compute_rho_u(f, rho, ux, uy);
+    compute_feq(rho, ux, uy, feq);
+    collide_srt_no_force(f, feq, d_params.omega);
+#endif
+#ifdef CELERIS_REFINEMENT_AUDIT
+    ++collision_counts[k];
+#endif
+}
+
+extern "C" __global__ void PatchExplode(
+    const float* coarse, float* fine, const uint32_t* targets,
+    const uint32_t* donors, unsigned int count, unsigned long coarse_cells)
+{
+    unsigned int row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= count) return;
+    for (unsigned int q = 0; q < NQ; ++q)
+        fine[index_f(targets[row], q)] = coarse[q * coarse_cells + donors[row]];
+#ifdef CELERIS_REFINEMENT_AUDIT
+    ++explosion_counts[targets[row]];
+#endif
+}
+
+extern "C" __global__ void PatchFineStep(const float* in, float* out, const uint16_t* flags) {
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x >= NX || y >= NY) return;
+    float f[NQ];
+    // Outer halo is invalid after EACH hop, including when reusing the buffer
+    // that held explosion data. Inner halo only streams; it never collides.
+    if (x < 2 || y < 2 || x >= NX - 2 || y >= NY - 2
+            || (flags != nullptr && is_solid(flags[k]))) {
+        fill_invalid(f);
+    } else {
+        unsigned long j[NQ];
+        compute_neighbors(x, y, j);  // These interior targets never wrap.
+        stream_pull_load(k, f, in, j);
+        if (x >= 4 && y >= 4 && x < NX - 4 && y < NY - 4)
+            patch_collide(f, k);
+    }
+    stream_pull_store(k, f, out);
+}
+
+extern "C" __global__ void PatchCoarseStep(
+    const float* in, float* out, const float* fine,
+    const uint64_t* targets, const uint64_t* children, unsigned int count,
+    unsigned int x0, unsigned int y0, unsigned int x1, unsigned int y1,
+    unsigned int wall_aware)
+{
+    unsigned int x, y; unsigned long k;
+    index_from_thread(x, y, k);
+    if (x >= NX || y >= NY) return;
+    // Both buffers were poisoned at initialization; covered cells stay invalid.
+    if (x >= x0 && x < x1 && y >= y0 && y < y1) return;
+    if (wall_aware && (y == 0 || y == NY - 1)) return;
+    unsigned long j[NQ];
+    compute_neighbors(x, y, j);
+    float f[NQ];
+    stream_pull_load(k, f, in, j);
+    if (x >= x0 - 1 && x <= x1 && y >= y0 - 1 && y <= y1) {
+        for (unsigned int q = 1; q < NQ; ++q) {
+            uint64_t slot = index_f(k, q);
+            unsigned int lo = 0, hi = count;
+            while (lo < hi) {
+                unsigned int mid = lo + (hi - lo) / 2;
+                if (targets[mid] < slot) lo = mid + 1;
+                else hi = mid;
+            }
+            if (lo < count && targets[lo] == slot) {
+                const uint64_t* s = children + 4ul * lo;
+                f[q] = 0.25f * (fine[s[0]] + fine[s[1]] + fine[s[2]] + fine[s[3]]);
+#ifdef CELERIS_REFINEMENT_AUDIT
+                ++incoming_counts[slot];
+#endif
+            }
+        }
+    }
+    // Includes the eight grazing-corner links with UNCOVERED coarse donors.
+    // All replacement arrivals must be in place before computing rho/u.
+    if (wall_aware) {
+        if (y == 1) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
+        else if (y == NY - 2) reconstruct_top_reservoir(f, RHO);
+    }
+    patch_collide(f, k);
+    stream_pull_store(k, f, out);
+}
@@ -21,6 +21,14 @@
 //   i.e. the population traveling in direction i at node n
 //   arrived from the node at n - c_i  = neighbor in direction opp(i)
 // ---------------------------------------------------------------------------
+// load_node: all NQ populations of node n as stored (no streaming).
+__device__ __forceinline__ void load_node(const fpxx* __restrict__ fi, unsigned long n,
+                                          float* __restrict__ f)
+{
+    for (int q = 0; q < NQ; ++q)
+        f[q] = load_ddf(fi, index_f(n, q));
+}
+
 __device__ __forceinline__ void stream_pull_load(unsigned long n,
                                                 float* __restrict__ f,
                                                 const fpxx* __restrict__ fi,
@@ -1,4 +1,4 @@
-// CelerisLab – kernel_v2.cu
+// CelerisLab – uniform.cu
 // ============================================================================
 // Compilation entry point (Stage 2 Architecture).
 //
@@ -7,7 +7,7 @@
 //
 // File layout (each <300 lines):
 //   operators/helpers.cuh  — bounce_back_swap, collide_dispatch
-//   step/init_flow.cu      — InitTubeFlow_v2, InitEsoPull
+//   step/init_flow.cu      — InitTubeFlow, InitEsoPull
 //   step/one_step_double.cu — OneStep (pull double-buffer)
 //   step/one_step_esopull.cu — EsoPullStep (single-buffer)
 //   step/aux_kernels.cu    — CurvedBoundaryKernel, SensorKernel
@@ -14,15 +14,16 @@ from threading import RLock
 import numpy as np
 import pycuda.driver as cuda

-from ..cuda import compiler_v2 as compiler
-from .blocks import BlockLayout2D
+from ..cuda import compiler
+from .blocks import (BlockLayout2D, launch_curved_boundary, launch_halo_copy, level_config,
+                     static_circle_tables)
 from .channel import channel_masks, validate_channel_config, validate_channel_transfers
 from ..body.jets import jet_link_velocity
-from .block_stepper import MAX_GRID_DIM
 from .hierarchy import build_hierarchy, is_rectangle
 from .refinement import ACTIVE, HALO, FixedPatch2D, FixedPatchSet2D, prepare_fine_boundary
 from .refinement_tables import (blocked_base_tables, blocked_fine_tables,  # noqa: F401 (re-export)
                                blocked_kind, projected_explosion_tables, _prepare_blocked_boundary,
+                                _BlockedBoundary,
                                BLOCKED_BOUNDARY_DEVICE_COLUMNS)
 from .descriptors import SOLID
 from .raw_exchange import summarize_exchange
@@ -65,17 +66,21 @@ def _level_name(level):
    return ("base", "fine", "child")[level] if level < 3 else f"level {level}"


+MAX_GRID_DIM = 65535  # CUDA gridDim.y/z limit (rows, blocks)
+
 def _check_levels(refinements, blocks, objects, jets):
    """Candidate limits of the level lists, checked before any device work."""
    depth = len(refinements) + 1
-    if depth < 2:
-        raise ValueError("at least one refinement level is required")
    for name, values in (("level_blocks", blocks), ("level_objects", objects), ("level_jets", jets)):
        if len(values) != depth:
            raise ValueError(f"{name} needs one entry per level ({depth})")
    for level, block in enumerate(blocks):
        if block is not None and len(block) != 2:
            raise ValueError(f"level_blocks[{level}] must be (by, bx)")
+    if depth == 1:  # one blocked level: the root carries the static circles
+        if blocks[0] is None or jets[0]:
+            raise ValueError("a single level must be blocked and has no Jets")
+        return
    if objects[0] or jets[0]:
        raise ValueError("base-level bodies/Jets are not supported")
    if any(blocks[level - 1] is not None and blocks[level] is None for level in range(1, depth)):
@@ -98,6 +103,9 @@ class FixedPatchStepper:
    level-wide grid, public fine arrays (2*ny, 2*nx), NaN outside ACTIVE cells);
    see lbm.hierarchy. ``level_blocks``, ``level_objects`` and ``level_jets``
    hold one entry per level, L0 included (L0 bodies/Jets are not supported).
+    ``refinements=[]`` is one blocked level: periodic SRT root without LES,
+    Jets or exchange observation; its ``level_objects[0]`` are stationary
+    passive circles (static BFL).

    ``coarse_boundary`` is ``periodic`` (default) or ``slot_reservoir``:
    a bottom on-site slot and top fixed-density reservoir from ``cfg``.
@@ -124,7 +132,7 @@ class FixedPatchStepper:
    stream. On a device/launch failure the instance becomes unusable: close and
    recreate it, rather than pretending a partially advanced buffer rolled back.

-    A ``level_blocks`` entry (by, bx) stores that level in blocks.cu storage:
+    A ``level_blocks`` entry (by, bx) stores that level in blocked.cu storage:
    the ENTIRE level shape, including the four-cell patch ghost band, is
    partitioned by a non-periodic BlockLayout2D with one-cell same-level halos.
    Patch kinds stay per cell (OUTER invalid, INNER stream only, ACTIVE
@@ -196,6 +204,11 @@ class FixedPatchStepper:
        self._ddf_shifting = bool(cfg.ddf_shifting)
        if self._ddf_shifting and any(b is None for b in level_blocks):
            raise ValueError("shifted populations require all-blocked levels; dense/mixed is unsupported")
+        if depth == 1 and (coarse_boundary != "periodic" or cfg.collision != "SRT" or cfg.les_enabled
+                           or explosion != "constant"
+                           or observe_raw_exchange or observe_window_exchange):
+            raise ValueError("a single level requires a periodic SRT root without LES, projection "
+                             "or exchange observation")
        self._channel = coarse_boundary == "channel"
        if self._channel:
            validate_channel_config(cfg, level_blocks)
@@ -229,12 +242,14 @@ class FixedPatchStepper:
            base = blocked_base_tables(pairs[0], self._layouts[0], np.argsort(pairs[0].coalesce_coarse))
            validate_channel_transfers(pairs[0], self._layouts[0], base[1])
        # Kernel configurations per level (geometry configs live in self.hierarchy).
-        fine_cfg = replace(cfg, nx=pairs[0].fine_shape[1], ny=pairs[0].fine_shape[0],
-                           viscosity=2 * cfg.viscosity)
-        if self._wall_aware:
-            fine_cfg = replace(fine_cfg, y_wall_bc="bounce_back", top_wall_bc="velocity",
-                               wall_slot_width=0, wall_normal_velocity=0.0, velocity=0.0)
-        self._level_cfgs = [replace(cfg), fine_cfg]
+        self._level_cfgs = [replace(cfg)]
+        if pairs:
+            fine_cfg = replace(cfg, nx=pairs[0].fine_shape[1], ny=pairs[0].fine_shape[0],
+                               viscosity=2 * cfg.viscosity)
+            if self._wall_aware:
+                fine_cfg = replace(fine_cfg, y_wall_bc="bounce_back", top_wall_bc="velocity",
+                                   wall_slot_width=0, wall_normal_velocity=0.0, velocity=0.0)
+            self._level_cfgs.append(fine_cfg)
        for pair in pairs[1:]:
            parent = self._level_cfgs[-1]
            self._level_cfgs.append(replace(parent, nx=pair.fine_shape[1], ny=pair.fine_shape[0],
@@ -275,12 +290,18 @@ class FixedPatchStepper:
        self._exchange_centers = tuple(tuple(tuple(o.geometry.center) for o in objects)
                                       for objects in level_objects)
        self._level_jets = level_jets
-        self._level_boundaries = [None] + [
+        root = None
+        if depth == 1 and level_objects[0]:
+            tables = static_circle_tables(self._layouts[0], level_objects[0])
+            root = _BlockedBoundary(None, tables, (), False)
+        self._level_boundaries = [root] + [
            _prepare_blocked_boundary(self._boundary_hosts[level], self._layouts[level],
                                      self._level_cfgs[level], level_jets[level],
                                      _shifted_blocked=self._ddf_shifting)
            for level in range(1, depth)]
        self._leaf = [kind == ACTIVE for kind in self.hierarchy.kinds]
+        if root is not None:
+            self._leaf[0] &= ~root.tables["solid"]
        for level in range(1, depth):
            host = self._boundary_hosts[level]
            if host is not None:
@@ -303,6 +324,9 @@ class FixedPatchStepper:
        self._block = (cfg.threads_per_block, 1, 1)
        with self.ctx.activate():
            try:
+                free, _ = cuda.mem_get_info()
+                if self._planned_bytes() > free:  # before any stream, module or buffer
+                    raise MemoryError("level buffers exceed free device memory")
                self._stream = cuda.Stream()
                for level in range(depth):
                    self._levels.append(self._build_level(level))
@@ -324,7 +348,9 @@ class FixedPatchStepper:
                                b.window_integral = self._allocate(count * 4 * 8)
                            b.accumulate = level.module.get_function("BlockAccumulateExchange")
                        b.kernel = level.module.get_function("BlockCurvedBoundary")
-                if self._level_boundaries[1] is not None:
+                if depth == 1:
+                    pass
+                elif self._level_boundaries[1] is not None:
                    b = self._level_boundaries[1]
                    self._fine_flags, self._link_tables, self._curved = b.flags, b.device, b.kernel
                elif self._boundary_hosts[1] is not None:
@@ -346,32 +372,47 @@ class FixedPatchStepper:
                    pass  # Preserve the original construction failure.
                raise

+    def _planned_bytes(self):
+        """Both G buffers, halo pairs and boundary columns of every level (transfer tables excluded).
+
+        A fail-fast lower bound only: passing it cannot guarantee that later
+        allocations succeed (other consumers, fragmentation, races).
+        """
+        total = 0
+        for index, layout in enumerate(self._layouts):
+            cells = int(np.prod(layout.blocked_shape if layout is not None
+                                else (9, *self.hierarchy.kinds[index].shape)))
+            total += 2 * 4 * cells
+            if layout is not None:
+                total += layout.halo_dst.nbytes + layout.halo_src.nbytes
+            b = self._level_boundaries[index]
+            if b is not None:
+                total += sum(np.asarray(v).nbytes for k, v in b.tables.items()
+                             if k == "flags" or k in BLOCKED_BOUNDARY_DEVICE_COLUMNS)
+        return total
+
    def _build_level(self, index):
        """Compile one level's step module and allocate its two G buffers."""
        layout, config = self._layouts[index], self._level_cfgs[index]
        pairs = self.hierarchy.pairs
-        tau = pairs[0].tau_coarse if index == 0 else pairs[index - 1].tau_fine
+        tau = (pairs[0].tau_coarse if pairs else 3 * config.viscosity + .5) if index == 0 \
+            else pairs[index - 1].tau_fine
        name, source = ("PatchCoarseStep" if index == 0 else "PatchFineStep"), "refinement.cu"
        if layout is not None:
-            config = replace(config, nx=layout.local_shape[1], ny=layout.local_shape[0],
-                             y_wall_bc="bounce_back", top_wall_bc="velocity",
-                             wall_slot_width=0, wall_normal_velocity=0.0, velocity=0.0)
-            name, source = ("BlockChannelStep" if index == 0 and self._channel else "BlockStep"), "blocks.cu"
-        config.validate()
-        if not np.isclose(config.omega, 1 / tau, rtol=1e-14, atol=0):
-            raise RuntimeError("level omega/tau mismatch")
+            config = level_config(config, layout, tau)
+            name, source = ("BlockChannelStep" if index == 0 and self._channel else "BlockStep"), "blocked.cu"
+        else:
+            config.validate()
+            if not np.isclose(config.omega, 1 / tau, rtol=1e-14, atol=0):
+                raise RuntimeError("level omega/tau mismatch")
        host = self._boundary_hosts[index]
        n_objects = len(host.objects) if name == "PatchFineStep" and host is not None else 0
-        build = compiler.generate_config(config, n_objects=n_objects)
-        try:
-            ptx = compiler.compile_kernel(self.ctx.sm_arch, source=source, build_dir=build)
-            module = compiler.load_module(ptx)
-        finally:
-            compiler.cleanup_build_dir(build)
-        ptr, size = module.get_global("d_params")
-        if size != struct.calcsize("=ffffI"):
-            raise RuntimeError("LBMParams host/device layout mismatch")
-        cuda.memcpy_htod(ptr, struct.pack("=ffffI", config.omega, 0, 0, 0, n_objects))
+        block_root_bc = 0
+        if index == 0 and layout is not None:
+            block_root_bc = 2 if self._channel else 1 if self._wall_aware else 0
+        module = compiler.build_module(config, source, arch=self.ctx.sm_arch, n_objects=n_objects,
+                                       block_root_bc=block_root_bc)
+        compiler.set_params(module, config.omega, n_objects=n_objects)
        shape = (9, *self.hierarchy.kinds[index].shape)
        size = int(np.prod(shape if layout is None else layout.blocked_shape)) * 4
        grid = ((config.nx + self._block[0] - 1) // self._block[0], config.ny,
@@ -516,7 +557,9 @@ class FixedPatchStepper:
            try:
                for level, array in zip(self._levels, arrays):
                    cuda.memcpy_htod(level.current, array)
-                    cuda.memcpy_htod(level.spare, array)
+                    # A single level starts its spare NaN so stale halos cannot hide.
+                    cuda.memcpy_htod(level.spare, np.full_like(array, np.nan)
+                                     if len(self._levels) == 1 else array)
                if self._fine_curved is not None and self._fine_curved.link_exchange_gpu is not None:
                    cuda.memset_d8(self._fine_curved.link_exchange_gpu, 0, self._fine_curved.count * 12)
                self._step_count = 0
@@ -680,15 +723,12 @@ class FixedPatchStepper:
            return
        if b.kernel is None or b.flags is None:
            raise RuntimeError("Boundary device buffers are missing")
-        t, count = b.device, len(b.tables["cell"])
+        t = b.device
        velocity = t.get("surface_velocity")
        if b.actuated and velocity is None:
            raise RuntimeError("Surface velocity device buffer is missing")
-        b.kernel(buffer, t["block"], t["cell"], t["dir"], t["q"], t["fallback"],
-                 np.uint32(count), velocity if velocity is not None else np.uintp(0),
-                 t.get("rx", np.uintp(0)), t.get("ry", np.uintp(0)),
-                 b.last_raw if b.last_raw is not None else np.uintp(0),
-                 grid=((count + self._block[0] - 1) // self._block[0], 1), **kw)
+        launch_curved_boundary(b.kernel, buffer, t, len(b.tables["cell"]), velocity=velocity,
+                               rx=t.get("rx"), ry=t.get("ry"), raw=b.last_raw, **kw)

    def _advance(self, index, parent_buffer, **kw):
        """Explode from the given parent buffer, then two hops of level ``index``."""
@@ -740,13 +780,13 @@ class FixedPatchStepper:
                             np.float64(h), np.float64(h),
                             grid=((count + self._block[0] - 1) // self._block[0], 1), **kw)
        flags = np.uintp(0) if level.boundary is None else level.boundary.flags
-        level.kernel(src, dst, np.uint32(level.layout.halo), flags, level.kind, grid=level.grid, **kw)
+        kind = np.uintp(0) if level.kind is None else level.kind  # single level: all ACTIVE
+        level.kernel(src, dst, np.uint32(level.layout.halo), flags, kind, grid=level.grid, **kw)

    def _gather_halo(self, level, buffer, **kw):
        pairs = len(level.layout.halo_dst)
        if pairs:
-            level.halo_copy(buffer, level.halo_dst, level.halo_src, np.uint32(pairs),
-                            grid=((pairs + self._block[0] - 1) // self._block[0], 1), **kw)
+            launch_halo_copy(level.halo_copy, buffer, level.halo_dst, level.halo_src, pairs, **kw)

    def step(self, n=1):
        """Advance G→G by n base intervals: level l takes 2**l hops per interval.
@@ -946,8 +986,8 @@ class FixedPatchStepper:
            if exc[0] is None:
                raise

-    # Read-only compatibility views over the level lists. Internal code must not use them;
-    # REMOVE AFTER the existing tests/docs migrate to hierarchy/_layouts/_level_* lists.
+    # Read-only compatibility views over the level lists, kept on purpose for the existing
+    # tests and docs. Internal code must not use them; it reads hierarchy/_layouts/_level_*.
    patch = property(lambda self: self.hierarchy.pairs[0])
    child_patch = property(lambda self: self.hierarchy.pairs[1] if self.hierarchy.depth > 2 else None)
    base_layout = property(lambda self: self._layouts[0])
@@ -28,7 +28,7 @@ class LBMStepper:
            self.boundary_values_gpu = cuda.mem_alloc((6 * cfg.nx + 2 * cfg.ny * cfg.nq) * 4)
        else:
            self.step_fn = module.get_function("OneStep")
-            self.init_fn = module.get_function("InitTubeFlow_v2")
+            self.init_fn = module.get_function("InitTubeFlow")

        self.curved_fn = module.get_function("CurvedBoundaryKernel")
        self.prepare_esopull_curved_fn = module.get_function("PrepareEsoPullCurvedKernel")
@@ -25,7 +25,7 @@ import pycuda.driver as cuda

 from .config import LBMConfig, BodyConfig, load_lbm_config, load_body_config
 from .cuda.context import CudaContext, context_bound
-from .cuda import compiler_v2 as compiler
+from .cuda import compiler
 from .lbm.field import LBMField
 from .lbm.stepper import LBMStepper
 from .body.objects import SimObject
@@ -46,7 +46,6 @@ class Simulation:
                 device_id: int = 0):
        # Internal stream for async DTOH coordination
        self._stream: Optional[cuda.Stream] = None
-        self._build_dir = None
        self._closed = False

        # Load configs
@@ -59,9 +58,7 @@ class Simulation:
        try:
            with self.ctx.activate():
                arch = self._resolve_compile_arch()
-                self._build_dir = compiler.generate_config(self.lbm_cfg, n_objects=0)
-                self._ptx_path = compiler.compile_kernel(arch=arch, build_dir=self._build_dir)
-                self._module = compiler.load_module(self._ptx_path)
+                self._module = compiler.build_module(self.lbm_cfg, "uniform.cu", arch=arch)
                self._compiled_n_objects = 0
                self._allocation_signature = self.lbm_cfg.allocation_signature()
                self._compiled_viscosity = self.lbm_cfg.viscosity
@@ -412,15 +409,8 @@ class Simulation:
        self.lbm_cfg.validate_runtime()
        self._assert_allocation_contract()
        arch = self._resolve_compile_arch()
-        next_dir = compiler.generate_config(self.lbm_cfg, n_objects=expected_n_objects)
-        try:
-            next_ptx = compiler.compile_kernel(arch=arch, build_dir=next_dir)
-            next_module = compiler.load_module(next_ptx)
-        except Exception:
-            compiler.cleanup_build_dir(next_dir)
-            raise
-        old_dir = self._build_dir
-        self._build_dir, self._ptx_path, self._module = next_dir, next_ptx, next_module
+        self._module = compiler.build_module(self.lbm_cfg, "uniform.cu", arch=arch,
+                                             n_objects=expected_n_objects)
        self._compiled_n_objects = int(expected_n_objects)

        # Reconnect field and stepper to new module
@@ -438,7 +428,6 @@ class Simulation:
        self.stepper._step_count = _prev_step_count

        self._assert_object_count_contract(expected_count=expected_n_objects)
-        compiler.cleanup_build_dir(old_dir)

    def recompile(self):
        """Re-generate config headers and recompile kernel.
@@ -715,8 +704,6 @@ class Simulation:
                self._module = None
                self._stream = None
            ctx.close()
-        compiler.cleanup_build_dir(getattr(self, "_build_dir", None))
-        self._build_dir = None
        self._closed = True

    def __enter__(self):
@@ -0,0 +1,66 @@
+"""Checkpoint save -> load round trip for a scene with passive bodies. Requires GPU.
+
+A fresh Simulation with the same scene, restored from the checkpoint must hold the same populations,
+flags, body count and step count, and must stay bitwise identical to the original
+after both continue for the same number of steps. Checkpoints with Jets are
+refused (documented limitation) and are not covered here.
+"""
+
+import os
+import tempfile
+import unittest
+
+import numpy as np
+
+from ._cases import make_sim
+
+# Three separated cylinders (pinball-like); cut-link receivers stay in bulk fluid.
+CIRCLES = (((40, 32), 6), ((56, 22), 6), ((56, 42), 6))
+
+
+def scene():
+    sim = make_sim()
+    sim.initialize()
+    for center, radius in CIRCLES:
+        sim.add_body("circle", center=center, radius=radius)
+    sim.sync_bodies()
+    return sim
+
+
+class CheckpointBodiesGPU(unittest.TestCase):
+    def test_round_trip_is_bitwise_and_continues_identically(self):
+        sim = scene()
+        sim.run(60)
+        with tempfile.NamedTemporaryFile(suffix=".h5", delete=False) as f:
+            path = f.name
+        restored = None
+        try:
+            saved = sim.save_checkpoint(path)
+            # The loader requires the same body count and geometry: rebuild the scene.
+            restored = scene()
+            restored.load_checkpoint(saved)
+
+            self.assertEqual(restored.bodies.count, len(CIRCLES))
+            self.assertEqual(restored.stepper.step_count, sim.stepper.step_count)
+            np.testing.assert_array_equal(restored.get_flags(), sim.get_flags())
+            a, b = sim.get_ddf(), restored.get_ddf()
+            self.assertEqual(a.dtype, b.dtype)
+            self.assertEqual(a.tobytes(), b.tobytes())
+
+            sim.run(40)
+            restored.run(40)
+            self.assertEqual(restored.stepper.step_count, sim.stepper.step_count)
+            a, b = sim.get_ddf(), restored.get_ddf()
+            self.assertTrue(np.isfinite(a).all())
+            self.assertEqual(a.tobytes(), b.tobytes())
+            for i in range(len(CIRCLES)):
+                self.assertTrue(np.all(np.isfinite(restored.read_force(i))))
+        finally:
+            if restored is not None:
+                restored.close()
+            os.unlink(path)
+            sim.close()
+
+
+if __name__ == "__main__":
+    unittest.main()
@@ -6,14 +6,14 @@ import pycuda.driver as cuda
 import pytest

 import CelerisLab
-from CelerisLab.cuda import compiler_v2
+from CelerisLab.cuda import compiler
 from ._cases import make_sim


 def test_source_provenance():
    root = Path(__file__).resolve().parents[2] / "src/CelerisLab"
    assert Path(CelerisLab.__file__).resolve().parent == root
-    assert Path(compiler_v2.kernel_path("kernel_v2.cu")).resolve() == root / "lbm/kernels/kernel_v2.cu"
+    assert Path(compiler.kernel_path("uniform.cu")).resolve() == root / "lbm/kernels/uniform.cu"


 def test_two_contexts_interleaved_and_non_lifo_close():
@@ -21,7 +21,7 @@ def test_two_contexts_interleaved_and_non_lifo_close():
    first = make_sim(nx=32, ny=16)
    second = make_sim(nx=48, ny=24, streaming="esopull", shifted=True)
    try:
-        assert first._build_dir != second._build_dir
+        assert first._module is not second._module
        assert cuda.Context.get_current() == before
        first.initialize()
        second.initialize()
@@ -16,6 +16,7 @@
 #define NZ 1
 #define NT 1
 #define COLLISION_MODEL 0
+#define BLOCK_ROOT_BC 2
 #define USE_LES 0
 #define RHO 1.f
 #define U0 .025f
@@ -40,7 +41,7 @@ inline float __fdiv_rn(float a,float b){return a/b;}
 #include "boundary/inlet_outlet.cuh"
 #include "step/one_step_double.cu"
 #else
-constexpr unsigned char KIND_OUTER=2;
+#include "core/cell_kind.cuh"
 #include "step/block_channel.cu"
 #endif
 int main(){
@@ -0,0 +1,46 @@
+"""Check historical source manifests against their own recorded commit, not today's checkout."""
+
+import io
+from pathlib import Path
+import subprocess
+import tarfile
+
+ROOT = Path(__file__).resolve().parents[2]
+
+
+def _require(commit):
+    """Missing recorded history is an explicit failure, never a skip or a repin."""
+    found = subprocess.run(
+        ["git", "-C", str(ROOT), "cat-file", "-e", commit + "^{commit}"],
+        capture_output=True,
+    )
+    if found.returncode:
+        raise AssertionError(
+            f"Required history missing: commit {commit} is not in this clone; "
+            "fetch it instead of repinning or substituting"
+        )
+
+
+def tracked(commit, prefix="src"):
+    """Tracked paths under prefix at the recorded commit."""
+    _require(commit)
+    return set(
+        subprocess.check_output(
+            ["git", "-C", str(ROOT), "ls-tree", "-r", "--name-only", commit, prefix],
+            text=True,
+        ).splitlines()
+    )
+
+
+def historical_tree(commit, dest):
+    """Extract the tracked tree of the recorded commit into dest."""
+    _require(commit)
+    archive = subprocess.check_output(
+        ["git", "-C", str(ROOT), "archive", "--format=tar", commit]
+    )
+    dest = Path(dest)
+    dest.mkdir(parents=True, exist_ok=True)
+    extra = {"filter": "data"} if hasattr(tarfile, "data_filter") else {}
+    with tarfile.open(fileobj=io.BytesIO(archive)) as tar:
+        tar.extractall(dest, **extra)
+    return dest
@@ -1,104 +0,0 @@
-"""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
-import unittest
-
-import numpy as np
-
-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."""
-    # 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)
-        or isinstance(node, ast.ImportFrom) and node.module == "cuda")]
-    module = types.ModuleType("CelerisLab.lbm._test_host_block_stepper")
-    module.__package__ = "CelerisLab.lbm"
-    module.__dict__.update(cuda=Mock(), compiler=Mock())
-    exec(compile(tree, str(path), "exec"), module.__dict__)
-    return module
-
-
-def backend():
-    host, log = host_module(), []
-    host.cuda.mem_get_info.return_value = (1 << 40, 1 << 40)
-    host.cuda.mem_alloc.side_effect = lambda size: Mock(size=size)  # distinct buffers
-    module = Mock()
-    module.get_global.return_value = (Mock(), 20)
-    module.get_function.side_effect = lambda name: Mock(
-        side_effect=lambda *a, **k: log.append((name, a[:2])))
-    host.compiler.load_module.return_value = module
-    cfg = LBMConfig(nx=16, ny=16, viscosity=.1)
-    run = host.PeriodicBlockStepper(cfg, BlockLayout2D.build(16, 16, 8, 8),
-                                    context=Mock(activate=nullcontext, sm_arch="sm_70"))
-    run._stream.synchronize.side_effect = lambda: log.append(("sync", ()))
-    run.initialize(np.full((9, 16, 16), 1 / 9))
-    log.clear()
-    return run, log
-
-
-class BlockStepSync(unittest.TestCase):
-    def test_step_n_waits_once_and_matches_n_single_steps(self):
-        many, log_many = backend()
-        many.step(3)
-        single, log_single = backend()
-        for _ in range(3):
-            single.step(1)
-        names = [name for name, _ in log_many]
-        self.assertEqual(names.count("sync"), 1)
-        self.assertEqual(names[-1], "sync")
-        self.assertEqual(names.count("BlockStep"), 3)
-        # Same launch order and buffer roles as three step(1) calls.
-        strip = lambda log: [(n, tuple(id(x) for x in a)) for n, a in log if n != "sync"]
-        roles = lambda run, log: [(n, tuple({id(run._current): "c", id(run._spare): "s"}.get(x, "-")
-                                            for x in ids)) for n, ids in strip(log)]
-        self.assertEqual([n for n, _ in strip(log_many)], [n for n, _ in strip(log_single)])
-        self.assertEqual(roles(many, log_many), roles(single, log_single))
-        self.assertIsNot(many._current, many._spare)
-        self.assertEqual((many.step_count, single.step_count), (3, 3))
-
-    def test_step_zero_launches_and_waits_for_nothing(self):
-        run, log = backend()
-        run.step(0)
-        self.assertEqual((log, run.step_count), ([], 0))
-
-    def test_failure_keeps_entry_step_count(self):
-        run, log = backend()
-        run.step(1)
-        run._stream.synchronize.side_effect = RuntimeError("device fault")
-        with self.assertRaises(RuntimeError):
-            run.step(2)
-        self.assertEqual((run.step_count, run._status), (1, "failed"))
-
-    def test_launch_failure_mid_batch_fails_closed_without_waiting(self):
-        run, log = backend()
-        calls = []
-        step = run._step
-
-        def failing(*args, **kwargs):
-            calls.append(1)
-            if len(calls) == 2:
-                raise RuntimeError("launch failure")
-            return step(*args, **kwargs)
-
-        run._step = failing
-        with self.assertRaises(RuntimeError):
-            run.step(3)
-        self.assertNotIn("sync", [name for name, _ in log])
-        self.assertEqual((run.step_count, run._status), (0, "failed"))
-        with self.assertRaises(RuntimeError):
-            run.download()
-
-
-if __name__ == "__main__":
-    unittest.main()
@@ -83,6 +83,32 @@ assert abs(sum(compute_feq_d2q9(1.,0.,0.)) - 1.) < 1e-6


 class BuildContracts(unittest.TestCase):
+    def test_block_root_modes_are_rendered_and_recorded(self):
+        with tempfile.TemporaryDirectory() as root:
+            for mode in (0, 1, 2):
+                with self.subTest(mode=mode):
+                    build = compiler.generate_config(LBMConfig(), build_root=root, block_root_bc=mode)
+                    try:
+                        method = Path(build, "config/config_method.h").read_text()
+                        self.assertIn(f"#define BLOCK_ROOT_BC {mode}\n", method)
+                        metadata = json.loads(Path(build, ".celeris-build").read_text())
+                        self.assertEqual(metadata["block_root_bc"], mode)
+                        self.assertNotIn("block_root_bc", metadata["config"])
+                    finally:
+                        compiler.cleanup_build_dir(build)
+            build = compiler.generate_config(LBMConfig(), build_root=root)
+            try:
+                self.assertEqual(json.loads(Path(build, ".celeris-build").read_text())["block_root_bc"], 0)
+            finally:
+                compiler.cleanup_build_dir(build)
+
+    def test_invalid_block_root_mode_fails_before_snapshot(self):
+        for mode in (-1, 3, True, False, 1.0, "1", None):
+            with self.subTest(mode=mode), patch.object(compiler.tempfile, "mkdtemp") as allocate:
+                with self.assertRaisesRegex(ValueError, "block_root_bc"):
+                    compiler.generate_config(LBMConfig(), block_root_bc=mode)
+                allocate.assert_not_called()
+
    @staticmethod
    def fake_compile(command, check):
        source = Path(command[-3])
@@ -100,8 +126,7 @@ class BuildContracts(unittest.TestCase):
            def work(nx, count):
                build = compiler.generate_config(LBMConfig(nx=nx), count, build_root=root)
                barrier.wait()
-                # Compatibility form is context-local even without explicit path.
-                output = compiler.compile_kernel()
+                output = compiler.compile_kernel(build_dir=build)
                data = json.loads(Path(output).read_text())
                self.assertEqual(data, {"nx": nx, "objects": count})
                metadata = json.loads(Path(output + ".json").read_text())
@@ -130,7 +155,7 @@ class BuildContracts(unittest.TestCase):
            compiler.cleanup_build_dir(a)
            compiler.cleanup_build_dir(b)
            with self.assertRaises(RuntimeError):
-                compiler.compile_kernel()
+                compiler.compile_kernel(build_dir=a)  # removed snapshot is never compiled

    def test_failed_compile_cannot_implicitly_load_previous_artifact(self):
        with tempfile.TemporaryDirectory() as root:
@@ -143,8 +168,7 @@ class BuildContracts(unittest.TestCase):
                with self.assertRaises(subprocess.CalledProcessError):
                    compiler.compile_kernel(build_dir=build)
            self.assertEqual(Path(previous).read_bytes(), old_bytes)
-            with self.assertRaisesRegex(RuntimeError, "No compiled PTX"):
-                compiler.load_module()
+            self.assertEqual(sorted(Path(build).glob("uniform-*.ptx")), [])  # no partial output
            compiler.cleanup_build_dir(build)


@@ -50,7 +50,7 @@ def test_masks_and_minimal_footprint():
 def test_actual_closure_read_conflicts():
    pair, layout, targets = transfer((6, 6, 18, 14))
    validate_channel_transfers(pair, layout, targets)
-    for target in (layout.slot(layout.owner[5, -1]-1, 2),
+    for target in (layout.slot(int(layout.owner[5, -1]) - 1, 2),  # integer under NumPy 1 and 2
                   layout.slot(layout.owner[1, 4], 4),
                   layout.slot(layout.owner[-2, 4], 3)):
        with pytest.raises(ValueError, match="immutable"):
@@ -74,7 +74,7 @@ def test_evolved_ghost_snapshot_and_physical_download():
    from .test_refinement_three_level_multipatch import host_module
    host = host_module()
    storage = {}
-    host.compiler.load_module.return_value.get_global.side_effect = lambda name: (name, 12 if name == "d_block_channel" else 20)
+    host.compiler.build_module.return_value.get_global.side_effect = lambda name: (name, 12 if name == "d_block_channel" else 20)
    host.cuda.mem_alloc.side_effect = lambda size: Mock()
    host.cuda.memcpy_htod.side_effect = lambda ptr, a: storage.__setitem__(id(ptr), np.array(a, copy=True))
    host.cuda.memcpy_dtoh.side_effect = lambda a, ptr: np.copyto(a, storage[id(ptr)])
@@ -96,6 +96,7 @@ def test_evolved_ghost_snapshot_and_physical_download():
        arrays[0][0, 0, 0] = np.nan
        with pytest.raises(ValueError, match="finite"):
            backend.initialize(*arrays)
-        assert backend._levels[0].kernel == host.compiler.load_module.return_value.get_function("BlockChannelStep")
+        assert backend._levels[0].kernel == host.compiler.build_module.return_value.get_function("BlockChannelStep")
+        assert [call.kwargs["block_root_bc"] for call in host.compiler.build_module.call_args_list] == [2, 0]
    finally:
        backend.close()
@@ -8,6 +8,7 @@ import numpy as np
 ROOT = Path(__file__).resolve().parents[2]
 REPO = ROOT
 sys.path.insert(0, str(ROOT / "tests/validation"))
+sys.path.insert(0, str(ROOT / "tests/unit"))
 sys.path.insert(0, str(REPO / "src"))
 from run_jet_circle_cv import (
    CX,
@@ -19,6 +20,7 @@ from run_jet_circle_cv import (
    cv_from_macros,
    aperture_probe,
    measured_aperture_flux,
+    PIN,
    verify,
 )

@@ -166,7 +168,13 @@ class Host(unittest.TestCase):
        self.assertEqual(a["steps"] * 4, b["steps"])

    def test_source_pin(self):
-        verify(REPO)
+        from tempfile import TemporaryDirectory
+        from historical_tree import historical_tree
+
+        with self.assertRaisesRegex(RuntimeError, "Source hash mismatch"):
+            verify(REPO)
+        with TemporaryDirectory() as directory:
+            verify(historical_tree(PIN, directory))


 if __name__ == "__main__":
@@ -15,6 +15,10 @@ sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
 from run_uniform_les_flow import source_guard, decode_macros, HERE, HEADER

 ROOT = HERE.parents[1]
+sys.path.insert(0, str(ROOT / "tests/unit"))
+from historical_tree import historical_tree, tracked
+
+A135 = "a135252ae78b47b57d7e77d19001700b3bedac8b"


 class GuardTests(unittest.TestCase):
@@ -42,19 +46,23 @@ class GuardTests(unittest.TestCase):
            decode_macros(invalid, w, cx, cy)

    def test_current_corrected_pins(self):
-        self.assertEqual(
-            source_guard(ROOT, "corrected", "a135252")["base_commit"],
-            "a135252ae78b47b57d7e77d19001700b3bedac8b",
-        )
+        with TemporaryDirectory() as directory:
+            root = historical_tree(A135, directory)
+            with patch("run_uniform_les_flow.subprocess.check_output", return_value="historical-tree\n"):
+                identity = source_guard(root, "corrected", "a135252")
+            # base_commit comes from the pins file; the mock only feeds checkout_commit.
+            self.assertEqual(identity["base_commit"], A135)
+            self.assertEqual(identity["checkout_commit"], "historical-tree")

    def test_corrected_and_tamper_guard(self):
        pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
        with TemporaryDirectory() as directory:
-            root = Path(directory)
+            origin = historical_tree(A135, Path(directory) / "origin")
+            root = Path(directory) / "fixture"
            for name in pins["sources"]:
                dest = root / name
                dest.parent.mkdir(parents=True, exist_ok=True)
-                shutil.copyfile(ROOT / name, dest)
+                shutil.copyfile(origin / name, dest)
            with patch(
                "run_uniform_les_flow.subprocess.check_output",
                return_value="isolated-fixture\n",
@@ -87,23 +95,26 @@ class GuardTests(unittest.TestCase):

    def test_current_manifest_pins_actual_corrected_tree(self):
        pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
-        tracked = subprocess.check_output(["git", "-C", str(ROOT), "ls-files", "src"], text=True)
-        self.assertEqual(set(pins["sources"]), set(tracked.splitlines()))
-        for name, expected in pins["sources"].items():
-            self.assertEqual(hashlib.sha256((ROOT / name).read_bytes()).hexdigest(), expected)
-        self.assertEqual(pins["sources"][HEADER], pins["candidate_header_sha256"])
-        self.assertNotEqual(pins["sources"][HEADER], pins["baseline_header_sha256"])
-        with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
-            source_guard(ROOT, "baseline", "a135252")
+        self.assertEqual(pins["base_commit"], A135)
+        self.assertEqual(set(pins["sources"]), tracked(A135))
+        with TemporaryDirectory() as directory:
+            root = historical_tree(A135, directory)
+            for name, expected in pins["sources"].items():
+                self.assertEqual(hashlib.sha256((root / name).read_bytes()).hexdigest(), expected)
+            self.assertEqual(pins["sources"][HEADER], pins["candidate_header_sha256"])
+            self.assertNotEqual(pins["sources"][HEADER], pins["baseline_header_sha256"])
+            with self.assertRaisesRegex(RuntimeError, "Source pin mismatch"):
+                source_guard(root, "baseline", "a135252")

    def test_named_baseline_accepts_only_reversed_header(self):
        pins = json.loads((HERE / "data/les_flow/source_pins_a135252.json").read_text())
        with TemporaryDirectory() as directory:
-            root = Path(directory)
+            origin = historical_tree(A135, Path(directory) / "origin")
+            root = Path(directory) / "fixture"
            for name in pins["sources"]:
                dest = root / name
                dest.parent.mkdir(parents=True, exist_ok=True)
-                shutil.copyfile(ROOT / name, dest)
+                shutil.copyfile(origin / name, dest)
            subprocess.run(["git", "apply", "--reverse", str(HERE / "les/data/smagorinsky_normalization.patch")], cwd=root, check=True)
            with patch("run_uniform_les_flow.subprocess.check_output", return_value="baseline-fixture\n"):
                source_guard(root, "baseline", "a135252")
@@ -135,6 +135,6 @@ def test_d2_entrypoint_compiles_in_isolation(tmp_path, model, shift, storage):
            source=re.sub(r"(#define\s+"+key+r"\s+)\S+",lambda match: match[1]+str(value),source)
        path.write_text(source)
    result=subprocess.run([nvcc,"-ptx","-arch=sm_70","-std=c++17",
-                           str(copy/"kernel_v2.cu"),"-o",str(tmp_path/"kernel.ptx")],
+                           str(copy/"uniform.cu"),"-o",str(tmp_path/"kernel.ptx")],
                          capture_output=True,text=True)
    assert result.returncode == 0, result.stderr
@@ -9,7 +9,7 @@ import pytest

 import CelerisLab
 from CelerisLab.config import LBMConfig, load_lbm_config
-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.common.checkpoint import save_checkpoint, load_checkpoint
 from .test_numerical_operators import compile_run
 from .test_state_regressions import device, make_field, manager
@@ -18,7 +18,7 @@ from .test_state_regressions import device, make_field, manager
 def test_planar_helper_moments_and_known_populations(tmp_path):
    root = Path(__file__).resolve().parents[2]
    assert Path(CelerisLab.__file__).resolve() == root/'src/CelerisLab/__init__.py'
-    assert Path(compiler.kernel_path('kernel_v2.cu')).resolve() == root/'src/CelerisLab/lbm/kernels/kernel_v2.cu'
+    assert Path(compiler.kernel_path('uniform.cu')).resolve() == root/'src/CelerisLab/lbm/kernels/uniform.cu'
    lines = compile_run(tmp_path, 'planar_velocity.cpp', dict(WALL_NORMAL_VELOCITY='0.0f'))
    cx = np.array([0,1,-1,0,0,1,-1,1,-1])
    cy = np.array([0,0,0,1,-1,1,-1,-1,1])
@@ -162,7 +162,7 @@ class ProjectedKernelTests(unittest.TestCase):
        targets = child.owner.ravel()[p.ghost_fine]
        donors = parent.owner.ravel()[p.ghost_coarse]
        source = (Path(__file__).resolve().parents[2] /
-                  "src/CelerisLab/lbm/kernels/blocks.cu").read_text()
+                  "src/CelerisLab/lbm/kernels/step/block_transfer.cu").read_text()
        kernel = source.split('extern "C" __global__ void BlockExplodeProjected(',1)[1]
        kernel = 'void BlockExplodeProjected(' + kernel.split('// Targets are receiver-local SOURCE slots',1)[0]
        def array(name, values, kind):
@@ -29,7 +29,7 @@ def backend(**options):
            name=name, side_effect=lambda *a, **kw: calls.append((name, a, kw)))
        return module

-    host.compiler.load_module.side_effect = load
+    host.compiler.build_module.side_effect = lambda *a, **k: load(None)
    host.cuda.memcpy_htod.side_effect = lambda ptr, a: storage.__setitem__(id(ptr), np.array(a, copy=True))
    host.cuda.memcpy_dtoh.side_effect = lambda a, ptr: np.copyto(a, storage[id(ptr)])
    cfg = options.pop('cfg', CFG)
@@ -25,7 +25,7 @@ def backend(observe=True, obj=None):
    module = Mock()
    module.get_global.return_value = (Mock(), 20)
    module.get_function.side_effect = lambda name: Mock(side_effect=lambda *a, **k: calls.append((name, a, k)))
-    host.compiler.load_module.return_value = module
+    host.compiler.build_module.return_value = module
    host.cuda.mem_alloc.side_effect = lambda n: Mock(size=n)
    host.cuda.memcpy_dtoh.side_effect = lambda array, pointer: array.fill(0)
    ctx = Mock(activate=nullcontext, sm_arch='sm_70')
@@ -5,6 +5,7 @@ import pytest

 ROOT = Path(__file__).resolve().parents[2]
 sys.path.insert(0, str(ROOT / "tests/validation"))
+sys.path.insert(0, str(ROOT / "tests/unit"))
 import re100_helpers as h
 import periodic as p

@@ -62,13 +63,16 @@ def test_periodic_synthetic():
    assert not p.assess(rows()[:100])["periodic"]


-def test_production_unmodified():
+def test_production_unmodified(tmp_path):
+    from historical_tree import historical_tree
+
    m = json.loads(
        (ROOT / "tests/validation/data/dfg_re100/source_manifest.json").read_text()
    )
    assert m["commit"] == h.PIN
+    origin = historical_tree(m["commit"], tmp_path)
    for name, digest in m["production_sha256"].items():
-        assert hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == digest, name
+        assert hashlib.sha256((origin / name).read_bytes()).hexdigest() == digest, name


 def test_reference():
@@ -7,7 +7,10 @@ import numpy as np
 import unittest

 sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "validation"))
+sys.path.insert(0, str(Path(__file__).resolve().parent))
 import run_refined_dfg as r
+from historical_tree import historical_tree
+from tempfile import TemporaryDirectory

 h, periodic = r.helpers(r.DEFAULT_SOURCE)

@@ -17,11 +20,24 @@ def test_source_pin():
    assert historical["commit"].startswith("cd3a657")
    mrt_history = json.loads((r.HERE / "data/refined_dfg/source_manifest_mrt.json").read_text())
    assert mrt_history["commit"].startswith("3ea131e")
-    assert r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")["commit"].startswith("a233c90")
+    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
+        r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")
    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
        r.verify_source(r.DEFAULT_SOURCE, "MRT")
    with unittest.TestCase().assertRaisesRegex(RuntimeError, "Source mismatch"):
        r.verify_source(r.DEFAULT_SOURCE, "SRT")
+    with TemporaryDirectory() as directory:
+        for collision, mode, prefix in (
+            ("SRT", "last_base_step", "cd3a657"),
+            ("MRT", "last_base_step", "3ea131e"),
+            ("MRT", "window", "a233c90"),
+        ):
+            name = {"SRT": "source_manifest.json", "MRT": "source_manifest_mrt.json"}[collision]
+            name = "source_manifest_window.json" if mode == "window" else name
+            commit = json.loads((r.HERE / "data/refined_dfg" / name).read_text())["commit"]
+            assert commit.startswith(prefix)
+            root = historical_tree(commit, Path(directory) / commit)
+            assert r.verify_source(root, collision, mode)["commit"] == commit


 def test_matched_units_and_planes():
@@ -141,16 +157,32 @@ def test_axis_specific_sampling_phase():
    np.testing.assert_allclose(fine_y - ((k + 43) - 0.5) * 0.00125, 0, atol=1e-15)


+def frozen_manifest(observation_mode):
+    """Frozen MRT manifest JSON as recorded, without hashing the working tree."""
+    name = "source_manifest_window.json" if observation_mode == "window" else "source_manifest_mrt.json"
+    return json.loads((r.HERE / "data/refined_dfg" / name).read_text())
+
+
+def selected_manifest(collision, observation_mode):
+    """The frozen manifest verify_source selects for (collision, observation_mode).
+
+    Source identity is owned by test_source_pin; this reads no checkout files.
+    """
+    return r.source_manifest(collision, observation_mode)
+
+
 def test_full_requires_completed_pilot():
    with unittest.TestCase().assertRaisesRegex(ValueError, "pilot-results"):
-        r.verify_pilot(None, r.verify_source(r.DEFAULT_SOURCE, "MRT", "window"))
+        r.verify_pilot(None, frozen_manifest("window"))


 def pilot_fixture(path, observation_mode="window", jet_scenario="none", case_selection="both", uniform_diameter=80, placement="wide"):
    """Synthetic contract fixture only; never a physical result or run receipt."""
+    # Receipt-logic fixture: the frozen manifest verify_source returns for a matching tree.
+    # Source identity itself is checked only by test_source_pin.
    manifest = (r.verify_source(r.DEFAULT_SOURCE, "MRT", "window", jet_scenario)
-                if observation_mode == "window" else
-                json.loads((r.HERE / "data/refined_dfg/source_manifest_mrt.json").read_text()))
+                if observation_mode == "window" and jet_scenario != "none" else
+                frozen_manifest(observation_mode))
    selection = dict(case_selection=case_selection, uniform_diameter=uniform_diameter, placement=placement)
    (path / "plan.json").write_text(
        json.dumps(
@@ -348,7 +380,7 @@ def test_record_actual_checkout_separately():
    actual = checkout_provenance(r.DEFAULT_SOURCE)
    assert len(actual["checkout_revision"]) == 40
    assert isinstance(actual["checkout_dirty"], bool)
-    assert r.verify_source(r.DEFAULT_SOURCE, "MRT", "window")["production_origin_commit"].startswith("d930")
+    assert selected_manifest("MRT", "window")["production_origin_commit"].startswith("d930")


 def test_window_plan_preserves_physics():
@@ -358,7 +390,7 @@ def test_window_plan_preserves_physics():
        assert baseline.pop("observation_mode") == "last_base_step"
        assert window.pop("observation_mode") == "window"
        assert baseline == window
-        assert r.verify_source(r.DEFAULT_SOURCE, collision, "window")["observation_mode"] == "window"
+        assert selected_manifest(collision, "window")["observation_mode"] == "window"
    with unittest.TestCase().assertRaises(ValueError):
        r.verify_source(r.DEFAULT_SOURCE, "MRT", "unknown")

@@ -22,13 +22,16 @@ class RefinementMRTHost(unittest.TestCase):

    def test_srt_statements_and_guards_preserved(self):
        """MRT is an additive compile-time branch; SRT statements and other guards remain."""
-        blocks = (KERNELS / "blocks.cu").read_text()
+        blocks = (KERNELS / "blocked.cu").read_text() + (KERNELS / "step/block_step.cu").read_text()
        channel = (KERNELS / "step/block_channel.cu").read_text()
-        refinement = (KERNELS / "refinement.cu").read_text()
+        block_collision = (KERNELS / "operators/block_collision.cuh").read_text()
+        self.assertIn('#include "operators/block_collision.cuh"', blocks)
+        self.assertEqual(blocks.count("block_collide(f);"), 1)
+        refinement = (KERNELS / "refinement.cu").read_text() + (KERNELS / "step/patch_step.cu").read_text()
        for text in ("compute_feq_from_momentum(drho, jx, jy, ux, uy, feq);",
                     "compute_feq(rho, ux, uy, feq);",
                     "collide_srt_no_force(f, feq, d_params.omega);"):
-            self.assertIn(text, blocks)
+            self.assertIn(text, block_collision)
        self.assertIn("collide_srt(f, feq, forcing, omega);", channel)
        self.assertIn("collide_srt_no_force(f, feq, d_params.omega);", refinement)
        for source in (blocks, refinement):
@@ -40,10 +43,11 @@ class RefinementMRTHost(unittest.TestCase):
        # 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):
+        for source in (blocks + "\n" + block_collision, channel, refinement):
            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)
+            self.assertNotRegex(source, r"compute_rho_u\(f, rho, ux, uy\);\s*collide_mrt")


 if __name__ == "__main__": unittest.main()
@@ -22,12 +22,13 @@ def host_module():
    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)
        or isinstance(node,ast.ImportFrom) and node.module in (
-            'cuda','block_stepper','curved_links'))]
+            'cuda','curved_links'))]
    module=types.ModuleType('CelerisLab.lbm._test_host_refinement')
    module.__package__='CelerisLab.lbm'
    from CelerisLab.body.jets import surface_link_velocity
-    module.__dict__.update(cuda=Mock(),compiler=Mock(),MAX_GRID_DIM=65535,
+    module.__dict__.update(cuda=Mock(),compiler=Mock(),
                           surface_link_velocity=surface_link_velocity)
+    module.cuda.mem_get_info.return_value=(1<<40,1<<40)  # free-memory preflight passes
    exec(compile(tree,str(path),'exec'),module.__dict__)
    return module

@@ -193,9 +194,9 @@ def mock_backend(*, children=CHILDREN, bodies=False, fail_compile=False, fail_up
        gpu.get_function.side_effect=lambda name: Mock(name=name)
        gpu.get_global.side_effect=lambda name: (Mock(),20 if name=='d_params' else 16)
        return gpu
-    module.compiler.load_module.side_effect=load
+    module.compiler.build_module.side_effect=lambda *a,**k: load(None)
    if fail_compile:
-        module.compiler.compile_kernel.side_effect=['first','second',RuntimeError('L2 compile failed')]
+        module.compiler.build_module.side_effect=[load(None),load(None),RuntimeError('L2 compile failed')]
    if fail_upload:
        def upload(*args):
            if len(allocations)>20:
@@ -86,7 +86,7 @@ class RuntimeChildRemap(unittest.TestCase):
            before = run.download()
            old_allocations = list(allocations)
            if failure == 'construction':
-                operation = patch.object(module.compiler, 'compile_kernel', side_effect=RuntimeError('candidate compile'))
+                operation = patch.object(module.compiler, 'build_module', side_effect=RuntimeError('candidate compile'))
            else:
                initialize = type(run).initialize
                def fail(candidate, *arrays):
@@ -62,7 +62,7 @@ class ShiftedJets(unittest.TestCase):
                    level_blocks=list(blocks),level_objects=[(),(),(body(),)],
                    level_jets=[(),(),(jet(),)])
        module.cuda.mem_alloc.assert_not_called()
-        module.compiler.compile_kernel.assert_not_called()
+        module.compiler.build_module.assert_not_called()

    def test_switch_upload_only_then_raw_remap_keeps_live_velocity(self):
        run,module,allocations,storage,objects=ready();clone=None
@@ -1,7 +1,4 @@
 """CPU host contracts; mocked CUDA executes real lifecycle, not numerical kernels."""
-import ast
-import types
-from pathlib import Path
 from contextlib import nullcontext
 from dataclasses import replace
 import unittest
@@ -78,13 +75,15 @@ class ShiftedStorage(unittest.TestCase):

    def test_shifted_density_and_tag_validation(self):
        old,new = topology(A),topology(B)
-        a = tuple(np.where(m[None], np.float32(-.001), np.nan).repeat(9,axis=0) for m in old.leaf_masks)
+        a = tuple(np.where(m[None], np.float32(-.001), np.float32(np.nan)).repeat(9,axis=0)
+                  for m in old.leaf_masks)  # FP32 under NumPy 1 and 2 promotion rules
        remap(old,new,a)
        with self.assertRaisesRegex(ValueError,'representation'): remap(old,new,a,False)
        with self.assertRaisesRegex(ValueError,'FP32'): remap(old,new,tuple(x.astype(np.float64) for x in a))
        a[0][:,old.leaf_masks[0]] = -1
        with self.assertRaisesRegex(ValueError,'positive density'): remap(old,new,a)
-        rest = tuple(np.where(m[None],np.float32(0),np.nan).repeat(9,axis=0) for m in old.leaf_masks)
+        rest = tuple(np.where(m[None],np.float32(0),np.float32(np.nan)).repeat(9,axis=0)
+                     for m in old.leaf_masks)
        remap(old,new,rest)
        with self.assertRaisesRegex(ValueError,'positive density'):
            remap_synchronized(old,new,rest,base_count=7,level_steps=(7,14,28),phase='postcollision_g')
@@ -133,27 +132,17 @@ class ShiftedStorage(unittest.TestCase):
            with self.assertRaisesRegex(ValueError,'all-blocked|blocked refined levels below'):
                module.FixedPatchStepper(cfg,context=None,refinements=[(4,4,20,28)],
                                         level_blocks=blocks)
-        module.compiler.compile_kernel.assert_not_called()
+        module.compiler.build_module.assert_not_called()

    def test_uniform_block_codec_and_raw_snapshot(self):
-        path = Path(backend.r.__file__).with_name('block_stepper.py')
-        tree = ast.parse(path.read_text())
-        tree.body = [n for n in tree.body if not (
-            isinstance(n, ast.Import) and any(a.name == 'pycuda.driver' for a in n.names)
-            or isinstance(n, ast.ImportFrom) and n.module == 'cuda')]
-        module = types.ModuleType('CelerisLab.lbm._test_block_host')
-        module.__package__ = 'CelerisLab.lbm'
-        cuda, compiler = Mock(), Mock()
-        module.__dict__.update(cuda=cuda,compiler=compiler)
-        exec(compile(tree,str(path),'exec'),module.__dict__)
-        cuda.mem_get_info.return_value=(10**9,10**9)
+        module = backend.host_module()  # single blocked level: FixedPatchStepper(refinements=[])
+        cuda, compiler = module.cuda, module.compiler
        cuda.mem_alloc.side_effect=lambda size: Mock()
        gpu=Mock();gpu.get_global.return_value=(Mock(),20)
-        compiler.load_module.return_value=gpu
+        compiler.build_module.return_value=gpu
        context=Mock(activate=nullcontext,sm_arch='sm_70')
-        layout=module.BlockLayout2D.build(16,16,4,4)
        cfg=replace(backend.CFG,nx=16,ny=16,ddf_shifting=True)
-        run=module.PeriodicBlockStepper(cfg,layout,context=context)
+        run=module.FixedPatchStepper(cfg,context=context,refinements=[],level_blocks=[(4,4)])
        storage={}
        cuda.memcpy_htod.side_effect=lambda p,a: storage.__setitem__(id(p),np.array(a,copy=True))
        cuda.memcpy_dtoh.side_effect=lambda a,p: np.copyto(a,storage[id(p)])
@@ -164,7 +153,8 @@ class ShiftedStorage(unittest.TestCase):
            self.assertTrue(raw.shifted)
            self.assertEqual(len(raw.values),1)
            self.assertTrue((raw.values[0][1]==np.float32(2**-35)).all())
-            np.testing.assert_array_equal(run.download(),physical)
+            (down,)=run.download()
+            np.testing.assert_array_equal(down,physical)
        finally: run.close()

    def test_passive_shifted_body_raw_remap_preserves_support(self):
@@ -36,12 +36,10 @@ def make_simulation(monkeypatch):
    sim = module.Simulation.__new__(module.Simulation)
    sim.ctx = Owner()
    sim._closed = False
-    sim._build_dir = "private-build"
    for name in ("stepper", "bodies", "field", "_module", "_stream"):
        setattr(sim, name, object())
    monkeypatch.setattr(module, "cuda", SimpleNamespace(
        Context=SimpleNamespace(synchronize=lambda: None)))
-    monkeypatch.setattr(module.compiler, "cleanup_build_dir", lambda path: None)
    return sim


@@ -57,7 +55,7 @@ def test_close_inside_activation_preserves_resources(monkeypatch):
    assert sim._closed and sim.ctx.is_closed


-@pytest.mark.parametrize("failure", ["synchronize", "detach", "directory"])
+@pytest.mark.parametrize("failure", ["synchronize", "detach"])
 def test_close_failure_is_retryable(monkeypatch, failure):
    sim = make_simulation(monkeypatch)
    if failure == "detach":
@@ -68,17 +66,12 @@ def test_close_failure_is_retryable(monkeypatch, failure):
            calls.append(args)
            if len(calls) == 1:
                raise RuntimeError(failure + " failed")
-        if failure == "synchronize":
-            monkeypatch.setattr(module.cuda.Context, "synchronize", fail_once)
-        else:
-            monkeypatch.setattr(module.compiler, "cleanup_build_dir", fail_once)
+        monkeypatch.setattr(module.cuda.Context, "synchronize", fail_once)
    with pytest.raises(RuntimeError, match="failed"):
        sim.close()
    assert not sim._closed
-    assert sim._build_dir == "private-build"
    sim.close()
    assert sim._closed and sim.ctx.is_closed
-    assert sim._build_dir is None
    sim.close()


@@ -0,0 +1,218 @@
+"""Host contract of FixedPatchStepper(refinements=[]): one blocked level (mocked driver).
+
+Per step: halo refresh -> static BFL (circles only) -> BlockStep with a null
+kind, one stream wait per step(n) call, NaN spare after initialize. An upload,
+launch, wait or download failure fails the instance and keeps its step count;
+close retains resources for retry when context activation fails.
+"""
+from contextlib import nullcontext
+from unittest.mock import MagicMock, Mock
+import unittest
+
+import numpy as np
+
+from CelerisLab.body.geometry.circle import CircleGeometry
+from CelerisLab.body.objects import SimObject
+from CelerisLab.config import LBMConfig
+from tests.unit.test_refinement_three_level_multipatch import host_module
+
+
+def circle(center=(15.5, 11.5), radius=4.3):
+    return SimObject(0, geometry=CircleGeometry(*center, radius), center=center, radius=radius)
+
+
+def backend(objects=(), failing_kernel=None, initialize=True):
+    host, log = host_module(), []
+    host.armed = False  # set True to make failing_kernel raise
+
+    def launch(name):
+        def run(*a, **k):
+            log.append((name, a))
+            if name == failing_kernel and host.armed:
+                raise RuntimeError(name + " launch failed")
+        return run
+
+    module = Mock()
+    module.get_global.return_value = (Mock(), 20)
+    module.get_function.side_effect = lambda name: Mock(side_effect=launch(name))
+    host.compiler.build_module.return_value = module
+    host.cuda.mem_alloc.side_effect = lambda size: Mock(size=size)
+    host.cuda.memcpy_htod.side_effect = lambda pointer, array: log.append(("upload", (pointer, array)))
+    run = host.FixedPatchStepper(LBMConfig(nx=32, ny=24, viscosity=.1),
+                                 context=Mock(activate=nullcontext, sm_arch="sm_70"),
+                                 refinements=[], level_blocks=[(12, 16)], level_objects=[objects])
+    run._stream.synchronize.side_effect = lambda: log.append(("sync", ()))
+    log.clear()
+    if initialize:
+        run.initialize(np.full(run._levels[0].shape, 1 / 9, np.float32))
+    return run, log, host
+
+
+def state():
+    return np.full((9, 24, 32), 1 / 9, np.float32)
+
+
+class SingleLevel(unittest.TestCase):
+    def test_spare_starts_nan_and_current_holds_state(self):
+        run, log, _ = backend()
+        level = run._levels[0]
+        uploads = {id(a[0]): a[1] for name, a in log if name == "upload"}
+        self.assertTrue(np.isnan(uploads[id(level.spare)]).all())
+        self.assertTrue(np.isfinite(level.layout.to_dense(uploads[id(level.current)])).all())
+
+    def test_launch_order_per_step(self):
+        for objects, order in (((), ["BlockHaloCopy", "BlockStep"]),
+                               ((circle(),), ["BlockHaloCopy", "BlockCurvedBoundary", "BlockStep"])):
+            with self.subTest(bodies=len(objects)):
+                run, log, _ = backend(objects)
+                log.clear()
+                run.step(2)
+                self.assertEqual([name for name, _ in log], order * 2 + ["sync"])
+                step = [a for name, a in log if name == "BlockStep"][0]
+                self.assertEqual(int(step[4]), 0)  # null kind: every core cell ACTIVE
+
+    def test_circle_solids_are_not_state(self):
+        run, _, _ = backend((circle(),))
+        solid = run._level_boundaries[0].tables["solid"]
+        self.assertTrue(solid.any())
+        np.testing.assert_array_equal(run._state[0], ~solid)
+
+    def test_initialize_upload_failure_fails_instance(self):
+        run, _, host = backend(initialize=False)
+        host.cuda.memcpy_htod.side_effect = RuntimeError("upload failed")
+        with self.assertRaisesRegex(RuntimeError, "upload failed"):
+            run.initialize(state())
+        with self.assertRaisesRegex(RuntimeError, "failed"):
+            run.step(1)
+
+    def test_launch_failure_fails_instance_and_keeps_step_count(self):
+        for kernel, objects in (("BlockHaloCopy", ()), ("BlockCurvedBoundary", (circle(),)),
+                                ("BlockStep", ())):
+            with self.subTest(kernel=kernel):
+                run, log, host = backend(objects, failing_kernel=kernel)
+                run.step(1)
+                host.armed = True
+                log.clear()
+                with self.assertRaisesRegex(RuntimeError, kernel + " launch failed"):
+                    run.step(2)
+                self.assertEqual(run.step_count, 1)
+                self.assertNotIn("sync", [name for name, _ in log])
+                with self.assertRaisesRegex(RuntimeError, "failed"):
+                    run.download()
+
+    def test_free_memory_preflight_before_stream_and_buffers(self):
+        for cfg, options in ((LBMConfig(nx=32, ny=24, viscosity=.1),
+                              dict(refinements=[], level_blocks=[(12, 16)])),
+                             (LBMConfig(nx=48, ny=48, viscosity=.02),  # every level counts
+                              dict(refinements=[[(6, 6, 26, 26)]], level_blocks=[(4, 4)] * 2))):
+            with self.subTest(depth=len(options["refinements"]) + 1):
+                host = host_module()
+                host.cuda.mem_get_info.return_value = (1024, 1 << 30)
+                with self.assertRaisesRegex(MemoryError, "exceed free device memory"):
+                    host.FixedPatchStepper(cfg, context=Mock(activate=nullcontext, sm_arch="sm_70"),
+                                           **options)
+                host.cuda.Stream.assert_not_called()
+                host.cuda.mem_alloc.assert_not_called()
+                host.compiler.build_module.assert_not_called()
+
+    def test_wait_failure_after_a_good_step_keeps_its_count(self):
+        run, _, _ = backend()
+        run.step(1)
+        run._stream.synchronize.side_effect = RuntimeError("device fault")
+        with self.assertRaises(RuntimeError):
+            run.step(2)
+        self.assertEqual((run.step_count, run._status), (1, "failed"))
+
+    def test_final_wait_failure_fails_instance_and_keeps_step_count(self):
+        run, _, _ = backend()
+        run._stream.synchronize.side_effect = RuntimeError("device error")
+        with self.assertRaisesRegex(RuntimeError, "device error"):
+            run.step(3)
+        self.assertEqual(run.step_count, 0)
+        with self.assertRaisesRegex(RuntimeError, "failed"):
+            run.step(1)
+
+    def test_step_n_waits_once_and_matches_n_single_steps(self):
+        many, log_many, _ = backend()
+        many.step(3)
+        single, log_single, _ = backend()
+        for _ in range(3):
+            single.step(1)
+        names = [name for name, _ in log_many]
+        self.assertEqual(names.count("sync"), 1)
+        self.assertEqual(names[-1], "sync")
+        self.assertEqual(names.count("BlockStep"), 3)
+
+        def roles(run, log):  # same launch order and buffer roles as three step(1) calls
+            level = run._levels[0]
+            tags = {id(level.current): "c", id(level.spare): "s"}
+            return [(n, tuple(tags.get(id(x), "-") for x in a)) for n, a in log if n != "sync"]
+        self.assertEqual(roles(many, log_many), roles(single, log_single))
+        self.assertIsNot(many._levels[0].current, many._levels[0].spare)
+        self.assertEqual((many.step_count, single.step_count), (3, 3))
+
+    def test_step_zero_launches_and_waits_for_nothing(self):
+        run, log, _ = backend()
+        log.clear()
+        run.step(0)
+        self.assertEqual((log, run.step_count), ([], 0))
+
+    def test_launch_failure_mid_batch_fails_without_waiting(self):
+        run, log, host = backend()
+        step = run._levels[0].kernel
+        calls = []
+
+        def failing(*args, **kwargs):
+            calls.append(1)
+            if len(calls) == 2:
+                raise RuntimeError("launch failure")
+            return step(*args, **kwargs)
+
+        run._levels[0].kernel = failing
+        log.clear()
+        with self.assertRaises(RuntimeError):
+            run.step(3)
+        self.assertNotIn("sync", [name for name, _ in log])
+        self.assertEqual((run.step_count, run._status), (0, "failed"))
+        with self.assertRaises(RuntimeError):
+            run.download()
+
+    def test_initialize_rejects_bad_state_before_upload(self):
+        run, log, _ = backend(initialize=False)
+        good = state()
+        for bad in (good[:, :-1], np.where(np.arange(32) == 3, np.nan, good), -good):
+            with self.assertRaises(ValueError):
+                run.initialize(bad)
+        self.assertNotIn("upload", [name for name, _ in log])
+        run.initialize(good)
+        self.assertEqual([name for name, _ in log].count("upload"), 2)
+
+    def test_download_failure_fails_instance(self):
+        run, _, host = backend()
+        host.cuda.memcpy_dtoh.side_effect = RuntimeError("read failed")
+        with self.assertRaisesRegex(RuntimeError, "read failed"):
+            run.download()
+        for operation in (run.download, run.step):
+            with self.assertRaisesRegex(RuntimeError, "failed"):
+                operation()
+
+    def test_activation_failure_retains_resources_for_retry(self):
+        run, _, _ = backend()
+        allocations, stream = run._allocations[:], run._stream
+        unavailable = MagicMock()
+        unavailable.__enter__.side_effect = RuntimeError("activation failed")
+        run.ctx.activate = Mock(side_effect=[unavailable, nullcontext()])
+        with self.assertRaisesRegex(RuntimeError, "activation failed"):
+            run.close()
+        self.assertEqual(run._status, "failed")
+        self.assertEqual(run._allocations, allocations)
+        self.assertIs(run._stream, stream)
+        run.close()
+        self.assertEqual(run._status, "closed")
+        self.assertEqual(run._levels, [])
+        for allocation in allocations:
+            allocation.free.assert_called_once()
+
+
+if __name__ == "__main__":
+    unittest.main()
@@ -31,7 +31,7 @@ def backend(window=False):
    module.get_global.return_value = (Mock(), 20)
    module.get_function.side_effect = lambda name: Mock(
        side_effect=lambda *a, **k: log.append((name, a, k)))
-    host.compiler.load_module.return_value = module
+    host.compiler.build_module.return_value = module
    host.cuda.mem_alloc.side_effect = allocate
    host.cuda.memcpy_dtoh.side_effect = lambda array, pointer: array.fill(0)
    host.cuda.memset_d8_async.side_effect = lambda *a: log.append(("memset", a, {}))
@@ -12,13 +12,15 @@ from source_provenance import checkout_provenance
 import re100_helpers
 import run_jet_circle_cv

+sys.path.insert(0, str(ROOT / "tests/unit"))
+from historical_tree import historical_tree, tracked
+
 ORIGIN = "47b05bc8abfc957fe876b0c63bdb386e3c6cbaa5"


-def test_manifests_cover_current_source_and_preserve_origin():
-    names = subprocess.check_output(
-        ["git", "-C", str(ROOT), "ls-files", "src"], text=True
-    ).splitlines()
+def test_manifests_cover_current_source_and_preserve_origin(tmp_path):
+    assert re100_helpers.PIN == run_jet_circle_cv.PIN == ORIGIN
+    manifests = {}
    for group, hashes_key, revision_key in (
        ("dfg_re100", "production_sha256", "commit"),
        ("jet_circle_cv", "files", "revision"),
@@ -29,18 +31,22 @@ def test_manifests_cover_current_source_and_preserve_origin():
            ).read_text()
        )
        assert manifest[revision_key] == ORIGIN
-        assert set(manifest[hashes_key]) == set(names)
        assert manifest["source_origin_commits"][-1] == ORIGIN
        assert (
            manifest["historical_checkpoints"]["production"]
            == "d79e3f6f0a424f5bcc3a1a6d610313e49daf97a2"
        )
        assert "not inherited from d79" in manifest["experiment"]
-        for name, expected in manifest[hashes_key].items():
+        manifests[group] = manifest[hashes_key]
+    # History-dependent checks last, so a missing commit masks nothing above.
+    names = tracked(ORIGIN)
+    origin = historical_tree(ORIGIN, tmp_path)
+    for group, hashes in manifests.items():
+        assert set(hashes) == set(names), group
+        for name, expected in hashes.items():
            assert (
-                hashlib.sha256((ROOT / name).read_bytes()).hexdigest() == expected
+                hashlib.sha256((origin / name).read_bytes()).hexdigest() == expected
            ), name
-    assert re100_helpers.PIN == run_jet_circle_cv.PIN == ORIGIN


 def test_runtime_checkout_identity_is_separate_from_source_origin(tmp_path):
@@ -30,7 +30,7 @@ original_module = gate.compile_module
 cache = {}
 def retain_compile(*args, **kwargs):
    result = original_compile(*args, **kwargs)
-    source = kwargs.get('source', 'kernel_v2.cu')
+    source = kwargs.get('source', 'uniform.cu')
    shutil.copyfile(result, a.output / (Path(source).stem + '.ptx'))
    metadata = Path(result + '.json')
    if metadata.exists(): shutil.copyfile(metadata, a.output / (Path(source).stem + '.ptx.json'))
@@ -1,4 +1,4 @@
-"""Periodic body-free block backend: CPU contract doubles plus one short GPU case.
+"""Single blocked level, FixedPatchStepper(refinements=[]): CPU contracts plus short GPU cases.

 Frozen GPU gates (32x24, 8 steps, positive low-Mach non-equilibrium state):
 finite cores; bit-identical 1x1, 2x2 and 4x2 splits; max |G - FP64 oracle|
@@ -11,9 +11,8 @@ Mass/momentum drift is reported, not gated.
 from contextlib import nullcontext
 from pathlib import Path
 import resource
-from threading import RLock
 import unittest
-from unittest.mock import MagicMock, Mock, patch
+from unittest.mock import Mock, patch

 import numpy as np
 import pycuda.driver as cuda
@@ -21,10 +20,10 @@ import pycuda.driver as cuda
 from CelerisLab.body.geometry.circle import CircleGeometry
 from CelerisLab.body.objects import SimObject
 from CelerisLab.config import LBMConfig
-from CelerisLab.cuda import compiler_v2 as compiler
-from CelerisLab.lbm import block_stepper as implementation
-from CelerisLab.lbm.blocks import NO_DONOR, BlockLayout2D
-from CelerisLab.lbm.block_stepper import PeriodicBlockStepper
+from CelerisLab.cuda import compiler
+from CelerisLab.lbm import refinement_stepper as implementation
+from CelerisLab.lbm.blocks import NO_DONOR, BlockLayout2D, static_circle_tables
+from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
 from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY

 ROOT = Path(__file__).resolve().parents[3]
@@ -112,72 +111,78 @@ def totals(g):
    return g.sum(), (CX[:, None, None] * g).sum(), (CY[:, None, None] * g).sum()


+def single(cfg=None, *, context, block=(12, 16), objects=(), **kw):
+    """One periodic blocked level (the former PeriodicBlockStepper capability)."""
+    return FixedPatchStepper(cfg or config(), context=context, refinements=[], level_blocks=[block],
+                             level_objects=[tuple(objects)], **kw)
+
+
 class CpuContracts(unittest.TestCase):
-    """No device: driver doubles prove rejection before allocation and lifecycle."""
+    """No device: driver doubles prove rejection before allocation."""

    def forbidden_context(self):
        return Mock(activate=Mock(side_effect=AssertionError("CUDA reached")), sm_arch="sm_70")

    def test_unsupported_inputs_rejected_before_cuda(self):
-        layout = BlockLayout2D.build(NY, NX, 12, 16)
-        cases = ((config(), BlockLayout2D.build(NY, NX, 12, 16, periodic=(True, False))),
-                 (config(), BlockLayout2D.build(NY, NX * 2, 12, 16)),
-                 (config(collision="MRT"), layout),
-                 (config(), "layout"))
-        for cfg, lay in cases:
-            with self.subTest(cfg=cfg.collision, layout=getattr(lay, "periodic", lay)):
-                with self.assertRaises(ValueError):
-                    PeriodicBlockStepper(cfg, lay, context=self.forbidden_context())
+        for cfg, kw in ((None, dict(block=None)),
+                        (None, dict(coarse_boundary="channel")),
+                        (None, dict(coarse_boundary="slot_reservoir")),
+                        (config(collision="MRT"), {}),
+                        (config(les_enabled=True), {}),
+                        (None, dict(explosion="projected_linear")),
+                        (None, dict(observe_raw_exchange=True)),
+                        (None, dict(observe_raw_exchange=True, observe_window_exchange=True)),
+                        (None, dict(level_jets=[(object(),)]))):
+            with self.subTest(kw=kw), self.assertRaises(ValueError):
+                single(cfg, context=self.forbidden_context(), **kw)
        moving = circle()
        moving.state.vx = .01
        build_links = CircleGeometry.build_cut_links
        for objects in ([moving], [circle((1.5, 11.5), 2.0)], [circle(), circle()]):
            with self.assertRaises(ValueError):
-                PeriodicBlockStepper(config(), layout, context=self.forbidden_context(), objects=objects)
+                single(context=self.forbidden_context(), objects=objects)
        for center, radius in (((16.0, 12.0), -.1), ((16.0, 12.0), -.5), ((15.5, 11.5), np.nan),
                               ((np.inf, 11.5), 4.3), ((15.5, np.nan), 4.3)):
            with self.subTest(center=center, radius=radius), self.assertRaises(ValueError):
-                PeriodicBlockStepper(config(), layout, context=self.forbidden_context(),
-                                     objects=[circle(center, radius)])
-        with patch.object(implementation, "_static_circles") as prepare:
+                single(context=self.forbidden_context(), objects=[circle(center, radius)])
+        with patch.object(implementation, "static_circle_tables") as prepare:
            with self.assertRaisesRegex(AssertionError, "CUDA reached"):  # body-free path
-                PeriodicBlockStepper(config(), layout, context=self.forbidden_context(), objects=iter(()))
+                single(context=self.forbidden_context(), objects=iter(()))
            prepare.assert_not_called()
+        with self.assertRaisesRegex(AssertionError, "CUDA reached"):
+            single(context=self.forbidden_context(), objects=[circle()])
        with patch.object(CircleGeometry, "build_cut_links",
                          lambda geometry, *a, **k: build_links(geometry, *a, **k)[1:]):
            with self.assertRaisesRegex(ValueError, "cover every fluid pull"):
-                PeriodicBlockStepper(config(), layout, context=self.forbidden_context(),
-                                     objects=[circle()])
+                single(context=self.forbidden_context(), objects=[circle()])
        with patch.object(implementation, "MAX_GRID_DIM", 3):
            with self.assertRaisesRegex(ValueError, "launch limits"):
-                PeriodicBlockStepper(config(), layout, context=self.forbidden_context())
+                single(context=self.forbidden_context())

    def test_memory_and_param_layout_rejected_before_allocation(self):
-        layout = BlockLayout2D.build(NY, NX, 12, 16)
        context = Mock(activate=nullcontext, sm_arch="sm_70")
        with patch.object(cuda, "mem_get_info", return_value=(1024, 1 << 30)), \
                patch.object(cuda, "Stream") as stream, patch.object(cuda, "mem_alloc") as alloc:
-            with self.assertRaises(MemoryError):
-                PeriodicBlockStepper(config(), layout, context=context)
+            with self.assertRaisesRegex(MemoryError, "exceed free device memory"):
+                single(context=context, objects=[circle()])
            stream.assert_not_called()
            alloc.assert_not_called()
        module = Mock()
        module.get_global.return_value = (Mock(), 16)
-        with patch.object(cuda, "mem_get_info", return_value=(1 << 30, 1 << 30)), \
-                patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
+        with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), patch.object(compiler, "generate_config"), \
                patch.object(compiler, "compile_kernel") as compile_kernel, \
                patch.object(compiler, "cleanup_build_dir"), \
                patch.object(compiler, "load_module", return_value=module), \
                patch.object(cuda, "mem_alloc") as alloc, patch.object(cuda, "memcpy_htod") as upload:
            with self.assertRaisesRegex(RuntimeError, "layout mismatch"):
-                PeriodicBlockStepper(config(), layout, context=context)
-            self.assertEqual(compile_kernel.call_args.kwargs["source"], "blocks.cu")
+                single(context=context)
+            self.assertEqual(compile_kernel.call_args.kwargs["source"], "blocked.cu")
            alloc.assert_not_called()
            upload.assert_not_called()

    def test_static_circle_tables(self):
        layout = BlockLayout2D.build(NY, NX, 12, 16)
-        tables = implementation._static_circles(layout, (circle(),))
+        tables = static_circle_tables(layout, (circle(),))
        links = CircleGeometry(*CENTER, RADIUS).build_cut_links(NX, NY, domain_flags=dense_flags())
        self.assertEqual(len(tables["cell"]), len(links))
        np.testing.assert_array_equal(tables["q"], np.array([l.q for l in links], np.float32))
@@ -185,84 +190,11 @@ class CpuContracts(unittest.TestCase):
        self.assertTrue((tables["flags"][solid_cells] & 2).all())
        self.assertEqual(int(np.count_nonzero(tables["flags"] & 2)), int(tables["solid"].sum()))

-    def fake_backend(self):
-        backend = PeriodicBlockStepper.__new__(PeriodicBlockStepper)
-        backend.layout = BlockLayout2D.build(NY, NX, 12, 16)
-        backend.ctx = Mock(activate=nullcontext)
-        backend._lock = RLock()
-        backend._status = "ready"
-        backend._step_count = 0
-        backend._stream = Mock()
-        backend._module = Mock()
-        backend._copy = backend._step = Mock()
-        backend.boundary, backend._curved, backend._tables, backend._links = None, None, {}, 0
-        backend._dst, backend._src = Mock(), Mock()
-        backend._block, backend._grid, backend._copy_grid = (32, 1, 1), (1, 1, 1), (1, 1)
-        backend._allocations = [Mock(), Mock()]
-        backend._current, backend._spare = backend._allocations
-        return backend

-    def test_initialize_rejects_bad_state_before_upload(self):
-        backend = self.fake_backend()
-        good = initial_state()
-        bad = (good[:, :-1], np.where(np.arange(NX) == 3, np.nan, good), -good)
-        with patch.object(cuda, "memcpy_htod") as upload:
-            for g in bad:
-                with self.assertRaises(ValueError):
-                    backend.initialize(g)
-            upload.assert_not_called()
-            backend.initialize(good)
-            self.assertEqual(upload.call_count, 2)
-            self.assertTrue(np.isnan(upload.call_args_list[1].args[1]).all())
-        backend.close()
-
-    def test_failed_launch_and_download_make_backend_unusable(self):
-        for failure in ("step", "download"):
-            backend = self.fake_backend()
-            if failure == "step":
-                backend._step = Mock(side_effect=RuntimeError("launch failed"))
-                with self.assertRaisesRegex(RuntimeError, "launch failed"):
-                    backend.step()
-            else:
-                with patch.object(cuda, "memcpy_dtoh", side_effect=RuntimeError("read failed")):
-                    with self.assertRaisesRegex(RuntimeError, "read failed"):
-                        backend.download()
-            for operation in (backend.download, backend.step):
-                with self.assertRaisesRegex(RuntimeError, "failed"):
-                    operation()
-            backend.close()
-            for allocation in backend_allocations(backend):
-                allocation.free.assert_called_once()
-
-    def test_activation_failure_retains_resources_for_retry(self):
-        backend = self.fake_backend()
-        allocations, stream, module = backend._allocations[:], backend._stream, backend._module
-        unavailable = MagicMock()
-        unavailable.__enter__.side_effect = RuntimeError("activation failed")
-        backend.ctx.activate = Mock(side_effect=[unavailable, nullcontext()])
-        with self.assertRaisesRegex(RuntimeError, "activation failed"):
-            backend.close()
-        self.assertEqual(backend._status, "failed")
-        self.assertEqual(backend._allocations, allocations)
-        self.assertIs(backend._stream, stream)
-        self.assertIs(backend._module, module)
-        backend.close()
-        self.assertEqual(backend._status, "closed")
-        self.assertIsNone(backend._module)
-        self.assertIsNone(backend._copy)
-        self.assertIsNone(backend._step)
-        for allocation in allocations:
-            allocation.free.assert_called_once()
-
-
-def backend_allocations(backend):
-    return (backend._current, backend._spare)
-
-
-class PeriodicBlockGPU(unittest.TestCase):
+class SingleLevelGPU(unittest.TestCase):
    def setUp(self):
        self.assertEqual(Path(implementation.__file__).resolve(),
-                         ROOT / "src/CelerisLab/lbm/block_stepper.py")
+                         ROOT / "src/CelerisLab/lbm/refinement_stepper.py")

    def tearDown(self):
        rss = resource.getrusage(resource.RUSAGE_SELF).ru_maxrss / 1024
@@ -273,13 +205,13 @@ class PeriodicBlockGPU(unittest.TestCase):
        g0 = initial_state()
        results = []
        with CudaContext(0) as context:
-            for by, bx in SPLITS:
-                layout = BlockLayout2D.build(NY, NX, by, bx)
-                with PeriodicBlockStepper(config(), layout, context=context) as backend:
+            for split in SPLITS:
+                with single(context=context, block=split) as backend:
                    backend.initialize(g0)
                    backend.step(STEPS)
                    self.assertEqual(backend.step_count, STEPS)
-                    results.append(backend.download())
+                    (g,) = backend.download()
+                    results.append(g)
        for (by, bx), g in zip(SPLITS, results):
            self.assertTrue(np.isfinite(g).all(), (by, bx))
            np.testing.assert_array_equal(g, results[0], err_msg=f"split {(by, bx)}")
@@ -306,16 +238,17 @@ class PeriodicBlockGPU(unittest.TestCase):
        g0 = initial_state()
        results = []
        with CudaContext(0) as context:
-            for by, bx in SPLITS:
-                layout = BlockLayout2D.build(NY, NX, by, bx)
-                with PeriodicBlockStepper(config(), layout, context=context, objects=[circle()]) as backend:
+            for split in SPLITS:
+                with single(context=context, block=split, objects=[circle()]) as backend:
+                    tables = backend._level_boundaries[0].tables
                    for key, values, dtype in (("q", [l.q for l in links], np.float32),
                                               ("dir", [l.dir for l in links], np.uint8),
                                               ("fallback", [l.fallback for l in links], np.uint8)):
-                        np.testing.assert_array_equal(backend.boundary[key], np.array(values, dtype))
+                        np.testing.assert_array_equal(tables[key], np.array(values, dtype))
                    backend.initialize(g0)
                    backend.step(STEPS)
-                    results.append(backend.download())
+                    (g,) = backend.download()
+                    results.append(g)
        for (by, bx), g in zip(SPLITS, results):
            self.assertTrue(np.isfinite(g[:, ~solid]).all(), (by, bx))
            self.assertTrue(np.isnan(g[:, solid]).all(), (by, bx))
@@ -11,7 +11,7 @@ import numpy as np
 import pycuda.driver as cuda

 from CelerisLab.config import LBMConfig
-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.cuda.context import CudaContext
 from CelerisLab.lbm.blocks import BlockLayout2D
 from CelerisLab.lbm.descriptors import FLUID, SOLID, BC_WALL, BC_INLET, BC_OUTLET
@@ -27,7 +27,7 @@ def state(ny, nx):

 def compile_module(cfg, context, source):
    cfg.validate()
-    build = compiler.generate_config(cfg)
+    build = compiler.generate_config(cfg, block_root_bc=2 if source == "blocked.cu" else 0)
    try:
        module = compiler.load_module(compiler.compile_kernel(context.sm_arch, source=source, build_dir=build))
    finally:
@@ -46,7 +46,7 @@ def run_root(cfg, context, initial, block=None, steps=8):
    try:
        layout = None if block is None else BlockLayout2D.build(cfg.ny, cfg.nx, *block)
        local = cfg if layout is None else replace(cfg, nx=layout.local_shape[1], ny=layout.local_shape[0])
-        module = compile_module(local, context, "kernel_v2.cu" if layout is None else "blocks.cu")
+        module = compile_module(local, context, "uniform.cu" if layout is None else "blocked.cu")
        data = initial if layout is None else layout.to_blocked(initial)
        current, spare = upload(data), upload(np.full_like(data, np.nan))
        tpb = cfg.threads_per_block
@@ -18,7 +18,7 @@ import CelerisLab
 from CelerisLab.body.geometry.circle import CircleGeometry
 from CelerisLab.body.objects import SimObject
 from CelerisLab.config import LBMConfig
-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.cuda.context import CudaContext
 from CelerisLab.lbm import refinement_stepper as implementation
 from CelerisLab.lbm.descriptors import FLUID, SOLID
@@ -282,7 +282,8 @@ def test_fine_surface_four_hops(speed):
@pytest.mark.parametrize("missing", ["surface_velocity_gpu", "link_exchange_gpu"])
 def test_missing_surface_device_buffer_rejects_before_launch(missing):
    backend = implementation.FixedPatchStepper.__new__(implementation.FixedPatchStepper)
-    backend.fine_boundary = SimpleNamespace(objects=[object_spec(.002)])
+    backend._boundary_hosts = [None, SimpleNamespace(objects=[object_spec(.002)])]
+    backend._level_jets = [(), ()]
    backend._fine_curved = SimpleNamespace(fluid_idx_gpu=object(),
                                          surface_velocity_gpu=object(), link_exchange_gpu=object())
    setattr(backend._fine_curved, missing, None)
@@ -224,14 +224,12 @@ class RawExchangeGPU(unittest.TestCase):
        Eight local hops are a short-time equivalence check, not convergence.
        The test adapter only enables the standalone kernel's optional inputs.
        """
-        from CelerisLab.lbm.blocks import BlockLayout2D
-        from CelerisLab.lbm.block_stepper import PeriodicBlockStepper
        fine_cfg = replace(CFG, nx=64, ny=48, viscosity=.2)
        obj = SimObject(0, geometry=CircleGeometry(31.5, 23.5, 4.3),
                        center=(31.5, 23.5), radius=4.3)
-        layout = BlockLayout2D.build(48, 64, 8, 8)
        with CudaContext(0) as ctx:
-            with PeriodicBlockStepper(fine_cfg, layout, context=ctx, objects=[obj]) as uniform, \
+            with FixedPatchStepper(fine_cfg, context=ctx, refinements=[], level_blocks=[(8, 8)],
+                                   level_objects=[(obj,)]) as uniform, \
                 FixedPatchStepper(CFG, context=ctx, **fine((16, 20)), observe_raw_exchange=True) as amr:
                amr.initialize(*(state(l.shape) for l in amr._levels))
                uniform.initialize(state((9, 48, 64)))
@@ -240,18 +238,21 @@ class RawExchangeGPU(unittest.TestCase):
                links = obj.geometry.build_cut_links(64, 48)
                np.testing.assert_array_equal([v.dir for v in links], b.tables['dir'])
                np.testing.assert_allclose([v.q for v in links], b.tables['q'], atol=2e-7, rtol=0)
-                # Same fine-grid force arms after translation to the same center.
-                rx = uniform._upload(b.tables['rx'])
-                ry = uniform._upload(b.tables['ry'])
+                # Same fine-grid force arms after translation to the same center: the
+                # single level's static BFL gets the AMR Jet column and force arms; the raw
+                # output pointer is substituted at launch (no stepper-side accumulation).
+                root = uniform._level_boundaries[0]
                velocity = uniform._upload(b.tables['surface_velocity'])
+                root.device.update(surface_velocity=velocity, rx=uniform._upload(b.tables['rx']),
+                                   ry=uniform._upload(b.tables['ry']))
                raw = uniform._allocate(count*16)
-                original = uniform._curved
+                original = root.kernel
                samples = []
                def measured_boundary(*args, **kw):
-                    original(*args[:7], velocity, rx, ry, raw, **kw)
+                    original(*args[:10], raw, **kw)  # BlockCurvedBoundary raw-exchange output
                    uniform._stream.synchronize()
                    samples.append(download(raw, (count, 4)))
-                uniform._curved = measured_boundary
+                root.kernel = measured_boundary
                for speed in (0., .002, -.001, 0.):
                    amr.set_jet_amplitude(0, speed)
                    cuda.memcpy_htod(velocity, b.tables['surface_velocity'])
@@ -2,6 +2,37 @@

 Run CPU and GPU classes in separate processes. GPU callers must first check
 resources/device availability; this file never selects or reserves a GPU.
+
+Dense/blocked route parity (DenseBlockedRouteParity): the same configuration,
+initial state, body/Jet schedule and tilings on three storage routes, only
+level_blocks varying: D/D [None, None], D/B [None, FINE_BLOCK] (dense root,
+the mixed pattern of current examples) and B/B [(12, 8), FINE_BLOCK]. Gate:
+every public and owner array identical in FP32 bits and NaN masks at steps
+0, 1, 2, 8 and after re-initialize; D/B and B/B raw-exchange ledgers equal.
+No tolerance. A mismatch is reported with route pair/level/step and diagnosed
+(boundary semantics vs arithmetic ordering), never widened.
+
+Route capabilities (rejections checked before CUDA by test_route_capability_matrix):
+
+  combination                                   D/D       D/B       B/B
+  periodic / slot_reservoir root, fine circle   yes       yes       yes   (parity gate)
+  actuated circle, velocity Jet + amplitude     yes       yes       yes   (parity gate)
+  raw/window exchange on a body level           rejected  yes       yes
+  shifted storage                               rejected  rejected  yes
+  projected explosion                           rejected  rejected  yes
+  channel root                                  rejected  rejected  yes
+  several L1 rectangles                         rejected  rejected  yes
+  depth >= 3 with a dense level                 rejected  rejected  yes
+  MRT, periodic root, circle/actuation/Jet      yes       yes       yes   (parity gate)
+  two fine bodies (Jet + passive, two actuated) yes       yes       yes   (parity gate)
+  MRT with slot/reservoir root                  rejected on every route (on_site_velocity is SRT-only)
+Slot/reservoir walls are compile-time WALL_* macros on a dense root and the
+d_block_wall constant on a blocked root: BC-encoding parity, bit-identical by
+this gate. Dense fine bodies use CurvedLinkSoA + CurvedBoundaryKernel
+(n_objects > 0), blocked ones BlockCurvedBoundary; same links and physics,
+compared through the population rows. MRT body-free is also covered by
+test_refinement_mrt_gpu::test_dense_versus_blocked_mrt_two_level. Parked:
+shifted and depth >= 3 (no dense route to compare).
 """
 from contextlib import ExitStack, contextmanager, nullcontext
 from dataclasses import replace
@@ -13,13 +44,14 @@ from unittest.mock import Mock, patch
 import numpy as np
 import pycuda.driver as cuda

-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.lbm import refinement_stepper
 from CelerisLab.lbm.blocks import BlockLayout2D
 from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
 from CelerisLab.lbm.refinement import ACTIVE, OUTER_GHOST, FixedPatch2D
 from CelerisLab.lbm.refinement_stepper import FixedPatchStepper, blocked_base_tables
 from test_refinement_blocked_fine import NX, NY, BOUNDS, config, state, circle, actuation
+from test_refinement_jets import SPEED, jet

 BASE_BLOCKS = ((24, 32), (12, 8), (3, 4), (1, 1))
 FINE_BLOCK = (16, 20)
@@ -73,6 +105,7 @@ def mocked_backend(mode="periodic", fail_alloc=None):
            stack.enter_context(patch.object(compiler, name))
        stack.enter_context(patch.object(compiler, "load_module", return_value=module))
        stack.enter_context(patch.object(cuda, "Stream", return_value=stream))
+        stack.enter_context(patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)))
        stack.enter_context(patch.object(cuda, "mem_alloc", side_effect=allocate))
        stack.enter_context(patch.object(cuda, "memcpy_htod",
                                        side_effect=lambda p, a: uploads.append((p, a))))
@@ -223,6 +256,8 @@ class BlockedBaseCPU(unittest.TestCase):
                    self.assertIs(collide[4], coarse.kind)
                walls = [a for ptr, a in m["uploads"] if isinstance(ptr, str) and ptr == "d_block_wall"]
                self.assertEqual(len(walls), 1)
+                self.assertEqual([call.kwargs["block_root_bc"] for call in
+                                  compiler.generate_config.call_args_list], [1, 0])
                width, height, slot, speed = struct.unpack("=IIIf", walls[0])
                self.assertEqual((width, height, slot), (NX, NY, 8))
                self.assertAlmostEqual(speed, .001)
@@ -321,5 +356,167 @@ class BlockedBaseGPU(unittest.TestCase):
                            print(f"PASS {label}: full finite bits/NaN masks at 0,1,2,8 + reset", flush=True)


+ROUTES = (("D/D", [None, None]), ("D/B", [None, FINE_BLOCK]), ("B/B", [(12, 8), FINE_BLOCK]))
+# Two-body rows: patch (4, 4, 28, 20) gives a 40 x 56 fine window, tiled (8, 8).
+WIDE_BOUNDS = (4, 4, 28, 20)
+WIDE_ROUTES = (("D/D", [None, None]), ("D/B", [None, (8, 8)]), ("B/B", [(12, 8), (8, 8)]))
+LEFT, RIGHT, RADIUS = (12.0, 12.0), (20.0, 12.0), 2.15
+
+
+def body(ident, center, spec=None):
+    from CelerisLab.body.geometry.circle import CircleGeometry
+    from CelerisLab.body.objects import SimObject
+    return SimObject(ident, geometry=CircleGeometry(*center, RADIUS), center=center, radius=RADIUS,
+                     surface_actuation=spec)
+
+
+TWO_BODY_CASES = (
+    ("jet + passive", lambda: [body(0, LEFT), body(1, RIGHT)],
+     lambda: [actuation(SPEED).as_jet(LEFT, RADIUS)], -SPEED),
+    ("actuated +/-", lambda: [body(0, LEFT, actuation(SPEED)), body(1, RIGHT, actuation(-SPEED))],
+     lambda: [], None),
+)
+
+
+def same_summary(a, b):
+    """Exact equality of nested exchange summaries (dicts, sequences, arrays, numbers)."""
+    if isinstance(a, dict):
+        return a.keys() == b.keys() and all(same_summary(a[k], b[k]) for k in a)
+    if isinstance(a, (list, tuple)):
+        return len(a) == len(b) and all(same_summary(x, y) for x, y in zip(a, b))
+    a, b = np.asarray(a), np.asarray(b)
+    if a.dtype.kind in "biufc" and b.dtype.kind in "biufc":
+        return a.shape == b.shape and np.array_equal(a, b, equal_nan=True)
+    return a.dtype == b.dtype and np.array_equal(a, b)  # labels such as the window name
+
+
+class DenseBlockedRouteCPU(unittest.TestCase):
+    def test_same_summary_compares_labels_and_every_number(self):
+        ledger = {"window": "last_base_step", "time_start": 3, "total": {"force": [1.0, -2.0],
+                  "torque": 0.5}, "aperture": {"force": np.array([np.nan, 1.0])}}
+        self.assertTrue(same_summary(ledger, replace_value(ledger, ())))
+        for path, value in ((("window",), "last_substep"), (("time_start",), 4),
+                            (("total", "force"), [1.0, -2.000001]), (("total", "torque"), 0.25),
+                            (("aperture", "force"), np.array([0.0, 1.0]))):
+            with self.subTest(path=path):
+                self.assertFalse(same_summary(ledger, replace_value(ledger, path, value)))
+                self.assertFalse(same_summary(replace_value(ledger, path, value), ledger))
+
+    def test_route_capability_matrix(self):
+        context = Mock(activate=Mock(side_effect=AssertionError("CUDA reached")))
+        periodic = config("periodic")
+        inner = (BOUNDS[0] + 2, BOUNDS[1] + 2, BOUNDS[2] - 2, BOUNDS[3] - 2)
+        second = (BOUNDS[0], BOUNDS[1], BOUNDS[0] + 2, BOUNDS[1] + 2)
+        cases = (
+            ("raw exchange on dense body level", periodic, dict(
+                level_blocks=[None, None], level_objects=[(), (circle(),)], observe_raw_exchange=True)),
+            ("shifted D/D", replace(periodic, ddf_shifting=True), dict(level_blocks=[None, None])),
+            ("shifted D/B", replace(periodic, ddf_shifting=True), dict(level_blocks=[None, FINE_BLOCK])),
+            ("projected D/D", periodic, dict(level_blocks=[None, None], explosion="projected_linear")),
+            ("projected D/B", periodic, dict(level_blocks=[None, FINE_BLOCK], explosion="projected_linear")),
+            ("channel D/B", periodic, dict(level_blocks=[None, FINE_BLOCK], coarse_boundary="channel")),
+            ("several L1 rectangles", periodic, dict(refinements=[[BOUNDS, second]], level_blocks=[None, None])),
+            ("depth 3 dense", periodic, dict(refinements=[BOUNDS, inner], level_blocks=[None, None, None])),
+        )
+        for label, cfg, kw in cases:
+            kw.setdefault("refinements", [BOUNDS])
+            with self.subTest(label), self.assertRaises(ValueError):
+                FixedPatchStepper(cfg, context=context, **kw)
+        context.activate.assert_not_called()
+
+
+def replace_value(tree, path, value=None):
+    """Deep copy of a nested dict with the item at ``path`` replaced (no change for ())."""
+    if isinstance(tree, dict):
+        return {k: (value if path == (k,) else replace_value(v, path[1:], value) if path[:1] == (k,)
+                    else replace_value(v, (), None)) for k, v in tree.items()}
+    return np.array(tree, copy=True) if isinstance(tree, np.ndarray) else tree
+
+
+class DenseBlockedRouteParity(unittest.TestCase):
+    """D/D vs D/B vs B/B, unshifted SRT, one circle or one velocity Jet; see module docstring."""
+
+    def run_routes(self, mode, objects=(), jets=(), switch=None, cfg=None, bounds=BOUNDS, routes=ROUTES):
+        from CelerisLab.cuda.context import CudaContext
+        cfg = cfg or config(mode)
+        p = FixedPatch2D.from_config(cfg, bounds)
+        arrays = state(NY, NX, 1), state(*p.fine_shape, 2)
+        bodies = sorted({obj.obj_id for obj in objects})
+        results = {}
+        with CudaContext(0) as context:
+            for name, blocks in routes:
+                observe = bool(objects) and blocks[1] is not None
+                snapshots, ledgers = [], []
+                with FixedPatchStepper(cfg, context=context, coarse_boundary=mode, refinements=[bounds],
+                                       level_blocks=blocks, level_objects=[(), tuple(objects)],
+                                       level_jets=[(), tuple(jets)],
+                                       observe_raw_exchange=observe) as backend:
+                    backend.initialize(*arrays)
+                    for step in (0, 1, 2, 8):
+                        if step == 8 and switch is not None:
+                            backend.set_jet_amplitude(0, switch)
+                        backend.step(step - backend.step_count)
+                        snapshots.append(backend.download() + raw(backend))
+                        if observe and step:
+                            ledgers.append(tuple(backend.read_raw_exchange(body, window=window)
+                                                 for body in bodies
+                                                 for window in ("last_base_step", "last_substep")))
+                    backend.initialize(*arrays)
+                    snapshots.append(backend.download() + raw(backend))
+                results[name] = snapshots, ledgers
+        return results
+
+    def check(self, label, results):
+        reference, _ = results["D/D"]
+        for name in ("D/B", "B/B"):
+            for step, actual, expected in zip((0, 1, 2, 8, "reset"), results[name][0], reference):
+                for index, (a, b) in enumerate(zip(actual, expected)):
+                    exact(self, a, b, f"{label} {name} vs D/D step={step} array={index}")
+        if results["D/B"][1]:
+            for step, a, b in zip((1, 2, 8), results["D/B"][1], results["B/B"][1]):
+                self.assertTrue(same_summary(a, b), f"{label} raw exchange D/B vs B/B step={step}")
+        for name, (snapshots, _) in results.items():
+            masses = [float(np.nansum(level.astype(np.float64))) for level in snapshots[3][:2]]
+            print(f"{label} {name}: step 8 mass by level {masses}", flush=True)
+        print(f"PASS {label}: D/D = D/B = B/B bits/NaN masks at 0,1,2,8 + reset", flush=True)
+
+    def test_body_free(self):
+        for mode in ("periodic", "slot_reservoir"):
+            self.check(f"{mode} body-free", self.run_routes(mode))
+
+    def test_passive_circle(self):
+        for mode in ("periodic", "slot_reservoir"):
+            self.check(f"{mode} circle", self.run_routes(mode, objects=[circle()]))
+
+    def test_actuated_circle(self):
+        for mode in ("periodic", "slot_reservoir"):
+            for speed in (SPEED, -SPEED):
+                self.check(f"{mode} actuated {speed}",
+                           self.run_routes(mode, objects=[circle(actuation(speed))]))
+
+    def test_velocity_jet_with_amplitude_switch(self):
+        for mode in ("periodic", "slot_reservoir"):
+            self.check(f"{mode} jet switch",
+                       self.run_routes(mode, objects=[circle()], jets=[jet()], switch=-SPEED))
+
+    def test_mrt_bodies(self):
+        """MRT, periodic root (on_site_velocity slot/reservoir is SRT-only on every route)."""
+        cfg = replace(config("periodic"), collision="MRT")
+        for label, objects, jets, switch in (("circle", [circle()], (), None),
+                                             ("actuated +", [circle(actuation(SPEED))], (), None),
+                                             ("actuated -", [circle(actuation(-SPEED))], (), None),
+                                             ("jet switch", [circle()], [jet()], -SPEED)):
+            self.check(f"MRT periodic {label}",
+                       self.run_routes("periodic", objects, jets, switch, cfg=cfg))
+
+    def test_two_bodies(self):
+        """Two circles on a wider patch (same tiles on every route): Jet + passive, two actuated."""
+        for mode in ("periodic", "slot_reservoir"):
+            for label, objects, jets, switch in TWO_BODY_CASES:
+                self.check(f"{mode} two bodies {label}",
+                           self.run_routes(mode, objects(), jets(), switch, bounds=WIDE_BOUNDS,
+                                           routes=WIDE_ROUTES))
+
+
 if __name__ == "__main__":
    unittest.main()
@@ -15,7 +15,7 @@ import pycuda.driver as cuda
 from CelerisLab.body.geometry.circle import CircleGeometry
 from CelerisLab.body.objects import SimObject, SurfaceActuation
 from CelerisLab.config import LBMConfig
-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.lbm import curved_links
 from CelerisLab.lbm.blocks import NO_DONOR, BlockLayout2D, blocked_link_tables
 from CelerisLab.lbm.descriptors import D2Q9_EX, D2Q9_EY
@@ -122,7 +122,7 @@ class BlockedFineCPU(unittest.TestCase):
                raise AssertionError("zero-byte device allocation")
            sizes.append(size)
            return Mock()
-        with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
+        with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), patch.object(compiler, "generate_config"), \
                patch.object(compiler, "compile_kernel"), patch.object(compiler, "cleanup_build_dir"), \
                patch.object(compiler, "load_module", return_value=module), \
                patch.object(cuda, "mem_alloc", side_effect=alloc), patch.object(cuda, "memcpy_htod"):
@@ -172,7 +172,7 @@ class BlockedFineCPU(unittest.TestCase):
            return functions.setdefault(name, Mock(side_effect=lambda *a, **k: calls.append((name, a))))
        module = Mock(get_function=Mock(side_effect=get_function))
        module.get_global.return_value = (Mock(), 20)
-        with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
+        with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), patch.object(compiler, "generate_config"), \
                patch.object(compiler, "compile_kernel"), patch.object(compiler, "cleanup_build_dir"), \
                patch.object(compiler, "load_module", return_value=module), \
                patch.object(cuda, "mem_alloc", side_effect=lambda size: Mock()), \
@@ -108,7 +108,9 @@ class FaultGuardsCPU(unittest.TestCase):

    def fake_backend(self):
        backend = PeriodicPatchStepper.__new__(PeriodicPatchStepper)
-        backend.fine_layout = None  # dense fine storage
+        backend._layouts = [None, None]  # dense storage on both levels
+        backend._ddf_shifting = False  # unshifted on every level, as _check_storage expects
+        backend._level_cfgs = [LBMConfig(nx=16, ny=16, viscosity=.1)] * 2  # ddf_shifting off
        backend.ctx = Mock(activate=nullcontext)
        backend._lock = RLock()
        backend._status = "ready"
@@ -158,13 +160,24 @@ class FaultGuardsCPU(unittest.TestCase):
            with self.assertRaisesRegex(RuntimeError, "four-cell"):
                PeriodicPatchStepper(cfg, context=context, refinements=[(4, 4, 12, 12)])
        p = implementation.FixedPatch2D.from_config(cfg, (4, 4, 12, 12))
-        with patch.object(cuda, "Stream"), patch.object(implementation.FixedPatch2D, "from_config",
-                                                       return_value=replace(p, tau_coarse=.9)):
-            with self.assertRaisesRegex(RuntimeError, "omega/tau"):
+        from CelerisLab.lbm import hierarchy
+        bad = patch.object(implementation.FixedPatch2D, "from_config",
+                           return_value=replace(p, tau_coarse=.9))
+        # A pair tau inconsistent with its level config is rejected twice over, both before
+        # any device allocation or upload: first by the hierarchy chain check, and, if that
+        # is bypassed, by the per-level omega/tau check in _build_level.
+        with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), bad, patch.object(cuda, "mem_alloc") as alloc, \
+                patch.object(cuda, "memcpy_htod") as upload:
+            with self.assertRaisesRegex(RuntimeError, "config chain is inconsistent"):
                PeriodicPatchStepper(cfg, context=context, refinements=[p.bounds])
+            with patch.object(hierarchy, "_check_chain"), \
+                    self.assertRaisesRegex(RuntimeError, "omega/tau"):
+                PeriodicPatchStepper(cfg, context=context, refinements=[p.bounds])
+            alloc.assert_not_called()
+            upload.assert_not_called()
        module = Mock()
        module.get_global.return_value = (Mock(), 16)
-        with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
+        with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), patch.object(compiler, "generate_config"), \
                patch.object(compiler, "compile_kernel"), patch.object(compiler, "cleanup_build_dir"), \
                patch.object(compiler, "load_module", return_value=module), \
                patch.object(cuda, "memcpy_htod") as upload:
@@ -20,7 +20,7 @@ import pycuda.driver as cuda

 import CelerisLab
 from CelerisLab.config import LBMConfig
-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.cuda.context import CudaContext
 from CelerisLab.lbm.refinement_stepper import FixedPatchStepper, PeriodicPatchStepper

@@ -16,7 +16,7 @@ from unittest.mock import Mock, patch
 import numpy as np
 import pycuda.driver as cuda

-from CelerisLab.cuda import compiler_v2 as compiler
+from CelerisLab.cuda import compiler
 from CelerisLab.lbm.descriptors import SOLID
 from CelerisLab.lbm.refinement import ACTIVE, FixedPatch2D, prepare_fine_boundary
 from CelerisLab.lbm.refinement_stepper import FixedPatchStepper
@@ -36,7 +36,7 @@ def mocked(objects, jets, fine_block):
    module = Mock(get_function=Mock(return_value=Mock()))
    module.get_global.return_value = (Mock(), 20)
    htod = Mock()
-    with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
+    with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), patch.object(compiler, "generate_config"), \
            patch.object(compiler, "compile_kernel"), patch.object(compiler, "cleanup_build_dir"), \
            patch.object(compiler, "load_module", return_value=module), \
            patch.object(cuda, "mem_alloc", side_effect=lambda size: Mock()), \
@@ -52,11 +52,11 @@ def child_jet(speed=.002):


 def mocked_child():
-    from CelerisLab.cuda import compiler_v2 as compiler
+    from CelerisLab.cuda import compiler
    context = Mock(activate=nullcontext, sm_arch="sm_70")
    module = Mock(get_function=Mock(side_effect=lambda name: Mock(name=name)))
    module.get_global.return_value = (Mock(), 20)
-    with patch.object(cuda, "Stream"), patch.object(compiler, "generate_config"), \
+    with patch.object(cuda, "Stream"), patch.object(cuda, "mem_get_info", return_value=(1 << 40, 1 << 40)), patch.object(compiler, "generate_config"), \
            patch.object(compiler, "compile_kernel"), patch.object(compiler, "cleanup_build_dir"), \
            patch.object(compiler, "load_module", return_value=module), \
            patch.object(cuda, "mem_alloc", side_effect=lambda size: Mock()), \
@@ -281,16 +281,21 @@ def host_geometry(source, h, collision="SRT", jet_scenario="none", placement="wi
    return cfg, pair, body, host


-def verify_source(source, collision="SRT", observation_mode="last_base_step", jet_scenario="none", observe_cv=False):
+def source_manifest(collision="SRT", observation_mode="last_base_step"):
+    """The frozen historical manifest for this collision/observation mode (no checkout read)."""
    if collision not in ("SRT", "MRT"):
        raise ValueError("Matched collision must be SRT or MRT")
    if observation_mode not in ("last_base_step", "window"):
        raise ValueError("Unknown refined observation mode")
-    if observe_cv and jet_scenario == "none":
-        raise ValueError("--observe-cv requires an explicit public Jet scenario")
    name = ("source_manifest_window.json" if observation_mode == "window" else
            "source_manifest.json" if collision == "SRT" else "source_manifest_mrt.json")
-    expected = json.loads((HERE / "data/refined_dfg" / name).read_text())
+    return json.loads((HERE / "data/refined_dfg" / name).read_text())
+
+
+def verify_source(source, collision="SRT", observation_mode="last_base_step", jet_scenario="none", observe_cv=False):
+    expected = source_manifest(collision, observation_mode)
+    if observe_cv and jet_scenario == "none":
+        raise ValueError("--observe-cv requires an explicit public Jet scenario")
    if jet_scenario != "none":
        if collision != "MRT" or observation_mode != "window":
            raise ValueError("Public Jet scenarios require MRT and --exchange-window")
@@ -298,7 +303,8 @@ def verify_source(source, collision="SRT", observation_mode="last_base_step", je
        sys.path.insert(0, str(HERE))
        from source_provenance import checkout_provenance

-        paths = set(expected["sha256"])
+        # Harness files from the manifest; production files as this checkout lists them.
+        paths = {name for name in expected["sha256"] if not name.startswith("src/")}
        if observe_cv:
            paths.add("tests/validation/run_jet_circle_cv.py")
        paths.update(subprocess.check_output(
@@ -310,7 +316,8 @@ def verify_source(source, collision="SRT", observation_mode="last_base_step", je
                    for name in sorted(paths)},
        )
    for name, digest in expected["sha256"].items():
-        if hashlib.sha256((source / name).read_bytes()).hexdigest() != digest:
+        path = source / name
+        if not path.is_file() or hashlib.sha256(path.read_bytes()).hexdigest() != digest:
            raise RuntimeError("Source mismatch: " + name)
    return expected

@@ -1,6 +1,6 @@
 """Host geometry, receipt-corruption tests and mocked scheduling; no CUDA/flow."""

-from contextlib import nullcontext, redirect_stdout
+from contextlib import contextmanager, nullcontext, redirect_stdout
 from copy import deepcopy
 from dataclasses import asdict, replace
 import hashlib
@@ -24,6 +24,44 @@ from oracle import LIMITS, TIMES, U, rectangles
 REPO = Path(__file__).resolve().parents[3]
 OLD_REVISION = "519473a61c7b319183aa54e55b9f097499dde0bf"
 MOCK_REVISION = "f" * 40
+
+
+@contextmanager
+def historical_package():
+    """Import CelerisLab from NUMERICAL_REVISION, the only tree run_case accepts.
+
+    The harness and its pinned build-capture helper are tooling for that revision
+    (cuda/compiler_v2.py); the current package is restored afterwards.
+    """
+    import tarfile
+    archive = subprocess.run(["git", "-C", str(REPO), "archive", "--format=tar",
+                              NUMERICAL_REVISION, "src/CelerisLab"],
+                             capture_output=True, check=True).stdout
+    before = set(sys.modules)
+    saved = {k: v for k, v in sys.modules.items() if k == "CelerisLab" or k.startswith("CelerisLab.")}
+    with tempfile.TemporaryDirectory() as tmp:
+        root = Path(tmp).resolve()
+        with tarfile.open(fileobj=io.BytesIO(archive)) as tar:
+            tar.extractall(root)
+        for name in saved:
+            del sys.modules[name]
+        sys.path.insert(0, str(root / "src"))
+        try:
+            yield root
+        finally:
+            sys.path.remove(str(root / "src"))
+            for name in set(sys.modules) - before | set(saved):  # nothing keeps the old tree
+                sys.modules.pop(name, None)
+            sys.modules.update(saved)
+
+
+def historical_runtime_source(repo):
+    """runtime_source as run_case records it for an accepted a135252 checkout."""
+    package = repo / "src/CelerisLab"
+    return dict(checkout=str(repo), module=str(package / "__init__.py"),
+                compiler=str(package / "cuda/compiler_v2.py"), kernel_root=str(package / "lbm/kernels"))
+
+
 RUNTIME = dict(device=dict(name="host mock", pci_bus_id="mock", capability=[7, 0],
                           driver=12040, runtime=[12, 4, 0]), nvcc_path="/mock/nvcc",
               nvcc_version="mock nvcc", python="mock Python", numpy=np.__version__, pycuda="mock")
@@ -60,7 +98,7 @@ def records(early=False):
            reference_harness=run_case.REFERENCE_HARNESS,
            actual_build_sidecars=[build_metadata(n * 2**level) for level in range(depth)],
            runtime_identity=deepcopy(RUNTIME), cuda_visible_devices="0",
-            runtime_source=run_case.runtime_source(REPO), completed=True,
+            runtime_source=historical_runtime_source(REPO), completed=True,
            workload=run_case.workload([k == ACTIVE for k in h.kinds], 16*n),
            rows=[dict(step=s, tstar=s*U/n, **{k: 0.0 for k in LIMITS}) for s in main_steps],
            early_pressure=dict(enabled=early, rows=[
@@ -125,6 +163,10 @@ class ProjectedPairsTests(unittest.TestCase):
                                 frozen[name])

    def test_actual_build_capture_rejects_tampered_ptx_and_sidecar(self):
+        with historical_package():
+            self._check_build_capture()
+
+    def _check_build_capture(self):
        from CelerisLab.cuda import compiler_v2
        from tests.validation.multigrid.check_projected_pairs import capture_builds

@@ -249,35 +291,41 @@ class ProjectedPairsTests(unittest.TestCase):
        self.assertFalse(assess(incomplete)["comparisons"][3]["passed"])

    def test_pinned_build_sources_match_host_generated_inputs(self):
-        """The pinned fingerprints are what today's compiler generates from the qualified
-        NUMERICAL_REVISION kernels; the current kernels may legitimately differ."""
+        """Check the pinned fingerprints using NUMERICAL_REVISION's renderer and kernels.
+        The imported LBMConfig is unchanged from that revision; this is a host
+        build-input check, not a reconstruction of the historical CUDA toolchain."""
+        import importlib.util
        import tarfile
-        from CelerisLab.cuda import compiler_v2

        kernels = "src/CelerisLab/lbm/kernels"
+        renderer = "src/CelerisLab/cuda/compiler_v2.py"
        if subprocess.run(["git", "-C", str(REPO), "cat-file", "-e",
                           f"{NUMERICAL_REVISION}^{{commit}}"], capture_output=True).returncode:
            self.fail(f"{NUMERICAL_REVISION} is not in this clone; the pinned build "
                      "sources cannot be checked")
        archive = subprocess.run(["git", "-C", str(REPO), "archive", "--format=tar",
-                                  NUMERICAL_REVISION, kernels],
+                                  NUMERICAL_REVISION, kernels, renderer],
                                 capture_output=True, check=True).stdout
        with tempfile.TemporaryDirectory() as snapshot:
            with tarfile.open(fileobj=io.BytesIO(archive)) as tar:
                tar.extractall(snapshot)
-            with patch.object(compiler_v2, "_KERNEL_DIR", str(Path(snapshot) / kernels)):
+            spec = importlib.util.spec_from_file_location(
+                "CelerisLab.cuda._compiler_a135252", Path(snapshot) / renderer)
+            compiler = importlib.util.module_from_spec(spec)
+            spec.loader.exec_module(compiler)
+            with patch.object(compiler, "_KERNEL_DIR", str(Path(snapshot) / kernels)):
                for size in (128, 256):
                    with self.subTest(size=size):
                        cfg = replace(run_case.config(size, 100), nx=10, ny=10)
                        cfg.validate()
-                        root = compiler_v2.generate_config(cfg)
+                        root = compiler.generate_config(cfg)
                        try:
                            actual = json.loads((Path(root) / ".celeris-build").read_text())
                            expected = build_metadata(size)
                            self.assertEqual(actual, {key: expected[key] for key in
                                                      ("source_sha256", "config", "n_objects")})
                        finally:
-                            compiler_v2.cleanup_build_dir(root)
+                            compiler.cleanup_build_dir(root)

    def test_source_preflight_requires_clean_exact_revision_and_tree(self):
        import oracle
@@ -332,6 +380,7 @@ class ProjectedPairsTests(unittest.TestCase):
            with patch.object(sys, "argv", pinned + ["--dry-run"]), \
                 patch.dict(sys.modules, cuda_modules), patch.dict("os.environ", CUDA_VISIBLE_DEVICES=""), \
                 patch.object(run_case, "verify_source", return_value=source()) as verify, \
+                 patch.object(run_case, "runtime_source", return_value=historical_runtime_source(REPO)), \
                 patch.object(run_case, "runtime_identity", side_effect=AssertionError("CUDA reached")), \
                 redirect_stdout(stdout):
                self.assertEqual(run_case.main(), 0)
@@ -340,10 +389,12 @@ class ProjectedPairsTests(unittest.TestCase):
            self.assertFalse(output.exists())

    def test_runtime_rejects_installed_kernel_tree(self):
-        from CelerisLab.cuda import compiler_v2
-        with patch.object(compiler_v2, "_KERNEL_DIR", "/other/CelerisLab/lbm/kernels"):
-            with self.assertRaisesRegex(RuntimeError, "CUDA source/include root"):
-                run_case.runtime_source(REPO)
+        with historical_package() as repo:
+            from CelerisLab.cuda import compiler_v2
+            self.assertEqual(run_case.runtime_source(repo)["checkout"], str(repo))
+            with patch.object(compiler_v2, "_KERNEL_DIR", "/other/CelerisLab/lbm/kernels"):
+                with self.assertRaisesRegex(RuntimeError, "CUDA source/include root"):
+                    run_case.runtime_source(repo)

    def test_existing_output_is_never_overwritten(self):
        with tempfile.TemporaryDirectory() as tmp:
@@ -419,6 +470,7 @@ class ProjectedPairsTests(unittest.TestCase):
                with patch.dict(sys.modules, modules), patch.object(sys, "argv", argv), \
                     patch.dict("os.environ", CUDA_VISIBLE_DEVICES="mock"), \
                     patch.object(run_case, "verify_source", return_value=source()) as verify, \
+                     patch.object(run_case, "runtime_source", return_value=historical_runtime_source(REPO)), \
                     patch.object(run_case, "runtime_identity", return_value=deepcopy(RUNTIME)), \
                     patch.object(run_case, "build_capture", return_value=nullcontext()), \
                     patch.object(run_case, "metrics", side_effect=metric), redirect_stdout(io.StringIO()):
@@ -38,7 +38,7 @@ VARIANTS = ("A", "W", "W1")
 # links). Engineering tolerance: 1e-9 of the interval's summed absolute per-step totals.
 INTERVAL_RTOL = 1e-9
 CANDIDATE_FILES = ("src/CelerisLab/lbm/refinement_stepper.py", "src/CelerisLab/lbm/raw_exchange.py",
-                   "src/CelerisLab/lbm/refinement_tables.py", "src/CelerisLab/lbm/kernels/blocks.cu")
+                   "src/CelerisLab/lbm/refinement_tables.py", "src/CelerisLab/lbm/kernels/blocked.cu")


 def components(obs):