//! БЭКЕНД ПОД ВИДЕОКАРТУ на wgpu (Vulkan / DX12 / Metal).
//!
//! Физика построчно повторяет `math.rs`, но на f32: в WGSL нет двойной точности. Это законно
//! именно потому, что порог вырожденности γ в схеме ОТНОСИТЕЛЬНЫЙ (`GREL`, доля от ⟨Δ|Δ⟩), а не
//! абсолютный — иначе в f32 он срабатывал бы на подавляющем большинстве узлов и схема молча
//! выродилась бы в LBGK. Топология задачи (маски, Bouzidi-линки, рамка и рестрикция патча) не
//! дублируется: она строится теми же процедурами из `cpu`, а сюда загружается готовыми списками.
//!
//! Раскладка полей — SoA: `f[i*n + node]`. В отличие от AoS на процессоре, здесь важен
//! коалесцированный доступ: соседние потоки читают соседние узлы одной и той же популяции.
//!
//! Ограничение реализации: за каждый шаг делается одно чтение 48-байтового буфера итогов
//! (сила, зонд, статистика). Это синхронизация с GPU на каждом шаге; она заметно ограничивает
//! частоту шагов на мелких сетках, зато `step()` возвращает те же величины, что и процессорный
//! бэкенд, шаг в шаг. На крупных сетках стоимость самой синхронизации теряется на фоне счёта.

use std::borrow::Cow;

use bytemuck::{Pod, Zeroable};
use wgpu::util::DeviceExt;

use crate::cpu;
use crate::math::{self, R};
use crate::{Collision, FieldKind, Spec, StepRec};

const WG: u32 = 64;

// ─────────────────────────────────────────────────────────────────────────────
//  Структуры, разделяемые с шейдером
// ─────────────────────────────────────────────────────────────────────────────

#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct LevelParams {
    n: u32,
    nx: u32,
    ny: u32,
    nlinks: u32,
    bcx: f32,
    bcy: f32,
    probe_node: u32,
    /// бит 0 — писать частичные суммы статистики; бит 1 — на этом уровне лежит зонд
    flags: u32,
}

#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Dyn {
    ux_in: f32,
    uy_in: f32,
    rho_out: f32,
    /// 0 — сдвиговая часть только девиатор (N1), 1 — девиатор со следом (N2)
    kbc_model: u32,
    outlet_extrap: u32,
    nparts: u32,
    refine: u32,
    collision: u32,
}

#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct Substep {
    idx: u32,
    _p: [u32; 3],
}

#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct AmrParams {
    nghost: u32,
    nrestrict: u32,
    ccount: u32,
    fcount: u32,
    r01: f32,
    rfc: f32,
    w: f32,
    _pad: f32,
}

#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct GLink {
    node: u32,
    far: u32,
    i: u32,
    ib: u32,
    kind: u32,
    q: f32,
    _p: [u32; 2],
}

#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable)]
struct GGhost {
    fine: u32,
    c00: u32,
    c10: u32,
    c01: u32,
    c11: u32,
    tx: f32,
    ty: f32,
    _p: u32,
}

/// Итоги шага, читаемые на хост. Порядок обязан совпадать с `results` в шейдере.
#[repr(C)]
#[derive(Clone, Copy, Pod, Zeroable, Default, Debug)]
struct Results {
    fx: f32,
    fy: f32,
    tz: f32,
    uy_probe: f32,
    rho_sum: f32,
    max_u: f32,
    g_sum: f32,
    g_min: f32,
    g_max: f32,
    cnt: f32,
    degen: f32,
    xineg: f32,
}

// ─────────────────────────────────────────────────────────────────────────────
//  Шейдер
// ─────────────────────────────────────────────────────────────────────────────

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 };
struct Dyn { ux_in:f32, uy_in:f32, rho_out:f32, kbc_model:u32, outlet_extrap:u32, nparts:u32, refine:u32, collision: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 GLink { node:u32, far:u32, i:u32, ib:u32, kind:u32, q:f32, p0: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;

@group(1) @binding(0) var<uniform>             D    : Dyn;
@group(1) @binding(1) var<uniform>             S    : Substep;
// results[0..11] — итоги шага, results[16 + s*4 + c] — сила подшага s.
// Сила и итоги держатся в ОДНОМ буфере намеренно: связка уровня использует 7 storage-биндингов,
// и отдельный буфер сил вывел бы конвейер за предел 8, гарантируемый лимитами по умолчанию.
@group(1) @binding(2) var<storage, read_write> results  : array<f32>;

@group(2) @binding(0) var<storage, read>       amr_src: array<f32>;
@group(2) @binding(1) var<storage, read_write> amr_dst: array<f32>;
@group(2) @binding(2) var<storage, read>       ghosts : array<GGhost>;
@group(2) @binding(3) var<storage, read>       gh_a   : array<f32>;
@group(2) @binding(4) var<storage, read>       gh_b   : array<f32>;
@group(2) @binding(5) var<storage, read>       rest   : array<vec2<u32>>;
@group(2) @binding(6) var<uniform>             A      : AmrParams;

const GREL: f32 = 1e-8;
const CS2 : f32 = 0.3333333333;

// ───── физика: построчный перенос math.rs ─────

// Энтропийное равновесие в product-form (точный максимизатор энтропии при заданных rho, rho*u)
fn feq9(rho: f32, ux0: f32, uy0: f32) -> array<f32,9> {
    let ux = clamp(ux0, -0.95, 0.95);
    let uy = clamp(uy0, -0.95, 0.95);
    let sx = sqrt(1.0 + 3.0*ux*ux);
    let sy = sqrt(1.0 + 3.0*uy*uy);
    let base = rho * (2.0 - sx) * (2.0 - sy);
    let qx = (2.0*ux + sx) / (1.0 - ux);
    let qy = (2.0*uy + sy) / (1.0 - uy);
    let ix = 1.0 / qx;
    let iy = 1.0 / qy;
    return array<f32,9>(
        base*(4.0/9.0),
        base*(1.0/9.0)*qx,  base*(1.0/9.0)*qy,  base*(1.0/9.0)*ix,  base*(1.0/9.0)*iy,
        base*(1.0/36.0)*qx*qy, base*(1.0/36.0)*ix*qy, base*(1.0/36.0)*ix*iy, base*(1.0/36.0)*qx*iy);
}

fn macros9(fv: array<f32,9>) -> vec3<f32> {
    let rho = fv[0]+fv[1]+fv[2]+fv[3]+fv[4]+fv[5]+fv[6]+fv[7]+fv[8];
    let mx = fv[1]+fv[5]+fv[8]-fv[3]-fv[6]-fv[7];
    let my = fv[2]+fv[5]+fv[6]-fv[4]-fv[7]-fv[8];
    return vec3<f32>(rho, mx/rho, my/rho);
}

// Проекция неравновесия на сдвиговую часть. model=0: {N, Pi_xy} (N1); model=1: плюс след T (N2)
fn shift9(d: array<f32,9>, model: u32) -> array<f32,9> {
    let a = 0.25*(d[1] + d[3] - d[2] - d[4]);
    let b = 0.25*(d[5] + d[7] - d[6] - d[8]);
    var s = array<f32,9>(0.0, a, -a, a, -a, b, -b, b, -b);
    if (model == 1u) {
        let dt = d[1] + d[2] + d[3] + d[4] + 2.0*(d[5] + d[6] + d[7] + d[8]);
        let q = 0.25*dt;
        s[0] = s[0] - dt;
        s[1] = s[1] + q; s[2] = s[2] + q; s[3] = s[3] + q; s[4] = s[4] + q;
    }
    return s;
}

fn load9(base: ptr<function, array<f32,9>>, off: u32, n: u32) {
    for (var i = 0u; i < 9u; i = i + 1u) { (*base)[i] = f[i*n + off]; }
}

// Возвращает (пост-столкновительные популяции, gamma, признак вырождения)
struct CollOut { fv: array<f32,9>, gamma: f32, degen: f32 };

fn collide9(fin: array<f32,9>, b: f32, op: u32, model: u32) -> CollOut {
    var out: CollOut;
    // WGSL разрешает переменный индекс только по памяти (var), а не по значению (let/параметр),
    // поэтому всё, что индексируется в цикле, кладётся в var
    var fv = fin;
    let m = macros9(fin);
    var fe = feq9(m.x, m.y, m.z);
    if (op == 1u) {                      // LBGK: то же самое при gamma = 2
        var g: array<f32,9>;
        for (var i = 0u; i < 9u; i = i + 1u) { g[i] = fv[i] + 2.0*b*(fe[i] - fv[i]); }
        out.fv = g; out.gamma = 2.0; out.degen = 0.0;
        return out;
    }
    var d: array<f32,9>;
    for (var i = 0u; i < 9u; i = i + 1u) { d[i] = fv[i] - fe[i]; }
    var ds = shift9(d, model);
    var num = 0.0; var den = 0.0; var nrm = 0.0;
    for (var i = 0u; i < 9u; i = i + 1u) {
        let inv = 1.0 / fe[i];
        let dh = d[i] - ds[i];
        num = num + ds[i]*dh*inv;
        den = den + dh*dh*inv;
        nrm = nrm + d[i]*d[i]*inv;
    }
    // относительный порог: den квадратична по неравновесию и физически мала
    let ok = den > GREL*nrm;
    let binv = 1.0 / b;
    var gm = 2.0;
    if (ok) { gm = binv - (2.0 - binv)*num/den; }
    var g: array<f32,9>;
    for (var i = 0u; i < 9u; i = i + 1u) {
        let dh = d[i] - ds[i];
        g[i] = fv[i] - b*(2.0*ds[i] + gm*dh);
    }
    out.fv = g; out.gamma = gm;
    out.degen = select(1.0, 0.0, ok);
    return out;
}

fn zou_he_inlet(fin: array<f32,9>, ux: f32, uy: f32) -> array<f32,9> {
    var g = fin;
    let rho = (g[0] + g[2] + g[4] + 2.0*(g[3] + g[6] + g[7])) / (1.0 - ux);
    let dd = 0.5*(g[2] - g[4]);
    g[1] = g[3] + (2.0/3.0)*rho*ux;
    g[5] = g[7] - dd + (1.0/6.0)*rho*ux + 0.5*rho*uy;
    g[8] = g[6] + dd + (1.0/6.0)*rho*ux - 0.5*rho*uy;
    return g;
}

fn zou_he_outlet(fin: array<f32,9>, rho_out: f32, uy_out: f32) -> array<f32,9> {
    var g = fin;
    let ux = (g[0] + g[2] + g[4] + 2.0*(g[1] + g[5] + g[8])) / rho_out - 1.0;
    let dd = 0.5*(g[2] - g[4]);
    g[3] = g[1] - (2.0/3.0)*rho_out*ux;
    g[6] = g[8] - dd - (1.0/6.0)*rho_out*ux + 0.5*rho_out*uy_out;
    g[7] = g[5] + dd - (1.0/6.0)*rho_out*ux - 0.5*rho_out*uy_out;
    return g;
}

// ───── ядра ─────

@compute @workgroup_size(64)
fn k_collide(@builtin(global_invocation_id) gid: vec3<u32>) {
    let nd = gid.x;
    if (nd >= P.n) { return; }
    var fv: array<f32,9>;
    load9(&fv, nd, P.n);
    if (solid[nd] != 0u) {
        // внутри тела не считаем: популяции там фиктивны, Bouzidi всё равно их перекрывает
        for (var i = 0u; i < 9u; i = i + 1u) { post[i*P.n + 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]; }
    gam[nd] = r.gamma;
}

@compute @workgroup_size(64)
fn k_stream(@builtin(global_invocation_id) gid: vec3<u32>) {
    let nd = gid.x;
    if (nd >= P.n) { return; }
    var cx = array<i32,9>(0, 1, 0, -1, 0, 1, -1, -1, 1);
    var cy = array<i32,9>(0, 0, 1, 0, -1, 1, 1, -1, -1);
    let x = i32(nd % P.nx);
    let y = i32(nd / P.nx);
    let nxi = i32(P.nx);
    let nyi = i32(P.ny);
    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)];
    }
}

@compute @workgroup_size(64)
fn k_bouzidi(@builtin(global_invocation_id) gid: vec3<u32>) {
    let k = gid.x;
    if (k >= P.nlinks) { return; }
    let L = links[k];
    let fi = post[L.i*P.n + L.node];
    var v = fi;
    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];
    } 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];
    }
    f[L.ib*P.n + L.node] = v;
}

@compute @workgroup_size(64)
fn k_walls(@builtin(global_invocation_id) gid: vec3<u32>) {
    let x = gid.x;
    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];
    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];
}

@compute @workgroup_size(64)
fn k_channel(@builtin(global_invocation_id) gid: vec3<u32>) {
    let y = gid.x;
    if (y >= P.ny) { return; }
    // вход: скоростной Zou-He по ВСЕМУ столбцу, включая угловые узлы
    let a = y*P.nx;
    var fv: array<f32,9>;
    load9(&fv, a, P.n);
    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]; }
    // выход: давление-Zou-He
    let b = y*P.nx + P.nx - 1u;
    var gv: array<f32,9>;
    load9(&gv, b, P.n);
    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);
        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]; }
}

// сила и момент по GMEM, редукция одной рабочей группой
var<workgroup> wfx: array<f32, 256>;
var<workgroup> wfy: array<f32, 256>;
var<workgroup> wtz: array<f32, 256>;

@compute @workgroup_size(256)
fn k_force(@builtin(local_invocation_id) lid: vec3<u32>) {
    let t = lid.x;
    var cx = array<f32,9>(0.0, 1.0, 0.0, -1.0, 0.0, 1.0, -1.0, -1.0, 1.0);
    var cy = array<f32,9>(0.0, 0.0, 1.0, 0.0, -1.0, 1.0, 1.0, -1.0, -1.0);
    var sx = 0.0; var sy = 0.0; var sz = 0.0;
    var k = t;
    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 dfx = cx[L.i]*fp + cx[L.i]*fb;
        let dfy = cy[L.i]*fp + cy[L.i]*fb;
        sx = sx + dfx;
        sy = sy + dfy;
        // плечо до точки пересечения линка со стенкой, а не до узла
        let rx = f32(L.node % P.nx) - P.bcx + L.q*cx[L.i];
        let ry = f32(L.node / P.nx) - P.bcy + L.q*cy[L.i];
        sz = sz + rx*dfy - ry*dfx;
        k = k + 256u;
    }
    wfx[t] = sx; wfy[t] = sy; wtz[t] = sz;
    workgroupBarrier();
    var s = 128u;
    loop {
        if (s == 0u) { break; }
        if (t < s) {
            wfx[t] = wfx[t] + wfx[t + s];
            wfy[t] = wfy[t] + wfy[t + s];
            wtz[t] = wtz[t] + wtz[t + s];
        }
        workgroupBarrier();
        s = s >> 1u;
    }
    if (t == 0u) {
        results[16u + S.idx*4u + 0u] = wfx[0];
        results[16u + S.idx*4u + 1u] = wfy[0];
        results[16u + S.idx*4u + 2u] = wtz[0];
    }
}

// частичные суммы диагностики: одна Partial на рабочую группу
var<workgroup> wp: array<Partial, 64>;

@compute @workgroup_size(64)
fn k_stats1(@builtin(global_invocation_id) gid: vec3<u32>,
            @builtin(local_invocation_id) lid: vec3<u32>,
            @builtin(workgroup_id) wid: vec3<u32>) {
    let nd = gid.x;
    let t = lid.x;
    var p: Partial;
    p.rho = 0.0; p.maxu = 0.0; p.gsum = 0.0; p.gmin = 1e30; p.gmax = -1e30;
    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);
        let m = macros9(fv);
        p.rho = m.x;
        p.maxu = sqrt(m.y*m.y + m.z*m.z);
        let g = gam[nd];
        p.gsum = g; p.gmin = g; p.gmax = g; p.cnt = 1.0;
        // вырожденный узел помечен ровно gamma = 2 (см. k_collide)
        p.degen = select(0.0, 1.0, g == 2.0);
        // объёмная вязкость модели D: xi = cs^2 (1/(gamma*beta) - 1/2)
        let b = beta[nd % P.nx];
        // при следе в сдвиговой части (N2) объёмная вязкость равна сдвиговой и всегда > 0
        var xi = CS2*(1.0/(b + b) - 0.5);
        if (D.kbc_model == 0u) { xi = CS2*(1.0/(g*b) - 0.5); }
        p.xineg = select(0.0, 1.0, xi < 0.0);
    }
    // зонд следа снимается с того уровня, на котором он лежит
    if ((P.flags & 2u) != 0u && nd == P.probe_node) {
        var fv: array<f32,9>;
        load9(&fv, nd, P.n);
        let m = macros9(fv);
        results[3] = m.z;
    }
    wp[t] = p;
    workgroupBarrier();
    var s = 32u;
    loop {
        if (s == 0u) { break; }
        if (t < s) {
            let o = wp[t + s];
            wp[t].rho = wp[t].rho + o.rho;
            wp[t].maxu = max(wp[t].maxu, o.maxu);
            wp[t].gsum = wp[t].gsum + o.gsum;
            wp[t].gmin = min(wp[t].gmin, o.gmin);
            wp[t].gmax = max(wp[t].gmax, o.gmax);
            wp[t].cnt = wp[t].cnt + o.cnt;
            wp[t].degen = wp[t].degen + o.degen;
            wp[t].xineg = wp[t].xineg + o.xineg;
        }
        workgroupBarrier();
        s = s >> 1u;
    }
    if (t == 0u && (P.flags & 1u) != 0u) { parts[wid.x] = wp[0]; }
}

// Размер группы обязан совпадать с длиной wp: при 256 потоках на 64 ячейки четыре потока
// писали бы в одну и ту же ячейку и три четверти частичных сумм терялись бы молча.
@compute @workgroup_size(64)
fn k_stats2(@builtin(local_invocation_id) lid: vec3<u32>) {
    let t = lid.x;
    var p: Partial;
    p.rho = 0.0; p.maxu = 0.0; p.gsum = 0.0; p.gmin = 1e30; p.gmax = -1e30;
    p.cnt = 0.0; p.degen = 0.0; p.xineg = 0.0;
    var k = t;
    loop {
        if (k >= D.nparts) { break; }
        let o = parts[k];
        p.rho = p.rho + o.rho;
        p.maxu = max(p.maxu, o.maxu);
        p.gsum = p.gsum + o.gsum;
        p.gmin = min(p.gmin, o.gmin);
        p.gmax = max(p.gmax, o.gmax);
        p.cnt = p.cnt + o.cnt;
        p.degen = p.degen + o.degen;
        p.xineg = p.xineg + o.xineg;
        k = k + 64u;
    }
    wp[t] = p;
    workgroupBarrier();
    // сведение 256 потоков через 64 ячейки делаем последовательно нулевым потоком:
    // объём данных крошечный, а корректность важнее пары микросекунд
    if (t == 0u) {
        var q: Partial;
        q.rho = 0.0; q.maxu = 0.0; q.gsum = 0.0; q.gmin = 1e30; q.gmax = -1e30;
        q.cnt = 0.0; q.degen = 0.0; q.xineg = 0.0;
        for (var j = 0u; j < 64u; j = j + 1u) {
            let o = wp[j];
            q.rho = q.rho + o.rho;
            q.maxu = max(q.maxu, o.maxu);
            q.gsum = q.gsum + o.gsum;
            q.gmin = min(q.gmin, o.gmin);
            q.gmax = max(q.gmax, o.gmax);
            q.cnt = q.cnt + o.cnt;
            q.degen = q.degen + o.degen;
            q.xineg = q.xineg + o.xineg;
        }
        var fx = 0.0; var fy = 0.0; var tz = 0.0;
        for (var s = 0u; s < D.refine; s = s + 1u) {
            fx = fx + results[16u + s*4u + 0u];
            fy = fy + results[16u + s*4u + 1u];
            tz = tz + results[16u + s*4u + 2u];
        }
        let inv = 1.0 / f32(D.refine);
        results[0] = fx*inv;
        results[1] = fy*inv;
        results[2] = tz*inv;
        results[4] = q.rho;
        results[5] = q.maxu;
        results[6] = q.gsum;
        results[7] = q.gmin;
        results[8] = q.gmax;
        results[9] = q.cnt;
        results[10] = q.degen;
        results[11] = q.xineg;
    }
}

// ───── AMR ─────

@compute @workgroup_size(64)
fn k_ghost(@builtin(global_invocation_id) gid: vec3<u32>) {
    let k = gid.x;
    if (k >= A.nghost) { return; }
    let g = ghosts[k];
    let w00 = (1.0 - g.tx)*(1.0 - g.ty);
    let w10 = g.tx*(1.0 - g.ty);
    let w01 = (1.0 - g.tx)*g.ty;
    let w11 = g.tx*g.ty;
    // отдельно интерполируются макропеременные, отдельно неравновесная часть (как в cpu.rs)
    var ws = array<f32,4>(w00, w10, w01, w11);
    var cs = array<u32,4>(g.c00, g.c10, g.c01, g.c11);
    var rho = 0.0; var ux = 0.0; var uy = 0.0;
    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]]; }
        let m = macros9(cf);
        rho = rho + ws[j]*m.x;
        ux  = ux  + ws[j]*m.y;
        uy  = uy  + ws[j]*m.z;
        var fe0 = feq9(m.x, m.y, m.z);
        for (var i = 0u; i < 9u; i = i + 1u) { neq[i] = neq[i] + ws[j]*(cf[i] - fe0[i]); }
    }
    var fe = feq9(rho, ux, uy);
    // неравновесная часть масштабируется: f^neq ~ tau*dt
    for (var i = 0u; i < 9u; i = i + 1u) { amr_dst[k*9u + i] = fe[i] + A.r01*neq[i]; }
}

@compute @workgroup_size(64)
fn k_fill(@builtin(global_invocation_id) gid: vec3<u32>) {
    let k = gid.x;
    if (k >= A.nghost) { return; }
    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];
    }
}

@compute @workgroup_size(64)
fn k_restrict(@builtin(global_invocation_id) gid: vec3<u32>) {
    let k = gid.x;
    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]; }
    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]);
    }
}
"#;

// ─────────────────────────────────────────────────────────────────────────────
//  Уровень на устройстве
// ─────────────────────────────────────────────────────────────────────────────

struct GpuLevel {
    nx: usize,
    ny: usize,
    n: usize,
    f: wgpu::Buffer,
    /// Стабилизатор γ поузлово — нужен для картинки; забирается с устройства как есть.
    gam: wgpu::Buffer,
    bind: wgpu::BindGroup,
    /// число рабочих групп редукции статистики (оно же длина буфера частичных сумм)
    parts_count: u32,
    nlinks: u32,
}

impl GpuLevel {
    fn link_groups(&self) -> u32 {
        ceil_div(self.nlinks.max(1), WG)
    }
}

fn soa_equilibrium(n: usize, u0: (R, R)) -> Vec<f32> {
    let fe = math::feq(1.0, u0.0, u0.1);
    let mut v = vec![0.0f32; 9 * n];
    for i in 0..9 {
        for k in 0..n {
            v[i * n + k] = fe[i] as f32;
        }
    }
    v
}

/// Залить уровень на устройство: поля, маски, линки, β и групповая привязка.
fn make_level(
    device: &wgpu::Device,
    layout: &wgpu::BindGroupLayout,
    nx: usize,
    ny: usize,
    geom: &cpu::Geom,
    beta: &[f32],
    probe_node: u32,
    flags: u32,
    u0: (R, R),
) -> GpuLevel {
    let n = nx * ny;
    let init = soa_equilibrium(n, u0);
    let mkf = |label: &str| {
        device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
            label: Some(label),
            contents: bytemuck::cast_slice(&init),
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC
                | wgpu::BufferUsages::COPY_DST,
        })
    };
    let f = mkf("f");
    let post = mkf("post");
    let gam = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
        label: Some("gamma"),
        contents: bytemuck::cast_slice(&vec![2.0f32; n]),
        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
    });
    let solid: Vec<u32> = geom.solid.iter().map(|s| *s as u32).collect();
    let solid_buf = storage_init(device, "solid", bytemuck::cast_slice(&solid));
    let links: Vec<GLink> = geom
        .links
        .iter()
        .map(|l| GLink {
            node: l.node,
            far: l.far,
            i: l.i as u32,
            ib: l.ib as u32,
            kind: match l.kind {
                math::LinkKind::Near => 0,
                math::LinkKind::Far => 1,
                math::LinkKind::Simple => 2,
            },
            q: l.q as f32,
            _p: [0; 2],
        })
        .collect();
    let links_buf = storage_init(device, "links", bytemuck::cast_slice(&links));
    let beta_buf = storage_init(device, "beta", bytemuck::cast_slice(beta));

    let parts_count = ceil_div(n as u32, WG);
    let parts = storage(device, "parts", (parts_count as u64) * 32);
    let params = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
        label: Some("level params"),
        contents: bytemuck::bytes_of(&LevelParams {
            n: n as u32,
            nx: nx as u32,
            ny: ny as u32,
            nlinks: links.len() as u32,
            bcx: geom.body_cx as f32,
            bcy: geom.body_cy as f32,
            probe_node,
            flags,
        }),
        usage: wgpu::BufferUsages::UNIFORM,
    });

    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, &params),
        ],
    });
    GpuLevel { nx, ny, n, f, gam, bind, parts_count, nlinks: links.len() as u32 }
}

// ─────────────────────────────────────────────────────────────────────────────
//  Симуляция
// ─────────────────────────────────────────────────────────────────────────────

pub struct Sim {
    spec: Spec,
    device: wgpu::Device,
    queue: wgpu::Queue,
    adapter_name: String,

    l0: GpuLevel,
    l1: Option<GpuLevel>,
    pre: wgpu::Buffer,

    dyn_buf: wgpu::Buffer,
    results_buf: wgpu::Buffer,
    staging: wgpu::Buffer,
    /// по одной группе на подшаг — различаются только индексом подшага в `Substep`
    dyn_binds: Vec<wgpu::BindGroup>,

    amr: Option<AmrRes>,

    pipes: Pipes,
    readback: wgpu::Buffer,

    fluid_count: R,
    step_index: u64,
    d_ref: R,
}

struct AmrRes {
    nghost: u32,
    nrestrict: u32,
    bg_ghost_old: wgpu::BindGroup,
    bg_ghost_new: wgpu::BindGroup,
    bg_fill: Vec<wgpu::BindGroup>,
    bg_restrict: wgpu::BindGroup,
}

struct Pipes {
    collide: wgpu::ComputePipeline,
    stream: wgpu::ComputePipeline,
    bouzidi: wgpu::ComputePipeline,
    walls: wgpu::ComputePipeline,
    channel: wgpu::ComputePipeline,
    force: wgpu::ComputePipeline,
    stats1: wgpu::ComputePipeline,
    stats2: wgpu::ComputePipeline,
    ghost: wgpu::ComputePipeline,
    fill: wgpu::ComputePipeline,
    restrict: wgpu::ComputePipeline,
    empty_bg: wgpu::BindGroup,
}

impl Sim {
    pub fn new(spec: Spec) -> Result<Sim, String> {
        let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
            backends: wgpu::Backends::PRIMARY,
            ..Default::default()
        });
        let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions {
            power_preference: wgpu::PowerPreference::HighPerformance,
            compatible_surface: None,
            force_fallback_adapter: false,
        }))
        .ok_or("подходящий GPU-адаптер не найден (нужен Vulkan, DX12 или Metal)")?;
        let info = adapter.get_info();
        let adapter_name = format!("{} ({:?})", info.name, info.backend);

        let (device, queue) = pollster::block_on(adapter.request_device(
            &wgpu::DeviceDescriptor {
                label: Some("kbc2d"),
                required_features: wgpu::Features::empty(),
                // берём то, что реально умеет адаптер: связке уровня нужно 8 storage-биндингов,
                // downlevel-профиль даёт всего 4
                required_limits: adapter.limits(),
                memory_hints: wgpu::MemoryHints::Performance,
            },
            None,
        ))
        .map_err(|e| format!("не удалось получить устройство: {e}"))?;

        // Условие Града на GPU пока не перенесено. Молча считать по Bouzidi нельзя — это
        // была бы другая физика под тем же ключом, поэтому отказываемся явно.
        if spec.wall == math::WallModel::Grad {
            return Err("--wall grad на GPU не реализован (перенесён только Bouzidi). Запусти с --backend cpu либо оставь --wall bouzidi"
                .into());
        }

        // ── топология берётся из процессорного бэкенда, а не строится заново ──
        let geom0 = cpu::Geom::build(spec.nx, spec.ny, &spec.body);
        let fluid_count = geom0.solid.iter().filter(|s| !**s).count() as R;

        let beta0: Vec<f32> = (0..spec.nx)
            .map(|x| {
                if spec.sponge_len == 0 {
                    return spec.beta0 as f32;
                }
                let start = spec.nx - 1 - spec.sponge_len;
                let s = math::smoothstep((x as R - start as R) / spec.sponge_len as R);
                let nu = math::nu_of_beta(spec.beta0);
                math::beta_of_nu(nu * (1.0 + (spec.sponge_mult - 1.0) * s)) as f32
            })
            .collect();

        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("kbc2d.wgsl"),
            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(SHADER)),
        });

        let bgl_level = level_layout(&device);
        let bgl_dyn = dyn_layout(&device);
        let bgl_amr = amr_layout(&device);
        let bgl_empty = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("empty"),
            entries: &[],
        });

        let pl_level = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
            label: Some("level"),
            bind_group_layouts: &[&bgl_level, &bgl_dyn],
            push_constant_ranges: &[],
        });
        let pl_amr = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
            label: Some("amr"),
            bind_group_layouts: &[&bgl_empty, &bgl_empty, &bgl_amr],
            push_constant_ranges: &[],
        });

        let mk = |layout: &wgpu::PipelineLayout, entry: &str| {
            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                label: Some(entry),
                layout: Some(layout),
                module: &module,
                entry_point: entry,
                compilation_options: Default::default(),
                cache: None,
            })
        };

        // ── общие буферы ──
        let dyn_buf = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("dyn"),
            size: std::mem::size_of::<Dyn>() as u64,
            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });
        // 16 слотов итогов + 8 подшагов по 4 числа на силу
        let results_buf = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("results"),
            size: (16 + 8 * 4) * 4,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
            mapped_at_creation: false,
        });
        let staging = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("staging"),
            size: std::mem::size_of::<Results>() as u64,
            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let nsub = spec.refine.max(1);
        let dyn_binds: Vec<wgpu::BindGroup> = (0..nsub)
            .map(|s| {
                let sb = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
                    label: Some("substep"),
                    contents: bytemuck::bytes_of(&Substep { idx: s as u32, _p: [0; 3] }),
                    usage: wgpu::BufferUsages::UNIFORM,
                });
                device.create_bind_group(&wgpu::BindGroupDescriptor {
                    label: Some("dyn"),
                    layout: &bgl_dyn,
                    entries: &[bind(0, &dyn_buf), bind(1, &sb), bind(2, &results_buf)],
                })
            })
            .collect();

        // стартовое поле: либо сразу набегающий поток, либо покой (см. cpu.rs)
        let u0 = if spec.init_uniform {
            let (sn, cs) = spec.flow_angle.sin_cos();
            (spec.units.u_lat * cs, spec.units.u_lat * sn)
        } else {
            (0.0, 0.0)
        };

        // ── зонд ──
        let (px, py) = spec.probe;
        let probe_on_fine = matches!(spec.patch, Some((ax, bx, ay, by))
            if px >= ax && px <= bx && py >= ay && py <= by)
            && spec.refine > 1;

        // ── уровень 0 ──
        let probe0 = if probe_on_fine { 0 } else { py * spec.nx + px };
        let flags0 = 1 | if probe_on_fine { 0 } else { 2 };
        let l0 = make_level(
            &device,
            &bgl_level,
            spec.nx,
            spec.ny,
            &geom0,
            &beta0,
            probe0 as u32,
            flags0,
            u0,
        );

        let pre = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("pre"),
            size: (9 * l0.n * 4) as u64,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        // ── уровень 1 и AMR ──
        let mut l1 = None;
        let mut amr = None;
        if spec.refine > 1 {
            let (ax, bx, ay, by) = spec.patch.ok_or("refine > 1 требует патч")?;
            let r = spec.refine;
            let body1 = spec.body.refined(r as R, ax as R, ay as R);
            let nfx = r * (bx - ax) + 1;
            let nfy = r * (by - ay) + 1;
            let geom1 = cpu::Geom::build(nfx, nfy, &body1);
            let tau0 = 1.0 / (2.0 * spec.beta0);
            let tau1 = r as R * (tau0 - 0.5) + 0.5;
            let beta1 = vec![(1.0 / (2.0 * tau1)) as f32; nfx];
            let r01 = tau1 / (r as R * tau0);
            let patch = cpu::Patch::new(&spec, &geom0, &geom1.solid, r01);

            let probe1 = if probe_on_fine { ((py - ay) * r) * nfx + (px - ax) * r } else { 0 };
            let flags1 = if probe_on_fine { 2 } else { 0 };
            let lvl1 =
                make_level(&device, &bgl_level, nfx, nfy, &geom1, &beta1, probe1 as u32, flags1, u0);

            let ghosts: Vec<GGhost> = patch
                .ghosts()
                .iter()
                .map(|g| GGhost {
                    fine: g.fine,
                    c00: g.c00,
                    c10: g.c10,
                    c01: g.c01,
                    c11: g.c11,
                    tx: g.tx as f32,
                    ty: g.ty as f32,
                    _p: 0,
                })
                .collect();
            let rest: Vec<[u32; 2]> =
                patch.restrict_pairs().iter().map(|&(c, fi)| [c, fi]).collect();

            let ghost_bytes = (ghosts.len() * 9 * 4) as u64;
            let gh_old = storage(&device, "gh_old", ghost_bytes);
            let gh_new = storage(&device, "gh_new", ghost_bytes);
            let ghosts_buf = storage_init(&device, "ghosts", bytemuck::cast_slice(&ghosts));
            let rest_buf = storage_init(&device, "rest", bytemuck::cast_slice(&rest));
            let dummy = storage(&device, "dummy", 16);

            let base = AmrParams {
                nghost: ghosts.len() as u32,
                nrestrict: rest.len() as u32,
                ccount: l0.n as u32,
                fcount: lvl1.n as u32,
                r01: r01 as f32,
                rfc: (1.0 / r01) as f32,
                w: 0.0,
                _pad: 0.0,
            };
            let mk_amr_bg = |src: &wgpu::Buffer,
                             dst: &wgpu::Buffer,
                             a: &wgpu::Buffer,
                             b: &wgpu::Buffer,
                             params: AmrParams| {
                let ub = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
                    label: Some("amr params"),
                    contents: bytemuck::bytes_of(&params),
                    usage: wgpu::BufferUsages::UNIFORM,
                });
                device.create_bind_group(&wgpu::BindGroupDescriptor {
                    label: Some("amr"),
                    layout: &bgl_amr,
                    entries: &[
                        bind(0, src),
                        bind(1, dst),
                        bind(2, &ghosts_buf),
                        bind(3, a),
                        bind(4, b),
                        bind(5, &rest_buf),
                        bind(6, &ub),
                    ],
                })
            };

            let bg_ghost_old = mk_amr_bg(&pre, &gh_old, &dummy, &dummy, base);
            let bg_ghost_new = mk_amr_bg(&l0.f, &gh_new, &dummy, &dummy, base);
            let bg_fill: Vec<wgpu::BindGroup> = (0..r)
                .map(|s| {
                    let mut p = base;
                    p.w = ((s + 1) as R / r as R) as f32;
                    mk_amr_bg(&dummy, &lvl1.f, &gh_old, &gh_new, p)
                })
                .collect();
            let bg_restrict = mk_amr_bg(&lvl1.f, &l0.f, &dummy, &dummy, base);

            amr = Some(AmrRes {
                nghost: ghosts.len() as u32,
                nrestrict: rest.len() as u32,
                bg_ghost_old,
                bg_ghost_new,
                bg_fill,
                bg_restrict,
            });
            l1 = Some(lvl1);
        }

        let pipes = Pipes {
            collide: mk(&pl_level, "k_collide"),
            stream: mk(&pl_level, "k_stream"),
            bouzidi: mk(&pl_level, "k_bouzidi"),
            walls: mk(&pl_level, "k_walls"),
            channel: mk(&pl_level, "k_channel"),
            force: mk(&pl_level, "k_force"),
            stats1: mk(&pl_level, "k_stats1"),
            stats2: mk(&pl_level, "k_stats2"),
            ghost: mk(&pl_amr, "k_ghost"),
            fill: mk(&pl_amr, "k_fill"),
            restrict: mk(&pl_amr, "k_restrict"),
            empty_bg: device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some("empty"),
                layout: &bgl_empty,
                entries: &[],
            }),
        };

        let readback = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("readback"),
            size: (9 * l0.n * 4) as u64,
            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let d_ref = match spec.patch {
            Some(_) if spec.refine > 1 => spec.body.d * spec.refine as R,
            _ => spec.body.d,
        };

        Ok(Sim {
            spec,
            device,
            queue,
            adapter_name,
            l0,
            l1,
            pre,
            dyn_buf,
            results_buf,
            staging,
            dyn_binds,
            amr,
            pipes,
            readback,
            fluid_count,
            step_index: 0,
            d_ref,
        })
    }

    pub fn name(&self) -> &'static str {
        // сигнатура бэкенда фиксирована; конкретный адаптер печатается отдельно
        Box::leak(format!("GPU: {} , f32", self.adapter_name).into_boxed_str())
    }

    pub fn force_ref_size(&self) -> R {
        self.d_ref
    }

    fn inlet(&self, t: u64) -> (R, R) {
        let sp = &self.spec;
        let ramp = if sp.init_uniform {
            1.0
        } else {
            math::smoothstep(t as R / sp.ramp.max(1) as R)
        };
        let u = sp.units.u_lat * ramp;
        let (s, c) = sp.flow_angle.sin_cos();
        let mut uy = u * s;
        if sp.pert_dur > 0 && t >= sp.ramp && t < sp.ramp + sp.pert_dur {
            let ph = (t - sp.ramp) as R / sp.pert_dur as R;
            uy += sp.pert_amp * sp.units.u_lat * (std::f64::consts::PI * ph).sin();
        }
        (u * c, uy)
    }

    pub fn step(&mut self) -> StepRec {
        let t = self.step_index;
        let (ux_in, uy_in) = self.inlet(t);
        let refine = self.spec.refine.max(1);
        self.queue.write_buffer(
            &self.dyn_buf,
            0,
            bytemuck::bytes_of(&Dyn {
                ux_in: ux_in as f32,
                uy_in: uy_in as f32,
                rho_out: 1.0,
                kbc_model: (self.spec.kbc_model == math::KbcModel::N2) as u32,
                outlet_extrap: self.spec.outlet_extrapolate as u32,
                nparts: self.l0.parts_count,
                refine: refine as u32,
                collision: (self.spec.collision == Collision::Bgk) as u32,
            }),
        );

        let mut enc = self
            .device
            .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("step") });

        // «до» — состояние 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);
        }

        {
            let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                label: Some("L0"),
                timestamp_writes: None,
            });
            p.set_bind_group(0, &self.l0.bind, &[]);
            p.set_bind_group(1, &self.dyn_binds[0], &[]);
            let ncell = ceil_div(self.l0.n as u32, WG);
            p.set_pipeline(&self.pipes.collide);
            p.dispatch_workgroups(ncell, 1, 1);
            p.set_pipeline(&self.pipes.stream);
            p.dispatch_workgroups(ncell, 1, 1);
            p.set_pipeline(&self.pipes.bouzidi);
            p.dispatch_workgroups(self.l0.link_groups(), 1, 1);
            // порядок обязателен: стенки снимают заворот по y, затем Zou-He — по x, на весь столбец
            p.set_pipeline(&self.pipes.walls);
            p.dispatch_workgroups(ceil_div(self.l0.nx as u32, WG), 1, 1);
            p.set_pipeline(&self.pipes.channel);
            p.dispatch_workgroups(ceil_div(self.l0.ny as u32, WG), 1, 1);
        }

        if let (Some(l1), Some(a)) = (self.l1.as_ref(), self.amr.as_ref()) {
            {
                let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                    label: Some("ghost"),
                    timestamp_writes: None,
                });
                p.set_bind_group(0, &self.pipes.empty_bg, &[]);
                p.set_bind_group(1, &self.pipes.empty_bg, &[]);
                p.set_pipeline(&self.pipes.ghost);
                p.set_bind_group(2, &a.bg_ghost_old, &[]);
                p.dispatch_workgroups(ceil_div(a.nghost, WG), 1, 1);
                p.set_bind_group(2, &a.bg_ghost_new, &[]);
                p.dispatch_workgroups(ceil_div(a.nghost, WG), 1, 1);
            }
            let ncell1 = ceil_div(l1.n as u32, WG);
            let nlink1 = l1.link_groups();
            for s in 0..refine {
                {
                    let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                        label: Some("L1"),
                        timestamp_writes: None,
                    });
                    p.set_bind_group(0, &l1.bind, &[]);
                    p.set_bind_group(1, &self.dyn_binds[s], &[]);
                    p.set_pipeline(&self.pipes.collide);
                    p.dispatch_workgroups(ncell1, 1, 1);
                    p.set_pipeline(&self.pipes.stream);
                    p.dispatch_workgroups(ncell1, 1, 1);
                    p.set_pipeline(&self.pipes.bouzidi);
                    p.dispatch_workgroups(nlink1, 1, 1);
                    // силу снимаем на каждом подшаге, усредняется она в k_stats2
                    p.set_pipeline(&self.pipes.force);
                    p.dispatch_workgroups(1, 1, 1);
                }
                {
                    let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                        label: Some("fill"),
                        timestamp_writes: None,
                    });
                    p.set_bind_group(0, &self.pipes.empty_bg, &[]);
                    p.set_bind_group(1, &self.pipes.empty_bg, &[]);
                    p.set_bind_group(2, &a.bg_fill[s], &[]);
                    p.set_pipeline(&self.pipes.fill);
                    p.dispatch_workgroups(ceil_div(a.nghost, WG), 1, 1);
                }
            }
            {
                let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                    label: Some("restrict"),
                    timestamp_writes: None,
                });
                p.set_bind_group(0, &self.pipes.empty_bg, &[]);
                p.set_bind_group(1, &self.pipes.empty_bg, &[]);
                p.set_bind_group(2, &a.bg_restrict, &[]);
                p.set_pipeline(&self.pipes.restrict);
                p.dispatch_workgroups(ceil_div(a.nrestrict, WG), 1, 1);
            }
        } else {
            let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                label: Some("force L0"),
                timestamp_writes: None,
            });
            p.set_bind_group(0, &self.l0.bind, &[]);
            p.set_bind_group(1, &self.dyn_binds[0], &[]);
            p.set_pipeline(&self.pipes.force);
            p.dispatch_workgroups(1, 1, 1);
        }

        {
            let mut p = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                label: Some("stats"),
                timestamp_writes: None,
            });
            p.set_bind_group(1, &self.dyn_binds[0], &[]);
            p.set_bind_group(0, &self.l0.bind, &[]);
            p.set_pipeline(&self.pipes.stats1);
            p.dispatch_workgroups(self.l0.parts_count, 1, 1);
            if let Some(l1) = self.l1.as_ref() {
                // на тонком уровне stats1 нужен только чтобы снять зонд (флаг записи сумм снят)
                p.set_bind_group(0, &l1.bind, &[]);
                p.dispatch_workgroups(ceil_div(l1.n as u32, WG), 1, 1);
                p.set_bind_group(0, &self.l0.bind, &[]);
            }
            p.set_pipeline(&self.pipes.stats2);
            p.dispatch_workgroups(1, 1, 1);
        }

        enc.copy_buffer_to_buffer(
            &self.results_buf,
            0,
            &self.staging,
            0,
            std::mem::size_of::<Results>() as u64,
        );
        self.queue.submit(Some(enc.finish()));

        let res: Results = self.read_staging();
        self.step_index += 1;

        let cnt = if res.cnt > 0.0 { res.cnt as R } else { self.fluid_count };
        StepRec {
            step: t,
            fx: res.fx as R,
            fy: res.fy as R,
            tz: res.tz as R,
            uy_probe: res.uy_probe as R,
            rho_mean: res.rho_sum as R / self.fluid_count,
            max_u: res.max_u as R,
            gamma_mean: res.g_sum as R / cnt,
            gamma_min: res.g_min as R,
            gamma_max: res.g_max as R,
            degenerate_frac: res.degen as R / cnt,
            xi_negative_frac: res.xineg as R / cnt,
        }
    }

    fn read_staging(&self) -> Results {
        let slice = self.staging.slice(..);
        let (tx, rx) = std::sync::mpsc::channel();
        slice.map_async(wgpu::MapMode::Read, move |r| {
            let _ = tx.send(r);
        });
        self.device.poll(wgpu::Maintain::Wait);
        let out = match rx.recv() {
            Ok(Ok(())) => {
                let data = slice.get_mapped_range();
                *bytemuck::from_bytes::<Results>(&data[..std::mem::size_of::<Results>()])
            }
            _ => Results::default(),
        };
        self.staging.unmap();
        out
    }

    /// Скачать популяции L0 на хост (нужно для кадров и проверки на NaN).
    fn download_l0(&self) -> Vec<f32> {
        self.download(&self.l0.f, (9 * self.l0.n * 4) as u64)
    }

    /// Скачать произвольный буфер уровня. Буфер приёма выделен один раз при сборке под самый
    /// большой запрос (популяции L0); для меньших читается только его начало.
    fn download(&self, src: &wgpu::Buffer, bytes: u64) -> Vec<f32> {
        let rb = &self.readback;
        let mut enc = self
            .device
            .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("dl") });
        enc.copy_buffer_to_buffer(src, 0, rb, 0, bytes);
        self.queue.submit(Some(enc.finish()));

        let slice = rb.slice(..);
        let (tx, rx) = std::sync::mpsc::channel();
        slice.map_async(wgpu::MapMode::Read, move |r| {
            let _ = tx.send(r);
        });
        self.device.poll(wgpu::Maintain::Wait);
        let out = match rx.recv() {
            Ok(Ok(())) => {
                let m = slice.get_mapped_range();
                bytemuck::cast_slice::<u8, f32>(&m[..bytes as usize]).to_vec()
            }
            _ => vec![f32::NAN; (bytes / 4) as usize],
        };
        rb.unmap();
        out
    }

    pub fn is_finite(&self) -> bool {
        self.download_l0().iter().all(|v| v.is_finite())
    }

    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 raw = self.download_l0();
        let node = |k: usize| -> [R; 9] {
            std::array::from_fn(|i| raw[i * n + k] as R)
        };
        let mut out = vec![0.0; n];
        match kind {
            FieldKind::Speed => {
                for k in 0..n {
                    let (_, ux, uy) = math::macros(&node(k));
                    out[k] = (ux * ux + uy * uy).sqrt();
                }
            }
            FieldKind::Density => {
                for k in 0..n {
                    out[k] = math::macros(&node(k)).0;
                }
            }
            FieldKind::Vorticity => {
                let mut ux = vec![0.0; n];
                let mut uy = vec![0.0; n];
                for k in 0..n {
                    let (_, a, b) = math::macros(&node(k));
                    ux[k] = a;
                    uy[k] = b;
                }
                for y in 0..ny {
                    for x in 0..nx {
                        let xp = (x + 1).min(nx - 1);
                        let xm = x.saturating_sub(1);
                        let yp = (y + 1).min(ny - 1);
                        let ym = y.saturating_sub(1);
                        let dvdx = (uy[y * nx + xp] - uy[y * nx + xm]) / (xp - xm).max(1) as R;
                        let dudy = (ux[yp * nx + x] - ux[ym * nx + x]) / (yp - ym).max(1) as R;
                        out[y * nx + x] = dvdx - dudy;
                    }
                }
            }
            FieldKind::Gamma => {
                // берём ровно то, что посчитало устройство: пересчёт на хосте с базовым β
                // соврал бы всюду, где β — поле (включённая губка)
                let g = self.download(&self.l0.gam, (n * 4) as u64);
                for k in 0..n {
                    out[k] = g[k] as R;
                }
            }
        }
        let solid = cpu::Geom::build(nx, ny, &self.spec.body).solid;
        (out, solid)
    }
}

// ─────────────────────────────────────────────────────────────────────────────
//  Мелкие помощники
// ─────────────────────────────────────────────────────────────────────────────

fn ceil_div(a: u32, b: u32) -> u32 {
    (a + b - 1) / b
}

fn bind(binding: u32, buf: &wgpu::Buffer) -> wgpu::BindGroupEntry<'_> {
    wgpu::BindGroupEntry { binding, resource: buf.as_entire_binding() }
}

fn storage(device: &wgpu::Device, label: &str, size: u64) -> wgpu::Buffer {
    device.create_buffer(&wgpu::BufferDescriptor {
        label: Some(label),
        size: size.max(16),
        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
        mapped_at_creation: false,
    })
}

fn storage_init(device: &wgpu::Device, label: &str, data: &[u8]) -> wgpu::Buffer {
    let pad;
    let data = if data.is_empty() {
        pad = [0u8; 16];
        &pad[..]
    } else {
        data
    };
    device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
        label: Some(label),
        contents: data,
        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
    })
}

fn ro(binding: u32) -> wgpu::BindGroupLayoutEntry {
    entry(binding, wgpu::BufferBindingType::Storage { read_only: true })
}
fn rw(binding: u32) -> wgpu::BindGroupLayoutEntry {
    entry(binding, wgpu::BufferBindingType::Storage { read_only: false })
}
fn un(binding: u32) -> wgpu::BindGroupLayoutEntry {
    entry(binding, wgpu::BufferBindingType::Uniform)
}
fn entry(binding: u32, ty: wgpu::BufferBindingType) -> wgpu::BindGroupLayoutEntry {
    wgpu::BindGroupLayoutEntry {
        binding,
        visibility: wgpu::ShaderStages::COMPUTE,
        ty: wgpu::BindingType::Buffer { ty, has_dynamic_offset: false, min_binding_size: None },
        count: None,
    }
}

fn level_layout(device: &wgpu::Device) -> wgpu::BindGroupLayout {
    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)],
    })
}
fn dyn_layout(device: &wgpu::Device) -> wgpu::BindGroupLayout {
    device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
        label: Some("dyn"),
        entries: &[un(0), un(1), rw(2)],
    })
}
fn amr_layout(device: &wgpu::Device) -> wgpu::BindGroupLayout {
    device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
        label: Some("amr"),
        entries: &[ro(0), rw(1), ro(2), ro(3), ro(4), ro(5), un(6)],
    })
}