Change API, add inputTexture

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<pc.AppBase|null>} 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<rgba8unorm, write>;
+            @group(0) @binding(0) var inputTexture: texture_2d<f32>;
+            // @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<rgba8unorm, write>;
 
-            @compute @workgroup_size(1, 1, 1)
+            @compute @workgroup_size(1)
             fn main(@builtin(global_invocation_id) global_id : vec3u) {
-                let clearColor: vec4<f32> = vec4<f32>(0.5);
-                textureStore(outputTexture, vec2<i32>(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);
             }
         `
     };

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<Uint8Array|undefined>} The mapped range.
+     */
+    async getMappedRange() {
+        return await this.impl?.getMappedRange?.();
+    }
+}
+
+export { Buffer };

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);

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.
      *

src/platform/graphics/webgpu/webgpu-buffer.js

@@ -87,7 +87,10 @@ class WebgpuBuffer {
     }
 
     /**
-     * @returns {Promise<Uint8Array>}
+     * Returns a mapped range of the underlying buffer.
+     * On WebGPU this will wait for the buffer to be copied to the CPU.
+     *
+     * @returns {Promise<Uint8Array>} The mapped range.
      */
     async getMappedRange() {
         await this.buffer.mapAsync(GPUMapMode.READ);

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();

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;