From 762317b683554aa8f8d5c1b21783d3c74172af79 Mon Sep 17 00:00:00 2001 From: Martin Valigursky <59932779+mvaligursky@users.noreply.github.com> Date: Tue, 2 Apr 2024 11:39:43 +0100 Subject: [PATCH] Implementation of StorageBuffer on WebGPU (#6201) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Implementation of StorageBuffer on WebGPU * comment * allow storage buffer data to be returned to a preallocated typed buffer * comment * Update src/platform/graphics/storage-buffer.js Co-authored-by: Will Eastcott * Update src/platform/graphics/webgpu/webgpu-buffer.js Co-authored-by: Will Eastcott * Update src/platform/graphics/webgpu/webgpu-graphics-device.js Co-authored-by: Will Eastcott * changing name to ‘Unnamed’ --------- Co-authored-by: Martin Valigursky Co-authored-by: Will Eastcott --- .../src/examples/compute/histogram/config.mjs | 30 +++ .../examples/compute/histogram/example.mjs | 178 ++++++++++++++++++ .../examples/compute/texture-gen/example.mjs | 2 +- src/core/constants.js | 8 + src/index.js | 3 +- src/platform/graphics/bind-group-format.js | 37 +++- src/platform/graphics/bind-group.js | 30 ++- src/platform/graphics/compute.js | 15 +- src/platform/graphics/constants.js | 53 ++++++ src/platform/graphics/graphics-device.js | 3 +- src/platform/graphics/storage-buffer.js | 69 +++++++ .../webgpu/webgpu-bind-group-format.js | 18 ++ .../graphics/webgpu/webgpu-bind-group.js | 18 ++ src/platform/graphics/webgpu/webgpu-buffer.js | 45 +++-- .../graphics/webgpu/webgpu-compute.js | 5 + .../graphics/webgpu/webgpu-graphics-device.js | 109 ++++++++++- .../graphics/webgpu/webgpu-index-buffer.js | 6 +- .../graphics/webgpu/webgpu-uniform-buffer.js | 13 +- .../graphics/webgpu/webgpu-vertex-buffer.js | 12 +- src/scene/renderer/render-pass-forward.js | 2 +- 20 files changed, 602 insertions(+), 54 deletions(-) create mode 100644 examples/src/examples/compute/histogram/config.mjs create mode 100644 examples/src/examples/compute/histogram/example.mjs create mode 100644 src/platform/graphics/storage-buffer.js diff --git a/examples/src/examples/compute/histogram/config.mjs b/examples/src/examples/compute/histogram/config.mjs new file mode 100644 index 00000000000..38039b991b1 --- /dev/null +++ b/examples/src/examples/compute/histogram/config.mjs @@ -0,0 +1,30 @@ +/** + * @type {import('../../../../types.mjs').ExampleConfig} + */ +export default { + HIDDEN: true, + WEBGPU_REQUIRED: true, + FILES: { + 'compute-shader.wgsl': /* wgsl */ ` + + @group(0) @binding(0) var inputTexture: texture_2d; + // @group(0) @binding(1) is a sampler of the inputTexture, but we don't need it in the shader. + @group(0) @binding(2) var bins: array>; + + fn luminance(color: vec3f) -> f32 { + return saturate(dot(color, vec3f(0.2126, 0.7152, 0.0722))); + } + + @compute @workgroup_size(1, 1, 1) + fn main(@builtin(global_invocation_id) global_invocation_id: vec3u) { + let numBins = f32(arrayLength(&bins)); + let lastBinIndex = u32(numBins - 1); + let position = global_invocation_id.xy; + let color = textureLoad(inputTexture, position, 0); + let v = luminance(color.rgb); + let bin = min(u32(v * numBins), lastBinIndex); + atomicAdd(&bins[bin], 1u); + } + ` + } +}; diff --git a/examples/src/examples/compute/histogram/example.mjs b/examples/src/examples/compute/histogram/example.mjs new file mode 100644 index 00000000000..62ec3a6dc7a --- /dev/null +++ b/examples/src/examples/compute/histogram/example.mjs @@ -0,0 +1,178 @@ +import * as pc from 'playcanvas'; +import { deviceType, rootPath } from '@examples/utils'; +import files from '@examples/files'; + +// Note: the example is based on this article: +// https://webgpufundamentals.org/webgpu/lessons/webgpu-compute-shaders-histogram.html +// A simpler but less performant version of the compute shader is used for simplicity. + +const canvas = document.getElementById('application-canvas'); +if (!(canvas instanceof HTMLCanvasElement)) { + throw new Error('No canvas found'); +} + +const assets = { + solid: new pc.Asset('solid', 'container', { url: rootPath + '/static/assets/models/icosahedron.glb' }), + helipad: new pc.Asset( + 'helipad-env-atlas', + 'texture', + { url: rootPath + '/static/assets/cubemaps/helipad-env-atlas.png' }, + { type: pc.TEXTURETYPE_RGBP, mipmaps: false } + ) +}; + +const gfxOptions = { + deviceTypes: [deviceType], + glslangUrl: rootPath + '/static/lib/glslang/glslang.js', + twgslUrl: rootPath + '/static/lib/twgsl/twgsl.js' +}; + +const device = await pc.createGraphicsDevice(canvas, gfxOptions); +const createOptions = new pc.AppOptions(); +createOptions.graphicsDevice = device; + +createOptions.componentSystems = [ + pc.RenderComponentSystem, + pc.CameraComponentSystem, + pc.LightComponentSystem, + pc.ScriptComponentSystem +]; +createOptions.resourceHandlers = [ + pc.TextureHandler, + pc.ContainerHandler +]; + +const app = new pc.AppBase(canvas); +app.init(createOptions); +app.start(); + +// Set the canvas to fill the window and automatically change resolution to be the same as the canvas size +app.setCanvasFillMode(pc.FILLMODE_FILL_WINDOW); +app.setCanvasResolution(pc.RESOLUTION_AUTO); + +// Ensure canvas is resized when window changes size +const resize = () => app.resizeCanvas(); +window.addEventListener('resize', resize); +app.on('destroy', () => { + window.removeEventListener('resize', resize); +}); + +const assetListLoader = new pc.AssetListLoader(Object.values(assets), app.assets); +assetListLoader.load(() => { + + // set up some general scene rendering properties + app.scene.toneMapping = pc.TONEMAP_ACES; + + // setup skydome + app.scene.skyboxMip = 2; + app.scene.skyboxIntensity = 0.3; + app.scene.envAtlas = assets.helipad.resource; + + // create camera entity + const camera = new pc.Entity('camera'); + camera.addComponent('camera'); + app.root.addChild(camera); + camera.setPosition(0, 0, 5); + + // Enable the camera to render the scene's color map, available as uSceneColorMap in the shaders. + // This allows us to use the rendered scene as an input for the histogram compute shader. + camera.camera.requestSceneColorMap(true); + + // create directional light entity + const light = new pc.Entity('light'); + light.addComponent('light', { + type: 'directional', + color: new pc.Color(1, 1, 1), + intensity: 15 + }); + app.root.addChild(light); + light.setEulerAngles(45, 0, 40); + + // a helper script that rotates the entity + const Rotator = pc.createScript('rotator'); + Rotator.prototype.update = function (/** @type {number} */ dt) { + this.entity.rotate(5 * dt, 10 * dt, -15 * dt); + }; + + // a compute shader that will compute the histogram of the input texture and write the result to the storage buffer + const shader = device.supportsCompute ? new pc.Shader(device, { + name: 'ComputeShader', + shaderLanguage: pc.SHADERLANGUAGE_WGSL, + cshader: files['compute-shader.wgsl'], + + // format of a bind group, providing resources for the compute shader + computeBindGroupFormat: new pc.BindGroupFormat(device, [ + // no uniform buffer + ], [ + // input texture - the scene color map + new pc.BindTextureFormat('uSceneColorMap', pc.SHADERSTAGE_COMPUTE) + ], [ + // no storage textures + ], [ + // output storage buffer + new pc.BindStorageBufferFormat('outBuffer', pc.SHADERSTAGE_COMPUTE) + ]) + }) : null; + + // Create a storage buffer to which the compute shader will write the histogram values. + const numBins = 256; + const histogramStorageBuffer = new pc.StorageBuffer( + device, numBins * 4, // 4 bytes per value, storing unsigned int + pc.BUFFERUSAGE_COPY_SRC | // needed for reading back the data to CPU + pc.BUFFERUSAGE_COPY_DST // needed for clearing the buffer + ); + + // Create an instance of the compute shader, and set the input and output data. Note that we do + // not provide a value for `uSceneColorMap` as this is done by the engine internally. + const compute = new pc.Compute(device, shader, 'ComputeHistogram'); + compute.setParameter('outBuffer', histogramStorageBuffer); + + // instantiate the spinning mesh + const solid = assets.solid.resource.instantiateRenderEntity(); + solid.addComponent('script'); + solid.script.create('rotator'); + solid.setLocalPosition(0, 0.4, 0); + solid.setLocalScale(0.35, 0.35, 0.35); + app.root.addChild(solid); + + let firstFrame = true; + app.on('update', function (/** @type {number} */ dt) { + + // The update function runs every frame before the frame gets rendered. On the first time it + // runs, the scene color map has not been rendered yet, so we skip the first frame. + if (firstFrame) { + firstFrame = false; + return; + } + + if (device.supportsCompute) { + + // clear the storage buffer, to avoid the accumulation buildup + histogramStorageBuffer.clear(); + + // dispatch the compute shader + compute.setupDispatch(app.graphicsDevice.width, app.graphicsDevice.height); + device.computeDispatch([compute]); + + // Read back the histogram data from the storage buffer. None that the returned promise + // will be resolved later, when the GPU is done running it, and so the histogram on the + // screen will be up to few frames behind. + const histogramData = new Uint32Array(numBins); + histogramStorageBuffer.read(0, undefined, histogramData).then( + (data) => { + // render the histogram using lines + const scale = 1 / 50000; + const positions = []; + for (let x = 0; x < data.length; x++) { + const value = pc.math.clamp(data[x] * scale, 0, 0.2); + positions.push(x * 0.001, -0.35, 4); + positions.push(x * 0.001, value - 0.35, 4); + } + app.drawLineArrays(positions, pc.Color.YELLOW); + } + ); + } + }); +}); + +export { app }; diff --git a/examples/src/examples/compute/texture-gen/example.mjs b/examples/src/examples/compute/texture-gen/example.mjs index 279f6959eea..5575328653d 100644 --- a/examples/src/examples/compute/texture-gen/example.mjs +++ b/examples/src/examples/compute/texture-gen/example.mjs @@ -136,7 +136,7 @@ assetListLoader.load(() => { }); // create an instance of the compute shader, and set the input and output textures - const compute = new pc.Compute(device, shader); + const compute = new pc.Compute(device, shader, 'ComputeModifyTexture'); compute.setParameter('inTexture', assets.texture.resource); compute.setParameter('outTexture', storageTexture); diff --git a/src/core/constants.js b/src/core/constants.js index 11911e4cb74..7dd5392065a 100644 --- a/src/core/constants.js +++ b/src/core/constants.js @@ -95,6 +95,14 @@ export const TRACEID_VRAM_VB = 'VRAM.Vb'; */ export const TRACEID_VRAM_IB = 'VRAM.Ib'; +/** + * Logs the vram use by the storage buffers. + * + * @type {string} + * @category Debug + */ +export const TRACEID_VRAM_SB = 'VRAM.Sb'; + /** * Logs the creation of bind groups. * diff --git a/src/index.js b/src/index.js index 16601db8251..c141a949bb2 100644 --- a/src/index.js +++ b/src/index.js @@ -72,7 +72,7 @@ export * from './platform/audio/constants.js'; // PLATFORM / GRAPHICS export * from './platform/graphics/constants.js'; export { createGraphicsDevice } from './platform/graphics/graphics-device-create.js'; -export { BindGroupFormat, BindBufferFormat, BindTextureFormat, BindStorageTextureFormat } from './platform/graphics/bind-group-format.js'; +export { BindGroupFormat, BindBufferFormat, BindTextureFormat, BindStorageTextureFormat, BindStorageBufferFormat } from './platform/graphics/bind-group-format.js'; export { BlendState } from './platform/graphics/blend-state.js'; export { Compute } from './platform/graphics/compute.js'; export { DepthState } from './platform/graphics/depth-state.js'; @@ -85,6 +85,7 @@ export { ScopeSpace } from './platform/graphics/scope-space.js'; export { Shader } from './platform/graphics/shader.js'; export { ShaderProcessorOptions } from './platform/graphics/shader-processor-options.js'; // used by splats in extras export { ShaderUtils } from './platform/graphics/shader-utils.js'; // used by splats in extras +export { StorageBuffer } from './platform/graphics/storage-buffer.js'; export { Texture } from './platform/graphics/texture.js'; export { TextureUtils } from './platform/graphics/texture-utils.js'; export { TransformFeedback } from './platform/graphics/transform-feedback.js'; diff --git a/src/platform/graphics/bind-group-format.js b/src/platform/graphics/bind-group-format.js index 8b634519f10..6dc1476681c 100644 --- a/src/platform/graphics/bind-group-format.js +++ b/src/platform/graphics/bind-group-format.js @@ -28,6 +28,22 @@ class BindBufferFormat { } } +/** + * @ignore + */ +class BindStorageBufferFormat { + /** @type {import('./scope-id.js').ScopeId} */ + scopeId; + + constructor(name, visibility) { + /** @type {string} */ + this.name = name; + + // SHADERSTAGE_VERTEX, SHADERSTAGE_FRAGMENT, SHADERSTAGE_COMPUTE + this.visibility = visibility; + } +} + /** * @ignore */ @@ -73,8 +89,6 @@ class BindStorageTextureFormat { * @ignore */ class BindGroupFormat { - compute = false; - /** * @param {import('./graphics-device.js').GraphicsDevice} graphicsDevice - The graphics device * used to manage this vertex format. @@ -84,8 +98,10 @@ class BindGroupFormat { * Defaults to an empty array. * @param {BindStorageTextureFormat[]} [storageTextureFormats] - An array of bind storage texture * formats (storage textures), used by the compute shader. Defaults to an empty array. + * @param {BindStorageBufferFormat[]} [storageBufferFormats] - An array of bind storage buffer + * formats. Defaults to an empty array. */ - constructor(graphicsDevice, bufferFormats = [], textureFormats = [], storageTextureFormats = []) { + constructor(graphicsDevice, bufferFormats = [], textureFormats = [], storageTextureFormats = [], storageBufferFormats = []) { this.id = id++; DebugHelper.setName(this, `BindGroupFormat_${this.id}`); @@ -127,6 +143,19 @@ class BindGroupFormat { tf.scopeId = scope.resolve(tf.name); }); + /** @type {BindStorageBufferFormat[]} */ + this.storageBufferFormats = storageBufferFormats; + + // maps a storage buffer format name to a slot index + /** @type {Map} */ + this.storageBufferFormatsMap = new Map(); + storageBufferFormats.forEach((bf, i) => { + this.storageBufferFormatsMap.set(bf.name, i); + + // resolve scope id + bf.scopeId = scope.resolve(bf.name); + }); + this.impl = graphicsDevice.createBindGroupFormatImpl(this); Debug.trace(TRACEID_BINDGROUPFORMAT_ALLOC, `Alloc: Id ${this.id}`, this); @@ -205,4 +234,4 @@ class BindGroupFormat { } } -export { BindBufferFormat, BindTextureFormat, BindGroupFormat, BindStorageTextureFormat }; +export { BindBufferFormat, BindTextureFormat, BindGroupFormat, BindStorageTextureFormat, BindStorageBufferFormat }; diff --git a/src/platform/graphics/bind-group.js b/src/platform/graphics/bind-group.js index a029dc9244f..87c56a3804b 100644 --- a/src/platform/graphics/bind-group.js +++ b/src/platform/graphics/bind-group.js @@ -6,8 +6,8 @@ import { DebugGraphics } from './debug-graphics.js'; let id = 0; /** - * A bind group represents an collection of {@link UniformBuffer} and {@link Texture} instance, - * which can be bind on a GPU for rendering. + * A bind group represents a collection of {@link UniformBuffer}, {@link Texture} and + * {@link StorageBuffer} instanced, which can be bind on a GPU for rendering. * * @ignore */ @@ -50,6 +50,7 @@ class BindGroup { this.textures = []; this.storageTextures = []; + this.storageBuffers = []; this.uniformBuffers = []; /** @type {import('./uniform-buffer.js').UniformBuffer} */ @@ -87,6 +88,22 @@ class BindGroup { } } + /** + * Assign a storage buffer to a slot. + * + * @param {string} name - The name of the storage buffer slot. + * @param {import('./storage-buffer.js').StorageBuffer} storageBuffer - The storage buffer to + * assign to the slot. + */ + setStorageBuffer(name, storageBuffer) { + const index = this.format.storageBufferFormatsMap.get(name); + Debug.assert(index !== undefined, `Setting a storage buffer [${name}] on a bind group with id: ${this.id} which does not contain in, while rendering [${DebugGraphics.toString()}]`, this); + if (this.storageBuffers[index] !== storageBuffer) { + this.storageBuffers[index] = storageBuffer; + this.dirty = true; + } + } + /** * Assign a texture to a named slot. * @@ -129,7 +146,7 @@ class BindGroup { update() { // TODO: implement faster version of this, which does not call SetTexture, which does a map lookup - const { textureFormats, storageTextureFormats } = this.format; + const { textureFormats, storageTextureFormats, storageBufferFormats } = this.format; for (let i = 0; i < textureFormats.length; i++) { const textureFormat = textureFormats[i]; @@ -145,6 +162,13 @@ class BindGroup { this.setStorageTexture(storageTextureFormat.name, value); } + for (let i = 0; i < storageBufferFormats.length; i++) { + const storageBufferFormat = storageBufferFormats[i]; + const value = storageBufferFormat.scopeId.value; + Debug.assert(value, `Value was not set when assigning storage buffer slot [${storageBufferFormat.name}] to a bind group, while rendering [${DebugGraphics.toString()}]`, this); + this.setStorageBuffer(storageBufferFormat.name, value); + } + // update uniform buffer offsets this.uniformBufferOffsets.length = this.uniformBuffers.length; for (let i = 0; i < this.uniformBuffers.length; i++) { diff --git a/src/platform/graphics/compute.js b/src/platform/graphics/compute.js index df8e827f774..e9f9a657018 100644 --- a/src/platform/graphics/compute.js +++ b/src/platform/graphics/compute.js @@ -24,6 +24,9 @@ class Compute { */ shader = null; + /** @type {string} */ + name; + /** * @type {Map} * @ignore @@ -55,10 +58,12 @@ class Compute { * @param {import('./graphics-device.js').GraphicsDevice} graphicsDevice - * The graphics device. * @param {import('./shader.js').Shader} shader - The compute shader. + * @param {string} [name] - The name of the compute instance, used for debugging only. */ - constructor(graphicsDevice, shader) { + constructor(graphicsDevice, shader, name = 'Unnamed') { this.device = graphicsDevice; this.shader = shader; + this.name = name; if (graphicsDevice.supportsCompute) { this.impl = graphicsDevice.createComputeImpl(this); @@ -69,8 +74,8 @@ class Compute { * Sets a shader parameter on a compute instance. * * @param {string} name - The name of the parameter to set. - * @param {number|number[]|Float32Array|import('./texture.js').Texture} value - The value for - * the specified parameter. + * @param {number|number[]|Float32Array|import('./texture.js').Texture|import('./storage-buffer.js').StorageBuffer} value + * - The value for the specified parameter. */ setParameter(name, value) { let param = this.parameters.get(name); @@ -86,8 +91,8 @@ class Compute { * Returns the value of a shader parameter from the compute instance. * * @param {string} name - The name of the parameter to get. - * @returns {number|number[]|Float32Array|import('./texture.js').Texture|undefined} The value of the - * specified parameter. + * @returns {number|number[]|Float32Array|import('./texture.js').Texture|import('./storage-buffer.js').StorageBuffer|undefined} + * The value of the specified parameter. */ getParameter(name) { return this.parameters.get(name)?.value; diff --git a/src/platform/graphics/constants.js b/src/platform/graphics/constants.js index 9f79399375c..8cea3dc2764 100644 --- a/src/platform/graphics/constants.js +++ b/src/platform/graphics/constants.js @@ -167,6 +167,59 @@ export const BLENDEQUATION_MIN = 3; */ export const BLENDEQUATION_MAX = 4; +/** + * A flag utilized during the construction of a {@link StorageBuffer} to make it available for read + * access by CPU. + * + * @type {number} + * @category Graphics + */ +export const BUFFERUSAGE_READ = 0x0001; + +/** + * A flag utilized during the construction of a {@link StorageBuffer} to make it available for write + * access by CPU. + * + * @type {number} + * @category Graphics + */ +export const BUFFERUSAGE_WRITE = 0x0002; + +/** + * A flag utilized during the construction of a {@link StorageBuffer} to ensure its compatibility + * when used as a source of a copy operation. + * + * @type {number} + * @category Graphics + */ +export const BUFFERUSAGE_COPY_SRC = 0x0004; + +/** + * A flag utilized during the construction of a {@link StorageBuffer} to ensure its compatibility + * when used as a destination of a copy operation, or as a target of a write operation. + * + * @type {number} + * @category Graphics + */ +export const BUFFERUSAGE_COPY_DST = 0x0008; + +// internal flags +export const BUFFERUSAGE_INDEX = 0x0010; +export const BUFFERUSAGE_VERTEX = 0x0020; +export const BUFFERUSAGE_UNIFORM = 0x0040; +export const BUFFERUSAGE_STORAGE = 0x0080; + +/** + * A flag utilized during the construction of a {@link StorageBuffer} to allow it to store indirect + * command arguments. + * TODO: This flag is hidden till the feature is implemented. + * + * @type {number} + * @category Graphics + * @ignore + */ +export const BUFFERUSAGE_INDIRECT = 0x0100; + /** * The data store contents will be modified once and used many times. * diff --git a/src/platform/graphics/graphics-device.js b/src/platform/graphics/graphics-device.js index 2ad5bc93b02..d70d37fa2ee 100644 --- a/src/platform/graphics/graphics-device.js +++ b/src/platform/graphics/graphics-device.js @@ -402,7 +402,8 @@ class GraphicsDevice extends EventHandler { tex: 0, vb: 0, ib: 0, - ub: 0 + ub: 0, + sb: 0 }; this._shaderStats = { diff --git a/src/platform/graphics/storage-buffer.js b/src/platform/graphics/storage-buffer.js new file mode 100644 index 00000000000..263af1329fb --- /dev/null +++ b/src/platform/graphics/storage-buffer.js @@ -0,0 +1,69 @@ +import { Debug } from '../../core/debug.js'; +import { TRACEID_VRAM_SB } from '../../core/constants.js'; +import { BUFFERUSAGE_STORAGE } from './constants.js'; + +let id = 0; + +/** + * A storage buffer represents a memory which both the CPU and the GPU can access. Typically it is + * used to provide data for compute shader, and to store the result of the computation. + * Note that this class is only supported on the WebGPU platform. + * + * @category Graphics + */ +class StorageBuffer { + id = id++; + + /** + * Create a new StorageBuffer instance. + * + * @param {import('./graphics-device.js').GraphicsDevice} graphicsDevice - The graphics device + * used to manage this storage buffer. + * @param {number} byteSize - The size of the storage buffer in bytes. + * @param {number} bufferUsage - The usage type of the storage buffer. Can be a combination of + * {@link BUFFERUSAGE_READ}, {@link BUFFERUSAGE_WRITE}, {@link BUFFERUSAGE_COPY_SRC} and + * {@link BUFFERUSAGE_COPY_DST} flags. + */ + constructor(graphicsDevice, byteSize, bufferUsage) { + this.device = graphicsDevice; + this.byteSize = byteSize; + this.bufferUsage = bufferUsage; + + this.impl = graphicsDevice.createBufferImpl(BUFFERUSAGE_STORAGE | bufferUsage); + this.impl.allocate(graphicsDevice, byteSize); + this.device.buffers.push(this); + + this.adjustVramSizeTracking(graphicsDevice._vram, this.byteSize); + } + + /** + * Frees resources associated with this storage buffer. + */ + destroy() { + + // stop tracking the buffer + const device = this.device; + const idx = device.buffers.indexOf(this); + if (idx !== -1) { + device.buffers.splice(idx, 1); + } + + this.adjustVramSizeTracking(device._vram, -this.byteSize); + this.impl.destroy(device); + } + + adjustVramSizeTracking(vram, size) { + Debug.trace(TRACEID_VRAM_SB, `${this.id} size: ${size} vram.sb: ${vram.sb} => ${vram.sb + size}`); + vram.sb += size; + } + + read(offset = 0, size = this.byteSize, data = null) { + return this.impl.read(this.device, offset, size, data); + } + + clear(offset = 0, size = this.byteSize) { + this.impl.clear(this.device, offset, size); + } +} + +export { StorageBuffer }; diff --git a/src/platform/graphics/webgpu/webgpu-bind-group-format.js b/src/platform/graphics/webgpu/webgpu-bind-group-format.js index 17446ceb53e..db086e244f0 100644 --- a/src/platform/graphics/webgpu/webgpu-bind-group-format.js +++ b/src/platform/graphics/webgpu/webgpu-bind-group-format.js @@ -199,6 +199,24 @@ class WebgpuBindGroupFormat { }); }); + // storage buffers + bindGroupFormat.storageBufferFormats.forEach((bufferFormat) => { + + const readOnly = false; + const visibility = WebgpuUtils.shaderStage(bufferFormat.visibility); + key += `#${index}SB:${visibility}-${readOnly ? 'ro' : 'rw'}`; + + entries.push({ + binding: index++, + visibility: visibility, + buffer: { + + // "storage", "read-only-storage" + type: readOnly ? 'read-only-storage' : 'storage' + } + }); + }); + /** @type {GPUBindGroupLayoutDescriptor} */ const descr = { entries: entries diff --git a/src/platform/graphics/webgpu/webgpu-bind-group.js b/src/platform/graphics/webgpu/webgpu-bind-group.js index 8c4f5f53705..15980414b91 100644 --- a/src/platform/graphics/webgpu/webgpu-bind-group.js +++ b/src/platform/graphics/webgpu/webgpu-bind-group.js @@ -128,6 +128,24 @@ class WebgpuBindGroup { }); }); + // storage buffers + bindGroup.storageBuffers.forEach((buffer, bufferIndex) => { + /** @type {GPUBuffer} */ + const wgpuBuffer = buffer.impl.buffer; + + Debug.assert(wgpuBuffer, 'NULL storage buffer cannot be used by the bind group'); + Debug.call(() => { + this.debugFormat += `${index}: SB\n`; + }); + + entries.push({ + binding: index++, + resource: { + buffer: wgpuBuffer + } + }); + }); + const descr = { layout: bindGroup.format.impl.bindGroupLayout, entries: entries diff --git a/src/platform/graphics/webgpu/webgpu-buffer.js b/src/platform/graphics/webgpu/webgpu-buffer.js index 994888cc9bf..20be5a88b6c 100644 --- a/src/platform/graphics/webgpu/webgpu-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-buffer.js @@ -8,11 +8,17 @@ import { Debug, DebugHelper } from '../../../core/debug.js'; */ class WebgpuBuffer { /** - * @type {GPUBuffer} + * @type {GPUBuffer|null} * @private */ buffer = null; + usageFlags = 0; + + constructor(usageFlags = 0) { + this.usageFlags = usageFlags; + } + destroy(device) { if (this.buffer) { this.buffer.destroy(); @@ -27,13 +33,19 @@ class WebgpuBuffer { loseContext() { } + allocate(device, size) { + Debug.assert(!this.buffer, "Buffer already allocated"); + this.buffer = device.wgpu.createBuffer({ + size, + usage: this.usageFlags + }); + } + /** * @param {import('./webgpu-graphics-device.js').WebgpuGraphicsDevice} device - Graphics device. - * @param {*} usage - - * @param {*} target - * @param {*} storage - */ - unlock(device, usage, target, storage) { + unlock(device, storage) { const wgpu = device.wgpu; @@ -42,18 +54,18 @@ class WebgpuBuffer { if (!this.buffer) { // size needs to be a multiple of 4 + // note: based on specs, descriptor.size must be a multiple of 4 if descriptor.mappedAtCreation is true const size = (storage.byteLength + 3) & ~3; - this.buffer = device.wgpu.createBuffer({ - size: size, - usage: target | GPUBufferUsage.COPY_DST - }); + this.usageFlags |= GPUBufferUsage.COPY_DST; + this.allocate(device, size); DebugHelper.setLabel(this.buffer, - target & GPUBufferUsage.VERTEX ? 'VertexBuffer' : - target & GPUBufferUsage.INDEX ? 'IndexBuffer' : - target & GPUBufferUsage.UNIFORM ? "UniformBuffer" : - '' + this.usageFlags & GPUBufferUsage.VERTEX ? 'VertexBuffer' : + this.usageFlags & GPUBufferUsage.INDEX ? 'IndexBuffer' : + this.usageFlags & GPUBufferUsage.UNIFORM ? "UniformBuffer" : + this.usageFlags & GPUBufferUsage.STORAGE ? "StorageBuffer" : + '' ); @@ -80,9 +92,14 @@ class WebgpuBuffer { // copy data to the gpu buffer Debug.trace(TRACEID_RENDER_QUEUE, `writeBuffer: ${this.buffer.label}`); wgpu.queue.writeBuffer(this.buffer, 0, data, 0, data.length); + } + + read(device, offset, size, data) { + return device.readStorageBuffer(this, offset, size, data); + } - // TODO: handle usage types: - // - BUFFER_STATIC, BUFFER_DYNAMIC, BUFFER_STREAM, BUFFER_GPUDYNAMIC + clear(device, offset, size) { + device.clearStorageBuffer(this, offset, size); } } diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index b7dd8be08de..161b0dd073b 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -1,5 +1,6 @@ import { Debug, DebugHelper } from "../../../core/debug.js"; import { BindGroup } from "../bind-group.js"; +import { DebugGraphics } from "../debug-graphics.js"; import { UniformBuffer } from "../uniform-buffer.js"; /** @@ -13,6 +14,8 @@ class WebgpuCompute { const { device, shader } = compute; + DebugGraphics.pushGpuMarker(device, `Compute:${compute.name}`); + // create bind group const { computeBindGroupFormat, computeUniformBufferFormat } = shader.impl; Debug.assert(computeBindGroupFormat, 'Compute shader does not have computeBindGroupFormat specified', shader); @@ -27,6 +30,8 @@ class WebgpuCompute { // pipeline this.pipeline = device.computePipeline.get(shader, computeBindGroupFormat); + + DebugGraphics.popGpuMarker(device); } updateBindGroup() { diff --git a/src/platform/graphics/webgpu/webgpu-graphics-device.js b/src/platform/graphics/webgpu/webgpu-graphics-device.js index 26098f4ec8e..75d1faeb4f2 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -3,7 +3,8 @@ import { Debug, DebugHelper } from '../../../core/debug.js'; import { path } from '../../../core/path.js'; import { - PIXELFORMAT_RGBA32F, PIXELFORMAT_RGBA8, PIXELFORMAT_BGRA8, DEVICETYPE_WEBGPU + PIXELFORMAT_RGBA32F, PIXELFORMAT_RGBA8, PIXELFORMAT_BGRA8, DEVICETYPE_WEBGPU, + BUFFERUSAGE_READ, BUFFERUSAGE_COPY_DST } from '../constants.js'; import { GraphicsDevice } from '../graphics-device.js'; import { DebugGraphics } from '../debug-graphics.js'; @@ -27,6 +28,7 @@ import { WebgpuDynamicBuffers } from './webgpu-dynamic-buffers.js'; import { WebgpuGpuProfiler } from './webgpu-gpu-profiler.js'; import { WebgpuResolver } from './webgpu-resolver.js'; import { WebgpuCompute } from './webgpu-compute.js'; +import { WebgpuBuffer } from './webgpu-buffer.js'; class WebgpuGraphicsDevice extends GraphicsDevice { /** @@ -377,6 +379,10 @@ class WebgpuGraphicsDevice extends GraphicsDevice { } } + createBufferImpl(usageFlags) { + return new WebgpuBuffer(usageFlags); + } + createUniformBufferImpl(uniformBuffer) { return new WebgpuUniformBuffer(uniformBuffer); } @@ -817,6 +823,104 @@ class WebgpuGraphicsDevice extends GraphicsDevice { } } + /** + * Clear the content of a storage buffer to 0. + * + * @param {import('./webgpu-buffer.js').WebgpuBuffer} storageBuffer - The storage buffer. + * @param {number} [offset] - The offset of data to clear. Defaults to 0. + * @param {number} [size] - The size of data to clear. Defaults to the full size of the buffer. + * @ignore + */ + clearStorageBuffer(storageBuffer, offset = 0, size = storageBuffer.byteSize) { + + // use existing or create new encoder + const commandEncoder = this.commandEncoder ?? this.wgpu.createCommandEncoder(); + + commandEncoder.clearBuffer(storageBuffer.buffer, offset, size); + + // if we created the encoder + if (!this.commandEncoder) { + DebugHelper.setLabel(commandEncoder, 'ReadStorageBuffer-Encoder'); + const cb = commandEncoder.finish(); + DebugHelper.setLabel(cb, 'ReadStorageBuffer-CommandBuffer'); + this.addCommandBuffer(cb); + } + } + + /** + * Read a content of a storage buffer. + * + * @param {import('./webgpu-buffer.js').WebgpuBuffer} storageBuffer - The storage buffer. + * @param {number} [offset] - The offset of data to read. Defaults to 0. + * @param {number} [size] - The size of data to read. Defaults to the full size of the buffer. + * @param {ArrayBufferView} [data] - Typed array to populate with the data read from the storage + * buffer. When typed array is supplied, enough space needs to be reserved, otherwise only + * partial data is copied. If not specified, the data is returned in an Uint8Array. Defaults to + * null. + * @param {boolean} [immediate] - If true, the read operation will be executed as soon as + * possible. This has a performance impact, so it should be used only when necessary. Defaults + * to false. + * @returns {Promise} A promise that resolves with the data read from the storage + * buffer. + * @ignore + */ + readStorageBuffer(storageBuffer, offset = 0, size = storageBuffer.byteSize, data = null, immediate = false) { + + // create a temporary staging buffer + const stagingBuffer = this.createBufferImpl(BUFFERUSAGE_READ | BUFFERUSAGE_COPY_DST); + stagingBuffer.allocate(this, size); + const destBuffer = stagingBuffer.buffer; + + // use existing or create new encoder + const commandEncoder = this.commandEncoder ?? this.wgpu.createCommandEncoder(); + + // copy the GPU buffer to the staging buffer + commandEncoder.copyBufferToBuffer(storageBuffer.buffer, offset, destBuffer, 0, size); + + // if we created new encoder + if (!this.commandEncoder) { + DebugHelper.setLabel(commandEncoder, 'ReadStorageBuffer-Encoder'); + const cb = commandEncoder.finish(); + DebugHelper.setLabel(cb, 'ReadStorageBuffer-CommandBuffer'); + this.addCommandBuffer(cb); + } + + // return a promise that resolves with the data + return new Promise((resolve, reject) => { + + const read = () => { + + destBuffer?.mapAsync(GPUMapMode.READ).then(() => { + + // copy data to a buffer + data ??= new Uint8Array(size); + const copySrc = destBuffer.getMappedRange(0, size); + + // use the same type as the target + const srcType = data.constructor; + data.set(new srcType(copySrc)); + + // release staging buffer + destBuffer.unmap(); + stagingBuffer.destroy(this); + + resolve(data); + }); + }; + + if (immediate) { + // submit the command buffer immediately + this.submit(); + read(); + } else { + // map the buffer during the next event handling cycle, when the command buffer is submitted + setTimeout(() => { + read(); + }); + } + }); + } + /** * Copies source render target into destination render target. Mostly used by post-effects. * @@ -837,7 +941,6 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // use existing or create new encoder if not in a render pass const commandEncoder = this.commandEncoder ?? this.wgpu.createCommandEncoder(); - DebugHelper.setLabel(commandEncoder, 'CopyRenderTarget-Encoder'); DebugGraphics.pushGpuMarker(this, 'COPY-RT'); @@ -900,6 +1003,8 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // if we created the encoder if (!this.commandEncoder) { + DebugHelper.setLabel(commandEncoder, 'CopyRenderTarget-Encoder'); + // copy operation runs next const cb = commandEncoder.finish(); DebugHelper.setLabel(cb, 'CopyRenderTarget-CommandBuffer'); diff --git a/src/platform/graphics/webgpu/webgpu-index-buffer.js b/src/platform/graphics/webgpu/webgpu-index-buffer.js index 7ffd076bfe5..8d24dc21041 100644 --- a/src/platform/graphics/webgpu/webgpu-index-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-index-buffer.js @@ -1,5 +1,5 @@ import { Debug } from '../../../core/debug.js'; -import { INDEXFORMAT_UINT8, INDEXFORMAT_UINT16 } from '../constants.js'; +import { INDEXFORMAT_UINT8, INDEXFORMAT_UINT16, BUFFERUSAGE_INDEX } from '../constants.js'; import { WebgpuBuffer } from "./webgpu-buffer.js"; /** @@ -11,7 +11,7 @@ class WebgpuIndexBuffer extends WebgpuBuffer { format = null; constructor(indexBuffer) { - super(); + super(BUFFERUSAGE_INDEX); Debug.assert(indexBuffer.format !== INDEXFORMAT_UINT8, "WebGPU does not support 8-bit index buffer format"); this.format = indexBuffer.format === INDEXFORMAT_UINT16 ? "uint16" : "uint32"; @@ -19,7 +19,7 @@ class WebgpuIndexBuffer extends WebgpuBuffer { unlock(indexBuffer) { const device = indexBuffer.device; - super.unlock(device, indexBuffer.usage, GPUBufferUsage.INDEX, indexBuffer.storage); + super.unlock(device, indexBuffer.storage); } } diff --git a/src/platform/graphics/webgpu/webgpu-uniform-buffer.js b/src/platform/graphics/webgpu/webgpu-uniform-buffer.js index 43db01b6762..766b28c4c33 100644 --- a/src/platform/graphics/webgpu/webgpu-uniform-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-uniform-buffer.js @@ -1,3 +1,4 @@ +import { BUFFERUSAGE_UNIFORM } from "../constants.js"; import { WebgpuBuffer } from "./webgpu-buffer.js"; /** @@ -7,21 +8,13 @@ import { WebgpuBuffer } from "./webgpu-buffer.js"; */ class WebgpuUniformBuffer extends WebgpuBuffer { constructor(uniformBuffer) { - super(); - } - - destroy(device) { - - super.destroy(device); - - - // TODO: clear up bound uniform buffers + super(BUFFERUSAGE_UNIFORM); } unlock(uniformBuffer) { const device = uniformBuffer.device; - super.unlock(device, undefined, GPUBufferUsage.UNIFORM, uniformBuffer.storageInt32.buffer); + super.unlock(device, uniformBuffer.storageInt32.buffer); } } diff --git a/src/platform/graphics/webgpu/webgpu-vertex-buffer.js b/src/platform/graphics/webgpu/webgpu-vertex-buffer.js index 73a8a6efcc7..cb7ba111a00 100644 --- a/src/platform/graphics/webgpu/webgpu-vertex-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-vertex-buffer.js @@ -1,3 +1,4 @@ +import { BUFFERUSAGE_VERTEX } from "../constants.js"; import { WebgpuBuffer } from "./webgpu-buffer.js"; /** @@ -7,20 +8,13 @@ import { WebgpuBuffer } from "./webgpu-buffer.js"; */ class WebgpuVertexBuffer extends WebgpuBuffer { constructor(vertexBuffer, format) { - super(); - } - - destroy(device) { - - super.destroy(device); - - // TODO: clear up bound vertex buffers + super(BUFFERUSAGE_VERTEX); } unlock(vertexBuffer) { const device = vertexBuffer.device; - super.unlock(device, vertexBuffer.usage, GPUBufferUsage.VERTEX, vertexBuffer.storage); + super.unlock(device, vertexBuffer.storage); } } diff --git a/src/scene/renderer/render-pass-forward.js b/src/scene/renderer/render-pass-forward.js index 0b04495adaf..5ff87e9737f 100644 --- a/src/scene/renderer/render-pass-forward.js +++ b/src/scene/renderer/render-pass-forward.js @@ -237,7 +237,7 @@ class RenderPassForward extends RenderPass { const { layer, transparent, camera } = renderAction; const cameraPass = layerComposition.camerasMap.get(camera); - DebugGraphics.pushGpuMarker(this.device, camera ? camera.entity.name : 'noname'); + DebugGraphics.pushGpuMarker(this.device, camera ? camera.entity.name : 'Unnamed'); DebugGraphics.pushGpuMarker(this.device, `${layer.name}(${transparent ? 'TRANSP' : 'OPAQUE'})`); // #if _PROFILER