File size: 31,834 Bytes

24c2ab9

<!DOCTYPE html>
<html lang="en">
<head>
<meta charset="utf-8">
<title>Neural Physics Engine — W8 GPU Dam Break</title>
<style>
  html,body{margin:0;height:100%;overflow:hidden;background:#06080f;
    font-family:'Segoe UI',system-ui,sans-serif;color:#cfe3ff}
  #app{position:fixed;inset:0}
  .hud{position:fixed;top:14px;left:14px;z-index:10;background:rgba(8,12,24,.72);
    border:1px solid rgba(90,160,255,.25);border-radius:12px;padding:14px 18px;
    backdrop-filter:blur(6px);min-width:240px;user-select:none}
  #fps{font-size:44px;font-weight:700;line-height:1;letter-spacing:-1px;
    font-variant-numeric:tabular-nums}
  #fps small{font-size:14px;font-weight:400;color:#7f9cc4;margin-left:4px}
  #spark{display:block;margin:6px 0 8px}
  .row{display:flex;justify-content:space-between;font-size:12.5px;
    color:#9db8dd;margin:2px 0;font-variant-numeric:tabular-nums}
  .row b{color:#e8f2ff;font-weight:600}
  .badge{display:inline-block;background:#0f3d2e;color:#7dffb0;
    border:1px solid #2a7a56;border-radius:5px;font-size:10px;padding:1px 6px;
    vertical-align:middle;margin-left:6px}
  .btns{margin-top:10px;display:flex;flex-wrap:wrap;gap:6px}
  button{background:#14213c;color:#cfe3ff;border:1px solid #33507f;
    border-radius:7px;padding:5px 10px;font-size:12px;cursor:pointer}
  button:hover{background:#1d3055}
  button.on{background:#1b4d3e;border-color:#37a06f;color:#b8ffdd}
  .title{position:fixed;top:16px;right:18px;text-align:right;z-index:10;
    pointer-events:none}
  .title h1{margin:0;font-size:17px;font-weight:600;color:#e8f2ff}
  .title p{margin:3px 0 0;font-size:11.5px;color:#7f9cc4}
  .foot{position:fixed;bottom:12px;right:18px;font-size:11px;color:#5f7ba3;
    z-index:10;pointer-events:none}
  #err{position:fixed;bottom:12px;left:14px;color:#ff7a7a;font-size:13px;
    max-width:70%;z-index:20;white-space:pre-wrap}
</style>
<script type="importmap">
{"imports":{
  "three":"https://unpkg.com/three@0.160.0/build/three.module.js",
  "three/addons/":"https://unpkg.com/three@0.160.0/examples/jsm/"
}}
</script>
</head>
<body>
<div id="app"></div>
<div class="hud">
  <div id="fps">--<small>fps</small><span class="badge">GPU compute</span></div>
  <canvas id="spark" width="200" height="30"></canvas>
  <div class="row"><span>frame CPU cost</span><b id="msTot">-- ms</b></div>
  <div class="row"><span>&nbsp;&nbsp;GPU encode+readback</span><b id="msFl">-- ms</b></div>
  <div class="row"><span>&nbsp;&nbsp;soft body (CPU)</span><b id="msSb">-- ms</b></div>
  <div class="row"><span>particles</span><b id="nPart">--</b></div>
  <div class="row"><span>tets</span><b id="nTet">--</b></div>
  <div class="row"><span>PD iters/step</span><b id="pdIt">--</b></div>
  <div class="btns">
    <button id="bReset">↺ Reset</button>
    <button id="bPause">⏸ Pause</button>
    <button id="bWarm" class="on">warm start: ON</button>
  </div>
  <div class="btns">
    <button id="b10k" class="on">~10k</button><button id="b25k">~25k</button>
    <button id="b50k">~50k</button>
  </div>
</div>
<div class="title">
  <h1>Neural Physics Engine — W8</h1>
  <p>GPU dam break on a soft body<br>
     WebGPU compute PBF · corotated PD jelly · two-way coupling<br>
     per-particle loop is embarrassingly parallel — roadmap §5.6</p>
</div>
<div class="foot">drag to orbit · scroll to zoom</div>
<div id="err"></div>

<script type="module">
import * as THREE from 'three';
import {OrbitControls} from 'three/addons/controls/OrbitControls.js';
import {RoomEnvironment} from 'three/addons/environments/RoomEnvironment.js';

const errBox=document.getElementById('err');
window.onerror=(m,s,l)=>errBox.textContent=m+' @'+l;

/* ================= parameters ================= */
const TANK={x:1.2,y:0.85,z:0.6};
// dt 1/120 + physical velocity cap: the old cap 0.45*H/DT shrank with
// particle spacing and choked the dam-break front (~4 m/s) to ~1.3 m/s
// at 10k particles — that is what made the water pour like syrup
const G=-9.81, DT=1/120, RHO0=1000, PBF_ITERS=3, XSPH=0.03, VCAP=5.0;
const PRESETS={b10k:0.0253, b25k:0.0186, b50k:0.0147};
let spacing=PRESETS.b10k;
const REACT_SCALE=1e6;

/* ================= WGSL ================= */
const WGSL=`
struct U {
  lo: vec4f, hi: vec4f, grav: vec4f,
  grid: vec4u,          // gx, gy, gz, cells
  p0: vec4f,            // H, H2, POLY6, SPIKY
  p1: vec4f,            // MASS, RHO0, EPSL, DT
  p2: vec4f,            // VMAX, DPMAX, RC, XSPH
  cnt: vec4u,           // N, NV, CAP, pad
};
@group(0) @binding(0) var<uniform> u: U;
@group(0) @binding(1) var<storage, read_write> P: array<vec4f>;
@group(0) @binding(2) var<storage, read_write> Q: array<vec4f>;
@group(0) @binding(3) var<storage, read_write> V: array<vec4f>;
@group(0) @binding(4) var<storage, read_write> DP: array<vec4f>;  // xyz dp, w lambda
@group(0) @binding(5) var<storage, read_write> cellCnt: array<atomic<u32>>;
@group(0) @binding(6) var<storage, read_write> cellPart: array<u32>;
@group(0) @binding(7) var<storage, read_write> jelly: array<vec4f>;
@group(0) @binding(8) var<storage, read_write> react: array<atomic<i32>>;

fn cellOf(p: vec3f) -> vec3u {
  let c = clamp(vec3i((p - u.lo.xyz)/u.p0.x),
                vec3i(0), vec3i(u.grid.xyz) - vec3i(1));
  return vec3u(c);
}
fn cellIdx(c: vec3u) -> u32 { return c.x + c.y*u.grid.x + c.z*u.grid.x*u.grid.y; }

@compute @workgroup_size(64) fn predict(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  var v = V[i].xyz + u.grav.xyz*u.p1.w;
  let s = length(v);
  if(s>u.p2.x){ v *= u.p2.x/s; }
  V[i] = vec4f(v,0);
  Q[i] = vec4f(clamp(P[i].xyz + v*u.p1.w, u.lo.xyz, u.hi.xyz), 0);
}
@compute @workgroup_size(64) fn clearGrid(@builtin(global_invocation_id) g: vec3u){
  if(g.x<u.grid.w){ atomicStore(&cellCnt[g.x],0u); }
}
@compute @workgroup_size(64) fn bin(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  let k=cellIdx(cellOf(Q[i].xyz));
  let slot=atomicAdd(&cellCnt[k],1u);
  if(slot<u.cnt.z){ cellPart[k*u.cnt.z+slot]=i; }
}
@compute @workgroup_size(64) fn densityLambda(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  let qi=Q[i].xyz;
  let H=u.p0.x; let H2=u.p0.y; let POLY6=u.p0.z; let SPIKY=u.p0.w;
  let MASS=u.p1.x; let mr=MASS/u.p1.y;
  var rho = MASS*POLY6*H2*H2*H2;
  var sum = vec3f(0); var g2 = 0.0;
  let c=vec3i(cellOf(qi));
  for(var a=-1;a<=1;a++){ for(var b=-1;b<=1;b++){ for(var d=-1;d<=1;d++){
    let cc=c+vec3i(a,b,d);
    if(any(cc<vec3i(0)) || any(cc>=vec3i(u.grid.xyz))){continue;}
    let k=cellIdx(vec3u(cc));
    let n=min(atomicLoad(&cellCnt[k]),u.cnt.z);
    for(var e=0u;e<n;e++){
      let j=cellPart[k*u.cnt.z+e];
      if(j==i){continue;}
      let dv=qi-Q[j].xyz;
      let r2=dot(dv,dv);
      if(r2>=H2){continue;}
      let t=H2-r2;
      rho += MASS*POLY6*t*t*t;
      if(r2<1e-12){continue;}
      let r=sqrt(r2); let hr=H-r;
      let gw=(SPIKY*hr*hr/r)*mr*dv;
      sum += gw; g2 += dot(gw,gw);
    }
  }}}
  g2 += dot(sum,sum);
  let lam = -(rho/u.p1.y - 1.0)/(g2 + u.p1.z);
  DP[i]=vec4f(0,0,0,lam);
}
@compute @workgroup_size(64) fn deltaP(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  let qi=Q[i].xyz; let li=DP[i].w;
  let H=u.p0.x; let H2=u.p0.y; let SPIKY=u.p0.w; let mr=u.p1.x/u.p1.y;
  var dp=vec3f(0);
  let c=vec3i(cellOf(qi));
  for(var a=-1;a<=1;a++){ for(var b=-1;b<=1;b++){ for(var d=-1;d<=1;d++){
    let cc=c+vec3i(a,b,d);
    if(any(cc<vec3i(0)) || any(cc>=vec3i(u.grid.xyz))){continue;}
    let k=cellIdx(vec3u(cc));
    let n=min(atomicLoad(&cellCnt[k]),u.cnt.z);
    for(var e=0u;e<n;e++){
      let j=cellPart[k*u.cnt.z+e];
      if(j==i){continue;}
      let dv=qi-Q[j].xyz;
      let r2=dot(dv,dv);
      if(r2>=H2 || r2<1e-12){continue;}
      let r=sqrt(r2); let hr=H-r;
      let gw=(SPIKY*hr*hr/r)*mr*dv;
      dp += (li+DP[j].w)*gw;
    }
  }}}
  let l=length(dp);
  if(l>u.p2.y){ dp *= u.p2.y/l; }   // overshoot guard (W7 lesson)
  DP[i]=vec4f(dp,li);
}
@compute @workgroup_size(64) fn apply(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  Q[i]=vec4f(clamp(Q[i].xyz+DP[i].xyz, u.lo.xyz, u.hi.xyz), 0);
}
@compute @workgroup_size(64) fn couple(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  var qi=Q[i].xyz;
  let rc=u.p2.z; let rc2=rc*rc;
  for(var v=0u;v<u.cnt.y;v++){
    let dv=qi-jelly[v].xyz;
    let d2=dot(dv,dv);
    if(d2>=rc2 || d2<1e-12){continue;}
    let dl=sqrt(d2); let ov=(rc-dl)/dl;
    qi += 0.62*ov*dv;
    if(jelly[v].w>0.5){continue;}   // pinned vertex: no reaction
    let r=-0.10*ov*dv*${REACT_SCALE}.0;
    atomicAdd(&react[v*3u+0u], i32(r.x));
    atomicAdd(&react[v*3u+1u], i32(r.y));
    atomicAdd(&react[v*3u+2u], i32(r.z));
  }
  Q[i]=vec4f(clamp(qi,u.lo.xyz,u.hi.xyz),0);
}
@compute @workgroup_size(64) fn velocity(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  V[i]=vec4f((Q[i].xyz-P[i].xyz)/u.p1.w, 0);
}
@compute @workgroup_size(64) fn xsph(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  let qi=Q[i].xyz; let vi=V[i].xyz;
  let H2=u.p0.y; let POLY6=u.p0.z; let mr=u.p1.x/u.p1.y;
  var acc=vec3f(0);
  let c=vec3i(cellOf(qi));
  for(var a=-1;a<=1;a++){ for(var b=-1;b<=1;b++){ for(var d=-1;d<=1;d++){
    let cc=c+vec3i(a,b,d);
    if(any(cc<vec3i(0)) || any(cc>=vec3i(u.grid.xyz))){continue;}
    let k=cellIdx(vec3u(cc));
    let n=min(atomicLoad(&cellCnt[k]),u.cnt.z);
    for(var e=0u;e<n;e++){
      let j=cellPart[k*u.cnt.z+e];
      if(j==i){continue;}
      let dv=qi-Q[j].xyz;
      let r2=dot(dv,dv);
      if(r2>=H2){continue;}
      let t=H2-r2;
      acc += (POLY6*t*t*t*mr)*(V[j].xyz-vi);
    }
  }}}
  DP[i]=vec4f(vi + u.p2.w*acc, 0);
}
@compute @workgroup_size(64) fn finalize(@builtin(global_invocation_id) g: vec3u){
  let i=g.x; if(i>=u.cnt.x){return;}
  let v=DP[i].xyz;
  V[i]=vec4f(v,0);
  P[i]=Q[i];
  Q[i]=vec4f(Q[i].xyz, length(v));   // pack speed for rendering
}
`;

/* ================= GPU setup ================= */
let device,pipes={},bindGroup,uniBuf,bufs={},NCELLS,CAP=24,N=0,NV=0;
let stagingQ=[],stagingR=[],stagingFree=[];
let renderPos;   // Float32Array view (N*4) fed to three.js

async function initGPU(){
  if(!navigator.gpu) throw new Error(
    'WebGPU not available in this browser — open dam_break_demo.html (WebGL/CPU version) instead.');
  const adapter=await navigator.gpu.requestAdapter();
  if(!adapter) throw new Error('No WebGPU adapter.');
  device=await adapter.requestDevice();
  device.lost.then(info=>{ errBox.textContent='GPU device lost: '+info.message; });
  device.addEventListener('uncapturederror',e=>{
    if(!errBox.textContent) errBox.textContent='WebGPU: '+e.error.message.slice(0,900);
  });
}

function makeFluidGPU(){
  const H=spacing*1.9, H2=H*H;
  const gx=Math.ceil(TANK.x/H), gy=Math.ceil(TANK.y/H), gz=Math.ceil(TANK.z/H);
  NCELLS=gx*gy*gz;
  // initial lattice
  const xs=[];
  for(let x=spacing/2;x<0.42;x+=spacing)
    for(let y=spacing/2;y<0.62;y+=spacing)
      for(let z=spacing/2;z<TANK.z;z+=spacing)
        xs.push(x+(Math.random()-.5)*1e-3, y+(Math.random()-.5)*1e-3,
                z+(Math.random()-.5)*1e-3, 0);
  N=xs.length/4;
  const pInit=new Float32Array(xs);
  // particle mass: crowded-lattice kernel sum (interior of a full lattice)
  const POLY6=315/(64*Math.PI*Math.pow(H,9));
  let smax=POLY6*H2*H2*H2;
  const reach=Math.ceil(H/spacing);
  for(let a=-reach;a<=reach;a++)for(let b=-reach;b<=reach;b++)
  for(let c=-reach;c<=reach;c++){
    if(!a&&!b&&!c) continue;
    const r2=(a*a+b*b+c*c)*spacing*spacing;
    if(r2<H2){ const t=H2-r2; smax+=POLY6*t*t*t; }
  }
  const MASS=RHO0/smax;
  const SPIKY=-45/(Math.PI*Math.pow(H,6));
  const EPSL=100*(MASS/RHO0)*(MASS/RHO0);
  const VMAX=VCAP, DPMAX=0.18*H;
  const RC=Math.max(0.06, 0.8*H);

  const mk=(size,usage)=>device.createBuffer({size,usage});
  const SU=GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_DST|GPUBufferUsage.COPY_SRC;
  bufs.P=mk(N*16,SU); bufs.Q=mk(N*16,SU); bufs.V=mk(N*16,SU); bufs.DP=mk(N*16,SU);
  bufs.cellCnt=mk(NCELLS*4,SU);
  bufs.cellPart=mk(NCELLS*CAP*4,SU);
  bufs.jelly=mk(Math.max(NV,1)*16,SU);
  bufs.react=mk(Math.max(NV,1)*12,SU);
  device.queue.writeBuffer(bufs.P,0,pInit);
  device.queue.writeBuffer(bufs.Q,0,pInit);
  device.queue.writeBuffer(bufs.V,0,new Float32Array(N*4));

  uniBuf=device.createBuffer({size:8*16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST});
  const uf=new Float32Array(32); const uu=new Uint32Array(uf.buffer);
  uf.set([0,0,0,0],0);                       // lo
  uf.set([TANK.x,TANK.y,TANK.z,0],4);        // hi
  uf.set([0,G,0,0],8);                       // grav
  uu.set([gx,gy,gz,NCELLS],12);              // grid
  uf.set([H,H2,POLY6,SPIKY],16);             // p0
  uf.set([MASS,RHO0,EPSL,DT],20);            // p1
  uf.set([VMAX,DPMAX,RC,XSPH],24);           // p2
  uu.set([N,NV,CAP,0],28);                   // cnt
  device.queue.writeBuffer(uniBuf,0,uf);

  const mod=device.createShaderModule({code:WGSL});
  // explicit shared layout: 'auto' would drop bindings unused by an entry
  // point and reject the common bind group
  const bgl=device.createBindGroupLayout({entries:[
    {binding:0,visibility:GPUShaderStage.COMPUTE,buffer:{type:'uniform'}},
    ...[1,2,3,4,5,6,7,8].map(b=>({binding:b,
      visibility:GPUShaderStage.COMPUTE,buffer:{type:'storage'}}))]});
  const pl=device.createPipelineLayout({bindGroupLayouts:[bgl]});
  for(const name of ['predict','clearGrid','bin','densityLambda','deltaP',
                     'apply','couple','velocity','xsph','finalize'])
    pipes[name]=device.createComputePipeline({layout:pl,
      compute:{module:mod,entryPoint:name}});
  const commonBG=device.createBindGroup({layout:bgl,
    entries:[{binding:0,resource:{buffer:uniBuf}},
      {binding:1,resource:{buffer:bufs.P}},{binding:2,resource:{buffer:bufs.Q}},
      {binding:3,resource:{buffer:bufs.V}},{binding:4,resource:{buffer:bufs.DP}},
      {binding:5,resource:{buffer:bufs.cellCnt}},
      {binding:6,resource:{buffer:bufs.cellPart}},
      {binding:7,resource:{buffer:bufs.jelly}},
      {binding:8,resource:{buffer:bufs.react}}]});
  bindGroup={};
  for(const name in pipes) bindGroup[name]=commonBG;

  stagingQ=[0,1,2].map(()=>mk(N*16,GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ));
  stagingR=[0,1,2].map(()=>mk(Math.max(NV,1)*12,GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ));
  stagingFree=[true,true,true];
  renderPos=new Float32Array(pInit);
}

/* ================= soft body (CPU, same as W7) ================= */
const SB={sp:0.075,nx:4,ny:5,nz:4,x0:0.78,z0:0.155,dens:400,
          alpha:0.32,tol:2e-4,maxIt:12};
let sbN,sbX,sbY,sbZ,sbVx,sbVy,sbVz,sbPin,sbM;
let tets=[],tetQ=[],tetQuat=[];
let sbC1x,sbC1y,sbC1z,sbC2x,sbC2y,sbC2z,sbFx,sbFy,sbFz,sbCorr,sbCnt,sbS;
let sbPx,sbPy,sbPz;   // previous Jacobi iterate (Chebyshev acceleration)
let warmStart=true,pdItersAvg=0;
let jellyUpload;

function vid(i,j,k){return (i*(SB.ny+1)+j)*(SB.nz+1)+k;}
function initSoft(){
  const {sp,nx,ny,nz,x0,z0}=SB;
  sbN=(nx+1)*(ny+1)*(nz+1); NV=sbN;
  sbX=new Float32Array(sbN);sbY=new Float32Array(sbN);sbZ=new Float32Array(sbN);
  for(let i=0;i<=nx;i++)for(let j=0;j<=ny;j++)for(let k=0;k<=nz;k++){
    const v=vid(i,j,k); sbX[v]=x0+i*sp; sbY[v]=j*sp; sbZ[v]=z0+k*sp;
  }
  sbVx=new Float32Array(sbN);sbVy=new Float32Array(sbN);sbVz=new Float32Array(sbN);
  sbPin=new Uint8Array(sbN);
  for(let i=0;i<=nx;i++)for(let k=0;k<=nz;k++) sbPin[vid(i,0,k)]=1;
  tets=[];
  for(let i=0;i<nx;i++)for(let j=0;j<ny;j++)for(let k=0;k<nz;k++){
    const c=[vid(i,j,k),vid(i+1,j,k),vid(i,j+1,k),vid(i,j,k+1),
             vid(i+1,j+1,k),vid(i+1,j,k+1),vid(i,j+1,k+1),vid(i+1,j+1,k+1)];
    const even=(i+j+k)%2===0;
    const T=even?[[0,1,2,3],[1,4,2,7],[1,3,5,7],[2,3,7,6],[1,2,3,7]]
                :[[1,0,4,5],[0,2,4,6],[0,3,5,6],[4,5,6,7],[0,4,6,5]];
    for(const t of T) tets.push([c[t[0]],c[t[1]],c[t[2]],c[t[3]]]);
  }
  sbM=new Float32Array(sbN); tetQ=[]; tetQuat=[];
  for(const t of tets){
    const cx=(sbX[t[0]]+sbX[t[1]]+sbX[t[2]]+sbX[t[3]])/4;
    const cy=(sbY[t[0]]+sbY[t[1]]+sbY[t[2]]+sbY[t[3]])/4;
    const cz=(sbZ[t[0]]+sbZ[t[1]]+sbZ[t[2]]+sbZ[t[3]])/4;
    tetQ.push(t.map(v=>[sbX[v]-cx,sbY[v]-cy,sbZ[v]-cz]));
    tetQuat.push([1,0,0,0]);
  }
  sbC1x=new Float32Array(sbN);sbC1y=new Float32Array(sbN);sbC1z=new Float32Array(sbN);
  sbC2x=new Float32Array(sbN);sbC2y=new Float32Array(sbN);sbC2z=new Float32Array(sbN);
  sbFx=new Float32Array(sbN);sbFy=new Float32Array(sbN);sbFz=new Float32Array(sbN);
  sbCorr=new Float32Array(sbN*3);sbCnt=new Float32Array(sbN);
  sbS=new Float32Array(sbN*3);
  sbPx=new Float32Array(sbN);sbPy=new Float32Array(sbN);sbPz=new Float32Array(sbN);
  jellyUpload=new Float32Array(sbN*4);
}
function extractR(A,q){
  for(let it=0;it<4;it++){
    const [w,x,y,z]=q;
    const R=[1-2*(y*y+z*z),2*(x*y-w*z),2*(x*z+w*y),
             2*(x*y+w*z),1-2*(x*x+z*z),2*(y*z-w*x),
             2*(x*z-w*y),2*(y*z+w*x),1-2*(x*x+y*y)];
    let ox=0,oy=0,oz=0,d=0;
    for(let c=0;c<3;c++){
      const rx=R[c],ry=R[3+c],rz=R[6+c],ax=A[c],ay=A[3+c],az=A[6+c];
      ox+=ry*az-rz*ay; oy+=rz*ax-rx*az; oz+=rx*ay-ry*ax;
      d+=rx*ax+ry*ay+rz*az;
    }
    const inv=1/(Math.abs(d)+1e-9);
    ox*=inv;oy*=inv;oz*=inv;
    const ang=Math.sqrt(ox*ox+oy*oy+oz*oz);
    if(ang<1e-7) break;
    const s=Math.sin(ang/2)/ang, cw=Math.cos(ang/2);
    const nw=cw*q[0]-s*(ox*q[1]+oy*q[2]+oz*q[3]);
    const nx2=cw*q[1]+s*(ox*q[0]+oy*q[3]-oz*q[2]);
    const ny2=cw*q[2]+s*(oy*q[0]+oz*q[1]-ox*q[3]);
    const nz2=cw*q[3]+s*(oz*q[0]+ox*q[2]-oy*q[1]);
    const n=1/Math.sqrt(nw*nw+nx2*nx2+ny2*ny2+nz2*nz2);
    q[0]=nw*n;q[1]=nx2*n;q[2]=ny2*n;q[3]=nz2*n;
  }
  const [w,x,y,z]=q;
  return [1-2*(y*y+z*z),2*(x*y-w*z),2*(x*z+w*y),
          2*(x*y+w*z),1-2*(x*x+z*z),2*(y*z-w*x),
          2*(x*z-w*y),2*(y*z+w*x),1-2*(x*x+y*y)];
}
function stepSoft(){
  const n=sbN;
  for(let v=0;v<n;v++){
    if(sbPin[v]){sbS[v*3]=sbX[v];sbS[v*3+1]=sbY[v];sbS[v*3+2]=sbZ[v];continue;}
    sbVy[v]+=G*DT;
    sbS[v*3]=sbX[v]+sbVx[v]*DT+sbFx[v];
    sbS[v*3+1]=sbY[v]+sbVy[v]*DT+sbFy[v];
    sbS[v*3+2]=sbZ[v]+sbVz[v]*DT+sbFz[v];
  }
  sbFx.fill(0);sbFy.fill(0);sbFz.fill(0);
  const ox=sbX.slice(),oy=sbY.slice(),oz=sbZ.slice();
  for(let v=0;v<n;v++){
    const wsx=warmStart?2*sbC1x[v]-sbC2x[v]:0;
    const wsy=warmStart?2*sbC1y[v]-sbC2y[v]:0;
    const wsz=warmStart?2*sbC1z[v]-sbC2z[v]:0;
    sbX[v]=sbS[v*3]+wsx;sbY[v]=sbS[v*3+1]+wsy;sbZ[v]=sbS[v*3+2]+wsz;
  }
  let iters=0;
  // Chebyshev-accelerated Jacobi (Wang 2015; same overshoot problem JGS2
  // arXiv:2506.06494 solves to second order — full JGS2 local ops are the
  // W9 candidate). rho: estimated Jacobi spectral radius; 2-iter delay.
  const RHO2=0.9*0.9; let om=1.0;
  sbPx.set(sbX);sbPy.set(sbY);sbPz.set(sbZ);
  for(let it=0;it<SB.maxIt;it++){
    iters++;
    if(it===2) om=2/(2-RHO2);
    else if(it>2) om=4/(4-RHO2*om);
    sbCorr.fill(0);sbCnt.fill(0);
    for(let ti=0;ti<tets.length;ti++){
      const t=tets[ti],Qo=tetQ[ti];
      const cx=(sbX[t[0]]+sbX[t[1]]+sbX[t[2]]+sbX[t[3]])/4;
      const cy=(sbY[t[0]]+sbY[t[1]]+sbY[t[2]]+sbY[t[3]])/4;
      const cz=(sbZ[t[0]]+sbZ[t[1]]+sbZ[t[2]]+sbZ[t[3]])/4;
      const A=[0,0,0,0,0,0,0,0,0];
      for(let vi=0;vi<4;vi++){
        const v=t[vi],q=Qo[vi];
        const dx=sbX[v]-cx,dy=sbY[v]-cy,dz=sbZ[v]-cz;
        A[0]+=dx*q[0];A[1]+=dx*q[1];A[2]+=dx*q[2];
        A[3]+=dy*q[0];A[4]+=dy*q[1];A[5]+=dy*q[2];
        A[6]+=dz*q[0];A[7]+=dz*q[1];A[8]+=dz*q[2];
      }
      const R=extractR(A,tetQuat[ti]);
      for(let vi=0;vi<4;vi++){
        const v=t[vi],q=Qo[vi];
        sbCorr[v*3]  +=cx+R[0]*q[0]+R[1]*q[1]+R[2]*q[2]-sbX[v];
        sbCorr[v*3+1]+=cy+R[3]*q[0]+R[4]*q[1]+R[5]*q[2]-sbY[v];
        sbCorr[v*3+2]+=cz+R[6]*q[0]+R[7]*q[1]+R[8]*q[2]-sbZ[v];
        sbCnt[v]++;
      }
    }
    let maxd=0;
    for(let v=0;v<n;v++){
      if(sbPin[v]||!sbCnt[v]) continue;
      const a=SB.alpha/sbCnt[v];
      const hx=sbX[v]+a*sbCorr[v*3]  +0.08*(sbS[v*3]  -sbX[v]);
      const hy=sbY[v]+a*sbCorr[v*3+1]+0.08*(sbS[v*3+1]-sbY[v]);
      const hz=sbZ[v]+a*sbCorr[v*3+2]+0.08*(sbS[v*3+2]-sbZ[v]);
      // Chebyshev blend of (xhat, x_k, x_{k-1})
      const nx2=om*(hx-sbX[v])+sbX[v]+(om-1)*(sbX[v]-sbPx[v]);
      const ny2=om*(hy-sbY[v])+sbY[v]+(om-1)*(sbY[v]-sbPy[v]);
      const nz2=om*(hz-sbZ[v])+sbZ[v]+(om-1)*(sbZ[v]-sbPz[v]);
      const dx2=nx2-sbX[v],dy2=ny2-sbY[v],dz2=nz2-sbZ[v];
      sbPx[v]=sbX[v];sbPy[v]=sbY[v];sbPz[v]=sbZ[v];
      sbX[v]=nx2;sbY[v]=Math.max(ny2,0);sbZ[v]=nz2;
      const d2=dx2*dx2+dy2*dy2+dz2*dz2;
      if(d2>maxd) maxd=d2;
    }
    if(Math.sqrt(maxd)<SB.tol) break;
  }
  pdItersAvg=pdItersAvg*0.95+iters*0.05;
  for(let v=0;v<n;v++){
    if(sbPin[v]) continue;
    sbC2x[v]=sbC1x[v];sbC2y[v]=sbC1y[v];sbC2z[v]=sbC1z[v];
    sbC1x[v]=sbX[v]-sbS[v*3];sbC1y[v]=sbY[v]-sbS[v*3+1];sbC1z[v]=sbZ[v]-sbS[v*3+2];
    sbVx[v]=0.998*(sbX[v]-ox[v])/DT;
    sbVy[v]=0.998*(sbY[v]-oy[v])/DT;
    sbVz[v]=0.998*(sbZ[v]-oz[v])/DT;
  }
}

/* ================= GPU frame ================= */
async function stepGPU(){
  // upload jelly vertices (w = pinned flag)
  for(let v=0;v<sbN;v++){
    jellyUpload[v*4]=sbX[v]; jellyUpload[v*4+1]=sbY[v];
    jellyUpload[v*4+2]=sbZ[v]; jellyUpload[v*4+3]=sbPin[v];
  }
  device.queue.writeBuffer(bufs.jelly,0,jellyUpload);
  device.queue.writeBuffer(bufs.react,0,new Int32Array(sbN*3));

  const enc=device.createCommandEncoder();
  const pass=enc.beginComputePass();
  const run=(name,count)=>{
    pass.setPipeline(pipes[name]); pass.setBindGroup(0,bindGroup[name]);
    pass.dispatchWorkgroups(Math.ceil(count/64));
  };
  run('predict',N);
  run('clearGrid',NCELLS);
  run('bin',N);
  for(let it=0;it<PBF_ITERS;it++){
    run('densityLambda',N); run('deltaP',N); run('apply',N);
  }
  run('couple',N);
  run('velocity',N);
  run('xsph',N);
  run('finalize',N);
  pass.end();
  // pipelined readback: never block the frame — map in the background and
  // render the freshest data that has arrived (1-2 frames of latency)
  const k=stagingFree.indexOf(true);
  if(k>=0){
    // capture THIS generation's free-array: after a reset, completions of
    // old maps must not free slots in the new staging set
    const myFree=stagingFree;
    myFree[k]=false;
    const sq=stagingQ[k], sr=stagingR[k], nHere=N, nvHere=sbN;
    enc.copyBufferToBuffer(bufs.Q,0,sq,0,nHere*16);
    enc.copyBufferToBuffer(bufs.react,0,sr,0,nvHere*12);
    device.queue.submit([enc.finish()]);
    Promise.all([sq.mapAsync(GPUMapMode.READ),sr.mapAsync(GPUMapMode.READ)])
      .then(()=>{
        if(renderPos.length===nHere*4)
          renderPos.set(new Float32Array(sq.getMappedRange()));
        const rr=new Int32Array(sr.getMappedRange());
        if(sbN===nvHere) for(let v=0;v<sbN;v++){
          sbFx[v]=rr[v*3]/REACT_SCALE;
          sbFy[v]=rr[v*3+1]/REACT_SCALE;
          sbFz[v]=rr[v*3+2]/REACT_SCALE;
        }
        sq.unmap(); sr.unmap();
        myFree[k]=true;
        if(interleaved) interleaved.needsUpdate=true;
      }).catch(()=>{ myFree[k]=true; });
  } else {
    device.queue.submit([enc.finish()]);
  }
  let bad=0,ymax=-9,ymin=9;
  for(let i=0;i<N;i++){
    const y=renderPos[i*4+1];
    if(!isFinite(y)) bad++;
    if(y>ymax) ymax=y; if(y<ymin) ymin=y;
  }
  window.__dbgGPU={bad,ymin,ymax,sample:Array.from(renderPos.slice(0,8))};
}

/* ================= rendering (three.js scene from W7) ================= */
const app=document.getElementById('app');
const renderer=new THREE.WebGLRenderer({antialias:true});
renderer.setPixelRatio(Math.min(devicePixelRatio,2));
renderer.setSize(innerWidth,innerHeight);
renderer.toneMapping=THREE.ACESFilmicToneMapping;
app.appendChild(renderer.domElement);
const scene=new THREE.Scene();
scene.background=new THREE.Color(0x06080f);
scene.fog=new THREE.Fog(0x06080f,3.5,8);
const pmrem=new THREE.PMREMGenerator(renderer);
scene.environment=pmrem.fromScene(new RoomEnvironment(),0.04).texture;
const camera=new THREE.PerspectiveCamera(40,innerWidth/innerHeight,0.05,50);
camera.position.set(2.15,1.15,1.9);
const controls=new OrbitControls(camera,renderer.domElement);
controls.target.set(TANK.x/2,0.28,TANK.z/2);
controls.enableDamping=true;controls.maxPolarAngle=1.52;
controls.minDistance=0.8;controls.maxDistance=6;
scene.add(new THREE.AmbientLight(0x334466,1.2));
const key=new THREE.DirectionalLight(0xbfd8ff,2.2);key.position.set(2,3,1.5);
scene.add(key);
const rim=new THREE.PointLight(0x2266ff,8,6);rim.position.set(-0.5,1.4,-0.8);
scene.add(rim);
const floor=new THREE.Mesh(new THREE.CircleGeometry(6,64),
  new THREE.MeshStandardMaterial({color:0x0a0f1e,roughness:.35,metalness:.6}));
floor.rotation.x=-Math.PI/2;floor.position.y=-0.002;scene.add(floor);
const grid=new THREE.GridHelper(8,40,0x1c2c50,0x101a33);
grid.position.y=0.001;scene.add(grid);
const tankGeo=new THREE.BoxGeometry(TANK.x,TANK.y,TANK.z);
const tankEdges=new THREE.LineSegments(new THREE.EdgesGeometry(tankGeo),
  new THREE.LineBasicMaterial({color:0x58a6ff,transparent:true,opacity:.85}));
tankEdges.position.set(TANK.x/2,TANK.y/2,TANK.z/2);scene.add(tankEdges);
const glass=new THREE.Mesh(tankGeo,new THREE.MeshPhysicalMaterial({
  color:0x88bbff,transparent:true,opacity:.05,roughness:.05,metalness:0,
  side:THREE.BackSide,depthWrite:false}));
glass.position.copy(tankEdges.position);scene.add(glass);

function makeSprite(){
  const c=document.createElement('canvas');c.width=c.height=64;
  const g=c.getContext('2d');
  const rg=g.createRadialGradient(28,26,4,32,32,30);
  rg.addColorStop(0,'rgba(255,255,255,1)');
  rg.addColorStop(.35,'rgba(255,255,255,.85)');
  rg.addColorStop(1,'rgba(255,255,255,0)');
  g.fillStyle=rg;g.fillRect(0,0,64,64);
  return new THREE.CanvasTexture(c);
}
let waterPts=null,waterGeo=null,interleaved=null;
function buildWater(){
  if(waterPts){scene.remove(waterPts);waterGeo.dispose();}
  waterGeo=new THREE.BufferGeometry();
  interleaved=new THREE.InterleavedBuffer(renderPos,4);
  interleaved.setUsage(THREE.DynamicDrawUsage);
  waterGeo.setAttribute('position',new THREE.InterleavedBufferAttribute(interleaved,3,0));
  waterGeo.setAttribute('aSpeed',new THREE.InterleavedBufferAttribute(interleaved,1,3));
  const mat=new THREE.ShaderMaterial({
    uniforms:{tex:{value:makeSprite()},size:{value:spacing*540}},
    vertexShader:`attribute float aSpeed; varying float vS;
      uniform float size;
      void main(){ vS=aSpeed;
        vec4 mv=modelViewMatrix*vec4(position,1.0);
        gl_PointSize=size/(-mv.z); gl_Position=projectionMatrix*mv; }`,
    fragmentShader:`uniform sampler2D tex; varying float vS;
      void main(){
        float a=texture2D(tex,gl_PointCoord).a; if(a<0.03) discard;
        vec3 deep=vec3(0.02,0.18,0.55), mid=vec3(0.05,0.55,0.95),
             foam=vec3(0.85,0.97,1.0);
        float s=clamp(vS*0.55,0.0,1.0);
        vec3 col=mix(deep,mid,smoothstep(0.0,0.55,s));
        col=mix(col,foam,smoothstep(0.55,1.0,s));
        gl_FragColor=vec4(col*a*1.25, a*0.9); }`,
    transparent:true,depthWrite:false,blending:THREE.AdditiveBlending});
  waterPts=new THREE.Points(waterGeo,mat);
  scene.add(waterPts);
}
let jellyMesh=null,jellyGeo=null;
function buildJelly(){
  if(jellyMesh){scene.remove(jellyMesh);jellyGeo.dispose();}
  const {nx,ny,nz}=SB;const idx=[];
  const quad=(a,b,c,d)=>idx.push(a,b,c,a,c,d);
  for(let i=0;i<nx;i++)for(let j=0;j<ny;j++){
    quad(vid(i,j,0),vid(i,j+1,0),vid(i+1,j+1,0),vid(i+1,j,0));
    quad(vid(i,j,nz),vid(i+1,j,nz),vid(i+1,j+1,nz),vid(i,j+1,nz));
  }
  for(let i=0;i<nx;i++)for(let k=0;k<nz;k++){
    quad(vid(i,0,k),vid(i+1,0,k),vid(i+1,0,k+1),vid(i,0,k+1));
    quad(vid(i,ny,k),vid(i,ny,k+1),vid(i+1,ny,k+1),vid(i+1,ny,k));
  }
  for(let j=0;j<ny;j++)for(let k=0;k<nz;k++){
    quad(vid(0,j,k),vid(0,j,k+1),vid(0,j+1,k+1),vid(0,j+1,k));
    quad(vid(nx,j,k),vid(nx,j+1,k),vid(nx,j+1,k+1),vid(nx,j,k+1));
  }
  jellyGeo=new THREE.BufferGeometry();
  jellyGeo.setAttribute('position',new THREE.BufferAttribute(new Float32Array(sbN*3),3));
  jellyGeo.setIndex(idx);
  jellyMesh=new THREE.Mesh(jellyGeo,new THREE.MeshPhysicalMaterial({
    color:0xff4f6e,roughness:.18,clearcoat:.8,clearcoatRoughness:.25,
    transmission:.25,thickness:.4,emissive:0x40060f,envMapIntensity:1.1}));
  scene.add(jellyMesh);
}
function syncMeshes(){
  interleaved.needsUpdate=true;
  const jp=jellyGeo.attributes.position.array;
  for(let v=0;v<sbN;v++){jp[v*3]=sbX[v];jp[v*3+1]=sbY[v];jp[v*3+2]=sbZ[v];}
  jellyGeo.attributes.position.needsUpdate=true;
  jellyGeo.computeVertexNormals();
}

/* ================= HUD + loop ================= */
const $=(id)=>document.getElementById(id);
let paused=false,msFl=0,msSb=0;
const frameTimes=new Array(90).fill(16.7);let ftIdx=0,lastT=performance.now();
const sparkCtx=$('spark').getContext('2d');
function drawSpark(){
  const c=sparkCtx;c.clearRect(0,0,200,30);
  c.beginPath();
  for(let i=0;i<90;i++){
    const t=frameTimes[(ftIdx+i)%90];
    const y=30-Math.min(28,t*(28/50));
    if(i===0)c.moveTo(i*200/89,y);else c.lineTo(i*200/89,y);
  }
  c.strokeStyle='#58a6ff';c.lineWidth=1.5;c.stroke();
  c.strokeStyle='rgba(120,255,170,.35)';
  const y60=30-16.7*(28/50);
  c.beginPath();c.moveTo(0,y60);c.lineTo(200,y60);c.stroke();
}
function resetAll(){
  window.__stage='reset-start';
  initSoft(); window.__stage='soft-done';
  makeFluidGPU(); window.__stage='fluid-done';
  buildWater();buildJelly();
  $('nPart').textContent=N.toLocaleString();
  $('nTet').textContent=tets.length;
  pdItersAvg=0;
}
$('bReset').onclick=resetAll;
$('bPause').onclick=()=>{paused=!paused;
  $('bPause').textContent=paused?'▶ Play':'⏸ Pause';
  $('bPause').classList.toggle('on',paused);};
$('bWarm').onclick=()=>{warmStart=!warmStart;
  $('bWarm').textContent='warm start: '+(warmStart?'ON':'OFF');
  $('bWarm').classList.toggle('on',warmStart);};
for(const id of ['b10k','b25k','b50k']) $(id).onclick=()=>{
  spacing=PRESETS[id];resetAll();
  for(const o of ['b10k','b25k','b50k']) $(o).classList.toggle('on',o===id);
};

let hudTick=0,stepping=false,simAcc=0;
async function animate(now){
  requestAnimationFrame(animate);
  const dt=now-lastT;lastT=now;
  frameTimes[ftIdx]=dt;ftIdx=(ftIdx+1)%90;
  if(!paused && !stepping){
    stepping=true;
    // real-time pacing: run as many fixed substeps as wall-clock time
    // demands (capped), so low FPS never turns into slow motion
    simAcc=Math.min(simAcc+dt,50);
    let nStep=0;
    const t0=performance.now();
    let tf=0;
    try{
      while(simAcc>=1000*DT && nStep<3){
        const a=performance.now();
        await stepGPU();
        tf+=performance.now()-a;
        stepSoft();
        simAcc-=1000*DT; nStep++;
      }
    }catch(e){ errBox.textContent=String(e); }
    const t2=performance.now();
    if(nStep){
      msFl=msFl*0.9+tf*0.1;
      msSb=msSb*0.9+(t2-t0-tf)*0.1;
      syncMeshes();
    }
    stepping=false;
  }
  controls.update();
  renderer.render(scene,camera);
  if(++hudTick%6===0){
    const avg=frameTimes.reduce((a,b)=>a+b,0)/90;
    const fps=1000/avg;
    const el=$('fps');
    el.firstChild.textContent=fps.toFixed(0);
    el.style.color=fps>48?'#7dffb0':(fps>28?'#ffd970':'#ff8080');
    $('msTot').textContent=(msFl+msSb).toFixed(1)+' ms';
    $('msFl').textContent=msFl.toFixed(1)+' ms';
    $('msSb').textContent=msSb.toFixed(1)+' ms';
    $('pdIt').textContent=pdItersAvg.toFixed(1)+(warmStart?' (warm)':' (cold)');
    drawSpark();
  }
}
addEventListener('resize',()=>{
  camera.aspect=innerWidth/innerHeight;camera.updateProjectionMatrix();
  renderer.setSize(innerWidth,innerHeight);
});

window.__stage='module-loaded';
initGPU().then(()=>{
  window.__stage='gpu-init';
  resetAll();
  window.__stage='reset-done';
  requestAnimationFrame(animate);
}).catch(e=>{errBox.textContent=String(e.message||e);});
</script>
</body>
</html>