475 lines
25 KiB
Diff
475 lines
25 KiB
Diff
--- a/src/CelerisLab/common/checkpoint.py
|
||
+++ b/src/CelerisLab/common/checkpoint.py
|
||
@@ -218,6 +218,8 @@
|
||
for key in ("y_wall_bc", "velocity", "wall_normal_velocity"):
|
||
if saved_cfg.get(key) != getattr(lbm_cfg, key):
|
||
raise ValueError(f"Checkpoint on-site boundary mismatch on '{key}'")
|
||
+ if saved_cfg.get("wall_slot_width", 0) != lbm_cfg.wall_slot_width:
|
||
+ raise ValueError("Checkpoint on-site boundary mismatch on 'wall_slot_width'")
|
||
for key in ("lattice_model", "nx", "ny", "nz", "store_precision",
|
||
"streaming", "ddf_shifting"):
|
||
saved_val = saved_cfg.get(key)
|
||
--- a/src/CelerisLab/config.py
|
||
+++ b/src/CelerisLab/config.py
|
||
@@ -87,6 +87,7 @@
|
||
outlet_srt_neq_damp: float = 0.5
|
||
y_wall_bc: str = "bounce_back"
|
||
wall_normal_velocity: float = 0.0
|
||
+ wall_slot_width: int = 0
|
||
omega_min: float = 0.01
|
||
omega_max: float = 1.99
|
||
|
||
@@ -151,6 +152,11 @@
|
||
)
|
||
if self.dim == 2 and self.nz != 1:
|
||
raise ValueError("nz must be 1 for 2D")
|
||
+ if (isinstance(self.wall_slot_width, bool) or not isinstance(self.wall_slot_width, int)
|
||
+ or self.wall_slot_width != 0 and not 2 <= self.wall_slot_width <= self.nx):
|
||
+ raise ValueError("wall_slot_width must be 0 (disabled) or an integer in [2, nx]")
|
||
+ if self.wall_slot_width and (self.y_wall_bc != "on_site_velocity" or self.velocity != 0):
|
||
+ raise ValueError("wall_slot_width requires on_site_velocity and velocity=0")
|
||
if self.y_wall_bc == "on_site_velocity":
|
||
if (not self.is_d2q9 or self.collision != "SRT"
|
||
or self.streaming != "double_buffer" or self.store_precision != "FP32"
|
||
@@ -196,8 +202,17 @@
|
||
self.validate()
|
||
return (self.lattice_model, self.nx, self.ny, self.nz, self.data_type,
|
||
self.store_precision, bool(self.ddf_shifting), self.streaming,
|
||
- (self.velocity, self.wall_normal_velocity)
|
||
+ (self.velocity, self.wall_normal_velocity, self.wall_slot_width)
|
||
if self.y_wall_bc == "on_site_velocity" else None)
|
||
+
|
||
+ @property
|
||
+ def top_wall_normal_velocity(self) -> float:
|
||
+ """Uniform top target balancing the bottom's discrete requested volume."""
|
||
+ if self.wall_slot_width in (0, self.nx):
|
||
+ return self.wall_normal_velocity
|
||
+ weights = (max(0.0, 1.0-((2*x-(self.nx-1))/self.wall_slot_width)**2)**2
|
||
+ for x in range(self.nx))
|
||
+ return self.wall_normal_velocity * math.fsum(weights) / self.nx
|
||
|
||
def to_macros(self) -> Dict[str, Any]:
|
||
"""Return flat dict of macro_name → value for header generation."""
|
||
@@ -227,6 +242,8 @@
|
||
"OUTLET_BACKFLOW_CLAMP": int(self.outlet_backflow_clamp),
|
||
"Y_WALL_BC": Y_WALL_BC_MAP[self.y_wall_bc],
|
||
"WALL_NORMAL_VELOCITY": _float_literal(self.wall_normal_velocity),
|
||
+ "WALL_SLOT_WIDTH": self.wall_slot_width,
|
||
+ "TOP_WALL_NORMAL_VELOCITY": _float_literal(self.top_wall_normal_velocity),
|
||
"OMEGA_COLLISION_MIN": _float_literal(self.omega_min),
|
||
"OMEGA_COLLISION_MAX": _float_literal(self.omega_max),
|
||
"TRT_MAGIC_PARAM": _float_literal(self.trt_magic_param),
|
||
@@ -296,6 +313,7 @@
|
||
outlet_srt_neq_damp=m["outlet"].get("srt_neq_damp", 0.5),
|
||
y_wall_bc=m.get("y_wall_bc", "bounce_back"),
|
||
wall_normal_velocity=m.get("wall_normal_velocity", 0.0),
|
||
+ wall_slot_width=m.get("wall_slot_width", 0),
|
||
omega_min=m["omega_guard"]["min"],
|
||
omega_max=m["omega_guard"]["max"],
|
||
threads_per_block=c["threads_per_block"],
|
||
--- a/src/CelerisLab/configs/CONFIG.md
|
||
+++ b/src/CelerisLab/configs/CONFIG.md
|
||
@@ -113,6 +113,33 @@
|
||
|
||
Supported only for D2Q9/SRT/FP32/unshifted/double_buffer without LES, objects
|
||
or body forces. Boundary targets and the mode are fixed at construction;
|
||
-checkpoint restart requires matching targets. This is an entire-plane
|
||
-Couette-throughflow reference, not a finite slot or curved jet API. The default
|
||
+checkpoint restart requires matching targets. With the slot disabled this is an
|
||
+entire-plane Couette-throughflow reference. The default
|
||
`bounce_back` mode and existing configurations retain their previous path.
|
||
+
|
||
+### Static centered bottom slot
|
||
+
|
||
+Optional `method.wall_slot_width` is an integer W: 0 (default) keeps the above
|
||
+uniform planes; 2≤W≤NX enables a centered bottom slot and requires
|
||
+`physics.velocity=0`. The same support restrictions apply. Width and amplitude
|
||
+are fixed at construction and must match on checkpoint restart; older v3
|
||
+checkpoints without this field mean W=0.
|
||
+
|
||
+For W<NX, let c=(NX−1)/2, s=2(x−c)/W at integer x=0..NX−1. The bottom target is
|
||
+`v=A*max(0,1-s*s)^2`, where A=`wall_normal_velocity` is the continuous peak,
|
||
+not a renormalized discrete peak. This profile and its first derivative vanish
|
||
+at the edges. W=NX explicitly means uniform v=A, not this shaped profile.
|
||
+With equal unit area per node, requested Q is the sum of these bottom targets;
|
||
+the top target is uniform Q/NX (`cfg.top_wall_normal_velocity`). The sum is
|
||
+computed in host double precision; device targets are FP32. This balances
|
||
+requested volume, not actual density-weighted mass exactly.
|
||
+
|
||
+Positive A injects from the bottom and exits through the top; negative A reverses
|
||
+both roles. Both are Cartesian y velocities; inward-normal velocity has the
|
||
+opposite sign at the top. The slot and its zero-velocity flanks share the same
|
||
+on-site y=1 plane. Initialization sets rho0, u=0 and linearly interpolates v
|
||
+between the two planes; it is not an assumed steady solution. This mode is
|
||
+static, single-level and flat; it does not support moving/curved slots.
|
||
+The finite-slot mode is experimental: the bounded W=16, A=±0.002 checks in
|
||
+`research/j2-slot-2026-10-12/REVIEW.md` exceeded the fixed late-time mass-drift
|
||
+tolerance. It is not yet validated for steady mass balance or long runs.
|
||
--- a/src/CelerisLab/cuda/compiler_v2.py
|
||
+++ b/src/CelerisLab/cuda/compiler_v2.py
|
||
@@ -176,6 +176,8 @@
|
||
#define OUTLET_BACKFLOW_CLAMP {m['OUTLET_BACKFLOW_CLAMP']}
|
||
#define Y_WALL_BC {m['Y_WALL_BC']}
|
||
#define WALL_NORMAL_VELOCITY {m['WALL_NORMAL_VELOCITY']}
|
||
+#define WALL_SLOT_WIDTH {m['WALL_SLOT_WIDTH']}
|
||
+#define TOP_WALL_NORMAL_VELOCITY {m['TOP_WALL_NORMAL_VELOCITY']}
|
||
|
||
#define OMEGA_COLLISION_MIN {m['OMEGA_COLLISION_MIN']}
|
||
#define OMEGA_COLLISION_MAX {m['OMEGA_COLLISION_MAX']}
|
||
--- a/src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh
|
||
+++ b/src/CelerisLab/lbm/kernels/boundary/planar_velocity.cuh
|
||
@@ -3,6 +3,18 @@
|
||
// the caller owns collision/store. No force, shifting or off-lattice walls.
|
||
#ifndef CELERIS_BOUNDARY_PLANAR_VELOCITY_CUH
|
||
#define CELERIS_BOUNDARY_PLANAR_VELOCITY_CUH
|
||
+
|
||
+__device__ __forceinline__ float bottom_wall_normal_velocity(unsigned int x)
|
||
+{
|
||
+#if WALL_SLOT_WIDTH > 0 && WALL_SLOT_WIDTH < NX
|
||
+ // C1 compact profile centered at (NX-1)/2; amplitude is its continuous peak.
|
||
+ const float s = (2.0f*(float)x-(float)(NX-1))/(float)WALL_SLOT_WIDTH;
|
||
+ const float w = fmaxf(0.0f, 1.0f-s*s);
|
||
+ return WALL_NORMAL_VELOCITY*w*w;
|
||
+#else
|
||
+ return WALL_NORMAL_VELOCITY; // Disabled or explicitly uniform full width.
|
||
+#endif
|
||
+}
|
||
|
||
__device__ __forceinline__ void reconstruct_bottom_velocity(float* f, float u, float v)
|
||
{
|
||
--- a/src/CelerisLab/lbm/kernels/step/init_flow.cu
|
||
+++ b/src/CelerisLab/lbm/kernels/step/init_flow.cu
|
||
@@ -91,7 +91,13 @@
|
||
// Initial g(0)=feq: both boundary velocities are already satisfied.
|
||
// Start from linear Couette, not the sought exponential steady profile.
|
||
float f[NQ];
|
||
+#if WALL_SLOT_WIDTH > 0 && WALL_SLOT_WIDTH < NX
|
||
+ const float t = (float)(y-1u)/(float)(NY-3);
|
||
+ const float v = (1.0f-t)*bottom_wall_normal_velocity(x)+t*TOP_WALL_NORMAL_VELOCITY;
|
||
+ compute_feq(RHO, 0.0f, v, f);
|
||
+#else
|
||
compute_feq(RHO, U0*(float)(y-1u)/(float)(NY-3), WALL_NORMAL_VELOCITY, f);
|
||
+#endif
|
||
stream_pull_store(k, f, fi);
|
||
#else
|
||
write_equilibrium(fi, k, inlet_target_u((float)y));
|
||
--- a/src/CelerisLab/lbm/kernels/step/one_step_double.cu
|
||
+++ b/src/CelerisLab/lbm/kernels/step/one_step_double.cu
|
||
@@ -91,8 +91,8 @@
|
||
#if Y_WALL_BC == 2
|
||
// Reconstruct f(t+dt) on the physical plane, then collide below.
|
||
// This replaces, rather than compounds, the halfway wall correction.
|
||
- if (y == 1u) reconstruct_bottom_velocity(f, 0.0f, WALL_NORMAL_VELOCITY);
|
||
- else reconstruct_top_velocity(f, U0, WALL_NORMAL_VELOCITY);
|
||
+ if (y == 1u) reconstruct_bottom_velocity(f, 0.0f, bottom_wall_normal_velocity(x));
|
||
+ else reconstruct_top_velocity(f, U0, TOP_WALL_NORMAL_VELOCITY);
|
||
#elif Y_WALL_BC == 1
|
||
apply_wall_freeslip_d2q9_y_pull(y, f, fi_in, k);
|
||
#else
|
||
--- /dev/null
|
||
+++ b/tests/unit/cpp/slot_profile.cpp
|
||
@@ -0,0 +1,10 @@
|
||
+#include <cmath>
|
||
+#include <cstdio>
|
||
+#define __device__
|
||
+#define __forceinline__ inline
|
||
+#include "boundary/planar_velocity.cuh"
|
||
+
|
||
+int main() {
|
||
+ for (unsigned x=0; x<NX; ++x)
|
||
+ printf("%.9g\n", bottom_wall_normal_velocity(x));
|
||
+}
|
||
--- a/tests/unit/test_planar_velocity.py
|
||
+++ b/tests/unit/test_planar_velocity.py
|
||
@@ -19,7 +19,7 @@
|
||
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'
|
||
- lines = compile_run(tmp_path, 'planar_velocity.cpp', {})
|
||
+ 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])
|
||
w = np.array([4/9]+[1/9]*4+[1/36]*4)
|
||
--- /dev/null
|
||
+++ b/tests/unit/test_slot_velocity.py
|
||
@@ -0,0 +1,81 @@
|
||
+"""Static slot targets and checkpoint compatibility, using the actual CUDA helper."""
|
||
+import json
|
||
+from types import SimpleNamespace
|
||
+
|
||
+import h5py
|
||
+import numpy as np
|
||
+import pytest
|
||
+
|
||
+from CelerisLab.config import LBMConfig, load_lbm_config
|
||
+from CelerisLab.common.checkpoint import save_checkpoint, load_checkpoint
|
||
+from .test_numerical_operators import compile_run, ROOT
|
||
+from .test_state_regressions import device, make_field, manager
|
||
+
|
||
+
|
||
+@pytest.mark.parametrize('width', [0, 16, 48])
|
||
+@pytest.mark.parametrize('amplitude', [0., .002, -.002])
|
||
+def test_actual_slot_targets_and_requested_volume(tmp_path, width, amplitude):
|
||
+ cfg = LBMConfig(nx=48, velocity=0., y_wall_bc='on_site_velocity',
|
||
+ wall_slot_width=width, wall_normal_velocity=amplitude)
|
||
+ macros = cfg.to_macros()
|
||
+ values = np.array(compile_run(tmp_path, 'slot_profile.cpp', {
|
||
+ k: macros[k] for k in ('NX', 'WALL_SLOT_WIDTH', 'WALL_NORMAL_VELOCITY')}), dtype=float)
|
||
+ # Independent node quadrature; includes both zero flanks and the full-width exception.
|
||
+ x = np.arange(48)-23.5
|
||
+ expected = amplitude*np.clip(1-(x/8)**2, 0, None)**2 if width == 16 else np.full(48, amplitude)
|
||
+ np.testing.assert_allclose(values, expected, atol=5e-10, rtol=0)
|
||
+ np.testing.assert_array_equal(values, values[::-1])
|
||
+ assert abs(values.sum()-48*cfg.top_wall_normal_velocity) <= 1e-8
|
||
+ assert abs(float(macros['TOP_WALL_NORMAL_VELOCITY'][:-1])-expected.mean()) < 1e-16
|
||
+
|
||
+
|
||
+@pytest.mark.parametrize('change', [dict(wall_slot_width=w) for w in (-1, 1, 49, 2.5, True)]
|
||
+ + [dict(y_wall_bc='bounce_back'), dict(velocity=.01)])
|
||
+def test_slot_invalid_configuration(change):
|
||
+ values = dict(nx=48, velocity=0., y_wall_bc='on_site_velocity', wall_slot_width=16)
|
||
+ values.update(change)
|
||
+ with pytest.raises(ValueError, match='wall_slot_width'):
|
||
+ LBMConfig(**values).validate_runtime()
|
||
+
|
||
+
|
||
+def test_slot_json_and_static_signature(tmp_path):
|
||
+ data = json.loads((ROOT/'src/CelerisLab/configs/config_lbm.json').read_text())
|
||
+ data['physics']['velocity'] = 0.
|
||
+ data['method'].update(y_wall_bc='on_site_velocity', wall_slot_width=16,
|
||
+ collision='SRT', streaming='double_buffer', store_precision='FP32',
|
||
+ ddf_shifting=False)
|
||
+ data['method']['les']['enabled'] = False
|
||
+ path = tmp_path/'config.json'
|
||
+ path.write_text(json.dumps(data))
|
||
+ cfg = load_lbm_config(path)
|
||
+ assert cfg.wall_slot_width == 16
|
||
+ before = cfg.allocation_signature()
|
||
+ cfg.wall_slot_width = 18
|
||
+ assert before != cfg.allocation_signature()
|
||
+
|
||
+
|
||
+@pytest.mark.parametrize('width, saved_width, accepted', [(6, 4, False), (6, None, False),
|
||
+ (0, None, True), (6, 6, True)])
|
||
+def test_slot_checkpoint_width_before_device_writes(device, tmp_path, width, saved_width, accepted):
|
||
+ field = make_field(device, velocity=0., y_wall_bc='on_site_velocity', wall_slot_width=width)
|
||
+ bodies = manager(field)
|
||
+ bodies.sync_to_gpu(field)
|
||
+ stepper = SimpleNamespace(_step_count=0)
|
||
+ path = save_checkpoint(field, SimpleNamespace(step_count=3), field.cfg, bodies,
|
||
+ str(tmp_path/'slot.h5'))
|
||
+ with h5py.File(path, 'a') as hf:
|
||
+ assert hf.attrs['ckpt_version'] == 3
|
||
+ config = json.loads(hf.attrs['config_json'])
|
||
+ config.pop('wall_slot_width')
|
||
+ if saved_width is not None:
|
||
+ config['wall_slot_width'] = saved_width
|
||
+ hf.attrs['config_json'] = json.dumps(config)
|
||
+ before, writes = field.ddf.copy(), device.writes
|
||
+ if accepted:
|
||
+ load_checkpoint(path, field, stepper, field.cfg, bodies)
|
||
+ np.testing.assert_array_equal(field.ddf, before)
|
||
+ assert stepper._step_count == 3
|
||
+ else:
|
||
+ with pytest.raises(ValueError, match='wall_slot_width'):
|
||
+ load_checkpoint(path, field, stepper, field.cfg, bodies)
|
||
+ assert device.writes == writes and stepper._step_count == 0
|
||
--- a/tests/validation/run_j1.py
|
||
+++ b/tests/validation/run_j1.py
|
||
@@ -28,13 +28,13 @@
|
||
u_change=2e-5, v_change=2e-7, budget_relative_residual=2e-7)
|
||
|
||
|
||
-def make_sim(v, mode='on_site_velocity'):
|
||
+def make_sim(v, mode='on_site_velocity', nx=16, top_u=.02, slot_width=0):
|
||
config = json.loads((ROOT/'src/CelerisLab/configs/config_lbm.json').read_text())
|
||
- config['grid'].update(nx=16, ny=35, nz=1, lattice_model='D2Q9')
|
||
- config['physics'].update(viscosity=.1, velocity=.02, rho=1.)
|
||
+ config['grid'].update(nx=nx, ny=35, nz=1, lattice_model='D2Q9')
|
||
+ config['physics'].update(viscosity=.1, velocity=top_u, rho=1.)
|
||
config['method'].update(collision='SRT', streaming='double_buffer',
|
||
store_precision='FP32', ddf_shifting=False,
|
||
- y_wall_bc=mode, wall_normal_velocity=v)
|
||
+ y_wall_bc=mode, wall_normal_velocity=v, wall_slot_width=slot_width)
|
||
config['method']['les']['enabled'] = False
|
||
with TemporaryDirectory(prefix='celeris-j1-config-') as folder:
|
||
lbm, body = Path(folder)/'lbm.json', Path(folder)/'body.json'
|
||
@@ -72,7 +72,7 @@
|
||
'''
|
||
|
||
|
||
-def population_budget(sim):
|
||
+def population_budget(sim, bottom_slot=None):
|
||
"""16 real steps, separate incoming/outgoing/collision records, no targets."""
|
||
saved = sim.get_ddf().copy()
|
||
count = sim.stepper.step_count
|
||
@@ -122,6 +122,14 @@
|
||
collision_roundoff=float(collision.sum()),
|
||
physical_residual=physical_residual,
|
||
budget_residual=residual))
|
||
+ if bottom_slot is not None:
|
||
+ outgoing = before[[4,6,7], 1, :].sum(axis=0)
|
||
+ rows[-1]['exchange'] = {
|
||
+ name: dict(incoming=float(incoming[:sim.field.nx][mask].sum()),
|
||
+ outgoing=float(outgoing[mask].sum()))
|
||
+ for name, mask in (('slot', bottom_slot), ('flanks', ~bottom_slot))}
|
||
+ rows[-1]['exchange']['top'] = dict(
|
||
+ incoming=float(incoming[sim.field.nx:].sum()), outgoing=float(outgoing_top))
|
||
identical = np.array_equal(sim.get_ddf(), plain)
|
||
return dict(steps=rows, instrumented_matches_plain=identical,
|
||
audit_source_sha256=audit_hash,
|
||
--- /dev/null
|
||
+++ b/tests/validation/run_j2_slot.py
|
||
@@ -0,0 +1,151 @@
|
||
+"""Bounded static-slot validation; frozen gates in research/j2-slot-2026-10-12/REVIEW.md."""
|
||
+import argparse
|
||
+import json
|
||
+import os
|
||
+from pathlib import Path
|
||
+import sys
|
||
+from tempfile import TemporaryDirectory
|
||
+from unittest.mock import patch
|
||
+
|
||
+import numpy as np
|
||
+import CelerisLab
|
||
+from CelerisLab.cuda import compiler_v2 as compiler
|
||
+from run_j1 import ROOT, make_sim, populations, observables, population_budget, off_path_check
|
||
+
|
||
+NX, WIDTH, CHUNK = 48, 16, 519
|
||
+GATES = dict(velocity_change=5e-5, mass_change=2e-6, mass_window_span=6e-6,
|
||
+ mass_offset=1e-3, rho_error=.005, speed=.01, zero_speed=2e-6,
|
||
+ flux_error=.005, flux_span=.002, boundary_error=2e-6,
|
||
+ rho_symmetry=2e-5, velocity_symmetry=.002, budget_residual=2e-7)
|
||
+
|
||
+
|
||
+def targets(amplitude):
|
||
+ # Independent discrete quadrature, including the zero-target flanks.
|
||
+ x = np.arange(NX)-(NX-1)/2
|
||
+ bottom = amplitude*np.clip(1-(x/(WIDTH/2))**2, 0, None)**2
|
||
+ return bottom, bottom.mean(), np.abs(x) < WIDTH/2
|
||
+
|
||
+
|
||
+def full_width_check(baseline):
|
||
+ snapshots = []
|
||
+ for source, width in ((baseline, 0), (Path(compiler.kernel_path('')), NX)):
|
||
+ with patch.object(compiler, '_KERNEL_DIR', str(source)):
|
||
+ with make_sim(.002, nx=NX, top_u=0., slot_width=width) as sim:
|
||
+ sim.initialize()
|
||
+ initial = sim.get_ddf().copy()
|
||
+ sim.run(64)
|
||
+ snapshots.append((initial, sim.get_ddf().copy(), sim.get_flags().copy()))
|
||
+ return bool(all(np.array_equal(a, b) for a, b in zip(*snapshots)))
|
||
+
|
||
+
|
||
+def restart_check(sim, amplitude):
|
||
+ with TemporaryDirectory(prefix='celeris-j2-checkpoint-') as folder:
|
||
+ path = sim.save_checkpoint(str(Path(folder)/'slot.h5'))
|
||
+ sim.run(7)
|
||
+ expected, flags, count = sim.get_ddf().copy(), sim.get_flags().copy(), sim.stepper.step_count
|
||
+ with make_sim(amplitude, nx=NX, top_u=0., slot_width=WIDTH) as target:
|
||
+ target.load_checkpoint(path)
|
||
+ target.run(7)
|
||
+ equal = (np.array_equal(expected, target.get_ddf())
|
||
+ and np.array_equal(flags, target.get_flags()) and count == target.stepper.step_count)
|
||
+ sim.load_checkpoint(path)
|
||
+ return bool(equal)
|
||
+
|
||
+
|
||
+def run_case(amplitude):
|
||
+ bottom, top, slot = targets(amplitude)
|
||
+ scale = max(abs(amplitude), .002)
|
||
+ with make_sim(amplitude, nx=NX, top_u=0., slot_width=WIDTH) as sim:
|
||
+ sim.initialize()
|
||
+ rho, ux, uy = observables(populations(sim))
|
||
+ initial_mass = float(rho.sum())
|
||
+ initial_error = float(max(np.abs(uy[0]-bottom).max(), np.abs(uy[-1]-top).max(), np.abs(ux).max()))
|
||
+ old_u, old_v, old_mass = ux.copy(), uy.copy(), initial_mass
|
||
+ history = [dict(step=0, mass=initial_mass)]
|
||
+ steady = False
|
||
+ for sample in range(1, 41):
|
||
+ sim.run(CHUNK)
|
||
+ f = populations(sim)
|
||
+ assert np.isfinite(f).all(), 'Nonfinite slot populations'
|
||
+ rho, ux, uy = observables(f)
|
||
+ mass = float(rho.sum())
|
||
+ history.append(dict(step=sim.stepper.step_count, mass=mass,
|
||
+ u_change=float(np.sqrt(np.mean((ux-old_u)**2))/scale),
|
||
+ v_change=float(np.sqrt(np.mean((uy-old_v)**2))/scale),
|
||
+ mass_change=abs(mass-old_mass)/initial_mass))
|
||
+ old_u, old_v, old_mass = ux.copy(), uy.copy(), mass
|
||
+ span = float(np.ptp([r['mass'] for r in history[-5:]])/initial_mass)
|
||
+ steady = bool(sample >= 8 and span <= GATES['mass_window_span'] and all(
|
||
+ r['u_change'] <= GATES['velocity_change'] and r['v_change'] <= GATES['velocity_change']
|
||
+ and r['mass_change'] <= GATES['mass_change'] for r in history[-4:]))
|
||
+ if steady:
|
||
+ break
|
||
+ sections = [8, 16, 24]
|
||
+ actual_mass_flux = (rho*uy).sum(axis=1)
|
||
+ actual_volume_flux = uy.sum(axis=1)
|
||
+ request_q = float(bottom.sum())
|
||
+ reference = float(rho.mean()*request_q)
|
||
+ flux_error = float(np.max(np.abs(actual_mass_flux[sections]-reference))/abs(reference)) if amplitude else 0.
|
||
+ flux_span = float(np.ptp(actual_mass_flux[sections])/abs(reference)) if amplitude else 0.
|
||
+ symmetry = dict(rho=float(np.max(np.abs(rho-rho[:, ::-1]))),
|
||
+ u=float(np.max(np.abs(ux+ux[:, ::-1]))/scale),
|
||
+ v=float(np.max(np.abs(uy-uy[:, ::-1]))/scale))
|
||
+ boundary_error = float(max(np.abs(uy[0]-bottom).max(), np.abs(uy[-1]-top).max(),
|
||
+ np.abs(ux[[0, -1]]).max()))
|
||
+ offset = abs(mass-initial_mass)/initial_mass
|
||
+ speed = float(np.hypot(ux, uy).max())
|
||
+ density_error = float(np.abs(rho-1).max())
|
||
+ nonuniform = bool(uy[8, NX//2]*amplitude > 0 and np.abs(ux).max() > 1e-5) if amplitude else True
|
||
+ restart = restart_check(sim, amplitude) if amplitude > 0 else None
|
||
+ budget = population_budget(sim, bottom_slot=slot)
|
||
+ checks = dict(steady=steady, initialization=initial_error <= GATES['boundary_error'],
|
||
+ boundary=boundary_error <= GATES['boundary_error'],
|
||
+ mass=offset <= GATES['mass_offset'] and span <= GATES['mass_window_span'],
|
||
+ bounded=density_error <= GATES['rho_error'] and speed <= GATES['speed'],
|
||
+ flux=flux_error <= GATES['flux_error'] and flux_span <= GATES['flux_span'],
|
||
+ symmetry=symmetry['rho'] <= GATES['rho_symmetry'] and
|
||
+ max(symmetry['u'], symmetry['v']) <= GATES['velocity_symmetry'],
|
||
+ zero=bool(amplitude or speed <= GATES['zero_speed']), nonuniform=nonuniform,
|
||
+ restart=restart is not False, audit=budget['instrumented_matches_plain'] and
|
||
+ budget['max_relative_residual'] <= GATES['budget_residual'])
|
||
+ checks = {name: bool(passed) for name, passed in checks.items()}
|
||
+ return dict(amplitude=amplitude, width=WIDTH, nx=NX, height=32, viscosity=.1, top_u=0.,
|
||
+ steps=sim.stepper.step_count-16, max_steps=40*CHUNK, checks=checks, passed=all(checks.values()),
|
||
+ requested_volume=request_q, top_target=top, bottom_targets=bottom.tolist(),
|
||
+ velocity_scale=scale, rho_scale=1., initial_boundary_error=initial_error,
|
||
+ boundary_error=boundary_error, mass_offset=offset, mass_window_span=span,
|
||
+ flux_error=flux_error, flux_span=flux_span, section_y=sections,
|
||
+ mass_flux_by_y=actual_mass_flux.tolist(), volume_flux_by_y=actual_volume_flux.tolist(),
|
||
+ symmetry=symmetry, max_speed=speed, rho_error=density_error,
|
||
+ rho=rho.tolist(), u=ux.tolist(), v=uy.tolist(), history=history,
|
||
+ checkpoint_bitwise_equal=restart, population_budget=budget)
|
||
+
|
||
+
|
||
+def main():
|
||
+ parser = argparse.ArgumentParser(description=__doc__)
|
||
+ parser.add_argument('--output', type=Path, required=True)
|
||
+ parser.add_argument('--baseline-kernels', type=Path, required=True,
|
||
+ help='Pre-J2 source for the required full-width and default-mode comparisons')
|
||
+ args = parser.parse_args()
|
||
+ assert Path.cwd() == ROOT == Path('/home/frank14f/CelerisLab')
|
||
+ assert sys.executable == '/home/frank14f/anaconda3/envs/pycuda_3_10/bin/python'
|
||
+ 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 os.environ.get('CUDA_VISIBLE_DEVICES') in ('0', '1')
|
||
+ report = dict(python=sys.executable, package=CelerisLab.__file__,
|
||
+ cuda_source=compiler.kernel_path('kernel_v2.cu'), physical_gpu=os.environ['CUDA_VISIBLE_DEVICES'],
|
||
+ gates=GATES, full_width_bitwise_equal=None, off_path_bitwise_equal=None, cases=[])
|
||
+ report['full_width_bitwise_equal'] = full_width_check(args.baseline_kernels)
|
||
+ report['off_path_bitwise_equal'] = off_path_check(args.baseline_kernels)
|
||
+ for amplitude in (0., .002, -.002):
|
||
+ case = run_case(amplitude)
|
||
+ report['cases'].append(case)
|
||
+ args.output.write_text(json.dumps(report, indent=2)+'\n')
|
||
+ print(json.dumps({k: case[k] for k in ('amplitude', 'steps', 'checks', 'mass_offset',
|
||
+ 'mass_window_span', 'flux_error', 'symmetry')}), flush=True)
|
||
+ return 0 if (report['full_width_bitwise_equal'] is True and report['off_path_bitwise_equal'] is True
|
||
+ and all(c['passed'] for c in report['cases'])) else 1
|
||
+
|
||
+
|
||
+if __name__ == '__main__':
|
||
+ raise SystemExit(main())
|