ClaudeandClaude Opus 5.5
733e4d1fe7
Describes current module/global, table, control and per-step state ownership for Simulation, PeriodicBlockStepper and FixedPatchStepper, and the window vs level-wide frame of single vs multi-rectangle levels. No code change. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2.6 KiB
2.6 KiB
Device state lifecycle (uniform Simulation, PeriodicBlockStepper, 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 | PTX module, bound kernel handles; __constant__ / __device__ globals d_params, d_w, d_cx, d_cy, d_block_channel, d_block_wall |
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. |
Every owner loads its own module (compiler_v2.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 (
kernel_v2.cu,Simulation): whole-domainNX/NY. - Blocked levels (
blocks.cu;PeriodicBlockStepperand every blockedFixedPatchStepperlevel):NX/NYare 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).