CLAUDE.md вне версионирования: один файл на все ветки

Файл описывает всё дерево целиком, а ветки содержат разные его части
(порт на Rust, исследование CFD, базовый C++-редактор). Отслеживаемая
копия при каждом переключении подменялась бы версией своей ветки, тогда
как нужна одна общая. Содержимое на всех ветках было идентично, так что
расхождений открепление не теряет.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01B9Gcr11JJJyf8NnWsXjDzZ

.gitignore

@@ -69,3 +69,6 @@ docs/theory/2d_solver/bench/out/
 .DS_Store
 Thumbs.db
 desktop.ini
+
+# Общие заметки Claude Code: один файл на все ветки, вне версионирования
+/CLAUDE.md

CLAUDE.md

@@ -1,202 +0,0 @@
-# CLAUDE.md
-
-This file provides guidance to Claude Code (claude.ai/code) when working with code in this repository.
-
-## Project Overview
-
-**SimVulcan** is a minimal **Vulkan 1.3 / C++20** 3D model editor. It renders a
-3D editor space — three reference grid planes (XY, XZ, YZ) through the origin plus
-coloured X/Y/Z axes — viewed through an orbit camera, and loads/displays `.obj`
-models with selectable display modes (solid, wireframe, solid + wireframe).
-
-The C++ side runs no simulation: it is a focused rendering skeleton with an ImGui
-interface (a Viewport control panel and a Mesh load panel). The repository *also*
-carries an unrelated Python research prototype under `docs/theory/` — see
-[Research prototype](#research-prototype-docstheory) below; it is not part of the
-CMake build.
-
-## Build
-
-Prerequisites: **Vulkan SDK 1.3.290+** (provides `glslangValidator` for offline
-shader compilation), **CMake 3.26+**, **Ninja**, a C++20 compiler (MSVC 19.36+,
-gcc 11+, clang 14+). On macOS, Vulkan is via MoltenVK (Apple Silicon only).
-
-```sh
-cmake --preset windows-msvc-release
-cmake --build --preset windows-msvc-release
-```
-
-Presets: `windows-msvc-debug`, `windows-msvc-release`, `linux-gcc-release`,
-`linux-clang-release`, `macos-arm64-release` (all Ninja, one dir per preset under
-`build/<presetName>/`). The first configure fetches dependencies via FetchContent
-(GLFW, GLM, volk, vk-bootstrap, VulkanMemoryAllocator, Dear ImGui, spdlog,
-tinyobjloader, Catch2) and needs network access. `VK_NO_PROTOTYPES` is set
-project-wide; Vulkan entry points load through **volk**.
-
-Warnings come from `simv_set_warnings` (`/W4 /permissive-`, or
-`-Wall -Wextra -Wpedantic -Wshadow -Wold-style-cast …`); they are **not** errors.
-`cmake/Sanitizers.cmake` defines `simv_enable_sanitizers` (Debug-only ASan/UBSan)
-but no target currently calls it — wire it in manually when chasing memory bugs.
-
-## Running
-
-The executable resolves SPIR-V relative to the working directory (`FindSpvPath`
-probes `spirv/<rel>` and `current_path()/spirv/<rel>`). The `SimVulcan` POST_BUILD
-step copies the compiled `spirv/` tree and `assets/` next to the executable, so
-**run from the executable's own directory** (`build/<preset>/src/app/`). `main.cpp`
-also probes several `../` ancestors for `assets/meshes`. Meshes load at runtime
-through the ImGui Mesh panel.
-
-Running writes two files into the CWD: `pipeline_cache.bin` (serialised
-`VkPipelineCache`, reloaded on the next start) and ImGui's `imgui.ini`. Both are
-disposable — delete them if pipeline creation or the panel layout misbehaves.
-
-`ContextOptions::enableValidation` / `enableDebugUtils` default to **true** in
-every build config, so the Khronos validation layer is requested even in Release;
-messages (error + warning severity) go through spdlog.
-
-`run.bat` at the repo root is **stale**: it launches
-`build/vs2022/src/app/Release/SimVulcan.exe`, a path the Ninja presets never
-produce. Do not point users at it without fixing the path first.
-
-## Tests
-
-Catch2 unit tests, pure CPU/math (mesh bounds + welding). No GPU required.
-
-```sh
-cmake --build --preset windows-msvc-debug --target simv_tests
-ctest --preset windows-msvc-debug
-```
-
-`windows-msvc-debug` is the only preset with a `testPreset` — for the other
-configs, invoke `ctest` in `build/<preset>/` directly.
-
-Single test / subset — either through CTest (each `TEST_CASE` is registered
-individually by `catch_discover_tests`):
-
-```sh
-ctest --preset windows-msvc-debug -R "WeldVertices" -V
-```
-
-or by running the binary with a Catch2 name or tag filter:
-
-```sh
-./build/windows-msvc-debug/tests/simv_tests.exe "[decimator]"
-./build/windows-msvc-debug/tests/simv_tests.exe --list-tests
-```
-
-## Architecture
-
-### Library / target map
-
-```
-SimVulcan (exe) → simv_core, simv_vk, simv_mesh, simv_editor
-simv_core   → simv_vk            (App owns Window + vk::Renderer; no Vulkan calls)
-simv_editor → simv_core, simv_mesh   (Camera, input, ImGui panels; no Vulkan)
-simv_mesh   → simv_core              (CPU mesh only; no Vulkan)
-simv_vk     → third-party (volk, vk-bootstrap, VMA, GLFW, GLM, spdlog, imgui)
-simv_shaders → glslangValidator (GLSL → SPIR-V)
-```
-
-Namespaces follow directories: `simv::core`, `simv::vk`, `simv::mesh`,
-`simv::editor`. Each library exports `src/` as its include root, so includes are
-written module-qualified (`#include "mesh/Mesh.h"`, `#include "vk/Renderer.h"`).
-
-### Frame loop
-
-`main.cpp` creates `core::App`, which owns the `core::Window` and a
-`vk::Renderer`, then runs the loop. Each frame `Renderer::DrawFrame` calls the UI
-callback (between ImGui NewFrame/Render), then records the scene: grid + mesh into
-one dynamic-rendering pass with a depth attachment, followed by ImGui, then
-presents (2 frames in flight, sync2 submits).
-
-`main.cpp` wires the editor via the UI callback: it draws the panels, applies
-mouse input to the `editor::Camera`, and pushes the resulting state into the
-renderer (`SetViewProj`, `SetRenderMode`, `SetGridVisible`). Mesh loads go through
-`MeshLoadPanel`'s callback → `Renderer::SetMeshCpu`.
-
-### Mesh load path
-
-`MeshLoadPanel` lists `*.obj` in the mesh directory and kicks off
-`mesh::LoadObjAsync` (worker thread, `std::future`). The future is **drained on
-the main thread** at the top of `MeshLoadPanel::Draw`, so the `OnLoaded` callback
-— and therefore the GPU upload — always runs on the render thread. Loader
-exceptions surface as the panel's status string.
-
-`main.cpp`'s `OnLoaded` welds the mesh (`WeldVertices`, tolerance `1e-4`), logs
-the counts, uploads via `SetMeshCpu`, and reframes the camera to the bbox radius.
-Despite the file name, `mesh/MeshDecimator.h` implements **only** spatial-hash
-vertex welding (collapse near-duplicates, drop degenerate triangles, recompute
-bounds) — there is no LOD/decimation.
-
-### Non-obvious invariants (read before editing)
-
-- **All Vulkan lives in `src/vk/`.** `core/`, `mesh/`, `editor/` and `app/` make no
-  Vulkan API calls. `vk::Renderer`'s public header is deliberately Vulkan-free
-  (pImpl + glm/`RenderMode` only) so `App` can own it without pulling in volk. Keep
-  it that way — do not leak `Vk*` types into the public interfaces of those modules.
-- **Single render pass with depth.** Scene (grid + mesh) and ImGui draw into one
-  `vkCmdBeginRendering` pass that has both a colour and a `D32_SFLOAT` depth
-  attachment. The ImGui backend is initialised with `depthAttachmentFormat` set so
-  its pipeline matches the pass; the depth buffer is recreated with the swapchain.
-- **Wireframe needs `fillModeNonSolid`.** `MeshRenderer` builds a fill pipeline and
-  a `VK_POLYGON_MODE_LINE` pipeline; the line pipeline uses a small depth bias so
-  the overlay sits on top of the fill. The device feature is requested in `Context`.
-  Culling is off (`VK_CULL_MODE_NONE`) — loaded models may have mixed winding.
-- **Camera is fixed on the origin.** `editor::Camera` orbits (yaw/pitch/distance)
-  the world origin; loaded models are recentred there via a translate-only model
-  matrix. Projection uses Vulkan clip space (`GLM_FORCE_DEPTH_ZERO_TO_ONE` + Y flip,
-  isolated to `Camera.cpp`).
-- **Buffers.** `vk::GpuMesh` owns the model's vertex/index buffers (staged upload on
-  the transfer queue). `GridRenderer` builds a static host-visible line buffer once.
-  Both scene renderers use only a push constant — no descriptor sets.
-- **Push-constant layout is a cross-file contract.** `MeshRenderer.cpp`'s anonymous
-  `MeshPC { mat4 mvp; vec4 color; }` must stay byte-identical to the
-  `push_constant` block in `mesh.vert`/`mesh.frag`; `color.a` is a *flag*, not
-  alpha (1 = flat-shaded, 0 = constant colour for wireframe). `GridRenderer` pushes
-  a bare `mat4` (vertex stage only). Change either side and you must change both.
-- **Swapchain recreation rebuilds sync objects.** `RecreateSwapDependent` recreates
-  the per-frame `imageAvailable` semaphores (a failed acquire can leave one
-  signalled) *and* the per-swapchain-image `renderFinished` semaphores (the image
-  count may change), then re-ensures both pipelines. Keep that ordering if you
-  touch resize handling.
-- **Mesh upload stalls the device.** `SetMeshCpu`/`ClearMesh` call `WaitIdle`
-  before touching `GpuMesh` — acceptable because loads are rare; do not copy that
-  pattern into per-frame paths.
-
-### Shaders
-
-GLSL under `shaders/editor/` (`mesh.{vert,frag}`, `grid.{vert,frag}`). `simv_shaders`
-compiles each to `build/<preset>/spirv/editor/<name>.spv` targeting `vulkan1.3`,
-with `shaders/` as the `-I` root (so `#include "common/foo.glsl"` would resolve).
-Adding a shader under that globbed dir is picked up automatically
-(`CONFIGURE_DEPENDS`). `mesh.frag` reconstructs a flat normal from screen-space
-derivatives, so the vertex stream carries only positions (tight `float3`, one
-binding, one attribute).
-
-## Logging
-
-`simv::core::Logger` (spdlog-backed, singleton) with category-based throttling.
-Log via `LogFmt(LogCategory, LogLevel, fmt, args...)`. Categories: `Core`,
-`Vulkan`, `MeshIO`, `UI`, `Test`. Per-category level, throttle interval and
-on/off are settable at runtime (`SetMinLevel` / `SetThrottle` / `SetEnabled`).
-Some low-level Vulkan code still calls `spdlog::` directly.
-
-## Research prototype (`docs/theory/`)
-
-Separate from the C++ application and from the CMake build: a Python **Lattice
-Boltzmann (D2Q9, KBC-N1)** research codebase — cylinder flow with a ×2 nested AMR
-patch and SDF+Bouzidi boundaries. **GPU/CuPy only, no CPU fallback.** Its
-documentation and the running experiment log are in Russian.
-
-- `docs/theory/solver_2x_sdf/` — the component-split solver (`run.py`,
-  `run_blockage.py`, `run_factors.py`); its `README.md` is the authoritative
-  status/experiment log and maps every file to the physics it owns.
-- `docs/theory/demos_gpu/` — demo runs that produce the notebook's figures/GIFs;
-  they import physics from `solver_2x_sdf`, never duplicate it.
-- `docs/theory/kbc_lbm.ipynb` — the write-up; it embeds pre-rendered artefacts and
-  does not execute the simulations.
-- `docs/origins/` — source PDFs behind the theory.
-
-Do not fold this into the CMake build, and do not treat it as dead code — it is
-active research with a documented experiment history.

docs/theory/2d_solver/Dockerfile

@@ -1,15 +1,23 @@
-# Образ для развёртывания решателя и валидационной кампании на сервере с NVIDIA.
+# Образ решателя и валидационной кампании.
 #
-# Сборка и запуск:
+# Поддерживает два способа добраться до видеокарты, и выбирать между ними не нужно —
+# доступным окажется ровно один:
+#
+#  1. НАСТОЯЩИЙ LINUX-СЕРВЕР с драйвером NVIDIA. Vulkan-ICD подкладывает внутрь
+#     NVIDIA Container Toolkit, и только если в NVIDIA_DRIVER_CAPABILITIES есть
+#     `graphics` — с одним `compute` wgpu не увидит ни одного адаптера. Запуск через
+#     docker-compose.server.yml.
+#
+#  2. WSL2. Драйвер NVIDIA для Linux там отсутствует как класс: карта отдаётся через
+#     /dev/dxg по протоколу WDDM, а нативный libGLX_nvidia про него не знает и
+#     enumerate возвращает ноль устройств. Зато с /dev/dxg умеет говорить dzn (Dozen) —
+#     драйвер Mesa, транслирующий Vulkan в D3D12. Он в образе и лежит. NVIDIA Container
+#     Toolkit при этом НЕ НУЖЕН: нужен проброс устройства и монтирование /usr/lib/wsl,
+#     где Microsoft держит libd3d12.so. Запуск через docker-compose.wsl.yml.
+#
+# Сборка и быстрая проверка:
 #   docker build -t kbc2d docs/theory/2d_solver
-#   docker run --rm --gpus all kbc2d --calibrate            # проверить, что GPU виден
+#   docker compose -f docker-compose.wsl.yml --profile check run --rm vulkan
-#   docker run -d  --gpus all -v "$PWD/out:/work/bench/out" --name kbc2d kbc2d --resume
-#
-# ГЛАВНАЯ ТОНКОСТЬ — Vulkan в контейнере. NVIDIA Container Toolkit подкладывает внутрь
-# Vulkan-ICD (nvidia_icd.json) только если в NVIDIA_DRIVER_CAPABILITIES есть `graphics`.
-# С одним `compute` wgpu не увидит НИ ОДНОГО адаптера и решатель откажется стартовать с
-# --backend gpu. Поэтому capabilities прописаны в образе, а vulkan-tools положен внутрь,
-# чтобы первым делом можно было выполнить `vulkaninfo` и убедиться, что карта видна.
 
 # ── сборка ───────────────────────────────────────────────────────────────────
 # Версия тулчейна закреплена: та же, на которой решатель собирался и проверялся.
@@ -18,24 +26,44 @@ WORKDIR /src
 
 # Сначала только манифесты — тогда слой с зависимостями переиспользуется, пока они не менялись.
 COPY Cargo.toml Cargo.lock ./
-RUN mkdir src && echo 'fn main() {}' > src/main.rs \
+RUN mkdir src && echo 'fn main() {}' > src/main.rs  && cargo build --release --locked 2>/dev/null || true  && rm -rf src
- && cargo build --release --locked 2>/dev/null || true \
- && rm -rf src
 
 COPY src ./src
 # touch нужен, чтобы cargo не спутал новые исходники с заглушкой из слоя выше
 RUN touch src/main.rs && cargo build --release --locked
 
 # ── исполнение ───────────────────────────────────────────────────────────────
-FROM debian:bookworm-slim
+# База Arch, а не debian-slim, по одной причине: dzn собирают не все дистрибутивы.
+# В mesa-vulkan-drivers Debian и Ubuntu его нет (проверено по списку файлов пакета),
+# в Arch он лежит отдельным пакетом vulkan-dzn той же версии Mesa, что и на хосте WSL.
+FROM archlinux:base
+
+ARG VERSION=dev
+ARG REVISION=unknown
 LABEL org.opencontainers.image.title="kbc2d"
 LABEL org.opencontainers.image.description="Двумерный решатель LBM D2Q9 с энтропийным столкновением KBC"
+LABEL org.opencontainers.image.version="${VERSION}"
+LABEL org.opencontainers.image.revision="${REVISION}"
+LABEL org.opencontainers.image.source="https://gitea.arseniev.info/NotBigGhost/TurbulenceCAD"
 
-RUN apt-get update && apt-get install -y --no-install-recommends \
+# Справка, руководства и переводы занимают в базе Arch заметно больше, чем сам
-        libvulkan1 vulkan-tools python3 ca-certificates \
+# драйвер, а в контейнере не нужны никому. NoExtract прописывается ДО установки.
- && rm -rf /var/lib/apt/lists/*
+RUN { echo 'NoExtract = usr/share/man/*'; \
+      echo 'NoExtract = usr/share/doc/*'; \
+      echo 'NoExtract = usr/share/info/*'; \
+      echo 'NoExtract = usr/share/locale/*'; \
+      echo 'NoExtract = usr/share/i18n/*'; } >> /etc/pacman.conf \
+ && pacman -Syu --noconfirm --needed \
+        vulkan-icd-loader vulkan-dzn vulkan-tools python \
+ && pacman -Scc --noconfirm \
+ && rm -rf /var/cache/pacman/pkg/* /var/lib/pacman/sync/* \
+           /usr/share/man /usr/share/doc /usr/share/info /usr/share/locale
 
-# Без graphics в списке возможностей Vulkan-ICD внутрь не попадёт — см. шапку.
+# Проверка на этапе сборки: без dzn образ бесполезен для WSL, и узнать об этом надо
+# здесь, а не через час после старта кампании.
+RUN ls /usr/share/vulkan/icd.d/ && ls /usr/share/vulkan/icd.d/dzn_icd*.json > /dev/null
+
+# Для пути 1 (настоящий сервер). На WSL эти переменные ни на что не влияют.
 ENV NVIDIA_VISIBLE_DEVICES=all
 ENV NVIDIA_DRIVER_CAPABILITIES=compute,utility,graphics
 

docs/theory/2d_solver/README.md

@@ -492,14 +492,37 @@ python preflight.py                # каждый сценарий старту
 
 ## Развёртывание (Docker)
 
+Собранный образ опубликован: **`notbigghost/kbc2d:1.2.0`** (он же `latest`, платформа `linux/amd64`).
+Исходники на сервере не нужны — достаточно перенести туда один файл
+`docker-compose.server.yml`:
+
+```sh
+mkdir -p ~/kbc2d && cd ~/kbc2d           # сюда же ляжет ./out с результатами
+# перенести docker-compose.server.yml
+
+docker compose -f docker-compose.server.yml --profile check run --rm vulkan     # карта видна?
+docker compose -f docker-compose.server.yml --profile check run --rm preflight  # сценарии стартуют?
+docker compose -f docker-compose.server.yml --profile check run --rm calibrate  # сколько MLUPS?
+docker compose -f docker-compose.server.yml up -d                               # кампания
+docker compose -f docker-compose.server.yml logs -f
+```
+
+Порядок именно такой: узнать, что карта не видна, лучше через минуту, чем через час. Замеренные
+`--calibrate` числа подставляются переменными `KBC2D_GPU_MLUPS` / `KBC2D_CPU_MLUPS` — только на
+оценки в часах, на счёт они не влияют.
+
+Собрать образ самому (`docker-compose.yml` рядом делает то же самое с `build:`):
+
 ```sh
 docker build -t kbc2d docs/theory/2d_solver
-docker run --rm --gpus all kbc2d --calibrate                        # проверить, что GPU виден
+docker run --rm --gpus all kbc2d --calibrate
-docker run -d --gpus all -v "$PWD/out:/work/bench/out" kbc2d --resume
 ```
 
 `ENTRYPOINT` — драйвер кампании, `CMD` по умолчанию `--dry-run`: случайный `docker run` покажет
-смету и выйдет, а не запустит сточасовую задачу.
+смету и выйдет, а не запустит сточасовую задачу. В серверном compose политика перезапуска —
+`on-failure`, а не `unless-stopped`: кампания завершается штатно с кодом 0, и «перезапускать
+всегда» крутило бы контейнер вхолостую по кругу, тогда как падение или перезагрузку хоста
+`on-failure` подхватывает, а `--resume` продолжает с места.
 
 **Главная тонкость — Vulkan внутри контейнера.** NVIDIA Container Toolkit подкладывает
 Vulkan-ICD (`nvidia_icd.json`) только если в `NVIDIA_DRIVER_CAPABILITIES` есть `graphics`;
@@ -510,8 +533,67 @@ Vulkan-ICD (`nvidia_icd.json`) только если в `NVIDIA_DRIVER_CAPABILIT
 Проверено локально: образ собирается, кампания внутри него проходит смоук с монтированием
 результатов на хост, физика совпадает с хостовой до последней цифры (ошибка Тейлора–Грина
 9.829e-3 при N=64 и 2.401e-3 при N=128 — те же значения, что вне контейнера), а `--backend gpu`
-без проброшенной карты отказывает явным сообщением, а не считает молча. GPU-путь в контейнере
+без проброшенной карты отказывает явным сообщением, а не считает молча.
-проверяется только на машине с картой.
+
+### WSL2 — отдельный путь
+
+В WSL2 всё вышеописанное не работает, и не из-за настроек. **Драйвера Vulkan для Linux у
+NVIDIA там нет**: карта отдаётся через `/dev/dxg` по протоколу WDDM, нативный
+`libGLX_nvidia` про него не знает и перечисляет ноль устройств. Container Toolkit
+подкладывать внутрь нечего, отсюда `could not select device driver "nvidia"`.
+
+Работает другое: **dzn** (Dozen) — драйвер Mesa, транслирующий Vulkan в D3D12, он умеет
+говорить с `/dev/dxg` напрямую. Он положен в образ (ради него база сменена с
+`debian:bookworm-slim` на `archlinux:base`: в пакетах Mesa у Debian и Ubuntu dzn не
+собирают). NVIDIA-runtime для этого пути не нужен вовсе — нужны проброс устройства и
+монтирование `/usr/lib/wsl`. Запуск через `docker-compose.wsl.yml`, подробности в
+[`bench/README.md`](bench/README.md).
+
+Три вещи, которые надо знать про этот путь.
+
+**wgpu по умолчанию прячет несоответствующие адаптеры.** dzn сообщает о себе
+`conformanceVersion = 0.0.0.0`, и wgpu молча его отбрасывает — решатель докладывает, что
+GPU не найден. Согласие даётся явно, переменной `WGPU_ALLOW_UNDERLYING_NONCOMPLIANT_ADAPTER=1`;
+чтобы она вообще читалась, в `gpu.rs` при создании инстанса стоит
+`InstanceFlags::from_build_config().with_env()`. Поведение по умолчанию не изменилось: без
+переменной несоответствующие адаптеры по-прежнему скрыты.
+
+**У dzn мал предел на размер одной привязки — 128 МиБ** против гигабайтов у нативных
+драйверов. Массив популяций занимает `nx · ny · 9 · 4` байта, поэтому при одной общей
+привязке потолок выходил 3.73 млн узлов, и сетки от 2048×2048 не запускались вовсе. Но
+ограничена именно привязка, а не буфер (`max_buffer_size` у dzn 2047 МиБ), так что тот же
+буфер теперь показывается **девятью привязками по одному направлению**: потолок
+поднимается до 33.5 млн узлов. Вариант выбирается по возможностям адаптера; где предела
+нет, собирается прежний общий, без `switch` в аксессорах. Оба дают одинаковые числа
+(Cd 2.39486 в обоих), раздельный стоит 5.7% пропускной способности. Принудительно
+включается переменной `KBC2D_SPLIT_POPULATIONS=1` — она нужна для сверки двух вариантов
+на одной карте.
+
+**Точность трансляция не портит.** Замерено на Intel Iris Xe одним и тем же прогоном
+(`bench/parity.py`, цилиндр Re=20, 8000 шагов):
+
+| | Cd | energy_end (Тейлор–Грин) |
+|---|---|---|
+| CPU, f64 | 2.39490 | 5.74754e-05 |
+| нативный Vulkan, f32 | 2.39486 | 5.74813e-05 |
+| dzn в контейнере, f32 | 2.39486 | 5.74810e-05 |
+
+Числа dzn и нативного драйвера сходятся до 5–6 значащих цифр, и разница между ними меньше,
+чем между любым из них и f64. Считает трансляция то же самое.
+
+**Скорость она портит, но тем меньше, чем крупнее сетка** — плата почти вся приходится на
+трансляцию вызова, а не счёта:
+
+| сетка | узлов | нативно | dzn в контейнере | плата |
+|---|---|---|---|---|
+| 320×192 | 61 тыс. | 188.8 MLUPS | 44.5 MLUPS | 4.2× |
+| 960×480 | 461 тыс. | 125.3 MLUPS | 86.0 MLUPS | 1.46× |
+| 1920×960 | 1.84 млн | 128.7 MLUPS | 98.7 MLUPS | 1.30× |
+
+Для кампании это решает дело: 95.1% её стоимости приходится на сетки крупнее 600 тыс.
+узлов, а на сетки мельче 150 тыс. — 0.0%. Ожидаемое удорожание всей кампании в контейнере —
+около трети, а не в разы. Проверяется профилем `calibrate`, который меряет в том числе
+1920×960.
 
 ## Дальше
 

docs/theory/2d_solver/bench/README.md

@@ -6,23 +6,158 @@
 
 ## Быстрый старт на сервере
 
-```sh
+Образ опубликован, собирать ничего не нужно: **`notbigghost/kbc2d:1.2.0`**. Исходники на
-docker build -t kbc2d docs/theory/2d_solver
+сервере тоже не нужны — переносится один файл `docker-compose.server.yml`.
-docker run --rm --gpus all --entrypoint python3 kbc2d preflight.py   # сценарии стартуют
-docker run --rm --gpus all kbc2d --calibrate                          # GPU виден
-docker run -d --gpus all -v "$PWD/out:/work/bench/out" --name kbc2d kbc2d --resume
-docker logs -f kbc2d
-```
-
-Первым делом стоит убедиться, что Vulkan внутри контейнера действительно видит карту:
 
 ```sh
-docker run --rm --gpus all --entrypoint vulkaninfo kbc2d --summary | head -20
+mkdir -p ~/kbc2d && cd ~/kbc2d          # сюда же ляжет ./out с результатами
+# перенести сюда docker-compose.server.yml
+
+C=docker-compose.server.yml
+docker compose -f $C --profile check run --rm vulkan      # 1. карта видна из контейнера?
+docker compose -f $C --profile check run --rm preflight   # 2. все 115 сценариев стартуют?
+docker compose -f $C --profile check run --rm calibrate   # 3. сколько MLUPS на этой машине?
+docker compose -f $C --profile check run --rm plan        # 4. смета в часах по замеренному
+docker compose -f $C up -d                                # 5. кампания
+docker compose -f $C logs -f
 ```
 
-Если адаптеров ноль — почти наверняка дело в `NVIDIA_DRIVER_CAPABILITIES`: Container Toolkit
+Порядок не случайный: узнать, что карта не видна, лучше на первом шаге, чем через час счёта.
-подкладывает Vulkan-ICD только при наличии `graphics` в списке. В образе это прописано, но
+
-может быть переопределено снаружи.
+Замеренные калибровкой числа подставляются переменными окружения — они влияют только на
+оценки в часах, не на счёт:
+
+```sh
+KBC2D_GPU_MLUPS=1450 KBC2D_CPU_MLUPS=40 docker compose -f $C --profile check run --rm plan
+```
+
+**Если `vulkaninfo` показывает ноль адаптеров** — почти наверняка дело в
+`NVIDIA_DRIVER_CAPABILITIES`: Container Toolkit подкладывает Vulkan-ICD только при наличии
+`graphics` в списке. В образе и в compose это прописано, но может быть переопределено снаружи.
+Проверить, что хост вообще умеет отдавать карту:
+
+```sh
+docker run --rm --gpus all nvidia/cuda:12.4.0-base-ubuntu22.04 nvidia-smi
+```
+
+## Запуск в WSL2
+
+Отдельный путь, потому что в WSL2 **драйвера Vulkan для Linux у NVIDIA не существует**.
+Карта отдаётся через `/dev/dxg` по протоколу WDDM; нативный `libGLX_nvidia` про него не
+знает и возвращает ноль устройств — загрузчик его выбрасывает. Отсюда и `could not select
+device driver "nvidia"`: NVIDIA Container Toolkit тут не поможет, потому что подкладывать
+внутрь нечего.
+
+Зато с `/dev/dxg` умеет говорить **dzn** (Dozen) — драйвер Mesa, транслирующий Vulkan в
+D3D12. Он лежит в образе. NVIDIA-runtime при этом **не нужен вовсе**: достаточно проброса
+устройства и монтирования `/usr/lib/wsl`, где Microsoft держит `libd3d12.so`.
+
+```sh
+C=docker-compose.wsl.yml
+docker compose -f $C --profile check run --rm vulkan      # 1. карта видна?
+docker compose -f $C --profile check run --rm parity      # 2. считает ли она правильно?
+docker compose -f $C --profile check run --rm calibrate   # 3. и с какой скоростью?
+docker compose -f $C --profile check run --rm preflight   # 4. все сценарии стартуют?
+docker compose -f $C --profile check run --rm plan        # 5. смета в часах
+docker compose -f $C up -d                                # 6. кампания
+```
+
+Если образа нет ни локально, ни в реестре — `docker compose -f $C build`, доступ к Docker
+Hub не обязателен.
+
+### Шаг parity обязателен
+
+dzn сообщает о себе `conformanceVersion = 0.0.0.0`: набор тестов соответствия Vulkan он не
+проходил. wgpu по этой причине по умолчанию **прячет** такие адаптеры, и решатель сообщает,
+что GPU не найден. Согласие считать на непроверенном драйвере даётся явно — переменной
+`WGPU_ALLOW_UNDERLYING_NONCOMPLIANT_ADAPTER=1`, она прописана в `docker-compose.wsl.yml`.
+
+Раз соответствие не проверено вендором, его проверяем сами. `bench/parity.py` гоняет два
+коротких эталона на CPU в f64 и на GPU и сверяет числа: вихрь Тейлора–Грина (есть точное
+решение) и стационарное обтекание цилиндра при Re=20. Течения выбраны намеренно не
+хаотические — там расхождение f32 и f64 не нарастает, поэтому заметное различие означает
+проблему драйвера, а не разрядности.
+
+Замерено на Intel Iris Xe: **dzn совпадает с нативным драйвером той же карты до 5–6
+значащих цифр** (`energy_end` 3.02277e-07 против 3.02283e-07, `cd` 2.39486 против 2.39486).
+Точность трансляция не портит. Но это замер на Intel; на своей карте прогоните сами — две
+минуты.
+
+### Предел 128 МиБ на привязку и как он снят
+
+dzn объявляет `max_storage_buffer_binding_size = 128 МиБ` (2²⁷ байт), тогда как нативные
+драйверы дают гигабайты. Массив популяций занимает `nx · ny · 9 · 4` байта, так что при
+одной общей привязке потолок выходил **3.73 млн узлов** (примерно 1920×1920), и 14
+прогонов кампании из 115 падали на создании bind group:
+
+```
+Buffer binding 0 range 150994944 exceeds `max_*_buffer_binding_size` limit 134217728
+```
+
+На эти 14 приходилось 70.7% стоимости кампании — дорогие прогоны как раз крупносеточные.
+
+Существенно, что ограничена только **привязка**: `max_buffer_size` у dzn 2047 МиБ, то есть
+буфер держать разрешено, нельзя лишь показать шейдеру его целиком. Поэтому решатель теперь
+умеет показывать тот же буфер **девятью привязками**, по одному направлению в каждой.
+Потолок поднимается в девять раз — до 33.5 млн узлов, чего хватает всей кампании с запасом
+(самая крупная сетка в ней 4096×4096 — 16.8 млн узлов).
+
+Вариант выбирается сам, по `max_storage_buffer_binding_size` адаптера: где предела нет,
+собирается прежний общий вариант без `switch` в аксессорах. Проверить оба на одной карте
+можно переменной `KBC2D_SPLIT_POPULATIONS=1` — она включает раздельные привязки
+принудительно.
+
+Замерено на Intel Iris Xe, один драйвер, два варианта привязки:
+
+| | Cd | energy_end | MLUPS (цилиндр) |
+|---|---|---|---|
+| общая привязка | 2.39486 | 5.74813e-05 | 169.7 |
+| девять привязок | 2.39486 | 5.74813e-05 | 160.0 |
+
+Числа совпадают полностью; раздельный вариант стоит **5.7%** пропускной способности, и
+включается только там, где без него счёт вообще невозможен.
+
+### Чего трансляция стоит по скорости
+
+Плата есть, но она почти вся — накладные расходы на вызов, а не на счёт, и потому падает
+с ростом сетки. Замерено на Intel Iris Xe, один и тот же решатель:
+
+| сетка | узлов | нативный Vulkan | dzn в контейнере | плата |
+|---|---|---|---|---|
+| 320×192 | 61 тыс. | 188.8 MLUPS | 44.5 MLUPS | **4.2×** |
+| 960×480 | 461 тыс. | 125.3 MLUPS | 86.0 MLUPS | **1.46×** |
+| 1920×960 | 1.84 млн | 128.7 MLUPS | 98.7 MLUPS | **1.30×** |
+| 2048×2048 | 4.19 млн | 111.8 MLUPS | 67.4 MLUPS | **1.66×** |
+
+Последняя строка стоит особняком: там уже раздельные привязки (см. ниже), и в плату
+входит их `switch` в аксессорах. Ожидаемый разброс по кампании — от 1.3× до 1.7×.
+
+Каждый шаг решателя — несколько отправок в очередь; на мелкой сетке трансляция вызова
+стоит дороже самого счёта, на крупной размазывается. Для кампании это решающее
+обстоятельство, потому что дорогие прогоны в ней как раз крупные:
+
+| узлов в сетке | прогонов | доля стоимости кампании |
+|---|---|---|
+| < 150 тыс. | 18 | 0.0% |
+| 150–600 тыс. | 47 | 4.9% |
+| > 600 тыс. | 50 | **95.1%** |
+
+То есть 95% времени кампания проводит там, где плата 1.3–1.5×, и почти не бывает там, где
+она четырёхкратна. Ожидаемое удорожание по всей кампании — около трети: 90 часов
+превращаются примерно в 115–120, а не в 400.
+
+Проверьте на своей карте: профиль `calibrate` меряет три сетки, и ориентироваться надо на
+**1920×960** — она представительна для кампании, а 240×120 показывает худший случай.
+
+## Своя сборка образа
+
+Если нужен образ из текущего состояния репозитория, а не опубликованный:
+
+```sh
+docker build -t kbc2d docs/theory/2d_solver   --build-arg VERSION=dev --build-arg REVISION=$(git rev-parse --short HEAD)
+```
+
+Рядом лежит `docker-compose.yml` — то же самое через `build:`, для локальной отладки.
 
 ## Без Docker
 

docs/theory/2d_solver/bench/parity.py

@@ -0,0 +1,188 @@
+#!/usr/bin/env python3
+"""Сверка GPU-пути с эталонным CPU-путём.
+
+Зачем. На настоящем сервере GPU-путь идёт через драйвер NVIDIA. В WSL такого драйвера
+для Linux нет, и единственный доступный Vulkan — dzn (Dozen), транслирующий вызовы в
+D3D12. Он честно сообщает о себе `conformanceVersion = 0.0.0.0`, то есть набор тестов
+соответствия не проходил. Арифметика f32 в обоих API описана одним и тем же IEEE 754,
+но точность трансцендентных функций (exp, log, sqrt — а они в энтропийном равновесии
+на каждом узле) стандартами задаётся с допуском в несколько ULP, и допуски эти разные.
+Плюс энтропийный стабилизатор сравнивает знаменатель с относительным порогом GREL=1e-8:
+достаточно мелкого расхождения, чтобы решение «вырожден или нет» поменялось.
+
+Поэтому пригодность dzn нельзя принимать на веру — её надо мерить. Скрипт гоняет два
+коротких эталона на CPU (f64) и на GPU и сверяет числа из summary.json.
+
+Прогоны подобраны намеренно НЕ хаотические: вихрь Тейлора–Грина затухает гладко и имеет
+точное решение, обтекание цилиндра при Re=20 стационарно. В таких течениях расхождение
+f32 и f64 не нарастает, поэтому любое заметное различие означает проблему драйвера, а не
+разрядности. На развитой дорожке (Re >= 100) сверять бессмысленно: там разойдётся любая
+пара запусков с разной арифметикой.
+
+Порог взят с запасом от факта: на нативном Vulkan разница Cd между CPU-f64 и GPU-f32
+измерена и составила 0.007%.
+
+  python3 parity.py                 # сверка
+  python3 parity.py --tol 2e-3      # свой порог
+  python3 parity.py --keep out/     # оставить сводки для разбора
+
+Код возврата 1, если хоть одно поле вышло за порог или прогон развалился.
+"""
+
+import argparse
+import json
+import os
+import shutil
+import subprocess
+import sys
+import tempfile
+
+HERE = os.path.dirname(os.path.abspath(__file__))
+BIN = os.environ.get("KBC2D_BIN") or os.path.join(
+    HERE, "..", "target", "release", "kbc2d" + (".exe" if os.name == "nt" else ""))
+
+# Поля сводки: "fields" сверяются с порогом, "report" только печатаются.
+#
+#  * energy_end, enstrophy_end — интегралы поля, ловят расхождение самого счёта;
+#  * cd — интегральная сила, накапливает ошибку по всей границе тела;
+#  * gamma_mean и degenerate_frac — поведение энтропийного стабилизатора, самый чуткий
+#    индикатор расхождений в exp/log;
+#  * rho_mean — сохранение массы, обязано совпадать почти побитово.
+#
+# Пороги откалиброваны по замеру на НАТИВНОМ драйвере Vulkan (Intel Iris Xe, Windows):
+# там же, где f64 и f32 сравниваются в отсутствие какой-либо трансляции, отклонения
+# составили cd 1.7e-5, rho_mean 1.5e-6, gamma_mean 6.5e-4, energy_end 3.4e-3 (вихрь
+# доведён до затухания). Пороги ниже взяты с запасом к этим числам, но заметно ниже
+# того, что дал бы по-настоящему сломанный драйвер.
+#
+# Прогоны намеренно НЕ доводятся до глубокого затухания: когда энергия падает до 1e-7,
+# f32 упирается в собственный шум, и сверять становится нечего — это свойство
+# разрядности, а не драйвера.
+CASES = [
+    {
+        "name": "Тейлор-Грин 128x128, 500 шагов",
+        "args": ["--case", "taylor-green", "--nx", "128", "--ny", "128",
+                 "--refine", "1", "--steps", "500", "--case-every", "100"],
+        # taylor_green_error намеренно НЕ в списке сверяемых, только в отчётных: это
+        # разность почти равных величин, и её относительное отклонение f32 от f64
+        # достигает десятков процентов на любом, в том числе нативном, драйвере.
+        # Замерено на нативном Vulkan: 0.0116 (f64) против 0.0218 (f32). Сверять по
+        # ней бессмысленно, а видеть её полезно.
+        "fields": {"energy_end": 1.0, "enstrophy_end": 1.0, "gamma_mean": 1.0},
+        "report": ["taylor_green_error"],
+    },
+    {
+        "name": "цилиндр Re=20, стационар, 8000 шагов",
+        "args": ["--shape", "cylinder", "--size", "16", "--nx", "320", "--ny", "192",
+                 "--re", "20", "--refine", "1", "--steps", "8000"],
+        "fields": {"cd": 1.0, "rho_mean": 0.05, "gamma_mean": 2.0,
+                   "degenerate_frac": 5.0},
+        "report": ["strouhal"],
+    },
+]
+
+
+def run(backend, args, summary_path):
+    """Один прогон. Возвращает разобранную сводку либо строку с ошибкой."""
+    cmd = [BIN] + args + ["--backend", backend, "--verbose", "quiet",
+                          "--report-every", "0", "--summary", summary_path]
+    proc = subprocess.run(cmd, capture_output=True, text=True, errors="replace")
+    if proc.returncode != 0:
+        tail = (proc.stderr or proc.stdout or "").strip().splitlines()
+        return None, "\n".join(tail[-6:]) or f"код возврата {proc.returncode}"
+    try:
+        with open(summary_path, encoding="utf-8") as f:
+            return json.load(f), None
+    except Exception as e:
+        return None, f"сводка не читается: {e}"
+
+
+def deviation(ref, got):
+    """Относительное отклонение; для околонулевых величин — абсолютное."""
+    if abs(ref) > 1e-12:
+        return abs(got - ref) / abs(ref)
+    return abs(got - ref)
+
+
+def main():
+    ap = argparse.ArgumentParser(description="Сверка GPU-пути с CPU-эталоном")
+    ap.add_argument("--tol", type=float, default=1e-3,
+                    help="базовый относительный порог (по умолчанию 1e-3)")
+    ap.add_argument("--keep", metavar="DIR", help="куда положить сводки прогонов")
+    args = ap.parse_args()
+
+    if not os.path.exists(BIN):
+        print(f"не найден бинарь решателя: {BIN}")
+        return 1
+
+    workdir = args.keep or tempfile.mkdtemp(prefix="kbc2d-parity-")
+    os.makedirs(workdir, exist_ok=True)
+    failures = []
+    speeds = []
+
+    for case in CASES:
+        print(f"\n=== {case['name']} ===")
+        summaries = {}
+        for backend in ("cpu", "gpu"):
+            path = os.path.join(workdir, f"{backend}.json")
+            data, err = run(backend, case["args"], path)
+            if err:
+                print(f"  {backend}: ПРОГОН НЕ УДАЛСЯ\n    {err}")
+                failures.append(f"{case['name']}: {backend} не отработал")
+                break
+            if data.get("blew_up"):
+                print(f"  {backend}: счёт РАЗВАЛИЛСЯ")
+                failures.append(f"{case['name']}: {backend} развалился")
+                break
+            summaries[backend] = data
+        if len(summaries) != 2:
+            continue
+
+        cpu, gpu = summaries["cpu"], summaries["gpu"]
+        m_cpu, m_gpu = cpu.get("mlups", 0.0), gpu.get("mlups", 0.0)
+        speeds.append((case["name"], m_cpu, m_gpu))
+        ratio = f", ускорение x{m_gpu / m_cpu:.1f}" if m_cpu > 0 else ""
+        print(f"  скорость: CPU {m_cpu:.1f} MLUPS, GPU {m_gpu:.1f} MLUPS{ratio}")
+
+        print(f"  {'поле':<22}{'CPU (f64)':>16}{'GPU':>16}{'откл.':>12}  порог")
+        for field in case.get("report", []):
+            if field in cpu and field in gpu:
+                ref, got = float(cpu[field]), float(gpu[field])
+                print(f"  {field:<22}{ref:>16.6g}{got:>16.6g}{deviation(ref, got):>11.2e}"
+                      f"  (только к сведению)")
+        for field, mult in case["fields"].items():
+            if field not in cpu or field not in gpu:
+                print(f"  {field:<22}{'нет в сводке':>16}")
+                continue
+            ref, got = float(cpu[field]), float(gpu[field])
+            dev = deviation(ref, got)
+            tol = args.tol * mult
+            mark = "" if dev <= tol else "   <-- ВЫШЕ ПОРОГА"
+            print(f"  {field:<22}{ref:>16.6g}{got:>16.6g}{dev:>11.2e}  {tol:.1e}{mark}")
+            if dev > tol:
+                failures.append(f"{case['name']}: {field} отклонилось на {dev:.2e} "
+                                f"при пороге {tol:.1e}")
+
+    print()
+    if speeds:
+        print("Скорость по прогонам:")
+        for name, m_cpu, m_gpu in speeds:
+            print(f"  {name}: CPU {m_cpu:.1f} MLUPS, GPU {m_gpu:.1f} MLUPS")
+        print()
+
+    if failures:
+        print("СВЕРКА НЕ ПРОЙДЕНА:")
+        for f in failures:
+            print(f"  * {f}")
+        print("\nGPU-путь на этом драйвере доверия не заслуживает — кампанию на нём "
+              "запускать нельзя.")
+    else:
+        print("Сверка пройдена: GPU считает то же, что CPU в двойной точности.")
+
+    if not args.keep:
+        shutil.rmtree(workdir, ignore_errors=True)
+    return 1 if failures else 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

docs/theory/2d_solver/bench/preflight.py

@@ -27,6 +27,39 @@ DROP = {"--gif", "--xt", "--case-csv", "--gif-field", "--gif-downsample", "--gif
         "--gif-every", "--gif-fps", "--series-every"}
 
 
+# Строки, которые ничего не объясняют и только вытесняют полезные.
+NOISE = (
+    "note: run with `RUST_BACKTRACE",
+    "note: Some details are omitted",
+    "WARNING: dzn is not a conformant",
+    "stack backtrace",
+    "Caused by:",
+)
+
+
+def explain(output):
+    """Достать из вывода упавшего прогона то, ради чего его вообще читают.
+
+    Раньше бралась последняя строка вывода — а у паники Rust последняя строка это
+    «note: run with RUST_BACKTRACE=1», то есть ровно ноль сведений о причине.
+    Полезное лежит либо в собственном сообщении решателя, либо на строке ПОСЛЕ
+    «panicked at»: там текст самой паники.
+    """
+    out = [l.rstrip() for l in output.splitlines() if l.strip()]
+    out = [l for l in out if not any(n in l for n in NOISE)]
+    if not out:
+        return "(без вывода)"
+    for l in out:
+        if l.lstrip().startswith("ОШИБКА"):
+            return l.strip()
+    for i, l in enumerate(out):
+        if "panicked at" in l:
+            rest = [x.strip() for x in out[i + 1:]]
+            return " | ".join(rest[:4]) if rest else l.strip()
+    return out[-1].strip()
+
+
+
 def main():
     ap = argparse.ArgumentParser()
     ap.add_argument("--group", help="только эта группа, например I")
@@ -64,14 +97,15 @@ def main():
         ok = p.returncode == 0
         print("." if ok else "X", end="", flush=True)
         if not ok:
-            tail = (p.stdout + p.stderr).decode("utf-8", "replace").strip().splitlines()
+            out = (p.stdout + p.stderr).decode("utf-8", "replace")
-            bad.append((r["id"], tail[-1] if tail else "(без вывода)"))
+            bad.append((r["id"], explain(out)))
         if i % 40 == 0:
             print(f" {i}/{len(runs)}", flush=True)
 
     print(f"\n\nпроверено сценариев: {len(runs)}, не запустились: {len(bad)}")
     for id_, msg in bad:
-        print(f"  {id_:<28} {msg}")
+        print(f"  {id_}")
+        print(f"      {msg}")
     return 1 if bad else 0
 
 

docs/theory/2d_solver/bench/run_campaign.ps1

@@ -30,7 +30,9 @@ if (-not (Test-Path $Scenarios)) { throw "не найден список сце
 
 if ($Calibrate) {
     "Калибровка на этой машине (три коротких прогона)…"
-    foreach ($c in @(@(240,120,'gpu'), @(960,480,'gpu'), @(480,240,'cpu'))) {
+    # 1920x960 добавлена не для красоты: на 95% стоимости кампании сетки крупнее
+    # 600 тыс. узлов, и оценивать по мелким - значит занижать пропускную способность.
+    foreach ($c in @(@(240,120,'gpu'), @(960,480,'gpu'), @(1920,960,'gpu'), @(480,240,'cpu'))) {
         $o = & $Bin --nx $c[0] --ny $c[1] --size 16 --refine 1 --steps 3000 `
                     --report-every 3000 --verbose full --backend $c[2] 2>$null
         $m = ($o | Select-String 'MLUPS' | Select-Object -Last 1)

docs/theory/2d_solver/bench/run_campaign.sh

@@ -48,7 +48,9 @@ command -v python3 >/dev/null 2>&1 && PY=python3 || PY=python
 # ── калибровка: три коротких прогона на месте, чтобы оценки в часах были не гаданием ──
 if [ "$CALIB" = 1 ]; then
   echo "Калибровка на этой машине (три коротких прогона)…"
-  for spec in "240 120 gpu" "960 480 gpu" "480 240 cpu"; do
+  # 1920x960 добавлена не для красоты: на 95% стоимости кампании сетки крупнее
+  # 600 тыс. узлов, и оценивать по мелким — значит занижать пропускную способность.
+  for spec in "240 120 gpu" "960 480 gpu" "1920 960 gpu" "480 240 cpu"; do
     set -- $spec
     m=$("$BIN" --nx "$1" --ny "$2" --size 16 --refine 1 --steps 3000 --report-every 3000 \
         --verbose full --backend "$3" 2>/dev/null | grep -o '[0-9.]* MLUPS' | tail -1)

docs/theory/2d_solver/docker-compose.server.yml

@@ -0,0 +1,97 @@
+# Развёртывание кампании на сервере с NVIDIA. Этот файл САМОДОСТАТОЧЕН: он тянет готовый образ
+# из реестра и исходников репозитория не требует. Скопировать на сервер достаточно его одного.
+#
+#   mkdir -p ~/kbc2d && cd ~/kbc2d
+#   curl -O <ссылка на этот файл>        # либо просто перенести файл руками
+#
+#   docker compose -f docker-compose.server.yml --profile check run --rm vulkan
+#   docker compose -f docker-compose.server.yml --profile check run --rm preflight
+#   docker compose -f docker-compose.server.yml --profile check run --rm calibrate
+#   docker compose -f docker-compose.server.yml --profile check run --rm plan
+#   docker compose -f docker-compose.server.yml up -d
+#   docker compose -f docker-compose.server.yml logs -f
+#
+# Порядок именно такой. Сначала vulkan: если карта не видна, кампания молча уйдёт считать
+# ничего — точнее, откажется стартовать на первом же прогоне, но узнать об этом через час
+# обиднее, чем через минуту. Потом preflight (каждый сценарий стартует на два шага), потом
+# calibrate (замер MLUPS этой машины), и только потом сама кампания.
+#
+# Результаты складываются в ./out на хосте — около 4.4 ГБ гифок и рядов за полный проход.
+
+name: kbc2d
+
+# ── общая часть всех сервисов ────────────────────────────────────────────────
+x-kbc2d: &kbc2d
+  image: notbigghost/kbc2d:1.2.0
+  pull_policy: missing
+  volumes:
+    - ./out:/work/bench/out
+  environment:
+    # ГЛАВНОЕ МЕСТО ВСЕГО ФАЙЛА. NVIDIA Container Toolkit подкладывает внутрь Vulkan-ICD
+    # (nvidia_icd.json) только если в этом списке есть `graphics`. С одним `compute` wgpu не
+    # увидит НИ ОДНОГО адаптера, и решатель откажется стартовать с --backend gpu. В образе
+    # значение уже прописано, здесь оно продублировано явно — чтобы его было видно тому, кто
+    # читает compose, а не Dockerfile.
+    NVIDIA_DRIVER_CAPABILITIES: compute,utility,graphics
+    NVIDIA_VISIBLE_DEVICES: all
+    # Оценки в часах драйвер считает из этих чисел. Умолчания взяты с другой машины —
+    # подставьте сюда то, что напечатает профиль calibrate.
+    KBC2D_GPU_MLUPS: ${KBC2D_GPU_MLUPS:-1200}
+    KBC2D_CPU_MLUPS: ${KBC2D_CPU_MLUPS:-22}
+  deploy:
+    resources:
+      reservations:
+        devices:
+          - driver: nvidia
+            count: all
+            capabilities: [gpu]
+
+services:
+  # ── сама кампания: единственный сервис, который поднимается по `up -d` ──────
+  campaign:
+    <<: *kbc2d
+    container_name: kbc2d-campaign
+    # --resume пропускает всё, у чего уже есть summary.json: перезапуск продолжает с места,
+    # а не начинает заново. Именно поэтому перезапуск здесь безопасен и дёшев.
+    command: ["--resume"]
+    # on-failure, а НЕ unless-stopped: кампания завершается штатно с кодом 0, и политика
+    # «перезапускать всегда» после её окончания крутила бы контейнер вхолостую по кругу.
+    # Падение (OOM, перезагрузка хоста) даёт ненулевой код и будет подхвачено.
+    restart: on-failure:5
+    stop_grace_period: 30s
+    # Девяносто часов вывода — это сотни мегабайт журнала; без ограничения он съест диск.
+    # Полные логи каждого прогона всё равно лежат в out/<id>/log.txt.
+    logging:
+      driver: json-file
+      options:
+        max-size: "50m"
+        max-file: "5"
+
+  # ── проверки перед запуском (профиль check, сами не поднимаются) ────────────
+  vulkan:
+    <<: *kbc2d
+    profiles: ["check"]
+    entrypoint: ["vulkaninfo"]
+    command: ["--summary"]
+
+  preflight:
+    <<: *kbc2d
+    profiles: ["check"]
+    entrypoint: ["python3"]
+    command: ["preflight.py"]
+
+  calibrate:
+    <<: *kbc2d
+    profiles: ["check"]
+    command: ["--calibrate"]
+
+  plan:
+    <<: *kbc2d
+    profiles: ["check"]
+    command: ["--dry-run"]
+# ── если результаты нужны не под root ────────────────────────────────────────
+# Контейнер работает от root, поэтому файлы в ./out окажутся root:root. Чтобы они
+# принадлежали вам, добавьте в x-kbc2d строку
+#     user: "${UID}:${GID}"
+# и запускайте как  UID=$(id -u) GID=$(id -g) docker compose -f … up -d
+# (каталог ./out при этом должен быть создан заранее и принадлежать вам).

docs/theory/2d_solver/docker-compose.wsl.yml

@@ -0,0 +1,114 @@
+# Запуск решателя и кампании в WSL2 — там, где драйвера NVIDIA для Linux не существует.
+#
+# ЧЕМ ЭТОТ ФАЙЛ ОТЛИЧАЕТСЯ ОТ docker-compose.server.yml. Тот рассчитан на настоящий
+# Linux-сервер: просит у Docker устройство через драйвер `nvidia`, а Vulkan-ICD внутрь
+# подкладывает NVIDIA Container Toolkit. В WSL2 этого сделать нельзя в принципе —
+# у NVIDIA там нет Linux-библиотеки Vulkan, карта отдаётся по протоколу WDDM через
+# /dev/dxg. Отсюда и ошибка `could not select device driver "nvidia"`.
+#
+# Здесь NVIDIA-runtime не запрашивается ВООБЩЕ. Нужно ровно две вещи:
+#   * проброс /dev/dxg — сама видеокарта;
+#   * монтирование /usr/lib/wsl — там Microsoft держит libd3d12.so и libdxcore.so.
+# Дальше работает dzn (Dozen) из образа: драйвер Mesa, транслирующий Vulkan в D3D12.
+# Ни nvidia-container-toolkit, ни nvidia-ctk, ни правки в daemon.json не требуются.
+#
+# ВАЖНО, ПРОЧТИТЕ ДО ЗАПУСКА КАМПАНИИ. dzn сообщает о себе conformanceVersion 0.0.0.0,
+# то есть набор тестов соответствия Vulkan не проходил. Поэтому порядок такой:
+#
+#   docker compose -f docker-compose.wsl.yml --profile check run --rm vulkan     # карта видна?
+#   docker compose -f docker-compose.wsl.yml --profile check run --rm parity     # числа те же?
+#   docker compose -f docker-compose.wsl.yml --profile check run --rm calibrate  # скорость какая?
+#   docker compose -f docker-compose.wsl.yml --profile check run --rm plan
+#   docker compose -f docker-compose.wsl.yml up -d
+#   docker compose -f docker-compose.wsl.yml logs -f
+#
+# Шаг parity — не формальность. Он гоняет вихрь Тейлора-Грина (есть точное решение) и
+# стационарное обтекание цилиндра на CPU в f64 и на GPU, и сверяет числа. Если он не
+# прошёл, кампанию запускать нельзя: результаты будет нечем защищать.
+#
+# Результаты складываются в ./out на хосте.
+
+name: kbc2d
+
+# ── общая часть всех сервисов ────────────────────────────────────────────────
+x-kbc2d: &kbc2d
+  image: notbigghost/kbc2d:1.2.0
+  # Контекст сборки — каталог с этим файлом. Если образа нет ни локально, ни в реестре,
+  # достаточно `docker compose -f docker-compose.wsl.yml build`: доступ к Docker Hub
+  # для запуска не обязателен.
+  build:
+    context: .
+    args:
+      VERSION: "1.2.0"
+  pull_policy: missing
+  devices:
+    - /dev/dxg:/dev/dxg
+  volumes:
+    - ./out:/work/bench/out
+    # Только на чтение: контейнеру нужны отсюда libd3d12.so, libd3d12core.so и
+    # libdxcore.so. Каталог наполняет сам WSL из C:\Windows\System32\lxss\lib.
+    - /usr/lib/wsl:/usr/lib/wsl:ro
+  environment:
+    # Без этого dzn не найдёт D3D12 и молча не даст ни одного адаптера.
+    LD_LIBRARY_PATH: /usr/lib/wsl/lib
+    # Осознанное согласие считать на драйвере, не прошедшем тесты соответствия Vulkan.
+    # Без этой переменной wgpu прячет dzn и решатель сообщает, что адаптер не найден.
+    # Прежде чем запускать кампанию, обязательно пройдите профиль parity.
+    WGPU_ALLOW_UNDERLYING_NONCOMPLIANT_ADAPTER: "1"
+    # Оценки в часах драйвер считает из этих чисел. Умолчания взяты с другой машины —
+    # подставьте сюда то, что напечатает профиль calibrate.
+    KBC2D_GPU_MLUPS: ${KBC2D_GPU_MLUPS:-1200}
+    KBC2D_CPU_MLUPS: ${KBC2D_CPU_MLUPS:-22}
+
+services:
+  # ── сама кампания: единственный сервис, который поднимается по `up -d` ──────
+  campaign:
+    <<: *kbc2d
+    container_name: kbc2d-campaign
+    # --resume пропускает всё, у чего уже есть summary.json: перезапуск продолжает с
+    # места, а не начинает заново.
+    command: ["--resume"]
+    # on-failure, а НЕ unless-stopped: кампания завершается штатно с кодом 0, и политика
+    # «перезапускать всегда» после её окончания крутила бы контейнер вхолостую.
+    restart: on-failure:5
+    stop_grace_period: 30s
+    logging:
+      driver: json-file
+      options:
+        max-size: "50m"
+        max-file: "5"
+
+  # ── проверки перед запуском (профиль check, сами не поднимаются) ────────────
+  vulkan:
+    <<: *kbc2d
+    profiles: ["check"]
+    entrypoint: ["vulkaninfo"]
+    command: ["--summary"]
+
+  parity:
+    <<: *kbc2d
+    profiles: ["check"]
+    entrypoint: ["python3"]
+    command: ["parity.py"]
+
+  preflight:
+    <<: *kbc2d
+    profiles: ["check"]
+    entrypoint: ["python3"]
+    command: ["preflight.py"]
+
+  calibrate:
+    <<: *kbc2d
+    profiles: ["check"]
+    command: ["--calibrate"]
+
+  plan:
+    <<: *kbc2d
+    profiles: ["check"]
+    command: ["--dry-run"]
+# ── если результаты нужны не под root ────────────────────────────────────────
+# Контейнер работает от root, поэтому файлы в ./out окажутся root:root. Чтобы они
+# принадлежали вам, добавьте в x-kbc2d строку
+#     user: "${UID}:${GID}"
+# и запускайте как  UID=$(id -u) GID=$(id -g) docker compose -f ... up -d
+# (каталог ./out при этом должен быть создан заранее и принадлежать вам).

docs/theory/2d_solver/src/gpu.rs

@@ -51,7 +51,9 @@ struct LevelParams {
     flags: u32,
     /// число граничных узлов (для моментных моделей стенки)
     nwall: u32,
-    _pad: [u32; 3],
+    /// шаг между направлениями в буфере популяций, во флоатах (см. dir_stride)
+    stride: u32,
+    _pad: [u32; 2],
 }
 
 #[repr(C)]
@@ -85,8 +87,9 @@ struct Substep {
 struct AmrParams {
     nghost: u32,
     nrestrict: u32,
-    ccount: u32,
+    /// шаги между направлениями в буферах популяций грубого и мелкого уровней
-    fcount: u32,
+    cstride: u32,
+    fstride: u32,
     r01: f32,
     rfc: f32,
     w: f32,
@@ -145,23 +148,22 @@ struct Results {
 
 const SHADER: &str = r#"
 // ───── структуры (обязаны совпадать с gpu.rs) ─────
-struct LevelParams { n:u32, nx:u32, ny:u32, nlinks:u32, bcx:f32, bcy:f32, probe_node:u32, flags:u32, nwall:u32, lp1:u32, lp2:u32, lp3:u32 };
+struct LevelParams { n:u32, nx:u32, ny:u32, nlinks:u32, bcx:f32, bcy:f32, probe_node:u32, flags:u32, nwall:u32, stride:u32, lp2:u32, lp3:u32 };
 // wall_mode: 0 — Bouzidi, 1 — Град, 2 — HRR, 3 — простой отскок
 struct Dyn { ux_in:f32, uy_in:f32, rho_out:f32, kbc_model:u32, outlet_extrap:u32, nparts:u32, refine:u32, collision:u32, slot:u32, wall_mode:u32, dp2:u32, dp3:u32 };
 struct Substep { idx:u32, p0:u32, p1:u32, p2:u32 };
-struct AmrParams { nghost:u32, nrestrict:u32, ccount:u32, fcount:u32, r01:f32, rfc:f32, w:f32, pad:f32 };
+struct AmrParams { nghost:u32, nrestrict:u32, cstride:u32, fstride:u32, r01:f32, rfc:f32, w:f32, pad:f32 };
 struct GLink { node:u32, far:u32, i:u32, ib:u32, kind:u32, q:f32, body:u32, p1:u32 };
 struct GGhost { fine:u32, c00:u32, c10:u32, c01:u32, c11:u32, tx:f32, ty:f32, p0:u32 };
 struct Partial { rho:f32, maxu:f32, gsum:f32, gmin:f32, gmax:f32, cnt:f32, degen:f32, xineg:f32 };
 
-@group(0) @binding(0) var<storage, read_write> f    : array<f32>;
-@group(0) @binding(1) var<storage, read_write> post : array<f32>;
 @group(0) @binding(2) var<storage, read_write> gam  : array<f32>;
 @group(0) @binding(3) var<storage, read>       solid: array<u32>;
 @group(0) @binding(4) var<storage, read>       links: array<GLink>;
 @group(0) @binding(5) var<storage, read>       beta : array<f32>;
 @group(0) @binding(6) var<storage, read_write> parts: array<Partial>;
 @group(0) @binding(7) var<uniform>             P    : LevelParams;
+//__POPULATIONS__
 
 @group(1) @binding(0) var<uniform>             D    : Dyn;
 @group(1) @binding(1) var<uniform>             S    : Substep;
@@ -242,8 +244,8 @@ fn shift9(d: array<f32,9>, model: u32) -> array<f32,9> {
     return s;
 }
 
-fn load9(base: ptr<function, array<f32,9>>, off: u32, n: u32) {
+fn load9(base: ptr<function, array<f32,9>>, off: u32) {
-    for (var i = 0u; i < 9u; i = i + 1u) { (*base)[i] = f[i*n + off]; }
+    for (var i = 0u; i < 9u; i = i + 1u) { (*base)[i] = fget(i, off); }
 }
 
 // Возвращает (пост-столкновительные популяции, gamma, признак вырождения)
@@ -324,15 +326,15 @@ fn k_collide(@builtin(global_invocation_id) gid: vec3<u32>,
     let nd = lin(gid, nwg);
     if (nd >= P.n) { return; }
     var fv: array<f32,9>;
-    load9(&fv, nd, P.n);
+    load9(&fv, nd);
     if (solid[nd] != 0u) {
         // внутри тела не считаем: популяции там фиктивны, Bouzidi всё равно их перекрывает
-        for (var i = 0u; i < 9u; i = i + 1u) { post[i*P.n + nd] = fv[i]; }
+        for (var i = 0u; i < 9u; i = i + 1u) { pset(i, nd, fv[i]); }
         gam[nd] = 2.0;
         return;
     }
     var r = collide9(fv, beta[nd % P.nx], D.collision, D.kbc_model);
-    for (var i = 0u; i < 9u; i = i + 1u) { post[i*P.n + nd] = r.fv[i]; }
+    for (var i = 0u; i < 9u; i = i + 1u) { pset(i, nd, r.fv[i]); }
     gam[nd] = r.gamma;
 }
 
@@ -350,7 +352,7 @@ fn k_stream(@builtin(global_invocation_id) gid: vec3<u32>,
     for (var i = 0u; i < 9u; i = i + 1u) {
         var xs = (x - cx[i]) % nxi; if (xs < 0) { xs = xs + nxi; }
         var ys = (y - cy[i]) % nyi; if (ys < 0) { ys = ys + nyi; }
-        f[i*P.n + nd] = post[i*P.n + u32(ys*nxi + xs)];
+        fset(i, nd, pget(i, u32(ys*nxi + xs)));
     }
 }
 
@@ -360,20 +362,20 @@ fn k_bouzidi(@builtin(global_invocation_id) gid: vec3<u32>,
     let k = lin(gid, nwg);
     if (k >= P.nlinks) { return; }
     let L = links[k];
-    let fi = post[L.i*P.n + L.node];
+    let fi = pget(L.i, L.node);
     var v = fi;
     // ступенчатая модель игнорирует долю пересечения: стенка ровно посередине между узлами
     if (D.wall_mode == 3u) {
-        f[L.ib*P.n + L.node] = fi;
+        fset(L.ib, L.node, fi);
         return;
     }
     if (L.kind == 0u) {                                  // q < 1/2, есть дальний жидкий сосед
-        v = 2.0*L.q*fi + (1.0 - 2.0*L.q)*post[L.i*P.n + L.far];
+        v = 2.0*L.q*fi + (1.0 - 2.0*L.q)*pget(L.i, L.far);
     } else if (L.kind == 1u) {                           // q >= 1/2
         let h = 1.0/(2.0*L.q);
-        v = h*fi + (1.0 - h)*post[L.ib*P.n + L.node];
+        v = h*fi + (1.0 - h)*pget(L.ib, L.node);
     }
-    f[L.ib*P.n + L.node] = v;
+    fset(L.ib, L.node, v);
 }
 
 @compute @workgroup_size(64)
@@ -382,13 +384,13 @@ fn k_walls(@builtin(global_invocation_id) gid: vec3<u32>,
     let x = lin(gid, nwg);
     if (x >= P.nx) { return; }
     // зеркальное отражение: касательный импульс сохраняется, нормальный заворачивается
-    f[2u*P.n + x] = f[4u*P.n + x];
+    fset(2u, x, fget(4u, x));
-    f[5u*P.n + x] = f[8u*P.n + x];
+    fset(5u, x, fget(8u, x));
-    f[6u*P.n + x] = f[7u*P.n + x];
+    fset(6u, x, fget(7u, x));
     let t = (P.ny - 1u)*P.nx + x;
-    f[4u*P.n + t] = f[2u*P.n + t];
+    fset(4u, t, fget(2u, t));
-    f[7u*P.n + t] = f[6u*P.n + t];
+    fset(7u, t, fget(6u, t));
-    f[8u*P.n + t] = f[5u*P.n + t];
+    fset(8u, t, fget(5u, t));
 }
 
 @compute @workgroup_size(64)
@@ -399,23 +401,23 @@ fn k_channel(@builtin(global_invocation_id) gid: vec3<u32>,
     // вход: скоростной Zou-He по ВСЕМУ столбцу, включая угловые узлы
     let a = y*P.nx;
     var fv: array<f32,9>;
-    load9(&fv, a, P.n);
+    load9(&fv, a);
     var gi = zou_he_inlet(fv, D.ux_in, D.uy_in);
-    for (var i = 0u; i < 9u; i = i + 1u) { f[i*P.n + a] = gi[i]; }
+    for (var i = 0u; i < 9u; i = i + 1u) { fset(i, a, gi[i]); }
     // выход: давление-Zou-He
     let b = y*P.nx + P.nx - 1u;
     var gv: array<f32,9>;
-    load9(&gv, b, P.n);
+    load9(&gv, b);
     var uyo = 0.0;
     if (D.outlet_extrap != 0u) {
         let c = y*P.nx + P.nx - 2u;
         var cv: array<f32,9>;
-        load9(&cv, c, P.n);
+        load9(&cv, c);
         let s = cv[0]+cv[1]+cv[2]+cv[3]+cv[4]+cv[5]+cv[6]+cv[7]+cv[8];
         uyo = ((cv[2]+cv[5]+cv[6]) - (cv[4]+cv[7]+cv[8])) / s;
     }
     var go = zou_he_outlet(gv, D.rho_out, uyo);
-    for (var i = 0u; i < 9u; i = i + 1u) { f[i*P.n + b] = go[i]; }
+    for (var i = 0u; i < 9u; i = i + 1u) { fset(i, b, go[i]); }
 }
 
 // сила и момент по GMEM, редукция одной рабочей группой
@@ -436,8 +438,8 @@ fn k_force(@builtin(local_invocation_id) lid: vec3<u32>) {
     loop {
         if (k >= P.nlinks) { break; }
         let L = links[k];
-        let fp = post[L.i*P.n + L.node];
+        let fp = pget(L.i, L.node);
-        let fb = f[L.ib*P.n + L.node];
+        let fb = fget(L.ib, L.node);
         let dfx = cx[L.i]*fp + cx[L.i]*fb;
         let dfy = cy[L.i]*fp + cy[L.i]*fb;
         // плечо до точки пересечения линка со стенкой, а не до узла
@@ -488,7 +490,7 @@ fn k_stats1(@builtin(global_invocation_id) gid: vec3<u32>,
     p.cnt = 0.0; p.degen = 0.0; p.xineg = 0.0;
     if (nd < P.n && solid[nd] == 0u) {
         var fv: array<f32,9>;
-        load9(&fv, nd, P.n);
+        load9(&fv, nd);
         let m = macros9(fv);
         p.rho = m.x;
         p.maxu = sqrt(m.y*m.y + m.z*m.z);
@@ -506,7 +508,7 @@ fn k_stats1(@builtin(global_invocation_id) gid: vec3<u32>,
     // зонд следа снимается с того уровня, на котором он лежит
     if ((P.flags & 2u) != 0u && nd == P.probe_node) {
         var fv: array<f32,9>;
-        load9(&fv, nd, P.n);
+        load9(&fv, nd);
         let m = macros9(fv);
         results[D.slot*SLOT + 3u] = m.z;
     }
@@ -626,7 +628,7 @@ fn k_stats2(@builtin(local_invocation_id) lid: vec3<u32>) {
 // хранилища прошлого шага.
 fn u_at_t(nd: u32) -> vec2<f32> {
     var ff: array<f32,9>;
-    for (var i = 0u; i < 9u; i = i + 1u) { ff[i] = post[i*P.n + nd]; }
+    for (var i = 0u; i < 9u; i = i + 1u) { ff[i] = pget(i, nd); }
     let m = macros9(ff);
     return vec2<f32>(m.y, m.z);
 }
@@ -733,15 +735,15 @@ fn k_moment_wall(@builtin(global_invocation_id) gid: vec3<u32>,
     var opp = array<u32,9>(0u, 3u, 4u, 1u, 2u, 7u, 8u, 5u, 6u);
     var rho = 0.0;
     for (var i = 0u; i < 9u; i = i + 1u) {
-        if (missing[i] == 1u) { rho = rho + post[opp[i]*P.n + node]; }
+        if (missing[i] == 1u) { rho = rho + pget(opp[i], node); }
-        else { rho = rho + f[i*P.n + node]; }
+        else { rho = rho + fget(i, node); }
     }
 
     let g4 = grad_u_at_t(node);
     var g = moment_wall9(rho, ux, uy, g4.x, g4.y, g4.z, g4.w,
                          beta[node % P.nx], select(0u, 1u, D.wall_mode == 2u));
     for (var i = 0u; i < 9u; i = i + 1u) {
-        if (missing[i] == 1u) { f[i*P.n + node] = g[i]; }
+        if (missing[i] == 1u) { fset(i, node, g[i]); }
     }
 }
 
@@ -764,7 +766,7 @@ fn k_ghost(@builtin(global_invocation_id) gid: vec3<u32>,
     var neq = array<f32,9>(0.0,0.0,0.0,0.0,0.0,0.0,0.0,0.0,0.0);
     for (var j = 0u; j < 4u; j = j + 1u) {
         var cf: array<f32,9>;
-        for (var i = 0u; i < 9u; i = i + 1u) { cf[i] = amr_src[i*A.ccount + cs[j]]; }
+        for (var i = 0u; i < 9u; i = i + 1u) { cf[i] = amr_src[i*A.cstride + cs[j]]; }
         let m = macros9(cf);
         rho = rho + ws[j]*m.x;
         ux  = ux  + ws[j]*m.y;
@@ -785,7 +787,7 @@ fn k_fill(@builtin(global_invocation_id) gid: vec3<u32>,
     let g = ghosts[k];
     // временная интерполяция рамки между состояниями L0 «до» и «после» шага
     for (var i = 0u; i < 9u; i = i + 1u) {
-        amr_dst[i*A.fcount + g.fine] = (1.0 - A.w)*gh_a[k*9u + i] + A.w*gh_b[k*9u + i];
+        amr_dst[i*A.fstride + g.fine] = (1.0 - A.w)*gh_a[k*9u + i] + A.w*gh_b[k*9u + i];
     }
 }
 
@@ -796,11 +798,11 @@ fn k_restrict(@builtin(global_invocation_id) gid: vec3<u32>,
     if (k >= A.nrestrict) { return; }
     let pr = rest[k];
     var ff: array<f32,9>;
-    for (var i = 0u; i < 9u; i = i + 1u) { ff[i] = amr_src[i*A.fcount + pr.y]; }
+    for (var i = 0u; i < 9u; i = i + 1u) { ff[i] = amr_src[i*A.fstride + pr.y]; }
     let m = macros9(ff);
     var fe = feq9(m.x, m.y, m.z);
     for (var i = 0u; i < 9u; i = i + 1u) {
-        amr_dst[i*A.ccount + pr.x] = fe[i] + A.rfc*(ff[i] - fe[i]);
+        amr_dst[i*A.cstride + pr.x] = fe[i] + A.rfc*(ff[i] - fe[i]);
     }
 }
 "#;
@@ -813,6 +815,8 @@ struct GpuLevel {
     nx: usize,
     ny: usize,
     n: usize,
+    /// шаг между направлениями в `f`; см. dir_stride
+    stride: usize,
     f: wgpu::Buffer,
     /// Стабилизатор γ поузлово — нужен для картинки; забирается с устройства как есть.
     gam: wgpu::Buffer,
@@ -833,12 +837,27 @@ impl GpuLevel {
 
 /// Переложить готовое стартовое поле (построенное общим кодом в `cpu::initial_field`)
 /// из AoS в раскладку SoA, которой пользуется GPU.
-fn to_soa(f: &[[R; math::Q]]) -> Vec<f32> {
+/// Выравнивание шага между направлениями, во ФЛОАТАХ.
+///
+/// Смещение каждой привязки обязано быть кратно `min_storage_buffer_offset_alignment`,
+/// а он равен 256 байтам у всех известных реализаций — это 64 значения f32. Шаг
+/// выравнивается всегда, даже когда привязка одна: так раскладка одна на оба варианта
+/// шейдера, и не надо помнить, какой из них сейчас собран. Перерасход памяти — меньше
+/// сотой доли процента.
+const DIR_ALIGN: usize = 64;
+
+/// Расстояние между началами соседних направлений в буфере популяций, во флоатах.
+fn dir_stride(n: usize) -> usize {
+    n.div_ceil(DIR_ALIGN) * DIR_ALIGN
+}
+
+fn to_soa(f: &[[R; math::Q]], stride: usize) -> Vec<f32> {
     let n = f.len();
-    let mut v = vec![0.0f32; 9 * n];
+    debug_assert!(stride >= n);
-    for (k, cell) in f.iter().enumerate() {
+    let mut v = vec![0.0f32; 9 * stride];
-        for i in 0..9 {
+    for (k, c) in f.iter().enumerate() {
-            v[i * n + k] = cell[i] as f32;
+        for i in 0..math::Q {
+            v[i * stride + k] = c[i] as f32;
         }
     }
     v
@@ -856,9 +875,11 @@ fn make_level(
     flags: u32,
     init_field: &[[R; math::Q]],
     wall_layout: &wgpu::BindGroupLayout,
+    split: bool,
 ) -> GpuLevel {
     let n = nx * ny;
-    let init = to_soa(init_field);
+    let stride = dir_stride(n);
+    let init = to_soa(init_field, stride);
     let mkf = |label: &str| {
         device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
             label: Some(label),
@@ -930,26 +951,41 @@ fn make_level(
             probe_node,
             flags,
             nwall,
-            _pad: [0; 3],
+            stride: stride as u32,
+            _pad: [0; 2],
         }),
         usage: wgpu::BufferUsages::UNIFORM,
     });
 
-    let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
+    // Общая часть связки одинакова в обоих вариантах; различаются только популяции.
-        label: Some("level"),
+    let mut entries = vec![
-        layout,
-        entries: &[
-            bind(0, &f),
-            bind(1, &post),
         bind(2, &gam),
         bind(3, &solid_buf),
         bind(4, &links_buf),
         bind(5, &beta_buf),
         bind(6, &parts),
         bind(7, &params),
-        ],
+    ];
+    if split {
+        // Тот же буфер, но показанный шейдеру по одному направлению: каждая привязка
+        // укладывается в n*4 байта вместо n*9*4, и предел драйвера на размер ОДНОЙ
+        // привязки перестаёт быть потолком для размера сетки.
+        for i in 0..math::Q {
+            let off = (i * stride * 4) as u64;
+            let len = (n * 4) as u64;
+            entries.push(bind_range(POP_F0 + i as u32, &f, off, len));
+            entries.push(bind_range(POP_P0 + i as u32, &post, off, len));
+        }
+    } else {
+        entries.push(bind(0, &f));
+        entries.push(bind(1, &post));
+    }
+    let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
+        label: Some("level"),
+        layout,
+        entries: &entries,
     });
-    GpuLevel { nx, ny, n, f, gam, bind, wall_bind, nwall, parts_count, nlinks: links.len() as u32 }
+    GpuLevel { nx, ny, n, stride, f, gam, bind, wall_bind, nwall, parts_count, nlinks: links.len() as u32 }
 }
 
 // ─────────────────────────────────────────────────────────────────────────────
@@ -1014,6 +1050,22 @@ impl Sim {
     pub fn new(spec: Spec) -> Result<Sim, String> {
         let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
             backends: wgpu::Backends::PRIMARY,
+            // Флаги читаются из окружения, и нужно это ровно ради одного —
+            // WGPU_ALLOW_UNDERLYING_NONCOMPLIANT_ADAPTER.
+            //
+            // По умолчанию wgpu ПРЯЧЕТ адаптеры, не прошедшие набор тестов соответствия
+            // Vulkan («Adapter is not Vulkan compliant, hiding adapter»), и делает это
+            // молча: request_adapter просто возвращает None. Под это правило попадает
+            // dzn (Dozen) — драйвер Mesa, транслирующий Vulkan в D3D12. В WSL2 он
+            // единственный способ добраться до карты: Linux-драйвера Vulkan у NVIDIA
+            // там нет, устройство отдаётся через /dev/dxg по протоколу WDDM.
+            //
+            // Поведение по умолчанию НЕ меняется: без переменной окружения
+            // несоответствующие адаптеры по-прежнему скрыты. Переменная — осознанное
+            // согласие считать на непроверенном драйвере, и прежде чем на нём считать,
+            // надо прогнать bench/parity.py: он сверяет GPU-путь с CPU в двойной
+            // точности на течениях с точным решением.
+            flags: wgpu::InstanceFlags::from_build_config().with_env(),
             ..Default::default()
         });
         let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions {
@@ -1051,6 +1103,45 @@ impl Sim {
             }
         }
 
+        // Влезает ли массив популяций в ОДНУ привязку. Считаем по самому крупному
+        // уровню: у патча измельчения своих узлов может оказаться больше, чем у L0.
+        let max_nodes = {
+            let n0 = spec.nx * spec.ny;
+            let n1 = match (spec.refine > 1, spec.patch) {
+                (true, Some((ax, bx, ay, by))) => {
+                    (spec.refine * (bx - ax) + 1) * (spec.refine * (by - ay) + 1)
+                }
+                _ => 0,
+            };
+            n0.max(n1)
+        };
+        let need = (math::Q * dir_stride(max_nodes) * 4) as u64;
+        // Переменная окружения нужна не для работы, а для проверки: без неё раздельный
+        // вариант включается только на драйверах с малым пределом, и сверить два
+        // варианта на одной карте было бы нечем.
+        let split = need > adapter.limits().max_storage_buffer_binding_size as u64
+            || std::env::var("KBC2D_SPLIT_POPULATIONS").is_ok_and(|v| v != "0");
+        if split {
+            // Шаг между направлениями выровнен на DIR_ALIGN значений f32. Если адаптер
+            // требует более крупного выравнивания смещений, раздельные привязки собрать
+            // нельзя — лучше сказать это прямо, чем ловить невнятную ошибку валидации.
+            let align = adapter.limits().min_storage_buffer_offset_alignment as usize;
+            if (DIR_ALIGN * 4) % align != 0 {
+                return Err(format!(
+                    "адаптер требует выравнивания смещений на {align} байт, а раскладка                      популяций рассчитана на {}. Считай на CPU либо на другом драйвере",
+                    DIR_ALIGN * 4
+                ));
+            }
+            let per = (dir_stride(max_nodes) * 4) as u64;
+            if per > adapter.limits().max_storage_buffer_binding_size as u64 {
+                return Err(format!(
+                    "адаптер ограничивает привязку {} МиБ, а одному направлению нужно {} МиБ.                      Сетка слишком крупная для этого драйвера — считай на CPU либо на драйвере                      без такого предела",
+                    adapter.limits().max_storage_buffer_binding_size / 1048576,
+                    per / 1048576
+                ));
+            }
+        }
+
         // ── топология берётся из процессорного бэкенда, а не строится заново ──
         let geom0 = cpu::Geom::build(spec.nx, spec.ny, &spec.scene);
         let fluid_count = geom0.solid.iter().filter(|s| !**s).count() as R;
@@ -1063,10 +1154,10 @@ impl Sim {
 
         let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("kbc2d.wgsl"),
-            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(SHADER)),
+            source: wgpu::ShaderSource::Wgsl(Cow::Owned(build_shader(split))),
         });
 
-        let bgl_level = level_layout(&device);
+        let bgl_level = level_layout(&device, split);
         let bgl_dyn = dyn_layout(&device);
         let bgl_amr = amr_layout(&device);
         let bgl_wall = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
@@ -1179,11 +1270,12 @@ impl Sim {
                 taper: if spec.init_uniform { spec.init_taper } else { 0.0 },
             }),
             &bgl_wall,
+            split,
         );
 
         let pre = device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("pre"),
-            size: (9 * l0.n * 4) as u64,
+            size: (math::Q * l0.stride * 4) as u64,
             usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
             mapped_at_creation: false,
         });
@@ -1219,6 +1311,7 @@ impl Sim {
                         taper: if spec.init_uniform { spec.init_taper * r as R } else { 0.0 },
                     }),
                     &bgl_wall,
+                    split,
                 );
 
             let ghosts: Vec<GGhost> = patch
@@ -1248,8 +1341,8 @@ impl Sim {
             let base = AmrParams {
                 nghost: ghosts.len() as u32,
                 nrestrict: rest.len() as u32,
-                ccount: l0.n as u32,
+                cstride: l0.stride as u32,
-                fcount: lvl1.n as u32,
+                fstride: lvl1.stride as u32,
                 r01: r01 as f32,
                 rfc: (1.0 / r01) as f32,
                 w: 0.0,
@@ -1324,7 +1417,7 @@ impl Sim {
 
         let readback = device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("readback"),
-            size: (9 * l0.n * 4) as u64,
+            size: (math::Q * l0.stride * 4) as u64,
             usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
             mapped_at_creation: false,
         });
@@ -1420,7 +1513,7 @@ impl Sim {
 
         // «до» — состояние L0 перед столкновением, нужно как старый край рамки патча
         if self.amr.is_some() {
-            enc.copy_buffer_to_buffer(&self.l0.f, 0, &self.pre, 0, (9 * self.l0.n * 4) as u64);
+            enc.copy_buffer_to_buffer(&self.l0.f, 0, &self.pre, 0, (math::Q * self.l0.stride * 4) as u64);
         }
 
         {
@@ -1620,7 +1713,7 @@ impl Sim {
 
     /// Скачать популяции L0 на хост (нужно для кадров и проверки на NaN).
     fn download_l0(&self) -> Vec<f32> {
-        self.download(&self.l0.f, (9 * self.l0.n * 4) as u64)
+        self.download(&self.l0.f, (math::Q * self.l0.stride * 4) as u64)
     }
 
     /// Скачать произвольный буфер уровня. Буфер приёма выделен один раз при сборке под самый
@@ -1653,11 +1746,12 @@ impl Sim {
     /// Полное поле скорости уровня L0 — для метрик эталонных течений и радиуса влияния.
     pub fn sample_velocity(&self) -> (Vec<R>, Vec<R>) {
         let n = self.l0.n;
+        let stride = self.l0.stride;
         let raw = self.download_l0();
         let mut ux = Vec::with_capacity(n);
         let mut uy = Vec::with_capacity(n);
         for k in 0..n {
-            let c: [R; 9] = std::array::from_fn(|i| raw[i * n + k] as R);
+            let c: [R; 9] = std::array::from_fn(|i| raw[i * stride + k] as R);
             let (_, a, b) = math::macros(&c);
             ux.push(a);
             uy.push(b);
@@ -1668,14 +1762,15 @@ impl Sim {
     /// Срез вдоль осевой линии: (ρ, u_x) по столбцам. На GPU это полное скачивание поля,
     /// поэтому x–t диагностика включается редким шагом и только там, где нужна.
     pub fn sample_centerline(&self) -> (Vec<R>, Vec<R>) {
-        let (nx, ny, n) = (self.l0.nx, self.l0.ny, self.l0.n);
+        let (nx, ny) = (self.l0.nx, self.l0.ny);
+        let stride = self.l0.stride;
         let raw = self.download_l0();
         let y = ny / 2;
         let mut rho = Vec::with_capacity(nx);
         let mut ux = Vec::with_capacity(nx);
         for x in 0..nx {
             let k = y * nx + x;
-            let c: [R; 9] = std::array::from_fn(|i| raw[i * n + k] as R);
+            let c: [R; 9] = std::array::from_fn(|i| raw[i * stride + k] as R);
             let (r, a, _) = math::macros(&c);
             rho.push(r);
             ux.push(a);
@@ -1689,9 +1784,10 @@ impl Sim {
 
     pub fn sample_field(&self, kind: FieldKind) -> (Vec<R>, Vec<bool>) {
         let (nx, ny, n) = (self.l0.nx, self.l0.ny, self.l0.n);
+        let stride = self.l0.stride;
         let raw = self.download_l0();
         let node = |k: usize| -> [R; 9] {
-            std::array::from_fn(|i| raw[i * n + k] as R)
+            std::array::from_fn(|i| raw[i * stride + k] as R)
         };
         let mut out = vec![0.0; n];
         match kind {
@@ -1766,6 +1862,19 @@ fn bind(binding: u32, buf: &wgpu::Buffer) -> wgpu::BindGroupEntry<'_> {
     wgpu::BindGroupEntry { binding, resource: buf.as_entire_binding() }
 }
 
+/// Привязать кусок буфера. Смещение обязано быть кратно
+/// `min_storage_buffer_offset_alignment` — этим и занят `dir_stride`.
+fn bind_range(binding: u32, buf: &wgpu::Buffer, offset: u64, size: u64) -> wgpu::BindGroupEntry<'_> {
+    wgpu::BindGroupEntry {
+        binding,
+        resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
+            buffer: buf,
+            offset,
+            size: std::num::NonZeroU64::new(size),
+        }),
+    }
+}
+
 fn storage(device: &wgpu::Device, label: &str, size: u64) -> wgpu::Buffer {
     device.create_buffer(&wgpu::BufferDescriptor {
         label: Some(label),
@@ -1808,10 +1917,111 @@ fn entry(binding: u32, ty: wgpu::BufferBindingType) -> wgpu::BindGroupLayoutEntr
     }
 }
 
-fn level_layout(device: &wgpu::Device) -> wgpu::BindGroupLayout {
+/// Номера привязок раздельных популяций. Взяты выше занятых (2..7), чтобы остальная
+/// часть связки не меняла нумерацию между вариантами шейдера.
+const POP_F0: u32 = 8;
+const POP_P0: u32 = 17;
+
+/// Собрать WGSL под конкретный адаптер.
+///
+/// `split` — раскладывать ли популяции по девяти привязкам вместо одной общей.
+///
+/// Зачем это вообще. Драйверы ограничивают размер ОДНОЙ привязки storage-буфера, и у
+/// разных он разный: нативные дают гигабайты, а dzn (трансляция Vulkan в D3D12 —
+/// единственный способ добраться до карты из WSL2, где Linux-драйвера Vulkan у NVIDIA
+/// нет) даёт 128 МиБ. Массив популяций занимает nx*ny*9*4 байта, так что на dzn потолок
+/// выходит 3.73 млн узлов — сетка 2048x2048 уже не проходит. При этом сам БУФЕР держать
+/// разрешено: max_buffer_size там 2 ГиБ. Ограничена только привязка.
+///
+/// Отсюда второй вариант: тот же буфер показывается девятью привязками, по одному
+/// направлению в каждой, и потолок поднимается в девять раз — до 33.5 млн узлов, чего
+/// хватает всей кампании. Платой служит switch в аксессорах, поэтому там, где предела
+/// нет, собирается прежний общий вариант без всякого switch.
+fn build_shader(split: bool) -> String {
+    let mut d = String::new();
+    if split {
+        for i in 0..math::Q {
+            d += &format!(
+                "@group(0) @binding({}) var<storage, read_write> f{i} : array<f32>;
+",
+                POP_F0 as usize + i
+            );
+            d += &format!(
+                "@group(0) @binding({}) var<storage, read_write> q{i} : array<f32>;
+",
+                POP_P0 as usize + i
+            );
+        }
+        let arm = |pfx: &str, body: &dyn Fn(usize) -> String| {
+            let mut t = String::new();
+            for i in 0..math::Q - 1 {
+                t += &format!("    case {i}u: {{ {} }}
+", body(i));
+            }
+            t += &format!("    default: {{ {} }}
+", body(math::Q - 1));
+            let _ = pfx;
+            t
+        };
+        d += "fn fget(i: u32, nd: u32) -> f32 {
+  switch i {
+";
+        d += &arm("f", &|i| format!("return f{i}[nd];"));
+        d += "  }
+}
+";
+        d += "fn fset(i: u32, nd: u32, v: f32) {
+  switch i {
+";
+        d += &arm("f", &|i| format!("f{i}[nd] = v;"));
+        d += "  }
+}
+";
+        d += "fn pget(i: u32, nd: u32) -> f32 {
+  switch i {
+";
+        d += &arm("q", &|i| format!("return q{i}[nd];"));
+        d += "  }
+}
+";
+        d += "fn pset(i: u32, nd: u32, v: f32) {
+  switch i {
+";
+        d += &arm("q", &|i| format!("q{i}[nd] = v;"));
+        d += "  }
+}
+";
+    } else {
+        d += "@group(0) @binding(0) var<storage, read_write> f    : array<f32>;
+";
+        d += "@group(0) @binding(1) var<storage, read_write> post : array<f32>;
+";
+        d += "fn fget(i: u32, nd: u32) -> f32 { return f[i*P.stride + nd]; }
+";
+        d += "fn fset(i: u32, nd: u32, v: f32) { f[i*P.stride + nd] = v; }
+";
+        d += "fn pget(i: u32, nd: u32) -> f32 { return post[i*P.stride + nd]; }
+";
+        d += "fn pset(i: u32, nd: u32, v: f32) { post[i*P.stride + nd] = v; }
+";
+    }
+    SHADER.replace("//__POPULATIONS__", &d)
+}
+
+fn level_layout(device: &wgpu::Device, split: bool) -> wgpu::BindGroupLayout {
+    let mut entries = vec![rw(2), ro(3), ro(4), ro(5), rw(6), un(7)];
+    if split {
+        for i in 0..math::Q as u32 {
+            entries.push(rw(POP_F0 + i));
+            entries.push(rw(POP_P0 + i));
+        }
+    } else {
+        entries.push(rw(0));
+        entries.push(rw(1));
+    }
     device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
         label: Some("level"),
-        entries: &[rw(0), rw(1), rw(2), ro(3), ro(4), ro(5), rw(6), un(7)],
+        entries: &entries,
     })
 }
 fn dyn_layout(device: &wgpu::Device) -> wgpu::BindGroupLayout {