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>

docs/lifecycle.md

@@ -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`.
+- Original uniform grid (`uniform.cu`, `Simulation`): whole-domain `NX`/`NY`.
-- Blocked levels (`blocks.cu`; `PeriodicBlockStepper` and every blocked
+- 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`).
 

src/CelerisLab/configs/CONFIG.md

@@ -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

src/CelerisLab/cuda/compiler_v2.py

@@ -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():

src/CelerisLab/lbm/block_stepper.py

@@ -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)

src/CelerisLab/lbm/kernels/blocked.cu

@@ -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

src/CelerisLab/lbm/kernels/step/aux_kernels.cu

@@ -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.

src/CelerisLab/lbm/kernels/step/block_boundary.cu

@@ -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

src/CelerisLab/lbm/kernels/step/block_step.cu

@@ -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

src/CelerisLab/lbm/kernels/step/block_transfer.cu

@@ -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(

src/CelerisLab/lbm/kernels/step/esopull_macro.cu

@@ -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).

src/CelerisLab/lbm/kernels/step/esopull_sync.cu

@@ -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

src/CelerisLab/lbm/kernels/step/init_flow.cu

@@ -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;

src/CelerisLab/lbm/kernels/step/one_step_double.cu

@@ -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

src/CelerisLab/lbm/kernels/step/one_step_esopull.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.

src/CelerisLab/lbm/kernels/uniform.cu

@@ -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

src/CelerisLab/lbm/refinement_stepper.py

@@ -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")

src/CelerisLab/lbm/stepper.py

@@ -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")

src/CelerisLab/simulation.py

@@ -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)
 

tests/integration/test_runtime_contracts_gpu.py

@@ -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():

tests/unit/test_build_contracts.py

@@ -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)
 
 

tests/unit/test_numerical_operators.py

@@ -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

tests/unit/test_planar_velocity.py

@@ -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])

tests/unit/test_refinement_mrt.py

@@ -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)

tests/validation/multigrid/channel_parity_diagnostic/replay.py

@@ -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'))

tests/validation/multigrid/test_block_stepper.py

@@ -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()
 

tests/validation/multigrid/test_channel_root_gpu.py

@@ -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