diff --git a/docs/theory/2d_solver/README.md b/docs/theory/2d_solver/README.md index b93e907..8525eb4 100644 --- a/docs/theory/2d_solver/README.md +++ b/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 шагов): diff --git a/docs/theory/2d_solver/bench/README.md b/docs/theory/2d_solver/bench/README.md index ed14c29..7575242 100644 --- a/docs/theory/2d_solver/bench/README.md +++ b/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, и оно жёсткое. Драйвер объявляет -`max_storage_buffer_binding_size = 128 МиБ` (2²⁷ байт), тогда как нативные драйверы дают -гигабайты. Сам буфер держать можно — `max_buffer_size` у dzn 2047 МиБ, — но показать -шейдеру одной привязкой больше 128 МиБ нельзя. - -Массив функций распределения занимает `nx · ny · 9 · 4` байт, значит потолок — -**3.73 млн узлов** (примерно 1920×1920). Всё, что крупнее, падает на создании bind group: +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 из 115**, и это не мелочь: на них приходится **70.7% -стоимости**, потому что дорогие прогоны — как раз крупносеточные. +На эти 14 приходилось 70.7% стоимости кампании — дорогие прогоны как раз крупносеточные. -| прогон | сетка | буфер | -|---|---|---| -| B10_cyl_d128 | 3.93 млн узлов | 135 МБ | -| A10, A11 (турбулентность N=2048) | 4.19 млн | 144 МБ | -| D14_naca_chord192 | 4.72 млн | 162 МБ | -| I01–I09 (сверхмелкие сетки) | 8.39 млн | 288 МБ | -| A12 (турбулентность N=4096) | 16.8 млн | 576 МБ | +Существенно, что ограничена только **привязка**: `max_buffer_size` у dzn 2047 МиБ, то есть +буфер держать разрешено, нельзя лишь показать шейдеру его целиком. Поэтому решатель теперь +умеет показывать тот же буфер **девятью привязками**, по одному направлению в каждой. +Потолок поднимается в девять раз — до 33.5 млн узлов, чего хватает всей кампании с запасом +(самая крупная сетка в ней 4096×4096 — 16.8 млн узлов). -`preflight` показывает их все с причиной — запускать его до кампании обязательно. +Вариант выбирается сам, по `max_storage_buffer_binding_size` адаптера: где предела нет, +собирается прежний общий вариант без `switch` в аксессорах. Проверить оба на одной карте +можно переменной `KBC2D_SPLIT_POPULATIONS=1` — она включает раздельные привязки +принудительно. -Что с этим делать — три возможности, и ни одна не бесплатна: +Замерено на Intel Iris Xe, один драйвер, два варианта привязки: -1. **Разделить кампанию.** 101 прогон в контейнере, 14 крупных — нативной сборкой (в - Windows или на настоящем Linux-сервере). Работает сегодня, кода не трогает, но это - два разных запуска. -2. **Перестроить буферы решателя** под привязку на направление: девять привязок по - `nx · ny · 4` байта вместо одной на `nx · ny · 9 · 4`. Тогда потолок поднимается до - 33.5 млн узлов и в контейнер влезает вся кампания. Цена — переделка раскладки в - ядрах WGSL (24 места обращения) и повторная валидация; на нативных драйверах смысла - в этом нет, так что понадобятся два варианта шейдера. -3. **Считать всю кампанию нативно.** Ни предела, ни платы за трансляцию — но и контейнера. +| | Cd | energy_end | MLUPS (цилиндр) | +|---|---|---|---| +| общая привязка | 2.39486 | 5.74813e-05 | 169.7 | +| девять привязок | 2.39486 | 5.74813e-05 | 160.0 | + +Числа совпадают полностью; раздельный вариант стоит **5.7%** пропускной способности, и +включается только там, где без него счёт вообще невозможен. ### Чего трансляция стоит по скорости @@ -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×. Каждый шаг решателя — несколько отправок в очередь; на мелкой сетке трансляция вызова стоит дороже самого счёта, на крупной размазывается. Для кампании это решающее diff --git a/docs/theory/2d_solver/docker-compose.server.yml b/docs/theory/2d_solver/docker-compose.server.yml index c6f85a5..c4ee1f8 100644 --- a/docs/theory/2d_solver/docker-compose.server.yml +++ b/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 diff --git a/docs/theory/2d_solver/docker-compose.wsl.yml b/docs/theory/2d_solver/docker-compose.wsl.yml index cbc4df9..1ca5ba3 100644 --- a/docs/theory/2d_solver/docker-compose.wsl.yml +++ b/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 diff --git a/docs/theory/2d_solver/src/gpu.rs b/docs/theory/2d_solver/src/gpu.rs index ac7541f..3bb15fc 100644 --- a/docs/theory/2d_solver/src/gpu.rs +++ b/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 f : array; -@group(0) @binding(1) var post : array; @group(0) @binding(2) var gam : array; @group(0) @binding(3) var solid: array; @group(0) @binding(4) var links: array; @group(0) @binding(5) var beta : array; @group(0) @binding(6) var parts: array; @group(0) @binding(7) var P : LevelParams; +//__POPULATIONS__ @group(1) @binding(0) var D : Dyn; @group(1) @binding(1) var S : Substep; @@ -242,8 +244,8 @@ fn shift9(d: array, model: u32) -> array { return s; } -fn load9(base: ptr>, off: u32, n: u32) { - for (var i = 0u; i < 9u; i = i + 1u) { (*base)[i] = f[i*n + off]; } +fn load9(base: ptr>, off: u32) { + 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, let nd = lin(gid, nwg); if (nd >= P.n) { return; } var fv: array; - 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, 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, 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, let x = lin(gid, nwg); if (x >= P.nx) { return; } // зеркальное отражение: касательный импульс сохраняется, нормальный заворачивается - f[2u*P.n + x] = f[4u*P.n + x]; - f[5u*P.n + x] = f[8u*P.n + x]; - f[6u*P.n + x] = f[7u*P.n + x]; + fset(2u, x, fget(4u, x)); + fset(5u, x, fget(8u, 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]; - f[7u*P.n + t] = f[6u*P.n + t]; - f[8u*P.n + t] = f[5u*P.n + t]; + fset(4u, t, fget(2u, t)); + fset(7u, t, fget(6u, t)); + fset(8u, t, fget(5u, t)); } @compute @workgroup_size(64) @@ -399,23 +401,23 @@ fn k_channel(@builtin(global_invocation_id) gid: vec3, // вход: скоростной Zou-He по ВСЕМУ столбцу, включая угловые узлы let a = y*P.nx; var fv: array; - 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; - 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; - 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) { loop { if (k >= P.nlinks) { break; } let L = links[k]; - let fp = post[L.i*P.n + L.node]; - let fb = f[L.ib*P.n + L.node]; + let fp = pget(L.i, 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, p.cnt = 0.0; p.degen = 0.0; p.xineg = 0.0; if (nd < P.n && solid[nd] == 0u) { var fv: array; - 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, // зонд следа снимается с того уровня, на котором он лежит if ((P.flags & 2u) != 0u && nd == P.probe_node) { var fv: array; - 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) { // хранилища прошлого шага. fn u_at_t(nd: u32) -> vec2 { var ff: array; - 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(m.y, m.z); } @@ -733,15 +735,15 @@ fn k_moment_wall(@builtin(global_invocation_id) gid: vec3, var opp = array(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]; } - else { rho = rho + f[i*P.n + node]; } + if (missing[i] == 1u) { rho = rho + pget(opp[i], 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, var neq = array(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; - 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, 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, if (k >= A.nrestrict) { return; } let pr = rest[k]; var ff: array; - 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 { +/// Выравнивание шага между направлениями, во ФЛОАТАХ. +/// +/// Смещение каждой привязки обязано быть кратно `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 { let n = f.len(); - let mut v = vec![0.0f32; 9 * n]; - for (k, cell) in f.iter().enumerate() { - for i in 0..9 { - v[i * n + k] = cell[i] as f32; + debug_assert!(stride >= n); + let mut v = vec![0.0f32; 9 * stride]; + for (k, c) in f.iter().enumerate() { + 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 mut entries = vec![ + bind(2, &gam), + bind(3, &solid_buf), + bind(4, &links_buf), + bind(5, &beta_buf), + bind(6, &parts), + bind(7, ¶ms), + ]; + 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: &[ - 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, ¶ms), - ], + 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 = 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, - fcount: lvl1.n as u32, + cstride: l0.stride 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 { - 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, Vec) { 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, Vec) { - 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, Vec) { 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 f{i} : array; +", + POP_F0 as usize + i + ); + d += &format!( + "@group(0) @binding({}) var q{i} : array; +", + 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 f : array; +"; + d += "@group(0) @binding(1) var post : array; +"; + 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 {