From c628f05c3c696d607545297e0bd080c98ebe5fe0 Mon Sep 17 00:00:00 2001 From: Erik Dubbelboer Date: Sun, 28 Jan 2024 15:41:09 +0100 Subject: [PATCH 1/4] First working version of the example --- .../examples/misc/compute-shader/config.mjs | 18 +++++ .../examples/misc/compute-shader/example.mjs | 72 +++++++++++++++++++ src/platform/graphics/compute.js | 8 +++ .../graphics/webgpu/webgpu-compute.js | 61 ++++++++++++++++ .../graphics/webgpu/webgpu-graphics-device.js | 12 +++- 5 files changed, 168 insertions(+), 3 deletions(-) create mode 100644 examples/src/examples/misc/compute-shader/config.mjs create mode 100644 examples/src/examples/misc/compute-shader/example.mjs diff --git a/examples/src/examples/misc/compute-shader/config.mjs b/examples/src/examples/misc/compute-shader/config.mjs new file mode 100644 index 00000000000..07ca00f7f33 --- /dev/null +++ b/examples/src/examples/misc/compute-shader/config.mjs @@ -0,0 +1,18 @@ +/** + * @type {import('../../../../types.mjs').ExampleConfig} + */ +export default { + WEBGPU_REQUIRED: true, + HIDDEN: true, + FILES: { + 'shader.wgsl': ` + @group(0) @binding(0) var outputTexture: texture_storage_2d; + + @compute @workgroup_size(1, 1, 1) + fn main(@builtin(global_invocation_id) global_id : vec3u) { + let clearColor: vec4 = vec4(0.5); + textureStore(outputTexture, vec2(global_id.xy), clearColor); + } + ` + } +}; diff --git a/examples/src/examples/misc/compute-shader/example.mjs b/examples/src/examples/misc/compute-shader/example.mjs new file mode 100644 index 00000000000..bad057a3ef8 --- /dev/null +++ b/examples/src/examples/misc/compute-shader/example.mjs @@ -0,0 +1,72 @@ +import * as pc from 'playcanvas'; +import files from '@examples/files'; +import { deviceType, rootPath } from '@examples/utils'; + +const canvas = document.getElementById('application-canvas'); +if (!(canvas instanceof HTMLCanvasElement)) { + throw new Error('No canvas found'); +} + +const gfxOptions = { + deviceTypes: [deviceType], + + // Even though we're using WGSL, we still need to provide glslang + // and twgsl to compile shaders used internally by the engine. + 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]; +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 texture = new pc.Texture(app.graphicsDevice, { + name: 'outputTexture', + width: 2, + height: 2, + format: pc.PIXELFORMAT_RGBA8, + mipmaps: false, + storage: true +}); + +app.graphicsDevice.scope.resolve("outputTexture").setValue(texture); + +const shaderDefinition = { + cshader: files['shader.wgsl'], + shaderLanguage: pc.SHADERLANGUAGE_WGSL, +}; +const shader = new pc.Shader(app.graphicsDevice, shaderDefinition); + +shader.impl.computeBindGroupFormat = new pc.BindGroupFormat(device,[], [], [ + new pc.BindStorageTextureFormat('outputTexture', pc.PIXELFORMAT_RGBA8, pc.TEXTUREDIMENSION_2D), +], { + compute: true +}); + +const compute = new pc.Compute(app.graphicsDevice, shader); + +compute.dispatch(texture.width, texture.height); + +const data = await compute.read(texture.impl.gpuTexture); + +console.log(data); + +export { app }; diff --git a/src/platform/graphics/compute.js b/src/platform/graphics/compute.js index 8f9f3b0ed30..3b18e766b92 100644 --- a/src/platform/graphics/compute.js +++ b/src/platform/graphics/compute.js @@ -39,6 +39,14 @@ class Compute { dispatch(x, y, z) { this.impl?.dispatch(x, y, z); } + + /** + * + * @param {GPUTexture} texture + */ + async read(texture) { + return this.impl?.read(texture); + } } export { Compute }; diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index 2a1d91d93f9..786b67d5d64 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -41,6 +41,67 @@ class WebgpuCompute { device.endComputePass(); } + + /** + * + * @param {GPUTexture} texture + */ + async read(texture) { + const device = this.compute.device; + device.startCompute(); + + // bind group data + const { bindGroup } = this; + bindGroup.update(); + device.setBindGroup(0, bindGroup); + + // Calculate bytes per pixel, assuming RGBA8 format (4 bytes per pixel) + const bytesPerPixel = 4; + + // Calculate bytes per row, ensuring it's a multiple of 256 + const bytesPerRow = Math.ceil((texture.width * bytesPerPixel) / 256) * 256; + + // Calculate the size of the buffer to hold the texture data + const bufferSize = bytesPerRow * texture.height; + + const gpuBuffer = device.wgpu.createBuffer({ + size: bufferSize, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ + }); + + const textureCopyView = { + texture, + origin: { x: 0, y: 0 }, + }; + const bufferCopyView = { + buffer: gpuBuffer, + bytesPerRow: bytesPerRow, + }; + const extent = { + width: texture.width, + height: texture.height, + }; + + // Encode command to copy from texture to buffer + device.commandEncoder.copyTextureToBuffer(textureCopyView, bufferCopyView, extent); + + device.endCompute(); + + await device.wgpu.queue.onSubmittedWorkDone(); + + // Ensure that the GPU operations are complete + await gpuBuffer.mapAsync(GPUMapMode.READ); + + // Read buffer contents + const arrayBuffer = gpuBuffer.getMappedRange(); + const data = new Uint8Array(arrayBuffer); // or another typed array based on the texture format + + // Cleanup + //gpuBuffer.unmap(); + //gpuBuffer.destroy(); + + return data; + } } export { WebgpuCompute }; diff --git a/src/platform/graphics/webgpu/webgpu-graphics-device.js b/src/platform/graphics/webgpu/webgpu-graphics-device.js index 4346064f3e7..6d3148d1940 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -654,8 +654,7 @@ class WebgpuGraphicsDevice extends GraphicsDevice { WebgpuDebug.end(this, { renderPass }); } - startComputePass() { - + startCompute() { WebgpuDebug.internal(this); WebgpuDebug.validate(this); @@ -668,6 +667,10 @@ class WebgpuGraphicsDevice extends GraphicsDevice { this.pipeline = null; // TODO: add performance queries to compute passes + } + + startComputePass() { + this.startCompute(); // start the pass this.passEncoder = this.commandEncoder.beginComputePass(); @@ -678,12 +681,15 @@ class WebgpuGraphicsDevice extends GraphicsDevice { } endComputePass() { - // end the compute pass this.passEncoder.end(); this.passEncoder = null; this.insideRenderPass = false; + this.endCompute(); + } + + endCompute() { // each render pass can use different number of bind groups this.bindGroupFormats.length = 0; From f6741f2f46d9826368016c3114d0e31149cb2f3f Mon Sep 17 00:00:00 2001 From: Erik Dubbelboer Date: Sun, 28 Jan 2024 16:53:37 +0100 Subject: [PATCH 2/4] Changed the API --- .../examples/misc/compute-shader/example.mjs | 7 ++- src/platform/graphics/compute.js | 7 +-- src/platform/graphics/webgpu/webgpu-buffer.js | 11 +++++ .../graphics/webgpu/webgpu-compute.js | 47 +++++-------------- .../graphics/webgpu/webgpu-graphics-device.js | 23 +++++---- 5 files changed, 46 insertions(+), 49 deletions(-) diff --git a/examples/src/examples/misc/compute-shader/example.mjs b/examples/src/examples/misc/compute-shader/example.mjs index bad057a3ef8..e0f709a9bd7 100644 --- a/examples/src/examples/misc/compute-shader/example.mjs +++ b/examples/src/examples/misc/compute-shader/example.mjs @@ -62,11 +62,16 @@ shader.impl.computeBindGroupFormat = new pc.BindGroupFormat(device,[], [], [ }); const compute = new pc.Compute(app.graphicsDevice, shader); +const buffer = compute.getBuffer(texture); +app.graphicsDevice.startComputePass(); compute.dispatch(texture.width, texture.height); +app.graphicsDevice.endComputePass(); -const data = await compute.read(texture.impl.gpuTexture); +const data = await buffer.getMappedRange(); console.log(data); +buffer.destroy(app.graphicsDevice); + export { app }; diff --git a/src/platform/graphics/compute.js b/src/platform/graphics/compute.js index 3b18e766b92..57c7ece2d9a 100644 --- a/src/platform/graphics/compute.js +++ b/src/platform/graphics/compute.js @@ -42,10 +42,11 @@ class Compute { /** * - * @param {GPUTexture} texture + * @param {import('./texture.js').Texture} texture + * @returns {import('./buffer.js').Buffer} */ - async read(texture) { - return this.impl?.read(texture); + getBuffer(texture) { + return this.impl?.getBuffer(texture); } } diff --git a/src/platform/graphics/webgpu/webgpu-buffer.js b/src/platform/graphics/webgpu/webgpu-buffer.js index 994888cc9bf..793b2d3fb3d 100644 --- a/src/platform/graphics/webgpu/webgpu-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-buffer.js @@ -15,6 +15,7 @@ class WebgpuBuffer { destroy(device) { if (this.buffer) { + this.buffer.unmap(); this.buffer.destroy(); this.buffer = null; } @@ -84,6 +85,16 @@ class WebgpuBuffer { // TODO: handle usage types: // - BUFFER_STATIC, BUFFER_DYNAMIC, BUFFER_STREAM, BUFFER_GPUDYNAMIC } + + /** + * @returns {Promise} + */ + async getMappedRange() { + await this.buffer.mapAsync(GPUMapMode.READ); + + const arrayBuffer = this.buffer.getMappedRange(); + return new Uint8Array(arrayBuffer); + } } export { WebgpuBuffer }; diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index 786b67d5d64..28f71d15ec0 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -1,5 +1,7 @@ import { Debug, DebugHelper } from "../../../core/debug.js"; import { BindGroup } from "../bind-group.js"; +import { Buffer } from "../buffer.js"; +import { WebgpuBuffer } from "./webgpu-buffer.js"; /** * A WebGPU implementation of the Compute. @@ -23,11 +25,7 @@ class WebgpuCompute { } dispatch(x, y, z) { - - // TODO: currently each dispatch is a separate compute pass, which is not optimal, and we should - // batch multiple dispatches into a single compute pass const device = this.compute.device; - device.startComputePass(); // bind group data const { bindGroup } = this; @@ -38,23 +36,14 @@ class WebgpuCompute { const passEncoder = device.passEncoder; passEncoder.setPipeline(this.pipeline); passEncoder.dispatchWorkgroups(x, y, z); - - device.endComputePass(); } /** * - * @param {GPUTexture} texture + * @param {import('../texture.js').Texture} texture + * @returns {import('../buffer.js').Buffer} */ - async read(texture) { - const device = this.compute.device; - device.startCompute(); - - // bind group data - const { bindGroup } = this; - bindGroup.update(); - device.setBindGroup(0, bindGroup); - + getBuffer(texture) { // Calculate bytes per pixel, assuming RGBA8 format (4 bytes per pixel) const bytesPerPixel = 4; @@ -64,13 +53,13 @@ class WebgpuCompute { // Calculate the size of the buffer to hold the texture data const bufferSize = bytesPerRow * texture.height; - const gpuBuffer = device.wgpu.createBuffer({ + const gpuBuffer = this.compute.device.wgpu.createBuffer({ size: bufferSize, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ }); const textureCopyView = { - texture, + texture: texture.impl.gpuTexture, origin: { x: 0, y: 0 }, }; const bufferCopyView = { @@ -82,25 +71,13 @@ class WebgpuCompute { height: texture.height, }; - // Encode command to copy from texture to buffer - device.commandEncoder.copyTextureToBuffer(textureCopyView, bufferCopyView, extent); - - device.endCompute(); - - await device.wgpu.queue.onSubmittedWorkDone(); - - // Ensure that the GPU operations are complete - await gpuBuffer.mapAsync(GPUMapMode.READ); - - // Read buffer contents - const arrayBuffer = gpuBuffer.getMappedRange(); - const data = new Uint8Array(arrayBuffer); // or another typed array based on the texture format + this.compute.device.copyTextureToBufferCommands.push([textureCopyView, bufferCopyView, extent]); - // Cleanup - //gpuBuffer.unmap(); - //gpuBuffer.destroy(); + const buffer = new Buffer(); + buffer.impl = new WebgpuBuffer(); + buffer.impl.buffer = gpuBuffer; - return data; + return buffer; } } diff --git a/src/platform/graphics/webgpu/webgpu-graphics-device.js b/src/platform/graphics/webgpu/webgpu-graphics-device.js index 6d3148d1940..141cdaf834d 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -89,6 +89,8 @@ class WebgpuGraphicsDevice extends GraphicsDevice { */ limits; + copyTextureToBufferCommands = []; + constructor(canvas, options = {}) { super(canvas, options); options = this.initOptions; @@ -654,7 +656,7 @@ class WebgpuGraphicsDevice extends GraphicsDevice { WebgpuDebug.end(this, { renderPass }); } - startCompute() { + startComputePass() { WebgpuDebug.internal(this); WebgpuDebug.validate(this); @@ -667,10 +669,6 @@ class WebgpuGraphicsDevice extends GraphicsDevice { this.pipeline = null; // TODO: add performance queries to compute passes - } - - startComputePass() { - this.startCompute(); // start the pass this.passEncoder = this.commandEncoder.beginComputePass(); @@ -684,12 +682,13 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // end the compute pass this.passEncoder.end(); this.passEncoder = null; - this.insideRenderPass = false; - this.endCompute(); - } + for (const [textureCopyView, bufferCopyView, extent] of this.copyTextureToBufferCommands) { + this.commandEncoder.copyTextureToBuffer(textureCopyView, bufferCopyView, extent); + } + + this.copyTextureToBufferCommands.length = 0; - endCompute() { // each render pass can use different number of bind groups this.bindGroupFormats.length = 0; @@ -698,9 +697,13 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // DebugHelper.setLabel(cb, `${renderPass.name}-CommandBuffer`); DebugHelper.setLabel(cb, 'ComputePass-CommandBuffer'); - this.addCommandBuffer(cb); + //this.addCommandBuffer(cb); + this.wgpu.queue.submit([cb]); + this.commandEncoder = null; + this.insideRenderPass = false; + WebgpuDebug.end(this); WebgpuDebug.end(this); } From 3c6cb4ed617537f545872df174bcf12735aad145 Mon Sep 17 00:00:00 2001 From: Erik Dubbelboer Date: Sun, 28 Jan 2024 17:56:25 +0100 Subject: [PATCH 3/4] Change API, add inputTexture --- .../examples/misc/compute-shader/config.mjs | 14 ++++++--- .../examples/misc/compute-shader/example.mjs | 29 +++++++++++++++---- src/platform/graphics/buffer.js | 28 ++++++++++++++++++ src/platform/graphics/compute.js | 6 ++-- src/platform/graphics/shader.js | 6 ++++ src/platform/graphics/webgpu/webgpu-buffer.js | 5 +++- .../graphics/webgpu/webgpu-compute.js | 28 +++++++++++++----- .../graphics/webgpu/webgpu-graphics-device.js | 9 ++++-- 8 files changed, 102 insertions(+), 23 deletions(-) create mode 100644 src/platform/graphics/buffer.js diff --git a/examples/src/examples/misc/compute-shader/config.mjs b/examples/src/examples/misc/compute-shader/config.mjs index 07ca00f7f33..6303b64ee44 100644 --- a/examples/src/examples/misc/compute-shader/config.mjs +++ b/examples/src/examples/misc/compute-shader/config.mjs @@ -6,12 +6,18 @@ export default { HIDDEN: true, FILES: { 'shader.wgsl': ` - @group(0) @binding(0) var outputTexture: texture_storage_2d; + @group(0) @binding(0) var inputTexture: texture_2d; + // @binding(1) is a sampler of the inputTexture, but we don't need it in the shader. + @group(0) @binding(2) var outputTexture: texture_storage_2d; - @compute @workgroup_size(1, 1, 1) + @compute @workgroup_size(1) fn main(@builtin(global_invocation_id) global_id : vec3u) { - let clearColor: vec4 = vec4(0.5); - textureStore(outputTexture, vec2(global_id.xy), clearColor); + let position : vec2i = vec2i(global_id.xy); + var color : vec4f = textureLoad(inputTexture, position, 0); + + color = vec4f(1.0) - color; + + textureStore(outputTexture, position, color); } ` } diff --git a/examples/src/examples/misc/compute-shader/example.mjs b/examples/src/examples/misc/compute-shader/example.mjs index e0f709a9bd7..3510ab4a8f2 100644 --- a/examples/src/examples/misc/compute-shader/example.mjs +++ b/examples/src/examples/misc/compute-shader/example.mjs @@ -7,6 +7,10 @@ if (!(canvas instanceof HTMLCanvasElement)) { throw new Error('No canvas found'); } +const assets = { + rocks: new pc.Asset('rocks', 'texture', { url: rootPath + '/static/assets/textures/seaside-rocks01-color.jpg' }) +}; + const gfxOptions = { deviceTypes: [deviceType], @@ -17,11 +21,16 @@ const gfxOptions = { }; const device = await pc.createGraphicsDevice(canvas, gfxOptions); + +if (!device.isWebGPU) { + throw new Error('WebGPU is required for this example.'); +} + const createOptions = new pc.AppOptions(); createOptions.graphicsDevice = device; -createOptions.componentSystems = [pc.RenderComponentSystem, pc.CameraComponentSystem, pc.LightComponentSystem]; -createOptions.resourceHandlers = [pc.TextureHandler, pc.ContainerHandler]; +createOptions.componentSystems = [pc.RenderComponentSystem]; +createOptions.resourceHandlers = [pc.TextureHandler]; const app = new pc.AppBase(canvas); app.init(createOptions); @@ -38,16 +47,21 @@ app.on('destroy', () => { window.removeEventListener('resize', resize); }); +const inputTexture = assets.rocks.resource; +const width = inputTexture.width; +const height = inputTexture.height; + const texture = new pc.Texture(app.graphicsDevice, { name: 'outputTexture', - width: 2, - height: 2, + width, + height, format: pc.PIXELFORMAT_RGBA8, mipmaps: false, storage: true }); app.graphicsDevice.scope.resolve("outputTexture").setValue(texture); +app.graphicsDevice.scope.resolve("inputTexture").setValue(inputTexture); const shaderDefinition = { cshader: files['shader.wgsl'], @@ -55,7 +69,9 @@ const shaderDefinition = { }; const shader = new pc.Shader(app.graphicsDevice, shaderDefinition); -shader.impl.computeBindGroupFormat = new pc.BindGroupFormat(device,[], [], [ +shader.computeBindGroupFormat = new pc.BindGroupFormat(device, [], [ + new pc.BindTextureFormat('inputTexture', pc.SHADERSTAGE_COMPUTE, pc.TEXTUREDIMENSION_2D, pc.SAMPLETYPE_FLOAT), +], [ new pc.BindStorageTextureFormat('outputTexture', pc.PIXELFORMAT_RGBA8, pc.TEXTUREDIMENSION_2D), ], { compute: true @@ -65,7 +81,8 @@ const compute = new pc.Compute(app.graphicsDevice, shader); const buffer = compute.getBuffer(texture); app.graphicsDevice.startComputePass(); -compute.dispatch(texture.width, texture.height); +compute.dispatch(width, height); +// TODO: potentially dispatch more compute work in the same pass. app.graphicsDevice.endComputePass(); const data = await buffer.getMappedRange(); diff --git a/src/platform/graphics/buffer.js b/src/platform/graphics/buffer.js new file mode 100644 index 00000000000..a59e381931e --- /dev/null +++ b/src/platform/graphics/buffer.js @@ -0,0 +1,28 @@ +/** + * ... + * + * @ignore + */ +class Buffer { + /** + * @type {import('./webgpu/webgpu-buffer.js').WebgpuBuffer|null} + * @private + */ + impl = null; + + destroy(device) { + this.impl?.destroy(device); + } + + /** + * Returns a mapped range of the underlying buffer. + * On WebGPU this will wait for the buffer to be copied to the CPU. + * + * @returns {Promise} The mapped range. + */ + async getMappedRange() { + return await this.impl?.getMappedRange?.(); + } +} + +export { Buffer }; diff --git a/src/platform/graphics/compute.js b/src/platform/graphics/compute.js index 57c7ece2d9a..eb198b5182e 100644 --- a/src/platform/graphics/compute.js +++ b/src/platform/graphics/compute.js @@ -41,9 +41,11 @@ class Compute { } /** + * Get a buffer that contains the data of the specified texture. + * This needs to be called before dispatch! But can be called before device.startComputePass(). * - * @param {import('./texture.js').Texture} texture - * @returns {import('./buffer.js').Buffer} + * @param {import('./texture.js').Texture} texture - The texture to get the buffer for. + * @returns {import('./buffer.js').Buffer} The buffer. */ getBuffer(texture) { return this.impl?.getBuffer(texture); diff --git a/src/platform/graphics/shader.js b/src/platform/graphics/shader.js index 19541310279..f0c4d19e6c7 100644 --- a/src/platform/graphics/shader.js +++ b/src/platform/graphics/shader.js @@ -33,6 +33,12 @@ class Shader { */ meshBindGroupFormat; + /** + * Format of the bind group for the compute bind group. + * @type {import('./bind-group-format.js').BindGroupFormat} + */ + computeBindGroupFormat; + /** * Creates a new Shader instance. * diff --git a/src/platform/graphics/webgpu/webgpu-buffer.js b/src/platform/graphics/webgpu/webgpu-buffer.js index 793b2d3fb3d..095c8e7f193 100644 --- a/src/platform/graphics/webgpu/webgpu-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-buffer.js @@ -87,7 +87,10 @@ class WebgpuBuffer { } /** - * @returns {Promise} + * Returns a mapped range of the underlying buffer. + * On WebGPU this will wait for the buffer to be copied to the CPU. + * + * @returns {Promise} The mapped range. */ async getMappedRange() { await this.buffer.mapAsync(GPUMapMode.READ); diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index 28f71d15ec0..73597b83bed 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -9,13 +9,15 @@ import { WebgpuBuffer } from "./webgpu-buffer.js"; * @ignore */ class WebgpuCompute { + copyTextureToBufferCommands = []; + constructor(compute) { this.compute = compute; const { device, shader } = compute; // create bind group - const { computeBindGroupFormat } = shader.impl; + const { computeBindGroupFormat } = shader; Debug.assert(computeBindGroupFormat, 'Compute shader does not have computeBindGroupFormat specified', shader); this.bindGroup = new BindGroup(device, computeBindGroupFormat); DebugHelper.setName(this.bindGroup, `Compute-BindGroup_${this.bindGroup.id}`); @@ -24,6 +26,13 @@ class WebgpuCompute { this.pipeline = device.computePipeline.get(shader, computeBindGroupFormat); } + /** + * Dispatch the compute work. + * + * @param {number} x - X dimension of the grid of work-groups to dispatch. + * @param {number} [y] - Y dimension of the grid of work-groups to dispatch. + * @param {number} [z] - Z dimension of the grid of work-groups to dispatch. + */ dispatch(x, y, z) { const device = this.compute.device; @@ -36,12 +45,17 @@ class WebgpuCompute { const passEncoder = device.passEncoder; passEncoder.setPipeline(this.pipeline); passEncoder.dispatchWorkgroups(x, y, z); + + this.compute.device.copyTextureToBufferCommands.push(...this.copyTextureToBufferCommands); + this.copyTextureToBufferCommands.length = 0; } /** + * Get a buffer that contains the data of the specified texture. + * This needs to be called before dispatch! But can be called before device.startComputePass(). * - * @param {import('../texture.js').Texture} texture - * @returns {import('../buffer.js').Buffer} + * @param {import('../texture.js').Texture} texture - The texture to get the buffer for. + * @returns {import('../buffer.js').Buffer} The buffer. */ getBuffer(texture) { // Calculate bytes per pixel, assuming RGBA8 format (4 bytes per pixel) @@ -60,18 +74,18 @@ class WebgpuCompute { const textureCopyView = { texture: texture.impl.gpuTexture, - origin: { x: 0, y: 0 }, + origin: { x: 0, y: 0 } }; const bufferCopyView = { buffer: gpuBuffer, - bytesPerRow: bytesPerRow, + bytesPerRow: bytesPerRow }; const extent = { width: texture.width, - height: texture.height, + height: texture.height }; - this.compute.device.copyTextureToBufferCommands.push([textureCopyView, bufferCopyView, extent]); + this.copyTextureToBufferCommands.push([textureCopyView, bufferCopyView, extent]); const buffer = new Buffer(); buffer.impl = new WebgpuBuffer(); diff --git a/src/platform/graphics/webgpu/webgpu-graphics-device.js b/src/platform/graphics/webgpu/webgpu-graphics-device.js index 141cdaf834d..31d4e8527f8 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -668,6 +668,8 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // clear cached encoder state this.pipeline = null; + this.copyTextureToBufferCommands.length = 0; + // TODO: add performance queries to compute passes // start the pass @@ -683,12 +685,11 @@ class WebgpuGraphicsDevice extends GraphicsDevice { this.passEncoder.end(); this.passEncoder = null; + // These commands can only be called outside of a compute pass. for (const [textureCopyView, bufferCopyView, extent] of this.copyTextureToBufferCommands) { this.commandEncoder.copyTextureToBuffer(textureCopyView, bufferCopyView, extent); } - this.copyTextureToBufferCommands.length = 0; - // each render pass can use different number of bind groups this.bindGroupFormats.length = 0; @@ -697,7 +698,9 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // DebugHelper.setLabel(cb, `${renderPass.name}-CommandBuffer`); DebugHelper.setLabel(cb, 'ComputePass-CommandBuffer'); - //this.addCommandBuffer(cb); + // Don't this.addCommandBuffer(cb) as that means we'll have to + // wait for the render pass to finish before this is submitted + // which isn't required. this.wgpu.queue.submit([cb]); this.commandEncoder = null; From b0926e2a0c23ec811fd1ec9dc723e0d2635b05b9 Mon Sep 17 00:00:00 2001 From: Erik Dubbelboer Date: Sat, 10 Feb 2024 15:37:40 +0100 Subject: [PATCH 4/4] Add GPUBuffer input and output --- .../examples/misc/compute-shader/config.mjs | 17 ++-- .../examples/misc/compute-shader/example.mjs | 93 +++++++++++-------- src/index.js | 1 + src/platform/graphics/bind-group-format.js | 25 +++-- src/platform/graphics/bind-group.js | 19 +++- src/platform/graphics/buffer.js | 50 +++++++++- src/platform/graphics/compute.js | 17 +++- src/platform/graphics/constants.js | 6 ++ .../graphics/null/null-graphics-device.js | 4 + .../graphics/webgl/webgl-graphics-device.js | 4 + .../webgpu/webgpu-bind-group-format.js | 15 +++ .../graphics/webgpu/webgpu-bind-group.js | 18 ++++ src/platform/graphics/webgpu/webgpu-buffer.js | 42 +++++++-- .../graphics/webgpu/webgpu-compute.js | 49 +++++++--- .../graphics/webgpu/webgpu-graphics-device.js | 19 ++++ 15 files changed, 301 insertions(+), 78 deletions(-) diff --git a/examples/src/examples/misc/compute-shader/config.mjs b/examples/src/examples/misc/compute-shader/config.mjs index 6303b64ee44..6a2d8017181 100644 --- a/examples/src/examples/misc/compute-shader/config.mjs +++ b/examples/src/examples/misc/compute-shader/config.mjs @@ -4,20 +4,25 @@ export default { WEBGPU_REQUIRED: true, HIDDEN: true, + NO_MINISTATS: true, FILES: { 'shader.wgsl': ` @group(0) @binding(0) var inputTexture: texture_2d; - // @binding(1) is a sampler of the inputTexture, but we don't need it in the shader. - @group(0) @binding(2) var outputTexture: texture_storage_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 inout: array>; @compute @workgroup_size(1) fn main(@builtin(global_invocation_id) global_id : vec3u) { - let position : vec2i = vec2i(global_id.xy); - var color : vec4f = textureLoad(inputTexture, position, 0); + let position = vec2i(global_id.xy); + var color = textureLoad(inputTexture, position, 0).rgb; + var input = vec3u(atomicLoad(&inout[0]), atomicLoad(&inout[1]), atomicLoad(&inout[2])); + var compare = vec3f(input) / 255.0; - color = vec4f(1.0) - color; + atomicAdd(&inout[3], 1u); - textureStore(outputTexture, position, color); + if (color.r >= compare.r && color.g >= compare.g && color.b >= compare.b) { + atomicAdd(&inout[4], 1u); + } } ` } diff --git a/examples/src/examples/misc/compute-shader/example.mjs b/examples/src/examples/misc/compute-shader/example.mjs index 3510ab4a8f2..b53bdb1c6cd 100644 --- a/examples/src/examples/misc/compute-shader/example.mjs +++ b/examples/src/examples/misc/compute-shader/example.mjs @@ -34,7 +34,6 @@ createOptions.resourceHandlers = [pc.TextureHandler]; 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); @@ -47,48 +46,68 @@ app.on('destroy', () => { window.removeEventListener('resize', resize); }); -const inputTexture = assets.rocks.resource; -const width = inputTexture.width; -const height = inputTexture.height; - -const texture = new pc.Texture(app.graphicsDevice, { - name: 'outputTexture', - width, - height, - format: pc.PIXELFORMAT_RGBA8, - mipmaps: false, - storage: true -}); +const assetListLoader = new pc.AssetListLoader(Object.values(assets), app.assets); +assetListLoader.load(async () => { + app.start(); -app.graphicsDevice.scope.resolve("outputTexture").setValue(texture); -app.graphicsDevice.scope.resolve("inputTexture").setValue(inputTexture); + // This example will use a compute shader to count the number of pixels brighter than + // a certain specified color. The result will be written to a buffer and read back to the CPU. + const inputTexture = assets.rocks.resource; + const compareColor = [0.5, 0.5, 0.5]; -const shaderDefinition = { - cshader: files['shader.wgsl'], - shaderLanguage: pc.SHADERLANGUAGE_WGSL, -}; -const shader = new pc.Shader(app.graphicsDevice, shaderDefinition); - -shader.computeBindGroupFormat = new pc.BindGroupFormat(device, [], [ - new pc.BindTextureFormat('inputTexture', pc.SHADERSTAGE_COMPUTE, pc.TEXTUREDIMENSION_2D, pc.SAMPLETYPE_FLOAT), -], [ - new pc.BindStorageTextureFormat('outputTexture', pc.PIXELFORMAT_RGBA8, pc.TEXTUREDIMENSION_2D), -], { - compute: true -}); + const width = inputTexture.width; + const height = inputTexture.height; + // The buffer we pass to the GPU need to be initialized with the compare color and counters of 0. + // Since the buffer is an uint32 array, we need to convert the color to 0-255 range. + const init = [compareColor[0] * 255, compareColor[1] * 255, compareColor[2] * 255, 0, 0]; + + const buffer = new pc.Buffer(app.graphicsDevice, { + size: 5 * 4, // 5 uint32s (3 color components, 1 counter, 1 counter for brighter pixels). + usage: pc.BUFFER_USAGE_STORAGE | pc.BUFFER_USAGE_COPY_SRC, + mappedAtCreation: true, + }); + + new Uint32Array(buffer.getMappedRange()).set(init); + buffer.unmap(); + + app.graphicsDevice.scope.resolve("inout").setValue(buffer); + app.graphicsDevice.scope.resolve("inputTexture").setValue(inputTexture); -const compute = new pc.Compute(app.graphicsDevice, shader); -const buffer = compute.getBuffer(texture); + const shaderDefinition = { + cshader: files['shader.wgsl'], + shaderLanguage: pc.SHADERLANGUAGE_WGSL, + }; + const shader = new pc.Shader(app.graphicsDevice, shaderDefinition); -app.graphicsDevice.startComputePass(); -compute.dispatch(width, height); -// TODO: potentially dispatch more compute work in the same pass. -app.graphicsDevice.endComputePass(); + shader.computeBindGroupFormat = new pc.BindGroupFormat(device, [], [ + new pc.BindTextureFormat('inputTexture', pc.SHADERSTAGE_COMPUTE, pc.TEXTUREDIMENSION_2D, pc.SAMPLETYPE_FLOAT), + ], [ + // No storage textures used. + ], [ + new pc.BindBufferFormat('inout', pc.SHADERSTAGE_COMPUTE), + ]); -const data = await buffer.getMappedRange(); + const compute = new pc.Compute(app.graphicsDevice, shader); -console.log(data); + // Get a buffer for the result of our compute work. + // This needs to be requested before dispatching the compute work. + const resultBuffer = compute.getBuffer(buffer); -buffer.destroy(app.graphicsDevice); + app.graphicsDevice.startComputePass(); + compute.dispatch(width, height); + app.graphicsDevice.endComputePass(); + + // Map the result buffer to the CPU to read the result. + await resultBuffer.mapAsync(); + + const data = new Uint32Array(resultBuffer.getMappedRange()); + + console.log('number of pixels:', data[3]); + console.log('number of pixels brighter than', compareColor, ':', data[4]); + + // Clean up the buffers we used. + buffer.destroy(app.graphicsDevice); + resultBuffer.destroy(app.graphicsDevice); +}); export { app }; diff --git a/src/index.js b/src/index.js index d0db3c2a810..0f5c417645d 100644 --- a/src/index.js +++ b/src/index.js @@ -91,6 +91,7 @@ export { UniformBufferFormat, UniformFormat } from './platform/graphics/uniform- export { VertexBuffer } from './platform/graphics/vertex-buffer.js'; export { VertexFormat } from './platform/graphics/vertex-format.js'; export { VertexIterator } from './platform/graphics/vertex-iterator.js'; +export { Buffer } from './platform/graphics/buffer.js'; // PLATFORM / GRAPHICS / webgl export { WebglGraphicsDevice } from './platform/graphics/webgl/webgl-graphics-device.js'; diff --git a/src/platform/graphics/bind-group-format.js b/src/platform/graphics/bind-group-format.js index b57b27cffae..04ace04b2e6 100644 --- a/src/platform/graphics/bind-group-format.js +++ b/src/platform/graphics/bind-group-format.js @@ -73,8 +73,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,17 +82,13 @@ 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 {object} [options] - Object for passing optional arguments. - * @param {boolean} [options.compute] - If true, this bind group format is used by the compute - * shader. + * @param {BindBufferFormat[]} [storageBufferFormats] - An array of bind storage buffer + * formats (storage buffers), used by the compute shader. Defaults to an empty array. */ - constructor(graphicsDevice, bufferFormats = [], textureFormats = [], storageTextureFormats = [], options = {}) { + constructor(graphicsDevice, bufferFormats = [], textureFormats = [], storageTextureFormats = [], storageBufferFormats = []) { this.id = id++; DebugHelper.setName(this, `BindGroupFormat_${this.id}`); - this.compute = options.compute ?? false; - Debug.assert(this.compute || storageTextureFormats.length === 0, "Storage textures can be specified only for compute"); - /** @type {import('./graphics-device.js').GraphicsDevice} */ this.device = graphicsDevice; const scope = graphicsDevice.scope; @@ -133,6 +127,19 @@ class BindGroupFormat { tf.scopeId = scope.resolve(tf.name); }); + /** @type {BindBufferFormat[]} */ + 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); diff --git a/src/platform/graphics/bind-group.js b/src/platform/graphics/bind-group.js index a029dc9244f..ed95468d7b9 100644 --- a/src/platform/graphics/bind-group.js +++ b/src/platform/graphics/bind-group.js @@ -50,6 +50,7 @@ class BindGroup { this.textures = []; this.storageTextures = []; + this.storageBuffers = []; this.uniformBuffers = []; /** @type {import('./uniform-buffer.js').UniformBuffer} */ @@ -123,13 +124,22 @@ class BindGroup { } } + setStorageBuffer(name, buffer) { + 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] !== buffer) { + this.storageBuffers[index] = buffer; + this.dirty = true; + } + } + /** * Applies any changes made to the bind group's properties. */ 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 +155,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/buffer.js b/src/platform/graphics/buffer.js index a59e381931e..ffb36f3dc93 100644 --- a/src/platform/graphics/buffer.js +++ b/src/platform/graphics/buffer.js @@ -1,3 +1,5 @@ +import { Debug } from '../../core/debug.js'; + /** * ... * @@ -10,18 +12,58 @@ class Buffer { */ impl = null; + /** + * @param {import('./graphics-device.js').GraphicsDevice} device - The graphics device + * used to manage this buffer. + * @param {object} options - The options for the buffer. + * @param {number} options.size - The size of the buffer in bytes. + * @param {number} [options.usage] - The usage of the buffer. One of: + * `pc.BUFFER_USAGE_STORAGE`, `pc.BUFFER_USAGE_COPY_SRC`, `pc.BUFFER_USAGE_COPY_DST`, + * `pc.BUFFER_USAGE_MAP_READ`. + * @param {boolean} [options.mappedAtCreation] - Whether the buffer is mapped at + * creation. Default is `false`. + */ + constructor(device, options) { + Debug.assert(device, "Texture constructor requires a graphicsDevice to be valid"); + Debug.assert(options.size, "Texture constructor requires a size to be valid"); + Debug.assert(options.usage, "Texture constructor requires a usage to be valid"); + + this.impl = device.createBufferImpl(options); + } + destroy(device) { this.impl?.destroy(device); } + get size() { + return this.impl?.size || 0; + } + + /** + * Map the buffer to CPU memory for reading or writing. After the promise is resolved, the buffer + * is mapped and can be accessed through the `getMappedRange` method. + * + * @param {boolean} write - Map for writing, otherwise map for reading, default is false. + * @returns {Promise} The mapped range. + */ + async mapAsync(write) { + await this.impl?.mapAsync(write); + } + + /** + * Unmap the buffer from CPU memory so it can be used by the GPU. + */ + unmap() { + this.impl?.unmap(); + } + /** * Returns a mapped range of the underlying buffer. - * On WebGPU this will wait for the buffer to be copied to the CPU. * - * @returns {Promise} The mapped range. + * @returns {ArrayBuffer|undefined} The mapped range. */ - async getMappedRange() { - return await this.impl?.getMappedRange?.(); + getMappedRange() { + return this.impl?.getMappedRange(); } } diff --git a/src/platform/graphics/compute.js b/src/platform/graphics/compute.js index eb198b5182e..a5f168e6047 100644 --- a/src/platform/graphics/compute.js +++ b/src/platform/graphics/compute.js @@ -41,14 +41,25 @@ class Compute { } /** - * Get a buffer that contains the data of the specified texture. + * Get a buffer that contains the data of the specified storage texture. * This needs to be called before dispatch! But can be called before device.startComputePass(). * * @param {import('./texture.js').Texture} texture - The texture to get the buffer for. * @returns {import('./buffer.js').Buffer} The buffer. */ - getBuffer(texture) { - return this.impl?.getBuffer(texture); + getTextureBuffer(texture) { + return this.impl?.getTextureBuffer(texture); + } + + /** + * Get a buffer that contains the data of the specified buffer. + * This needs to be called before dispatch! But can be called before device.startComputePass(). + * + * @param {import('./buffer.js').Buffer} buffer - The buffer to get the data from. + * @returns {import('./buffer.js').Buffer} The buffer. + */ + getBuffer(buffer) { + return this.impl?.getBuffer(buffer); } } diff --git a/src/platform/graphics/constants.js b/src/platform/graphics/constants.js index 1612c4c9cc2..6f700d33ab5 100644 --- a/src/platform/graphics/constants.js +++ b/src/platform/graphics/constants.js @@ -1810,3 +1810,9 @@ export const CHUNKAPI_1_58 = '1.58'; export const CHUNKAPI_1_60 = '1.60'; export const CHUNKAPI_1_62 = '1.62'; export const CHUNKAPI_1_65 = '1.65'; + +// These should be the same as their GPUBufferUsage counterparts. +export const BUFFER_USAGE_MAP_READ = 1; +export const BUFFER_USAGE_COPY_SRC = 4; +export const BUFFER_USAGE_COPY_DST = 8; +export const BUFFER_USAGE_STORAGE = 128; diff --git a/src/platform/graphics/null/null-graphics-device.js b/src/platform/graphics/null/null-graphics-device.js index 96616cd76a3..bb3b583a848 100644 --- a/src/platform/graphics/null/null-graphics-device.js +++ b/src/platform/graphics/null/null-graphics-device.js @@ -91,6 +91,10 @@ class NullGraphicsDevice extends GraphicsDevice { return new NullIndexBuffer(indexBuffer); } + createBufferImpl(options) { + return null; + } + createShaderImpl(shader) { return new NullShader(shader); } diff --git a/src/platform/graphics/webgl/webgl-graphics-device.js b/src/platform/graphics/webgl/webgl-graphics-device.js index d3a69e5f1c3..2234e6e52e4 100644 --- a/src/platform/graphics/webgl/webgl-graphics-device.js +++ b/src/platform/graphics/webgl/webgl-graphics-device.js @@ -955,6 +955,10 @@ class WebglGraphicsDevice extends GraphicsDevice { return new WebglIndexBuffer(indexBuffer); } + createBufferImpl(options) { + return null; + } + createShaderImpl(shader) { return new WebglShader(shader); } diff --git a/src/platform/graphics/webgpu/webgpu-bind-group-format.js b/src/platform/graphics/webgpu/webgpu-bind-group-format.js index 17446ceb53e..0343f764dcf 100644 --- a/src/platform/graphics/webgpu/webgpu-bind-group-format.js +++ b/src/platform/graphics/webgpu/webgpu-bind-group-format.js @@ -93,6 +93,7 @@ class WebgpuBindGroupFormat { // - sampler: GPUSamplerBindingLayout, resource type is GPUSampler // - texture: GPUTextureBindingLayout, resource type is GPUTextureView // - storageTexture: GPUStorageTextureBindingLayout, resource type is GPUTextureView + // - storageBuffer: GPUStorageBufferBindingLayout, resource type is GPUBufferBinding // - externalTexture: GPUExternalTextureBindingLayout, resource type is GPUExternalTexture const entries = []; @@ -199,6 +200,20 @@ class WebgpuBindGroupFormat { }); }); + // storage buffers + bindGroupFormat.storageBufferFormats.forEach((bufferFormat) => { + + key += `#${index}SB:`; + + entries.push({ + binding: index++, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: '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 095c8e7f193..6632d202ac3 100644 --- a/src/platform/graphics/webgpu/webgpu-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-buffer.js @@ -13,6 +13,10 @@ class WebgpuBuffer { */ buffer = null; + init(device, options) { + this.buffer = device.wgpu.createBuffer(options); + } + destroy(device) { if (this.buffer) { this.buffer.unmap(); @@ -21,6 +25,10 @@ class WebgpuBuffer { } } + get size() { + return this.buffer ? this.buffer.size : 0; + } + get initialized() { return !!this.buffer; } @@ -86,17 +94,39 @@ class WebgpuBuffer { // - BUFFER_STATIC, BUFFER_DYNAMIC, BUFFER_STREAM, BUFFER_GPUDYNAMIC } + /** + * Map the buffer to CPU memory for reading or writing. After the promise is resolved, the buffer + * is mapped and can be accessed through the `getMappedRange` method. + * + * @param {boolean} write - Map for writing, otherwise map for reading, default is false. + * @returns {Promise} The mapped range. + */ + async mapAsync(write) { + if (this.buffer) { + await this.buffer.mapAsync(write ? GPUMapMode.WRITE : GPUMapMode.READ); + } + } + + /** + * Unmap the buffer from CPU memory so it can be used by the GPU. + */ + unmap() { + if (this.buffer) { + this.buffer.unmap(); + } + } + /** * Returns a mapped range of the underlying buffer. - * On WebGPU this will wait for the buffer to be copied to the CPU. * - * @returns {Promise} The mapped range. + * @returns {ArrayBuffer|undefined} The mapped range. */ - async getMappedRange() { - await this.buffer.mapAsync(GPUMapMode.READ); + getMappedRange() { + if (!this.buffer) { + return; + } - const arrayBuffer = this.buffer.getMappedRange(); - return new Uint8Array(arrayBuffer); + return this.buffer.getMappedRange(); } } diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index 73597b83bed..167a2646e72 100644 --- a/src/platform/graphics/webgpu/webgpu-compute.js +++ b/src/platform/graphics/webgpu/webgpu-compute.js @@ -1,7 +1,7 @@ import { Debug, DebugHelper } from "../../../core/debug.js"; import { BindGroup } from "../bind-group.js"; import { Buffer } from "../buffer.js"; -import { WebgpuBuffer } from "./webgpu-buffer.js"; +import { pixelFormatInfo, BUFFER_USAGE_COPY_DST, BUFFER_USAGE_MAP_READ } from "../constants.js"; /** * A WebGPU implementation of the Compute. @@ -9,8 +9,16 @@ import { WebgpuBuffer } from "./webgpu-buffer.js"; * @ignore */ class WebgpuCompute { + /** + * @ignore + */ copyTextureToBufferCommands = []; + /** + * @ignore + */ + copyBufferToBufferCommands = []; + constructor(compute) { this.compute = compute; @@ -48,18 +56,21 @@ class WebgpuCompute { this.compute.device.copyTextureToBufferCommands.push(...this.copyTextureToBufferCommands); this.copyTextureToBufferCommands.length = 0; + + this.compute.device.copyBufferToBufferCommands.push(...this.copyBufferToBufferCommands); + this.copyBufferToBufferCommands.length = 0; } /** - * Get a buffer that contains the data of the specified texture. + * Get a buffer that contains the data of the specified storage texture. * This needs to be called before dispatch! But can be called before device.startComputePass(). * * @param {import('../texture.js').Texture} texture - The texture to get the buffer for. * @returns {import('../buffer.js').Buffer} The buffer. */ - getBuffer(texture) { - // Calculate bytes per pixel, assuming RGBA8 format (4 bytes per pixel) - const bytesPerPixel = 4; + getTextureBuffer(texture) { + const formatInfo = pixelFormatInfo.get(texture.format); + const bytesPerPixel = formatInfo.size; // Calculate bytes per row, ensuring it's a multiple of 256 const bytesPerRow = Math.ceil((texture.width * bytesPerPixel) / 256) * 256; @@ -67,9 +78,9 @@ class WebgpuCompute { // Calculate the size of the buffer to hold the texture data const bufferSize = bytesPerRow * texture.height; - const gpuBuffer = this.compute.device.wgpu.createBuffer({ + const buffer = new Buffer(this.compute.device, { size: bufferSize, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ + usage: BUFFER_USAGE_COPY_DST | BUFFER_USAGE_MAP_READ }); const textureCopyView = { @@ -77,7 +88,7 @@ class WebgpuCompute { origin: { x: 0, y: 0 } }; const bufferCopyView = { - buffer: gpuBuffer, + buffer: buffer.impl.buffer, bytesPerRow: bytesPerRow }; const extent = { @@ -87,12 +98,26 @@ class WebgpuCompute { this.copyTextureToBufferCommands.push([textureCopyView, bufferCopyView, extent]); - const buffer = new Buffer(); - buffer.impl = new WebgpuBuffer(); - buffer.impl.buffer = gpuBuffer; - return buffer; } + + /** + * Get a buffer that contains the data of the specified buffer. + * This needs to be called before dispatch! But can be called before device.startComputePass(). + * + * @param {import('../buffer.js').Buffer} buffer - The buffer to get the data from. + * @returns {import('../buffer.js').Buffer} The buffer. + */ + getBuffer(buffer) { + const gpuBuffer = new Buffer(this.compute.device, { + size: buffer.size, + usage: BUFFER_USAGE_COPY_DST | BUFFER_USAGE_MAP_READ + }); + + this.copyBufferToBufferCommands.push([buffer.impl.buffer, gpuBuffer.impl.buffer, buffer.size]); + + return gpuBuffer; + } } export { WebgpuCompute }; diff --git a/src/platform/graphics/webgpu/webgpu-graphics-device.js b/src/platform/graphics/webgpu/webgpu-graphics-device.js index 31d4e8527f8..ec44bd092ad 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -26,6 +26,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 { /** @@ -89,8 +90,16 @@ class WebgpuGraphicsDevice extends GraphicsDevice { */ limits; + /** + * @ignore + */ copyTextureToBufferCommands = []; + /** + * @ignore + */ + copyBufferToBufferCommands = []; + constructor(canvas, options = {}) { super(canvas, options); options = this.initOptions; @@ -375,6 +384,12 @@ class WebgpuGraphicsDevice extends GraphicsDevice { return new WebgpuIndexBuffer(indexBuffer); } + createBufferImpl(options) { + const buffer = new WebgpuBuffer(); + buffer.init(this, options); + return buffer; + } + createShaderImpl(shader) { return new WebgpuShader(shader); } @@ -669,6 +684,7 @@ class WebgpuGraphicsDevice extends GraphicsDevice { this.pipeline = null; this.copyTextureToBufferCommands.length = 0; + this.copyBufferToBufferCommands.length = 0; // TODO: add performance queries to compute passes @@ -689,6 +705,9 @@ class WebgpuGraphicsDevice extends GraphicsDevice { for (const [textureCopyView, bufferCopyView, extent] of this.copyTextureToBufferCommands) { this.commandEncoder.copyTextureToBuffer(textureCopyView, bufferCopyView, extent); } + for (const [srcBuffer, dstBuffer, size] of this.copyBufferToBufferCommands) { + this.commandEncoder.copyBufferToBuffer(srcBuffer, 0, dstBuffer, 0, size); + } // each render pass can use different number of bind groups this.bindGroupFormats.length = 0;