Hugging Face's logo

NeuralVerified
/
neural-physics-engine

PyTorch
physics-simulation
projective-dynamics
neural-physics
graphics
fluid-simulation
soft-body
from-scratch
research
Model card Files
xet
 Community
neural-physics-engine / dam_break_gpu.html
Quazim0t0's picture
Quazim0t0
Import from Quazim0t0/neural-physics-engine; repoint refs to NeuralVerified
24c2ab9 verified
Raw
History Blame Contribute Delete
31.8 kB
<!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>