diff --git a/docs/ENGINE.md b/docs/ENGINE.md index 35658f2508..458c5603fd 100644 --- a/docs/ENGINE.md +++ b/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 diff --git a/packages/sdk-browser/src/backend/types.ts b/packages/sdk-browser/src/backend/types.ts index af6ab13958..de0b094f60 100644 --- a/packages/sdk-browser/src/backend/types.ts +++ b/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()`). */ diff --git a/packages/sdk-browser/src/gpu/core/engineShaders.fixture.ts b/packages/sdk-browser/src/gpu/core/engineShaders.fixture.ts index ec054e97ca..6dee7fea23 100644 --- a/packages/sdk-browser/src/gpu/core/engineShaders.fixture.ts +++ b/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) => Object.fromEntries(Object.entries(sources).map(([input, code]) => [`${label}_${input}`, code])); @@ -88,4 +89,5 @@ export const ENGINE_SHADERS: Record = { REDUCE_WGSL, TRANSPARENT_COMPACT_SHADER, WATER_COMPOSITE_SHADER, + PARTICLES_WGSL, }; diff --git a/packages/sdk-browser/src/measurement/measurement.ts b/packages/sdk-browser/src/measurement/measurement.ts index 20008e57e8..6999f867ef 100644 --- a/packages/sdk-browser/src/measurement/measurement.ts +++ b/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( diff --git a/packages/sdk-browser/src/particles/particles.test.ts b/packages/sdk-browser/src/particles/particles.test.ts new file mode 100644 index 0000000000..522cfd295a --- /dev/null +++ b/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); +}); diff --git a/packages/sdk-browser/src/particles/webgpuParticles.ts b/packages/sdk-browser/src/particles/webgpuParticles.ts new file mode 100644 index 0000000000..7b4bdef99a --- /dev/null +++ b/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 step: Step; +@group(0) @binding(1) var staged: array; +@group(0) @binding(2) var particles: array; +@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) => { + 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(); + 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; + +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); +} diff --git a/packages/sdk-browser/src/stage/mapping.ts b/packages/sdk-browser/src/stage/mapping.ts index aa3834852c..525d5873af 100644 --- a/packages/sdk-browser/src/stage/mapping.ts +++ b/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> = 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'], diff --git a/packages/sdk-browser/src/webgpu/frame/hold.ts b/packages/sdk-browser/src/webgpu/frame/hold.ts index c86d308aa6..b8bad9faab 100644 --- a/packages/sdk-browser/src/webgpu/frame/hold.ts +++ b/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; } diff --git a/packages/sdk-browser/src/webgpu/pages/io/metrics.ts b/packages/sdk-browser/src/webgpu/pages/io/metrics.ts index ea6230b64f..483f5a19e9 100644 --- a/packages/sdk-browser/src/webgpu/pages/io/metrics.ts +++ b/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(); diff --git a/packages/sdk-browser/src/webgpu/pages/render/encodeBlend.ts b/packages/sdk-browser/src/webgpu/pages/render/encodeBlend.ts index b5fee87be7..accd1f07d8 100644 --- a/packages/sdk-browser/src/webgpu/pages/render/encodeBlend.ts +++ b/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 || diff --git a/packages/sdk-browser/src/webgpu/pages/state/gpu.ts b/packages/sdk-browser/src/webgpu/pages/state/gpu.ts index 91f94a7fdf..2769143e36 100644 --- a/packages/sdk-browser/src/webgpu/pages/state/gpu.ts +++ b/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, }; } diff --git a/packages/sdk-browser/src/world/core/world.ts b/packages/sdk-browser/src/world/core/world.ts index 295f8628c7..cfb976aaa8 100644 --- a/packages/sdk-browser/src/world/core/world.ts +++ b/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; } diff --git a/packages/sdk-browser/src/world/core/worldSession.test.ts b/packages/sdk-browser/src/world/core/worldSession.test.ts index 79e1d20970..412f1aabb0 100644 --- a/packages/sdk-browser/src/world/core/worldSession.test.ts +++ b/packages/sdk-browser/src/world/core/worldSession.test.ts @@ -1,7 +1,9 @@ import test from 'node:test'; import assert from 'node:assert/strict'; -import { awaitViewPages } from './worldSession.ts'; +import { attachParticles, awaitViewPages, registerWorld } from './worldSession.ts'; import type { MeasuredWorld } from '../session/explorer.ts'; +import type { BeforeFrameInfo, World } from './world.ts'; +import { ParticlePool } from '../../../../sdk-core/src/fluids/particles.ts'; test('world.awaitPages waits for the scene, then for pages alone, never an image, heard (#408)', async () => { const steps: unknown[] = []; @@ -13,3 +15,27 @@ test('world.awaitPages waits for the scene, then for pages alone, never an image await awaitViewPages(runtime, () => session, onProgress); assert.deepEqual(steps, ['settled', { image: false, onProgress }]); }); + +test('attachParticles gives the pool to every session and the frames the world draws (#420)', () => { + type Hook = (frame: BeforeFrameInfo) => void; + const hooks = new Set(); + let asked = 0; + const add = (hook: Hook) => (hooks.add(hook), () => hooks.delete(hook)); + const world = { beforeFrame: add, onFrame: add, invalidate: () => asked++ } as unknown as World; + const held = { particles: [] as ParticlePool[] }; + registerWorld(world, { session: () => null, last: () => null }, held); + const pool = new ParticlePool({ capacity: 8 }); + const remove = attachParticles(world, pool); + assert.deepEqual([held.particles, asked], [[pool], 1], 'held, and a frame asked'); + for (const again of [world, {} as World]) + assert.throws(() => attachParticles(again, pool), /^Error: PARTICLES_ATTACH/); + const frame = () => hooks.forEach((hook) => hook({ delta: 0.02, time: 0.02 })); + frame(); + assert.equal(asked, 1, 'an idle pool asks for no frame'); + pool.emit(0, 0, 0, 0, 1, 0, 2); + frame(); + assert.equal(asked, 2, 'a moving one asks for the next'); + assert.equal(pool.flush().dt, 0.02, "the world's frame time is the step's, the idle one not"); + remove(); + assert.deepEqual([held.particles, hooks.size, asked], [[], 0, 3], 'and an image to release it'); +}); diff --git a/packages/sdk-browser/src/world/core/worldSession.ts b/packages/sdk-browser/src/world/core/worldSession.ts index 76348ca4ed..bc6446debe 100644 --- a/packages/sdk-browser/src/world/core/worldSession.ts +++ b/packages/sdk-browser/src/world/core/worldSession.ts @@ -1,13 +1,17 @@ import type { MeasuredWorld } from '../session/explorer.ts'; import type { FrameMetrics, JobProgress } from '../../../../sdk-core/src/index.ts'; +import type { ParticlePool } from '../../../../sdk-core/src/fluids/particles.ts'; +import type { World } from './world.ts'; export type { JobProgress }; /** What the families that read a world's engine reach it by, without the page holding it. */ type Access = { session: () => MeasuredWorld | null; last: () => FrameMetrics | null }; -const worlds = new WeakMap(); +/** What the world holds for every session it opens (`worldSwitches.ts`). */ +type Held = { particles: ParticlePool[] }; +const worlds = new WeakMap(); -export const registerWorld = (world: object, access: Access) => { - worlds.set(world, access); +export const registerWorld = (world: object, access: Access, held: Held) => { + worlds.set(world, { ...access, particles: held.particles }); }; /** The session drawing `world` now; a world that draws nothing yet is refused by name. */ @@ -18,6 +22,26 @@ export function sessionOf(world: object): MeasuredWorld { return session; } +/** Steps `pool` on the GPU at every frame `world` draws, its time advanced by the world's loop, + * which draws on while the pool moves: the measurement entry's way in (#420) until particles + * have a public face (#423). Returns the remover. A pool is attached once, to a made world. */ +export function attachParticles(world: World, pool: ParticlePool) { + const pools = worlds.get(world)?.particles; + if (!pools || pools.includes(pool)) throw new Error('PARTICLES_ATTACH: attached, or no world'); + pools.push(pool); + const stops = [ + world.beforeFrame(({ delta }) => pool.advance(delta)), + world.onFrame(() => pool.moving && world.invalidate()), // after every hook's emission + ]; + world.invalidate(); + return () => { + for (const stop of stops) stop(); + const at = pools.indexOf(pool); + if (at >= 0) pools.splice(at, 1); + world.invalidate(); // the next image gives the pool's buffers back + }; +} + /** The metrics of the last frame `world` drew, null before its first. */ export const lastFrameOf = (world: object) => worlds.get(world)?.last() ?? null; diff --git a/packages/sdk-browser/src/world/core/worldSwitches.ts b/packages/sdk-browser/src/world/core/worldSwitches.ts index 2501633aaa..65d429052d 100644 --- a/packages/sdk-browser/src/world/core/worldSwitches.ts +++ b/packages/sdk-browser/src/world/core/worldSwitches.ts @@ -3,6 +3,7 @@ import type { WorldRenderer } from '../capability/worldReady.ts'; import type { WorldOptions } from './worldOptions.ts'; import { EffectChain } from '../../../../sdk-core/src/world/effect/chain.ts'; import { createGuideSet, type Guides } from '../../guides/guideSet.ts'; +import type { ParticlePool } from '../../../../sdk-core/src/fluids/particles.ts'; /** What of the world's runtime the switches reach: its open session, and its reopening. */ interface SwitchedRuntime { @@ -30,6 +31,8 @@ export function worldSwitches( // One chain for the world's life: every session draws it, a change asks for a frame. effects: new EffectChain(invalidate), guides: createGuideSet(invalidate), + // The particle pools the measurement entry attaches (`attachParticles`); none by default. + particles: [] as ParticlePool[], }; return { held, diff --git a/packages/sdk-browser/src/world/render/compose.test.ts b/packages/sdk-browser/src/world/render/compose.test.ts index b1dbadfdc3..dbacd7990b 100644 --- a/packages/sdk-browser/src/world/render/compose.test.ts +++ b/packages/sdk-browser/src/world/render/compose.test.ts @@ -10,6 +10,7 @@ import type { HostDrawOutput, RenderBackend } from '../../backend/types.ts'; import { createFrameComposer } from './compose.ts'; import { createWebglRenderTarget } from '../../webgl/core/renderTarget.ts'; import { createTestContext } from '../../webgl/core/testContext.fixture.ts'; +import { ParticlePool } from '../../../../sdk-core/src/fluids/particles.ts'; const camera = G.perspectiveCamera(); @@ -118,3 +119,12 @@ test('an engine that draws nothing on the host surface is refused by name', () = const backend = { id: 'mute', scene: {} } as unknown as RenderBackend; assert.throws(() => compose(backend, null), /HOST_DRAW_UNSUPPORTED:mute/); }); + +test('a world with a particle pool is refused by name on WebGL2, never drawn without it', () => { + const { gl } = createTestContext(); + const particles = [new ParticlePool({ capacity: 8 })]; + const compose = createFrameComposer(gl, camera, { particles }); + const { backend, outputs } = engine(); + assert.throws(() => compose(backend, null), /^Error: PARTICLES_UNSUPPORTED/); + assert.deepEqual([outputs.length, particles[0].moving], [0, false], 'refused, it asks no frame'); +}); diff --git a/packages/sdk-browser/src/world/render/compose.ts b/packages/sdk-browser/src/world/render/compose.ts index 0f717c09f2..9d3a19e5c1 100644 --- a/packages/sdk-browser/src/world/render/compose.ts +++ b/packages/sdk-browser/src/world/render/compose.ts @@ -17,6 +17,7 @@ import { type WebglRenderTarget, } from '../../webgl/core/renderTarget.ts'; import { createWebglEffects, type WebglEffectOutput } from '../../effects/webglEffects.ts'; +import type { ParticlePool } from '../../../../sdk-core/src/fluids/particles.ts'; const NONE: readonly EffectPass[] = []; @@ -35,16 +36,17 @@ export type ComposedChain = { chain: EffectChain; shown: () => boolean }; * since it was kept is drawn again. The page's `guides` are drawn over the image the destination * got, the chain's included, before that copy is kept, at the host's `pixelRatio`; a change to * them spares no redraw. + * No engine drawn here steps particles yet (#759): a world with a pool is refused by name. * Nothing here belongs to a rendering library. */ export function createFrameComposer( gl: WebGL2RenderingContext, camera: HostCamera, - layers: { effects?: ComposedChain } & ( + layers: { effects?: ComposedChain; particles?: readonly ParticlePool[] } & ( { guides?: undefined } | { guides: GuideSet; pixelRatio: () => number } ) = {}, ) { - const { effects: composed } = layers; + const { effects: composed, particles = [] } = layers; const heldFrame = createHeldFrame(gl); const guideDraw = createWebglGuideDraw(gl); let guidesDrawn = layers.guides?.revision ?? 0; @@ -106,6 +108,9 @@ export function createFrameComposer( ) => { const { width, height } = bindWebglTarget(gl, target); if (present(backend)) return; + for (const pool of particles) pool.refused = true; // it asks no frame of its own + if (particles.length) + throw new Error('PARTICLES_UNSUPPORTED: WebGL2 does not step particle pools yet (#759)'); const revision = composed?.chain.revision ?? 0; const guidesHeld = !layers.guides || layers.guides.revision === guidesDrawn; if ( diff --git a/packages/sdk-browser/src/world/render/hostState.ts b/packages/sdk-browser/src/world/render/hostState.ts index d1354e2b55..7a1532f226 100644 --- a/packages/sdk-browser/src/world/render/hostState.ts +++ b/packages/sdk-browser/src/world/render/hostState.ts @@ -93,6 +93,7 @@ export function createExplorerHostState( }, guides: options.guides, pixelRatio: () => pixelRatioOf(options), + particles: options.particles, }), compositor: createComparisonCompositor(gl), } diff --git a/packages/sdk-browser/src/world/session/backends.ts b/packages/sdk-browser/src/world/session/backends.ts index 0293dcc8ce..d77a6cc9c6 100644 --- a/packages/sdk-browser/src/world/session/backends.ts +++ b/packages/sdk-browser/src/world/session/backends.ts @@ -113,6 +113,7 @@ export async function prepareExplorerBackends(session: ExplorerSession, inputs: shadowPageInvalidation: options.shadowPageInvalidation, sceneLighting: sceneLightingSource, guides: options.guides, + particles: options.particles, // Bounced light stays off unless asked: its step holds 1.1 to 1.3 ms on Emerald, above 1 ms. bounce: options.bounce, bounceBudgetMs: options.bounceBudgetMs, diff --git a/packages/sdk-browser/src/world/session/options.ts b/packages/sdk-browser/src/world/session/options.ts index 490ab9c90c..39e5a1b978 100644 --- a/packages/sdk-browser/src/world/session/options.ts +++ b/packages/sdk-browser/src/world/session/options.ts @@ -127,6 +127,8 @@ export interface MeasuredWorldOptions { sceneLighting?: Object3D; /** Lines and points drawn over the image, held by the world (`world.guides`). */ guides?: import('../../guides/guideSet.ts').GuideSet; + /** Particle pools, held by the world (`attachParticles`). */ + particles?: readonly import('../../../../sdk-core/src/fluids/particles.ts').ParticlePool[]; /** Bounced light. Off by default; `true` turns it on for the whole session. */ bounce?: boolean; /** Target duration of the "Bounce" step per frame, in milliseconds. 0.8 ms by default. */ diff --git a/packages/sdk-core/src/fluids/particles.test.ts b/packages/sdk-core/src/fluids/particles.test.ts new file mode 100644 index 0000000000..36d5dce690 --- /dev/null +++ b/packages/sdk-core/src/fluids/particles.test.ts @@ -0,0 +1,74 @@ +import test from 'node:test'; +import assert from 'node:assert/strict'; +import { PARTICLE_FLOATS, ParticlePool } from './particles.ts'; + +const emitMany = (pool: ParticlePool, n: number) => { + for (let i = 0; i < n; i++) pool.emit(i, 0, 0, 0, 1, 0, 2); +}; + +test('a record is the particle as the GPU keeps it: position, age zero, velocity, lifetime', () => { + const pool = new ParticlePool({ capacity: 8, emitPerFrame: 4 }); + assert.equal(pool.emit(1, 2, 3, 4, 5, 6, 7), true); + assert.deepEqual([...pool.staging.subarray(0, PARTICLE_FLOATS)], [1, 2, 3, 0, 4, 5, 6, 7]); + pool.advance(0.01); + assert.deepEqual({ ...pool.flush() }, { first: 0, count: 1, dt: 0.01 }); +}); + +test('the ring wraps past capacity onto the oldest slots, and allocates nothing after creation', () => { + const pool = new ParticlePool({ capacity: 8, emitPerFrame: 4 }); + const staging = pool.staging, + step = pool.flush(), + slots = []; + for (const n of [4, 4, 3, 2]) { + emitMany(pool, n); + assert.equal(pool.flush(), step, 'the same step object every image'); + slots.push(`${step.first}+${step.count}`); + } + // Slots 0..3, 4..7, then 0..2 overwrite the first image's particles, then 3..4. + assert.deepEqual(slots, ['0+4', '4+4', '0+3', '3+2']); + assert.equal(pool.staging, staging, 'the same staging'); + assert.equal(pool.emitted, 13); +}); + +test('a full staging refuses and counts; the step clamps its time and consumes it', () => { + const pool = new ParticlePool({ capacity: 8, emitPerFrame: 2 }); + emitMany(pool, 3); + assert.equal(pool.dropped, 1); + pool.advance(1); + assert.deepEqual({ ...pool.flush() }, { first: 0, count: 2, dt: 1 / 15 }); + assert.deepEqual({ ...pool.flush() }, { first: 2, count: 0, dt: 0 }, 'nothing twice'); +}); + +test('a pool out of bounds is refused by name', () => { + assert.throws(() => new ParticlePool({ capacity: 0 }), /^Error: PARTICLE_CAPACITY/); + assert.equal(new ParticlePool({ capacity: 64 }).emitPerFrame, 64, 'the default, at most it'); + assert.throws( + () => new ParticlePool({ capacity: 4, emitPerFrame: 5 }), + /^Error: PARTICLE_EMISSION/, + ); +}); + +test('a particle ten kilometres out keeps its sub-millimetre steps: positions are from the origin', () => { + const far = 10_000, + pool = new ParticlePool({ capacity: 8, origin: [far, 0, far] }); + pool.emit(far + 0.5, 2, far, 0.024, 0, 0, 4); + assert.deepEqual([...pool.staging.subarray(0, 3)], [0.5, 2, 0]); + // A drift of 0.4 mm in one step of 1/60 s, in the GPU step's 32-bit arithmetic: x += vx × dt. + const step = (x: number) => Math.fround(x + Math.fround(0.024 * Math.fround(1 / 60))); + assert.ok(step(pool.staging[0]) > pool.staging[0], 'from the origin, the step is kept'); + assert.equal(step(Math.fround(far + 0.5)), Math.fround(far + 0.5), 'in world floats, lost'); +}); + +test('an idle pool stops moving: its step takes no time once its last particle is dead', () => { + const pool = new ParticlePool({ capacity: 8 }); + assert.equal(pool.moving, false, 'nothing emitted'); + pool.emit(0, 0, 0, 0, 0, 0, 0.1); + assert.equal(pool.moving, true); + for (const dt of [0.05, 0.05, 0.05, 0.05]) { + pool.advance(dt); + assert.equal(pool.flush().dt, dt); + } + assert.equal(pool.moving, false, 'the lifetime has run out'); + pool.advance(0.05); + assert.equal(pool.flush().dt, 0); +}); diff --git a/packages/sdk-core/src/fluids/particles.ts b/packages/sdk-core/src/fluids/particles.ts new file mode 100644 index 0000000000..f93750a073 --- /dev/null +++ b/packages/sdk-core/src/fluids/particles.ts @@ -0,0 +1,116 @@ +/** + * The engine's particle pool (#420), stepped on the GPU (`sdk-browser/src/particles/`): a fixed + * capacity, a ring where emission writes the oldest slot, the image's records staged in one buffer + * made at creation. Nothing is compacted or read back: a particle past its lifetime is dead, and + * the GPU skips it. Positions are from the pool's origin, its emitter's place: ten kilometres out a + * particle still moves by a fraction of a millimetre, which 32-bit world floats round away. The + * step never reads the origin; drawing adds it back (#755). + */ +import { GRAVITY_PRESETS } from '../physics/options.ts'; + +/** Floats of one particle and of one emission record, the same eight words: position from the + * origin then age, velocity then lifetime. A record's age is zero: the GPU copies it as it is. */ +export const PARTICLE_FLOATS = 8; +/** The largest pool: 32 MiB of state on WebGPU, a 1024 × 1024 texture pair on WebGL2. */ +const MAX_CAPACITY = 1 << 20; +/** The longest step an image takes, seconds: a stalled tab does not fling its particles away. */ +const MAX_STEP = 1 / 15; + +/** How a pool is made; the capacity is fixed for its life. */ +export interface ParticlePoolSpec { + /** Particles the pool holds; emission past it overwrites the oldest. */ + capacity: number; + /** Records one image may stage: a sixty-fourth of the capacity by default, 256 to capacity. */ + emitPerFrame?: number; + /** Metres per second squared on every live particle; gravity by default. */ + acceleration?: readonly [number, number, number]; + /** The emitter's place in the world, metres, fixed for the pool's life; the origin by default. */ + origin?: readonly [number, number, number]; +} + +/** One image's step: `count` records land from ring slot `first`, then live ones move `dt` s. */ +export type ParticleStep = { first: number; count: number; dt: number }; + +export class ParticlePool { + readonly capacity: number; + readonly emitPerFrame: number; + readonly acceleration: Float32Array; + /** The origin in double precision: emission subtracts it before rounding to 32 bits. */ + readonly origin: Float64Array; + /** The staged records, `emitPerFrame` of them, read by the renderer up to `step.count`. */ + readonly staging: Float32Array; + /** Records staged since creation, and those refused: the image's staging full, or the pool + * `refused` by a renderer that cannot step it, which it then no longer asks frames for. */ + emitted = 0; + dropped = 0; + refused = false; + private readonly step: ParticleStep = { first: 0, count: 0, dt: 0 }; + private cursor = 0; + private staged = 0; + private pending = 0; + /** Seconds the longest-lived particle may have, one step over: the GPU's 32-bit age lags. */ + private liveFor = 0; + + constructor({ capacity, emitPerFrame, acceleration, origin }: ParticlePoolSpec) { + const perFrame = emitPerFrame ?? Math.min(capacity, Math.max(256, capacity >> 6)); + if (!Number.isInteger(capacity) || capacity < 1 || capacity > MAX_CAPACITY) + throw new Error(`PARTICLE_CAPACITY: a pool holds 1 to ${MAX_CAPACITY} particles`); + if (!Number.isInteger(perFrame) || perFrame < 1 || perFrame > capacity) + throw new Error('PARTICLE_EMISSION: a pool stages 1 to `capacity` records a frame'); + this.capacity = capacity; + this.emitPerFrame = perFrame; + this.acceleration = Float32Array.from(acceleration ?? [0, -GRAVITY_PRESETS.earth, 0]); + this.origin = Float64Array.from(origin ?? [0, 0, 0]); + this.staging = new Float32Array(perFrame * PARTICLE_FLOATS); + } + + /** Stages one particle at world position `x, y, z`, born at the next image; false, and + * counted, when the image's staging is full. */ + emit(x: number, y: number, z: number, vx: number, vy: number, vz: number, lifetime: number) { + if (this.refused || this.staged === this.emitPerFrame) { + this.dropped++; + return false; + } + const at = this.staged++ * PARTICLE_FLOATS, + words = this.staging, + origin = this.origin; + words[at] = x - origin[0]; + words[at + 1] = y - origin[1]; + words[at + 2] = z - origin[2]; + words[at + 3] = 0; + words[at + 4] = vx; + words[at + 5] = vy; + words[at + 6] = vz; + words[at + 7] = lifetime; + this.liveFor = Math.max(this.liveFor, lifetime + MAX_STEP); + this.emitted++; + return true; + } + + /** Whether the next step changes anything: a record staged, or a particle still alive. An idle + * pool neither dispatches nor keeps the image from being held. */ + get moving() { + return !this.refused && (this.staged > 0 || this.liveFor > 0); + } + + /** Adds `seconds` to the time the next image steps; an idle pool lets them pass untaken, so + * a held image's time never flings the next newborn particles. */ + advance(seconds: number) { + if (this.moving) this.pending += seconds; + } + + /** The step of the image being encoded, always the same object: the staged records take the + * ring's next slots, and the time advanced since the last step, clamped, is consumed; an idle + * pool's step takes no time. */ + flush(): Readonly { + const step = this.step; + step.first = this.cursor; + step.count = this.staged; + step.dt = Math.min(this.pending, MAX_STEP); + this.liveFor -= step.dt; + this.cursor = (this.cursor + this.staged) % this.capacity; + this.staged = 0; + this.pending = 0; + return step; + } +} diff --git a/tests/browser/probes/particles-step-gpu.ts b/tests/browser/probes/particles-step-gpu.ts new file mode 100644 index 0000000000..74d137443d --- /dev/null +++ b/tests/browser/probes/particles-step-gpu.ts @@ -0,0 +1,75 @@ +// The engine's particle step (#420), actually run in Chromium WebGPU: the committed WGSL compiles, +// and one step of a pool ten kilometres from the world origin moves a newborn particle by its +// 0.4 mm drift, leaves one born dead and a slot nobody emitted into as they are. The words are the +// engine's: the pool stages the records, `createStepWords` writes the uniform; only this proof +// reads the state back. +// +// node --experimental-strip-types tests/browser/probes/particles-step-gpu.ts +import assert from 'node:assert/strict'; +import { dansPageWebgpu } from './pageWebgpu.ts'; +import { ParticlePool } from '../../../packages/sdk-core/src/fluids/particles.ts'; +import * as step from '../../../packages/sdk-browser/src/particles/webgpuParticles.ts'; + +type Args = { shader: string; words: number[]; staged: number[]; bytes: number; groups: number }; + +/** Run in the page: the step's three bindings, one dispatch, the state read back. */ +async function executer({ shader, words, staged, bytes, groups }: Args) { + const appareil = await globalThis.ouvrirAppareil(); + if (!appareil) return { indisponible: 'no WebGPU adapter' }; + const { device, erreurs } = appareil; + const { module, compilation } = await appareil.compile(shader); + if (compilation.length) return { compilation, erreurs }; + const pipeline = device.createComputePipeline({ layout: 'auto', compute: { module } }); + const { UNIFORM, STORAGE, COPY_DST, COPY_SRC, MAP_READ } = GPUBufferUsage; + const buffer = (size: number, usage: number) => device.createBuffer({ size, usage }); + const [uniform, records, state, read] = [ + buffer(32, UNIFORM | COPY_DST), + buffer(staged.length * 4, STORAGE | COPY_DST), + buffer(bytes, STORAGE | COPY_SRC), + buffer(bytes, MAP_READ | COPY_DST), + ]; + device.queue.writeBuffer(uniform, 0, new Uint32Array(words)); + device.queue.writeBuffer(records, 0, new Float32Array(staged)); + const entries = [uniform, records, state].map((b, binding) => ({ + binding, + resource: { buffer: b }, + })); + const encoder = device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries })); + pass.dispatchWorkgroups(groups); + pass.end(); + encoder.copyBufferToBuffer(state, 0, read, 0, bytes); + device.queue.submit([encoder.finish()]); + await read.mapAsync(GPUMapMode.READ); + const particles = Array.from(new Float32Array(read.getMappedRange(), 0, 24)); + read.unmap(); + return { adaptateur: (await appareil.fermer()).court, particles, erreurs }; +} + +const far = 10_000, + dt = 1 / 60, + pool = new ParticlePool({ capacity: 100, emitPerFrame: 4, origin: [far, 0, far] }); +pool.emit(far + 0.5, 2, far, 0.024, 1, 0, 4); // drifts 0.4 mm along x in one step +pool.emit(far, 3, far, 5, 5, 5, 0); // born dead: kept as staged, never moved +pool.advance(dt); +const words = step.createStepWords(); +words.write(pool, pool.flush()); +const resultat = await dansPageWebgpu(executer, { + shader: step.PARTICLES_WGSL, + words: Array.from(words.uints), + staged: Array.from(pool.staging), + bytes: pool.capacity * 32, + groups: Math.ceil(pool.capacity / step.PARTICLE_WORKGROUP), +}); +console.log(JSON.stringify(resultat, null, 2)); +assert.equal(resultat.indisponible ?? null, null, String(resultat.indisponible)); +assert.deepEqual([resultat.compilation ?? [], resultat.erreurs], [[], []]); +const particles = resultat.particles!, + [x, y, , age] = particles; +assert.ok(Math.abs(x - (0.5 + 0.024 * dt)) < 1e-6 && x > 0.5, `drifted 0.4 mm: x = ${x}`); +assert.ok(y > 2 && age === Math.fround(dt), 'rose, and aged by the step'); +assert.deepEqual(particles.slice(8, 16), [0, 3, 0, 0, 5, 5, 5, 0], 'born dead: as staged'); +assert.deepEqual(particles.slice(16, 24), [0, 0, 0, 0, 0, 0, 0, 0], 'never emitted: untouched'); +console.log('OK: the particle step actually run — see particles-step-gpu.ts');