diff --git a/docs/lifecycle.md b/docs/lifecycle.md index 615c878..c3e5c79 100644 --- a/docs/lifecycle.md +++ b/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`. -- 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`). diff --git a/src/CelerisLab/configs/CONFIG.md b/src/CelerisLab/configs/CONFIG.md index c50411e..3724e18 100644 --- a/src/CelerisLab/configs/CONFIG.md +++ b/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 diff --git a/src/CelerisLab/cuda/compiler_v2.py b/src/CelerisLab/cuda/compiler_v2.py index ec22b02..6916d1d 100644 --- a/src/CelerisLab/cuda/compiler_v2.py +++ b/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(): diff --git a/src/CelerisLab/lbm/block_stepper.py b/src/CelerisLab/lbm/block_stepper.py index 6260aa1..1b7ce9e 100644 --- a/src/CelerisLab/lbm/block_stepper.py +++ b/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) diff --git a/src/CelerisLab/lbm/kernels/blocks.cu b/src/CelerisLab/lbm/kernels/blocked.cu similarity index 97% rename from src/CelerisLab/lbm/kernels/blocks.cu rename to src/CelerisLab/lbm/kernels/blocked.cu index 65d1320..18f7e42 100644 --- a/src/CelerisLab/lbm/kernels/blocks.cu +++ b/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 diff --git a/src/CelerisLab/lbm/kernels/step/aux_kernels.cu b/src/CelerisLab/lbm/kernels/step/aux_kernels.cu index 2b985a3..025d099 100644 --- a/src/CelerisLab/lbm/kernels/step/aux_kernels.cu +++ b/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. diff --git a/src/CelerisLab/lbm/kernels/step/block_boundary.cu b/src/CelerisLab/lbm/kernels/step/block_boundary.cu index 9b89451..387a68e 100644 --- a/src/CelerisLab/lbm/kernels/step/block_boundary.cu +++ b/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 diff --git a/src/CelerisLab/lbm/kernels/step/block_step.cu b/src/CelerisLab/lbm/kernels/step/block_step.cu index dc1a38c..0207c7b 100644 --- a/src/CelerisLab/lbm/kernels/step/block_step.cu +++ b/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 diff --git a/src/CelerisLab/lbm/kernels/step/block_transfer.cu b/src/CelerisLab/lbm/kernels/step/block_transfer.cu index b2b09cb..9b14fed 100644 --- a/src/CelerisLab/lbm/kernels/step/block_transfer.cu +++ b/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( diff --git a/src/CelerisLab/lbm/kernels/step/esopull_macro.cu b/src/CelerisLab/lbm/kernels/step/esopull_macro.cu index dc91964..07a52f6 100644 --- a/src/CelerisLab/lbm/kernels/step/esopull_macro.cu +++ b/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). diff --git a/src/CelerisLab/lbm/kernels/step/esopull_sync.cu b/src/CelerisLab/lbm/kernels/step/esopull_sync.cu index 30ce615..6ac4316 100644 --- a/src/CelerisLab/lbm/kernels/step/esopull_sync.cu +++ b/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 diff --git a/src/CelerisLab/lbm/kernels/step/init_flow.cu b/src/CelerisLab/lbm/kernels/step/init_flow.cu index a31072e..3478b37 100644 --- a/src/CelerisLab/lbm/kernels/step/init_flow.cu +++ b/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; diff --git a/src/CelerisLab/lbm/kernels/step/one_step_double.cu b/src/CelerisLab/lbm/kernels/step/one_step_double.cu index ef582b9..00e361e 100644 --- a/src/CelerisLab/lbm/kernels/step/one_step_double.cu +++ b/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 diff --git a/src/CelerisLab/lbm/kernels/step/one_step_esopull.cu b/src/CelerisLab/lbm/kernels/step/one_step_esopull.cu index b853382..3e0881d 100644 --- a/src/CelerisLab/lbm/kernels/step/one_step_esopull.cu +++ b/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. diff --git a/src/CelerisLab/lbm/kernels/kernel_v2.cu b/src/CelerisLab/lbm/kernels/uniform.cu similarity index 96% rename from src/CelerisLab/lbm/kernels/kernel_v2.cu rename to src/CelerisLab/lbm/kernels/uniform.cu index 5637974..051dd18 100644 --- a/src/CelerisLab/lbm/kernels/kernel_v2.cu +++ b/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 diff --git a/src/CelerisLab/lbm/refinement_stepper.py b/src/CelerisLab/lbm/refinement_stepper.py index 2a387fc..45e87da 100644 --- a/src/CelerisLab/lbm/refinement_stepper.py +++ b/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") diff --git a/src/CelerisLab/lbm/stepper.py b/src/CelerisLab/lbm/stepper.py index 940e62d..d141688 100644 --- a/src/CelerisLab/lbm/stepper.py +++ b/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") diff --git a/src/CelerisLab/simulation.py b/src/CelerisLab/simulation.py index c2cfc95..4b47062 100644 --- a/src/CelerisLab/simulation.py +++ b/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) diff --git a/tests/integration/test_runtime_contracts_gpu.py b/tests/integration/test_runtime_contracts_gpu.py index e39d38e..5692c7f 100644 --- a/tests/integration/test_runtime_contracts_gpu.py +++ b/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(): diff --git a/tests/unit/test_build_contracts.py b/tests/unit/test_build_contracts.py index 6b5fa16..5c1c12a 100644 --- a/tests/unit/test_build_contracts.py +++ b/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) diff --git a/tests/unit/test_numerical_operators.py b/tests/unit/test_numerical_operators.py index ebf261f..c76d8b6 100644 --- a/tests/unit/test_numerical_operators.py +++ b/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 diff --git a/tests/unit/test_planar_velocity.py b/tests/unit/test_planar_velocity.py index d007caf..a91cb59 100644 --- a/tests/unit/test_planar_velocity.py +++ b/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]) diff --git a/tests/unit/test_refinement_mrt.py b/tests/unit/test_refinement_mrt.py index dd2f93a..e0de52f 100644 --- a/tests/unit/test_refinement_mrt.py +++ b/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) diff --git a/tests/validation/multigrid/channel_parity_diagnostic/replay.py b/tests/validation/multigrid/channel_parity_diagnostic/replay.py index 695a078..7bf31d6 100644 --- a/tests/validation/multigrid/channel_parity_diagnostic/replay.py +++ b/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')) diff --git a/tests/validation/multigrid/test_block_stepper.py b/tests/validation/multigrid/test_block_stepper.py index 4ec7abc..c6eb987 100644 --- a/tests/validation/multigrid/test_block_stepper.py +++ b/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() diff --git a/tests/validation/multigrid/test_channel_root_gpu.py b/tests/validation/multigrid/test_channel_root_gpu.py index 905b720..fccdbcb 100644 --- a/tests/validation/multigrid/test_channel_root_gpu.py +++ b/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