feat(fluids): one GPU particle pool, stepped by a WebGPU compute pass; WebGL2 refused by name until #759 (#420)

docs/ENGINE.md

@@ -769,6 +769,7 @@ writes a pose buffer and an event buffer. No emscripten glue is kept; the engine
 - **Timing.** The `physics` stage of `WEBGPU_STAGES` / `WEBGL_STAGES` (host step `physicsMs`) is the
   page's share; the worker's per-step time is reported apart, in `world.physics.stats.stepMs`
   (the module's step alone, the clock `scripts/bench-physics.ts` reads in Node).
+  Its GPU column is the particle step (`Trillion3D particles`, `particles/webgpuParticles.ts`).
 - **Threads.** On a cross-origin isolated page the page loads `joltPhysicsThreads.wasm` (atomics,
   bulk memory, shared memory) and Jolt's own thread pool steps it: each pool thread starts in C
   through `pthread_create`, which the loader (`physics/joltThreads.ts`) answers with a worker that

packages/sdk-browser/src/backend/types.ts

@@ -172,6 +172,8 @@ export interface BackendContext {
   sceneLighting?: Object3D;
   /** The page's guides, held by its world (`guides/guideSet.ts`): drawn over the image. */
   guides?: import('../guides/guideSet.ts').GuideSet;
+  /** The world's particle pools, stepped once per image (`../particles/`). */
+  particles?: readonly import('../../../sdk-core/src/fluids/particles.ts').ParticlePool[];
   /** Contract lights, owned by the host and shared by every engine of the session. */
   sceneLights?: SceneLightStore;
   /** Imported light ids, in cache order: the host sets or removes them (`importedLights()`). */

packages/sdk-browser/src/gpu/core/engineShaders.fixture.ts

@@ -43,6 +43,7 @@ import { REDUCE_WGSL } from '../../webgpu/tile/reduce.ts';
 import { TRANSPARENT_COMPACT_SHADER } from '../../webgpu/transparent/shader.ts';
 import { WATER_COMPOSITE_SHADER } from '../../webgpu/water/compositeWgsl.ts';
 import { WATER_SURFACE_WGSL } from '../../webgpu/water/surfaceWgsl.ts';
+import { PARTICLES_WGSL } from '../../particles/webgpuParticles.ts';
 
 const compositions = (label: string, sources: Record<string, string>) =>
   Object.fromEntries(Object.entries(sources).map(([input, code]) => [`${label}_${input}`, code]));
@@ -88,4 +89,5 @@ export const ENGINE_SHADERS: Record<string, string> = {
   REDUCE_WGSL,
   TRANSPARENT_COMPACT_SHADER,
   WATER_COMPOSITE_SHADER,
+  PARTICLES_WGSL,
 };

packages/sdk-browser/src/measurement/measurement.ts

@@ -19,6 +19,8 @@ export type { RenderBackend, BackendFactory, MeasuredWorldOptions } from '../bac
 export { replicateInstances } from '../scene/replicateInstances.ts';
 export { autonomousPagesBackend } from '../backend/autonomous/pages.ts';
 export { webgpuPagesBackend } from '../webgpu/pages/pages.ts';
+export { attachParticles } from '../world/core/worldSession.ts';
+export { ParticlePool } from '../../../sdk-core/src/fluids/particles.ts';
 
 /** Browser job adapter. A completed session is owned by the caller; cancel/fail after construct disposes it. */
 export async function createMeasuredWorldJob(

packages/sdk-browser/src/particles/particles.test.ts

@@ -0,0 +1,80 @@
+// The CPU half of the WebGPU particle step (#420): the words it hands the GPU for a pool, and a
+// frame with a pool never held. What the GPU does with them is the measurer's.
+import test from 'node:test';
+import assert from 'node:assert/strict';
+import { setImmediate as tick } from 'node:timers/promises';
+import { fakeDevice, written } from '../../../../tests/kit/gpu/fakeDevice.ts';
+import { holdWebgpuFrame, keepWebgpuFrame } from '../webgpu/frame/hold.ts';
+import { settledRt } from '../webgpu/frame/hold.fixture.ts';
+import { ParticlePool } from '../../../sdk-core/src/fluids/particles.ts';
+import { PARTICLES_PASS, createWebgpuParticles } from './webgpuParticles.ts';
+
+/** An encoder that records its compute passes and their dispatches. */
+function computeRecorder() {
+  const passes: { label?: string; dispatches: number[] }[] = [];
+  const encoder = {
+    beginComputePass: ({ label }: GPUComputePassDescriptor) => {
+      const pass = { label, dispatches: [] as number[] };
+      passes.push(pass);
+      return {
+        setPipeline() {},
+        setBindGroup() {},
+        dispatchWorkgroups: (x: number) => void pass.dispatches.push(x),
+        end() {},
+      };
+    },
+  } as unknown as GPUCommandEncoder;
+  return { encoder, passes };
+}
+
+test('WebGPU: one timed pass writes the step words and the staged records, once', async () => {
+  const gpu = fakeDevice();
+  const particles = createWebgpuParticles(gpu.device, (error) => assert.fail(String(error)));
+  const pool = new ParticlePool({ capacity: 1000, emitPerFrame: 1000 });
+  for (let i = 0; i < 3; i++) pool.emit(i, 1, 2, 3, 4, 5, 6);
+  pool.advance(0.01);
+  const { encoder, passes } = computeRecorder();
+  assert.equal(particles.run([pool], encoder), 0, 'compiling: the pool waits, its records kept');
+  await tick();
+  assert.equal(particles.run([pool], encoder), 1);
+  assert.deepEqual(passes, [{ label: PARTICLES_PASS, dispatches: [1] }], 'the emitted slots only');
+  const [step, records] = gpu.writes;
+  const words = new Uint8Array(written(step)).buffer;
+  assert.deepEqual(
+    [...new Float32Array(words, 0, 4)],
+    [0, Math.fround(-9.81), 0, 0.01].map(Math.fround),
+  );
+  assert.deepEqual([...new Uint32Array(words, 16, 3)], [0, 3, 1000], 'first slot, count, capacity');
+  assert.deepEqual([...written(records)], [...pool.staging.subarray(0, 24)]);
+  assert.equal(particles.run([pool], encoder), 0, 'nothing staged, no time: no pass');
+  assert.equal(gpu.buffers.length, 3, 'state, staging and step made once');
+  particles.run([], encoder);
+  assert.equal(gpu.destroyed.length, 3, 'a pool the world let go of gives its buffers back');
+});
+
+test('WebGPU: a step that cannot compile is heard, and its pools stop asking frames', async () => {
+  const heard: unknown[] = [],
+    { device } = fakeDevice({ compute: false });
+  const particles = createWebgpuParticles(device, (error) => heard.push(error));
+  const pool = new ParticlePool({ capacity: 8 });
+  pool.emit(0, 0, 0, 0, 1, 0, 2);
+  await tick();
+  assert.equal(particles.run([pool], computeRecorder().encoder), 0);
+  assert.deepEqual([heard.length, pool.moving, pool.emit(0, 0, 0, 0, 1, 0, 2)], [1, false, false]);
+});
+
+test("WebGPU: a still frame is held until one of the world's pools moves", () => {
+  const rt = settledRt(),
+    { device } = fakeDevice();
+  for (let i = 0; i < 2; i++) {
+    rt.run.frame++;
+    keepWebgpuFrame(rt);
+  }
+  assert.equal(holdWebgpuFrame(rt, device), true, 'still, and no pool: held');
+  const idle = new ParticlePool({ capacity: 8 });
+  Object.assign(rt.context, { particles: [idle] });
+  assert.equal(holdWebgpuFrame(rt, device), true, 'an idle pool changes nothing');
+  idle.emit(0, 0, 0, 0, 1, 0, 2);
+  assert.equal(holdWebgpuFrame(rt, device), false, 'a moving one does');
+  assert.equal(rt.run.frameHeld, false);
+});

packages/sdk-browser/src/particles/webgpuParticles.ts

@@ -0,0 +1,153 @@
+import {
+  PARTICLE_FLOATS,
+  type ParticlePool,
+  type ParticleStep,
+} from '../../../sdk-core/src/fluids/particles.ts';
+import type { WebgpuPagesRuntime } from '../webgpu/pages/runtime.ts';
+import { createCheckedShaderModule } from '../gpu/core/shaderModule.ts';
+import { bounceGroup, bounceLayout } from '../bounce/bindings.ts';
+
+/** The pass label the GPU timings name the particle step by (`passesGpu`). */
+export const PARTICLES_PASS = 'Trillion3D particles';
+/** Slots one workgroup steps. */
+export const PARTICLE_WORKGROUP = 64;
+
+/** One invocation per slot: the record the ring gives it this image replaces it, then a live
+ *  particle moves, its position counted from the pool's origin; a dead one nobody emitted into
+ *  is left as it is. */
+export const PARTICLES_WGSL = /* wgsl */ `
+struct Particle { position: vec4f, velocity: vec4f }
+struct Step { acceleration: vec3f, dt: f32, first: u32, count: u32, capacity: u32, pad: u32 }
+@group(0) @binding(0) var<uniform> step: Step;
+@group(0) @binding(1) var<storage, read> staged: array<Particle>;
+@group(0) @binding(2) var<storage, read_write> particles: array<Particle>;
+@compute @workgroup_size(${PARTICLE_WORKGROUP})
+fn main(@builtin(global_invocation_id) id: vec3u) {
+  let i = id.x;
+  if (i >= step.capacity) { return; }
+  let k = (i + step.capacity - step.first) % step.capacity;
+  var p = particles[i];
+  if (k < step.count) { p = staged[k]; } else if (p.position.w >= p.velocity.w) { return; }
+  if (p.position.w < p.velocity.w) {
+    let velocity = p.velocity.xyz + step.acceleration * step.dt;
+    p.velocity = vec4f(velocity, p.velocity.w);
+    p.position = vec4f(p.position.xyz + velocity * step.dt, p.position.w + step.dt);
+  }
+  particles[i] = p;
+}`;
+
+/** The step uniform, the WGSL `Step`: acceleration and `dt`, then the ring's first slot, count
+ *  and capacity. Made once; `write` rewrites it for one pool's step. */
+export function createStepWords() {
+  const buffer = new ArrayBuffer(32),
+    floats = new Float32Array(buffer),
+    uints = new Uint32Array(buffer);
+  const write = (pool: ParticlePool, { first, count, dt }: Readonly<ParticleStep>) => {
+    floats.set(pool.acceleration);
+    floats[3] = dt;
+    uints[4] = first;
+    uints[5] = count;
+    uints[6] = pool.capacity;
+  };
+  return { buffer, uints, write };
+}
+
+type PoolState = { step: GPUBuffer; staged: GPUBuffer; state: GPUBuffer; group: GPUBindGroup };
+
+/**
+ * The WebGPU particle step: one compute pass, timed under `PARTICLES_PASS`, one dispatch per pool
+ * that has records or time to take. A pool's state is one storage buffer made the first time it
+ * is stepped; its records ride in a staging buffer of the pool's size, written up to the image's
+ * count. The pipeline compiles in the background; until it arrives no pool is taken, so what they
+ * stage waits. `fail` hears a pipeline that could not be made, and every pool is then `refused`.
+ */
+export function createWebgpuParticles(device: GPUDevice, fail: (error: unknown) => void) {
+  const layout = bounceLayout(device, ['uniform', 'read-only-storage', 'storage']);
+  let pipeline: GPUComputePipeline | null | undefined;
+  createCheckedShaderModule(device, PARTICLES_WGSL, 'PARTICLES')
+    .then((module) =>
+      device.createComputePipelineAsync({
+        label: PARTICLES_PASS,
+        layout: device.createPipelineLayout({ bindGroupLayouts: [layout] }),
+        compute: { module, entryPoint: 'main' },
+      }),
+    )
+    .then((made) => (pipeline = made))
+    .catch((error) => ((pipeline = null), fail(error)));
+  const words = createStepWords();
+  const pass: GPUComputePassDescriptor = { label: PARTICLES_PASS },
+    made = new Map<ParticlePool, PoolState>();
+  const buffer = (name: string, size: number, usage: number) =>
+    device.createBuffer({ label: `${PARTICLES_PASS} ${name}`, size, usage });
+  const release = (kept: readonly ParticlePool[]) => {
+    for (const [pool, { step, staged, state }] of made)
+      if (!kept.includes(pool)) {
+        for (const gone of [step, staged, state]) gone.destroy();
+        made.delete(pool);
+      }
+  };
+  const make = (pool: ParticlePool) => {
+    const { STORAGE, UNIFORM, COPY_DST } = GPUBufferUsage;
+    const step = buffer('step', words.buffer.byteLength, UNIFORM | COPY_DST),
+      staged = buffer('staging', pool.staging.byteLength, STORAGE | COPY_DST),
+      state = buffer('state', pool.capacity * PARTICLE_FLOATS * 4, STORAGE);
+    const kept = { step, staged, state, group: bounceGroup(device, layout, [step, staged, state]) };
+    made.set(pool, kept);
+    return kept;
+  };
+  return {
+    /** Steps `pools` in `encoder`; returns the dispatches encoded. */
+    run(pools: readonly ParticlePool[], encoder: GPUCommandEncoder) {
+      if (pipeline === undefined) return 0;
+      let computing: GPUComputePassEncoder | undefined,
+        dispatches = 0,
+        held = 0;
+      for (const pool of pools) {
+        if (made.has(pool)) held++;
+        pool.refused = !pipeline;
+        const step = pool.flush(),
+          { count } = step;
+        if (!pipeline || (!count && !step.dt)) continue;
+        const kept = made.get(pool) ?? make(pool);
+        words.write(pool, step);
+        device.queue.writeBuffer(kept.step, 0, words.buffer);
+        if (count)
+          device.queue.writeBuffer(kept.staged, 0, pool.staging, 0, count * PARTICLE_FLOATS);
+        if (!computing) {
+          computing = encoder.beginComputePass(pass);
+          computing.setPipeline(pipeline);
+        }
+        computing.setBindGroup(0, kept.group);
+        const slots = Math.min(pool.capacity, pool.emitted); // past them, nothing was emitted
+        computing.dispatchWorkgroups(Math.ceil(slots / PARTICLE_WORKGROUP));
+        dispatches++;
+      }
+      computing?.end();
+      if (made.size > held) release(pools); // a pool the world let go of
+      return dispatches;
+    },
+    dispose: () => release([]),
+  };
+}
+
+export type WebgpuParticles = ReturnType<typeof createWebgpuParticles>;
+
+const moving = (pool: ParticlePool) => pool.moving;
+/** True while one of the world's pools moves: the image changes, and is not held. */
+export const particlesMoved = (rt: WebgpuPagesRuntime) => !!rt.context.particles?.some(moving);
+
+/** The world's pools on this image, stepped in the image's command buffer ahead of its
+ *  transparent stage, which draws them (#755). */
+export function encodeParticles(
+  rt: WebgpuPagesRuntime,
+  device: GPUDevice,
+  encoder: GPUCommandEncoder,
+) {
+  const pools = rt.context.particles;
+  // Once made, the step runs with no pool left too: it gives a released pool's buffers back.
+  if (!pools || (!pools.length && !rt.gpu.particles)) return;
+  rt.gpu.particles ??= createWebgpuParticles(device, (error) =>
+    rt.diag.diagnosticFailure('particles-unavailable', error),
+  );
+  rt.run.gpuComputeDispatches += rt.gpu.particles.run(pools, encoder);
+}

packages/sdk-browser/src/stage/mapping.ts

@@ -10,6 +10,7 @@ import { SHADOW_LAYER_PASS } from '../gpu/shadow/staticLayer.ts';
 import { SHADOW_TRANSMITTANCE_PASS } from '../gpu/shadow/transmittance.ts';
 import { LIGHT_CUT_PASS } from '../gpu/dag/encode.ts';
 import { MATERIAL_DEPTH_PASS, MATERIAL_SURFACES_PASS } from '../webgpu/core/materialPasses.ts';
+import { PARTICLES_PASS } from '../particles/webgpuParticles.ts';
 import type { StageAdd } from './profiler.ts';
 
 /**
@@ -75,6 +76,7 @@ const PASSES: Readonly<Record<string, PassRow>> = Object.freeze({
   [LIGHT_TILES_PASS]: ['lightLists', 'other'],
   [BOUNCE_SURFACE_PASS]: ['bounce', 'other'],
   [BOUNCE_PROBE_PASS]: ['bounce', 'other'],
+  [PARTICLES_PASS]: ['physics', 'other'],
   [DEFERRED_LIGHTING_PASS]: ['lighting', 'other'],
   [TAA_PASS]: ['antialiasing', 'other'],
   'Trillion3D HDR composition': ['present', 'other'],

packages/sdk-browser/src/webgpu/frame/hold.ts

@@ -5,6 +5,7 @@ import type { WebgpuPagesRuntime } from '../pages/runtime.ts';
 import { shadowsUnsettled } from '../pages/state/lights.ts';
 import { effectsMoved } from '../pages/render/encodeEffects.ts';
 import { guidesMoved } from '../pages/render/encodeGuides.ts';
+import { particlesMoved } from '../../particles/webgpuParticles.ts';
 import { frameTargetsAwaited } from '../pages/prepare/targetGrant.ts';
 import { deviceAnswering } from './deviceAnswer.ts';
 
@@ -154,7 +155,7 @@ export function holdWebgpuFrame(rt: WebgpuPagesRuntime, device: GPUDevice) {
     beginTaaFrame(rt, run.gate.cam, quiet);
     // Guides or an effect chain the page changed, or a chain the last image lacked while its
     // programs compiled, are drawn by a full image; the accumulation stays still for it.
-    if (!quiet || !taaSettled(rt) || guidesMoved(rt) || effectsMoved(rt)) {
+    if (!quiet || !taaSettled(rt) || guidesMoved(rt) || effectsMoved(rt) || particlesMoved(rt)) {
       run.frameHeld = false;
       return false;
     }

packages/sdk-browser/src/webgpu/pages/io/metrics.ts

@@ -149,6 +149,8 @@ export function disposeWebgpuPages(rt: WebgpuPagesRuntime) {
   gpu.effects = undefined;
   gpu.guides?.dispose();
   gpu.guides = undefined;
+  gpu.particles?.dispose();
+  gpu.particles = undefined;
   rt.lights.tiles?.dispose();
   rt.lights.shadows?.dispose();
   rt.lights.pageRequests?.dispose();

packages/sdk-browser/src/webgpu/pages/render/encodeBlend.ts

@@ -7,6 +7,7 @@ import { orderBlendPasses, orderVisibleBlend } from '../../blend/order.ts';
 import { drawFallbackBlendPass, writeFallbackBlendUniforms } from '../../blend/fallback.ts';
 import { encodeTransparentInstances } from '../../transparent/draw.ts';
 import { encodeWaterPass } from '../../water/pass.ts';
+import { encodeParticles } from '../../../particles/webgpuParticles.ts';
 import { blendLightResources } from '../../blend/lighting.ts';
 import { voidStaleBlendGroups } from '../../blend/identity.ts';
 import { viewProj } from '../helpers.ts';
@@ -32,6 +33,8 @@ export function encodeBlend(
   uniformBase: number,
 ) {
   const { gpu, vis, run, timing, blendState, diag } = rt;
+  // Every image path reaches this stage: the particles step here, beside the water.
+  encodeParticles(rt, device, encoder);
   if (
     !gpu.pipelineBlend ||
     !blendState.blendGpu.length ||

packages/sdk-browser/src/webgpu/pages/state/gpu.ts

@@ -12,6 +12,7 @@ import type { WebgpuEffects } from '../../../effects/webgpuEffects.ts';
 import { UNIFORM_STRIDE } from '../../blend/uniforms.ts';
 import type { ModePipelines } from '../../blend/stagePipelines.ts';
 import type { WebgpuGuidePass } from '../../../guides/guidePass.ts';
+import type { WebgpuParticles } from '../../../particles/webgpuParticles.ts';
 import type { DeviceGrant } from '../../../gpu/core/errorScope.ts';
 
 /** GPU resources of the forward path: page cache, pipelines, frame targets and presentation. */
@@ -89,6 +90,8 @@ export interface WebgpuGpuState {
   guides: WebgpuGuidePass | undefined;
   /** Revision of the page's guides the last encoded image drew (`encodeGuides.ts`). */
   guideRevision: number;
+  /** The particle step, made by the first image with a pool (`../../../particles/`). */
+  particles: WebgpuParticles | undefined;
 }
 
 /** The frozen colour the water composite rereads, and the depth its surface stage tests and
@@ -149,5 +152,6 @@ export function createWebgpuGpuState(viewport: readonly [number, number]): Webgp
     effectsRevision: 0,
     guides: undefined,
     guideRevision: 0,
+    particles: undefined,
   };
 }

packages/sdk-browser/src/world/core/world.ts

@@ -192,7 +192,7 @@ export function createWorld(target: WorldTarget, options: WorldOptions = {}) {
     },
   };
   frames.add(noticeEffectBudget(world.budget, canvas, world.effects, diagnostic.notices));
-  registerWorld(world, { session: () => runtime.explorer, last: () => frames.last });
+  registerWorld(world, { session: () => runtime.explorer, last: () => frames.last }, switches.held);
   return world;
 }