--- 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 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 +#include +#define __device__ +#define __forceinline__ inline +#include "boundary/planar_velocity.cuh" + +int main() { + for (unsigned x=0; x= 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())