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..6a2d8017181 --- /dev/null +++ b/examples/src/examples/misc/compute-shader/config.mjs @@ -0,0 +1,29 @@ +/** + * @type {import('../../../../types.mjs').ExampleConfig} + */ +export default { + WEBGPU_REQUIRED: true, + HIDDEN: true, + NO_MINISTATS: true, + FILES: { + 'shader.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 inout: array>; + + @compute @workgroup_size(1) + fn main(@builtin(global_invocation_id) global_id : vec3u) { + 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; + + atomicAdd(&inout[3], 1u); + + 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 new file mode 100644 index 00000000000..b53bdb1c6cd --- /dev/null +++ b/examples/src/examples/misc/compute-shader/example.mjs @@ -0,0 +1,113 @@ +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 assets = { + rocks: new pc.Asset('rocks', 'texture', { url: rootPath + '/static/assets/textures/seaside-rocks01-color.jpg' }) +}; + +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); + +if (!device.isWebGPU) { + throw new Error('WebGPU is required for this example.'); +} + +const createOptions = new pc.AppOptions(); +createOptions.graphicsDevice = device; + +createOptions.componentSystems = [pc.RenderComponentSystem]; +createOptions.resourceHandlers = [pc.TextureHandler]; + +const app = new pc.AppBase(canvas); +app.init(createOptions); + +// 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(async () => { + app.start(); + + // 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 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 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), + ], [ + // No storage textures used. + ], [ + new pc.BindBufferFormat('inout', pc.SHADERSTAGE_COMPUTE), + ]); + + const compute = new pc.Compute(app.graphicsDevice, shader); + + // 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); + + 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 new file mode 100644 index 00000000000..ffb36f3dc93 --- /dev/null +++ b/src/platform/graphics/buffer.js @@ -0,0 +1,70 @@ +import { Debug } from '../../core/debug.js'; + +/** + * ... + * + * @ignore + */ +class Buffer { + /** + * @type {import('./webgpu/webgpu-buffer.js').WebgpuBuffer|null} + * @private + */ + 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. + * + * @returns {ArrayBuffer|undefined} The mapped range. + */ + getMappedRange() { + return this.impl?.getMappedRange(); + } +} + +export { Buffer }; diff --git a/src/platform/graphics/compute.js b/src/platform/graphics/compute.js index 8f9f3b0ed30..a5f168e6047 100644 --- a/src/platform/graphics/compute.js +++ b/src/platform/graphics/compute.js @@ -39,6 +39,28 @@ class Compute { dispatch(x, y, z) { this.impl?.dispatch(x, y, z); } + + /** + * 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. + */ + 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); + } } export { Compute }; diff --git a/src/platform/graphics/constants.js b/src/platform/graphics/constants.js index 9f79399375c..1f46995460e 100644 --- a/src/platform/graphics/constants.js +++ b/src/platform/graphics/constants.js @@ -1831,3 +1831,9 @@ export const CHUNKAPI_1_60 = '1.60'; export const CHUNKAPI_1_62 = '1.62'; export const CHUNKAPI_1_65 = '1.65'; export const CHUNKAPI_1_70 = '1.70'; + +// 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/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/webgl/webgl-graphics-device.js b/src/platform/graphics/webgl/webgl-graphics-device.js index cfab7e88043..e3eef133bef 100644 --- a/src/platform/graphics/webgl/webgl-graphics-device.js +++ b/src/platform/graphics/webgl/webgl-graphics-device.js @@ -953,6 +953,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 994888cc9bf..6632d202ac3 100644 --- a/src/platform/graphics/webgpu/webgpu-buffer.js +++ b/src/platform/graphics/webgpu/webgpu-buffer.js @@ -13,13 +13,22 @@ class WebgpuBuffer { */ buffer = null; + init(device, options) { + this.buffer = device.wgpu.createBuffer(options); + } + destroy(device) { if (this.buffer) { + this.buffer.unmap(); this.buffer.destroy(); this.buffer = null; } } + get size() { + return this.buffer ? this.buffer.size : 0; + } + get initialized() { return !!this.buffer; } @@ -84,6 +93,41 @@ class WebgpuBuffer { // TODO: handle usage types: // - 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. + * + * @returns {ArrayBuffer|undefined} The mapped range. + */ + getMappedRange() { + if (!this.buffer) { + return; + } + + return this.buffer.getMappedRange(); + } } export { WebgpuBuffer }; diff --git a/src/platform/graphics/webgpu/webgpu-compute.js b/src/platform/graphics/webgpu/webgpu-compute.js index 2a1d91d93f9..167a2646e72 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 { pixelFormatInfo, BUFFER_USAGE_COPY_DST, BUFFER_USAGE_MAP_READ } from "../constants.js"; /** * A WebGPU implementation of the Compute. @@ -7,13 +9,23 @@ import { BindGroup } from "../bind-group.js"; * @ignore */ class WebgpuCompute { + /** + * @ignore + */ + copyTextureToBufferCommands = []; + + /** + * @ignore + */ + copyBufferToBufferCommands = []; + 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}`); @@ -22,12 +34,15 @@ 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) { - - // 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; @@ -39,7 +54,69 @@ class WebgpuCompute { passEncoder.setPipeline(this.pipeline); passEncoder.dispatchWorkgroups(x, y, z); - device.endComputePass(); + 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 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. + */ + 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; + + // Calculate the size of the buffer to hold the texture data + const bufferSize = bytesPerRow * texture.height; + + const buffer = new Buffer(this.compute.device, { + size: bufferSize, + usage: BUFFER_USAGE_COPY_DST | BUFFER_USAGE_MAP_READ + }); + + const textureCopyView = { + texture: texture.impl.gpuTexture, + origin: { x: 0, y: 0 } + }; + const bufferCopyView = { + buffer: buffer.impl.buffer, + bytesPerRow: bytesPerRow + }; + const extent = { + width: texture.width, + height: texture.height + }; + + this.copyTextureToBufferCommands.push([textureCopyView, bufferCopyView, extent]); + + 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; } } diff --git a/src/platform/graphics/webgpu/webgpu-graphics-device.js b/src/platform/graphics/webgpu/webgpu-graphics-device.js index 85dca2c2435..706fc016291 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -27,6 +27,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 { /** @@ -90,6 +91,16 @@ class WebgpuGraphicsDevice extends GraphicsDevice { */ limits; + /** + * @ignore + */ + copyTextureToBufferCommands = []; + + /** + * @ignore + */ + copyBufferToBufferCommands = []; + constructor(canvas, options = {}) { super(canvas, options); options = this.initOptions; @@ -378,6 +389,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); } @@ -666,7 +683,6 @@ class WebgpuGraphicsDevice extends GraphicsDevice { } startComputePass() { - WebgpuDebug.internal(this); WebgpuDebug.validate(this); @@ -678,6 +694,9 @@ class WebgpuGraphicsDevice extends GraphicsDevice { // clear cached encoder state this.pipeline = null; + this.copyTextureToBufferCommands.length = 0; + this.copyBufferToBufferCommands.length = 0; + // TODO: add performance queries to compute passes // start the pass @@ -689,11 +708,17 @@ class WebgpuGraphicsDevice extends GraphicsDevice { } endComputePass() { - // end the compute pass this.passEncoder.end(); this.passEncoder = null; - this.insideRenderPass = false; + + // 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); + } + 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; @@ -703,9 +728,15 @@ 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; + this.insideRenderPass = false; + WebgpuDebug.end(this); WebgpuDebug.end(this); }