Compare commits
3
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 ClaudeandClaude Opus 5.5
|
0f8fdbc30e | ||
|
ClaudeandClaude Opus 5.5
|
e4538219b2 | ||
|
ClaudeandClaude Opus 5.5
|
888c80381c |
+27
-26
@@ -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` |
|
||||
| Uniform blocked backend `PeriodicBlockStepper` | doubly periodic, SRT only, stationary passive circles | `lbm/block_stepper.py` `PeriodicBlockStepper.__init__`, `_static_circles` |
|
||||
|
||||
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.
|
||||
|
||||
@@ -36,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]))) {
|
||||
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);
|
||||
}
|
||||
// Kernels: patch collision helper, PatchExplode, PatchFineStep, PatchCoarseStep.
|
||||
#include "step/patch_step.cu"
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
@@ -949,8 +949,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])
|
||||
|
||||
@@ -27,7 +27,7 @@ class RefinementMRTHost(unittest.TestCase):
|
||||
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()
|
||||
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);"):
|
||||
|
||||
Reference in New Issue
Block a user