kernels: name compilation entries by storage route (uniform.cu, blocked.cu)

kernel_v2.cu -> uniform.cu (whole-domain NX/NY, Simulation) and
blocks.cu -> blocked.cu (one block window, every blocked level), with
InitTubeFlow_v2 -> InitTubeFlow. All live callers, include comments,
CONFIG.md, lifecycle doc and current-tree tests move together; no alias.
refinement.cu keeps its name until the dense-route decision.
Historical-revision tooling that names the old files (LES source pins,
timing/J campaign runners, TG run_case, supported_settings citations at
888c8038, development notes) is unchanged on purpose.
No arithmetic change; the PTX symbol InitTubeFlow_v2 is renamed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This commit is contained in:
Claude
2026-10-26 14:07:34 +00:00
co-authored by Claude Opus 5.5
parent dbcdb87081
commit 2b3c2a7529
26 changed files with 34 additions and 34 deletions
+2 -2
View File
@@ -12,8 +12,8 @@ it is; it adds no framework. Line references are to candidate 0f8fdbc3.
Storage routes are separate concerns:
- Original uniform grid (`kernel_v2.cu`, `Simulation`): whole-domain `NX`/`NY`.
- Blocked levels (`blocks.cu`; `PeriodicBlockStepper` and every blocked
- Original uniform grid (`uniform.cu`, `Simulation`): whole-domain `NX`/`NY`.
- Blocked levels (`blocked.cu`; `PeriodicBlockStepper` and every blocked
`FixedPatchStepper` level): `NX`/`NY` are one block window including halo.
- Dense refinement levels (`refinement.cu`, `level_blocks[l] is None`).
+1 -1
View File
@@ -60,7 +60,7 @@ 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_v2.generate_config` 会复制 kernel 源码到独立临时构建目录,并在副本中生成配置;不会改写包内头文件或 PTX。返回目录须显式传给 `compile_kernel(build_dir=...)`,编译结果与配置摘要均保存在该目录。`Simulation` 负责目录清理;直接使用编译器的调用者须调用 `cleanup_build_dir()`。跳过生成配置直接编译仓库中的 `uniform.cu` 属不受支持用法。
```
config_lbm.json
+2 -2
View File
@@ -7,7 +7,7 @@ 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, source=entry, build_dir=build_dir) — nvcc -ptx
3. load_module(ptx_path) — PyCUDA module from PTX
Entries: kernel_v2.cu (uniform grid), blocks.cu (blocked level), refinement.cu
Entries: uniform.cu (uniform grid), blocked.cu (blocked level), refinement.cu
(dense refinement level).
"""
@@ -241,7 +241,7 @@ def _render_config(m: dict, n_objects: int, cdir: str, *, block_root_bc: int = 0
# ---------------------------------------------------------------------------
# Compilation
# ---------------------------------------------------------------------------
def compile_kernel(arch: str = "sm_70", source: str = "kernel_v2.cu",
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():
+1 -1
View File
@@ -128,7 +128,7 @@ class PeriodicBlockStepper:
config.validate()
if not np.isclose(config.omega, 1 / (3 * checked.viscosity + .5), rtol=1e-14, atol=0):
raise RuntimeError("block omega/tau mismatch")
module = compiler.build_module(config, "blocks.cu", arch=self.ctx.sm_arch)
module = compiler.build_module(config, "blocked.cu", arch=self.ctx.sm_arch)
compiler.set_params(module, config.omega)
self._module = module
self._current = self._allocate(self._bytes)
@@ -1,4 +1,4 @@
// CelerisLab/lbm/kernels/blocks.cu
// 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
@@ -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.
@@ -1,6 +1,6 @@
// CelerisLab/lbm/kernels/step/block_boundary.cu
// Receiver-local curved-boundary staging and its raw-exchange accumulation.
// Part of the blocks.cu compilation entry; relies on its includes and guards.
// 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
@@ -1,6 +1,6 @@
// CelerisLab/lbm/kernels/step/block_step.cu
// Generic blocked pull/collide/store for periodic/fine and slot/reservoir levels.
// Part of the blocks.cu compilation entry; relies on its includes and guards.
// 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
@@ -1,6 +1,6 @@
// CelerisLab/lbm/kernels/step/block_transfer.cu
// Same-level halo refresh and 2:1 explosion/coalescence between blocked levels.
// Part of the blocks.cu compilation entry; relies on its includes and guards.
// 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(
@@ -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
+3 -3
View File
@@ -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.
@@ -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
+2 -2
View File
@@ -124,7 +124,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
@@ -356,7 +356,7 @@ class FixedPatchStepper:
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"
name, source = ("BlockChannelStep" if index == 0 and self._channel else "BlockStep"), "blocked.cu"
config.validate()
if not np.isclose(config.omega, 1 / tau, rtol=1e-14, atol=0):
raise RuntimeError("level omega/tau mismatch")
+1 -1
View File
@@ -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")
+2 -2
View File
@@ -58,7 +58,7 @@ class Simulation:
try:
with self.ctx.activate():
arch = self._resolve_compile_arch()
self._module = compiler.build_module(self.lbm_cfg, "kernel_v2.cu", arch=arch)
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
@@ -409,7 +409,7 @@ class Simulation:
self.lbm_cfg.validate_runtime()
self._assert_allocation_contract()
arch = self._resolve_compile_arch()
self._module = compiler.build_module(self.lbm_cfg, "kernel_v2.cu", arch=arch,
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)
@@ -13,7 +13,7 @@ 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_v2.kernel_path("uniform.cu")).resolve() == root / "lbm/kernels/uniform.cu"
def test_two_contexts_interleaved_and_non_lifo_close():
+1 -1
View File
@@ -168,7 +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)
self.assertEqual(sorted(Path(build).glob("kernel_v2-*.ptx")), []) # no partial output
self.assertEqual(sorted(Path(build).glob("uniform-*.ptx")), []) # no partial output
compiler.cleanup_build_dir(build)
+1 -1
View File
@@ -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
+1 -1
View File
@@ -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])
+1 -1
View File
@@ -22,7 +22,7 @@ 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() + (KERNELS / "step/block_step.cu").read_text()
blocks = (KERNELS / "blocked.cu").read_text() + (KERNELS / "step/block_step.cu").read_text()
channel = (KERNELS / "step/block_channel.cu").read_text()
block_collision = (KERNELS / "operators/block_collision.cuh").read_text()
self.assertIn('#include "operators/block_collision.cuh"', blocks)
@@ -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'))
@@ -171,7 +171,7 @@ class CpuContracts(unittest.TestCase):
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")
self.assertEqual(compile_kernel.call_args.kwargs["source"], "blocked.cu")
alloc.assert_not_called()
upload.assert_not_called()
@@ -27,7 +27,7 @@ def state(ny, nx):
def compile_module(cfg, context, source):
cfg.validate()
build = compiler.generate_config(cfg, block_root_bc=2 if source == "blocks.cu" else 0)
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