Раздельные привязки популяций: снят предел dzn на размер сетки

Массив популяций занимает nx*ny*9*4 байта и показывался шейдеру одной привязкой.
dzn объявляет max_storage_buffer_binding_size = 128 МиБ, поэтому потолок выходил
3.73 млн узлов: 14 прогонов кампании из 115 падали на создании bind group, а на
них приходится 70.7% её стоимости.

Ограничена при этом ровно привязка: max_buffer_size у dzn 2047 МиБ. Поэтому тот
же буфер теперь показывается девятью привязками, по одному направлению в каждой,
и потолок поднимается до 33.5 млн узлов — самая крупная сетка кампании
(4096x4096, 16.8 млн) проходит с запасом.

Как устроено:
 * шаг между направлениями выровнен на 256 байт (dir_stride), потому что смещение
   привязки обязано быть кратно min_storage_buffer_offset_alignment;
 * обращения к популяциям в WGSL идут через fget/fset/pget/pset, а их тело
   генерируется под вариант (build_shader);
 * вариант выбирается по max_storage_buffer_binding_size адаптера — где предела
   нет, собирается прежний общий, без switch в аксессорах;
 * KBC2D_SPLIT_POPULATIONS=1 включает раздельные принудительно: иначе сверить два
   варианта на одной карте нечем.

Проверено:
 * на одном драйвере оба варианта дают одно и то же — Cd 2.39486, energy_end
   5.74813e-05; раздельный стоит 5.7% пропускной способности;
 * на сетке 2048x2048, где раздельные привязки и нужны, нативный прогон против
   контейнерного: energy_end расходится на 2.4e-07, enstrophy_end на 1.3e-07;
 * preflight в контейнере: 115 сценариев из 115, ни одного отказа (было 14);
 * 29 собственных тестов решателя зелёные.

Образ опубликован как notbigghost/kbc2d:1.2.0 (он же latest).

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

docs/theory/2d_solver/README.md

@@ -492,7 +492,7 @@ python preflight.py                # каждый сценарий старту
 
 ## Развёртывание (Docker)
 
-Собранный образ опубликован: **`notbigghost/kbc2d:1.1.0`** (он же `latest`, платформа `linux/amd64`).
+Собранный образ опубликован: **`notbigghost/kbc2d:1.2.0`** (он же `latest`, платформа `linux/amd64`).
 Исходники на сервере не нужны — достаточно перенести туда один файл
 `docker-compose.server.yml`:
 
@@ -549,7 +549,7 @@ NVIDIA там нет**: карта отдаётся через `/dev/dxg` по
 монтирование `/usr/lib/wsl`. Запуск через `docker-compose.wsl.yml`, подробности в
 [`bench/README.md`](bench/README.md).
 
-Две вещи, которые надо знать про этот путь.
+Три вещи, которые надо знать про этот путь.
 
 **wgpu по умолчанию прячет несоответствующие адаптеры.** dzn сообщает о себе
 `conformanceVersion = 0.0.0.0`, и wgpu молча его отбрасывает — решатель докладывает, что
@@ -558,6 +558,17 @@ GPU не найден. Согласие даётся явно, переменн
 `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 шагов):
 

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

@@ -6,7 +6,7 @@
 
 ## Быстрый старт на сервере
 
-Образ опубликован, собирать ничего не нужно: **`notbigghost/kbc2d:1.1.0`**. Исходники на
+Образ опубликован, собирать ничего не нужно: **`notbigghost/kbc2d:1.2.0`**. Исходники на
 сервере тоже не нужны — переносится один файл `docker-compose.server.yml`.
 
 ```sh
@@ -83,44 +83,39 @@ dzn сообщает о себе `conformanceVersion = 0.0.0.0`: набор те
 Точность трансляция не портит. Но это замер на Intel; на своей карте прогоните сами — две
 минуты.
 
-### Чего этот путь НЕ может: предел 128 МиБ на привязку
+### Предел 128 МиБ на привязку и как он снят
 
-Главное ограничение dzn, и оно жёсткое. Драйвер объявляет
+dzn объявляет `max_storage_buffer_binding_size = 128 МиБ` (2²⁷ байт), тогда как нативные
-`max_storage_buffer_binding_size = 128 МиБ` (2²⁷ байт), тогда как нативные драйверы дают
+драйверы дают гигабайты. Массив популяций занимает `nx · ny · 9 · 4` байта, так что при
-гигабайты. Сам буфер держать можно — `max_buffer_size` у dzn 2047 МиБ, — но показать
+одной общей привязке потолок выходил **3.73 млн узлов** (примерно 1920×1920), и 14
-шейдеру одной привязкой больше 128 МиБ нельзя.
+прогонов кампании из 115 падали на создании bind group:
-
-Массив функций распределения занимает `nx · ny · 9 · 4` байт, значит потолок —
-**3.73 млн узлов** (примерно 1920×1920). Всё, что крупнее, падает на создании bind group:
 
 ```
 Buffer binding 0 range 150994944 exceeds `max_*_buffer_binding_size` limit 134217728
 ```
 
-В кампании таких прогонов **14 из 115**, и это не мелочь: на них приходится **70.7%
+На эти 14 приходилось 70.7% стоимости кампании — дорогие прогоны как раз крупносеточные.
-стоимости**, потому что дорогие прогоны — как раз крупносеточные.
 
-| прогон | сетка | буфер |
+Существенно, что ограничена только **привязка**: `max_buffer_size` у dzn 2047 МиБ, то есть
-|---|---|---|
+буфер держать разрешено, нельзя лишь показать шейдеру его целиком. Поэтому решатель теперь
-| B10_cyl_d128 | 3.93 млн узлов | 135 МБ |
+умеет показывать тот же буфер **девятью привязками**, по одному направлению в каждой.
-| A10, A11 (турбулентность N=2048) | 4.19 млн | 144 МБ |
+Потолок поднимается в девять раз — до 33.5 млн узлов, чего хватает всей кампании с запасом
-| D14_naca_chord192 | 4.72 млн | 162 МБ |
+(самая крупная сетка в ней 4096×4096 — 16.8 млн узлов).
-| I01–I09 (сверхмелкие сетки) | 8.39 млн | 288 МБ |
-| A12 (турбулентность N=4096) | 16.8 млн | 576 МБ |
 
-`preflight` показывает их все с причиной — запускать его до кампании обязательно.
+Вариант выбирается сам, по `max_storage_buffer_binding_size` адаптера: где предела нет,
+собирается прежний общий вариант без `switch` в аксессорах. Проверить оба на одной карте
+можно переменной `KBC2D_SPLIT_POPULATIONS=1` — она включает раздельные привязки
+принудительно.
 
-Что с этим делать — три возможности, и ни одна не бесплатна:
+Замерено на Intel Iris Xe, один драйвер, два варианта привязки:
 
-1. **Разделить кампанию.** 101 прогон в контейнере, 14 крупных — нативной сборкой (в
+| | Cd | energy_end | MLUPS (цилиндр) |
-   Windows или на настоящем Linux-сервере). Работает сегодня, кода не трогает, но это
+|---|---|---|---|
-   два разных запуска.
+| общая привязка | 2.39486 | 5.74813e-05 | 169.7 |
-2. **Перестроить буферы решателя** под привязку на направление: девять привязок по
+| девять привязок | 2.39486 | 5.74813e-05 | 160.0 |
-   `nx · ny · 4` байта вместо одной на `nx · ny · 9 · 4`. Тогда потолок поднимается до
+
-   33.5 млн узлов и в контейнер влезает вся кампания. Цена — переделка раскладки в
+Числа совпадают полностью; раздельный вариант стоит **5.7%** пропускной способности, и
-   ядрах WGSL (24 места обращения) и повторная валидация; на нативных драйверах смысла
+включается только там, где без него счёт вообще невозможен.
-   в этом нет, так что понадобятся два варианта шейдера.
-3. **Считать всю кампанию нативно.** Ни предела, ни платы за трансляцию — но и контейнера.
 
 ### Чего трансляция стоит по скорости
 
@@ -132,6 +127,10 @@ Buffer binding 0 range 150994944 exceeds `max_*_buffer_binding_size` limit 13421
 | 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×.
 
 Каждый шаг решателя — несколько отправок в очередь; на мелкой сетке трансляция вызова
 стоит дороже самого счёта, на крупной размазывается. Для кампании это решающее

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

@@ -22,7 +22,7 @@ name: kbc2d
 
 # ── общая часть всех сервисов ────────────────────────────────────────────────
 x-kbc2d: &kbc2d
-  image: notbigghost/kbc2d:1.1.0
+  image: notbigghost/kbc2d:1.2.0
   pull_policy: missing
   volumes:
     - ./out:/work/bench/out

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

@@ -32,14 +32,14 @@ name: kbc2d
 
 # ── общая часть всех сервисов ────────────────────────────────────────────────
 x-kbc2d: &kbc2d
-  image: notbigghost/kbc2d:1.1.0
+  image: notbigghost/kbc2d:1.2.0
   # Контекст сборки — каталог с этим файлом. Если образа нет ни локально, ни в реестре,
   # достаточно `docker compose -f docker-compose.wsl.yml build`: доступ к Docker Hub
   # для запуска не обязателен.
   build:
     context: .
     args:
-      VERSION: "1.1.0"
+      VERSION: "1.2.0"
   pull_policy: missing
   devices:
     - /dev/dxg:/dev/dxg

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 }
 }
 
 // ─────────────────────────────────────────────────────────────────────────────
@@ -1067,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;
@@ -1079,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 {
@@ -1195,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,
         });
@@ -1235,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
@@ -1264,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,
@@ -1340,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,
         });
@@ -1436,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);
         }
 
         {
@@ -1636,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)
     }
 
     /// Скачать произвольный буфер уровня. Буфер приёма выделен один раз при сборке под самый
@@ -1669,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);
@@ -1684,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);
@@ -1705,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 {
@@ -1782,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),
@@ -1824,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 {