WebGPU time loop for the cymbal2 shell model, with numbl setup and output stages
17 changed files+4504−0
.gitignoreadded+2−0View file
@@ -0,0 +1,2 @@
1+node_modules/
2+*.wav
README.mdadded+71−0View file
@@ -0,0 +1,71 @@
1+# cymbal2-webgpu
2+
3+A limitation of the [cymbal2](https://vault1.magland.org/simulated_instruments/cymbal2) instrument is render time: the sound is a nonlinear time-domain integration of a von Karman shell, several hundred thousand strictly sequential steps, and the browser app executes those steps through numbl's interpreter at roughly 1.9 ms per step, which puts a full default render near twenty minutes. Here we explore moving the time loop, and only the time loop, onto the GPU through WebGPU, keeping numbl as the runtime for everything around it. On the modest integrated GPU of a laptop (Intel Iris Xe, running at its 300 MHz idle clock throughout, see below), the same render takes 83 microseconds per step: 48 seconds for the full six-second crash, about 4.4x faster than MATLAB R2026b, 7.5x faster than the numbl CLI with its native addon, and 23x faster than the numbl-in-browser path the webapp actually uses. The audio is single precision throughout; the contact diagnostics match the f64 reference to 0.03%, first contact lands on exactly the same step, and the acoustic rms agrees within 3%, which is within the tolerance the cymbal2 validation suite already grants the nonlinear trajectory.
4+
5+## Division of labor
6+
7+The design principle is that the MATLAB code in `simulated_instruments/cymbal2/matlab` remains the single source of truth, and the GPU replaces only `cymbal2_sim.m`, the step loop. numbl executes the unmodified setup stage (`cymbal2_system`, `cymbal2_mallet`, and the stroke schedule block copied verbatim from `cymbal2_sim`) and hands the resulting tensors to the host, which packs them into one read-only storage buffer; numbl likewise executes the unmodified output stage (`cymbal2_output`) on the acceleration signal the GPU returns. The WGSL mirrors the step structure of `cymbal2_sim.m` line for line: the same implicit per-wavenumber update, the same boundary rows (already baked into the stored inverses), the same Hertz contact, hand state machine, and beam update, with the sparse radial and azimuthal difference matrices applied as stencils and the dense transforms, per-wavenumber solves, and beam matrices applied as stored matrices. Scalar parameters, node indices, and the launch-step thresholds of the stroke schedule are baked into the generated shader source, and the thresholds are computed on the host in f64 with exactly the comparison `cymbal2_sim` uses, so the gesture is step-for-step aligned with the CPU reference.
8+
9+```
10+numbl (setup: grid, per-wavenumber inverses, mallet, schedule)
11+ -> WebGPU (the time loop: 8 compute dispatches per step)
12+ -> numbl (output: DC blocker, decimation, normalize)
13+```
14+
15+Usage, from this directory (`npm install` first; the `webgpu` package is prebuilt Google Dawn):
16+
17+```
18+node src/run.mjs # full default render -> cymbal2-gpu.wav
19+node src/run.mjs --pinned # the validation suite's reduced configuration
20+node src/run.mjs --compare # GPU vs the unmodified cymbal2_sim through numbl
21+node src/run.mjs --mode mega # the single-workgroup variant, for comparison
22+node src/run.mjs --velocity 4 --count 3 --rate 2 # parameter overrides
23+```
24+
25+## The step as eight dispatches
26+
27+One time step is eight compute dispatches at full GPU width, with the dispatch boundaries providing the global synchronization the stages need:
28+
29+| stage | work |
30+|---|---|
31+| AB | hand state machine, Hertz contact, and beam update (workgroup 0); shell derivative stencils and the Airy right-hand side (all workgroups) |
32+| C | azimuthal analysis of the Airy right-hand side |
33+| D | Airy solve, one small radial system per wavenumber |
34+| E | azimuthal synthesis of the stress function |
35+| FG | the von Karman bracket and forcing, stress-function stencils recomputed inline |
36+| H | azimuthal analysis of the nonlinear load |
37+| I | the implicit shell update, three small radial products per wavenumber |
38+| J | azimuthal synthesis of the new displacement; pickup acceleration |
39+
40+Because every recorded dispatch must be identical, the step index k is the sum of a chunk base written by the host into the state buffer once per 1024 steps and a within-chunk offset read from a dynamic-offset uniform, and all slot rotation is a pure function of k: the three-generation displacement and wavenumber-state slots rotate every step whether or not the shell is active, which is harmless before first contact because every slot is still zero, and exact afterwards because every active step writes. The shell-active flag latches inside AB; the stages of the same step read it one dispatch later, so they always observe the latch, and the one benign race inside AB itself can only skip work on a field that is still identically zero. The nonlinear trajectory of a cymbal is chaotic, so f32 cannot track the f64 reference for long; what the comparison pins down is the deterministic part (contact step, peak force, contact speed, the first few hundred steps of the acceleration) and the statistics of the rest.
41+
42+## What we measured, including what failed
43+
44+The exploration tried three architectures, and the measurements on this device (Intel Iris Xe integrated, Mesa/Dawn, one workgroup occupying at most one of six subslices) are worth recording plainly.
45+
46+**A single-workgroup megakernel** (`--mode mega`), which loops over thousands of steps inside one dispatch with workgroup barriers between stages, was the first correct implementation and the natural design, since it puts nothing on the CPU and nothing in the dispatch queue. It runs at 396 us/step at workgroup size 256 and 151 us/step at 1024, and profiling showed why: barriers cost almost nothing (about 0.03 us each), but a single workgroup carries too few hardware threads to hide storage-load latency, so the dense 60-point reductions run at about 6 GFLOP/s no matter how they are written. We tried register-accumulator chains, fully unrolled scalar accumulators, vec4 loads with padded layouts, and shared-memory staging; none moved the number, because the bottleneck is latency times too-few-threads, not instruction count. Notably the megakernel keeps the GPU at its full 1300 MHz clock, being continuously busy, and is still slower than the multi-dispatch backend running at 300 MHz.
47+
48+**The multi-dispatch backend** (the default) is the opposite trade: full width, but a dispatch boundary between stages. Per-dispatch overhead is about 2 us, and the eight stages sum to 83 us/step. The costs are now spread across the whole GPU, and the same reduction that was latency-bound in one workgroup is unremarkable across twenty-five.
49+
50+**A persistent kernel** (`--mode persistent`), thirteen workgroups launched once and synchronized between stages by a global spin barrier on storage atomics, would remove the dispatch boundaries entirely and was the most interesting candidate. It runs without deadlock on this device, and it is wrong: outputs are nondeterministic and diverge shortly after contact. The reason is fundamental rather than a bug in the barrier: WGSL provides no device-scope memory fence, its atomics are relaxed, and `storageBarrier()` is workgroup-scoped, so non-atomic data written by one workgroup is not guaranteed visible to another inside a dispatch, and on this driver it demonstrably is not. Passing every array through atomics would make it correct and slow. The mode is kept as a documented negative result.
51+
52+**The frequency governor matters more than any of this.** Throughout the multi-dispatch runs the GPU sits at its 300 MHz minimum clock; the stream of small dispatches separated by hazard barriers never looks busy enough to the i915 governor to ramp, while the megakernel's long dispatches do (1300 MHz), and pinning the clock requires root. Attempts to keep the clock up from inside the application, a concurrent second-device load and a spinner workgroup inside each dispatch, both made things slower: the GPU time-slices between contexts too coarsely, and the spinner extends every dispatch since the barriers wait for it. The practical conclusion is that the 83 us/step figure is a floor set by the governor, not by the work: the arithmetic itself would take roughly a quarter as long at full clock, and a discrete GPU, or a browser whose compositor keeps clocks up, should land well below this number without any code change.
53+
54+Timings for the full default render (578k steps at 96 kHz, 18-inch crash, 5.75 m/s stroke), all on the same laptop:
55+
56+| path | per step | full render |
57+|---|---|---|
58+| numbl in Node/browser (the webapp today) | ~1900 us | ~18 min |
59+| numbl CLI with native addon | ~620 us | ~6 min |
60+| MATLAB R2026b | ~370 us | ~3.6 min |
61+| WebGPU multi-dispatch (this, at 300 MHz) | 83 us | 48 s |
62+
63+The `src/probe*.mjs` scripts are the microbenchmarks behind these statements. One caution learned the hard way: short-lived timing harnesses on this stack mislead, both from startup transients and from what appears to be early resolution of completion promises under deep queues, and the numbers above come from the instrumented production runner on full-length renders, not from the probes.
64+
65+## Correctness
66+
67+`--compare` runs the unmodified `cymbal2_sim` through numbl on the reduced pinned configuration (44.1 kHz, no oversampling, 1.5 s tail) and checks the GPU trace against it: first contact on the identical step (1375), peak force 311.23 N against 311.14, tip contact speed -8.224 m/s against -8.223, peak indentation 5.644 mm against 5.643, four contact events in both, and relative L2 of the acceleration over the first 256 steps after contact of about 3e-2, which is f32-versus-f64 on a trajectory with positive Lyapunov exponent, not an indexing error (an indexing error produces order-one differences immediately). The output-stage rms lands at 0.1001 against the suite's pinned 0.0973 with its percent-level tolerance. The megakernel and multi-dispatch backends agree with each other to the same order, differing only in floating-point evaluation order.
68+
69+## What this means for the webapp
70+
71+Nothing here is Node-specific: the host code speaks standard WebGPU, numbl already runs in the render worker, and the tensor marshalling in `src/setup.mjs` is the same `executeCode` pattern `renderWorker.ts` uses. Integrating this into the cymbal2 webapp would mean adding the kernel generator and the multi-dispatch runner to the worker, using it when `navigator.gpu` yields a device, and keeping the pure-numbl path as the fallback; the preview/full tier distinction would survive unchanged, though at these speeds a preview is nearly interactive (the pinned-size render takes 5.6 s here) and a full render drops from twenty minutes to under a minute. Two caveats carry over from the findings above: the browser path needs the same watchdog-conscious chunking this runner already does (1024 steps per submit, two in flight), and edited MATLAB code in the app's editor would still render through the interpreter, since the WGSL mirrors `cymbal2_sim.m` rather than compiling it; a numbl-to-WGSL compilation of the step, the way turing-surface compiles its models, is the natural next step but a substantially larger one, since this loop is matmul- and branch-shaped rather than elementwise.
package-lock.jsonadded+1798−0View file
@@ -0,0 +1,1798 @@
1+{
2+ "name": "cymbal2-webgpu",
3+ "version": "0.1.0",
4+ "lockfileVersion": 3,
5+ "requires": true,
6+ "packages": {
7+ "": {
8+ "name": "cymbal2-webgpu",
9+ "version": "0.1.0",
10+ "dependencies": {
11+ "numbl": "^0.4.19",
12+ "webgpu": "^0.4.0"
13+ }
14+ },
15+ "node_modules/@babel/runtime": {
16+ "version": "7.29.7",
17+ "resolved": "https://registry.npmjs.org/@babel/runtime/-/runtime-7.29.7.tgz",
18+ "integrity": "sha512-Nq8OhGWiZIZGV6hLHoyAKLLcJihP/xFeBMGJoUrxTX2psI8dCifzLhZISFb+VWS3wFMRDmCGw5R+dOySCqPLhw==",
19+ "license": "MIT",
20+ "engines": {
21+ "node": ">=6.9.0"
22+ }
23+ },
24+ "node_modules/@types/debug": {
25+ "version": "4.1.13",
26+ "resolved": "https://registry.npmjs.org/@types/debug/-/debug-4.1.13.tgz",
27+ "integrity": "sha512-KSVgmQmzMwPlmtljOomayoR89W4FynCAi3E8PPs7vmDVPe84hT+vGPKkJfThkmXs0x0jAaa9U8uW8bbfyS2fWw==",
28+ "license": "MIT",
29+ "dependencies": {
30+ "@types/ms": "*"
31+ }
32+ },
33+ "node_modules/@types/estree": {
34+ "version": "1.0.9",
35+ "resolved": "https://registry.npmjs.org/@types/estree/-/estree-1.0.9.tgz",
36+ "integrity": "sha512-GhdPgy1el4/ImP05X05Uw4cw2/M93BCUmnEvWZNStlCzEKME4Fkk+YpoA5OiHNQmoS7Cafb8Xa3Pya8m1Qrzeg==",
37+ "license": "MIT"
38+ },
39+ "node_modules/@types/estree-jsx": {
40+ "version": "1.0.5",
41+ "resolved": "https://registry.npmjs.org/@types/estree-jsx/-/estree-jsx-1.0.5.tgz",
42+ "integrity": "sha512-52CcUVNFyfb1A2ALocQw/Dd1BQFNmSdkuC3BkZ6iqhdMfQz7JWOFRuJFloOzjk+6WijU56m9oKXFAXc7o3Towg==",
43+ "license": "MIT",
44+ "dependencies": {
45+ "@types/estree": "*"
46+ }
47+ },
48+ "node_modules/@types/hast": {
49+ "version": "3.0.5",
50+ "resolved": "https://registry.npmjs.org/@types/hast/-/hast-3.0.5.tgz",
51+ "integrity": "sha512-rp/ezSWaD1m44dPKICGhiskI13nVr7qTloFwDa/IYkhhf5nzwP+zIQcIJh3WIFSBOy/H1PzB40jPjMDksN4F+g==",
52+ "license": "MIT",
53+ "dependencies": {
54+ "@types/unist": "*"
55+ }
56+ },
57+ "node_modules/@types/mdast": {
58+ "version": "4.0.4",
59+ "resolved": "https://registry.npmjs.org/@types/mdast/-/mdast-4.0.4.tgz",
60+ "integrity": "sha512-kGaNbPh1k7AFzgpud/gMdvIm5xuECykRR+JnWKQno9TAXVa6WIVCGTPvYGekIDL4uwCZQSYbUxNBSb1aUo79oA==",
61+ "license": "MIT",
62+ "dependencies": {
63+ "@types/unist": "*"
64+ }
65+ },
66+ "node_modules/@types/ms": {
67+ "version": "2.1.0",
68+ "resolved": "https://registry.npmjs.org/@types/ms/-/ms-2.1.0.tgz",
69+ "integrity": "sha512-GsCCIZDE/p3i96vtEqx+7dBUGXrc7zeSK3wwPHIaRThS+9OhWIXRqzs4d6k1SVU8g91DrNRWxWUGhp5KXQb2VA==",
70+ "license": "MIT"
71+ },
72+ "node_modules/@types/prismjs": {
73+ "version": "1.26.6",
74+ "resolved": "https://registry.npmjs.org/@types/prismjs/-/prismjs-1.26.6.tgz",
75+ "integrity": "sha512-vqlvI7qlMvcCBbVe0AKAb4f97//Hy0EBTaiW8AalRnG/xAN5zOiWWyrNqNXeq8+KAuvRewjCVY1+IPxk4RdNYw==",
76+ "license": "MIT"
77+ },
78+ "node_modules/@types/react": {
79+ "version": "19.2.18",
80+ "resolved": "https://registry.npmjs.org/@types/react/-/react-19.2.18.tgz",
81+ "integrity": "sha512-AnzbBERsrLKtk2XSfTbYRLjQPdy116Sty4q+T+Bp3IC4l6jNBvreVPAHmpq9qhXQM7CXZPjLVmGMw9sy+hxQ3w==",
82+ "license": "MIT",
83+ "peer": true,
84+ "dependencies": {
85+ "csstype": "^3.2.2"
86+ }
87+ },
88+ "node_modules/@types/unist": {
89+ "version": "3.0.3",
90+ "resolved": "https://registry.npmjs.org/@types/unist/-/unist-3.0.3.tgz",
91+ "integrity": "sha512-ko/gIFJRv177XgZsZcBwnqJN5x/Gien8qNOn0D5bQU/zAzVf9Zt3BlcUiLqhV9y4ARk0GbT3tnUiPNgnTXzc/Q==",
92+ "license": "MIT"
93+ },
94+ "node_modules/@ungap/structured-clone": {
95+ "version": "1.4.0",
96+ "resolved": "https://registry.npmjs.org/@ungap/structured-clone/-/structured-clone-1.4.0.tgz",
97+ "integrity": "sha512-1mEZtMKPM09vDmQt5y7YvmN2+DFTP7Tg0EWXdic8/C6VRnpb33e4ghisCIE3WZjsE2N8mf+QV1Zqh7ZFYLWInQ==",
98+ "license": "ISC"
99+ },
100+ "node_modules/@webgpu/types": {
101+ "version": "0.1.72",
102+ "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.72.tgz",
103+ "integrity": "sha512-0cF7RFM2edNoiIS1ODJp0/Gzv4/xSXhwoR0YCza+OWpJWtn4wmo9DvK91aLlH9+uUnwIriP7ZiC3WitmyhuzBw==",
104+ "license": "BSD-3-Clause"
105+ },
106+ "node_modules/bail": {
107+ "version": "2.0.2",
108+ "resolved": "https://registry.npmjs.org/bail/-/bail-2.0.2.tgz",
109+ "integrity": "sha512-0xO6mYd7JB2YesxDKplafRpsiOzPt9V02ddPCLbY1xYGPOX24NTyN50qnUxgCPcSoYMhKpAuBTjQoRZCAkUDRw==",
110+ "license": "MIT",
111+ "funding": {
112+ "type": "github",
113+ "url": "https://github.com/sponsors/wooorm"
114+ }
115+ },
116+ "node_modules/ccount": {
117+ "version": "2.0.1",
118+ "resolved": "https://registry.npmjs.org/ccount/-/ccount-2.0.1.tgz",
119+ "integrity": "sha512-eyrF0jiFpY+3drT6383f1qhkbGsLSifNAjA61IUjZjmLCWjItY6LB9ft9YhoDgwfmclB2zhu51Lc7+95b8NRAg==",
120+ "license": "MIT",
121+ "funding": {
122+ "type": "github",
123+ "url": "https://github.com/sponsors/wooorm"
124+ }
125+ },
126+ "node_modules/character-entities": {
127+ "version": "2.0.2",
128+ "resolved": "https://registry.npmjs.org/character-entities/-/character-entities-2.0.2.tgz",
129+ "integrity": "sha512-shx7oQ0Awen/BRIdkjkvz54PnEEI/EjwXDSIZp86/KKdbafHh1Df/RYGBhn4hbe2+uKC9FnT5UCEdyPz3ai9hQ==",
130+ "license": "MIT",
131+ "funding": {
132+ "type": "github",
133+ "url": "https://github.com/sponsors/wooorm"
134+ }
135+ },
136+ "node_modules/character-entities-html4": {
137+ "version": "2.1.0",
138+ "resolved": "https://registry.npmjs.org/character-entities-html4/-/character-entities-html4-2.1.0.tgz",
139+ "integrity": "sha512-1v7fgQRj6hnSwFpq1Eu0ynr/CDEw0rXo2B61qXrLNdHZmPKgb7fqS1a2JwF0rISo9q77jDI8VMEHoApn8qDoZA==",
140+ "license": "MIT",
141+ "funding": {
142+ "type": "github",
143+ "url": "https://github.com/sponsors/wooorm"
144+ }
145+ },
146+ "node_modules/character-entities-legacy": {
147+ "version": "3.0.0",
148+ "resolved": "https://registry.npmjs.org/character-entities-legacy/-/character-entities-legacy-3.0.0.tgz",
149+ "integrity": "sha512-RpPp0asT/6ufRm//AJVwpViZbGM/MkjQFxJccQRHmISF/22NBtsHqAWmL+/pmkPWoIUJdWyeVleTl1wydHATVQ==",
150+ "license": "MIT",
151+ "funding": {
152+ "type": "github",
153+ "url": "https://github.com/sponsors/wooorm"
154+ }
155+ },
156+ "node_modules/character-reference-invalid": {
157+ "version": "2.0.1",
158+ "resolved": "https://registry.npmjs.org/character-reference-invalid/-/character-reference-invalid-2.0.1.tgz",
159+ "integrity": "sha512-iBZ4F4wRbyORVsu0jPV7gXkOsGYjGHPmAyv+HiHG8gi5PtC9KI2j1+v8/tlibRvjoWX027ypmG/n0HtO5t7unw==",
160+ "license": "MIT",
161+ "funding": {
162+ "type": "github",
163+ "url": "https://github.com/sponsors/wooorm"
164+ }
165+ },
166+ "node_modules/comma-separated-tokens": {
167+ "version": "2.0.3",
168+ "resolved": "https://registry.npmjs.org/comma-separated-tokens/-/comma-separated-tokens-2.0.3.tgz",
169+ "integrity": "sha512-Fu4hJdvzeylCfQPp9SGWidpzrMs7tTrlu6Vb8XGaRGck8QSNZJJp538Wrb60Lax4fPwR64ViY468OIUTbRlGZg==",
170+ "license": "MIT",
171+ "funding": {
172+ "type": "github",
173+ "url": "https://github.com/sponsors/wooorm"
174+ }
175+ },
176+ "node_modules/csstype": {
177+ "version": "3.2.3",
178+ "resolved": "https://registry.npmjs.org/csstype/-/csstype-3.2.3.tgz",
179+ "integrity": "sha512-z1HGKcYy2xA8AGQfwrn0PAy+PB7X/GSj3UVJW9qKyn43xWa+gl5nXmU4qqLMRzWVLFC8KusUX8T/0kCiOYpAIQ==",
180+ "license": "MIT",
181+ "peer": true
182+ },
183+ "node_modules/debug": {
184+ "version": "4.4.3",
185+ "resolved": "https://registry.npmjs.org/debug/-/debug-4.4.3.tgz",
186+ "integrity": "sha512-RGwwWnwQvkVfavKVt22FGLw+xYSdzARwm0ru6DhTVA3umU5hZc28V3kO4stgYryrTlLpuvgI9GiijltAjNbcqA==",
187+ "license": "MIT",
188+ "dependencies": {
189+ "ms": "^2.1.3"
190+ },
191+ "engines": {
192+ "node": ">=6.0"
193+ },
194+ "peerDependenciesMeta": {
195+ "supports-color": {
196+ "optional": true
197+ }
198+ }
199+ },
200+ "node_modules/decode-named-character-reference": {
201+ "version": "1.3.0",
202+ "resolved": "https://registry.npmjs.org/decode-named-character-reference/-/decode-named-character-reference-1.3.0.tgz",
203+ "integrity": "sha512-GtpQYB283KrPp6nRw50q3U9/VfOutZOe103qlN7BPP6Ad27xYnOIWv4lPzo8HCAL+mMZofJ9KEy30fq6MfaK6Q==",
204+ "license": "MIT",
205+ "dependencies": {
206+ "character-entities": "^2.0.0"
207+ },
208+ "funding": {
209+ "type": "github",
210+ "url": "https://github.com/sponsors/wooorm"
211+ }
212+ },
213+ "node_modules/dequal": {
214+ "version": "2.0.3",
215+ "resolved": "https://registry.npmjs.org/dequal/-/dequal-2.0.3.tgz",
216+ "integrity": "sha512-0je+qPKHEMohvfRTCEo3CrPG6cAzAYgmzKyxRiYSSDkS6eGJdyVJm7WaYA5ECaAD9wLB2T4EEeymA5aFVcYXCA==",
217+ "license": "MIT",
218+ "engines": {
219+ "node": ">=6"
220+ }
221+ },
222+ "node_modules/devlop": {
223+ "version": "1.1.0",
224+ "resolved": "https://registry.npmjs.org/devlop/-/devlop-1.1.0.tgz",
225+ "integrity": "sha512-RWmIqhcFf1lRYBvNmr7qTNuyCt/7/ns2jbpp1+PalgE/rDQcBT0fioSMUpJ93irlUhC5hrg4cYqe6U+0ImW0rA==",
226+ "license": "MIT",
227+ "dependencies": {
228+ "dequal": "^2.0.0"
229+ },
230+ "funding": {
231+ "type": "github",
232+ "url": "https://github.com/sponsors/wooorm"
233+ }
234+ },
235+ "node_modules/escape-string-regexp": {
236+ "version": "5.0.0",
237+ "resolved": "https://registry.npmjs.org/escape-string-regexp/-/escape-string-regexp-5.0.0.tgz",
238+ "integrity": "sha512-/veY75JbMK4j1yjvuUxuVsiS/hr/4iHs9FTT6cgTexxdE0Ly/glccBAkloH/DofkjRbZU3bnoj38mOmhkZ0lHw==",
239+ "license": "MIT",
240+ "engines": {
241+ "node": ">=12"
242+ },
243+ "funding": {
244+ "url": "https://github.com/sponsors/sindresorhus"
245+ }
246+ },
247+ "node_modules/estree-util-is-identifier-name": {
248+ "version": "3.0.0",
249+ "resolved": "https://registry.npmjs.org/estree-util-is-identifier-name/-/estree-util-is-identifier-name-3.0.0.tgz",
250+ "integrity": "sha512-hFtqIDZTIUZ9BXLb8y4pYGyk6+wekIivNVTcmvk8NoOh+VeRn5y6cEHzbURrWbfp1fIqdVipilzj+lfaadNZmg==",
251+ "license": "MIT",
252+ "funding": {
253+ "type": "opencollective",
254+ "url": "https://opencollective.com/unified"
255+ }
256+ },
257+ "node_modules/extend": {
258+ "version": "3.0.2",
259+ "resolved": "https://registry.npmjs.org/extend/-/extend-3.0.2.tgz",
260+ "integrity": "sha512-fjquC59cD7CyW6urNXK0FBufkZcoiGG80wTuPujX590cB5Ttln20E2UB4S/WARVqhXffZl2LNgS+gQdPIIim/g==",
261+ "license": "MIT"
262+ },
263+ "node_modules/fault": {
264+ "version": "1.0.4",
265+ "resolved": "https://registry.npmjs.org/fault/-/fault-1.0.4.tgz",
266+ "integrity": "sha512-CJ0HCB5tL5fYTEA7ToAq5+kTwd++Borf1/bifxd9iT70QcXr4MRrO3Llf8Ifs70q+SJcGHFtnIE/Nw6giCtECA==",
267+ "license": "MIT",
268+ "dependencies": {
269+ "format": "^0.2.0"
270+ },
271+ "funding": {
272+ "type": "github",
273+ "url": "https://github.com/sponsors/wooorm"
274+ }
275+ },
276+ "node_modules/fflate": {
277+ "version": "0.8.3",
278+ "resolved": "https://registry.npmjs.org/fflate/-/fflate-0.8.3.tgz",
279+ "integrity": "sha512-tbZNuJrLwGUp3zshBtdy4W+ORxZuIh8a5ilyIEQDC5rY1f3U20JMry0Ll3WBzU58EZKsEuJFXhb5gwv8CsPvgA==",
280+ "license": "MIT"
281+ },
282+ "node_modules/format": {
283+ "version": "0.2.2",
284+ "resolved": "https://registry.npmjs.org/format/-/format-0.2.2.tgz",
285+ "integrity": "sha512-wzsgA6WOq+09wrU1tsJ09udeR/YZRaeArL9e1wPbFg3GG2yDnC2ldKpxs4xunpFF9DgqCqOIra3bc1HWrJ37Ww==",
286+ "engines": {
287+ "node": ">=0.4.x"
288+ }
289+ },
290+ "node_modules/h5wasm": {
291+ "version": "0.10.3",
292+ "resolved": "https://registry.npmjs.org/h5wasm/-/h5wasm-0.10.3.tgz",
293+ "integrity": "sha512-W4Jy5ExtX/VNbyD8GdOBckDuj6AL16TemppVNxZsV3rJZEWCv2sxlCzOttZLer3zkMttbDYsWHl0qt1z3Bln+Q==",
294+ "license": "SEE LICENSE IN LICENSE.txt"
295+ },
296+ "node_modules/hast-util-parse-selector": {
297+ "version": "4.0.0",
298+ "resolved": "https://registry.npmjs.org/hast-util-parse-selector/-/hast-util-parse-selector-4.0.0.tgz",
299+ "integrity": "sha512-wkQCkSYoOGCRKERFWcxMVMOcYE2K1AaNLU8DXS9arxnLOUEWbOXKXiJUNzEpqZ3JOKpnha3jkFrumEjVliDe7A==",
300+ "license": "MIT",
301+ "dependencies": {
302+ "@types/hast": "^3.0.0"
303+ },
304+ "funding": {
305+ "type": "opencollective",
306+ "url": "https://opencollective.com/unified"
307+ }
308+ },
309+ "node_modules/hast-util-to-jsx-runtime": {
310+ "version": "2.3.6",
311+ "resolved": "https://registry.npmjs.org/hast-util-to-jsx-runtime/-/hast-util-to-jsx-runtime-2.3.6.tgz",
312+ "integrity": "sha512-zl6s8LwNyo1P9uw+XJGvZtdFF1GdAkOg8ujOw+4Pyb76874fLps4ueHXDhXWdk6YHQ6OgUtinliG7RsYvCbbBg==",
313+ "license": "MIT",
314+ "dependencies": {
315+ "@types/estree": "^1.0.0",
316+ "@types/hast": "^3.0.0",
317+ "@types/unist": "^3.0.0",
318+ "comma-separated-tokens": "^2.0.0",
319+ "devlop": "^1.0.0",
320+ "estree-util-is-identifier-name": "^3.0.0",
321+ "hast-util-whitespace": "^3.0.0",
322+ "mdast-util-mdx-expression": "^2.0.0",
323+ "mdast-util-mdx-jsx": "^3.0.0",
324+ "mdast-util-mdxjs-esm": "^2.0.0",
325+ "property-information": "^7.0.0",
326+ "space-separated-tokens": "^2.0.0",
327+ "style-to-js": "^1.0.0",
328+ "unist-util-position": "^5.0.0",
329+ "vfile-message": "^4.0.0"
330+ },
331+ "funding": {
332+ "type": "opencollective",
333+ "url": "https://opencollective.com/unified"
334+ }
335+ },
336+ "node_modules/hast-util-whitespace": {
337+ "version": "3.0.0",
338+ "resolved": "https://registry.npmjs.org/hast-util-whitespace/-/hast-util-whitespace-3.0.0.tgz",
339+ "integrity": "sha512-88JUN06ipLwsnv+dVn+OIYOvAuvBMy/Qoi6O7mQHxdPXpjy+Cd6xRkWwux7DKO+4sYILtLBRIKgsdpS2gQc7qw==",
340+ "license": "MIT",
341+ "dependencies": {
342+ "@types/hast": "^3.0.0"
343+ },
344+ "funding": {
345+ "type": "opencollective",
346+ "url": "https://opencollective.com/unified"
347+ }
348+ },
349+ "node_modules/hastscript": {
350+ "version": "9.0.1",
351+ "resolved": "https://registry.npmjs.org/hastscript/-/hastscript-9.0.1.tgz",
352+ "integrity": "sha512-g7df9rMFX/SPi34tyGCyUBREQoKkapwdY/T04Qn9TDWfHhAYt4/I0gMVirzK5wEzeUqIjEB+LXC/ypb7Aqno5w==",
353+ "license": "MIT",
354+ "dependencies": {
355+ "@types/hast": "^3.0.0",
356+ "comma-separated-tokens": "^2.0.0",
357+ "hast-util-parse-selector": "^4.0.0",
358+ "property-information": "^7.0.0",
359+ "space-separated-tokens": "^2.0.0"
360+ },
361+ "funding": {
362+ "type": "opencollective",
363+ "url": "https://opencollective.com/unified"
364+ }
365+ },
366+ "node_modules/highlight.js": {
367+ "version": "10.7.3",
368+ "resolved": "https://registry.npmjs.org/highlight.js/-/highlight.js-10.7.3.tgz",
369+ "integrity": "sha512-tzcUFauisWKNHaRkN4Wjl/ZA07gENAjFl3J/c480dprkGTg5EQstgaNFqBfUqCq54kZRIEcreTsAgF/m2quD7A==",
370+ "license": "BSD-3-Clause",
371+ "engines": {
372+ "node": "*"
373+ }
374+ },
375+ "node_modules/highlightjs-vue": {
376+ "version": "1.0.0",
377+ "resolved": "https://registry.npmjs.org/highlightjs-vue/-/highlightjs-vue-1.0.0.tgz",
378+ "integrity": "sha512-PDEfEF102G23vHmPhLyPboFCD+BkMGu+GuJe2d9/eH4FsCwvgBpnc9n0pGE+ffKdph38s6foEZiEjdgHdzp+IA==",
379+ "license": "CC0-1.0"
380+ },
381+ "node_modules/html-url-attributes": {
382+ "version": "3.0.1",
383+ "resolved": "https://registry.npmjs.org/html-url-attributes/-/html-url-attributes-3.0.1.tgz",
384+ "integrity": "sha512-ol6UPyBWqsrO6EJySPz2O7ZSr856WDrEzM5zMqp+FJJLGMW35cLYmmZnl0vztAZxRUoNZJFTCohfjuIJ8I4QBQ==",
385+ "license": "MIT",
386+ "funding": {
387+ "type": "opencollective",
388+ "url": "https://opencollective.com/unified"
389+ }
390+ },
391+ "node_modules/inline-style-parser": {
392+ "version": "0.2.7",
393+ "resolved": "https://registry.npmjs.org/inline-style-parser/-/inline-style-parser-0.2.7.tgz",
394+ "integrity": "sha512-Nb2ctOyNR8DqQoR0OwRG95uNWIC0C1lCgf5Naz5H6Ji72KZ8OcFZLz2P5sNgwlyoJ8Yif11oMuYs5pBQa86csA==",
395+ "license": "MIT"
396+ },
397+ "node_modules/is-alphabetical": {
398+ "version": "2.0.1",
399+ "resolved": "https://registry.npmjs.org/is-alphabetical/-/is-alphabetical-2.0.1.tgz",
400+ "integrity": "sha512-FWyyY60MeTNyeSRpkM2Iry0G9hpr7/9kD40mD/cGQEuilcZYS4okz8SN2Q6rLCJ8gbCt6fN+rC+6tMGS99LaxQ==",
401+ "license": "MIT",
402+ "funding": {
403+ "type": "github",
404+ "url": "https://github.com/sponsors/wooorm"
405+ }
406+ },
407+ "node_modules/is-alphanumerical": {
408+ "version": "2.0.1",
409+ "resolved": "https://registry.npmjs.org/is-alphanumerical/-/is-alphanumerical-2.0.1.tgz",
410+ "integrity": "sha512-hmbYhX/9MUMF5uh7tOXyK/n0ZvWpad5caBA17GsC6vyuCqaWliRG5K1qS9inmUhEMaOBIW7/whAnSwveW/LtZw==",
411+ "license": "MIT",
412+ "dependencies": {
413+ "is-alphabetical": "^2.0.0",
414+ "is-decimal": "^2.0.0"
415+ },
416+ "funding": {
417+ "type": "github",
418+ "url": "https://github.com/sponsors/wooorm"
419+ }
420+ },
421+ "node_modules/is-decimal": {
422+ "version": "2.0.1",
423+ "resolved": "https://registry.npmjs.org/is-decimal/-/is-decimal-2.0.1.tgz",
424+ "integrity": "sha512-AAB9hiomQs5DXWcRB1rqsxGUstbRroFOPPVAomNk/3XHR5JyEZChOyTWe2oayKnsSsr/kcGqF+z6yuH6HHpN0A==",
425+ "license": "MIT",
426+ "funding": {
427+ "type": "github",
428+ "url": "https://github.com/sponsors/wooorm"
429+ }
430+ },
431+ "node_modules/is-hexadecimal": {
432+ "version": "2.0.1",
433+ "resolved": "https://registry.npmjs.org/is-hexadecimal/-/is-hexadecimal-2.0.1.tgz",
434+ "integrity": "sha512-DgZQp241c8oO6cA1SbTEWiXeoxV42vlcJxgH+B3hi1AiqqKruZR3ZGF8In3fj4+/y/7rHvlOZLZtgJ/4ttYGZg==",
435+ "license": "MIT",
436+ "funding": {
437+ "type": "github",
438+ "url": "https://github.com/sponsors/wooorm"
439+ }
440+ },
441+ "node_modules/is-plain-obj": {
442+ "version": "4.1.0",
443+ "resolved": "https://registry.npmjs.org/is-plain-obj/-/is-plain-obj-4.1.0.tgz",
444+ "integrity": "sha512-+Pgi+vMuUNkJyExiMBt5IlFoMyKnr5zhJ4Uspz58WOhBF5QoIZkFyNHIbBAtHwzVAgk5RtndVNsDRN61/mmDqg==",
445+ "license": "MIT",
446+ "engines": {
447+ "node": ">=12"
448+ },
449+ "funding": {
450+ "url": "https://github.com/sponsors/sindresorhus"
451+ }
452+ },
453+ "node_modules/koffi": {
454+ "version": "2.16.3",
455+ "resolved": "https://registry.npmjs.org/koffi/-/koffi-2.16.3.tgz",
456+ "integrity": "sha512-E9y1AsgYGlaxMhcZzHr8y96QF2U5XzA12GGVAfbWqIubTwPNMXQarfBzePNXHe0xtIEtNd6ifAv3GAKYGUeBAQ==",
457+ "hasInstallScript": true,
458+ "license": "MIT",
459+ "optional": true,
460+ "funding": {
461+ "url": "https://liberapay.com/Koromix"
462+ }
463+ },
464+ "node_modules/longest-streak": {
465+ "version": "3.1.0",
466+ "resolved": "https://registry.npmjs.org/longest-streak/-/longest-streak-3.1.0.tgz",
467+ "integrity": "sha512-9Ri+o0JYgehTaVBBDoMqIl8GXtbWg711O3srftcHhZ0dqnETqLaoIK0x17fUw9rFSlK/0NlsKe0Ahhyl5pXE2g==",
468+ "license": "MIT",
469+ "funding": {
470+ "type": "github",
471+ "url": "https://github.com/sponsors/wooorm"
472+ }
473+ },
474+ "node_modules/lowlight": {
475+ "version": "1.20.0",
476+ "resolved": "https://registry.npmjs.org/lowlight/-/lowlight-1.20.0.tgz",
477+ "integrity": "sha512-8Ktj+prEb1RoCPkEOrPMYUN/nCggB7qAWe3a7OpMjWQkh3l2RD5wKRQ+o8Q8YuI9RG/xs95waaI/E6ym/7NsTw==",
478+ "license": "MIT",
479+ "dependencies": {
480+ "fault": "^1.0.0",
481+ "highlight.js": "~10.7.0"
482+ },
483+ "funding": {
484+ "type": "github",
485+ "url": "https://github.com/sponsors/wooorm"
486+ }
487+ },
488+ "node_modules/markdown-table": {
489+ "version": "3.0.4",
490+ "resolved": "https://registry.npmjs.org/markdown-table/-/markdown-table-3.0.4.tgz",
491+ "integrity": "sha512-wiYz4+JrLyb/DqW2hkFJxP7Vd7JuTDm77fvbM8VfEQdmSMqcImWeeRbHwZjBjIFki/VaMK2BhFi7oUUZeM5bqw==",
492+ "license": "MIT",
493+ "funding": {
494+ "type": "github",
495+ "url": "https://github.com/sponsors/wooorm"
496+ }
497+ },
498+ "node_modules/mdast-util-find-and-replace": {
499+ "version": "3.0.2",
500+ "resolved": "https://registry.npmjs.org/mdast-util-find-and-replace/-/mdast-util-find-and-replace-3.0.2.tgz",
501+ "integrity": "sha512-Tmd1Vg/m3Xz43afeNxDIhWRtFZgM2VLyaf4vSTYwudTyeuTneoL3qtWMA5jeLyz/O1vDJmmV4QuScFCA2tBPwg==",
502+ "license": "MIT",
503+ "dependencies": {
504+ "@types/mdast": "^4.0.0",
505+ "escape-string-regexp": "^5.0.0",
506+ "unist-util-is": "^6.0.0",
507+ "unist-util-visit-parents": "^6.0.0"
508+ },
509+ "funding": {
510+ "type": "opencollective",
511+ "url": "https://opencollective.com/unified"
512+ }
513+ },
514+ "node_modules/mdast-util-from-markdown": {
515+ "version": "2.0.3",
516+ "resolved": "https://registry.npmjs.org/mdast-util-from-markdown/-/mdast-util-from-markdown-2.0.3.tgz",
517+ "integrity": "sha512-W4mAWTvSlKvf8L6J+VN9yLSqQ9AOAAvHuoDAmPkz4dHf553m5gVj2ejadHJhoJmcmxEnOv6Pa8XJhpxE93kb8Q==",
518+ "license": "MIT",
519+ "dependencies": {
520+ "@types/mdast": "^4.0.0",
521+ "@types/unist": "^3.0.0",
522+ "decode-named-character-reference": "^1.0.0",
523+ "devlop": "^1.0.0",
524+ "mdast-util-to-string": "^4.0.0",
525+ "micromark": "^4.0.0",
526+ "micromark-util-decode-numeric-character-reference": "^2.0.0",
527+ "micromark-util-decode-string": "^2.0.0",
528+ "micromark-util-normalize-identifier": "^2.0.0",
529+ "micromark-util-symbol": "^2.0.0",
530+ "micromark-util-types": "^2.0.0",
531+ "unist-util-stringify-position": "^4.0.0"
532+ },
533+ "funding": {
534+ "type": "opencollective",
535+ "url": "https://opencollective.com/unified"
536+ }
537+ },
538+ "node_modules/mdast-util-gfm": {
539+ "version": "3.1.0",
540+ "resolved": "https://registry.npmjs.org/mdast-util-gfm/-/mdast-util-gfm-3.1.0.tgz",
541+ "integrity": "sha512-0ulfdQOM3ysHhCJ1p06l0b0VKlhU0wuQs3thxZQagjcjPrlFRqY215uZGHHJan9GEAXd9MbfPjFJz+qMkVR6zQ==",
542+ "license": "MIT",
543+ "dependencies": {
544+ "mdast-util-from-markdown": "^2.0.0",
545+ "mdast-util-gfm-autolink-literal": "^2.0.0",
546+ "mdast-util-gfm-footnote": "^2.0.0",
547+ "mdast-util-gfm-strikethrough": "^2.0.0",
548+ "mdast-util-gfm-table": "^2.0.0",
549+ "mdast-util-gfm-task-list-item": "^2.0.0",
550+ "mdast-util-to-markdown": "^2.0.0"
551+ },
552+ "funding": {
553+ "type": "opencollective",
554+ "url": "https://opencollective.com/unified"
555+ }
556+ },
557+ "node_modules/mdast-util-gfm-autolink-literal": {
558+ "version": "2.0.1",
559+ "resolved": "https://registry.npmjs.org/mdast-util-gfm-autolink-literal/-/mdast-util-gfm-autolink-literal-2.0.1.tgz",
560+ "integrity": "sha512-5HVP2MKaP6L+G6YaxPNjuL0BPrq9orG3TsrZ9YXbA3vDw/ACI4MEsnoDpn6ZNm7GnZgtAcONJyPhOP8tNJQavQ==",
561+ "license": "MIT",
562+ "dependencies": {
563+ "@types/mdast": "^4.0.0",
564+ "ccount": "^2.0.0",
565+ "devlop": "^1.0.0",
566+ "mdast-util-find-and-replace": "^3.0.0",
567+ "micromark-util-character": "^2.0.0"
568+ },
569+ "funding": {
570+ "type": "opencollective",
571+ "url": "https://opencollective.com/unified"
572+ }
573+ },
574+ "node_modules/mdast-util-gfm-footnote": {
575+ "version": "2.1.0",
576+ "resolved": "https://registry.npmjs.org/mdast-util-gfm-footnote/-/mdast-util-gfm-footnote-2.1.0.tgz",
577+ "integrity": "sha512-sqpDWlsHn7Ac9GNZQMeUzPQSMzR6Wv0WKRNvQRg0KqHh02fpTz69Qc1QSseNX29bhz1ROIyNyxExfawVKTm1GQ==",
578+ "license": "MIT",
579+ "dependencies": {
580+ "@types/mdast": "^4.0.0",
581+ "devlop": "^1.1.0",
582+ "mdast-util-from-markdown": "^2.0.0",
583+ "mdast-util-to-markdown": "^2.0.0",
584+ "micromark-util-normalize-identifier": "^2.0.0"
585+ },
586+ "funding": {
587+ "type": "opencollective",
588+ "url": "https://opencollective.com/unified"
589+ }
590+ },
591+ "node_modules/mdast-util-gfm-strikethrough": {
592+ "version": "2.0.0",
593+ "resolved": "https://registry.npmjs.org/mdast-util-gfm-strikethrough/-/mdast-util-gfm-strikethrough-2.0.0.tgz",
594+ "integrity": "sha512-mKKb915TF+OC5ptj5bJ7WFRPdYtuHv0yTRxK2tJvi+BDqbkiG7h7u/9SI89nRAYcmap2xHQL9D+QG/6wSrTtXg==",
595+ "license": "MIT",
596+ "dependencies": {
597+ "@types/mdast": "^4.0.0",
598+ "mdast-util-from-markdown": "^2.0.0",
599+ "mdast-util-to-markdown": "^2.0.0"
600+ },
601+ "funding": {
602+ "type": "opencollective",
603+ "url": "https://opencollective.com/unified"
604+ }
605+ },
606+ "node_modules/mdast-util-gfm-table": {
607+ "version": "2.0.0",
608+ "resolved": "https://registry.npmjs.org/mdast-util-gfm-table/-/mdast-util-gfm-table-2.0.0.tgz",
609+ "integrity": "sha512-78UEvebzz/rJIxLvE7ZtDd/vIQ0RHv+3Mh5DR96p7cS7HsBhYIICDBCu8csTNWNO6tBWfqXPWekRuj2FNOGOZg==",
610+ "license": "MIT",
611+ "dependencies": {
612+ "@types/mdast": "^4.0.0",
613+ "devlop": "^1.0.0",
614+ "markdown-table": "^3.0.0",
615+ "mdast-util-from-markdown": "^2.0.0",
616+ "mdast-util-to-markdown": "^2.0.0"
617+ },
618+ "funding": {
619+ "type": "opencollective",
620+ "url": "https://opencollective.com/unified"
621+ }
622+ },
623+ "node_modules/mdast-util-gfm-task-list-item": {
624+ "version": "2.0.0",
625+ "resolved": "https://registry.npmjs.org/mdast-util-gfm-task-list-item/-/mdast-util-gfm-task-list-item-2.0.0.tgz",
626+ "integrity": "sha512-IrtvNvjxC1o06taBAVJznEnkiHxLFTzgonUdy8hzFVeDun0uTjxxrRGVaNFqkU1wJR3RBPEfsxmU6jDWPofrTQ==",
627+ "license": "MIT",
628+ "dependencies": {
629+ "@types/mdast": "^4.0.0",
630+ "devlop": "^1.0.0",
631+ "mdast-util-from-markdown": "^2.0.0",
632+ "mdast-util-to-markdown": "^2.0.0"
633+ },
634+ "funding": {
635+ "type": "opencollective",
636+ "url": "https://opencollective.com/unified"
637+ }
638+ },
639+ "node_modules/mdast-util-mdx-expression": {
640+ "version": "2.0.1",
641+ "resolved": "https://registry.npmjs.org/mdast-util-mdx-expression/-/mdast-util-mdx-expression-2.0.1.tgz",
642+ "integrity": "sha512-J6f+9hUp+ldTZqKRSg7Vw5V6MqjATc+3E4gf3CFNcuZNWD8XdyI6zQ8GqH7f8169MM6P7hMBRDVGnn7oHB9kXQ==",
643+ "license": "MIT",
644+ "dependencies": {
645+ "@types/estree-jsx": "^1.0.0",
646+ "@types/hast": "^3.0.0",
647+ "@types/mdast": "^4.0.0",
648+ "devlop": "^1.0.0",
649+ "mdast-util-from-markdown": "^2.0.0",
650+ "mdast-util-to-markdown": "^2.0.0"
651+ },
652+ "funding": {
653+ "type": "opencollective",
654+ "url": "https://opencollective.com/unified"
655+ }
656+ },
657+ "node_modules/mdast-util-mdx-jsx": {
658+ "version": "3.2.0",
659+ "resolved": "https://registry.npmjs.org/mdast-util-mdx-jsx/-/mdast-util-mdx-jsx-3.2.0.tgz",
660+ "integrity": "sha512-lj/z8v0r6ZtsN/cGNNtemmmfoLAFZnjMbNyLzBafjzikOM+glrjNHPlf6lQDOTccj9n5b0PPihEBbhneMyGs1Q==",
661+ "license": "MIT",
662+ "dependencies": {
663+ "@types/estree-jsx": "^1.0.0",
664+ "@types/hast": "^3.0.0",
665+ "@types/mdast": "^4.0.0",
666+ "@types/unist": "^3.0.0",
667+ "ccount": "^2.0.0",
668+ "devlop": "^1.1.0",
669+ "mdast-util-from-markdown": "^2.0.0",
670+ "mdast-util-to-markdown": "^2.0.0",
671+ "parse-entities": "^4.0.0",
672+ "stringify-entities": "^4.0.0",
673+ "unist-util-stringify-position": "^4.0.0",
674+ "vfile-message": "^4.0.0"
675+ },
676+ "funding": {
677+ "type": "opencollective",
678+ "url": "https://opencollective.com/unified"
679+ }
680+ },
681+ "node_modules/mdast-util-mdxjs-esm": {
682+ "version": "2.0.1",
683+ "resolved": "https://registry.npmjs.org/mdast-util-mdxjs-esm/-/mdast-util-mdxjs-esm-2.0.1.tgz",
684+ "integrity": "sha512-EcmOpxsZ96CvlP03NghtH1EsLtr0n9Tm4lPUJUBccV9RwUOneqSycg19n5HGzCf+10LozMRSObtVr3ee1WoHtg==",
685+ "license": "MIT",
686+ "dependencies": {
687+ "@types/estree-jsx": "^1.0.0",
688+ "@types/hast": "^3.0.0",
689+ "@types/mdast": "^4.0.0",
690+ "devlop": "^1.0.0",
691+ "mdast-util-from-markdown": "^2.0.0",
692+ "mdast-util-to-markdown": "^2.0.0"
693+ },
694+ "funding": {
695+ "type": "opencollective",
696+ "url": "https://opencollective.com/unified"
697+ }
698+ },
699+ "node_modules/mdast-util-phrasing": {
700+ "version": "4.1.0",
701+ "resolved": "https://registry.npmjs.org/mdast-util-phrasing/-/mdast-util-phrasing-4.1.0.tgz",
702+ "integrity": "sha512-TqICwyvJJpBwvGAMZjj4J2n0X8QWp21b9l0o7eXyVJ25YNWYbJDVIyD1bZXE6WtV6RmKJVYmQAKWa0zWOABz2w==",
703+ "license": "MIT",
704+ "dependencies": {
705+ "@types/mdast": "^4.0.0",
706+ "unist-util-is": "^6.0.0"
707+ },
708+ "funding": {
709+ "type": "opencollective",
710+ "url": "https://opencollective.com/unified"
711+ }
712+ },
713+ "node_modules/mdast-util-to-hast": {
714+ "version": "13.2.1",
715+ "resolved": "https://registry.npmjs.org/mdast-util-to-hast/-/mdast-util-to-hast-13.2.1.tgz",
716+ "integrity": "sha512-cctsq2wp5vTsLIcaymblUriiTcZd0CwWtCbLvrOzYCDZoWyMNV8sZ7krj09FSnsiJi3WVsHLM4k6Dq/yaPyCXA==",
717+ "license": "MIT",
718+ "dependencies": {
719+ "@types/hast": "^3.0.0",
720+ "@types/mdast": "^4.0.0",
721+ "@ungap/structured-clone": "^1.0.0",
722+ "devlop": "^1.0.0",
723+ "micromark-util-sanitize-uri": "^2.0.0",
724+ "trim-lines": "^3.0.0",
725+ "unist-util-position": "^5.0.0",
726+ "unist-util-visit": "^5.0.0",
727+ "vfile": "^6.0.0"
728+ },
729+ "funding": {
730+ "type": "opencollective",
731+ "url": "https://opencollective.com/unified"
732+ }
733+ },
734+ "node_modules/mdast-util-to-markdown": {
735+ "version": "2.1.2",
736+ "resolved": "https://registry.npmjs.org/mdast-util-to-markdown/-/mdast-util-to-markdown-2.1.2.tgz",
737+ "integrity": "sha512-xj68wMTvGXVOKonmog6LwyJKrYXZPvlwabaryTjLh9LuvovB/KAH+kvi8Gjj+7rJjsFi23nkUxRQv1KqSroMqA==",
738+ "license": "MIT",
739+ "dependencies": {
740+ "@types/mdast": "^4.0.0",
741+ "@types/unist": "^3.0.0",
742+ "longest-streak": "^3.0.0",
743+ "mdast-util-phrasing": "^4.0.0",
744+ "mdast-util-to-string": "^4.0.0",
745+ "micromark-util-classify-character": "^2.0.0",
746+ "micromark-util-decode-string": "^2.0.0",
747+ "unist-util-visit": "^5.0.0",
748+ "zwitch": "^2.0.0"
749+ },
750+ "funding": {
751+ "type": "opencollective",
752+ "url": "https://opencollective.com/unified"
753+ }
754+ },
755+ "node_modules/mdast-util-to-string": {
756+ "version": "4.0.0",
757+ "resolved": "https://registry.npmjs.org/mdast-util-to-string/-/mdast-util-to-string-4.0.0.tgz",
758+ "integrity": "sha512-0H44vDimn51F0YwvxSJSm0eCDOJTRlmN0R1yBh4HLj9wiV1Dn0QoXGbvFAWj2hSItVTlCmBF1hqKlIyUBVFLPg==",
759+ "license": "MIT",
760+ "dependencies": {
761+ "@types/mdast": "^4.0.0"
762+ },
763+ "funding": {
764+ "type": "opencollective",
765+ "url": "https://opencollective.com/unified"
766+ }
767+ },
768+ "node_modules/micromark": {
769+ "version": "4.0.2",
770+ "resolved": "https://registry.npmjs.org/micromark/-/micromark-4.0.2.tgz",
771+ "integrity": "sha512-zpe98Q6kvavpCr1NPVSCMebCKfD7CA2NqZ+rykeNhONIJBpc1tFKt9hucLGwha3jNTNI8lHpctWJWoimVF4PfA==",
772+ "funding": [
773+ {
774+ "type": "GitHub Sponsors",
775+ "url": "https://github.com/sponsors/unifiedjs"
776+ },
777+ {
778+ "type": "OpenCollective",
779+ "url": "https://opencollective.com/unified"
780+ }
781+ ],
782+ "license": "MIT",
783+ "dependencies": {
784+ "@types/debug": "^4.0.0",
785+ "debug": "^4.0.0",
786+ "decode-named-character-reference": "^1.0.0",
787+ "devlop": "^1.0.0",
788+ "micromark-core-commonmark": "^2.0.0",
789+ "micromark-factory-space": "^2.0.0",
790+ "micromark-util-character": "^2.0.0",
791+ "micromark-util-chunked": "^2.0.0",
792+ "micromark-util-combine-extensions": "^2.0.0",
793+ "micromark-util-decode-numeric-character-reference": "^2.0.0",
794+ "micromark-util-encode": "^2.0.0",
795+ "micromark-util-normalize-identifier": "^2.0.0",
796+ "micromark-util-resolve-all": "^2.0.0",
797+ "micromark-util-sanitize-uri": "^2.0.0",
798+ "micromark-util-subtokenize": "^2.0.0",
799+ "micromark-util-symbol": "^2.0.0",
800+ "micromark-util-types": "^2.0.0"
801+ }
802+ },
803+ "node_modules/micromark-core-commonmark": {
804+ "version": "2.0.3",
805+ "resolved": "https://registry.npmjs.org/micromark-core-commonmark/-/micromark-core-commonmark-2.0.3.tgz",
806+ "integrity": "sha512-RDBrHEMSxVFLg6xvnXmb1Ayr2WzLAWjeSATAoxwKYJV94TeNavgoIdA0a9ytzDSVzBy2YKFK+emCPOEibLeCrg==",
807+ "funding": [
808+ {
809+ "type": "GitHub Sponsors",
810+ "url": "https://github.com/sponsors/unifiedjs"
811+ },
812+ {
813+ "type": "OpenCollective",
814+ "url": "https://opencollective.com/unified"
815+ }
816+ ],
817+ "license": "MIT",
818+ "dependencies": {
819+ "decode-named-character-reference": "^1.0.0",
820+ "devlop": "^1.0.0",
821+ "micromark-factory-destination": "^2.0.0",
822+ "micromark-factory-label": "^2.0.0",
823+ "micromark-factory-space": "^2.0.0",
824+ "micromark-factory-title": "^2.0.0",
825+ "micromark-factory-whitespace": "^2.0.0",
826+ "micromark-util-character": "^2.0.0",
827+ "micromark-util-chunked": "^2.0.0",
828+ "micromark-util-classify-character": "^2.0.0",
829+ "micromark-util-html-tag-name": "^2.0.0",
830+ "micromark-util-normalize-identifier": "^2.0.0",
831+ "micromark-util-resolve-all": "^2.0.0",
832+ "micromark-util-subtokenize": "^2.0.0",
833+ "micromark-util-symbol": "^2.0.0",
834+ "micromark-util-types": "^2.0.0"
835+ }
836+ },
837+ "node_modules/micromark-extension-gfm": {
838+ "version": "3.0.0",
839+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm/-/micromark-extension-gfm-3.0.0.tgz",
840+ "integrity": "sha512-vsKArQsicm7t0z2GugkCKtZehqUm31oeGBV/KVSorWSy8ZlNAv7ytjFhvaryUiCUJYqs+NoE6AFhpQvBTM6Q4w==",
841+ "license": "MIT",
842+ "dependencies": {
843+ "micromark-extension-gfm-autolink-literal": "^2.0.0",
844+ "micromark-extension-gfm-footnote": "^2.0.0",
845+ "micromark-extension-gfm-strikethrough": "^2.0.0",
846+ "micromark-extension-gfm-table": "^2.0.0",
847+ "micromark-extension-gfm-tagfilter": "^2.0.0",
848+ "micromark-extension-gfm-task-list-item": "^2.0.0",
849+ "micromark-util-combine-extensions": "^2.0.0",
850+ "micromark-util-types": "^2.0.0"
851+ },
852+ "funding": {
853+ "type": "opencollective",
854+ "url": "https://opencollective.com/unified"
855+ }
856+ },
857+ "node_modules/micromark-extension-gfm-autolink-literal": {
858+ "version": "2.1.0",
859+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm-autolink-literal/-/micromark-extension-gfm-autolink-literal-2.1.0.tgz",
860+ "integrity": "sha512-oOg7knzhicgQ3t4QCjCWgTmfNhvQbDDnJeVu9v81r7NltNCVmhPy1fJRX27pISafdjL+SVc4d3l48Gb6pbRypw==",
861+ "license": "MIT",
862+ "dependencies": {
863+ "micromark-util-character": "^2.0.0",
864+ "micromark-util-sanitize-uri": "^2.0.0",
865+ "micromark-util-symbol": "^2.0.0",
866+ "micromark-util-types": "^2.0.0"
867+ },
868+ "funding": {
869+ "type": "opencollective",
870+ "url": "https://opencollective.com/unified"
871+ }
872+ },
873+ "node_modules/micromark-extension-gfm-footnote": {
874+ "version": "2.1.0",
875+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm-footnote/-/micromark-extension-gfm-footnote-2.1.0.tgz",
876+ "integrity": "sha512-/yPhxI1ntnDNsiHtzLKYnE3vf9JZ6cAisqVDauhp4CEHxlb4uoOTxOCJ+9s51bIB8U1N1FJ1RXOKTIlD5B/gqw==",
877+ "license": "MIT",
878+ "dependencies": {
879+ "devlop": "^1.0.0",
880+ "micromark-core-commonmark": "^2.0.0",
881+ "micromark-factory-space": "^2.0.0",
882+ "micromark-util-character": "^2.0.0",
883+ "micromark-util-normalize-identifier": "^2.0.0",
884+ "micromark-util-sanitize-uri": "^2.0.0",
885+ "micromark-util-symbol": "^2.0.0",
886+ "micromark-util-types": "^2.0.0"
887+ },
888+ "funding": {
889+ "type": "opencollective",
890+ "url": "https://opencollective.com/unified"
891+ }
892+ },
893+ "node_modules/micromark-extension-gfm-strikethrough": {
894+ "version": "2.1.0",
895+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm-strikethrough/-/micromark-extension-gfm-strikethrough-2.1.0.tgz",
896+ "integrity": "sha512-ADVjpOOkjz1hhkZLlBiYA9cR2Anf8F4HqZUO6e5eDcPQd0Txw5fxLzzxnEkSkfnD0wziSGiv7sYhk/ktvbf1uw==",
897+ "license": "MIT",
898+ "dependencies": {
899+ "devlop": "^1.0.0",
900+ "micromark-util-chunked": "^2.0.0",
901+ "micromark-util-classify-character": "^2.0.0",
902+ "micromark-util-resolve-all": "^2.0.0",
903+ "micromark-util-symbol": "^2.0.0",
904+ "micromark-util-types": "^2.0.0"
905+ },
906+ "funding": {
907+ "type": "opencollective",
908+ "url": "https://opencollective.com/unified"
909+ }
910+ },
911+ "node_modules/micromark-extension-gfm-table": {
912+ "version": "2.1.1",
913+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm-table/-/micromark-extension-gfm-table-2.1.1.tgz",
914+ "integrity": "sha512-t2OU/dXXioARrC6yWfJ4hqB7rct14e8f7m0cbI5hUmDyyIlwv5vEtooptH8INkbLzOatzKuVbQmAYcbWoyz6Dg==",
915+ "license": "MIT",
916+ "dependencies": {
917+ "devlop": "^1.0.0",
918+ "micromark-factory-space": "^2.0.0",
919+ "micromark-util-character": "^2.0.0",
920+ "micromark-util-symbol": "^2.0.0",
921+ "micromark-util-types": "^2.0.0"
922+ },
923+ "funding": {
924+ "type": "opencollective",
925+ "url": "https://opencollective.com/unified"
926+ }
927+ },
928+ "node_modules/micromark-extension-gfm-tagfilter": {
929+ "version": "2.0.0",
930+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm-tagfilter/-/micromark-extension-gfm-tagfilter-2.0.0.tgz",
931+ "integrity": "sha512-xHlTOmuCSotIA8TW1mDIM6X2O1SiX5P9IuDtqGonFhEK0qgRI4yeC6vMxEV2dgyr2TiD+2PQ10o+cOhdVAcwfg==",
932+ "license": "MIT",
933+ "dependencies": {
934+ "micromark-util-types": "^2.0.0"
935+ },
936+ "funding": {
937+ "type": "opencollective",
938+ "url": "https://opencollective.com/unified"
939+ }
940+ },
941+ "node_modules/micromark-extension-gfm-task-list-item": {
942+ "version": "2.1.0",
943+ "resolved": "https://registry.npmjs.org/micromark-extension-gfm-task-list-item/-/micromark-extension-gfm-task-list-item-2.1.0.tgz",
944+ "integrity": "sha512-qIBZhqxqI6fjLDYFTBIa4eivDMnP+OZqsNwmQ3xNLE4Cxwc+zfQEfbs6tzAo2Hjq+bh6q5F+Z8/cksrLFYWQQw==",
945+ "license": "MIT",
946+ "dependencies": {
947+ "devlop": "^1.0.0",
948+ "micromark-factory-space": "^2.0.0",
949+ "micromark-util-character": "^2.0.0",
950+ "micromark-util-symbol": "^2.0.0",
951+ "micromark-util-types": "^2.0.0"
952+ },
953+ "funding": {
954+ "type": "opencollective",
955+ "url": "https://opencollective.com/unified"
956+ }
957+ },
958+ "node_modules/micromark-factory-destination": {
959+ "version": "2.0.1",
960+ "resolved": "https://registry.npmjs.org/micromark-factory-destination/-/micromark-factory-destination-2.0.1.tgz",
961+ "integrity": "sha512-Xe6rDdJlkmbFRExpTOmRj9N3MaWmbAgdpSrBQvCFqhezUn4AHqJHbaEnfbVYYiexVSs//tqOdY/DxhjdCiJnIA==",
962+ "funding": [
963+ {
964+ "type": "GitHub Sponsors",
965+ "url": "https://github.com/sponsors/unifiedjs"
966+ },
967+ {
968+ "type": "OpenCollective",
969+ "url": "https://opencollective.com/unified"
970+ }
971+ ],
972+ "license": "MIT",
973+ "dependencies": {
974+ "micromark-util-character": "^2.0.0",
975+ "micromark-util-symbol": "^2.0.0",
976+ "micromark-util-types": "^2.0.0"
977+ }
978+ },
979+ "node_modules/micromark-factory-label": {
980+ "version": "2.0.1",
981+ "resolved": "https://registry.npmjs.org/micromark-factory-label/-/micromark-factory-label-2.0.1.tgz",
982+ "integrity": "sha512-VFMekyQExqIW7xIChcXn4ok29YE3rnuyveW3wZQWWqF4Nv9Wk5rgJ99KzPvHjkmPXF93FXIbBp6YdW3t71/7Vg==",
983+ "funding": [
984+ {
985+ "type": "GitHub Sponsors",
986+ "url": "https://github.com/sponsors/unifiedjs"
987+ },
988+ {
989+ "type": "OpenCollective",
990+ "url": "https://opencollective.com/unified"
991+ }
992+ ],
993+ "license": "MIT",
994+ "dependencies": {
995+ "devlop": "^1.0.0",
996+ "micromark-util-character": "^2.0.0",
997+ "micromark-util-symbol": "^2.0.0",
998+ "micromark-util-types": "^2.0.0"
999+ }
1000+ },
1001+ "node_modules/micromark-factory-space": {
1002+ "version": "2.0.1",
1003+ "resolved": "https://registry.npmjs.org/micromark-factory-space/-/micromark-factory-space-2.0.1.tgz",
1004+ "integrity": "sha512-zRkxjtBxxLd2Sc0d+fbnEunsTj46SWXgXciZmHq0kDYGnck/ZSGj9/wULTV95uoeYiK5hRXP2mJ98Uo4cq/LQg==",
1005+ "funding": [
1006+ {
1007+ "type": "GitHub Sponsors",
1008+ "url": "https://github.com/sponsors/unifiedjs"
1009+ },
1010+ {
1011+ "type": "OpenCollective",
1012+ "url": "https://opencollective.com/unified"
1013+ }
1014+ ],
1015+ "license": "MIT",
1016+ "dependencies": {
1017+ "micromark-util-character": "^2.0.0",
1018+ "micromark-util-types": "^2.0.0"
1019+ }
1020+ },
1021+ "node_modules/micromark-factory-title": {
1022+ "version": "2.0.1",
1023+ "resolved": "https://registry.npmjs.org/micromark-factory-title/-/micromark-factory-title-2.0.1.tgz",
1024+ "integrity": "sha512-5bZ+3CjhAd9eChYTHsjy6TGxpOFSKgKKJPJxr293jTbfry2KDoWkhBb6TcPVB4NmzaPhMs1Frm9AZH7OD4Cjzw==",
1025+ "funding": [
1026+ {
1027+ "type": "GitHub Sponsors",
1028+ "url": "https://github.com/sponsors/unifiedjs"
1029+ },
1030+ {
1031+ "type": "OpenCollective",
1032+ "url": "https://opencollective.com/unified"
1033+ }
1034+ ],
1035+ "license": "MIT",
1036+ "dependencies": {
1037+ "micromark-factory-space": "^2.0.0",
1038+ "micromark-util-character": "^2.0.0",
1039+ "micromark-util-symbol": "^2.0.0",
1040+ "micromark-util-types": "^2.0.0"
1041+ }
1042+ },
1043+ "node_modules/micromark-factory-whitespace": {
1044+ "version": "2.0.1",
1045+ "resolved": "https://registry.npmjs.org/micromark-factory-whitespace/-/micromark-factory-whitespace-2.0.1.tgz",
1046+ "integrity": "sha512-Ob0nuZ3PKt/n0hORHyvoD9uZhr+Za8sFoP+OnMcnWK5lngSzALgQYKMr9RJVOWLqQYuyn6ulqGWSXdwf6F80lQ==",
1047+ "funding": [
1048+ {
1049+ "type": "GitHub Sponsors",
1050+ "url": "https://github.com/sponsors/unifiedjs"
1051+ },
1052+ {
1053+ "type": "OpenCollective",
1054+ "url": "https://opencollective.com/unified"
1055+ }
1056+ ],
1057+ "license": "MIT",
1058+ "dependencies": {
1059+ "micromark-factory-space": "^2.0.0",
1060+ "micromark-util-character": "^2.0.0",
1061+ "micromark-util-symbol": "^2.0.0",
1062+ "micromark-util-types": "^2.0.0"
1063+ }
1064+ },
1065+ "node_modules/micromark-util-character": {
1066+ "version": "2.1.1",
1067+ "resolved": "https://registry.npmjs.org/micromark-util-character/-/micromark-util-character-2.1.1.tgz",
1068+ "integrity": "sha512-wv8tdUTJ3thSFFFJKtpYKOYiGP2+v96Hvk4Tu8KpCAsTMs6yi+nVmGh1syvSCsaxz45J6Jbw+9DD6g97+NV67Q==",
1069+ "funding": [
1070+ {
1071+ "type": "GitHub Sponsors",
1072+ "url": "https://github.com/sponsors/unifiedjs"
1073+ },
1074+ {
1075+ "type": "OpenCollective",
1076+ "url": "https://opencollective.com/unified"
1077+ }
1078+ ],
1079+ "license": "MIT",
1080+ "dependencies": {
1081+ "micromark-util-symbol": "^2.0.0",
1082+ "micromark-util-types": "^2.0.0"
1083+ }
1084+ },
1085+ "node_modules/micromark-util-chunked": {
1086+ "version": "2.0.1",
1087+ "resolved": "https://registry.npmjs.org/micromark-util-chunked/-/micromark-util-chunked-2.0.1.tgz",
1088+ "integrity": "sha512-QUNFEOPELfmvv+4xiNg2sRYeS/P84pTW0TCgP5zc9FpXetHY0ab7SxKyAQCNCc1eK0459uoLI1y5oO5Vc1dbhA==",
1089+ "funding": [
1090+ {
1091+ "type": "GitHub Sponsors",
1092+ "url": "https://github.com/sponsors/unifiedjs"
1093+ },
1094+ {
1095+ "type": "OpenCollective",
1096+ "url": "https://opencollective.com/unified"
1097+ }
1098+ ],
1099+ "license": "MIT",
1100+ "dependencies": {
1101+ "micromark-util-symbol": "^2.0.0"
1102+ }
1103+ },
1104+ "node_modules/micromark-util-classify-character": {
1105+ "version": "2.0.1",
1106+ "resolved": "https://registry.npmjs.org/micromark-util-classify-character/-/micromark-util-classify-character-2.0.1.tgz",
1107+ "integrity": "sha512-K0kHzM6afW/MbeWYWLjoHQv1sgg2Q9EccHEDzSkxiP/EaagNzCm7T/WMKZ3rjMbvIpvBiZgwR3dKMygtA4mG1Q==",
1108+ "funding": [
1109+ {
1110+ "type": "GitHub Sponsors",
1111+ "url": "https://github.com/sponsors/unifiedjs"
1112+ },
1113+ {
1114+ "type": "OpenCollective",
1115+ "url": "https://opencollective.com/unified"
1116+ }
1117+ ],
1118+ "license": "MIT",
1119+ "dependencies": {
1120+ "micromark-util-character": "^2.0.0",
1121+ "micromark-util-symbol": "^2.0.0",
1122+ "micromark-util-types": "^2.0.0"
1123+ }
1124+ },
1125+ "node_modules/micromark-util-combine-extensions": {
1126+ "version": "2.0.1",
1127+ "resolved": "https://registry.npmjs.org/micromark-util-combine-extensions/-/micromark-util-combine-extensions-2.0.1.tgz",
1128+ "integrity": "sha512-OnAnH8Ujmy59JcyZw8JSbK9cGpdVY44NKgSM7E9Eh7DiLS2E9RNQf0dONaGDzEG9yjEl5hcqeIsj4hfRkLH/Bg==",
1129+ "funding": [
1130+ {
1131+ "type": "GitHub Sponsors",
1132+ "url": "https://github.com/sponsors/unifiedjs"
1133+ },
1134+ {
1135+ "type": "OpenCollective",
1136+ "url": "https://opencollective.com/unified"
1137+ }
1138+ ],
1139+ "license": "MIT",
1140+ "dependencies": {
1141+ "micromark-util-chunked": "^2.0.0",
1142+ "micromark-util-types": "^2.0.0"
1143+ }
1144+ },
1145+ "node_modules/micromark-util-decode-numeric-character-reference": {
1146+ "version": "2.0.2",
1147+ "resolved": "https://registry.npmjs.org/micromark-util-decode-numeric-character-reference/-/micromark-util-decode-numeric-character-reference-2.0.2.tgz",
1148+ "integrity": "sha512-ccUbYk6CwVdkmCQMyr64dXz42EfHGkPQlBj5p7YVGzq8I7CtjXZJrubAYezf7Rp+bjPseiROqe7G6foFd+lEuw==",
1149+ "funding": [
1150+ {
1151+ "type": "GitHub Sponsors",
1152+ "url": "https://github.com/sponsors/unifiedjs"
1153+ },
1154+ {
1155+ "type": "OpenCollective",
1156+ "url": "https://opencollective.com/unified"
1157+ }
1158+ ],
1159+ "license": "MIT",
1160+ "dependencies": {
1161+ "micromark-util-symbol": "^2.0.0"
1162+ }
1163+ },
1164+ "node_modules/micromark-util-decode-string": {
1165+ "version": "2.0.1",
1166+ "resolved": "https://registry.npmjs.org/micromark-util-decode-string/-/micromark-util-decode-string-2.0.1.tgz",
1167+ "integrity": "sha512-nDV/77Fj6eH1ynwscYTOsbK7rR//Uj0bZXBwJZRfaLEJ1iGBR6kIfNmlNqaqJf649EP0F3NWNdeJi03elllNUQ==",
1168+ "funding": [
1169+ {
1170+ "type": "GitHub Sponsors",
1171+ "url": "https://github.com/sponsors/unifiedjs"
1172+ },
1173+ {
1174+ "type": "OpenCollective",
1175+ "url": "https://opencollective.com/unified"
1176+ }
1177+ ],
1178+ "license": "MIT",
1179+ "dependencies": {
1180+ "decode-named-character-reference": "^1.0.0",
1181+ "micromark-util-character": "^2.0.0",
1182+ "micromark-util-decode-numeric-character-reference": "^2.0.0",
1183+ "micromark-util-symbol": "^2.0.0"
1184+ }
1185+ },
1186+ "node_modules/micromark-util-encode": {
1187+ "version": "2.0.1",
1188+ "resolved": "https://registry.npmjs.org/micromark-util-encode/-/micromark-util-encode-2.0.1.tgz",
1189+ "integrity": "sha512-c3cVx2y4KqUnwopcO9b/SCdo2O67LwJJ/UyqGfbigahfegL9myoEFoDYZgkT7f36T0bLrM9hZTAaAyH+PCAXjw==",
1190+ "funding": [
1191+ {
1192+ "type": "GitHub Sponsors",
1193+ "url": "https://github.com/sponsors/unifiedjs"
1194+ },
1195+ {
1196+ "type": "OpenCollective",
1197+ "url": "https://opencollective.com/unified"
1198+ }
1199+ ],
1200+ "license": "MIT"
1201+ },
1202+ "node_modules/micromark-util-html-tag-name": {
1203+ "version": "2.0.1",
1204+ "resolved": "https://registry.npmjs.org/micromark-util-html-tag-name/-/micromark-util-html-tag-name-2.0.1.tgz",
1205+ "integrity": "sha512-2cNEiYDhCWKI+Gs9T0Tiysk136SnR13hhO8yW6BGNyhOC4qYFnwF1nKfD3HFAIXA5c45RrIG1ub11GiXeYd1xA==",
1206+ "funding": [
1207+ {
1208+ "type": "GitHub Sponsors",
1209+ "url": "https://github.com/sponsors/unifiedjs"
1210+ },
1211+ {
1212+ "type": "OpenCollective",
1213+ "url": "https://opencollective.com/unified"
1214+ }
1215+ ],
1216+ "license": "MIT"
1217+ },
1218+ "node_modules/micromark-util-normalize-identifier": {
1219+ "version": "2.0.1",
1220+ "resolved": "https://registry.npmjs.org/micromark-util-normalize-identifier/-/micromark-util-normalize-identifier-2.0.1.tgz",
1221+ "integrity": "sha512-sxPqmo70LyARJs0w2UclACPUUEqltCkJ6PhKdMIDuJ3gSf/Q+/GIe3WKl0Ijb/GyH9lOpUkRAO2wp0GVkLvS9Q==",
1222+ "funding": [
1223+ {
1224+ "type": "GitHub Sponsors",
1225+ "url": "https://github.com/sponsors/unifiedjs"
1226+ },
1227+ {
1228+ "type": "OpenCollective",
1229+ "url": "https://opencollective.com/unified"
1230+ }
1231+ ],
1232+ "license": "MIT",
1233+ "dependencies": {
1234+ "micromark-util-symbol": "^2.0.0"
1235+ }
1236+ },
1237+ "node_modules/micromark-util-resolve-all": {
1238+ "version": "2.0.1",
1239+ "resolved": "https://registry.npmjs.org/micromark-util-resolve-all/-/micromark-util-resolve-all-2.0.1.tgz",
1240+ "integrity": "sha512-VdQyxFWFT2/FGJgwQnJYbe1jjQoNTS4RjglmSjTUlpUMa95Htx9NHeYW4rGDJzbjvCsl9eLjMQwGeElsqmzcHg==",
1241+ "funding": [
1242+ {
1243+ "type": "GitHub Sponsors",
1244+ "url": "https://github.com/sponsors/unifiedjs"
1245+ },
1246+ {
1247+ "type": "OpenCollective",
1248+ "url": "https://opencollective.com/unified"
1249+ }
1250+ ],
1251+ "license": "MIT",
1252+ "dependencies": {
1253+ "micromark-util-types": "^2.0.0"
1254+ }
1255+ },
1256+ "node_modules/micromark-util-sanitize-uri": {
1257+ "version": "2.0.1",
1258+ "resolved": "https://registry.npmjs.org/micromark-util-sanitize-uri/-/micromark-util-sanitize-uri-2.0.1.tgz",
1259+ "integrity": "sha512-9N9IomZ/YuGGZZmQec1MbgxtlgougxTodVwDzzEouPKo3qFWvymFHWcnDi2vzV1ff6kas9ucW+o3yzJK9YB1AQ==",
1260+ "funding": [
1261+ {
1262+ "type": "GitHub Sponsors",
1263+ "url": "https://github.com/sponsors/unifiedjs"
1264+ },
1265+ {
1266+ "type": "OpenCollective",
1267+ "url": "https://opencollective.com/unified"
1268+ }
1269+ ],
1270+ "license": "MIT",
1271+ "dependencies": {
1272+ "micromark-util-character": "^2.0.0",
1273+ "micromark-util-encode": "^2.0.0",
1274+ "micromark-util-symbol": "^2.0.0"
1275+ }
1276+ },
1277+ "node_modules/micromark-util-subtokenize": {
1278+ "version": "2.1.0",
1279+ "resolved": "https://registry.npmjs.org/micromark-util-subtokenize/-/micromark-util-subtokenize-2.1.0.tgz",
1280+ "integrity": "sha512-XQLu552iSctvnEcgXw6+Sx75GflAPNED1qx7eBJ+wydBb2KCbRZe+NwvIEEMM83uml1+2WSXpBAcp9IUCgCYWA==",
1281+ "funding": [
1282+ {
1283+ "type": "GitHub Sponsors",
1284+ "url": "https://github.com/sponsors/unifiedjs"
1285+ },
1286+ {
1287+ "type": "OpenCollective",
1288+ "url": "https://opencollective.com/unified"
1289+ }
1290+ ],
1291+ "license": "MIT",
1292+ "dependencies": {
1293+ "devlop": "^1.0.0",
1294+ "micromark-util-chunked": "^2.0.0",
1295+ "micromark-util-symbol": "^2.0.0",
1296+ "micromark-util-types": "^2.0.0"
1297+ }
1298+ },
1299+ "node_modules/micromark-util-symbol": {
1300+ "version": "2.0.1",
1301+ "resolved": "https://registry.npmjs.org/micromark-util-symbol/-/micromark-util-symbol-2.0.1.tgz",
1302+ "integrity": "sha512-vs5t8Apaud9N28kgCrRUdEed4UJ+wWNvicHLPxCa9ENlYuAY31M0ETy5y1vA33YoNPDFTghEbnh6efaE8h4x0Q==",
1303+ "funding": [
1304+ {
1305+ "type": "GitHub Sponsors",
1306+ "url": "https://github.com/sponsors/unifiedjs"
1307+ },
1308+ {
1309+ "type": "OpenCollective",
1310+ "url": "https://opencollective.com/unified"
1311+ }
1312+ ],
1313+ "license": "MIT"
1314+ },
1315+ "node_modules/micromark-util-types": {
1316+ "version": "2.0.2",
1317+ "resolved": "https://registry.npmjs.org/micromark-util-types/-/micromark-util-types-2.0.2.tgz",
1318+ "integrity": "sha512-Yw0ECSpJoViF1qTU4DC6NwtC4aWGt1EkzaQB8KPPyCRR8z9TWeV0HbEFGTO+ZY1wB22zmxnJqhPyTpOVCpeHTA==",
1319+ "funding": [
1320+ {
1321+ "type": "GitHub Sponsors",
1322+ "url": "https://github.com/sponsors/unifiedjs"
1323+ },
1324+ {
1325+ "type": "OpenCollective",
1326+ "url": "https://opencollective.com/unified"
1327+ }
1328+ ],
1329+ "license": "MIT"
1330+ },
1331+ "node_modules/ms": {
1332+ "version": "2.1.3",
1333+ "resolved": "https://registry.npmjs.org/ms/-/ms-2.1.3.tgz",
1334+ "integrity": "sha512-6FlzubTLZG3J2a/NVCAleEhjzq5oxgHyaCU9yYXvcLsvoVaHJq/s5xXI6/XXP6tz7R9xAOtHnSO/tXtF3WRTlA==",
1335+ "license": "MIT"
1336+ },
1337+ "node_modules/node-addon-api": {
1338+ "version": "8.9.2",
1339+ "resolved": "https://registry.npmjs.org/node-addon-api/-/node-addon-api-8.9.2.tgz",
1340+ "integrity": "sha512-VijLXbi3UACN69I0JVXJsX4tjACjNoQDgv2gTF6sx2wWEi8tkSg2eX8p5gSIFi8z2+DL3oHmY6OyKce38SDolg==",
1341+ "license": "MIT",
1342+ "engines": {
1343+ "node": "^18 || ^20 || >= 21"
1344+ }
1345+ },
1346+ "node_modules/numbl": {
1347+ "version": "0.4.19",
1348+ "resolved": "https://registry.npmjs.org/numbl/-/numbl-0.4.19.tgz",
1349+ "integrity": "sha512-wfTAhrdOSuzUoiglglTwykwEguU0Vl5qBJNhj2hyP8FzHACd/w1qX69OhEC7sbyE+gGEH9gjPjFsBkoi9UTGqw==",
1350+ "hasInstallScript": true,
1351+ "license": "Apache-2.0",
1352+ "dependencies": {
1353+ "fflate": "^0.8.2",
1354+ "h5wasm": "^0.10.3",
1355+ "node-addon-api": "^8.3.0",
1356+ "pako": "^2.1.0",
1357+ "qhull-wasm": "^0.0.1",
1358+ "react-markdown": "^10.1.0",
1359+ "react-syntax-highlighter": "^16.1.1",
1360+ "remark-gfm": "^4.0.1",
1361+ "three": "^0.183.2"
1362+ },
1363+ "bin": {
1364+ "numbl": "dist-cli/cli.js"
1365+ },
1366+ "engines": {
1367+ "node": ">=18"
1368+ },
1369+ "optionalDependencies": {
1370+ "koffi": "^2.15.2"
1371+ },
1372+ "peerDependencies": {
1373+ "react": ">=18",
1374+ "react-dom": ">=18"
1375+ },
1376+ "peerDependenciesMeta": {
1377+ "react": {
1378+ "optional": true
1379+ },
1380+ "react-dom": {
1381+ "optional": true
1382+ }
1383+ }
1384+ },
1385+ "node_modules/pako": {
1386+ "version": "2.2.0",
1387+ "resolved": "https://registry.npmjs.org/pako/-/pako-2.2.0.tgz",
1388+ "integrity": "sha512-zJq6RP/5q+TO2OpFV3FHzlPnFjmkb7Nc99a5SNjJE+uu/PkpChs+NIZSSzbBoD+6kjiISXjfYdwj1ZRQ81dz/w==",
1389+ "funding": [
1390+ {
1391+ "type": "github",
1392+ "url": "https://github.com/sponsors/puzrin"
1393+ },
1394+ {
1395+ "type": "github",
1396+ "url": "https://github.com/sponsors/nodeca"
1397+ }
1398+ ],
1399+ "license": "(MIT AND Zlib)"
1400+ },
1401+ "node_modules/parse-entities": {
1402+ "version": "4.0.2",
1403+ "resolved": "https://registry.npmjs.org/parse-entities/-/parse-entities-4.0.2.tgz",
1404+ "integrity": "sha512-GG2AQYWoLgL877gQIKeRPGO1xF9+eG1ujIb5soS5gPvLQ1y2o8FL90w2QWNdf9I361Mpp7726c+lj3U0qK1uGw==",
1405+ "license": "MIT",
1406+ "dependencies": {
1407+ "@types/unist": "^2.0.0",
1408+ "character-entities-legacy": "^3.0.0",
1409+ "character-reference-invalid": "^2.0.0",
1410+ "decode-named-character-reference": "^1.0.0",
1411+ "is-alphanumerical": "^2.0.0",
1412+ "is-decimal": "^2.0.0",
1413+ "is-hexadecimal": "^2.0.0"
1414+ },
1415+ "funding": {
1416+ "type": "github",
1417+ "url": "https://github.com/sponsors/wooorm"
1418+ }
1419+ },
1420+ "node_modules/parse-entities/node_modules/@types/unist": {
1421+ "version": "2.0.11",
1422+ "resolved": "https://registry.npmjs.org/@types/unist/-/unist-2.0.11.tgz",
1423+ "integrity": "sha512-CmBKiL6NNo/OqgmMn95Fk9Whlp2mtvIv+KNpQKN2F4SjvrEesubTRWGYSg+BnWZOnlCaSTU1sMpsBOzgbYhnsA==",
1424+ "license": "MIT"
1425+ },
1426+ "node_modules/prismjs": {
1427+ "version": "1.30.0",
1428+ "resolved": "https://registry.npmjs.org/prismjs/-/prismjs-1.30.0.tgz",
1429+ "integrity": "sha512-DEvV2ZF2r2/63V+tK8hQvrR2ZGn10srHbXviTlcv7Kpzw8jWiNTqbVgjO3IY8RxrrOUF8VPMQQFysYYYv0YZxw==",
1430+ "license": "MIT",
1431+ "engines": {
1432+ "node": ">=6"
1433+ }
1434+ },
1435+ "node_modules/property-information": {
1436+ "version": "7.2.0",
1437+ "resolved": "https://registry.npmjs.org/property-information/-/property-information-7.2.0.tgz",
1438+ "integrity": "sha512-IAtzIB6sUiWaJYrX9smp3V46pBGbBeLFRGdh25kg1334VcBlD8HzhPeNIWQH9zhGmo2itIe25EHt9dQP7G5hmg==",
1439+ "license": "MIT",
1440+ "funding": {
1441+ "type": "github",
1442+ "url": "https://github.com/sponsors/wooorm"
1443+ }
1444+ },
1445+ "node_modules/qhull-wasm": {
1446+ "version": "0.0.1",
1447+ "resolved": "https://registry.npmjs.org/qhull-wasm/-/qhull-wasm-0.0.1.tgz",
1448+ "integrity": "sha512-A09kumKhdRIX0YNR+Y2G0bEuqCJ3xFOH61+x6uajHqDyvRydlZKWhH1AGGsM4PCyGm/cPFVbqOUrPJpq5n8HeA==",
1449+ "license": "(MIT AND Qhull)",
1450+ "engines": {
1451+ "node": ">=18"
1452+ }
1453+ },
1454+ "node_modules/react": {
1455+ "version": "19.2.8",
1456+ "resolved": "https://registry.npmjs.org/react/-/react-19.2.8.tgz",
1457+ "integrity": "sha512-PWaYA1L/q9u2u7xYQi+Y3L3Yfnie7XyLeaJICV1MGD6LprsBxcAqGjYyr0eY3p+QdsA+x/Irkt4Qif8D63+Sbw==",
1458+ "license": "MIT",
1459+ "peer": true,
1460+ "engines": {
1461+ "node": ">=0.10.0"
1462+ }
1463+ },
1464+ "node_modules/react-markdown": {
1465+ "version": "10.1.0",
1466+ "resolved": "https://registry.npmjs.org/react-markdown/-/react-markdown-10.1.0.tgz",
1467+ "integrity": "sha512-qKxVopLT/TyA6BX3Ue5NwabOsAzm0Q7kAPwq6L+wWDwisYs7R8vZ0nRXqq6rkueboxpkjvLGU9fWifiX/ZZFxQ==",
1468+ "license": "MIT",
1469+ "dependencies": {
1470+ "@types/hast": "^3.0.0",
1471+ "@types/mdast": "^4.0.0",
1472+ "devlop": "^1.0.0",
1473+ "hast-util-to-jsx-runtime": "^2.0.0",
1474+ "html-url-attributes": "^3.0.0",
1475+ "mdast-util-to-hast": "^13.0.0",
1476+ "remark-parse": "^11.0.0",
1477+ "remark-rehype": "^11.0.0",
1478+ "unified": "^11.0.0",
1479+ "unist-util-visit": "^5.0.0",
1480+ "vfile": "^6.0.0"
1481+ },
1482+ "funding": {
1483+ "type": "opencollective",
1484+ "url": "https://opencollective.com/unified"
1485+ },
1486+ "peerDependencies": {
1487+ "@types/react": ">=18",
1488+ "react": ">=18"
1489+ }
1490+ },
1491+ "node_modules/react-syntax-highlighter": {
1492+ "version": "16.1.1",
1493+ "resolved": "https://registry.npmjs.org/react-syntax-highlighter/-/react-syntax-highlighter-16.1.1.tgz",
1494+ "integrity": "sha512-PjVawBGy80C6YbC5DDZJeUjBmC7skaoEUdvfFQediQHgCL7aKyVHe57SaJGfQsloGDac+gCpTfRdtxzWWKmCXA==",
1495+ "license": "MIT",
1496+ "dependencies": {
1497+ "@babel/runtime": "^7.28.4",
1498+ "highlight.js": "^10.4.1",
1499+ "highlightjs-vue": "^1.0.0",
1500+ "lowlight": "^1.17.0",
1501+ "prismjs": "^1.30.0",
1502+ "refractor": "^5.0.0"
1503+ },
1504+ "engines": {
1505+ "node": ">= 16.20.2"
1506+ },
1507+ "peerDependencies": {
1508+ "react": ">= 0.14.0"
1509+ }
1510+ },
1511+ "node_modules/refractor": {
1512+ "version": "5.0.0",
1513+ "resolved": "https://registry.npmjs.org/refractor/-/refractor-5.0.0.tgz",
1514+ "integrity": "sha512-QXOrHQF5jOpjjLfiNk5GFnWhRXvxjUVnlFxkeDmewR5sXkr3iM46Zo+CnRR8B+MDVqkULW4EcLVcRBNOPXHosw==",
1515+ "license": "MIT",
1516+ "dependencies": {
1517+ "@types/hast": "^3.0.0",
1518+ "@types/prismjs": "^1.0.0",
1519+ "hastscript": "^9.0.0",
1520+ "parse-entities": "^4.0.0"
1521+ },
1522+ "funding": {
1523+ "type": "github",
1524+ "url": "https://github.com/sponsors/wooorm"
1525+ }
1526+ },
1527+ "node_modules/remark-gfm": {
1528+ "version": "4.0.1",
1529+ "resolved": "https://registry.npmjs.org/remark-gfm/-/remark-gfm-4.0.1.tgz",
1530+ "integrity": "sha512-1quofZ2RQ9EWdeN34S79+KExV1764+wCUGop5CPL1WGdD0ocPpu91lzPGbwWMECpEpd42kJGQwzRfyov9j4yNg==",
1531+ "license": "MIT",
1532+ "dependencies": {
1533+ "@types/mdast": "^4.0.0",
1534+ "mdast-util-gfm": "^3.0.0",
1535+ "micromark-extension-gfm": "^3.0.0",
1536+ "remark-parse": "^11.0.0",
1537+ "remark-stringify": "^11.0.0",
1538+ "unified": "^11.0.0"
1539+ },
1540+ "funding": {
1541+ "type": "opencollective",
1542+ "url": "https://opencollective.com/unified"
1543+ }
1544+ },
1545+ "node_modules/remark-parse": {
1546+ "version": "11.0.0",
1547+ "resolved": "https://registry.npmjs.org/remark-parse/-/remark-parse-11.0.0.tgz",
1548+ "integrity": "sha512-FCxlKLNGknS5ba/1lmpYijMUzX2esxW5xQqjWxw2eHFfS2MSdaHVINFmhjo+qN1WhZhNimq0dZATN9pH0IDrpA==",
1549+ "license": "MIT",
1550+ "dependencies": {
1551+ "@types/mdast": "^4.0.0",
1552+ "mdast-util-from-markdown": "^2.0.0",
1553+ "micromark-util-types": "^2.0.0",
1554+ "unified": "^11.0.0"
1555+ },
1556+ "funding": {
1557+ "type": "opencollective",
1558+ "url": "https://opencollective.com/unified"
1559+ }
1560+ },
1561+ "node_modules/remark-rehype": {
1562+ "version": "11.1.2",
1563+ "resolved": "https://registry.npmjs.org/remark-rehype/-/remark-rehype-11.1.2.tgz",
1564+ "integrity": "sha512-Dh7l57ianaEoIpzbp0PC9UKAdCSVklD8E5Rpw7ETfbTl3FqcOOgq5q2LVDhgGCkaBv7p24JXikPdvhhmHvKMsw==",
1565+ "license": "MIT",
1566+ "dependencies": {
1567+ "@types/hast": "^3.0.0",
1568+ "@types/mdast": "^4.0.0",
1569+ "mdast-util-to-hast": "^13.0.0",
1570+ "unified": "^11.0.0",
1571+ "vfile": "^6.0.0"
1572+ },
1573+ "funding": {
1574+ "type": "opencollective",
1575+ "url": "https://opencollective.com/unified"
1576+ }
1577+ },
1578+ "node_modules/remark-stringify": {
1579+ "version": "11.0.0",
1580+ "resolved": "https://registry.npmjs.org/remark-stringify/-/remark-stringify-11.0.0.tgz",
1581+ "integrity": "sha512-1OSmLd3awB/t8qdoEOMazZkNsfVTeY4fTsgzcQFdXNq8ToTN4ZGwrMnlda4K6smTFKD+GRV6O48i6Z4iKgPPpw==",
1582+ "license": "MIT",
1583+ "dependencies": {
1584+ "@types/mdast": "^4.0.0",
1585+ "mdast-util-to-markdown": "^2.0.0",
1586+ "unified": "^11.0.0"
1587+ },
1588+ "funding": {
1589+ "type": "opencollective",
1590+ "url": "https://opencollective.com/unified"
1591+ }
1592+ },
1593+ "node_modules/space-separated-tokens": {
1594+ "version": "2.0.2",
1595+ "resolved": "https://registry.npmjs.org/space-separated-tokens/-/space-separated-tokens-2.0.2.tgz",
1596+ "integrity": "sha512-PEGlAwrG8yXGXRjW32fGbg66JAlOAwbObuqVoJpv/mRgoWDQfgH1wDPvtzWyUSNAXBGSk8h755YDbbcEy3SH2Q==",
1597+ "license": "MIT",
1598+ "funding": {
1599+ "type": "github",
1600+ "url": "https://github.com/sponsors/wooorm"
1601+ }
1602+ },
1603+ "node_modules/stringify-entities": {
1604+ "version": "4.0.4",
1605+ "resolved": "https://registry.npmjs.org/stringify-entities/-/stringify-entities-4.0.4.tgz",
1606+ "integrity": "sha512-IwfBptatlO+QCJUo19AqvrPNqlVMpW9YEL2LIVY+Rpv2qsjCGxaDLNRgeGsQWJhfItebuJhsGSLjaBbNSQ+ieg==",
1607+ "license": "MIT",
1608+ "dependencies": {
1609+ "character-entities-html4": "^2.0.0",
1610+ "character-entities-legacy": "^3.0.0"
1611+ },
1612+ "funding": {
1613+ "type": "github",
1614+ "url": "https://github.com/sponsors/wooorm"
1615+ }
1616+ },
1617+ "node_modules/style-to-js": {
1618+ "version": "1.1.21",
1619+ "resolved": "https://registry.npmjs.org/style-to-js/-/style-to-js-1.1.21.tgz",
1620+ "integrity": "sha512-RjQetxJrrUJLQPHbLku6U/ocGtzyjbJMP9lCNK7Ag0CNh690nSH8woqWH9u16nMjYBAok+i7JO1NP2pOy8IsPQ==",
1621+ "license": "MIT",
1622+ "dependencies": {
1623+ "style-to-object": "1.0.14"
1624+ }
1625+ },
1626+ "node_modules/style-to-object": {
1627+ "version": "1.0.14",
1628+ "resolved": "https://registry.npmjs.org/style-to-object/-/style-to-object-1.0.14.tgz",
1629+ "integrity": "sha512-LIN7rULI0jBscWQYaSswptyderlarFkjQ+t79nzty8tcIAceVomEVlLzH5VP4Cmsv6MtKhs7qaAiwlcp+Mgaxw==",
1630+ "license": "MIT",
1631+ "dependencies": {
1632+ "inline-style-parser": "0.2.7"
1633+ }
1634+ },
1635+ "node_modules/three": {
1636+ "version": "0.183.2",
1637+ "resolved": "https://registry.npmjs.org/three/-/three-0.183.2.tgz",
1638+ "integrity": "sha512-di3BsL2FEQ1PA7Hcvn4fyJOlxRRgFYBpMTcyOgkwJIaDOdJMebEFPA+t98EvjuljDx4hNulAGwF6KIjtwI5jgQ==",
1639+ "license": "MIT"
1640+ },
1641+ "node_modules/trim-lines": {
1642+ "version": "3.0.1",
1643+ "resolved": "https://registry.npmjs.org/trim-lines/-/trim-lines-3.0.1.tgz",
1644+ "integrity": "sha512-kRj8B+YHZCc9kQYdWfJB2/oUl9rA99qbowYYBtr4ui4mZyAQ2JpvVBd/6U2YloATfqBhBTSMhTpgBHtU0Mf3Rg==",
1645+ "license": "MIT",
1646+ "funding": {
1647+ "type": "github",
1648+ "url": "https://github.com/sponsors/wooorm"
1649+ }
1650+ },
1651+ "node_modules/trough": {
1652+ "version": "2.2.0",
1653+ "resolved": "https://registry.npmjs.org/trough/-/trough-2.2.0.tgz",
1654+ "integrity": "sha512-tmMpK00BjZiUyVyvrBK7knerNgmgvcV/KLVyuma/SC+TQN167GrMRciANTz09+k3zW8L8t60jWO1GpfkZdjTaw==",
1655+ "license": "MIT",
1656+ "funding": {
1657+ "type": "github",
1658+ "url": "https://github.com/sponsors/wooorm"
1659+ }
1660+ },
1661+ "node_modules/unified": {
1662+ "version": "11.0.5",
1663+ "resolved": "https://registry.npmjs.org/unified/-/unified-11.0.5.tgz",
1664+ "integrity": "sha512-xKvGhPWw3k84Qjh8bI3ZeJjqnyadK+GEFtazSfZv/rKeTkTjOJho6mFqh2SM96iIcZokxiOpg78GazTSg8+KHA==",
1665+ "license": "MIT",
1666+ "dependencies": {
1667+ "@types/unist": "^3.0.0",
1668+ "bail": "^2.0.0",
1669+ "devlop": "^1.0.0",
1670+ "extend": "^3.0.0",
1671+ "is-plain-obj": "^4.0.0",
1672+ "trough": "^2.0.0",
1673+ "vfile": "^6.0.0"
1674+ },
1675+ "funding": {
1676+ "type": "opencollective",
1677+ "url": "https://opencollective.com/unified"
1678+ }
1679+ },
1680+ "node_modules/unist-util-is": {
1681+ "version": "6.0.1",
1682+ "resolved": "https://registry.npmjs.org/unist-util-is/-/unist-util-is-6.0.1.tgz",
1683+ "integrity": "sha512-LsiILbtBETkDz8I9p1dQ0uyRUWuaQzd/cuEeS1hoRSyW5E5XGmTzlwY1OrNzzakGowI9Dr/I8HVaw4hTtnxy8g==",
1684+ "license": "MIT",
1685+ "dependencies": {
1686+ "@types/unist": "^3.0.0"
1687+ },
1688+ "funding": {
1689+ "type": "opencollective",
1690+ "url": "https://opencollective.com/unified"
1691+ }
1692+ },
1693+ "node_modules/unist-util-position": {
1694+ "version": "5.0.0",
1695+ "resolved": "https://registry.npmjs.org/unist-util-position/-/unist-util-position-5.0.0.tgz",
1696+ "integrity": "sha512-fucsC7HjXvkB5R3kTCO7kUjRdrS0BJt3M/FPxmHMBOm8JQi2BsHAHFsy27E0EolP8rp0NzXsJ+jNPyDWvOJZPA==",
1697+ "license": "MIT",
1698+ "dependencies": {
1699+ "@types/unist": "^3.0.0"
1700+ },
1701+ "funding": {
1702+ "type": "opencollective",
1703+ "url": "https://opencollective.com/unified"
1704+ }
1705+ },
1706+ "node_modules/unist-util-stringify-position": {
1707+ "version": "4.0.0",
1708+ "resolved": "https://registry.npmjs.org/unist-util-stringify-position/-/unist-util-stringify-position-4.0.0.tgz",
1709+ "integrity": "sha512-0ASV06AAoKCDkS2+xw5RXJywruurpbC4JZSm7nr7MOt1ojAzvyyaO+UxZf18j8FCF6kmzCZKcAgN/yu2gm2XgQ==",
1710+ "license": "MIT",
1711+ "dependencies": {
1712+ "@types/unist": "^3.0.0"
1713+ },
1714+ "funding": {
1715+ "type": "opencollective",
1716+ "url": "https://opencollective.com/unified"
1717+ }
1718+ },
1719+ "node_modules/unist-util-visit": {
1720+ "version": "5.1.0",
1721+ "resolved": "https://registry.npmjs.org/unist-util-visit/-/unist-util-visit-5.1.0.tgz",
1722+ "integrity": "sha512-m+vIdyeCOpdr/QeQCu2EzxX/ohgS8KbnPDgFni4dQsfSCtpz8UqDyY5GjRru8PDKuYn7Fq19j1CQ+nJSsGKOzg==",
1723+ "license": "MIT",
1724+ "dependencies": {
1725+ "@types/unist": "^3.0.0",
1726+ "unist-util-is": "^6.0.0",
1727+ "unist-util-visit-parents": "^6.0.0"
1728+ },
1729+ "funding": {
1730+ "type": "opencollective",
1731+ "url": "https://opencollective.com/unified"
1732+ }
1733+ },
1734+ "node_modules/unist-util-visit-parents": {
1735+ "version": "6.0.2",
1736+ "resolved": "https://registry.npmjs.org/unist-util-visit-parents/-/unist-util-visit-parents-6.0.2.tgz",
1737+ "integrity": "sha512-goh1s1TBrqSqukSc8wrjwWhL0hiJxgA8m4kFxGlQ+8FYQ3C/m11FcTs4YYem7V664AhHVvgoQLk890Ssdsr2IQ==",
1738+ "license": "MIT",
1739+ "dependencies": {
1740+ "@types/unist": "^3.0.0",
1741+ "unist-util-is": "^6.0.0"
1742+ },
1743+ "funding": {
1744+ "type": "opencollective",
1745+ "url": "https://opencollective.com/unified"
1746+ }
1747+ },
1748+ "node_modules/vfile": {
1749+ "version": "6.0.3",
1750+ "resolved": "https://registry.npmjs.org/vfile/-/vfile-6.0.3.tgz",
1751+ "integrity": "sha512-KzIbH/9tXat2u30jf+smMwFCsno4wHVdNmzFyL+T/L3UGqqk6JKfVqOFOZEpZSHADH1k40ab6NUIXZq422ov3Q==",
1752+ "license": "MIT",
1753+ "dependencies": {
1754+ "@types/unist": "^3.0.0",
1755+ "vfile-message": "^4.0.0"
1756+ },
1757+ "funding": {
1758+ "type": "opencollective",
1759+ "url": "https://opencollective.com/unified"
1760+ }
1761+ },
1762+ "node_modules/vfile-message": {
1763+ "version": "4.0.3",
1764+ "resolved": "https://registry.npmjs.org/vfile-message/-/vfile-message-4.0.3.tgz",
1765+ "integrity": "sha512-QTHzsGd1EhbZs4AsQ20JX1rC3cOlt/IWJruk893DfLRr57lcnOeMaWG4K0JrRta4mIJZKth2Au3mM3u03/JWKw==",
1766+ "license": "MIT",
1767+ "dependencies": {
1768+ "@types/unist": "^3.0.0",
1769+ "unist-util-stringify-position": "^4.0.0"
1770+ },
1771+ "funding": {
1772+ "type": "opencollective",
1773+ "url": "https://opencollective.com/unified"
1774+ }
1775+ },
1776+ "node_modules/webgpu": {
1777+ "version": "0.4.0",
1778+ "resolved": "https://registry.npmjs.org/webgpu/-/webgpu-0.4.0.tgz",
1779+ "integrity": "sha512-F5pimn3Aoi0zWjuRdiVs5TnrUwSzD2lESBohsIUsqyitWkGRQlXU2fhV6ycXlQTa1bvAf3sjqiUpBEpmSQ5ptA==",
1780+ "hasInstallScript": true,
1781+ "license": "MIT",
1782+ "dependencies": {
1783+ "@webgpu/types": "^0.1.69",
1784+ "debug": "^4.4.0"
1785+ }
1786+ },
1787+ "node_modules/zwitch": {
1788+ "version": "2.0.4",
1789+ "resolved": "https://registry.npmjs.org/zwitch/-/zwitch-2.0.4.tgz",
1790+ "integrity": "sha512-bXE4cR/kVZhKZX/RjPEflHaKVhUVl85noU3v6b8apfQEc1x4A+zBxjZ4lN8LqGd6WZ3dl98pY4o717VFmoPp+A==",
1791+ "license": "MIT",
1792+ "funding": {
1793+ "type": "github",
1794+ "url": "https://github.com/sponsors/wooorm"
1795+ }
1796+ }
1797+ }
1798+}
package.jsonadded+15−0View file
@@ -0,0 +1,15 @@
1+{
2+ "name": "cymbal2-webgpu",
3+ "private": true,
4+ "version": "0.1.0",
5+ "type": "module",
6+ "scripts": {
7+ "render": "node src/run.mjs",
8+ "compare": "node src/run.mjs --compare",
9+ "bench": "node src/run.mjs --bench"
10+ },
11+ "dependencies": {
12+ "numbl": "^0.4.19",
13+ "webgpu": "^0.4.0"
14+ }
15+}
src/gpu.mjsadded+289−0View file
@@ -0,0 +1,289 @@
1+// WebGPU host for the megakernel: device setup (Dawn in Node via the
2+// optional `webgpu` package; in a browser this maps onto navigator.gpu),
3+// buffer upload, the dispatch queue, and readback.
4+
5+export async function installNodeWebGpu() {
6+ const mod = await import('webgpu');
7+ Object.assign(globalThis, mod.globals);
8+ const dawnFlags = process.env.DAWN_FLAGS?.split(';').filter(Boolean) ?? [];
9+ Object.defineProperty(globalThis, 'navigator', {
10+ value: { gpu: mod.create(dawnFlags) },
11+ configurable: true,
12+ writable: true,
13+ });
14+}
15+
16+export async function requestDevice() {
17+ const adapter = await navigator.gpu.requestAdapter();
18+ if (!adapter) throw new Error('no WebGPU adapter');
19+ const L = adapter.limits;
20+ const device = await adapter.requestDevice({
21+ requiredLimits: {
22+ maxComputeInvocationsPerWorkgroup: L.maxComputeInvocationsPerWorkgroup,
23+ maxComputeWorkgroupSizeX: L.maxComputeWorkgroupSizeX,
24+ maxComputeWorkgroupStorageSize: L.maxComputeWorkgroupStorageSize,
25+ },
26+ });
27+ device.addEventListener?.('uncapturederror', (e) => {
28+ console.error('WebGPU uncaptured error:', e.error?.message ?? e);
29+ });
30+ const info = adapter.info
31+ ? `${adapter.info.vendor ?? '?'} ${adapter.info.architecture ?? ''} ${adapter.info.description ?? ''}`.trim()
32+ : 'unknown adapter';
33+ return { device, info };
34+}
35+
36+// Multi-dispatch backend: 8 pipelines, one bind group with a dynamic-offset
37+// uniform carrying the within-chunk step index; the host writes the chunk
38+// base into state cell 0 before each chunk and keeps two chunks in flight.
39+export async function runMultiDispatch(device, L, stages, onProgress) {
40+ const nsteps = L.nsteps;
41+ const bgl = device.createBindGroupLayout({
42+ entries: [
43+ { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
44+ { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
45+ { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
46+ { binding: 3, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform', hasDynamicOffset: true, minBindingSize: 16 } },
47+ ],
48+ });
49+ const layout = device.createPipelineLayout({ bindGroupLayouts: [bgl] });
50+
51+ const pipelines = [];
52+ const simpleBGs = [];
53+ for (const name of stages.order) {
54+ const module = device.createShaderModule({ code: stages.sources[name] });
55+ const errs = (await module.getCompilationInfo()).messages.filter((m) => m.type === 'error');
56+ if (errs.length) {
57+ throw new Error(`stage ${name} WGSL errors:\n` +
58+ errs.map((m) => ` ${m.lineNum}:${m.linePos} ${m.message}`).join('\n'));
59+ }
60+ pipelines.push(device.createComputePipeline({ layout, compute: { module, entryPoint: 'main' } }));
61+ }
62+
63+ const chunk = Number(process.env.CHUNK ?? 1024);
64+ const mkBuf = (size, usage) => device.createBuffer({ size, usage });
65+ const opsBuf = mkBuf(L.opsData.byteLength, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
66+ device.queue.writeBuffer(opsBuf, 0, L.opsData);
67+ const stBuf = mkBuf(L.stLen * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST);
68+ const outBuf = mkBuf(2 * nsteps * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
69+ const uniBuf = mkBuf(chunk * 256, GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST);
70+ {
71+ const uni = new Uint32Array(chunk * 64);
72+ for (let s = 0; s < chunk; s++) uni[s * 64] = s;
73+ device.queue.writeBuffer(uniBuf, 0, uni);
74+ }
75+ const readOut = mkBuf(2 * nsteps * 4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
76+ const readSt = mkBuf(L.stLen * 4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
77+
78+ const bindGroup = device.createBindGroup({
79+ layout: bgl,
80+ entries: [
81+ { binding: 0, resource: { buffer: opsBuf } },
82+ { binding: 1, resource: { buffer: stBuf } },
83+ { binding: 2, resource: { buffer: outBuf } },
84+ { binding: 3, resource: { buffer: uniBuf, offset: 0, size: 16 } },
85+ ],
86+ });
87+
88+ const t0 = performance.now();
89+ const pending = [];
90+ let kbase = 0;
91+ let recMs = 0;
92+ const inflight = Number(process.env.INFLIGHT ?? 2);
93+ while (kbase < nsteps) {
94+ const n = Math.min(chunk, nsteps - kbase);
95+ device.queue.writeBuffer(stBuf, 0, new Float32Array([kbase]));
96+ const tr = performance.now();
97+ const enc = device.createCommandEncoder();
98+ const pass = enc.beginComputePass();
99+
100+ for (let s = 0; s < n; s++) {
101+ pass.setBindGroup(0, bindGroup, [s * 256]);
102+ for (const p of pipelines) {
103+ pass.setPipeline(p);
104+ pass.dispatchWorkgroups(stages.nwg);
105+ }
106+ }
107+ pass.end();
108+ const cb = enc.finish();
109+ recMs += performance.now() - tr;
110+ device.queue.submit([cb]);
111+ pending.push(device.queue.onSubmittedWorkDone());
112+ if (pending.length >= inflight) await pending.shift();
113+ kbase += n;
114+ if (onProgress) onProgress(kbase, nsteps, performance.now() - t0);
115+ }
116+ while (pending.length) await pending.shift();
117+ const gpuMs = performance.now() - t0;
118+ console.log(`\n [recording ${(1000 * recMs / nsteps).toFixed(1)} us/step of ${(1000 * gpuMs / nsteps).toFixed(1)} wall]`);
119+
120+ const enc = device.createCommandEncoder();
121+ enc.copyBufferToBuffer(outBuf, 0, readOut, 0, 2 * nsteps * 4);
122+ enc.copyBufferToBuffer(stBuf, 0, readSt, 0, L.stLen * 4);
123+ device.queue.submit([enc.finish()]);
124+ await readOut.mapAsync(GPUMapMode.READ);
125+ await readSt.mapAsync(GPUMapMode.READ);
126+ const outData = new Float32Array(readOut.getMappedRange().slice(0));
127+ const stData = new Float32Array(readSt.getMappedRange().slice(0));
128+ readOut.unmap();
129+ readSt.unmap();
130+ for (const b of [opsBuf, stBuf, outBuf, uniBuf, readOut, readSt]) b.destroy();
131+
132+ return {
133+ acc: outData.subarray(0, nsteps),
134+ fhist: outData.subarray(nsteps, 2 * nsteps),
135+ state: stData,
136+ gpuMs,
137+ };
138+}
139+
140+
141+// Persistent-kernel backend: identical dispatches, each integrating CHUNK
142+// steps with global spin barriers between stages. sync[] holds the barrier
143+// atomics; state cell 0 (kbase) advances GPU-side at the end of each chunk.
144+export async function runPersistent(device, L, pk, onProgress) {
145+ const nsteps = L.nsteps;
146+ const module = device.createShaderModule({ code: pk.code });
147+ const errs = (await module.getCompilationInfo()).messages.filter((m) => m.type === 'error');
148+ if (errs.length) {
149+ throw new Error('persistent kernel WGSL errors:\n' +
150+ errs.map((m) => ` ${m.lineNum}:${m.linePos} ${m.message}`).join('\n'));
151+ }
152+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
153+
154+ const mkBuf = (size, usage) => device.createBuffer({ size, usage });
155+ const opsBuf = mkBuf(L.opsData.byteLength, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
156+ device.queue.writeBuffer(opsBuf, 0, L.opsData);
157+ const stBuf = mkBuf(L.stLen * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
158+ const outBuf = mkBuf(2 * nsteps * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
159+ const syncBuf = mkBuf(16, GPUBufferUsage.STORAGE);
160+ const readOut = mkBuf(2 * nsteps * 4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
161+ const readSt = mkBuf(L.stLen * 4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
162+
163+ const bindGroup = device.createBindGroup({
164+ layout: pipeline.getBindGroupLayout(0),
165+ entries: [
166+ { binding: 0, resource: { buffer: opsBuf } },
167+ { binding: 1, resource: { buffer: stBuf } },
168+ { binding: 2, resource: { buffer: outBuf } },
169+ { binding: 3, resource: { buffer: syncBuf } },
170+ ],
171+ });
172+
173+ const totalDispatches = Math.ceil(nsteps / pk.chunk);
174+ const perSubmit = 8;
175+ const t0 = performance.now();
176+ let done = 0;
177+ const pending = [];
178+ while (done < totalDispatches) {
179+ const n = Math.min(perSubmit, totalDispatches - done);
180+ const enc = device.createCommandEncoder();
181+ for (let i = 0; i < n; i++) {
182+ const pass = enc.beginComputePass();
183+ pass.setPipeline(pipeline);
184+ pass.setBindGroup(0, bindGroup);
185+ pass.dispatchWorkgroups(pk.nwg);
186+ pass.end();
187+ }
188+ device.queue.submit([enc.finish()]);
189+ pending.push(device.queue.onSubmittedWorkDone());
190+ if (pending.length >= 2) await pending.shift();
191+ done += n;
192+ if (onProgress) onProgress(Math.min(done * pk.chunk, nsteps), nsteps, performance.now() - t0);
193+ }
194+ while (pending.length) await pending.shift();
195+ const gpuMs = performance.now() - t0;
196+
197+ const enc = device.createCommandEncoder();
198+ enc.copyBufferToBuffer(outBuf, 0, readOut, 0, 2 * nsteps * 4);
199+ enc.copyBufferToBuffer(stBuf, 0, readSt, 0, L.stLen * 4);
200+ device.queue.submit([enc.finish()]);
201+ await readOut.mapAsync(GPUMapMode.READ);
202+ await readSt.mapAsync(GPUMapMode.READ);
203+ const outData = new Float32Array(readOut.getMappedRange().slice(0));
204+ const stData = new Float32Array(readSt.getMappedRange().slice(0));
205+ readOut.unmap();
206+ readSt.unmap();
207+ for (const b of [opsBuf, stBuf, outBuf, syncBuf, readOut, readSt]) b.destroy();
208+
209+ return {
210+ acc: outData.subarray(0, nsteps),
211+ fhist: outData.subarray(nsteps, 2 * nsteps),
212+ state: stData,
213+ gpuMs,
214+ };
215+}
216+
217+export async function runMegakernel(device, kernel, nsteps, onProgress) {
218+ const module = device.createShaderModule({ code: kernel.code });
219+ const compInfo = await module.getCompilationInfo();
220+ const errors = compInfo.messages.filter((m) => m.type === 'error');
221+ if (errors.length) {
222+ throw new Error(
223+ 'WGSL compile errors:\n' +
224+ errors.map((m) => ` ${m.lineNum}:${m.linePos} ${m.message}`).join('\n'),
225+ );
226+ }
227+
228+ const pipeline = device.createComputePipeline({
229+ layout: 'auto',
230+ compute: { module, entryPoint: 'main' },
231+ });
232+
233+ const mkBuf = (size, usage) => device.createBuffer({ size, usage });
234+ const opsBuf = mkBuf(kernel.opsData.byteLength, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
235+ device.queue.writeBuffer(opsBuf, 0, kernel.opsData);
236+ const stBuf = mkBuf(kernel.stLen * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
237+ const outBuf = mkBuf(2 * nsteps * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
238+ const readOut = mkBuf(2 * nsteps * 4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
239+ const readSt = mkBuf(kernel.stLen * 4, GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST);
240+
241+ const bindGroup = device.createBindGroup({
242+ layout: pipeline.getBindGroupLayout(0),
243+ entries: [
244+ { binding: 0, resource: { buffer: opsBuf } },
245+ { binding: 1, resource: { buffer: stBuf } },
246+ { binding: 2, resource: { buffer: outBuf } },
247+ ],
248+ });
249+
250+ const totalDispatches = Math.ceil(nsteps / kernel.chunk);
251+ const perSubmit = 16;
252+ const t0 = performance.now();
253+ let done = 0;
254+ while (done < totalDispatches) {
255+ const n = Math.min(perSubmit, totalDispatches - done);
256+ const enc = device.createCommandEncoder();
257+ for (let i = 0; i < n; i++) {
258+ const pass = enc.beginComputePass();
259+ pass.setPipeline(pipeline);
260+ pass.setBindGroup(0, bindGroup);
261+ pass.dispatchWorkgroups(1);
262+ pass.end();
263+ }
264+ device.queue.submit([enc.finish()]);
265+ await device.queue.onSubmittedWorkDone();
266+ done += n;
267+ if (onProgress) onProgress(Math.min(done * kernel.chunk, nsteps), nsteps, performance.now() - t0);
268+ }
269+ const gpuMs = performance.now() - t0;
270+
271+ const enc = device.createCommandEncoder();
272+ enc.copyBufferToBuffer(outBuf, 0, readOut, 0, 2 * nsteps * 4);
273+ enc.copyBufferToBuffer(stBuf, 0, readSt, 0, kernel.stLen * 4);
274+ device.queue.submit([enc.finish()]);
275+ await readOut.mapAsync(GPUMapMode.READ);
276+ await readSt.mapAsync(GPUMapMode.READ);
277+ const outData = new Float32Array(readOut.getMappedRange().slice(0));
278+ const stData = new Float32Array(readSt.getMappedRange().slice(0));
279+ readOut.unmap();
280+ readSt.unmap();
281+ for (const b of [opsBuf, stBuf, outBuf, readOut, readSt]) b.destroy();
282+
283+ return {
284+ acc: outData.subarray(0, nsteps),
285+ fhist: outData.subarray(nsteps, 2 * nsteps),
286+ state: stData,
287+ gpuMs,
288+ };
289+}
src/heater.mjsadded+41−0View file
@@ -0,0 +1,41 @@
1+// Experiment: keep the GPU frequency governor boosted by running a small
2+// continuous compute load (one workgroup) from a separate device, to
3+// quantify how much of the multi-dispatch backend's time is the 300 MHz
4+// idle clock rather than the work.
5+import { installNodeWebGpu, requestDevice } from './gpu.mjs';
6+
7+async function main() {
8+ await installNodeWebGpu();
9+ const { device } = await requestDevice();
10+ const code = `
11+@group(0) @binding(0) var<storage, read_write> st: array<f32>;
12+@compute @workgroup_size(64)
13+fn main(@builtin(local_invocation_index) lid: u32) {
14+ var x: f32 = st[lid];
15+ for (var i: u32 = 0u; i < 15000u; i = i + 1u) {
16+ x = x * 1.0000001 + 1e-8;
17+ }
18+ st[lid] = x;
19+}`;
20+ const module = device.createShaderModule({ code });
21+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
22+ const buf = device.createBuffer({ size: 256, usage: GPUBufferUsage.STORAGE });
23+ const bg = device.createBindGroup({
24+ layout: pipeline.getBindGroupLayout(0),
25+ entries: [{ binding: 0, resource: { buffer: buf } }],
26+ });
27+ const secs = Number(process.argv[2] ?? 30);
28+ const t0 = performance.now();
29+ while (performance.now() - t0 < secs * 1000) {
30+ const enc = device.createCommandEncoder();
31+ const pass = enc.beginComputePass();
32+ pass.setPipeline(pipeline);
33+ pass.setBindGroup(0, bg);
34+ pass.dispatchWorkgroups(1);
35+ pass.end();
36+ device.queue.submit([enc.finish()]);
37+ await device.queue.onSubmittedWorkDone();
38+ }
39+}
40+
41+main().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); });
src/kernel.mjsadded+538−0View file
@@ -0,0 +1,538 @@
1+// WGSL megakernel generator for the cymbal2 time loop. One workgroup loops
2+// over CHUNK time steps inside a single dispatch, with barriers between the
3+// stages of each step; every dispatch is identical, so a whole render is a
4+// queue of identical dispatches with no CPU in the loop. All constants
5+// (grid sizes, node indices, physics scalars, the stroke schedule thresholds)
6+// are baked into the shader source; the operators live in one read-only
7+// storage buffer, the evolving state and scratch in a read-write one, the
8+// acceleration and force histories in a third.
9+//
10+// The stage structure mirrors cymbal2_sim.m line for line, with the sparse
11+// radial (D1, D2) and azimuthal (T1t, T2) difference matrices applied as
12+// stencils and the dense transforms (F, Gi), the per-wavenumber page solves
13+// (Fi, Ai, AiB, AiC) and the beam matrices applied as stored matrices.
14+
15+export const fmt = (x) => {
16+ if (!Number.isFinite(x)) throw new Error('non-finite kernel constant');
17+ let s = x.toPrecision(9);
18+ if (!/[.e]/.test(s)) s += '.0';
19+ return s + 'f';
20+};
21+
22+// Shared between the megakernel and the multi-dispatch backend: the packed
23+// operator buffer, the state-buffer layout, and the baked constants.
24+export function buildLayout(setup) {
25+ const nrt = setup.nrt | 0;
26+ const nth = setup.nth | 0;
27+ const nf = setup.nfree | 0;
28+ const N = nrt * nth;
29+ const K = setup.K | 0;
30+
31+ const o = setup.ops;
32+ const parts = [
33+ ['D1', o.D1, nrt * nrt],
34+ ['D2', o.D2, nrt * nrt],
35+ ['F', o.F, nth * nth],
36+ ['GI', o.Gi, nth * nth],
37+ ['RD', o.Rd, nrt],
38+ ['R2D', o.R2d, nrt],
39+ ['AIB', o.AiB, nrt * nrt * nth],
40+ ['AIC', o.AiC, nrt * nrt * nth],
41+ ['AI', o.Ai, nrt * nrt * nth],
42+ ['FI', o.Fi, nrt * nrt * nth],
43+ ['ASIB', o.ASiB, nf * nf],
44+ ['ASIC', o.ASiC, nf * nf],
45+ ['VM', o.vM, nf],
46+ ['VR', o.vR, nf],
47+ ['VK', o.vK, nf],
48+ ['TIPCOL', o.tipcol, nf],
49+ ];
50+ const ht = setup.ht;
51+ const kgo = setup.tlaunch.map((tl) => {
52+ let g = Math.max(0, Math.ceil(tl / ht) - 1);
53+ while (g > 0 && g * ht >= tl) g -= 1;
54+ while ((g + 1) * ht < tl) g += 1;
55+ return g;
56+ });
57+ parts.push(['KGO', Float32Array.from(kgo), K]);
58+ parts.push(['VSTR', Float32Array.from(setup.vstr), K]);
59+ // Transposed beam matrices: each beam row thread streams its own row
60+ // sequentially instead of striding across columns.
61+ const tr = (m) => {
62+ const t = new Float32Array(nf * nf);
63+ for (let r = 0; r < nf; r++) for (let p = 0; p < nf; p++) t[p + r * nf] = m[r + p * nf];
64+ return t;
65+ };
66+ parts.push(['ASIBT', tr(o.ASiB), nf * nf]);
67+ parts.push(['ASICT', tr(o.ASiC), nf * nf]);
68+
69+ const off = {};
70+ let cursor = 0;
71+ for (const [name, arr, len] of parts) {
72+ if (arr.length !== len) throw new Error(`ops part ${name}: ${arr.length} != ${len}`);
73+ off[name] = cursor;
74+ cursor += len;
75+ }
76+ const opsData = new Float32Array(cursor);
77+ for (const [name, arr] of parts) opsData.set(arr, off[name]);
78+
79+ const SC = 0;
80+ const ST_CVEL = 16;
81+ const ST_D = ST_CVEL + K;
82+ const ST_U = ST_D + 3 * nf;
83+ const ST_UH = ST_U + 3 * N;
84+ const scratchNames = [
85+ 'UR', 'URR', 'V1', 'UT1', 'V2', 'RHSA', 'RHSAF', 'PH3',
86+ 'PHI', 'PR', 'PRR', 'W1', 'PHT1', 'NL', 'NLF',
87+ ];
88+ const A = {};
89+ let scur = ST_UH + 3 * N;
90+ for (const n of scratchNames) {
91+ A[n] = scur;
92+ scur += N;
93+ }
94+ const stLen = scur;
95+
96+ const c = {
97+ HT: ht,
98+ HT2I: 1 / (ht * ht),
99+ GAP: setup.gap,
100+ KM: setup.KM,
101+ ALM: setup.alM,
102+ PS: setup.ps,
103+ AH: setup.aH,
104+ AHR: setup.aHr,
105+ MISSD: setup.missDepth,
106+ EHRS: setup.Eh / setup.Rs,
107+ EH2: setup.Eh / 2,
108+ RSI: 1 / setup.Rs,
109+ SNL: (ht * ht) / setup.rhoh,
110+ FSC: setup.fsc,
111+ C1R: 1 / (2 * setup.hr),
112+ C2R: 1 / (setup.hr * setup.hr),
113+ C1T: 1 / (2 * setup.hth),
114+ C2T: 1 / (setup.hth * setup.hth),
115+ };
116+
117+ return {
118+ nrt, nth, Nr: setup.Nr | 0, nf, N, K,
119+ nsteps: setup.nsteps | 0,
120+ opsData, off, stLen,
121+ layout: { SC, ST_CVEL, ST_D, ST_U, ST_UH },
122+ A, scratchNames, c,
123+ LS0: setup.ls - 1, JS0: setup.js - 1,
124+ LP0: setup.lp - 1, JP0: setup.jp - 1,
125+ TIPI0: setup.tipi - 1,
126+ NLR0: setup.nlr0 - 1,
127+ HOLDSTEPS: setup.holdSteps | 0,
128+ };
129+}
130+
131+export function buildKernel(setup, opts = {}) {
132+ const ws = opts.ws ?? 256;
133+ const chunk = opts.chunk ?? 2048;
134+
135+ const nrt = setup.nrt | 0;
136+ const nth = setup.nth | 0;
137+ const Nr = setup.Nr | 0;
138+ const nf = setup.nfree | 0;
139+ const N = nrt * nth;
140+ const K = setup.K | 0;
141+ const nsteps = setup.nsteps | 0;
142+
143+ // ---- ops buffer packing ----
144+ const o = setup.ops;
145+ const parts = [
146+ ['D1', o.D1, nrt * nrt],
147+ ['D2', o.D2, nrt * nrt],
148+ ['F', o.F, nth * nth],
149+ ['GI', o.Gi, nth * nth],
150+ ['RD', o.Rd, nrt],
151+ ['R2D', o.R2d, nrt],
152+ ['AIB', o.AiB, nrt * nrt * nth],
153+ ['AIC', o.AiC, nrt * nrt * nth],
154+ ['AI', o.Ai, nrt * nrt * nth],
155+ ['FI', o.Fi, nrt * nrt * nth],
156+ ['ASIB', o.ASiB, nf * nf],
157+ ['ASIC', o.ASiC, nf * nf],
158+ ['VM', o.vM, nf],
159+ ['VR', o.vR, nf],
160+ ['VK', o.vK, nf],
161+ ['TIPCOL', o.tipcol, nf],
162+ ];
163+ // Launch-step thresholds: smallest 0-based kk with (kk+1)*ht >= tlaunch,
164+ // decided in f64 exactly as MATLAB's k*ht >= tlaunch decides it.
165+ const ht = setup.ht;
166+ const kgo = setup.tlaunch.map((tl) => {
167+ let g = Math.max(0, Math.ceil(tl / ht) - 1);
168+ while (g > 0 && g * ht >= tl) g -= 1;
169+ while ((g + 1) * ht < tl) g += 1;
170+ return g;
171+ });
172+ parts.push(['KGO', Float32Array.from(kgo), K]);
173+ parts.push(['VSTR', Float32Array.from(setup.vstr), K]);
174+
175+ const off = {};
176+ let cursor = 0;
177+ for (const [name, arr, len] of parts) {
178+ if (arr.length !== len) throw new Error(`ops part ${name}: ${arr.length} != ${len}`);
179+ off[name] = cursor;
180+ cursor += len;
181+ }
182+ const opsData = new Float32Array(cursor);
183+ for (const [name, arr] of parts) opsData.set(arr, off[name]);
184+
185+ // ---- state buffer layout ----
186+ const SC = 0; // 16 scalars
187+ const ST_CVEL = 16;
188+ const ST_D = ST_CVEL + K;
189+ const ST_U = ST_D + 3 * nf;
190+ const ST_UH = ST_U + 3 * N;
191+ const scratchNames = [
192+ 'UR', 'URR', 'V1', 'UT1', 'V2', 'RHSA', 'RHSAF', 'PH3',
193+ 'PHI', 'PR', 'PRR', 'W1', 'PHT1', 'NL', 'NLF',
194+ ];
195+ const A = {};
196+ let scur = ST_UH + 3 * N;
197+ for (const n of scratchNames) {
198+ A[n] = scur;
199+ scur += N;
200+ }
201+ const stLen = scur;
202+
203+ // ---- baked constants ----
204+ const LS0 = setup.ls - 1, JS0 = setup.js - 1;
205+ const LP0 = setup.lp - 1, JP0 = setup.jp - 1;
206+ const TIPI0 = setup.tipi - 1;
207+ const NLR0 = setup.nlr0 - 1; // 0-based first kept NL row
208+ const c = {
209+ HT: ht,
210+ HT2I: 1 / (ht * ht),
211+ GAP: setup.gap,
212+ KM: setup.KM,
213+ ALM: setup.alM,
214+ PS: setup.ps,
215+ AH: setup.aH,
216+ AHR: setup.aHr,
217+ MISSD: setup.missDepth,
218+ EHRS: setup.Eh / setup.Rs,
219+ EH2: setup.Eh / 2,
220+ RSI: 1 / setup.Rs,
221+ SNL: (ht * ht) / setup.rhoh,
222+ FSC: setup.fsc,
223+ C1R: 1 / (2 * setup.hr),
224+ C2R: 1 / (setup.hr * setup.hr),
225+ C1T: 1 / (2 * setup.hth),
226+ C2T: 1 / (setup.hth * setup.hth),
227+ };
228+ const constLines = Object.entries(c)
229+ .map(([k, v]) => `const ${k}: f32 = ${fmt(v)};`)
230+ .join('\n');
231+
232+ const code = /* wgsl */ `
233+const NRT = ${nrt}u;
234+const NTH = ${nth}u;
235+const N = ${N}u;
236+const NF = ${nf}u;
237+const WS = ${ws}u;
238+const CHUNK = ${chunk}u;
239+const NSTEPS = ${nsteps}u;
240+const KS = ${K}u;
241+const NRU = ${Nr}u;
242+const LS0 = ${LS0}u; const JS0 = ${JS0}u;
243+const LP0 = ${LP0}u; const JP0 = ${JP0}u;
244+const TIPI = ${TIPI0}u;
245+const NLR0 = ${NLR0}u;
246+const HOLDSTEPS = ${setup.holdSteps | 0};
247+${constLines}
248+${Object.entries(off).map(([k, v]) => `const OP_${k} = ${v}u;`).join('\n')}
249+const SC = ${SC}u;
250+const ST_CVEL = ${ST_CVEL}u;
251+const ST_D = ${ST_D}u;
252+const ST_U = ${ST_U}u;
253+const ST_UH = ${ST_UH}u;
254+${scratchNames.map((n) => `const A_${n} = ${A[n]}u;`).join('\n')}
255+
256+@group(0) @binding(0) var<storage, read> ops: array<f32>;
257+@group(0) @binding(1) var<storage, read_write> st: array<f32>;
258+@group(0) @binding(2) var<storage, read_write> outb: array<f32>;
259+
260+var<workgroup> wgKbase: u32;
261+var<workgroup> wgF: f32;
262+var<workgroup> wgHA: f32;
263+var<workgroup> wgHB: f32;
264+var<workgroup> wgHC: f32;
265+var<workgroup> wgShell: u32;
266+var<workgroup> wgRot: u32;
267+
268+fn thm(j: u32) -> u32 { return (j + NTH - 1u) % NTH; }
269+fn thp(j: u32) -> u32 { return (j + 1u) % NTH; }
270+
271+@compute @workgroup_size(${ws})
272+fn main(@builtin(local_invocation_index) lid: u32) {
273+ var shellA: u32 = 0u; var inC: u32 = 0u; var phase: u32 = 0u;
274+ var strokeIdx: u32 = 0u; var holdLeft: i32 = 0; var vtar: f32 = 0.0;
275+ var vH: f32 = 0.0; var uHm: f32 = 0.0; var uH0: f32 = 0.0; var uHp: f32 = 0.0;
276+ var peakXi: f32 = 0.0; var peakDisp: f32 = 0.0;
277+ if (lid == 0u) {
278+ wgKbase = u32(st[SC + 0u]);
279+ shellA = u32(st[SC + 1u]); inC = u32(st[SC + 2u]); phase = u32(st[SC + 3u]);
280+ strokeIdx = u32(st[SC + 4u]); holdLeft = i32(st[SC + 5u]); vtar = st[SC + 6u];
281+ vH = st[SC + 7u]; uHm = st[SC + 8u]; uH0 = st[SC + 9u]; uHp = st[SC + 10u];
282+ peakXi = st[SC + 12u]; peakDisp = st[SC + 13u];
283+ wgRot = u32(st[SC + 11u]);
284+ wgShell = shellA;
285+ }
286+ let kbase = workgroupUniformLoad(&wgKbase);
287+
288+ for (var s: u32 = 0u; s < CHUNK; s = s + 1u) {
289+ let k = kbase + s;
290+ if (k >= NSTEPS) { break; }
291+ let dNew = ST_D + (k % 3u) * NF;
292+ let dCur = ST_D + ((k + 2u) % 3u) * NF;
293+ let dOld = ST_D + ((k + 1u) % 3u) * NF;
294+
295+ // ---- stage 1: contact force + hand state machine (lane 0) ----
296+ if (lid == 0u) {
297+ let xi = st[ST_U + wgRot * N + LS0 + JS0 * NRT] - st[dCur + TIPI] - GAP;
298+ var fM: f32 = 0.0;
299+ if (xi > 0.0) {
300+ fM = KM * pow(xi, ALM);
301+ shellA = 1u;
302+ if (inC == 0u && strokeIdx < KS && st[ST_CVEL + strokeIdx] == 0.0) {
303+ st[ST_CVEL + strokeIdx] = (st[dCur + TIPI] - st[dOld + TIPI]) / HT;
304+ }
305+ inC = 1u;
306+ if (xi > peakXi) { peakXi = xi; }
307+ } else {
308+ inC = 0u;
309+ }
310+ outb[NSTEPS + k] = fM;
311+ wgF = fM;
312+ wgHA = -(uHp - 2.0 * uH0 + uHm);
313+ wgHB = -(HT * 0.5) * (uHp - uHm);
314+ wgHC = -HT * HT * (PS * uH0 + ((1.0 - PS) * 0.5) * (uHp + uHm));
315+ if (phase == 0u) {
316+ if (strokeIdx < KS && f32(k) >= ops[OP_KGO + strokeIdx]) {
317+ phase = 1u; vtar = ops[OP_VSTR + strokeIdx];
318+ }
319+ } else if (phase == 1u) {
320+ vH = max(vH - AH * HT, -vtar);
321+ if (inC != 0u) { phase = 3u; holdLeft = HOLDSTEPS; }
322+ else if (vH <= -vtar) { phase = 2u; }
323+ } else if (phase == 2u) {
324+ if (inC != 0u) { phase = 3u; holdLeft = HOLDSTEPS; }
325+ else if (uHp <= -MISSD) { phase = 4u; }
326+ } else if (phase == 3u) {
327+ holdLeft = holdLeft - 1;
328+ if (holdLeft <= 0) { phase = 4u; }
329+ } else if (phase == 4u) {
330+ vH = min(vH + AHR * HT, 0.5 * vtar);
331+ if (vH > 0.0 && uHp >= -vH * vH / (2.0 * AHR)) { phase = 5u; }
332+ } else {
333+ vH = vH - AHR * HT;
334+ if (vH <= 0.0 || uHp >= 0.0) { vH = 0.0; phase = 0u; strokeIdx = strokeIdx + 1u; }
335+ }
336+ uHm = uH0; uH0 = uHp; uHp = uHp + vH * HT;
337+ wgShell = shellA;
338+ }
339+ workgroupBarrier();
340+ storageBarrier();
341+ let fM1 = wgF;
342+ let hA = wgHA; let hB = wgHB; let hC = wgHC;
343+
344+ // ---- stage 2: beam update ----
345+ for (var r: u32 = lid; r < NF; r = r + WS) {
346+ var sum: f32 = ops[OP_VM + r] * hA + ops[OP_VR + r] * hB + ops[OP_VK + r] * hC
347+ + (HT * HT * fM1) * ops[OP_TIPCOL + r];
348+ for (var p: u32 = 0u; p < NF; p = p + 1u) {
349+ sum = sum + ops[OP_ASIB + r + p * NF] * st[dCur + p]
350+ + ops[OP_ASIC + r + p * NF] * st[dOld + p];
351+ }
352+ st[dNew + r] = sum;
353+ }
354+
355+ let sh = workgroupUniformLoad(&wgShell);
356+ if (sh != 0u) {
357+ let rot = workgroupUniformLoad(&wgRot);
358+ let uS = ST_U + rot * N;
359+ let u1S = ST_U + ((rot + 2u) % 3u) * N;
360+ let unS = ST_U + ((rot + 1u) % 3u) * N;
361+ let h3S = ST_UH + rot * N;
362+ let h13S = ST_UH + ((rot + 2u) % 3u) * N;
363+ let h3nS = ST_UH + ((rot + 1u) % 3u) * N;
364+
365+ // ---- A: radial/azimuthal derivatives of U ----
366+ for (var i: u32 = lid; i < N; i = i + WS) {
367+ let l = i % NRT; let j = i / NRT;
368+ var ur: f32 = 0.0; var urr: f32 = 0.0;
369+ if (l >= 1u && l <= NRT - 2u) {
370+ let um = st[uS + (l - 1u) + j * NRT];
371+ let u0 = st[uS + l + j * NRT];
372+ let up = st[uS + (l + 1u) + j * NRT];
373+ ur = (up - um) * C1R;
374+ urr = (um - 2.0 * u0 + up) * C2R;
375+ }
376+ st[A_UR + i] = ur;
377+ st[A_URR + i] = urr;
378+ let ujm = st[uS + l + thm(j) * NRT];
379+ let uj0 = st[uS + l + j * NRT];
380+ let ujp = st[uS + l + thp(j) * NRT];
381+ st[A_V1 + i] = ops[OP_RD + l] * ur + ops[OP_R2D + l] * ((ujm - 2.0 * uj0 + ujp) * C2T);
382+ st[A_UT1 + i] = (ujp - ujm) * C1T;
383+ }
384+ storageBarrier();
385+
386+ // ---- B: bracket pieces and Airy right-hand side ----
387+ for (var i: u32 = lid; i < N; i = i + WS) {
388+ let l = i % NRT; let j = i / NRT;
389+ let urt = (st[A_UR + l + thp(j) * NRT] - st[A_UR + l + thm(j) * NRT]) * C1T;
390+ let v2 = ops[OP_RD + l] * urt - ops[OP_R2D + l] * st[A_UT1 + i];
391+ st[A_V2 + i] = v2;
392+ var rv: f32 = 0.0;
393+ if (l >= 3u && l <= NRU - 2u) {
394+ let urr = st[A_URR + i]; let v1 = st[A_V1 + i];
395+ rv = EHRS * (urr + v1) - EH2 * 2.0 * (urr * v1 - v2 * v2);
396+ }
397+ st[A_RHSA + i] = rv;
398+ }
399+ storageBarrier();
400+
401+ // ---- C: transform the Airy rhs ----
402+ for (var i: u32 = lid; i < N; i = i + WS) {
403+ let l = i % NRT; let j = i / NRT;
404+ var sum: f32 = 0.0;
405+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
406+ sum = sum + st[A_RHSA + l + q * NRT] * ops[OP_F + q + j * NTH];
407+ }
408+ st[A_RHSAF + i] = sum;
409+ }
410+ storageBarrier();
411+
412+ // ---- D: Airy solve per wavenumber column ----
413+ for (var i: u32 = lid; i < N; i = i + WS) {
414+ let l = i % NRT; let j = i / NRT;
415+ let base = OP_FI + j * NRT * NRT;
416+ var sum: f32 = 0.0;
417+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
418+ sum = sum + ops[base + l + p * NRT] * st[A_RHSAF + p + j * NRT];
419+ }
420+ st[A_PH3 + i] = sum;
421+ }
422+ storageBarrier();
423+
424+ // ---- E: stress function back to physical space ----
425+ for (var i: u32 = lid; i < N; i = i + WS) {
426+ let l = i % NRT; let j = i / NRT;
427+ var sum: f32 = 0.0;
428+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
429+ sum = sum + st[A_PH3 + l + q * NRT] * ops[OP_GI + q + j * NTH];
430+ }
431+ st[A_PHI + i] = sum;
432+ }
433+ storageBarrier();
434+
435+ // ---- F: derivatives of the stress function ----
436+ for (var i: u32 = lid; i < N; i = i + WS) {
437+ let l = i % NRT; let j = i / NRT;
438+ var pr: f32 = 0.0; var prr: f32 = 0.0;
439+ if (l >= 1u && l <= NRT - 2u) {
440+ let pm = st[A_PHI + (l - 1u) + j * NRT];
441+ let p0 = st[A_PHI + l + j * NRT];
442+ let pp = st[A_PHI + (l + 1u) + j * NRT];
443+ pr = (pp - pm) * C1R;
444+ prr = (pm - 2.0 * p0 + pp) * C2R;
445+ }
446+ st[A_PR + i] = pr;
447+ st[A_PRR + i] = prr;
448+ let pjm = st[A_PHI + l + thm(j) * NRT];
449+ let pj0 = st[A_PHI + l + j * NRT];
450+ let pjp = st[A_PHI + l + thp(j) * NRT];
451+ st[A_W1 + i] = ops[OP_RD + l] * pr + ops[OP_R2D + l] * ((pjm - 2.0 * pj0 + pjp) * C2T);
452+ st[A_PHT1 + i] = (pjp - pjm) * C1T;
453+ }
454+ storageBarrier();
455+
456+ // ---- G: von Karman bracket, forcing, membrane term ----
457+ for (var i: u32 = lid; i < N; i = i + WS) {
458+ let l = i % NRT; let j = i / NRT;
459+ var nl: f32 = 0.0;
460+ if (l >= NLR0 && l < NRU) {
461+ let prt = (st[A_PR + l + thp(j) * NRT] - st[A_PR + l + thm(j) * NRT]) * C1T;
462+ let w2 = ops[OP_RD + l] * prt - ops[OP_R2D + l] * st[A_PHT1 + i];
463+ let prr = st[A_PRR + i]; let w1 = st[A_W1 + i];
464+ nl = prr * st[A_V1 + i] + st[A_URR + i] * w1 - 2.0 * w2 * st[A_V2 + i]
465+ - (prr + w1) * RSI;
466+ if (l == LS0 && j == JS0) { nl = nl - fM1 * FSC; }
467+ nl = nl * SNL;
468+ }
469+ st[A_NL + i] = nl;
470+ }
471+ storageBarrier();
472+
473+ // ---- H: transform the nonlinear load ----
474+ for (var i: u32 = lid; i < N; i = i + WS) {
475+ let l = i % NRT; let j = i / NRT;
476+ var sum: f32 = 0.0;
477+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
478+ sum = sum + st[A_NL + l + q * NRT] * ops[OP_F + q + j * NTH];
479+ }
480+ st[A_NLF + i] = sum;
481+ }
482+ storageBarrier();
483+
484+ // ---- I: implicit shell update per wavenumber column ----
485+ for (var i: u32 = lid; i < N; i = i + WS) {
486+ let l = i % NRT; let j = i / NRT;
487+ let base = j * NRT * NRT;
488+ var sum: f32 = 0.0;
489+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
490+ let pj = p + j * NRT;
491+ sum = sum + ops[OP_AIB + base + l + p * NRT] * st[h3S + pj]
492+ + ops[OP_AIC + base + l + p * NRT] * st[h13S + pj]
493+ + ops[OP_AI + base + l + p * NRT] * st[A_NLF + pj];
494+ }
495+ st[h3nS + i] = sum;
496+ }
497+ storageBarrier();
498+
499+ // ---- J: back to physical space; pickup acceleration ----
500+ for (var i: u32 = lid; i < N; i = i + WS) {
501+ let l = i % NRT; let j = i / NRT;
502+ var sum: f32 = 0.0;
503+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
504+ sum = sum + st[h3nS + l + q * NRT] * ops[OP_GI + q + j * NTH];
505+ }
506+ st[unS + i] = sum;
507+ if (l == LP0 && j == JP0) {
508+ outb[k] = (sum - 2.0 * st[uS + i] + st[u1S + i]) * HT2I;
509+ }
510+ }
511+ storageBarrier();
512+ if (lid == 0u) {
513+ wgRot = (rot + 1u) % 3u;
514+ if ((k + 1u) % 512u == 0u) {
515+ var pd: f32 = 0.0;
516+ for (var l2: u32 = 0u; l2 < NRT; l2 = l2 + 1u) {
517+ pd = max(pd, abs(st[unS + l2 + JS0 * NRT]));
518+ }
519+ if (pd > peakDisp) { peakDisp = pd; }
520+ }
521+ }
522+ }
523+ workgroupBarrier();
524+ storageBarrier();
525+ }
526+
527+ if (lid == 0u) {
528+ st[SC + 0u] = f32(min(kbase + CHUNK, NSTEPS));
529+ st[SC + 1u] = f32(shellA); st[SC + 2u] = f32(inC); st[SC + 3u] = f32(phase);
530+ st[SC + 4u] = f32(strokeIdx); st[SC + 5u] = f32(holdLeft); st[SC + 6u] = vtar;
531+ st[SC + 7u] = vH; st[SC + 8u] = uHm; st[SC + 9u] = uH0; st[SC + 10u] = uHp;
532+ st[SC + 11u] = f32(wgRot); st[SC + 12u] = peakXi; st[SC + 13u] = peakDisp;
533+ }
534+}
535+`;
536+
537+ return { code, opsData, stLen, N, chunk, ws, off, layout: { SC, ST_CVEL, ST_D, ST_U, ST_UH } };
538+}
src/kernels2.mjsadded+306−0View file
@@ -0,0 +1,306 @@
1+// Multi-dispatch backend: each stage of a time step is its own compute
2+// dispatch at full GPU width, and one step is 8 dispatches:
3+// AB hand state machine + contact + beam (workgroup 0), and the shell
4+// derivative stencils / Airy right-hand side (all workgroups)
5+// C RHSAF = RHSA * F (azimuthal analysis)
6+// D PH3 = Fi_j RHSAF_j (Airy solve per wavenumber)
7+// E PHI = PH3 * Gi (azimuthal synthesis)
8+// FG NL from PHI (stress-function stencils recomputed inline), forcing
9+// H NLF = NL * F
10+// I UH3N = AiB UH3 + AiC UH13 + Ai NLF (implicit update per column)
11+// J UNEW = UH3N * Gi, pickup acceleration
12+//
13+// The step index k is kbase (state cell 0, written by the host once per
14+// chunk) plus s from a dynamic-offset uniform, so every recorded dispatch is
15+// identical up to that offset and slot rotation is a pure function of k.
16+// U/Uh slots rotate every step whether or not the shell is active: before
17+// first contact all slots are zero, so the rotation is harmless, and after
18+// first contact every step writes. The shell-active flag latches inside AB;
19+// stages of the same step read it a dispatch later, so they always see the
20+// latch, and the benign race inside AB itself can only cause zero-work on a
21+// zero field.
22+
23+import { fmt } from './kernel.mjs';
24+
25+export function buildStages(L, opts = {}) {
26+ const wg = opts.wg ?? Number(process.env.WG ?? 64);
27+ const { nrt, nth, N, nf, K, nsteps, off, A, c } = L;
28+ const { SC, ST_CVEL, ST_D, ST_U, ST_UH } = L.layout;
29+
30+ const constLines = Object.entries(c)
31+ .map(([k, v]) => `const ${k}: f32 = ${fmt(v)};`)
32+ .join('\n');
33+
34+ const header = `
35+const NRT = ${nrt}u;
36+const NTH = ${nth}u;
37+const N = ${N}u;
38+const NF = ${nf}u;
39+const NSTEPS = ${nsteps}u;
40+const KS = ${K}u;
41+const NRU = ${L.Nr}u;
42+const LS0 = ${L.LS0}u; const JS0 = ${L.JS0}u;
43+const LP0 = ${L.LP0}u; const JP0 = ${L.JP0}u;
44+const TIPI = ${L.TIPI0}u;
45+const NLR0 = ${L.NLR0}u;
46+const HOLDSTEPS = ${L.HOLDSTEPS};
47+${constLines}
48+${Object.entries(off).map(([k, v]) => `const OP_${k} = ${v}u;`).join('\n')}
49+const ST_CVEL = ${ST_CVEL}u;
50+const ST_D = ${ST_D}u;
51+const ST_U = ${ST_U}u;
52+const ST_UH = ${ST_UH}u;
53+${L.scratchNames.map((n) => `const A_${n} = ${A[n]}u;`).join('\n')}
54+
55+@group(0) @binding(0) var<storage, read> ops: array<f32>;
56+@group(0) @binding(1) var<storage, read_write> st: array<f32>;
57+@group(0) @binding(2) var<storage, read_write> outb: array<f32>;
58+@group(0) @binding(3) var<uniform> ctl: vec4<u32>;
59+
60+fn thm(j: u32) -> u32 { return (j + NTH - 1u) % NTH; }
61+fn thp(j: u32) -> u32 { return (j + 1u) % NTH; }
62+`;
63+
64+
65+ const stages = {};
66+
67+ stages.AB = `${header}
68+var<workgroup> wgF: f32;
69+var<workgroup> wgHA: f32;
70+var<workgroup> wgHB: f32;
71+var<workgroup> wgHC: f32;
72+var<workgroup> shD: array<f32, ${nf}>;
73+var<workgroup> shD1: array<f32, ${nf}>;
74+
75+fn u_dr(b: u32, l: u32, j: u32) -> f32 {
76+ if (l < 1u || l > NRT - 2u) { return 0.0; }
77+ return (st[b + (l + 1u) + j * NRT] - st[b + (l - 1u) + j * NRT]) * C1R;
78+}
79+
80+@compute @workgroup_size(${wg})
81+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
82+ @builtin(local_invocation_index) lid: u32,
83+ @builtin(workgroup_id) wgid: vec3<u32>) {
84+ // No early-out on k here: the host records exactly nsteps steps, and the
85+ // workgroupBarrier below requires uniform control flow up to this point.
86+ let k = u32(st[0u]) + ctl.x;
87+ let uS = ST_U + (k % 3u) * N;
88+ if (wgid.x == 0u) {
89+ if (lid == 0u) {
90+ var shellA = u32(st[1u]); var inC = u32(st[2u]); var phase = u32(st[3u]);
91+ var strokeIdx = u32(st[4u]); var holdLeft = i32(st[5u]); var vtar = st[6u];
92+ var vH = st[7u]; var uHm = st[8u]; var uH0 = st[9u]; var uHp = st[10u];
93+ var peakXi = st[12u]; var peakDisp = st[13u];
94+ let dCur = ST_D + ((k + 2u) % 3u) * NF;
95+ let dOld = ST_D + ((k + 1u) % 3u) * NF;
96+ if (shellA != 0u && k > 0u && (k % 512u) == 0u) {
97+ var pd: f32 = 0.0;
98+ for (var l: u32 = 0u; l < NRT; l = l + 1u) {
99+ pd = max(pd, abs(st[uS + l + JS0 * NRT]));
100+ }
101+ if (pd > peakDisp) { peakDisp = pd; }
102+ }
103+ let xi = st[uS + LS0 + JS0 * NRT] - st[dCur + TIPI] - GAP;
104+ var fM: f32 = 0.0;
105+ if (xi > 0.0) {
106+ fM = KM * pow(xi, ALM);
107+ shellA = 1u;
108+ if (inC == 0u && strokeIdx < KS && st[ST_CVEL + strokeIdx] == 0.0) {
109+ st[ST_CVEL + strokeIdx] = (st[dCur + TIPI] - st[dOld + TIPI]) / HT;
110+ }
111+ inC = 1u;
112+ if (xi > peakXi) { peakXi = xi; }
113+ } else {
114+ inC = 0u;
115+ }
116+ outb[NSTEPS + k] = fM;
117+ wgF = fM;
118+ wgHA = -(uHp - 2.0 * uH0 + uHm);
119+ wgHB = -(HT * 0.5) * (uHp - uHm);
120+ wgHC = -HT * HT * (PS * uH0 + ((1.0 - PS) * 0.5) * (uHp + uHm));
121+ if (phase == 0u) {
122+ if (strokeIdx < KS && f32(k) >= ops[OP_KGO + strokeIdx]) {
123+ phase = 1u; vtar = ops[OP_VSTR + strokeIdx];
124+ }
125+ } else if (phase == 1u) {
126+ vH = max(vH - AH * HT, -vtar);
127+ if (inC != 0u) { phase = 3u; holdLeft = HOLDSTEPS; }
128+ else if (vH <= -vtar) { phase = 2u; }
129+ } else if (phase == 2u) {
130+ if (inC != 0u) { phase = 3u; holdLeft = HOLDSTEPS; }
131+ else if (uHp <= -MISSD) { phase = 4u; }
132+ } else if (phase == 3u) {
133+ holdLeft = holdLeft - 1;
134+ if (holdLeft <= 0) { phase = 4u; }
135+ } else if (phase == 4u) {
136+ vH = min(vH + AHR * HT, 0.5 * vtar);
137+ if (vH > 0.0 && uHp >= -vH * vH / (2.0 * AHR)) { phase = 5u; }
138+ } else {
139+ vH = vH - AHR * HT;
140+ if (vH <= 0.0 || uHp >= 0.0) { vH = 0.0; phase = 0u; strokeIdx = strokeIdx + 1u; }
141+ }
142+ uHm = uH0; uH0 = uHp; uHp = uHp + vH * HT;
143+ st[1u] = f32(shellA); st[2u] = f32(inC); st[3u] = f32(phase);
144+ st[4u] = f32(strokeIdx); st[5u] = f32(holdLeft); st[6u] = vtar;
145+ st[7u] = vH; st[8u] = uHm; st[9u] = uH0; st[10u] = uHp;
146+ st[12u] = peakXi; st[13u] = peakDisp; st[14u] = fM;
147+ }
148+ if (lid < NF) {
149+ shD[lid] = st[ST_D + ((k + 2u) % 3u) * NF + lid];
150+ shD1[lid] = st[ST_D + ((k + 1u) % 3u) * NF + lid];
151+ }
152+ workgroupBarrier();
153+ if (lid < NF) {
154+ let r = lid;
155+ var sum: f32 = ops[OP_VM + r] * wgHA + ops[OP_VR + r] * wgHB + ops[OP_VK + r] * wgHC
156+ + (HT * HT * wgF) * ops[OP_TIPCOL + r];
157+ for (var p: u32 = 0u; p < NF; p = p + 1u) {
158+ sum = sum + ops[OP_ASIBT + p + r * NF] * shD[p]
159+ + ops[OP_ASICT + p + r * NF] * shD1[p];
160+ }
161+ st[ST_D + (k % 3u) * NF + r] = sum;
162+ }
163+ }
164+ let i = gid.x;
165+ if (i >= N) { return; }
166+ if (st[1u] == 0.0) { return; }
167+ let l = i % NRT; let j = i / NRT;
168+ var ur: f32 = 0.0; var urr: f32 = 0.0;
169+ if (l >= 1u && l <= NRT - 2u) {
170+ let um = st[uS + (l - 1u) + j * NRT];
171+ let u0 = st[uS + l + j * NRT];
172+ let up = st[uS + (l + 1u) + j * NRT];
173+ ur = (up - um) * C1R;
174+ urr = (um - 2.0 * u0 + up) * C2R;
175+ }
176+ let ujm = st[uS + l + thm(j) * NRT];
177+ let uj0 = st[uS + l + j * NRT];
178+ let ujp = st[uS + l + thp(j) * NRT];
179+ let v1 = ops[OP_RD + l] * ur + ops[OP_R2D + l] * ((ujm - 2.0 * uj0 + ujp) * C2T);
180+ let urt = (u_dr(uS, l, thp(j)) - u_dr(uS, l, thm(j))) * C1T;
181+ let v2 = ops[OP_RD + l] * urt - ops[OP_R2D + l] * ((ujp - ujm) * C1T);
182+ st[A_URR + i] = urr;
183+ st[A_V1 + i] = v1;
184+ st[A_V2 + i] = v2;
185+ var rv: f32 = 0.0;
186+ if (l >= 3u && l <= NRU - 2u) {
187+ rv = EHRS * (urr + v1) - EH2 * 2.0 * (urr * v1 - v2 * v2);
188+ }
189+ st[A_RHSA + i] = rv;
190+}
191+`;
192+
193+ const gate = `
194+ let k = u32(st[0u]) + ctl.x;
195+ if (k >= NSTEPS) { return; }
196+ let i = gid.x;
197+ if (i >= N) { return; }
198+ if (st[1u] == 0.0) { return; }
199+ let l = i % NRT; let j = i / NRT;`;
200+
201+ const transform = (src, dst, mat) => `${header}
202+@compute @workgroup_size(${wg})
203+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
204+${gate}
205+ var sum: f32 = 0.0;
206+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
207+ sum = sum + st[${src} + l + q * NRT] * ops[${mat} + q + j * NTH];
208+ }
209+ st[${dst} + i] = sum;
210+}
211+`;
212+
213+ stages.C = transform('A_RHSA', 'A_RHSAF', 'OP_F');
214+
215+ stages.D = `${header}
216+@compute @workgroup_size(${wg})
217+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
218+${gate}
219+ let base = OP_FI + j * NRT * NRT;
220+ var sum: f32 = 0.0;
221+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
222+ sum = sum + ops[base + l + p * NRT] * st[A_RHSAF + p + j * NRT];
223+ }
224+ st[A_PH3 + i] = sum;
225+}
226+`;
227+
228+ stages.E = transform('A_PH3', 'A_PHI', 'OP_GI');
229+
230+ stages.FG = `${header}
231+fn p_dr(l: u32, j: u32) -> f32 {
232+ if (l < 1u || l > NRT - 2u) { return 0.0; }
233+ return (st[A_PHI + (l + 1u) + j * NRT] - st[A_PHI + (l - 1u) + j * NRT]) * C1R;
234+}
235+@compute @workgroup_size(${wg})
236+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
237+${gate}
238+ var nl: f32 = 0.0;
239+ if (l >= NLR0 && l < NRU) {
240+ let pm = st[A_PHI + (l - 1u) + j * NRT];
241+ let p0 = st[A_PHI + l + j * NRT];
242+ let pp = st[A_PHI + (l + 1u) + j * NRT];
243+ let pr = (pp - pm) * C1R;
244+ let prr = (pm - 2.0 * p0 + pp) * C2R;
245+ let pjm = st[A_PHI + l + thm(j) * NRT];
246+ let pjp = st[A_PHI + l + thp(j) * NRT];
247+ let w1 = ops[OP_RD + l] * pr + ops[OP_R2D + l] * ((pjm - 2.0 * p0 + pjp) * C2T);
248+ let prt = (p_dr(l, thp(j)) - p_dr(l, thm(j))) * C1T;
249+ let w2 = ops[OP_RD + l] * prt - ops[OP_R2D + l] * ((pjp - pjm) * C1T);
250+ nl = prr * st[A_V1 + i] + st[A_URR + i] * w1 - 2.0 * w2 * st[A_V2 + i]
251+ - (prr + w1) * RSI;
252+ if (l == LS0 && j == JS0) { nl = nl - st[14u] * FSC; }
253+ nl = nl * SNL;
254+ }
255+ st[A_NL + i] = nl;
256+}
257+`;
258+
259+ stages.H = transform('A_NL', 'A_NLF', 'OP_F');
260+
261+ stages.I = `${header}
262+@compute @workgroup_size(${wg})
263+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
264+${gate}
265+ let h3S = ST_UH + (k % 3u) * N;
266+ let h13S = ST_UH + ((k + 2u) % 3u) * N;
267+ let h3nS = ST_UH + ((k + 1u) % 3u) * N;
268+ let base = j * NRT * NRT;
269+ var sum: f32 = 0.0;
270+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
271+ let pj = p + j * NRT;
272+ sum = sum + ops[OP_AIB + base + l + p * NRT] * st[h3S + pj]
273+ + ops[OP_AIC + base + l + p * NRT] * st[h13S + pj]
274+ + ops[OP_AI + base + l + p * NRT] * st[A_NLF + pj];
275+ }
276+ st[h3nS + i] = sum;
277+}
278+`;
279+
280+ stages.J = `${header}
281+@compute @workgroup_size(${wg})
282+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
283+${gate}
284+ let h3nS = ST_UH + ((k + 1u) % 3u) * N;
285+ let uS = ST_U + (k % 3u) * N;
286+ let u1S = ST_U + ((k + 2u) % 3u) * N;
287+ let unS = ST_U + ((k + 1u) % 3u) * N;
288+ var sum: f32 = 0.0;
289+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
290+ sum = sum + st[h3nS + l + q * NRT] * ops[OP_GI + q + j * NTH];
291+ }
292+ st[unS + i] = sum;
293+ if (l == LP0 && j == JP0) {
294+ outb[k] = (sum - 2.0 * st[uS + i] + st[u1S + i]) * HT2I;
295+ }
296+}
297+`;
298+
299+ const nwgMain = Math.ceil(N / wg);
300+ return {
301+ order: ['AB', 'C', 'D', 'E', 'FG', 'H', 'I', 'J'],
302+ sources: stages,
303+ wg,
304+ nwg: nwgMain,
305+ };
306+}
src/kernels3.mjsadded+329−0View file
@@ -0,0 +1,329 @@
1+// Persistent-kernel backend: ONE dispatch of NWG workgroups runs CHUNK time
2+// steps, with a global spin barrier on storage atomics between stages in
3+// place of dispatch boundaries. This removes the per-dispatch launch bubble
4+// that keeps an idle-clocked GPU at its minimum frequency, and it uses the
5+// whole GPU, unlike the one-workgroup megakernel.
6+//
7+// WebGPU does not guarantee cross-workgroup forward progress, so this mode
8+// assumes all NWG workgroups are co-resident (13 workgroups of 128 lanes is
9+// far below any real device's residency) and is validated against the
10+// multi-dispatch backend, which remains the portable fallback. No thread
11+// ever returns early: every guard is non-returning so all threads reach
12+// every barrier.
13+
14+import { fmt } from './kernel.mjs';
15+
16+export function buildPersistent(L, opts = {}) {
17+ const wg = opts.wg ?? 128;
18+ const chunk = opts.chunk ?? 512;
19+ const { nrt, nth, N, nf, K, nsteps, off, A, c } = L;
20+ const { ST_CVEL, ST_D, ST_U, ST_UH } = L.layout;
21+ const nwg = Math.ceil(N / wg);
22+
23+ const constLines = Object.entries(c)
24+ .map(([k, v]) => `const ${k}: f32 = ${fmt(v)};`)
25+ .join('\n');
26+
27+ const code = /* wgsl */ `
28+const NRT = ${nrt}u;
29+const NTH = ${nth}u;
30+const N = ${N}u;
31+const NF = ${nf}u;
32+const NSTEPS = ${nsteps}u;
33+const KS = ${K}u;
34+const NRU = ${L.Nr}u;
35+const NWG = ${nwg}u;
36+const CHUNK = ${chunk}u;
37+const LS0 = ${L.LS0}u; const JS0 = ${L.JS0}u;
38+const LP0 = ${L.LP0}u; const JP0 = ${L.JP0}u;
39+const TIPI = ${L.TIPI0}u;
40+const NLR0 = ${L.NLR0}u;
41+const HOLDSTEPS = ${L.HOLDSTEPS};
42+${constLines}
43+${Object.entries(off).map(([k, v]) => `const OP_${k} = ${v}u;`).join('\n')}
44+const ST_CVEL = ${ST_CVEL}u;
45+const ST_D = ${ST_D}u;
46+const ST_U = ${ST_U}u;
47+const ST_UH = ${ST_UH}u;
48+${L.scratchNames.map((n) => `const A_${n} = ${A[n]}u;`).join('\n')}
49+
50+@group(0) @binding(0) var<storage, read> ops: array<f32>;
51+@group(0) @binding(1) var<storage, read_write> st: array<f32>;
52+@group(0) @binding(2) var<storage, read_write> outb: array<f32>;
53+@group(0) @binding(3) var<storage, read_write> sync: array<atomic<u32>, 2>;
54+
55+var<workgroup> wgF: f32;
56+var<workgroup> wgHA: f32;
57+var<workgroup> wgHB: f32;
58+var<workgroup> wgHC: f32;
59+var<workgroup> shD: array<f32, ${nf}>;
60+var<workgroup> shD1: array<f32, ${nf}>;
61+
62+fn thm(j: u32) -> u32 { return (j + NTH - 1u) % NTH; }
63+fn thp(j: u32) -> u32 { return (j + 1u) % NTH; }
64+
65+var<private> P_lid: u32;
66+
67+// Classic two-phase global barrier: sync[0] arrive counter, sync[1] epoch.
68+fn globalBarrier() {
69+ workgroupBarrier();
70+ storageBarrier();
71+ if (P_lid == 0u) {
72+ let gen = atomicLoad(&sync[1]);
73+ let n = atomicAdd(&sync[0], 1u);
74+ if (n == NWG - 1u) {
75+ atomicStore(&sync[0], 0u);
76+ atomicAdd(&sync[1], 1u);
77+ } else {
78+ loop {
79+ if (atomicLoad(&sync[1]) != gen) { break; }
80+ }
81+ }
82+ }
83+ workgroupBarrier();
84+ storageBarrier();
85+}
86+
87+fn u_dr(b: u32, l: u32, j: u32) -> f32 {
88+ if (l < 1u || l > NRT - 2u) { return 0.0; }
89+ return (st[b + (l + 1u) + j * NRT] - st[b + (l - 1u) + j * NRT]) * C1R;
90+}
91+
92+fn p_dr(l: u32, j: u32) -> f32 {
93+ if (l < 1u || l > NRT - 2u) { return 0.0; }
94+ return (st[A_PHI + (l + 1u) + j * NRT] - st[A_PHI + (l - 1u) + j * NRT]) * C1R;
95+}
96+
97+@compute @workgroup_size(${wg})
98+fn main(@builtin(global_invocation_id) gid: vec3<u32>,
99+ @builtin(local_invocation_index) lid: u32,
100+ @builtin(workgroup_id) wgid: vec3<u32>) {
101+ P_lid = lid;
102+ let i = gid.x;
103+ let l = i % NRT;
104+ let j = i / NRT;
105+ let onGrid = i < N;
106+ let kbase = u32(st[0u]);
107+
108+ for (var s: u32 = 0u; s < CHUNK; s = s + 1u) {
109+ let k = kbase + s;
110+ let live = k < NSTEPS;
111+ let uS = ST_U + (k % 3u) * N;
112+
113+ // ---- AB: hand + contact (wg0 lane0), beam (wg0), stencils (all) ----
114+ if (wgid.x == 0u) {
115+ if (lid == 0u && live) {
116+ var shellA = u32(st[1u]); var inC = u32(st[2u]); var phase = u32(st[3u]);
117+ var strokeIdx = u32(st[4u]); var holdLeft = i32(st[5u]); var vtar = st[6u];
118+ var vH = st[7u]; var uHm = st[8u]; var uH0 = st[9u]; var uHp = st[10u];
119+ var peakXi = st[12u]; var peakDisp = st[13u];
120+ let dCur = ST_D + ((k + 2u) % 3u) * NF;
121+ let dOld = ST_D + ((k + 1u) % 3u) * NF;
122+ if (shellA != 0u && k > 0u && (k % 512u) == 0u) {
123+ var pd: f32 = 0.0;
124+ for (var l2: u32 = 0u; l2 < NRT; l2 = l2 + 1u) {
125+ pd = max(pd, abs(st[uS + l2 + JS0 * NRT]));
126+ }
127+ if (pd > peakDisp) { peakDisp = pd; }
128+ }
129+ let xi = st[uS + LS0 + JS0 * NRT] - st[dCur + TIPI] - GAP;
130+ var fM: f32 = 0.0;
131+ if (xi > 0.0) {
132+ fM = KM * pow(xi, ALM);
133+ shellA = 1u;
134+ if (inC == 0u && strokeIdx < KS && st[ST_CVEL + strokeIdx] == 0.0) {
135+ st[ST_CVEL + strokeIdx] = (st[dCur + TIPI] - st[dOld + TIPI]) / HT;
136+ }
137+ inC = 1u;
138+ if (xi > peakXi) { peakXi = xi; }
139+ } else {
140+ inC = 0u;
141+ }
142+ outb[NSTEPS + k] = fM;
143+ wgF = fM;
144+ wgHA = -(uHp - 2.0 * uH0 + uHm);
145+ wgHB = -(HT * 0.5) * (uHp - uHm);
146+ wgHC = -HT * HT * (PS * uH0 + ((1.0 - PS) * 0.5) * (uHp + uHm));
147+ if (phase == 0u) {
148+ if (strokeIdx < KS && f32(k) >= ops[OP_KGO + strokeIdx]) {
149+ phase = 1u; vtar = ops[OP_VSTR + strokeIdx];
150+ }
151+ } else if (phase == 1u) {
152+ vH = max(vH - AH * HT, -vtar);
153+ if (inC != 0u) { phase = 3u; holdLeft = HOLDSTEPS; }
154+ else if (vH <= -vtar) { phase = 2u; }
155+ } else if (phase == 2u) {
156+ if (inC != 0u) { phase = 3u; holdLeft = HOLDSTEPS; }
157+ else if (uHp <= -MISSD) { phase = 4u; }
158+ } else if (phase == 3u) {
159+ holdLeft = holdLeft - 1;
160+ if (holdLeft <= 0) { phase = 4u; }
161+ } else if (phase == 4u) {
162+ vH = min(vH + AHR * HT, 0.5 * vtar);
163+ if (vH > 0.0 && uHp >= -vH * vH / (2.0 * AHR)) { phase = 5u; }
164+ } else {
165+ vH = vH - AHR * HT;
166+ if (vH <= 0.0 || uHp >= 0.0) { vH = 0.0; phase = 0u; strokeIdx = strokeIdx + 1u; }
167+ }
168+ uHm = uH0; uH0 = uHp; uHp = uHp + vH * HT;
169+ st[1u] = f32(shellA); st[2u] = f32(inC); st[3u] = f32(phase);
170+ st[4u] = f32(strokeIdx); st[5u] = f32(holdLeft); st[6u] = vtar;
171+ st[7u] = vH; st[8u] = uHm; st[9u] = uH0; st[10u] = uHp;
172+ st[12u] = peakXi; st[13u] = peakDisp; st[14u] = fM;
173+ }
174+ if (lid < NF && live) {
175+ shD[lid] = st[ST_D + ((k + 2u) % 3u) * NF + lid];
176+ shD1[lid] = st[ST_D + ((k + 1u) % 3u) * NF + lid];
177+ }
178+ workgroupBarrier();
179+ if (lid < NF && live) {
180+ let r = lid;
181+ var sum: f32 = ops[OP_VM + r] * wgHA + ops[OP_VR + r] * wgHB + ops[OP_VK + r] * wgHC
182+ + (HT * HT * wgF) * ops[OP_TIPCOL + r];
183+ for (var p: u32 = 0u; p < NF; p = p + 1u) {
184+ sum = sum + ops[OP_ASIBT + p + r * NF] * shD[p]
185+ + ops[OP_ASICT + p + r * NF] * shD1[p];
186+ }
187+ st[ST_D + (k % 3u) * NF + r] = sum;
188+ }
189+ } else {
190+ workgroupBarrier(); // match wg0's barrier count
191+ }
192+ // The shell-active latch for THIS step is decided by wg0 above; other
193+ // workgroups may see the pre-latch zero here, which only skips work on a
194+ // still-zero field, exactly as in the multi-dispatch backend.
195+ var act = onGrid && live && st[1u] != 0.0;
196+ if (act) {
197+ var ur: f32 = 0.0; var urr: f32 = 0.0;
198+ if (l >= 1u && l <= NRT - 2u) {
199+ let um = st[uS + (l - 1u) + j * NRT];
200+ let u0 = st[uS + l + j * NRT];
201+ let up = st[uS + (l + 1u) + j * NRT];
202+ ur = (up - um) * C1R;
203+ urr = (um - 2.0 * u0 + up) * C2R;
204+ }
205+ let ujm = st[uS + l + thm(j) * NRT];
206+ let uj0 = st[uS + l + j * NRT];
207+ let ujp = st[uS + l + thp(j) * NRT];
208+ let v1 = ops[OP_RD + l] * ur + ops[OP_R2D + l] * ((ujm - 2.0 * uj0 + ujp) * C2T);
209+ let urt = (u_dr(uS, l, thp(j)) - u_dr(uS, l, thm(j))) * C1T;
210+ let v2 = ops[OP_RD + l] * urt - ops[OP_R2D + l] * ((ujp - ujm) * C1T);
211+ st[A_URR + i] = urr;
212+ st[A_V1 + i] = v1;
213+ st[A_V2 + i] = v2;
214+ var rv: f32 = 0.0;
215+ if (l >= 3u && l <= NRU - 2u) {
216+ rv = EHRS * (urr + v1) - EH2 * 2.0 * (urr * v1 - v2 * v2);
217+ }
218+ st[A_RHSA + i] = rv;
219+ }
220+ globalBarrier();
221+ // After the barrier every workgroup sees the latch for this step.
222+ act = onGrid && live && st[1u] != 0.0;
223+
224+ // ---- C: RHSAF = RHSA * F ----
225+ if (act) {
226+ var sum: f32 = 0.0;
227+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
228+ sum = sum + st[A_RHSA + l + q * NRT] * ops[OP_F + q + j * NTH];
229+ }
230+ st[A_RHSAF + i] = sum;
231+ }
232+ globalBarrier();
233+
234+ // ---- D: Airy solve per wavenumber ----
235+ if (act) {
236+ let base = OP_FI + j * NRT * NRT;
237+ var sum: f32 = 0.0;
238+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
239+ sum = sum + ops[base + l + p * NRT] * st[A_RHSAF + p + j * NRT];
240+ }
241+ st[A_PH3 + i] = sum;
242+ }
243+ globalBarrier();
244+
245+ // ---- E: PHI = PH3 * Gi ----
246+ if (act) {
247+ var sum: f32 = 0.0;
248+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
249+ sum = sum + st[A_PH3 + l + q * NRT] * ops[OP_GI + q + j * NTH];
250+ }
251+ st[A_PHI + i] = sum;
252+ }
253+ globalBarrier();
254+
255+ // ---- FG: nonlinear load ----
256+ if (act) {
257+ var nl: f32 = 0.0;
258+ if (l >= NLR0 && l < NRU) {
259+ let pm = st[A_PHI + (l - 1u) + j * NRT];
260+ let p0 = st[A_PHI + l + j * NRT];
261+ let pp = st[A_PHI + (l + 1u) + j * NRT];
262+ let pr = (pp - pm) * C1R;
263+ let prr = (pm - 2.0 * p0 + pp) * C2R;
264+ let pjm = st[A_PHI + l + thm(j) * NRT];
265+ let pjp = st[A_PHI + l + thp(j) * NRT];
266+ let w1 = ops[OP_RD + l] * pr + ops[OP_R2D + l] * ((pjm - 2.0 * p0 + pjp) * C2T);
267+ let prt = (p_dr(l, thp(j)) - p_dr(l, thm(j))) * C1T;
268+ let w2 = ops[OP_RD + l] * prt - ops[OP_R2D + l] * ((pjp - pjm) * C1T);
269+ nl = prr * st[A_V1 + i] + st[A_URR + i] * w1 - 2.0 * w2 * st[A_V2 + i]
270+ - (prr + w1) * RSI;
271+ if (l == LS0 && j == JS0) { nl = nl - st[14u] * FSC; }
272+ nl = nl * SNL;
273+ }
274+ st[A_NL + i] = nl;
275+ }
276+ globalBarrier();
277+
278+ // ---- H: NLF = NL * F ----
279+ if (act) {
280+ var sum: f32 = 0.0;
281+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
282+ sum = sum + st[A_NL + l + q * NRT] * ops[OP_F + q + j * NTH];
283+ }
284+ st[A_NLF + i] = sum;
285+ }
286+ globalBarrier();
287+
288+ // ---- I: implicit update per wavenumber ----
289+ if (act) {
290+ let h3S = ST_UH + (k % 3u) * N;
291+ let h13S = ST_UH + ((k + 2u) % 3u) * N;
292+ let h3nS = ST_UH + ((k + 1u) % 3u) * N;
293+ let base = j * NRT * NRT;
294+ var sum: f32 = 0.0;
295+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
296+ let pj = p + j * NRT;
297+ sum = sum + ops[OP_AIB + base + l + p * NRT] * st[h3S + pj]
298+ + ops[OP_AIC + base + l + p * NRT] * st[h13S + pj]
299+ + ops[OP_AI + base + l + p * NRT] * st[A_NLF + pj];
300+ }
301+ st[h3nS + i] = sum;
302+ }
303+ globalBarrier();
304+
305+ // ---- J: synthesis, pickup acceleration ----
306+ if (act) {
307+ let h3nS = ST_UH + ((k + 1u) % 3u) * N;
308+ let u1S = ST_U + ((k + 2u) % 3u) * N;
309+ let unS = ST_U + ((k + 1u) % 3u) * N;
310+ var sum: f32 = 0.0;
311+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
312+ sum = sum + st[h3nS + l + q * NRT] * ops[OP_GI + q + j * NTH];
313+ }
314+ st[unS + i] = sum;
315+ if (l == LP0 && j == JP0) {
316+ outb[k] = (sum - 2.0 * st[uS + i] + st[u1S + i]) * HT2I;
317+ }
318+ }
319+ globalBarrier();
320+ }
321+
322+ if (wgid.x == 0u && lid == 0u) {
323+ st[0u] = f32(min(kbase + CHUNK, NSTEPS));
324+ }
325+}
326+`;
327+
328+ return { code, chunk, wg, nwg };
329+}
src/probe.mjsadded+77−0View file
@@ -0,0 +1,77 @@
1+// Microbenchmark: what do workgroup/storage barriers and the stage loops cost
2+// on this device, per simulated step, for one workgroup?
3+import { installNodeWebGpu, requestDevice } from './gpu.mjs';
4+
5+const N = 1560;
6+const WS = Number(process.env.WS ?? 256);
7+const STEPS = 2048;
8+
9+function kernel(body, nbar) {
10+ return `
11+const N = ${N}u; const WS = ${WS}u; const STEPS = ${STEPS}u;
12+@group(0) @binding(0) var<storage, read_write> st: array<f32>;
13+@compute @workgroup_size(${WS})
14+fn main(@builtin(local_invocation_index) lid: u32) {
15+ for (var s: u32 = 0u; s < STEPS; s = s + 1u) {
16+ ${Array.from({ length: nbar }, () => body).join('\n')}
17+ }
18+ if (lid == 0u) { st[0] = f32(STEPS); }
19+}
20+`;
21+}
22+
23+async function time(device, code, label) {
24+ const module = device.createShaderModule({ code });
25+ const info = await module.getCompilationInfo();
26+ const errs = info.messages.filter((m) => m.type === 'error');
27+ if (errs.length) throw new Error(errs.map((m) => m.message).join('\n'));
28+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
29+ const buf = device.createBuffer({ size: 4 * N * 20, usage: GPUBufferUsage.STORAGE });
30+ const bg = device.createBindGroup({
31+ layout: pipeline.getBindGroupLayout(0),
32+ entries: [{ binding: 0, resource: { buffer: buf } }],
33+ });
34+ // warmup
35+ let enc = device.createCommandEncoder();
36+ let pass = enc.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bg); pass.dispatchWorkgroups(1); pass.end();
37+ device.queue.submit([enc.finish()]);
38+ await device.queue.onSubmittedWorkDone();
39+ const t0 = performance.now();
40+ const reps = 4;
41+ enc = device.createCommandEncoder();
42+ for (let r = 0; r < reps; r++) {
43+ pass = enc.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bg); pass.dispatchWorkgroups(1); pass.end();
44+ }
45+ device.queue.submit([enc.finish()]);
46+ await device.queue.onSubmittedWorkDone();
47+ const ms = (performance.now() - t0) / reps;
48+ console.log(`${label}: ${(1000 * ms / STEPS).toFixed(2)} us/step`);
49+ buf.destroy();
50+}
51+
52+async function main() {
53+ await installNodeWebGpu();
54+ const { device, info } = await requestDevice();
55+ console.log(`device: ${info}, WS=${WS}`);
56+
57+ await time(device, kernel('', 0), 'empty loop');
58+ await time(device, kernel('workgroupBarrier();', 14), '14 workgroupBarrier');
59+ await time(device, kernel('storageBarrier();', 14), '14 storageBarrier');
60+ await time(device, kernel(`
61+ for (var i: u32 = lid; i < N; i = i + WS) {
62+ st[N + i] = st[i] * 1.0001 + 0.5;
63+ }
64+ storageBarrier();`, 10), '10 x (elementwise N + storageBarrier)');
65+ await time(device, kernel(`
66+ for (var i: u32 = lid; i < N; i = i + WS) {
67+ var sum: f32 = 0.0;
68+ let l = i % 26u; let j = i / 26u;
69+ for (var q: u32 = 0u; q < 60u; q = q + 1u) {
70+ sum = sum + st[l + q * 26u] * st[2u * N + q + j * 60u];
71+ }
72+ st[N + i] = sum;
73+ }
74+ storageBarrier();`, 4), '4 x (60-MAC matmul N + storageBarrier)');
75+}
76+
77+main().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); });
src/probe2.mjsadded+246−0View file
@@ -0,0 +1,246 @@
1+// Which matmul formulation is fast on this device, for X (26x60) * M (60x60),
2+// one workgroup? Candidates: per-output serial reduction (baseline), shared-
3+// memory staged input, column-per-thread with register accumulators, and the
4+// combination. Also the 26x26 page solve (pagemtimes) both ways.
5+import { installNodeWebGpu, requestDevice } from './gpu.mjs';
6+
7+const NRT = 26, NTH = 60, N = NRT * NTH;
8+const WS = Number(process.env.WS ?? 256);
9+const STEPS = 2048;
10+
11+function wrap(body, decl = '') {
12+ return `
13+const NRT = ${NRT}u; const NTH = ${NTH}u; const N = ${N}u; const WS = ${WS}u; const STEPS = ${STEPS}u;
14+@group(0) @binding(0) var<storage, read> ops: array<f32>;
15+@group(0) @binding(1) var<storage, read_write> st: array<f32>;
16+${decl}
17+@compute @workgroup_size(${WS})
18+fn main(@builtin(local_invocation_index) lid: u32) {
19+ for (var s: u32 = 0u; s < STEPS; s = s + 1u) {
20+${body}
21+ workgroupBarrier();
22+ storageBarrier();
23+ }
24+ if (lid == 0u) { st[0] = st[1] + 1.0; }
25+}
26+`;
27+}
28+
29+const CASES = {
30+ baseline: wrap(`
31+ for (var i: u32 = lid; i < N; i = i + WS) {
32+ let l = i % NRT; let j = i / NRT;
33+ var sum: f32 = 0.0;
34+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
35+ sum = sum + st[N + l + q * NRT] * ops[q + j * NTH];
36+ }
37+ st[2u * N + i] = sum;
38+ }`),
39+
40+ sharedX: wrap(`
41+ for (var i: u32 = lid; i < N; i = i + WS) { xsh[i] = st[N + i]; }
42+ workgroupBarrier();
43+ for (var i: u32 = lid; i < N; i = i + WS) {
44+ let l = i % NRT; let j = i / NRT;
45+ var sum: f32 = 0.0;
46+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
47+ sum = sum + xsh[l + q * NRT] * ops[q + j * NTH];
48+ }
49+ st[2u * N + i] = sum;
50+ }`, 'var<workgroup> xsh: array<f32, 1560>;'),
51+
52+ colRegs: wrap(`
53+ if (lid < NTH) {
54+ let j = lid;
55+ var acc: array<f32, ${NRT}>;
56+ for (var l: u32 = 0u; l < NRT; l = l + 1u) { acc[l] = 0.0; }
57+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
58+ let mq = ops[q + j * NTH];
59+ for (var l: u32 = 0u; l < NRT; l = l + 1u) {
60+ acc[l] = acc[l] + st[N + l + q * NRT] * mq;
61+ }
62+ }
63+ for (var l: u32 = 0u; l < NRT; l = l + 1u) { st[2u * N + l + j * NRT] = acc[l]; }
64+ }`),
65+
66+ colRegsShared: wrap(`
67+ for (var i: u32 = lid; i < N; i = i + WS) { xsh[i] = st[N + i]; }
68+ workgroupBarrier();
69+ if (lid < NTH) {
70+ let j = lid;
71+ var acc: array<f32, ${NRT}>;
72+ for (var l: u32 = 0u; l < NRT; l = l + 1u) { acc[l] = 0.0; }
73+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
74+ let mq = ops[q + j * NTH];
75+ for (var l: u32 = 0u; l < NRT; l = l + 1u) {
76+ acc[l] = acc[l] + xsh[l + q * NRT] * mq;
77+ }
78+ }
79+ for (var l: u32 = 0u; l < NRT; l = l + 1u) { st[2u * N + l + j * NRT] = acc[l]; }
80+ }`, 'var<workgroup> xsh: array<f32, 1560>;'),
81+
82+ colBlockShared: wrap(`
83+ for (var i: u32 = lid; i < N; i = i + WS) { xsh[i] = st[N + i]; }
84+ workgroupBarrier();
85+ if (lid < 240u) {
86+ let j = lid % NTH; let b = lid / NTH; // 4 blocks of 7 rows (last short)
87+ let l0 = b * 7u;
88+ let nl = min(7u, NRT - l0);
89+ var acc: array<f32, 7>;
90+ for (var l: u32 = 0u; l < 7u; l = l + 1u) { acc[l] = 0.0; }
91+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
92+ let mq = ops[q + j * NTH];
93+ for (var l: u32 = 0u; l < nl; l = l + 1u) {
94+ acc[l] = acc[l] + xsh[l0 + l + q * NRT] * mq;
95+ }
96+ }
97+ for (var l: u32 = 0u; l < nl; l = l + 1u) { st[2u * N + l0 + l + j * NRT] = acc[l]; }
98+ }`, 'var<workgroup> xsh: array<f32, 1560>;'),
99+
100+ pageBaseline: wrap(`
101+ for (var i: u32 = lid; i < N; i = i + WS) {
102+ let l = i % NRT; let j = i / NRT;
103+ let base = j * NRT * NRT;
104+ var sum: f32 = 0.0;
105+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
106+ sum = sum + ops[base + l + p * NRT] * st[N + p + j * NRT];
107+ }
108+ st[2u * N + i] = sum;
109+ }`),
110+
111+ pageColRegsShared: wrap(`
112+ for (var i: u32 = lid; i < N; i = i + WS) { xsh[i] = st[N + i]; }
113+ workgroupBarrier();
114+ if (lid < 240u) {
115+ let j = lid % NTH; let b = lid / NTH;
116+ let l0 = b * 7u;
117+ let nl = min(7u, NRT - l0);
118+ let base = j * NRT * NRT;
119+ var acc: array<f32, 7>;
120+ for (var l: u32 = 0u; l < 7u; l = l + 1u) { acc[l] = 0.0; }
121+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
122+ let xq = xsh[p + j * NRT];
123+ for (var l: u32 = 0u; l < nl; l = l + 1u) {
124+ acc[l] = acc[l] + ops[base + l0 + l + p * NRT] * xq;
125+ }
126+ }
127+ for (var l: u32 = 0u; l < nl; l = l + 1u) { st[2u * N + l0 + l + j * NRT] = acc[l]; }
128+ }`, 'var<workgroup> xsh: array<f32, 1560>;'),
129+};
130+
131+async function time(device, code, label) {
132+ const module = device.createShaderModule({ code });
133+ const info = await module.getCompilationInfo();
134+ const errs = info.messages.filter((m) => m.type === 'error');
135+ if (errs.length) throw new Error(label + ':\n' + errs.map((m) => `${m.lineNum}: ${m.message}`).join('\n'));
136+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
137+ const ops = device.createBuffer({ size: 4 * NRT * NRT * NTH, usage: GPUBufferUsage.STORAGE });
138+ const stb = device.createBuffer({ size: 4 * N * 4, usage: GPUBufferUsage.STORAGE });
139+ const bg = device.createBindGroup({
140+ layout: pipeline.getBindGroupLayout(0),
141+ entries: [
142+ { binding: 0, resource: { buffer: ops } },
143+ { binding: 1, resource: { buffer: stb } },
144+ ],
145+ });
146+ const run = async (reps) => {
147+ const enc = device.createCommandEncoder();
148+ for (let r = 0; r < reps; r++) {
149+ const pass = enc.beginComputePass();
150+ pass.setPipeline(pipeline); pass.setBindGroup(0, bg); pass.dispatchWorkgroups(1); pass.end();
151+ }
152+ device.queue.submit([enc.finish()]);
153+ await device.queue.onSubmittedWorkDone();
154+ };
155+ await run(1);
156+ const t0 = performance.now();
157+ await run(4);
158+ const ms = (performance.now() - t0) / 4;
159+ console.log(`${label}: ${(1000 * ms / STEPS).toFixed(2)} us per stage-step`);
160+ ops.destroy(); stb.destroy();
161+}
162+
163+async function main() {
164+ await installNodeWebGpu();
165+ const { device, info } = await requestDevice();
166+ console.log(`device: ${info}, WS=${WS}`);
167+ for (const [label, code] of Object.entries(CASES)) await time(device, code, label);
168+}
169+
170+if (!process.env.MORE) { main().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); }); }
171+
172+// --- appended variants ---
173+export const MORE = {
174+ transposedBaseline: wrap(`
175+ for (var i: u32 = lid; i < N; i = i + WS) {
176+ let l = i % NRT; let j = i / NRT;
177+ var sum: f32 = 0.0;
178+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
179+ sum = sum + st[N + q + l * NTH] * ops[q + j * NTH];
180+ }
181+ st[2u * N + i] = sum;
182+ }`),
183+ colRegs13x120: wrap(`
184+ if (lid < 120u) {
185+ let j = lid % NTH; let b = lid / NTH;
186+ let l0 = b * 13u;
187+ var acc: array<f32, 13>;
188+ for (var l: u32 = 0u; l < 13u; l = l + 1u) { acc[l] = 0.0; }
189+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
190+ let mq = ops[q + j * NTH];
191+ for (var l: u32 = 0u; l < 13u; l = l + 1u) {
192+ acc[l] = acc[l] + st[N + l0 + l + q * NRT] * mq;
193+ }
194+ }
195+ for (var l: u32 = 0u; l < 13u; l = l + 1u) { st[2u * N + l0 + l + j * NRT] = acc[l]; }
196+ }`),
197+ page3colRegs: wrap(`
198+ if (lid < 120u) {
199+ let j = lid % NTH; let b = lid / NTH;
200+ let l0 = b * 13u;
201+ let base = j * NRT * NRT;
202+ var acc: array<f32, 13>;
203+ for (var l: u32 = 0u; l < 13u; l = l + 1u) { acc[l] = 0.0; }
204+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
205+ let x1 = st[N + p + j * NRT];
206+ let x2 = st[2u * N + p + j * NRT];
207+ let x3 = st[3u * N + p + j * NRT];
208+ for (var l: u32 = 0u; l < 13u; l = l + 1u) {
209+ acc[l] = acc[l] + ops[base + l0 + l + p * NRT] * x1
210+ + ops[base + l0 + l + p * NRT + 676u] * x2
211+ + ops[base + l0 + l + p * NRT + 1352u] * x3;
212+ }
213+ }
214+ for (var l: u32 = 0u; l < 13u; l = l + 1u) { st[l0 + l + j * NRT] = acc[l]; }
215+ }`),
216+};
217+
218+async function main2() {
219+ await installNodeWebGpu();
220+ const { device, info } = await requestDevice();
221+ console.log(`device: ${info}, WS=${WS}`);
222+ for (const [label, code] of Object.entries(MORE)) await time(device, code, label);
223+ // dispatch overhead: many tiny dispatches in one submit
224+ {
225+ const code = `
226+@group(0) @binding(0) var<storage, read_write> st: array<f32>;
227+@compute @workgroup_size(64)
228+fn main(@builtin(local_invocation_index) lid: u32) { if (lid == 0u) { st[0] = st[0] + 1.0; } }
229+`;
230+ const module = device.createShaderModule({ code });
231+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
232+ const buf = device.createBuffer({ size: 16, usage: GPUBufferUsage.STORAGE });
233+ const bg = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: buf } }] });
234+ const enc = device.createCommandEncoder();
235+ const pass = enc.beginComputePass();
236+ pass.setPipeline(pipeline); pass.setBindGroup(0, bg);
237+ for (let i = 0; i < 8192; i++) pass.dispatchWorkgroups(1);
238+ pass.end();
239+ const cb = enc.finish();
240+ const t0 = performance.now();
241+ device.queue.submit([cb]);
242+ await device.queue.onSubmittedWorkDone();
243+ console.log(`dispatch overhead: ${((performance.now() - t0) * 1000 / 8192).toFixed(2)} us/dispatch`);
244+ }
245+}
246+if (process.env.MORE) { main2().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); }); }
src/probe3.mjsadded+149−0View file
@@ -0,0 +1,149 @@
1+// vec4 variants with the radial dimension padded to 28 (7 vec4 lanes).
2+// Transform: out[l,j] = sum_q X[l,q] M[q,j]; X stored 28 x 60, M 60 x 60.
3+// Page: out[l,j] = sum_p A[l,p,j] X[p,j]; A pages 28 x 28.
4+import { installNodeWebGpu, requestDevice } from './gpu.mjs';
5+
6+const PADL = 28, NTH = 60, NP = PADL * NTH; // 1680
7+const NPV = NP / 4;
8+const WS = 256;
9+const STEPS = 2048;
10+
11+function wrap(body) {
12+ return `
13+const PADL = ${PADL}u; const NTH = ${NTH}u; const NP = ${NP}u; const WS = ${WS}u; const STEPS = ${STEPS}u;
14+@group(0) @binding(0) var<storage, read> ops: array<f32>;
15+@group(0) @binding(1) var<storage, read> opsv: array<vec4<f32>>;
16+@group(0) @binding(2) var<storage, read_write> stv: array<vec4<f32>>;
17+fn ld(i: u32) -> f32 { return stv[i / 4u][i % 4u]; }
18+@compute @workgroup_size(${WS})
19+fn main(@builtin(local_invocation_index) lid: u32) {
20+ for (var s: u32 = 0u; s < STEPS; s = s + 1u) {
21+${body}
22+ workgroupBarrier();
23+ storageBarrier();
24+ }
25+ if (lid == 0u) { stv[0].x = stv[0].y + 1.0; }
26+}
27+`;
28+}
29+
30+const CASES = {
31+ // 60 threads, 7 vec4 accumulators, X read as vec4, M scalar broadcast
32+ colRegsVec4: wrap(`
33+ if (lid < NTH) {
34+ let j = lid;
35+ var acc: array<vec4<f32>, 7>;
36+ for (var v: u32 = 0u; v < 7u; v = v + 1u) { acc[v] = vec4<f32>(0.0); }
37+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
38+ let mq = ops[q + j * NTH];
39+ let xb = (NP + q * PADL) / 4u;
40+ for (var v: u32 = 0u; v < 7u; v = v + 1u) {
41+ acc[v] = acc[v] + stv[xb + v] * mq;
42+ }
43+ }
44+ let ob = (2u * NP + j * PADL) / 4u;
45+ for (var v: u32 = 0u; v < 7u; v = v + 1u) { stv[ob + v] = acc[v]; }
46+ }`),
47+
48+ // 120 threads: j column plus upper/lower half (4 vec4s h=0, 3 vec4s h=1)
49+ colRegsVec4x120: wrap(`
50+ if (lid < 120u) {
51+ let j = lid % NTH; let h = lid / NTH;
52+ let v0 = h * 4u;
53+ let nv = 4u - h; // 4 or 3
54+ var acc: array<vec4<f32>, 4>;
55+ for (var v: u32 = 0u; v < 4u; v = v + 1u) { acc[v] = vec4<f32>(0.0); }
56+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
57+ let mq = ops[q + j * NTH];
58+ let xb = (NP + q * PADL) / 4u + v0;
59+ for (var v: u32 = 0u; v < 4u; v = v + 1u) {
60+ if (v < nv) { acc[v] = acc[v] + stv[xb + v] * mq; }
61+ }
62+ }
63+ let ob = (2u * NP + j * PADL) / 4u + v0;
64+ for (var v: u32 = 0u; v < 4u; v = v + 1u) {
65+ if (v < nv) { stv[ob + v] = acc[v]; }
66+ }
67+ }`),
68+
69+ // page solve: 60 threads, A columns read as vec4 (28-padded), X broadcast
70+ pageVec4: wrap(`
71+ if (lid < NTH) {
72+ let j = lid;
73+ let base = j * PADL * PADL;
74+ var acc: array<vec4<f32>, 7>;
75+ for (var v: u32 = 0u; v < 7u; v = v + 1u) { acc[v] = vec4<f32>(0.0); }
76+ for (var p: u32 = 0u; p < PADL; p = p + 1u) {
77+ let xq = ld(NP + p + j * PADL);
78+ let ab = (base + p * PADL) / 4u;
79+ for (var v: u32 = 0u; v < 7u; v = v + 1u) {
80+ acc[v] = acc[v] + opsv[ab + v] * xq;
81+ }
82+ }
83+ let ob = (2u * NP + j * PADL) / 4u;
84+ for (var v: u32 = 0u; v < 7u; v = v + 1u) { stv[ob + v] = acc[v]; }
85+ }`),
86+
87+ // triple page solve (the shell update): same but three A blocks, three X
88+ page3Vec4: wrap(`
89+ if (lid < NTH) {
90+ let j = lid;
91+ let base = j * PADL * PADL;
92+ var acc: array<vec4<f32>, 7>;
93+ for (var v: u32 = 0u; v < 7u; v = v + 1u) { acc[v] = vec4<f32>(0.0); }
94+ for (var p: u32 = 0u; p < PADL; p = p + 1u) {
95+ let x1 = ld(NP + p + j * PADL);
96+ let x2 = ld(2u * NP + p + j * PADL);
97+ let x3 = ld(3u * NP + p + j * PADL);
98+ let ab = (base + p * PADL) / 4u;
99+ let s2 = ${PADL * PADL * NTH / 4}u;
100+ for (var v: u32 = 0u; v < 7u; v = v + 1u) {
101+ acc[v] = acc[v] + opsv[ab + v] * x1 + opsv[ab + v + s2] * x2 + opsv[ab + v + 2u * s2] * x3;
102+ }
103+ }
104+ let ob = (j * PADL) / 4u;
105+ for (var v: u32 = 0u; v < 7u; v = v + 1u) { stv[ob + v] = acc[v]; }
106+ }`),
107+};
108+
109+async function time(device, code, label) {
110+ const module = device.createShaderModule({ code });
111+ const info = await module.getCompilationInfo();
112+ const errs = info.messages.filter((m) => m.type === 'error');
113+ if (errs.length) throw new Error(label + ':\n' + errs.map((m) => `${m.lineNum}: ${m.message}`).join('\n'));
114+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
115+ const ops = device.createBuffer({ size: 4 * PADL * PADL * NTH * 3, usage: GPUBufferUsage.STORAGE });
116+ const stb = device.createBuffer({ size: 4 * NP * 4, usage: GPUBufferUsage.STORAGE });
117+ const entries = [];
118+ if (/\bops\[/.test(code)) entries.push({ binding: 0, resource: { buffer: ops } });
119+ if (/\bopsv\[/.test(code)) entries.push({ binding: 1, resource: { buffer: ops } });
120+ entries.push({ binding: 2, resource: { buffer: stb } });
121+ const bg = device.createBindGroup({
122+ layout: pipeline.getBindGroupLayout(0),
123+ entries,
124+ });
125+ const run = async (reps) => {
126+ const enc = device.createCommandEncoder();
127+ for (let r = 0; r < reps; r++) {
128+ const pass = enc.beginComputePass();
129+ pass.setPipeline(pipeline); pass.setBindGroup(0, bg); pass.dispatchWorkgroups(1); pass.end();
130+ }
131+ device.queue.submit([enc.finish()]);
132+ await device.queue.onSubmittedWorkDone();
133+ };
134+ await run(1);
135+ const t0 = performance.now();
136+ await run(4);
137+ const ms = (performance.now() - t0) / 4;
138+ console.log(`${label}: ${(1000 * ms / STEPS).toFixed(2)} us per stage-step`);
139+ ops.destroy(); stb.destroy();
140+}
141+
142+async function main() {
143+ await installNodeWebGpu();
144+ const { device, info } = await requestDevice();
145+ console.log(`device: ${info}`);
146+ for (const [label, code] of Object.entries(CASES)) await time(device, code, label);
147+}
148+
149+main().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); });
src/probe4.mjsadded+131−0View file
@@ -0,0 +1,131 @@
1+// Fully unrolled register-accumulator variants: no local arrays anywhere,
2+// so nothing can spill to scratch memory; accumulators are named scalars.
3+import { installNodeWebGpu, requestDevice } from './gpu.mjs';
4+
5+const NRT = 26, NTH = 60, N = NRT * NTH;
6+const WS = 256;
7+const STEPS = 2048;
8+
9+function wrap(body) {
10+ return `
11+const NRT = ${NRT}u; const NTH = ${NTH}u; const N = ${N}u; const WS = ${WS}u; const STEPS = ${STEPS}u;
12+@group(0) @binding(0) var<storage, read> ops: array<f32>;
13+@group(0) @binding(2) var<storage, read_write> st: array<f32>;
14+@compute @workgroup_size(${WS})
15+fn main(@builtin(local_invocation_index) lid: u32) {
16+ for (var s: u32 = 0u; s < STEPS; s = s + 1u) {
17+${body}
18+ workgroupBarrier();
19+ storageBarrier();
20+ }
21+ if (lid == 0u) { st[0] = st[1] + 1.0; }
22+}
23+`;
24+}
25+
26+// transform: 60 threads, 26 scalar accumulators, unrolled inner block
27+function transformUnrolled() {
28+ const acc = Array.from({ length: NRT }, (_, l) => `var a${l}: f32 = 0.0;`).join(' ');
29+ const fmas = Array.from({ length: NRT }, (_, l) => `a${l} = a${l} + st[xb + ${l}u] * mq;`).join('\n ');
30+ const stores = Array.from({ length: NRT }, (_, l) => `st[ob + ${l}u] = a${l};`).join(' ');
31+ return wrap(`
32+ if (lid < NTH) {
33+ let j = lid;
34+ ${acc}
35+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
36+ let mq = ops[q + j * NTH];
37+ let xb = N + q * NRT;
38+ ${fmas}
39+ }
40+ let ob = 2u * N + j * NRT;
41+ ${stores}
42+ }`);
43+}
44+
45+// transform with two columns per thread (120 threads), 26 accs each... too many regs;
46+// instead: 13 accumulators x 120 threads (j, half), unrolled
47+function transformUnrolledHalf() {
48+ const acc = Array.from({ length: 13 }, (_, l) => `var a${l}: f32 = 0.0;`).join(' ');
49+ const fmas = Array.from({ length: 13 }, (_, l) => `a${l} = a${l} + st[xb + ${l}u] * mq;`).join('\n ');
50+ const stores = Array.from({ length: 13 }, (_, l) => `st[ob + ${l}u] = a${l};`).join(' ');
51+ return wrap(`
52+ if (lid < 120u) {
53+ let j = lid % NTH; let h = lid / NTH;
54+ let l0 = h * 13u;
55+ ${acc}
56+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
57+ let mq = ops[q + j * NTH];
58+ let xb = N + q * NRT + l0;
59+ ${fmas}
60+ }
61+ let ob = 2u * N + j * NRT + l0;
62+ ${stores}
63+ }`);
64+}
65+
66+// page solve unrolled the same way: 60 threads, 26 accs, A read coalesced
67+function pageUnrolled() {
68+ const acc = Array.from({ length: NRT }, (_, l) => `var a${l}: f32 = 0.0;`).join(' ');
69+ const fmas = Array.from({ length: NRT }, (_, l) => `a${l} = a${l} + ops[ab + ${l}u] * xq;`).join('\n ');
70+ const stores = Array.from({ length: NRT }, (_, l) => `st[ob + ${l}u] = a${l};`).join(' ');
71+ return wrap(`
72+ if (lid < NTH) {
73+ let j = lid;
74+ let base = j * NRT * NRT;
75+ ${acc}
76+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
77+ let xq = st[N + p + j * NRT];
78+ let ab = base + p * NRT;
79+ ${fmas}
80+ }
81+ let ob = 2u * N + j * NRT;
82+ ${stores}
83+ }`);
84+}
85+
86+const CASES = {
87+ transformUnrolled: transformUnrolled(),
88+ transformUnrolledHalf: transformUnrolledHalf(),
89+ pageUnrolled: pageUnrolled(),
90+};
91+
92+async function time(device, code, label) {
93+ const module = device.createShaderModule({ code });
94+ const info = await module.getCompilationInfo();
95+ const errs = info.messages.filter((m) => m.type === 'error');
96+ if (errs.length) throw new Error(label + ':\n' + errs.map((m) => `${m.lineNum}: ${m.message}`).join('\n'));
97+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
98+ const ops = device.createBuffer({ size: 4 * NRT * NRT * NTH, usage: GPUBufferUsage.STORAGE });
99+ const stb = device.createBuffer({ size: 4 * N * 4, usage: GPUBufferUsage.STORAGE });
100+ const bg = device.createBindGroup({
101+ layout: pipeline.getBindGroupLayout(0),
102+ entries: [
103+ { binding: 0, resource: { buffer: ops } },
104+ { binding: 2, resource: { buffer: stb } },
105+ ],
106+ });
107+ const run = async (reps) => {
108+ const enc = device.createCommandEncoder();
109+ for (let r = 0; r < reps; r++) {
110+ const pass = enc.beginComputePass();
111+ pass.setPipeline(pipeline); pass.setBindGroup(0, bg); pass.dispatchWorkgroups(1); pass.end();
112+ }
113+ device.queue.submit([enc.finish()]);
114+ await device.queue.onSubmittedWorkDone();
115+ };
116+ await run(1);
117+ const t0 = performance.now();
118+ await run(4);
119+ const ms = (performance.now() - t0) / 4;
120+ console.log(`${label}: ${(1000 * ms / STEPS).toFixed(2)} us per stage-step`);
121+ ops.destroy(); stb.destroy();
122+}
123+
124+async function main() {
125+ await installNodeWebGpu();
126+ const { device, info } = await requestDevice();
127+ console.log(`device: ${info}`);
128+ for (const [label, code] of Object.entries(CASES)) await time(device, code, label);
129+}
130+
131+main().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); });
src/probe5.mjsadded+130−0View file
@@ -0,0 +1,130 @@
1+// Multi-dispatch step: each stage its own dispatch at full GPU width.
2+// Representative step = 2 stencil stages + 4 transforms + 2 page solves,
3+// recorded 512 steps per command buffer.
4+import { installNodeWebGpu, requestDevice } from './gpu.mjs';
5+
6+const NRT = 26, NTH = 60, N = NRT * NTH;
7+const WG = Number(process.env.WG ?? 128);
8+const NWG = Math.ceil(N / WG);
9+const STEPS = 512;
10+
11+const DYNOFF = !!process.env.DYNOFF;
12+const GATE = !!process.env.GATE;
13+const HDR = `
14+const NRT = ${NRT}u; const NTH = ${NTH}u; const N = ${N}u;
15+@group(0) @binding(0) var<storage, read> ops: array<f32>;
16+@group(0) @binding(1) var<storage, read_write> st: array<f32>;
17+${DYNOFF ? '@group(0) @binding(2) var<uniform> ctl: vec4<u32>;' : ''}
18+`;
19+
20+const KERNELS = {
21+ stencil: `${HDR}
22+@compute @workgroup_size(${WG})
23+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
24+ let i = gid.x;
25+ if (i >= N) { return; }
26+ ${GATE ? 'if (st[1] == 0.0) { return; }' : ''}
27+ let l = i % NRT; let j = i / NRT;
28+ let jm = (j + NTH - 1u) % NTH; let jp = (j + 1u) % NTH;
29+ var v: f32 = 0.0;
30+ if (l >= 1u && l <= NRT - 2u) {
31+ v = (st[N + (l - 1u) + j * NRT] - 2.0 * st[N + l + j * NRT] + st[N + (l + 1u) + j * NRT]) * 0.5
32+ + (st[N + l + jm * NRT] - st[N + l + jp * NRT]) * 0.25;
33+ }
34+ st[2u * N + i] = v;
35+}`,
36+ transform: `${HDR}
37+@compute @workgroup_size(${WG})
38+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
39+ let i = gid.x;
40+ if (i >= N) { return; }
41+ ${GATE ? 'if (st[1] == 0.0) { return; }' : ''}
42+ let l = i % NRT; let j = i / NRT;
43+ var sum: f32 = 0.0;
44+ for (var q: u32 = 0u; q < NTH; q = q + 1u) {
45+ sum = sum + st[N + l + q * NRT] * ops[q + j * NTH];
46+ }
47+ st[2u * N + i] = sum + f32(${DYNOFF ? 'ctl.x' : '0u'});
48+}`,
49+ page: `${HDR}
50+@compute @workgroup_size(${WG})
51+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
52+ let i = gid.x;
53+ if (i >= N) { return; }
54+ ${GATE ? 'if (st[1] == 0.0) { return; }' : ''}
55+ let l = i % NRT; let j = i / NRT;
56+ let base = j * NRT * NRT;
57+ var sum: f32 = 0.0;
58+ for (var p: u32 = 0u; p < NRT; p = p + 1u) {
59+ sum = sum + ops[base + l + p * NRT] * st[N + p + j * NRT];
60+ }
61+ st[2u * N + i] = sum;
62+}`,
63+};
64+
65+async function main() {
66+ await installNodeWebGpu();
67+ const { device, info } = await requestDevice();
68+ console.log(`device: ${info}, wg size ${WG} x ${NWG} workgroups`);
69+
70+ const BIG = Number(process.env.BIG ?? 1);
71+ const ops = device.createBuffer({ size: 4 * NRT * NRT * NTH * BIG, usage: GPUBufferUsage.STORAGE });
72+ const stb = device.createBuffer({ size: 4 * N * 4 * BIG, usage: GPUBufferUsage.STORAGE });
73+ const uni = device.createBuffer({ size: STEPS * 256, usage: GPUBufferUsage.UNIFORM });
74+
75+ const pipes = {};
76+ for (const [name, code] of Object.entries(KERNELS)) {
77+ const module = device.createShaderModule({ code });
78+ const errs = (await module.getCompilationInfo()).messages.filter((m) => m.type === 'error');
79+ if (errs.length) throw new Error(name + ': ' + errs.map((m) => m.message).join('\n'));
80+ const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'main' } });
81+ const entries = [
82+ { binding: 0, resource: { buffer: ops } },
83+ { binding: 1, resource: { buffer: stb } },
84+ ];
85+ if (DYNOFF && name === 'transform') entries.push({ binding: 2, resource: { buffer: uni, offset: 0, size: 16 } });
86+ const bg = device.createBindGroup({
87+ layout: pipeline.getBindGroupLayout(0),
88+ entries,
89+ });
90+ pipes[name] = { pipeline, bg };
91+ }
92+
93+ let stepSeq = ['stencil', 'transform', 'page', 'transform', 'stencil', 'transform', 'page', 'transform'];
94+ if (process.env.SEQ1) stepSeq = ['transform'];
95+ if (process.env.SEQN) stepSeq = Array(Number(process.env.SEQN)).fill('transform');
96+ const record = () => {
97+ const enc = device.createCommandEncoder();
98+ const pass = enc.beginComputePass();
99+ for (let s = 0; s < STEPS; s++) {
100+ for (const name of stepSeq) {
101+ const p = pipes[name];
102+ pass.setPipeline(p.pipeline);
103+ pass.setBindGroup(0, p.bg);
104+ pass.dispatchWorkgroups(NWG);
105+ }
106+ }
107+ pass.end();
108+ return enc.finish();
109+ };
110+
111+ device.queue.writeBuffer(stb, 4, new Float32Array([1]));
112+ // warmup
113+ device.queue.submit([record()]);
114+ await device.queue.onSubmittedWorkDone();
115+
116+ let tRec = 0;
117+ const t0 = performance.now();
118+ const reps = Number(process.env.REPS ?? 4);
119+ for (let r = 0; r < reps; r++) {
120+ const tr = performance.now();
121+ const cb = record();
122+ tRec += performance.now() - tr;
123+ device.queue.submit([cb]);
124+ }
125+ await device.queue.onSubmittedWorkDone();
126+ const ms = (performance.now() - t0) / reps;
127+ console.log(`8-dispatch step: ${(1000 * ms / STEPS).toFixed(2)} us/step total (cpu recording ${(1000 * tRec / reps / STEPS).toFixed(2)} us/step)`);
128+}
129+
130+main().then(() => process.exit(0)).catch((e) => { console.error(e); process.exit(1); });
src/run.mjsadded+161−0View file
@@ -0,0 +1,161 @@
1+// Driver: numbl setup stage -> WebGPU time loop -> numbl output stage.
2+//
3+// node src/run.mjs full default render, writes cymbal2-gpu.wav
4+// node src/run.mjs --pinned the validate suite's reduced configuration
5+// node src/run.mjs --compare GPU vs CPU (numbl cymbal2_sim) on the pinned config
6+// node src/run.mjs --mode mega|multi|persistent backend (default multi)
7+// node src/run.mjs --velocity 4 ... parameter overrides
8+
9+import { writeFileSync } from 'node:fs';
10+import { runSetup, runOutput, runCpuSim, loadMFiles } from './setup.mjs';
11+import { buildKernel, buildLayout } from './kernel.mjs';
12+import { buildStages } from './kernels2.mjs';
13+import { buildPersistent } from './kernels3.mjs';
14+import { installNodeWebGpu, requestDevice, runMegakernel, runMultiDispatch, runPersistent } from './gpu.mjs';
15+import { encodeWav } from './wav.mjs';
16+
17+function parseArgs(argv) {
18+ const a = { params: {}, chunk: 2048, ws: 256, wav: 'cymbal2-gpu.wav', compare: false, bench: false, mode: 'multi' };
19+ for (let i = 2; i < argv.length; i++) {
20+ const t = argv[i];
21+ if (t === '--pinned') Object.assign(a.params, { sample_rate: 44100, oversample: 1, tail_s: 1.5 });
22+ else if (t === '--compare') { a.compare = true; Object.assign(a.params, { sample_rate: 44100, oversample: 1, tail_s: 1.5 }); }
23+ else if (t === '--bench') a.bench = true;
24+ else if (t === '--velocity') a.params.strike_velocity = Number(argv[++i]);
25+ else if (t === '--count') a.params.strike_count = Number(argv[++i]);
26+ else if (t === '--rate') a.params.strike_rate_hz = Number(argv[++i]);
27+ else if (t === '--oversample') a.params.oversample = Number(argv[++i]);
28+ else if (t === '--tail') a.params.tail_s = Number(argv[++i]);
29+ else if (t === '--chunk') a.chunk = Number(argv[++i]);
30+ else if (t === '--ws') a.ws = Number(argv[++i]);
31+ else if (t === '--wav') a.wav = argv[++i];
32+ else if (t === '--mode') a.mode = argv[++i];
33+ else if (t === '--set') { const [k, v] = argv[++i].split('='); a.params[k] = Number(v); }
34+ else throw new Error(`unknown argument ${t}`);
35+ }
36+ return a;
37+}
38+
39+function hostDiagnostics(fhist, state, kernel, setup) {
40+ let peakF = 0, contactSteps = 0, events = 0, prev = 0;
41+ for (let i = 0; i < fhist.length; i++) {
42+ const f = fhist[i];
43+ if (f > peakF) peakF = f;
44+ if (f > 0) { contactSteps++; if (prev === 0) events++; }
45+ prev = f;
46+ }
47+ const L = kernel.layout;
48+ return {
49+ peakForceN: peakF,
50+ contactMs: (1000 * contactSteps) / setup.simRate,
51+ contactEvents: events,
52+ peakXiMm: state[L.SC + 12] * 1000,
53+ peakDispThickness: state[L.SC + 13] / setup.thickness,
54+ contactVel: Array.from(state.subarray(L.ST_CVEL, L.ST_CVEL + setup.K)),
55+ strokesPlayed: Math.min(state[L.SC + 4] + (state[L.SC + 3] > 0 ? 1 : 0), setup.K),
56+ };
57+}
58+
59+const rms = (x) => {
60+ let s = 0;
61+ for (let i = 0; i < x.length; i++) s += x[i] * x[i];
62+ return Math.sqrt(s / x.length);
63+};
64+
65+async function main() {
66+ const args = parseArgs(process.argv);
67+
68+ console.log('setup (numbl): grid, per-wavenumber inverses, mallet, schedule ...');
69+ let t0 = performance.now();
70+ const setup = runSetup(args.params);
71+ console.log(` ${(performance.now() - t0).toFixed(0)} ms; grid ${setup.nrt}x${setup.nth}, nfree ${setup.nfree}, nsteps ${setup.nsteps}, simRate ${setup.simRate}`);
72+
73+ await installNodeWebGpu();
74+ const { device, info } = await requestDevice();
75+ console.log(`device: ${info}`);
76+
77+ let kernel;
78+ let gpu;
79+ t0 = performance.now();
80+ const progress = (k, n, ms) => {
81+ if (k === n || k % (1024 * 64) === 0) {
82+ process.stdout.write(`\r ${k}/${n} steps, ${(1000 * ms / k).toFixed(1)} us/step`);
83+ }
84+ };
85+ if (args.mode === 'persistent') {
86+ const L = buildLayout(setup);
87+ kernel = L;
88+ const pk = buildPersistent(L, { chunk: args.chunk >= 4096 ? 512 : Math.min(args.chunk, 2048) });
89+ console.log(`persistent kernel: ${pk.nwg} workgroups of ${pk.wg}, ${pk.chunk} steps/dispatch, ops ${(L.opsData.byteLength / 1024).toFixed(0)} KB`);
90+ gpu = await runPersistent(device, L, pk, progress);
91+ } else if (args.mode === 'mega') {
92+ kernel = buildKernel(setup, { ws: args.ws, chunk: args.chunk });
93+ console.log(`megakernel: ws ${args.ws}, ${args.chunk} steps/dispatch, ops ${(kernel.opsData.byteLength / 1024).toFixed(0)} KB`);
94+ gpu = await runMegakernel(device, kernel, setup.nsteps, progress);
95+ } else {
96+ const L = buildLayout(setup);
97+ kernel = L; // shares .layout for diagnostics
98+ const stages = buildStages(L, {});
99+ console.log(`multi-dispatch: ${stages.order.length} dispatches/step, wg ${stages.wg} x ${stages.nwg}, ops ${(L.opsData.byteLength / 1024).toFixed(0)} KB`);
100+ gpu = await runMultiDispatch(device, L, stages, progress);
101+ }
102+ console.log('');
103+ const usPerStep = (1000 * gpu.gpuMs) / setup.nsteps;
104+ console.log(`gpu integration: ${(gpu.gpuMs / 1000).toFixed(2)} s total, ${usPerStep.toFixed(1)} us/step`);
105+
106+ let bad = 0;
107+ let firstBad = -1;
108+ for (let i = 0; i < gpu.acc.length; i++) if (!Number.isFinite(gpu.acc[i])) { bad++; if (firstBad < 0) firstBad = i; }
109+ if (process.env.DUMP_ACC) {
110+ const { writeFileSync: wf } = await import('node:fs');
111+ wf(process.env.DUMP_ACC, Buffer.from(gpu.acc.buffer, gpu.acc.byteOffset, gpu.acc.byteLength));
112+ console.log(`acc dumped; nonfinite ${bad}${bad ? ' first at ' + firstBad : ''}`);
113+ }
114+ if (bad > 0 && !process.env.DUMP_ACC) throw new Error(`simulation diverged: ${bad} non-finite acceleration samples (first at ${firstBad})`);
115+ if (bad > 0 && process.env.DUMP_ACC) { console.log('skipping output stage (diverged)'); return; }
116+
117+ const d = hostDiagnostics(gpu.fhist, gpu.state, kernel, setup);
118+ console.log(`contact: peakF ${d.peakForceN.toFixed(2)} N, cvel ${d.contactVel[0].toFixed(3)} m/s, ` +
119+ `${d.contactMs.toFixed(2)} ms over ${d.contactEvents} events, peakXi ${d.peakXiMm.toFixed(3)} mm, ` +
120+ `peakDisp ${d.peakDispThickness.toFixed(1)} thicknesses`);
121+
122+ if (args.compare) {
123+ console.log('cpu reference (numbl, unmodified cymbal2_sim) ...');
124+ const cpu = runCpuSim(setup.mfiles, args.params);
125+ console.log(` ${(cpu.ms / 1000).toFixed(1)} s total, ${(1000 * cpu.ms / cpu.nsteps).toFixed(1)} us/step`);
126+ console.log(` cpu: peakF ${cpu.peakForceN.toFixed(2)} N, cvel ${cpu.contactVel.toFixed(3)}, contactMs ${cpu.contactMs.toFixed(2)}, events ${cpu.contactEvents}, peakXi ${cpu.peakXiMm.toFixed(3)} mm`);
127+
128+ // Contact window alignment and force agreement
129+ const firstC = (f) => f.findIndex((x) => x > 0);
130+ const kg = firstC(gpu.fhist), kc = firstC(cpu.fhist);
131+ console.log(` first contact step: gpu ${kg}, cpu ${kc}${kg === kc ? '' : ' << MISALIGNED'}`);
132+ let femax = 0;
133+ for (let i = 0; i < cpu.fhist.length; i++) {
134+ const e = Math.abs(gpu.fhist[i] - cpu.fhist[i]);
135+ if (e > femax) femax = e;
136+ }
137+ console.log(` max |force diff| ${femax.toExponential(2)} N (peak ${cpu.peakForceN.toFixed(1)} N)`);
138+
139+ // Short-horizon acceleration agreement (before chaotic divergence)
140+ for (const w of [256, 1024, 4096]) {
141+ let num = 0, den = 0;
142+ for (let i = kc; i < Math.min(kc + w, cpu.acc.length); i++) {
143+ num += (gpu.acc[i] - cpu.acc[i]) ** 2;
144+ den += cpu.acc[i] ** 2;
145+ }
146+ console.log(` acc rel-L2 over ${w} steps after contact: ${Math.sqrt(num / Math.max(den, 1e-30)).toExponential(2)}`);
147+ }
148+ console.log(` acc rms: gpu ${rms(gpu.acc).toExponential(4)}, cpu ${rms(cpu.acc).toExponential(4)}`);
149+ }
150+
151+ console.log('output stage (numbl cymbal2_output) ...');
152+ const out = runOutput(loadMFiles(), args.params, gpu.acc);
153+ console.log(` y: ${out.y.length} samples, rms ${rms(out.y).toFixed(4)}, peakAccel ${out.peakAccel.toExponential(3)} m/s^2`);
154+ writeFileSync(args.wav, encodeWav(out.y, args.params.sample_rate ?? 48000));
155+ console.log(`wrote ${args.wav}`);
156+}
157+
158+main().then(() => process.exit(0)).catch((e) => {
159+ console.error(e);
160+ process.exit(1);
161+});
src/setup.mjsadded+203−0View file
@@ -0,0 +1,203 @@
1+// Setup and output stages of the cymbal2 render, run through numbl exactly as
2+// the webapp runs them. The GPU path replaces only cymbal2_sim: everything the
3+// time loop needs (grid operators, per-wavenumber inverses, mallet matrices,
4+// stroke schedule) is computed here by the unmodified MATLAB code and pulled
5+// out as tensors.
6+
7+import { readFileSync, readdirSync } from 'node:fs';
8+import { fileURLToPath } from 'node:url';
9+import path from 'node:path';
10+import { executeCode, RTV } from 'numbl';
11+
12+const MATLAB_DIR = path.resolve(
13+ path.dirname(fileURLToPath(import.meta.url)),
14+ '../../../simulated_instruments/cymbal2/matlab',
15+);
16+
17+export function loadMFiles() {
18+ return readdirSync(MATLAB_DIR)
19+ .filter((f) => f.endsWith('.m'))
20+ .sort()
21+ .map((name) => ({
22+ name,
23+ source: readFileSync(path.join(MATLAB_DIR, name), 'utf8'),
24+ }));
25+}
26+
27+// MATLAB lines overriding defaults from a plain params object.
28+function pOverrides(params) {
29+ const lines = [];
30+ for (const [k, v] of Object.entries(params ?? {})) {
31+ if (typeof v === 'number') lines.push(`P.${k} = ${v};`);
32+ else if (typeof v === 'boolean') lines.push(`P.${k} = ${v ? 1 : 0};`);
33+ }
34+ return lines.join('\n');
35+}
36+
37+function stage(mfiles, source, inputs) {
38+ return executeCode(
39+ source,
40+ { initialVariableValues: inputs, displayResults: false },
41+ mfiles,
42+ 'stage.m',
43+ );
44+}
45+
46+function tensor(v, name) {
47+ if (!v || v.kind !== 'tensor') throw new Error(`expected tensor for ${name}`);
48+ return v; // {kind, shape, data: Float64Array} column-major
49+}
50+
51+function scalar(v, name) {
52+ if (typeof v === 'number') return v;
53+ if (v && v.kind === 'tensor' && v.data.length === 1) return v.data[0];
54+ throw new Error(`expected scalar for ${name}`);
55+}
56+
57+function f32(v) {
58+ const t = v.data;
59+ const out = new Float32Array(t.length);
60+ out.set(t);
61+ return out;
62+}
63+
64+// The stroke-schedule block is the one in cymbal2_sim.m, verbatim, so the GPU
65+// render integrates exactly the schedule the CPU render would.
66+const SCHEDULE_SRC = `
67+K = max(1, round(P.strike_count));
68+tstr = P.pre_delay_ms * 1e-3 + (0:K - 1) / max(1e-3, P.strike_rate_hz);
69+vstr = zeros(1, K);
70+state = P.seed * 2654435761;
71+for i = 1:K
72+ state = mod(state * 1664525 + 1013904223, 4294967296);
73+ x = state / 4294967296;
74+ vstr(i) = max(0.05, P.strike_velocity * (1 + P.velocity_spread * (2 * x - 1)));
75+end
76+gap = P.strike_gap_m;
77+if gap <= 0
78+ gap = 1.05 * max(vstr)^2 / (2 * P.strike_accel) + 0.005;
79+end
80+travel = gap ./ vstr + vstr / (2 * P.strike_accel);
81+tlaunch = tstr - travel;
82+off = -min(0, min(tlaunch));
83+tlaunch = tlaunch + off;
84+duration = off + (K - 1) / max(1e-3, P.strike_rate_hz) + P.tail_s;
85+nsteps = floor(duration * P.sample_rate * P.oversample);
86+`;
87+
88+export function runSetup(params) {
89+ const mfiles = loadMFiles();
90+ const src = `
91+P = cymbal2_defaults();
92+${pOverrides(params)}
93+S = cymbal2_system(P);
94+MB = cymbal2_mallet(P);
95+G = S.G;
96+o_D1 = G.D1; o_D2 = G.D2;
97+o_F = G.F; o_Gi = G.Gi;
98+o_Rd = G.Rd; o_R2d = G.R2d;
99+o_AiB = S.AiB; o_AiC = S.AiC; o_Ai = S.Ai; o_Fi = S.Fi;
100+o_ASiB = MB.ASiB; o_ASiC = MB.ASiC;
101+o_vM = MB.vM; o_vR = MB.vR; o_vK = MB.vK; o_tipcol = MB.tipcol;
102+${SCHEDULE_SRC}
103+o_tlaunch = tlaunch; o_vstr = vstr;
104+c_nrt = G.nrt; c_Nr = G.Nr; c_nth = G.nth; c_hr = G.hr; c_hth = G.hth;
105+c_ht = S.ht; c_Eh = S.Eh; c_Rs = S.Rs; c_rhoh = S.rhoh;
106+c_fsc = S.forceScale; c_nlr0 = S.eomRow0; c_nDropped = S.nDropped;
107+c_ls = S.ls; c_js = S.js; c_lp = S.lp; c_jp = S.jp;
108+c_strikeR = S.strikeRadiusM; c_pickupR = S.pickupRadiusM;
109+c_tipi = MB.tipi; c_nfree = MB.nfree; c_ps = MB.ps;
110+c_KM = P.head_stiffness; c_alM = P.head_alpha;
111+c_aH = P.strike_accel; c_aHr = P.retract_accel;
112+c_holdSteps = max(1, round(P.strike_hold_ms * 1e-3 / S.ht));
113+c_missDepth = gap + 0.06;
114+c_gap = gap; c_K = K; c_nsteps = nsteps;
115+c_simRate = P.sample_rate * P.oversample;
116+c_thickness = P.thickness_m;
117+`;
118+ const r = stage(mfiles, src);
119+ const v = r.variableValues;
120+ const g = (n) => tensor(v[n], n);
121+ const s = (n) => scalar(v[n], n);
122+
123+ const setup = {
124+ mfiles,
125+ ops: {
126+ D1: f32(g('o_D1')), D2: f32(g('o_D2')),
127+ F: f32(g('o_F')), Gi: f32(g('o_Gi')),
128+ Rd: f32(g('o_Rd')), R2d: f32(g('o_R2d')),
129+ AiB: f32(g('o_AiB')), AiC: f32(g('o_AiC')),
130+ Ai: f32(g('o_Ai')), Fi: f32(g('o_Fi')),
131+ ASiB: f32(g('o_ASiB')), ASiC: f32(g('o_ASiC')),
132+ vM: f32(g('o_vM')), vR: f32(g('o_vR')),
133+ vK: f32(g('o_vK')), tipcol: f32(g('o_tipcol')),
134+ },
135+ tlaunch: Array.from(g('o_tlaunch').data),
136+ vstr: Array.from(g('o_vstr').data),
137+ nrt: s('c_nrt'), Nr: s('c_Nr'), nth: s('c_nth'),
138+ hr: s('c_hr'), hth: s('c_hth'),
139+ ht: s('c_ht'), Eh: s('c_Eh'), Rs: s('c_Rs'), rhoh: s('c_rhoh'),
140+ fsc: s('c_fsc'), nlr0: s('c_nlr0'), nDropped: s('c_nDropped'),
141+ ls: s('c_ls'), js: s('c_js'), lp: s('c_lp'), jp: s('c_jp'),
142+ strikeR: s('c_strikeR'), pickupR: s('c_pickupR'),
143+ tipi: s('c_tipi'), nfree: s('c_nfree'), ps: s('c_ps'),
144+ KM: s('c_KM'), alM: s('c_alM'), aH: s('c_aH'), aHr: s('c_aHr'),
145+ holdSteps: s('c_holdSteps'), missDepth: s('c_missDepth'),
146+ gap: s('c_gap'), K: s('c_K'), nsteps: s('c_nsteps'),
147+ simRate: s('c_simRate'), thickness: s('c_thickness'),
148+ };
149+ if (setup.vstr.length !== setup.K) throw new Error('schedule length mismatch');
150+ return setup;
151+}
152+
153+// The output stage (DC blocker, decimation, fade, normalize), run through the
154+// unmodified cymbal2_output.m with the GPU-computed acceleration handed in.
155+export function runOutput(mfiles, params, acc) {
156+ const accT = RTV.tensor(Float64Array.from(acc), [1, acc.length]);
157+ const r = stage(
158+ mfiles,
159+ `
160+P = cymbal2_defaults();
161+${pOverrides(params)}
162+[y, peakAccel] = cymbal2_output(P, acc);
163+`,
164+ { acc: accT },
165+ );
166+ const y = tensor(r.variableValues.y, 'y');
167+ return { y: f32(y), peakAccel: scalar(r.variableValues.peakAccel, 'peakAccel') };
168+}
169+
170+// CPU reference: the unmodified cymbal2_sim through numbl (slow; for
171+// validation of the GPU path on reduced configurations).
172+export function runCpuSim(mfiles, params) {
173+ const t0 = performance.now();
174+ const r = stage(
175+ mfiles,
176+ `
177+P = cymbal2_defaults();
178+${pOverrides(params)}
179+S = cymbal2_system(P);
180+MB = cymbal2_mallet(P);
181+[acc, dinfo, fhist] = cymbal2_sim(P, S, MB);
182+d_peakF = dinfo.peakForceN;
183+d_cvel = dinfo.contactVel(1);
184+d_contactMs = dinfo.contactMs;
185+d_events = dinfo.contactEvents;
186+d_xi = dinfo.peakXiMm;
187+d_nsteps = dinfo.nsteps;
188+`,
189+ );
190+ const ms = performance.now() - t0;
191+ const v = r.variableValues;
192+ return {
193+ acc: f32(tensor(v.acc, 'acc')),
194+ fhist: f32(tensor(v.fhist, 'fhist')),
195+ peakForceN: scalar(v.d_peakF, 'd_peakF'),
196+ contactVel: scalar(v.d_cvel, 'd_cvel'),
197+ contactMs: scalar(v.d_contactMs, 'd_contactMs'),
198+ contactEvents: scalar(v.d_events, 'd_events'),
199+ peakXiMm: scalar(v.d_xi, 'd_xi'),
200+ nsteps: scalar(v.d_nsteps, 'd_nsteps'),
201+ ms,
202+ };
203+}
src/wav.mjsadded+18−0View file
@@ -0,0 +1,18 @@
1+// Minimal 16-bit PCM WAV writer.
2+export function encodeWav(y, sampleRate) {
3+ const n = y.length;
4+ const buf = new ArrayBuffer(44 + 2 * n);
5+ const dv = new DataView(buf);
6+ const wstr = (o, s) => { for (let i = 0; i < s.length; i++) dv.setUint8(o + i, s.charCodeAt(i)); };
7+ wstr(0, 'RIFF'); dv.setUint32(4, 36 + 2 * n, true); wstr(8, 'WAVE');
8+ wstr(12, 'fmt '); dv.setUint32(16, 16, true);
9+ dv.setUint16(20, 1, true); dv.setUint16(22, 1, true);
10+ dv.setUint32(24, sampleRate, true); dv.setUint32(28, sampleRate * 2, true);
11+ dv.setUint16(32, 2, true); dv.setUint16(34, 16, true);
12+ wstr(36, 'data'); dv.setUint32(40, 2 * n, true);
13+ for (let i = 0; i < n; i++) {
14+ const v = Math.max(-1, Math.min(1, y[i]));
15+ dv.setInt16(44 + 2 * i, Math.round(v * 32767), true);
16+ }
17+ return Buffer.from(buf);
18+}