/ proof-of-concept / steelsim41
Sign in
proof-of-concept / steelsim41
steelsim41 / src / sim / shaders / step.wgsl
318 lines · 15.0 KBBlameHistoryRaw
1// One explicit time step (symplectic Euler) of the von Karman plate with a
2// thickness field.
3//
4// State layout (structure of arrays, n = nx * ny nodes):
5// pos: [u | v | w | wLo] wLo carries the rounding error of the w update
6// vel: [du | dv | dw]
7// A workgroup of 16x16 threads advances a 12x12 tile of nodes in three
8// phases separated by barriers. (1) Every thread loads one entry of the tile
9// plus its 2-node halo (16x16 values of w, u, v, w0, the thickness, and the
10// node weights) into shared memory. (2) The threads compute, once each
11// and again into shared memory, the membrane forces of the 13x13 cells
12// around the tile and the weighted bending moments and mixed differences
13// that the bending stencil gathers. (3) Each of the 144 tile nodes evaluates
14// its update from shared memory alone. On grids of a few thousand nodes the
15// kernel is bound by the latency of one workgroup rather than by throughput,
16// so the shared phases, which replace a tenfold per-node recomputation of
17// the stencil ingredients, are what make a step cheap.
18//
19// The thickness enters as the local bending stiffness D(h) at the nodes and
20// the cells, the local membrane stiffnesses at the cells (mean thickness of
21// the four corners), and the node mass. Nodes outside the plate (weight 0)
22// are skipped: a second difference exists only where the node and both
23// neighbors belong to the plate, and a cell contributes only when its four
24// corners do. That reproduces the free edge of the rectangle and gives a
25// staircase free edge for any other outline, including a center hole.
27@group(0) @binding(0) var<uniform> P: Params;
28@group(0) @binding(1) var<uniform> SP: StepParams;
29@group(0) @binding(2) var<storage, read> posIn: array<f32>;
30@group(0) @binding(3) var<storage, read> velIn: array<f32>;
31@group(0) @binding(4) var<storage, read_write> posOut: array<f32>;
32@group(0) @binding(5) var<storage, read_write> velOut: array<f32>;
33@group(0) @binding(6) var<storage, read_write> hist: array<f32>; // [slot * n + node]: transverse velocity at the audio sample instants
34// Two 4-float slots of drumstick state (position, velocity, force, max force).
35@group(0) @binding(7) var<storage, read_write> stick: array<f32>;
37const TILE = 12; // nodes per workgroup along each axis (must match simulation.ts)
38const HALO = 2;
39const TW = 16; // TILE + 2 * HALO: the loaded tile, one thread per entry
40const CT = 13; // TILE + 1: cells adjacent to the tile's nodes
41const MT = 14; // TILE + 2: nodes at which the bending moments are needed
42// Tile with halo, local index (j - oy) * TW + (i - ox).
43var<workgroup> tw: array<f32, 256>; // w
44var<workgroup> tu: array<f32, 256>; // u
45var<workgroup> tv: array<f32, 256>; // v
46var<workgroup> tw0: array<f32, 256>; // stress-free shape
47var<workgroup> th: array<f32, 256>; // thickness
48var<workgroup> tn: array<u32, 64>; // node weights in quarters (0..4), four to a word
49// Membrane force contributions of the CT x CT cells, 4 nodes x 3 components each.
50var<workgroup> cellF: array<f32, 2028>;
51// Weighted bending moments at the MT x MT inner nodes and mixed differences at the CT x CT cells.
52var<workgroup> ma: array<f32, 196>;
53var<workgroup> mb: array<f32, 196>;
54var<workgroup> cc: array<f32, 169>;
56// Tile origin in global node coordinates (set once per invocation).
57var<private> ox: i32;
58var<private> oy: i32;
59var<private> ihx2: f32;
60var<private> ihy2: f32;
61var<private> ihxy: f32;
63fn id(i: i32, j: i32) -> u32 { return u32(j) * P.nx + u32(i); }
64fn lidx(i: i32, j: i32) -> u32 { return u32(j - oy) * u32(TW) + u32(i - ox); }
65fn midx(i: i32, j: i32) -> u32 { return u32(j - oy - 1) * u32(MT) + u32(i - ox - 1); }
66fn cidx(ci: i32, cj: i32) -> u32 { return u32(cj - oy - 1) * u32(CT) + u32(ci - ox - 1); }
67// Tile lookups with global indices; the tile covers [ox, ox + TW) x [oy, oy + TW).
68fn W(i: i32, j: i32) -> f32 { return tw[lidx(i, j)]; }
69fn U(i: i32, j: i32) -> f32 { return tu[lidx(i, j)]; }
70fn V(i: i32, j: i32) -> f32 { return tv[lidx(i, j)]; }
71fn W0(i: i32, j: i32) -> f32 { return tw0[lidx(i, j)]; }
72// Transverse velocity, read from global memory (only the loss Laplacian needs it,
73// and keeping it out of shared memory keeps the workgroup within 16 KB).
74fn DW(i: i32, j: i32) -> f32 { return velIn[2u * P.n + id(i, j)]; }
75fn TH(i: i32, j: i32) -> f32 { return th[lidx(i, j)]; }
76fn NWq(i: i32, j: i32) -> u32 { let l = lidx(i, j); return (tn[l >> 2u] >> ((l & 3u) * 8u)) & 0xFFu; }
77fn NW(i: i32, j: i32) -> f32 { return 0.25 * f32(NWq(i, j)); }
78fn INS(i: i32, j: i32) -> bool { return NWq(i, j) != 0u; }
79fn insCell(ci: i32, cj: i32) -> bool { return INS(ci, cj) && INS(ci + 1, cj) && INS(ci + 1, cj + 1) && INS(ci, cj + 1); }
80// Bending stiffness at a node and at a cell (mean thickness of the four corners).
81fn Dn(i: i32, j: i32) -> f32 { let h = TH(i, j); return P.Db * h * h * h; }
82fn cellH(ci: i32, cj: i32) -> f32 { return 0.25 * (TH(ci, cj) + TH(ci + 1, cj) + TH(ci + 1, cj + 1) + TH(ci, cj + 1)); }
84// Graph Laplacian of the transverse velocity at node (i, j), over neighbors in the plate.
85fn lapDW(i: i32, j: i32) -> f32 {
86 let c = DW(i, j);
87 var lap = 0.0;
88 if (INS(i - 1, j)) { lap += (DW(i - 1, j) - c) * ihx2; }
89 if (INS(i + 1, j)) { lap += (DW(i + 1, j) - c) * ihx2; }
90 if (INS(i, j - 1)) { lap += (DW(i, j - 1) - c) * ihy2; }
91 if (INS(i, j + 1)) { lap += (DW(i, j + 1) - c) * ihy2; }
92 return lap;
95// Weighted bending moments at nodes; zero where the second difference is undefined.
96fn Ma(i: i32, j: i32) -> f32 {
97 if (!(INS(i - 1, j) && INS(i, j) && INS(i + 1, j))) { return 0.0; }
98 let a = (W(i + 1, j) - 2.0 * W(i, j) + W(i - 1, j)) * ihx2;
99 var b = 0.0;
100 if (INS(i, j - 1) && INS(i, j + 1)) { b = (W(i, j + 1) - 2.0 * W(i, j) + W(i, j - 1)) * ihy2; }
101 return NW(i, j) * Dn(i, j) * (a + P.nu * b);
103fn Mb(i: i32, j: i32) -> f32 {
104 if (!(INS(i, j - 1) && INS(i, j) && INS(i, j + 1))) { return 0.0; }
105 let b = (W(i, j + 1) - 2.0 * W(i, j) + W(i, j - 1)) * ihy2;
106 var a = 0.0;
107 if (INS(i - 1, j) && INS(i + 1, j)) { a = (W(i + 1, j) - 2.0 * W(i, j) + W(i - 1, j)) * ihx2; }
108 return NW(i, j) * Dn(i, j) * (b + P.nu * a);
110// Weighted mixed second difference at a cell center; zero outside the plate.
111fn Cc(ci: i32, cj: i32) -> f32 {
112 if (!insCell(ci, cj)) { return 0.0; }
113 let h = cellH(ci, cj);
114 return P.Db * h * h * h * (W(ci + 1, cj + 1) - W(ci + 1, cj) - W(ci, cj + 1) + W(ci, cj)) * ihxy;
117// Membrane forces of the bilinear cell (ci, cj) on its four nodes (2x2 Gauss),
118// written to cellF at local cell index lc.
119fn cellForces(ci: i32, cj: i32, lc: u32) {
120 let base = lc * 12u;
121 if (!insCell(ci, cj)) {
122 for (var q = 0u; q < 12u; q++) { cellF[base + q] = 0.0; }
123 return;
124 }
125 let hc = cellH(ci, cj);
126 let C = P.Cm * hc;
127 let G = P.Gm * hc;
128 let Un = vec4f(U(ci, cj), U(ci + 1, cj), U(ci + 1, cj + 1), U(ci, cj + 1));
129 let Vn = vec4f(V(ci, cj), V(ci + 1, cj), V(ci + 1, cj + 1), V(ci, cj + 1));
130 let W0n = vec4f(W0(ci, cj), W0(ci + 1, cj), W0(ci + 1, cj + 1), W0(ci, cj + 1));
131 let Wn = vec4f(W(ci, cj), W(ci + 1, cj), W(ci + 1, cj + 1), W(ci, cj + 1)) + W0n; // total shape
132 let q = 0.25 * P.hx * P.hy;
133 let ihx = 1.0 / P.hx; let ihy = 1.0 / P.hy;
134 let g3 = 0.57735026918962576;
135 var Fu = vec4f(0.0); var Fv = vec4f(0.0); var Fw = vec4f(0.0);
136 for (var gp = 0; gp < 4; gp++) {
137 var xi = -g3; if ((gp & 1) != 0) { xi = g3; }
138 var eta = -g3; if ((gp & 2) != 0) { eta = g3; }
139 let dNdx = vec4f(-(1.0 - eta), 1.0 - eta, 1.0 + eta, -(1.0 + eta)) * (0.5 * ihx);
140 let dNdy = vec4f(-(1.0 - xi), -(1.0 + xi), 1.0 + xi, 1.0 - xi) * (0.5 * ihy);
141 let ux = dot(dNdx, Un); let uy = dot(dNdy, Un);
142 let vx = dot(dNdx, Vn); let vy = dot(dNdy, Vn);
143 let wx = dot(dNdx, Wn); let wy = dot(dNdy, Wn);
144 let w0x = dot(dNdx, W0n); let w0y = dot(dNdy, W0n);
145 // Marguerre strains: total slope minus the stress-free slope.
146 let exx = ux + 0.5 * (wx * wx - w0x * w0x);
147 let eyy = vy + 0.5 * (wy * wy - w0y * w0y);
148 let gxy = uy + vx + wx * wy - w0x * w0y;
149 let Nxx = C * (exx + P.nu * eyy);
150 let Nyy = C * (eyy + P.nu * exx);
151 let Nxy = G * gxy;
152 Fu -= q * (Nxx * dNdx + Nxy * dNdy);
153 Fv -= q * (Nyy * dNdy + Nxy * dNdx);
154 Fw -= q * ((Nxx * wx + Nxy * wy) * dNdx + (Nxy * wx + Nyy * wy) * dNdy);
155 }
156 for (var a = 0u; a < 4u; a++) {
157 cellF[base + a * 3u] = Fu[a];
158 cellF[base + a * 3u + 1u] = Fv[a];
159 cellF[base + a * 3u + 2u] = Fw[a];
160 }
163// True for the plate nodes of the drumstick contact patch (same rule as patchAround in params.ts).
164fn inContactPatch(i: i32, j: i32) -> bool {
165 if (nodeWeight(i, j) == 0.0) { return false; }
166 let dxs = nodeX(i) - P.strikeX; let dys = nodeY(j) - P.strikeY;
167 return dxs * dxs + dys * dys <= P.strikeR2 || id(i, j) == P.strikeNode;
169// True for the plate nodes of the mount patch.
170fn inMount(i: i32, j: i32) -> bool {
171 if (P.mountOn == 0u || nodeWeight(i, j) == 0.0) { return false; }
172 let dxm = nodeX(i) - P.mountX; let dym = nodeY(j) - P.mountY;
173 return dxm * dxm + dym * dym <= P.mountR2 || id(i, j) == P.mountNode;
176// Stored membrane force on local node `a` of cell (ci, cj).
177fn cellForce(ci: i32, cj: i32, a: u32) -> vec3f {
178 let base = cidx(ci, cj) * 12u + a * 3u;
179 return vec3f(cellF[base], cellF[base + 1u], cellF[base + 2u]);
182@compute @workgroup_size(16, 16)
183fn main(@builtin(local_invocation_id) lid: vec3u, @builtin(workgroup_id) wid: vec3u) {
184 let n = P.n;
185 ox = i32(wid.x) * TILE - HALO;
186 oy = i32(wid.y) * TILE - HALO;
187 ihx2 = 1.0 / (P.hx * P.hx); ihy2 = 1.0 / (P.hy * P.hy); ihxy = 1.0 / (P.hx * P.hy);
188 let lt = lid.y * 16u + lid.x;
190 // Phase 1: one entry of the tile with halo per thread.
191 {
192 let gi = ox + i32(lid.x); let gj = oy + i32(lid.y);
193 var ww = 0.0; var uu = 0.0; var vv = 0.0; var z0 = 0.0; var hh = 0.0;
194 if (gi >= 0 && gi < i32(P.nx) && gj >= 0 && gj < i32(P.ny)) {
195 let k = id(gi, gj);
196 let st = nd[k];
197 ww = posIn[2u * n + k]; uu = posIn[k]; vv = posIn[n + k]; z0 = st.x; hh = st.z;
198 }
199 tw[lt] = ww; tu[lt] = uu; tv[lt] = vv; tw0[lt] = z0; th[lt] = hh;
200 if (lt < 64u) {
201 var word = 0u;
202 for (var q = 0u; q < 4u; q++) {
203 let e = lt * 4u + q;
204 let ei = ox + i32(e % u32(TW)); let ej = oy + i32(e / u32(TW));
205 var wq = 0u;
206 if (ei >= 0 && ei < i32(P.nx) && ej >= 0 && ej < i32(P.ny)) { wq = u32(round(4.0 * nd[id(ei, ej)].y)); }
207 word |= wq << (q * 8u);
208 }
209 tn[lt] = word;
210 }
211 }
212 workgroupBarrier();
214 // Phase 2: membrane forces of the CT x CT cells, bending moments at the MT x MT
215 // inner nodes, mixed differences at the cells; one item per thread each.
216 if (lt < 169u) {
217 let ci = ox + 1 + i32(lt % u32(CT)); let cj = oy + 1 + i32(lt / u32(CT));
218 cellForces(ci, cj, lt);
219 cc[lt] = Cc(ci, cj);
220 }
221 if (lt < 196u) {
222 let mi = ox + 1 + i32(lt % u32(MT)); let mj = oy + 1 + i32(lt / u32(MT));
223 ma[lt] = Ma(mi, mj);
224 mb[lt] = Mb(mi, mj);
225 }
226 workgroupBarrier();
228 // Phase 3: the update of the TILE x TILE nodes (no barriers past this point).
229 if (lt >= 144u) { return; }
230 let i = ox + HALO + i32(lt % u32(TILE)); let j = oy + HALO + i32(lt / u32(TILE));
231 if (i >= i32(P.nx) || j >= i32(P.ny)) { return; }
232 let k = id(i, j);
233 let weight = NW(i, j);
234 if (weight == 0.0) {
235 // Outside the plate: the node stays at rest.
236 posOut[k] = 0.0; posOut[n + k] = 0.0; posOut[2u * n + k] = 0.0; posOut[3u * n + k] = 0.0;
237 velOut[k] = 0.0; velOut[n + k] = 0.0; velOut[2u * n + k] = 0.0;
238 return;
239 }
241 // Elastic forces: bending from the discrete Kirchhoff energy, membrane from the cells.
242 let bend = (ma[midx(i - 1, j)] - 2.0 * ma[midx(i, j)] + ma[midx(i + 1, j)]) * ihx2
243 + (mb[midx(i, j - 1)] - 2.0 * mb[midx(i, j)] + mb[midx(i, j + 1)]) * ihy2
244 + 2.0 * (1.0 - P.nu) * (cc[cidx(i, j)] - cc[cidx(i - 1, j)] - cc[cidx(i, j - 1)] + cc[cidx(i - 1, j - 1)]) * ihxy;
245 var F = vec3f(0.0, 0.0, -P.hx * P.hy * bend);
246 F += cellForce(i, j, 0u);
247 F += cellForce(i - 1, j, 1u);
248 F += cellForce(i - 1, j - 1, 2u);
249 F += cellForce(i, j - 1, 3u);
250 let pOld = vec3f(posIn[k], posIn[n + k], posIn[2u * n + k]);
251 let vel = vec3f(velIn[k], velIn[n + k], velIn[2u * n + k]);
252 // The mount: a spring and a dashpot to ground over the mount patch, on all three components.
253 let mounted = inMount(i, j);
254 if (mounted) { F -= P.mountKw * weight * pOld + P.mountCw * weight * vel; }
255 let mass = P.rho * TH(i, j) * P.hx * P.hy * weight;
256 var acc = F * (1.0 / mass);
257 // Optional scaling of the in-plane inertia (lengthens the stable time step).
258 acc.x *= P.invInPlaneMass; acc.y *= P.invInPlaneMass;
260 // Losses: frequency-independent on all components, the graph Laplacian of the
261 // transverse velocity (~k^2), and an extra frequency-independent in-plane loss.
262 acc -= 2.0 * SP.sigma0 * vel;
263 acc.z += 2.0 * SP.sigma1 * lapDW(i, j);
264 acc.x -= 2.0 * SP.sigmaIn * vel.x;
265 acc.y -= 2.0 * SP.sigmaIn * vel.y;
267 // Drumstick: Hertz contact between a free mass and the contact patch (the
268 // plate nodes within the tip radius, plus the nearest one). The contact
269 // coordinate is the mass-weighted mean displacement of the patch, and the
270 // force is spread over the patch in proportion to the node masses, so the
271 // pair is conjugate and the contact has a potential energy.
272 if ((SP.flags & 1u) != 0u) {
273 let sIn = 4u * SP.parity;
274 if (P.mountOn == 0u) {
275 // Free plate: inertia relief. The force of the previous step is applied
276 // to the whole plate as the rigid acceleration field that balances the
277 // patch load and its moments about the center, so the plate gains no net
278 // momentum (the thread would take it up in reality). The one-step lag
279 // integrates to zero. A mounted plate takes the reaction through the mount.
280 let Fprev = stick[sIn + 2u];
281 acc.z += Fprev * (1.0 / P.mass + nodeY(j) * P.patchCY / P.Ix + nodeX(i) * P.patchCX / P.Iy);
282 }
283 if (inContactPatch(i, j)) {
284 let sx = stick[sIn]; let sv = stick[sIn + 1u]; let sMax = stick[sIn + 3u];
285 var wc = 0.0;
286 for (var pj = i32(P.patchJ0); pj <= i32(P.patchJ1); pj++) {
287 for (var pi = i32(P.patchI0); pi <= i32(P.patchI1); pi++) {
288 if (inContactPatch(pi, pj)) { wc += nodeMass(pi, pj) * posIn[2u * n + id(pi, pj)]; }
289 }
290 }
291 wc /= P.patchMass;
292 let pen = wc - sx; // the stick sits on the +z side and moves toward -z
293 var Fc = 0.0;
294 if (pen > 0.0) { Fc = P.K * pow(pen, P.alpha); }
295 acc.z -= Fc / P.patchMass;
296 if (k == P.strikeNode) {
297 let sOut = 4u * (1u - SP.parity);
298 let svn = sv + P.dt * Fc / P.stickMass;
299 stick[sOut] = sx + P.dt * svn;
300 stick[sOut + 1u] = svn;
301 stick[sOut + 2u] = Fc;
302 stick[sOut + 3u] = max(sMax, Fc);
303 }
304 }
305 }
307 let vNew = vel + P.dt * acc;
308 if (SP.audioSlot >= 0) { hist[u32(SP.audioSlot) * n + k] = vNew.z; }
309 var pNew = pOld + P.dt * vNew;
310 // Compensated update of w: the increment dt * v can be far below the
311 // float32 resolution of w, so the rounding error is carried separately.
312 let y = P.dt * vNew.z + posIn[3u * n + k];
313 let t = pOld.z + y;
314 let lo = (pOld.z - t) + y;
315 pNew.z = t;
316 posOut[k] = pNew.x; posOut[n + k] = pNew.y; posOut[2u * n + k] = pNew.z; posOut[3u * n + k] = lo;
317 velOut[k] = vNew.x; velOut[n + k] = vNew.y; velOut[2u * n + k] = vNew.z;
↑↓ move↵openescclose