From 000bfd63bfec639e654c0396036c1d4f221e2d73 Mon Sep 17 00:00:00 2001 From: Amp Date: Thu, 20 Aug 2026 11:48:06 +0000 Subject: [PATCH] Optimize render graph and add GPU profiling Amp-Thread-ID: https://ampcode.com/threads/T-01a01380-b478-77d0-84a0-102880a5c5ae Co-authored-by: Heaust Azure --- addons/handles/src/ComputePass.ts | 3 + addons/handles/src/Scene.ts | 349 ++++++++++++++++--- addons/handles/src/importers/gltf.ts | 33 +- addons/handles/src/importers/worker.ts | 27 +- addons/handles/src/materials/PBRMaterial.ts | 9 +- addons/handles/src/materials/Texture.ts | 44 ++- core/Cargo.toml | 2 +- core/index.js | 34 +- core/lib.rs | 29 ++ core/render_graph/gpu_resource.rs | 347 ++++++++++++++++--- core/render_graph/render_graph.rs | 103 +++++- core/render_graph/store.rs | 48 ++- core/renderer/render.rs | 144 +++++++- core/renderer/wgpu.rs | 8 +- core/worker.js | 15 + docs/.vitepress/Playground.vue | 362 +++++++++++++++++--- docs/.vitepress/playgrounds.js | 31 +- docs/guide/core.md | 17 +- docs/guide/materials.md | 6 +- docs/playground.md | 8 +- package-lock.json | 203 +++++++++++ package.json | 3 + 22 files changed, 1622 insertions(+), 203 deletions(-) diff --git a/addons/handles/src/ComputePass.ts b/addons/handles/src/ComputePass.ts index fc7c727..ebefc86 100644 --- a/addons/handles/src/ComputePass.ts +++ b/addons/handles/src/ComputePass.ts @@ -4,6 +4,7 @@ export type GraphBuffer = { id: string; array: string; usage?: string[]; + sync?: "frame" | "loadout"; }; export type GraphTexture = { @@ -18,6 +19,8 @@ export type GraphSampler = { id: string; magFilter?: "nearest" | "linear"; minFilter?: "nearest" | "linear"; + addressModeU?: "clamp-to-edge" | "repeat" | "mirror-repeat"; + addressModeV?: "clamp-to-edge" | "repeat" | "mirror-repeat"; }; export type GraphBinding = { diff --git a/addons/handles/src/Scene.ts b/addons/handles/src/Scene.ts index 5a0d227..6fc00e6 100644 --- a/addons/handles/src/Scene.ts +++ b/addons/handles/src/Scene.ts @@ -91,7 +91,7 @@ fn main(@builtin(global_invocation_id) id: vec3) { clusters[3] = bitcast(light.b); }`; -const forwardShader = /* wgsl */ ` +const basicForwardShader = /* wgsl */ ` struct Accent { color: vec4 } struct VertexOutput { @builtin(position) position: vec4, @@ -153,6 +153,103 @@ fn fragment(input: VertexOutput) -> @location(0) vec4 { return vec4(color, base.a); }`; +function pbrShader(mask: number) { + const baseTexture = mask & 1; + const materialTexture = mask & 2; + const normalTexture = mask & 4; + return /* wgsl */ ` +struct VertexOutput { + @builtin(position) position: vec4, + @location(0) world: vec3, + @location(1) normal: vec3, + @location(2) uv: vec2, + @location(3) tangent: vec4, + @location(4) @interpolate(flat) mesh: u32, + @location(5) @interpolate(flat) material: u32, +} + +@group(0) @binding(0) var clusters: array; +@group(0) @binding(1) var accent: vec4; +@group(0) @binding(2) var positions: array>; +@group(0) @binding(3) var quaternions: array>; +@group(0) @binding(4) var scales: array>; +@group(0) @binding(5) var meshInfo: array; +@group(0) @binding(6) var materials: array>; +@group(0) @binding(7) var cameras: array>; +${mask ? "@group(1) @binding(0) var materialSampler: sampler;" : ""} +${baseTexture ? "@group(1) @binding(1) var baseTexture: texture_2d;" : ""} +${materialTexture ? "@group(1) @binding(2) var materialTexture: texture_2d;" : ""} +${normalTexture ? "@group(1) @binding(3) var normalTexture: texture_2d;" : ""} + +fn rotate(q: vec4, value: vec3) -> vec3 { + return value + 2.0 * cross(q.xyz, cross(q.xyz, value) + q.w * value); +} + +@vertex +fn vertex( + @location(0) point: vec3, + @location(1) localNormal: vec3, + @location(2) uv: vec2, + ${normalTexture ? "@location(3) localTangent: vec4," : ""} + @builtin(instance_index) packed: u32, +) -> VertexOutput { + let instance = packed & 65535u; + let scale = scales[instance].xyz; + let world = rotate(quaternions[instance], point * scale) + positions[instance].xyz; + var clip = vec4(world, 1.0); + if (cameras[2].w != 0.0) { + let cameraNode = u32(cameras[1].x); + let inverse = vec4(-quaternions[cameraNode].xyz, quaternions[cameraNode].w); + let view = rotate(inverse, world - positions[cameraNode].xyz); + if (cameras[1].y == 1.0) { + let size = max(cameras[1].z, 0.0001); + clip = vec4(view.x / (size * cameras[0].y * 0.5), view.y / (size * 0.5), -view.z / cameras[0].w, 1.0); + } else { + let focal = 1.0 / tan(cameras[0].x * 0.5); + let depth = (-view.z * cameras[0].w - cameras[0].z * cameras[0].w) / (cameras[0].w - cameras[0].z); + clip = vec4(view.x * focal / cameras[0].y, view.y * focal, depth, -view.z); + } + } + var output: VertexOutput; + output.position = select(vec4(2.0, 2.0, 2.0, 1.0), clip, meshInfo[instance * 4u + 2u] != 0u); + output.world = world; + output.normal = normalize(rotate(quaternions[instance], localNormal / scale)); + output.uv = uv; + output.tangent = ${normalTexture ? "vec4(normalize(rotate(quaternions[instance], localTangent.xyz * scale)), localTangent.w)" : "vec4(1.0, 0.0, 0.0, 1.0)"}; + output.mesh = instance; + output.material = packed >> 16u; + return output; +} + +@fragment +fn fragment(input: VertexOutput) -> @location(0) vec4 { + let fallback = meshInfo[input.mesh * 4u + 1u]; + let material = select(fallback, input.material - 1u, input.material != 0u); + let factor = materials[material * 3u]; + let properties = materials[material * 3u + 1u]; + let extra = materials[material * 3u + 2u]; + let base = factor * ${baseTexture ? "textureSample(baseTexture, materialSampler, input.uv)" : "vec4(1.0)"}; + let packedMaterial = ${materialTexture ? "textureSample(materialTexture, materialSampler, input.uv)" : "vec4(1.0)"}; + let metallic = clamp(properties.x * ${materialTexture ? "packedMaterial.b" : "1.0"}, 0.0, 1.0); + let roughness = clamp(properties.y * ${materialTexture ? "packedMaterial.g" : "1.0"}, 0.04, 1.0); + var normal = normalize(input.normal); + ${normalTexture ? "let tangent = normalize(input.tangent.xyz); let bitangent = normalize(cross(normal, tangent)) * input.tangent.w; let mapped = textureSample(normalTexture, materialSampler, input.uv).xyz * 2.0 - 1.0; normal = normalize(mat3x3(tangent, bitangent, normal) * vec3(mapped.xy * extra.y, mapped.z));" : ""} + let cameraNode = u32(cameras[1].x); + let view = normalize(positions[cameraNode].xyz - input.world); + let lightDirection = normalize(vec3(0.4, 0.7, 0.6)); + let halfVector = normalize(lightDirection + view); + let nDotL = max(dot(normal, lightDirection), 0.0); + let nDotH = max(dot(normal, halfVector), 0.0); + let f0 = mix(vec3(0.04), base.rgb, metallic); + let specular = f0 * pow(nDotH, max(2.0, 2.0 / (roughness * roughness) - 2.0)); + let clusterLight = vec3(bitcast(clusters[1]), bitcast(clusters[2]), bitcast(clusters[3])); + let ambient = vec3(0.08) + clusterLight; + let diffuse = base.rgb * (1.0 - metallic) * nDotL; + let emissive = vec3(properties.z, properties.w, extra.x); + return vec4((base.rgb * ambient + diffuse + specular * nDotL + emissive) * accent.a, base.a); +}`; +} + const emptyForwardShader = /* wgsl */ ` struct Accent { color: vec4 } @group(0) @binding(0) var accent: Accent; @@ -257,10 +354,13 @@ export class Scene { constructor( canvas: HTMLCanvasElement, - options: { arenaBytes?: number; fps?: number; hdr?: boolean } = {}, + options: { arenaBytes?: number; debug?: boolean; fps?: number; hdr?: boolean } = {}, ) { this.hdr = options.hdr ?? true; - this.core = new YawnCore(canvas, { arenaBytes: options.arenaBytes }); + this.core = new YawnCore(canvas, { + arenaBytes: options.arenaBytes, + debug: options.debug, + }); this.ready = this.#initialize(options.fps ?? 60); } @@ -635,7 +735,23 @@ export class Scene { usage: ["render", "sampled"], transient: false, }); + addTexture({ + id: "depth", + format: "depth24plus", + size: ["canvas", "canvas", 1], + usage: ["render"], + transient: true, + }); addSampler({ id: "linear", magFilter: "linear", minFilter: "linear" }); + addSampler({ + id: "material-linear", + magFilter: "linear", + minFilter: "linear", + addressModeU: "repeat", + addressModeV: "repeat", + }); + for (const { number: _, source: __, ...texture } of this.#textures.values()) + addTexture(texture); let previous = computeIds.length ? computeIds : ["cluster-lights"]; if (!renderedMeshes.length) { @@ -665,22 +781,7 @@ export class Scene { ["cameras", "cameras"], ]) addBuffer({ id, array, usage: ["storage"] }); - renderPipelines.push({ - id: "forward-pbr", - code: forwardShader, - vertex: { - entry: "vertex", - buffers: [ - { - arrayStride: 16, - attributes: [ - { format: "float32x3", offset: 0, shaderLocation: 0 }, - ], - }, - ], - }, - fragment: { entry: "fragment", targets: [{ format: hdrFormat }] }, - }); + const forwardPipelines = new Set(); let firstRender = true; for (const mesh of renderedMeshes) { const vertex = `geometry-${mesh.geometryId}-positions`; @@ -688,13 +789,33 @@ export class Scene { id: vertex, array: `geometry.${mesh.geometryId}.positions`, usage: ["vertex"], + sync: "loadout", }); + const normal = `geometry-${mesh.geometryId}-normals`; + const uv = `geometry-${mesh.geometryId}-uvs`; + const tangent = `geometry-${mesh.geometryId}-tangents`; + const hasNormals = !!this.geometryData(mesh.geometryId, "normals"); + const hasUvs = !!this.geometryData(mesh.geometryId, "uvs"); + const hasTangents = !!this.geometryData(mesh.geometryId, "tangents"); + for (const [present, id, kind] of [ + [hasNormals, normal, "normals"], + [hasUvs, uv, "uvs"], + [hasTangents, tangent, "tangents"], + ] as const) + if (present) + addBuffer({ + id, + array: `geometry.${mesh.geometryId}.${kind}`, + usage: ["vertex"], + sync: "loadout", + }); const indexed = mesh.indexCount > 0; if (indexed) addBuffer({ id: `geometry-${mesh.geometryId}-indices`, array: `geometry.${mesh.geometryId}.indices`, usage: ["index"], + sync: "loadout", }); const draws = indexed && mesh.faceMaterials.size @@ -717,8 +838,103 @@ export class Scene { ]; for (const draw of draws) { const id = `forward-${mesh.id}-${draw.face}`; - if (mesh.id > 65535 || (draw.material ?? 0) > 65534) + const material = + draw.material ?? Number(this.array("meshInfo").row(mesh.id)[1]); + if (mesh.id > 65535 || material > 65534) throw new RangeError("Scene handle limit"); + const pointers = this.array("materialTextures").row(material); + const texture = (lane: number) => + pointers[lane] + ? this.#textures.get(Number(pointers[lane]) - 1) + : undefined; + const baseTexture = texture(0); + const materialTexture = texture(1); + const normalTexture = hasTangents ? texture(2) : undefined; + const detailed = hasNormals && hasUvs; + const mask = detailed + ? (baseTexture ? 1 : 0) | + (materialTexture ? 2 : 0) | + (normalTexture ? 4 : 0) + : 0; + const pipeline = detailed ? `forward-pbr-${mask}` : "forward-basic"; + if (!forwardPipelines.has(pipeline)) { + forwardPipelines.add(pipeline); + renderPipelines.push({ + id: pipeline, + code: detailed ? pbrShader(mask) : basicForwardShader, + vertex: { + entry: "vertex", + buffers: detailed + ? [ + { + arrayStride: 16, + attributes: [ + { + format: "float32x3", + offset: 0, + shaderLocation: 0, + }, + ], + }, + { + arrayStride: 16, + attributes: [ + { + format: "float32x3", + offset: 0, + shaderLocation: 1, + }, + ], + }, + { + arrayStride: 16, + attributes: [ + { + format: "float32x2", + offset: 0, + shaderLocation: 2, + }, + ], + }, + ...(normalTexture + ? [ + { + arrayStride: 16, + attributes: [ + { + format: "float32x4", + offset: 0, + shaderLocation: 3, + }, + ], + }, + ] + : []), + ] + : [ + { + arrayStride: 16, + attributes: [ + { + format: "float32x3", + offset: 0, + shaderLocation: 0, + }, + ], + }, + ], + }, + fragment: { + entry: "fragment", + targets: [{ format: hdrFormat }], + }, + depthStencil: { + format: "depth24plus", + depth_write_enabled: true, + depth_compare: "less", + }, + }); + } const instance = (mesh.id + (draw.material === undefined @@ -728,18 +944,56 @@ export class Scene { passes.push({ id, type: "render", - pipeline: "forward-pbr", + pipeline, after: previous, bindings: [ - "clusters", - "accent", - "node-positions", - "node-quaternions", - "node-scales", - "mesh-info", - "materials", - "cameras", - ].map((resource, binding) => ({ group: 0, binding, resource })), + ...[ + "clusters", + "accent", + "node-positions", + "node-quaternions", + "node-scales", + "mesh-info", + "materials", + "cameras", + ].map((resource, binding) => ({ + group: 0, + binding, + resource, + })), + ...(detailed && mask + ? [ + { group: 1, binding: 0, resource: "material-linear" }, + ...(baseTexture + ? [ + { + group: 1, + binding: 1, + resource: baseTexture.id, + }, + ] + : []), + ...(materialTexture + ? [ + { + group: 1, + binding: 2, + resource: materialTexture.id, + }, + ] + : []), + ...(normalTexture + ? [ + { + group: 1, + binding: 3, + resource: normalTexture.id, + }, + ] + : []), + ] + : []), + ], color: [ { resource: "hdr", @@ -748,7 +1002,20 @@ export class Scene { : { load: "load" }), }, ], - vertexBuffers: [{ slot: 0, resource: vertex }], + depth: { + resource: "depth", + ...(firstRender ? { clear: 1 } : { load: "load" }), + }, + vertexBuffers: [ + { slot: 0, resource: vertex }, + ...(detailed + ? [ + { slot: 1, resource: normal }, + { slot: 2, resource: uv }, + ...(normalTexture ? [{ slot: 3, resource: tangent }] : []), + ] + : []), + ], ...(indexed ? { indexBuffer: { @@ -827,26 +1094,6 @@ export class Scene { previous = [pipeline]; } - for (const { number, source: _, ...texture } of this.#textures.values()) { - addTexture(texture); - const id = `retain-texture-${number}`; - computePipelines.push({ - id, - code: "@group(0) @binding(0) var source: texture_2d; @group(0) @binding(1) var output: array; @compute @workgroup_size(1) fn main() { output[1] = textureDimensions(source).x; }", - entry: "main", - }); - passes.push({ - id, - type: "compute", - pipeline: id, - after: ["cluster-lights"], - dispatch: [1, 1, 1], - bindings: [ - { group: 0, binding: 0, resource: texture.id }, - { group: 0, binding: 1, resource: "clusters" }, - ], - }); - } const toneMap = String( [...this.#effects.values()].find( (effect) => effect.kind === "colorGrading", diff --git a/addons/handles/src/importers/gltf.ts b/addons/handles/src/importers/gltf.ts index 66fff50..520865a 100644 --- a/addons/handles/src/importers/gltf.ts +++ b/addons/handles/src/importers/gltf.ts @@ -1,6 +1,7 @@ import { Mesh } from "../Mesh"; import type { Scene } from "../Scene"; import { PBRMaterial } from "../materials/PBRMaterial"; +import { Texture } from "../materials/Texture"; let worker: Worker | undefined; let nextRequest = 1; @@ -42,9 +43,31 @@ export async function importGltf(scene: Scene, url: string | URL) { nodes: result.primitives.length, materials: result.materials.length, }); - return scene.batchGraphUpdates(async () => { + let textures: Texture[] = []; + const imported = await scene.batchGraphUpdates(async () => { + const srgb = new Set(); + for (const material of result.materials) { + if (material.baseColorTexture >= 0) srgb.add(material.baseColorTexture); + if (material.emissiveTexture >= 0) srgb.add(material.emissiveTexture); + } + textures = result.textures.map( + ({ image }: { image: ImageBitmap }, index: number) => + new Texture(scene, { + source: image, + format: srgb.has(index) ? "rgba8unorm-srgb" : "rgba8unorm", + }), + ); + await Promise.all(textures.map((texture: Texture) => texture.ready)); const materials = result.materials.map( - (options: any) => new PBRMaterial(scene, options), + (options: any) => + new PBRMaterial(scene, { + ...options, + baseColorTexture: textures[options.baseColorTexture], + metallicRoughnessTexture: + textures[options.metallicRoughnessTexture], + normalTexture: textures[options.normalTexture], + emissiveTexture: textures[options.emissiveTexture], + }), ); await Promise.all(materials.map((material: PBRMaterial) => material.ready)); const meshes: Mesh[] = result.primitives.map( @@ -67,4 +90,10 @@ export async function importGltf(scene: Scene, url: string | URL) { await Promise.all(meshes.map((mesh) => mesh.ready)); return meshes; }); + await Promise.all( + result.textures.map((_: unknown, index: number) => + scene.core.deleteTexture(textures[index].resource), + ), + ); + return imported; } diff --git a/addons/handles/src/importers/worker.ts b/addons/handles/src/importers/worker.ts index d3a7225..5987c9e 100644 --- a/addons/handles/src/importers/worker.ts +++ b/addons/handles/src/importers/worker.ts @@ -144,6 +144,26 @@ async function load(url: string) { return new Uint8Array(await result.arrayBuffer()); }), ); + const imageBlobs = await Promise.all( + (document.images ?? []).map(async (image: any) => { + if (image.bufferView !== undefined) { + const view = document.bufferViews[image.bufferView]; + const bytes = buffers[view.buffer]; + const start = view.byteOffset ?? 0; + return new Blob([bytes.slice(start, start + view.byteLength)], { + type: image.mimeType, + }); + } + const result = await fetch(new URL(image.uri, url)); + if (!result.ok) throw new Error(`HTTP_${result.status}`); + return result.blob(); + }), + ); + const textures = await Promise.all( + (document.textures ?? []).map(async (texture: any) => ({ + image: await createImageBitmap(imageBlobs[texture.source]), + })), + ); const accessor = (id: number, integer = false) => { const source = document.accessors[id]; @@ -199,6 +219,10 @@ async function load(url: string) { roughness: pbr.roughnessFactor ?? 0.7, emissive: material.emissiveFactor ?? [0, 0, 0], alphaCutoff: material.alphaCutoff ?? 0.5, + baseColorTexture: pbr.baseColorTexture?.index ?? -1, + metallicRoughnessTexture: pbr.metallicRoughnessTexture?.index ?? -1, + normalTexture: material.normalTexture?.index ?? -1, + emissiveTexture: material.emissiveTexture?.index ?? -1, }; }); const primitives: any[] = []; @@ -264,7 +288,7 @@ async function load(url: string) { if (!nodes.length) for (let id = 0; id < (document.meshes ?? []).length; id++) emitMesh(id, {}); - return { materials, primitives }; + return { materials, primitives, textures }; } addEventListener("message", async ({ data }) => { @@ -275,6 +299,7 @@ addEventListener("message", async ({ data }) => { .map((name) => primitive[name]?.buffer) .filter(Boolean), ); + transfers.push(...result.textures.map((texture: any) => texture.image)); (postMessage as any)({ request: data.request, result }, transfers); } catch (error) { postMessage({ diff --git a/addons/handles/src/materials/PBRMaterial.ts b/addons/handles/src/materials/PBRMaterial.ts index 3005b77..d868472 100644 --- a/addons/handles/src/materials/PBRMaterial.ts +++ b/addons/handles/src/materials/PBRMaterial.ts @@ -23,7 +23,14 @@ export class PBRMaterial { constructor(scene: Scene, options: PBRMaterialOptions = {}) { this.scene = scene; - this.ready = scene.allocateMaterial().then((id) => { + const textures = [ + options.baseColorTexture, + options.metallicRoughnessTexture, + options.normalTexture, + options.emissiveTexture, + ].filter((texture): texture is Texture => texture !== undefined); + this.ready = Promise.all(textures.map((texture) => texture.ready)).then(async () => { + const id = await scene.allocateMaterial(); this.id = id; const color = Array.from(options.baseColor ?? [1, 1, 1, 1]); const emissive = Array.from(options.emissive ?? [0, 0, 0]); diff --git a/addons/handles/src/materials/Texture.ts b/addons/handles/src/materials/Texture.ts index cb2850e..5381ada 100644 --- a/addons/handles/src/materials/Texture.ts +++ b/addons/handles/src/materials/Texture.ts @@ -13,7 +13,7 @@ let nextTextureName = 1; /** A graph texture resource handle; image decoding/upload policy remains outside core. */ export class Texture { readonly scene: Scene; - readonly id: number; + id = -1; readonly resource: string; readonly source?: string | ImageBitmap; readonly ready: Promise; @@ -21,24 +21,44 @@ export class Texture { constructor(scene: Scene, options: TextureOptions = {}) { this.scene = scene; - this.resource = `texture-${nextTextureName++}`; + const resource = `texture-${nextTextureName++}`; + this.resource = resource; this.source = options.source; - const registration = scene.registerTexture({ - id: this.resource, - source: options.source, - size: options.size ?? [1, 1, 1], - format: options.format ?? "rgba8unorm", - usage: [...new Set([...(options.usage ?? ["copyDst"]), "sampled"])], - transient: options.transient ?? false, - }); - this.id = registration.number; - this.ready = registration.ready; + this.ready = (async () => { + const image = + typeof options.source === "string" + ? await createImageBitmap(await (await fetch(options.source)).blob()) + : options.source; + const registration = scene.registerTexture({ + id: resource, + source: image, + size: + options.size ?? + (image instanceof ImageBitmap + ? [image.width, image.height, 1] + : [1, 1, 1]), + format: options.format ?? "rgba8unorm", + usage: [ + ...new Set([ + ...(options.usage ?? ["copyDst"]), + "sampled", + ...(image instanceof ImageBitmap ? ["render"] : []), + ]), + ], + transient: options.transient ?? false, + }); + this.id = registration.number; + await registration.ready; + if (image instanceof ImageBitmap) + await scene.core.uploadTexture(resource, image); + })(); } async dispose() { await this.ready; if (this.#disposed) return; this.#disposed = true; + await this.scene.core.deleteTexture(this.resource); await this.scene.unregisterTexture(this.id); } } diff --git a/core/Cargo.toml b/core/Cargo.toml index da0c9c4..e073b10 100644 --- a/core/Cargo.toml +++ b/core/Cargo.toml @@ -15,5 +15,5 @@ serde = { version = "1", features = ["derive"] } serde_json = "1" wasm-bindgen = "0.2" wasm-bindgen-futures = "0.4" -web-sys = { version = "0.3", features = ["OffscreenCanvas"] } +web-sys = { version = "0.3", features = ["ImageBitmap", "OffscreenCanvas"] } wgpu = { version = "26", default-features = false, features = ["serde", "webgpu", "wgsl"] } diff --git a/core/index.js b/core/index.js index 66fda8c..9d110d9 100644 --- a/core/index.js +++ b/core/index.js @@ -46,9 +46,10 @@ export class YawnCore { #buffer; #arrays = new Map(); #pending = new Map(); + #profileListeners = new Set(); #next = 1; - constructor(canvas, { arenaBytes = 64 * 1024 * 1024, workerFactory } = {}) { + constructor(canvas, { arenaBytes = 64 * 1024 * 1024, debug = false, workerFactory } = {}) { if (!canvas) throw new TypeError("canvas is required"); this.#worker = workerFactory?.() ?? new Worker(new URL("./worker.js", import.meta.url), { type: "module", @@ -59,12 +60,13 @@ export class YawnCore { this.#worker.addEventListener("messageerror", () => this.#fail("WORKER_ERROR")); this.#worker.start?.(); const offscreen = canvas.transferControlToOffscreen?.() ?? canvas; - this.ready = this.#request("init", { canvas: offscreen, arenaBytes }, [offscreen]).then(result => { + this.ready = this.#request("init", { canvas: offscreen, arenaBytes }, [offscreen]).then(async result => { this.#buffer = result.buffer; for (const descriptor of result.rows) this.#arrays.set( descriptor.name, new SharedRows(this.#buffer, descriptor), ); + if (debug) await this.#request("set-profiler", { enabled: true }); }); } @@ -117,6 +119,18 @@ export class YawnCore { return this.#request("switch-loadout", { id }); } + async uploadTexture(name, image) { + await this.ready; + if (typeof name !== "string" || !(image instanceof ImageBitmap)) + throw new TypeError("TEXTURE_SOURCE"); + return this.#request("upload-texture", { name, image }, [image]); + } + + async deleteTexture(name) { + await this.ready; + return this.#request("delete-texture", { name }); + } + async play() { await this.ready; return this.#request("play"); @@ -132,6 +146,17 @@ export class YawnCore { return this.#request("set-fps", { fps }); } + async setProfiler(enabled) { + await this.ready; + return this.#request("set-profiler", { enabled: Boolean(enabled) }); + } + + onProfile(listener) { + if (typeof listener !== "function") throw new TypeError("PROFILE_LISTENER"); + this.#profileListeners.add(listener); + return () => this.#profileListeners.delete(listener); + } + array(name) { const array = this.#arrays.get(name); if (!array) throw new Error(`UNKNOWN_ARRAY: ${name}`); @@ -147,6 +172,10 @@ export class YawnCore { } #message(message) { + if (message?.type === "profile") { + for (const listener of this.#profileListeners) listener(message.stats); + return; + } const pending = this.#pending.get(message?.request); if (!pending) return; this.#pending.delete(message.request); @@ -161,6 +190,7 @@ export class YawnCore { dispose() { this.#fail("DISPOSED"); + this.#profileListeners.clear(); this.#worker.terminate(); } } diff --git a/core/lib.rs b/core/lib.rs index dd2cf7e..e0ee149 100644 --- a/core/lib.rs +++ b/core/lib.rs @@ -174,6 +174,21 @@ impl Core { .map_err(|error| JsError::new(&error)) } + pub fn upload_texture(&self, name: String, image: web_sys::ImageBitmap) -> Result<(), JsError> { + let gpu = self.gpu.borrow(); + let gpu = gpu + .as_ref() + .ok_or_else(|| JsError::new("WEBGPU_UNINITIALIZED"))?; + self.store + .borrow_mut() + .upload_texture(name, image, gpu) + .map_err(|error| JsError::new(&error)) + } + + pub fn delete_texture(&self, name: &str) { + self.store.borrow_mut().delete_texture(name); + } + pub fn play(&self) { self.render.play(); } @@ -185,4 +200,18 @@ impl Core { pub fn set_fps(&self, fps: u32) -> Result<(), JsError> { self.render.set_fps(fps).map_err(JsError::new) } + + pub fn set_profiler(&self, enabled: bool) -> bool { + let supported = self + .gpu + .borrow() + .as_ref() + .is_some_and(|gpu| gpu.timestamp_queries); + self.render.set_profiling(enabled && supported); + supported + } + + pub fn take_profile(&self) -> Option { + self.render.take_profile() + } } diff --git a/core/render_graph/gpu_resource.rs b/core/render_graph/gpu_resource.rs index a058063..6d6c812 100644 --- a/core/render_graph/gpu_resource.rs +++ b/core/render_graph/gpu_resource.rs @@ -1,11 +1,13 @@ use std::collections::{BTreeMap, HashMap}; use std::num::NonZeroU64; +use std::sync::atomic::{AtomicU8, Ordering}; +use std::sync::Arc; use serde::de::DeserializeOwned; use serde_json::Value; use crate::gpu::Wgpu; -use crate::graph::{Binding, Extent, Pass, RenderGraph, RenderPipeline, Texture}; +use crate::graph::{Binding, Execution, Extent, Pass, RenderGraph, RenderPipeline, Texture}; use crate::render_data::RenderData; pub struct GpuResources { @@ -16,28 +18,59 @@ pub struct GpuResources { pub render_pipelines: HashMap, pub compute_pipelines: HashMap, pub passes: Vec, + pub profiler: Option, +} + +pub struct GpuProfiler { + query_set: wgpu::QuerySet, + resolve: wgpu::Buffer, + readback: Arc, + query_count: u32, +} + +pub struct ProfileMap { + readback: Arc, + state: Arc, + labels: Vec, + timestamp_period: f32, } pub struct GpuBuffer { pub buffer: wgpu::Buffer, pub source: String, + pub sync_each_frame: bool, } +#[derive(Clone)] pub struct GpuTexture { - pub _texture: wgpu::Texture, + pub texture: wgpu::Texture, pub view: wgpu::TextureView, + key: String, + uploaded: bool, } pub enum GpuPass { - Render(wgpu::RenderBundle), + Render { + label: String, + first: usize, + last: usize, + bundle: wgpu::RenderBundle, + }, Compute { + label: String, + pass: usize, pipeline: wgpu::ComputePipeline, bind_groups: Vec<(u32, wgpu::BindGroup)>, }, } impl GpuResources { - pub fn activate(graph: &RenderGraph, gpu: &Wgpu, data: &RenderData) -> Result { + pub fn activate( + graph: &RenderGraph, + gpu: &Wgpu, + data: &RenderData, + previous: Option<&Self>, + ) -> Result { let mut buffers = HashMap::new(); for source in &graph.resources.buffers { let rows = data.rows(&source.array).ok_or("GRAPH_ARRAY_UNKNOWN")?; @@ -54,6 +87,7 @@ impl GpuResources { GpuBuffer { buffer, source: source.array.clone(), + sync_each_frame: source.sync == "frame", }, ); } @@ -65,12 +99,32 @@ impl GpuResources { let physical = match physical_slots.get(&source.slot) { Some(&physical) => physical, None => { + let key = source.key()?; + if !source.transient { + if let Some(texture) = previous + .and_then(|resources| { + resources + .texture_slots + .get(&source.id) + .map(|slot| &resources.textures[*slot]) + }) + .filter(|texture| texture.key == key) + { + let physical = textures.len(); + textures.push(texture.clone()); + physical_slots.insert(source.slot, physical); + texture_slots.insert(source.id.clone(), physical); + continue; + } + } let descriptor = texture_descriptor(source, gpu.width, gpu.height)?; let texture = gpu.device.create_texture(&descriptor); let physical = textures.len(); textures.push(GpuTexture { view: texture.create_view(&wgpu::TextureViewDescriptor::default()), - _texture: texture, + texture, + key, + uploaded: false, }); physical_slots.insert(source.slot, physical); physical @@ -121,6 +175,30 @@ impl GpuResources { }) .collect::, _>>()?; + let profiler = (gpu.timestamp_queries && !graph.executions.is_empty()).then(|| { + let query_count = graph.executions.len() as u32 * 2; + let bytes = u64::from(query_count) * 8; + GpuProfiler { + query_set: gpu.device.create_query_set(&wgpu::QuerySetDescriptor { + label: Some("frame-profile"), + ty: wgpu::QueryType::Timestamp, + count: query_count, + }), + resolve: gpu.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("frame-profile-resolve"), + size: bytes, + usage: wgpu::BufferUsages::QUERY_RESOLVE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }), + readback: Arc::new(gpu.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("frame-profile-readback"), + size: bytes, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + mapped_at_creation: false, + })), + query_count, + } + }); let mut resources = Self { buffers, textures, @@ -129,11 +207,27 @@ impl GpuResources { render_pipelines, compute_pipelines, passes: Vec::new(), + profiler, }; - for pass in &graph.passes { - let compiled = match pass.kind.as_str() { - "render" => GpuPass::Render(resources.render_bundle(graph, pass, gpu)?), - "compute" => { + for execution in &graph.executions { + let compiled = match execution { + Execution::Render(passes) => GpuPass::Render { + label: if passes.len() == 1 { + graph.passes[passes[0]].id.clone() + } else if passes + .iter() + .all(|index| graph.passes[*index].id.starts_with("forward-")) + { + format!("Forward ({} draws)", passes.len()) + } else { + format!("Render ({} draws)", passes.len()) + }, + first: passes[0], + last: *passes.last().unwrap(), + bundle: resources.render_bundle(graph, passes, gpu)?, + }, + Execution::Compute(index) => { + let pass = &graph.passes[*index]; let pipeline = resources .compute_pipelines .get(&pass.pipeline) @@ -145,11 +239,12 @@ impl GpuResources { gpu, )?; GpuPass::Compute { + label: pass.id.clone(), + pass: *index, pipeline, bind_groups, } } - _ => return Err("GRAPH_PASS".into()), }; resources.passes.push(compiled); } @@ -163,16 +258,146 @@ impl GpuResources { .map(|texture| &texture.view) } + pub fn upload_texture( + &mut self, + id: &str, + image: &web_sys::ImageBitmap, + gpu: &Wgpu, + ) -> Result<(), String> { + let texture = self + .texture_slots + .get(id) + .and_then(|slot| self.textures.get_mut(*slot)) + .ok_or("GRAPH_TEXTURE_UNKNOWN")?; + gpu.queue.copy_external_image_to_texture( + &wgpu::CopyExternalImageSourceInfo { + source: wgpu::ExternalImageSource::ImageBitmap(image.clone()), + origin: wgpu::Origin2d::ZERO, + flip_y: false, + }, + wgpu::TexelCopyTextureInfo { + texture: &texture.texture, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + } + .to_tagged(wgpu::PredefinedColorSpace::Srgb, false), + wgpu::Extent3d { + width: image.width(), + height: image.height(), + depth_or_array_layers: 1, + }, + ); + texture.uploaded = true; + Ok(()) + } + + pub fn needs_upload(&self, id: &str) -> bool { + self.texture_slots + .get(id) + .and_then(|slot| self.textures.get(*slot)) + .is_some_and(|texture| !texture.uploaded) + } + + pub fn render_timestamps(&self, pass: usize) -> Option> { + self.profiler + .as_ref() + .map(|profiler| wgpu::RenderPassTimestampWrites { + query_set: &profiler.query_set, + beginning_of_pass_write_index: Some(pass as u32 * 2), + end_of_pass_write_index: Some(pass as u32 * 2 + 1), + }) + } + + pub fn compute_timestamps(&self, pass: usize) -> Option> { + self.profiler + .as_ref() + .map(|profiler| wgpu::ComputePassTimestampWrites { + query_set: &profiler.query_set, + beginning_of_pass_write_index: Some(pass as u32 * 2), + end_of_pass_write_index: Some(pass as u32 * 2 + 1), + }) + } + + pub fn resolve_profile(&self, encoder: &mut wgpu::CommandEncoder) { + let Some(profiler) = &self.profiler else { + return; + }; + encoder.resolve_query_set( + &profiler.query_set, + 0..profiler.query_count, + &profiler.resolve, + 0, + ); + encoder.copy_buffer_to_buffer( + &profiler.resolve, + 0, + &profiler.readback, + 0, + u64::from(profiler.query_count) * 8, + ); + } + + pub fn map_profile(&self, timestamp_period: f32) -> Option { + let profiler = self.profiler.as_ref()?; + let state = Arc::new(AtomicU8::new(0)); + let callback = state.clone(); + profiler + .readback + .map_async(wgpu::MapMode::Read, .., move |result| { + callback.store(if result.is_ok() { 1 } else { 2 }, Ordering::Release); + }); + Some(ProfileMap { + readback: profiler.readback.clone(), + state, + labels: self + .passes + .iter() + .map(|pass| pass.label().to_owned()) + .collect(), + timestamp_period, + }) + } +} + +impl ProfileMap { + pub fn state(&self) -> u8 { + self.state.load(Ordering::Acquire) + } + + pub fn read(self) -> Vec<(String, f64)> { + let bytes = self.readback.get_mapped_range(..); + let values = bytes + .chunks_exact(8) + .map(|bytes| u64::from_le_bytes(bytes.try_into().unwrap())) + .collect::>(); + let profile = self + .labels + .into_iter() + .zip(values.chunks_exact(2)) + .map(|(label, timestamps)| { + ( + label, + timestamps[1].saturating_sub(timestamps[0]) as f64 + * f64::from(self.timestamp_period) + / 1_000_000.0, + ) + }) + .collect(); + drop(bytes); + self.readback.unmap(); + profile + } +} + +impl GpuResources { fn render_bundle( &self, graph: &RenderGraph, - pass: &Pass, + passes: &[usize], gpu: &Wgpu, ) -> Result { - let pipeline = self - .render_pipelines - .get(&pass.pipeline) - .ok_or("GRAPH_PIPELINE")?; + let pass = &graph.passes[passes[0]]; let declaration = graph .pipelines .render @@ -204,40 +429,56 @@ impl GpuResources { sample_count: multisample(&declaration.multisample)?.count, multiview: None, }); - encoder.set_pipeline(pipeline); - for (group, bind_group) in - self.bind_groups(pass, |group| pipeline.get_bind_group_layout(group), gpu)? - { - encoder.set_bind_group(group, &bind_group, &[]); - } - for binding in &pass.vertex_buffers { - let buffer = &self - .buffers - .get(&binding.resource) - .ok_or("GRAPH_RESOURCE_UNKNOWN")? - .buffer; - encoder.set_vertex_buffer(binding.slot, buffer.slice(binding.offset..)); - } - if let Some(binding) = &pass.index_buffer { - let buffer = &self - .buffers - .get(&binding.resource) - .ok_or("GRAPH_RESOURCE_UNKNOWN")? - .buffer; - encoder.set_index_buffer( - buffer.slice(binding.offset..), - parse(&binding.format, "GRAPH_INDEX_FORMAT")?, - ); - encoder.draw_indexed( - pass.draw.first_index..pass.draw.first_index + pass.draw.indices, - pass.draw.base_vertex, - pass.draw.first_instance..pass.draw.first_instance + pass.draw.instances, - ); - } else { - encoder.draw( - pass.draw.first_vertex..pass.draw.first_vertex + pass.draw.vertices, - pass.draw.first_instance..pass.draw.first_instance + pass.draw.instances, - ); + let mut previous_pipeline = None; + let mut previous_bindings: Option<&[Binding]> = None; + for index in passes { + let pass = &graph.passes[*index]; + let pipeline = self + .render_pipelines + .get(&pass.pipeline) + .ok_or("GRAPH_PIPELINE")?; + let pipeline_changed = previous_pipeline != Some(pass.pipeline.as_str()); + if pipeline_changed { + encoder.set_pipeline(pipeline); + previous_pipeline = Some(pass.pipeline.as_str()); + } + if pipeline_changed || previous_bindings != Some(pass.bindings.as_slice()) { + for (group, bind_group) in + self.bind_groups(pass, |group| pipeline.get_bind_group_layout(group), gpu)? + { + encoder.set_bind_group(group, &bind_group, &[]); + } + previous_bindings = Some(&pass.bindings); + } + for binding in &pass.vertex_buffers { + let buffer = &self + .buffers + .get(&binding.resource) + .ok_or("GRAPH_RESOURCE_UNKNOWN")? + .buffer; + encoder.set_vertex_buffer(binding.slot, buffer.slice(binding.offset..)); + } + if let Some(binding) = &pass.index_buffer { + let buffer = &self + .buffers + .get(&binding.resource) + .ok_or("GRAPH_RESOURCE_UNKNOWN")? + .buffer; + encoder.set_index_buffer( + buffer.slice(binding.offset..), + parse(&binding.format, "GRAPH_INDEX_FORMAT")?, + ); + encoder.draw_indexed( + pass.draw.first_index..pass.draw.first_index + pass.draw.indices, + pass.draw.base_vertex, + pass.draw.first_instance..pass.draw.first_instance + pass.draw.instances, + ); + } else { + encoder.draw( + pass.draw.first_vertex..pass.draw.first_vertex + pass.draw.vertices, + pass.draw.first_instance..pass.draw.first_instance + pass.draw.instances, + ); + } } Ok(encoder.finish(&wgpu::RenderBundleDescriptor { label: Some(&pass.id), @@ -290,6 +531,14 @@ impl GpuResources { } } +impl GpuPass { + fn label(&self) -> &str { + match self { + Self::Render { label, .. } | Self::Compute { label, .. } => label, + } + } +} + fn create_render_pipeline( source: &RenderPipeline, gpu: &Wgpu, diff --git a/core/render_graph/render_graph.rs b/core/render_graph/render_graph.rs index b2efe27..90a81df 100644 --- a/core/render_graph/render_graph.rs +++ b/core/render_graph/render_graph.rs @@ -12,6 +12,14 @@ pub struct RenderGraph { #[serde(default)] pub pipelines: PipelineDeclarations, pub passes: Vec, + #[serde(skip)] + pub executions: Vec, +} + +#[derive(Clone)] +pub enum Execution { + Compute(usize), + Render(Vec), } #[derive(Clone, Default, Deserialize)] @@ -30,6 +38,8 @@ pub struct Buffer { pub array: String, #[serde(default)] pub usage: Vec, + #[serde(default = "frame_sync")] + pub sync: String, } #[derive(Clone, Deserialize, Serialize)] @@ -186,7 +196,7 @@ pub struct Pass { pub dispatch: [u32; 3], } -#[derive(Clone, Deserialize)] +#[derive(Clone, Deserialize, PartialEq, Eq)] pub struct Binding { #[serde(default)] pub group: u32, @@ -318,7 +328,9 @@ impl RenderGraph { for pass in &mut self.passes { pass.dependencies = pass.after.iter().map(|id| sorted_ids[id]).collect(); } - self.plan_resources() + self.plan_resources()?; + self.plan_execution(); + Ok(()) } fn plan_resources(&mut self) -> Result<(), &'static str> { @@ -391,6 +403,80 @@ impl RenderGraph { self.validate_ids() } + fn plan_execution(&mut self) { + let mut executions = Vec::new(); + for index in 0..self.passes.len() { + let merge = executions.last().is_some_and(|execution| match execution { + Execution::Render(passes) => self.can_merge_render(*passes.last().unwrap(), index), + Execution::Compute(_) => false, + }); + if merge { + let Some(Execution::Render(passes)) = executions.last_mut() else { + unreachable!() + }; + passes.push(index); + } else if self.passes[index].kind == "render" { + executions.push(Execution::Render(vec![index])); + } else { + executions.push(Execution::Compute(index)); + } + } + self.executions = executions; + } + + fn can_merge_render(&self, previous: usize, next: usize) -> bool { + let previous = &self.passes[previous]; + let next = &self.passes[next]; + if next.kind != "render" + || previous.color.len() != next.color.len() + || self.sample_count(&previous.pipeline) != self.sample_count(&next.pipeline) + || previous + .color + .iter() + .zip(&next.color) + .any(|(previous, next)| { + previous.resource != next.resource + || previous.store != "store" + || next.load != "load" + }) + { + return false; + } + let same_depth = match (&previous.depth, &next.depth) { + (None, None) => true, + (Some(previous), Some(next)) => { + previous.resource == next.resource + && previous.store == "store" + && next.load == "load" + } + _ => false, + }; + same_depth + && !next.bindings.iter().any(|binding| { + next.color + .iter() + .any(|attachment| attachment.resource == binding.resource) + || next + .depth + .as_ref() + .is_some_and(|attachment| attachment.resource == binding.resource) + }) + } + + fn sample_count(&self, pipeline: &str) -> Option { + self.pipelines + .render + .iter() + .find(|value| value.id == pipeline) + .map(|value| { + value + .multisample + .get("count") + .and_then(Value::as_u64) + .unwrap_or(1) + }) + } + fn validate_ids(&self) -> Result<(), &'static str> { let mut resources = HashSet::new(); for id in self @@ -405,6 +491,14 @@ impl RenderGraph { return Err("GRAPH_RESOURCE"); } } + if self + .resources + .buffers + .iter() + .any(|buffer| !matches!(buffer.sync.as_str(), "frame" | "loadout")) + { + return Err("GRAPH_BUFFER_SYNC"); + } let mut pipelines = HashSet::new(); for id in self .pipelines @@ -443,7 +537,7 @@ impl Pass { } impl Texture { - fn key(&self) -> Result { + pub(crate) fn key(&self) -> Result { let mut usage = self.usage.clone(); usage.sort_unstable(); usage.dedup(); @@ -498,3 +592,6 @@ fn index_format() -> String { fn dispatch() -> [u32; 3] { [1, 1, 1] } +fn frame_sync() -> String { + "frame".into() +} diff --git a/core/render_graph/store.rs b/core/render_graph/store.rs index 3a74a93..b27f60b 100644 --- a/core/render_graph/store.rs +++ b/core/render_graph/store.rs @@ -13,6 +13,7 @@ pub struct Loadout { #[derive(Default)] pub struct Store { graphs: HashMap, + texture_sources: HashMap, active: Option, } @@ -25,11 +26,44 @@ impl Store { pub fn switch(&mut self, id: &str, gpu: &Wgpu, data: &RenderData) -> Result<(), String> { let graph = self.graphs.get(id).ok_or("GRAPH_UNKNOWN")?.clone(); - let resources = GpuResources::activate(&graph, gpu, data)?; + let mut resources = GpuResources::activate( + &graph, + gpu, + data, + self.active.as_ref().map(|active| &active.resources), + )?; + for (name, image) in &self.texture_sources { + if resources.needs_upload(name) { + resources.upload_texture(name, image, gpu)?; + } + } self.active = Some(Loadout { graph, resources }); Ok(()) } + pub fn upload_texture( + &mut self, + name: String, + image: web_sys::ImageBitmap, + gpu: &Wgpu, + ) -> Result<(), String> { + if let Some(active) = &mut self.active { + if active.resources.texture_slots.contains_key(&name) { + active.resources.upload_texture(&name, &image, gpu)?; + } + } + if let Some(previous) = self.texture_sources.insert(name, image) { + previous.close(); + } + Ok(()) + } + + pub fn delete_texture(&mut self, name: &str) { + if let Some(image) = self.texture_sources.remove(name) { + image.close(); + } + } + pub fn active_mut(&mut self) -> Option<&mut Loadout> { self.active.as_mut() } @@ -62,7 +96,17 @@ impl Store { return Ok(()); }; let graph = active.graph.clone(); - let resources = GpuResources::activate(&graph, gpu, data)?; + let mut resources = GpuResources::activate( + &graph, + gpu, + data, + self.active.as_ref().map(|active| &active.resources), + )?; + for (name, image) in &self.texture_sources { + if resources.needs_upload(name) { + resources.upload_texture(name, image, gpu)?; + } + } self.active = Some(Loadout { graph, resources }); Ok(()) } diff --git a/core/renderer/render.rs b/core/renderer/render.rs index a85b5e0..e841ee9 100644 --- a/core/renderer/render.rs +++ b/core/renderer/render.rs @@ -1,10 +1,12 @@ use std::cell::{Cell, RefCell}; use std::rc::Rc; +use std::sync::atomic::{AtomicBool, Ordering}; +use std::sync::Arc; use gloo_timers::future::TimeoutFuture; use crate::gpu::Wgpu; -use crate::gpu_resource::GpuPass; +use crate::gpu_resource::{GpuPass, ProfileMap}; use crate::graph::{ColorAttachment, DepthAttachment}; use crate::render_data::RenderData; use crate::store::{Loadout, Store}; @@ -16,6 +18,13 @@ pub struct RenderLoop { frame: Cell, elapsed: Cell, last: Cell, + profiling: Cell, + profile: RefCell>, +} + +struct Submission { + completed: Arc, + profile: Option, } impl RenderLoop { @@ -27,6 +36,8 @@ impl RenderLoop { frame: Cell::new(0), elapsed: Cell::new(0.0), last: Cell::new(js_sys::Date::now()), + profiling: Cell::new(false), + profile: RefCell::new(None), } } @@ -47,6 +58,17 @@ impl RenderLoop { Ok(()) } + pub fn set_profiling(&self, enabled: bool) { + self.profiling.set(enabled); + if !enabled { + self.profile.borrow_mut().take(); + } + } + + pub fn take_profile(&self) -> Option { + self.profile.borrow_mut().take() + } + pub fn start( self: &Rc, gpu: Rc>>, @@ -71,11 +93,51 @@ impl RenderLoop { data.update_info(delta as f32, frame, elapsed as f32, control.fps.get()); data.skip_render() }; - if !skip { + let submission = if !skip { if let (Some(gpu), Some(loadout)) = (gpu.borrow_mut().as_mut(), store.borrow_mut().active_mut()) { - let _ = gpu.render(loadout, &data.borrow()); + gpu.render(loadout, &data.borrow(), control.profiling.get()) + .ok() + .flatten() + } else { + None + } + } else { + None + }; + if let Some(submission) = submission { + while !submission.completed.load(Ordering::Acquire) { + TimeoutFuture::new(0).await; + } + if let Some(profile) = submission.profile { + while profile.state() == 0 { + TimeoutFuture::new(0).await; + } + if profile.state() == 1 { + let passes = profile + .read() + .into_iter() + .map(|(name, milliseconds)| { + serde_json::json!({ + "name": name, + "milliseconds": milliseconds, + }) + }) + .collect::>(); + let milliseconds = passes + .iter() + .filter_map(|pass| pass["milliseconds"].as_f64()) + .sum::(); + *control.profile.borrow_mut() = Some( + serde_json::json!({ + "frame": frame, + "milliseconds": milliseconds, + "passes": passes, + }) + .to_string(), + ); + } } } } @@ -88,8 +150,16 @@ impl RenderLoop { } impl Wgpu { - fn render(&mut self, loadout: &mut Loadout, data: &RenderData) -> Result<(), String> { + fn render( + &mut self, + loadout: &mut Loadout, + data: &RenderData, + profile: bool, + ) -> Result, String> { for buffer in loadout.resources.buffers.values() { + if !buffer.sync_each_frame { + continue; + } self.queue.write_buffer( &buffer.buffer, 0, @@ -102,7 +172,7 @@ impl Wgpu { self.surface.configure(&self.device, &self.config); self.surface.get_current_texture().map_err(|_| "SURFACE")? } - Err(wgpu::SurfaceError::Timeout) => return Ok(()), + Err(wgpu::SurfaceError::Timeout) => return Ok(None), Err(_) => return Err("SURFACE".into()), }; let surface_view = output @@ -113,15 +183,21 @@ impl Wgpu { .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("frame"), }); - for (pass, compiled) in loadout.graph.passes.iter().zip(&loadout.resources.passes) { + for (pass_index, compiled) in loadout.resources.passes.iter().enumerate() { match compiled { GpuPass::Compute { + pass, pipeline, bind_groups, + .. } => { + let pass = &loadout.graph.passes[*pass]; + let timestamp_writes = profile + .then(|| loadout.resources.compute_timestamps(pass_index)) + .flatten(); let mut command = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some(&pass.id), - timestamp_writes: None, + timestamp_writes, }); command.set_pipeline(pipeline); for (group, bind_group) in bind_groups { @@ -133,11 +209,19 @@ impl Wgpu { pass.dispatch[2], ); } - GpuPass::Render(bundle) => { + GpuPass::Render { + first, + last, + bundle, + .. + } => { + let pass = &loadout.graph.passes[*first]; + let last = &loadout.graph.passes[*last]; let colors = pass .color .iter() - .map(|attachment| { + .zip(&last.color) + .map(|(attachment, final_attachment)| { Ok(Some(wgpu::RenderPassColorAttachment { view: view( &loadout.resources, @@ -146,39 +230,57 @@ impl Wgpu { )?, depth_slice: None, resolve_target: None, - ops: color_ops(attachment)?, + ops: color_ops(attachment, final_attachment)?, })) }) .collect::, String>>()?; let depth = pass .depth .as_ref() - .map(|attachment| { + .zip(last.depth.as_ref()) + .map(|(attachment, final_attachment)| { Ok::<_, String>(wgpu::RenderPassDepthStencilAttachment { view: view( &loadout.resources, &surface_view, &attachment.resource, )?, - depth_ops: Some(depth_ops(attachment)?), + depth_ops: Some(depth_ops(attachment, final_attachment)?), stencil_ops: None, }) }) .transpose()?; + let timestamp_writes = profile + .then(|| loadout.resources.render_timestamps(pass_index)) + .flatten(); let mut command = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some(&pass.id), color_attachments: &colors, depth_stencil_attachment: depth, - timestamp_writes: None, + timestamp_writes, occlusion_query_set: None, }); command.execute_bundles([bundle]); } } } + if profile { + loadout.resources.resolve_profile(&mut encoder); + } self.queue.submit([encoder.finish()]); + let profile = profile + .then(|| { + loadout + .resources + .map_profile(self.queue.get_timestamp_period()) + }) + .flatten(); output.present(); - Ok(()) + let completed = Arc::new(AtomicBool::new(false)); + let callback = completed.clone(); + self.queue + .on_submitted_work_done(move || callback.store(true, Ordering::Release)); + Ok(Some(Submission { completed, profile })) } } @@ -196,7 +298,10 @@ fn view<'a>( } } -fn color_ops(attachment: &ColorAttachment) -> Result, String> { +fn color_ops( + attachment: &ColorAttachment, + final_attachment: &ColorAttachment, +) -> Result, String> { let clear = attachment.clear.as_slice(); Ok(wgpu::Operations { load: match attachment.load.as_str() { @@ -209,18 +314,21 @@ fn color_ops(attachment: &ColorAttachment) -> Result return Err("GRAPH_LOAD_OP".into()), }, - store: store_op(&attachment.store)?, + store: store_op(&final_attachment.store)?, }) } -fn depth_ops(attachment: &DepthAttachment) -> Result, String> { +fn depth_ops( + attachment: &DepthAttachment, + final_attachment: &DepthAttachment, +) -> Result, String> { Ok(wgpu::Operations { load: match attachment.load.as_str() { "load" => wgpu::LoadOp::Load, "clear" => wgpu::LoadOp::Clear(attachment.clear), _ => return Err("GRAPH_LOAD_OP".into()), }, - store: store_op(&attachment.store)?, + store: store_op(&final_attachment.store)?, }) } diff --git a/core/renderer/wgpu.rs b/core/renderer/wgpu.rs index ca69f05..ce742f3 100644 --- a/core/renderer/wgpu.rs +++ b/core/renderer/wgpu.rs @@ -7,6 +7,7 @@ pub struct Wgpu { pub format: wgpu::TextureFormat, pub width: u32, pub height: u32, + pub timestamp_queries: bool, } impl Wgpu { @@ -30,8 +31,12 @@ impl Wgpu { }) .await .map_err(|_| "WEBGPU_UNAVAILABLE")?; + let required_features = adapter.features() & wgpu::Features::TIMESTAMP_QUERY; let (device, queue) = adapter - .request_device(&wgpu::DeviceDescriptor::default()) + .request_device(&wgpu::DeviceDescriptor { + required_features, + ..Default::default() + }) .await .map_err(|_| "DEVICE")?; let config = surface @@ -48,6 +53,7 @@ impl Wgpu { format, width, height, + timestamp_queries: required_features.contains(wgpu::Features::TIMESTAMP_QUERY), }) } } diff --git a/core/worker.js b/core/worker.js index c64d7ce..eb1c420 100644 --- a/core/worker.js +++ b/core/worker.js @@ -1,6 +1,7 @@ import initWasm, { Core } from "./pkg/yawn_core.js"; let core; +let profileTimer; const fail = code => { throw new Error(code); }; @@ -39,6 +40,12 @@ addEventListener("message", async ({ data: message }) => { case "switch-loadout": core.switch_loadout(message.id); break; + case "upload-texture": + core.upload_texture(message.name, message.image); + break; + case "delete-texture": + core.delete_texture(message.name); + break; case "play": core.play(); break; @@ -48,6 +55,14 @@ addEventListener("message", async ({ data: message }) => { case "set-fps": core.set_fps(message.fps); break; + case "set-profiler": + result = core.set_profiler(Boolean(message.enabled)); + clearInterval(profileTimer); + profileTimer = result && message.enabled ? setInterval(() => { + const stats = core.take_profile(); + if (stats) postMessage({ type: "profile", stats: JSON.parse(stats) }); + }, 250) : undefined; + break; default: fail("MESSAGE"); } diff --git a/docs/.vitepress/Playground.vue b/docs/.vitepress/Playground.vue index 31bc9b7..9f3c01f 100644 --- a/docs/.vitepress/Playground.vue +++ b/docs/.vitepress/Playground.vue @@ -1,36 +1,117 @@