From 0709cfad6d3f4100cd7746745a522cfdb2fef3ff Mon Sep 17 00:00:00 2001 From: Erik Dubbelboer Date: Sun, 28 Jan 2024 17:56:25 +0100 Subject: [PATCH] Change API, add inputTexture --- examples/src/examples/misc/compute-shader.mjs | 119 ++++++++++-------- 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 +- 7 files changed, 138 insertions(+), 63 deletions(-) create mode 100644 src/platform/graphics/buffer.js diff --git a/examples/src/examples/misc/compute-shader.mjs b/examples/src/examples/misc/compute-shader.mjs index 8003ba87b35..ef31d4c0913 100644 --- a/examples/src/examples/misc/compute-shader.mjs +++ b/examples/src/examples/misc/compute-shader.mjs @@ -4,7 +4,11 @@ import * as pc from 'playcanvas'; * @param {import('../../options.mjs').ExampleOptions} options - The example options. * @returns {Promise} The example application. */ -async function example({ canvas, deviceType, files, glslangPath, twgslPath }) { +async function example({ canvas, deviceType, files, glslangPath, twgslPath, assetPath }) { + const assets = { + rocks: new pc.Asset('rocks', 'texture', { url: assetPath + 'textures/seaside-rocks01-color.jpg' }), + }; + const gfxOptions = { deviceTypes: [deviceType], @@ -26,66 +30,75 @@ async function example({ canvas, deviceType, files, glslangPath, twgslPath }) { createOptions.componentSystems = [ pc.RenderComponentSystem, - pc.CameraComponentSystem ]; createOptions.resourceHandlers = [ // @ts-ignore pc.TextureHandler, - // @ts-ignore - 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 + const assetListLoader = new pc.AssetListLoader(Object.values(assets), app.assets); + assetListLoader.load(async () => { + app.start(); + + // Ensure canvas is resized when window changes size + const resize = () => app.resizeCanvas(); + window.addEventListener('resize', resize); + 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 + }); + + app.graphicsDevice.scope.resolve("outputTexture").setValue(texture); + 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), + ], [ + new pc.BindStorageTextureFormat('outputTexture', pc.PIXELFORMAT_RGBA8, pc.TEXTUREDIMENSION_2D), + ], { + compute: true + }); + + const compute = new pc.Compute(app.graphicsDevice, shader); + const buffer = compute.getBuffer(texture); + + app.graphicsDevice.startComputePass(); + compute.dispatch(width, height); + // TODO: potentially dispatch more compute work in the same pass. + app.graphicsDevice.endComputePass(); + + const data = await buffer.getMappedRange(); + + console.log(data); + + buffer.destroy(app.graphicsDevice); }); - 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); - const buffer = compute.getBuffer(texture); - - app.graphicsDevice.startComputePass(); - compute.dispatch(texture.width, texture.height); - app.graphicsDevice.endComputePass(); - - const data = await buffer.getMappedRange(); - - console.log(data); - - buffer.destroy(app.graphicsDevice); - return app; } @@ -94,12 +107,18 @@ class ComputeShaderExample { static WEBGPU_ENABLED = 'force'; static 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/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 db21985af5a..8781d701154 100644 --- a/src/platform/graphics/shader.js +++ b/src/platform/graphics/shader.js @@ -32,6 +32,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 299bffe5d43..602348a4bec 100644 --- a/src/platform/graphics/webgpu/webgpu-graphics-device.js +++ b/src/platform/graphics/webgpu/webgpu-graphics-device.js @@ -671,6 +671,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 @@ -686,12 +688,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; @@ -700,7 +701,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;