Support for indirect draw calls for WebGPU

examples/src/examples/compute/indirect-draw.compute-shader.wgsl

@@ -0,0 +1,37 @@
+// Indexed draw call parameters
+struct DrawIndexedIndirectArgs {
+    indexCount: u32,
+    instanceCount: u32,
+    firstIndex: u32,
+    baseVertex: i32,
+    firstInstance: u32
+};
+
+// Binding 0: uniform buffer holding draw call metadata and runtime config
+struct Uniforms {
+    indirectMetaData: vec4i,      // .x = indexCount, .y = firstIndex, .z = baseVertex
+    time: f32,                    // current time in seconds
+    maxInstanceCount: u32,        // max number of instances
+    indirectSlot: u32             // index into indirectDrawBuffer
+};
+@group(0) @binding(0) var<uniform> uniforms: Uniforms;
+
+// Binding 1: storage buffer to write draw args
+@group(0) @binding(1) var<storage, read_write> indirectDrawBuffer: array<DrawIndexedIndirectArgs>;
+
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) gid: vec3u) {
+    let metaData = uniforms.indirectMetaData;
+
+    // Generate oscillating instance count using a sine wave
+    let wave = abs(sin(uniforms.time));
+    let visibleCount = u32(wave * f32(uniforms.maxInstanceCount));
+
+    // generate draw call parameters based on metadata. Supply computed number of instances.
+    let index = uniforms.indirectSlot;
+    indirectDrawBuffer[index].indexCount = u32(metaData.x);
+    indirectDrawBuffer[index].instanceCount = visibleCount;
+    indirectDrawBuffer[index].firstIndex = u32(metaData.y);
+    indirectDrawBuffer[index].baseVertex = metaData.z;
+    indirectDrawBuffer[index].firstInstance = 0u;
+}

examples/src/examples/compute/indirect-draw.example.mjs

@@ -0,0 +1,190 @@
+// @config DESCRIPTION This example shows a basic usage of indirect drawing, and the compute shader changes the number of instances that are rendered.
+// @config WEBGL_DISABLED
+import files from 'examples/files';
+import { deviceType, rootPath } from 'examples/utils';
+import * as pc from 'playcanvas';
+
+const canvas = /** @type {HTMLCanvasElement} */ (document.getElementById('application-canvas'));
+window.focus();
+
+const assets = {
+    helipad: new pc.Asset(
+        'helipad-env-atlas',
+        'texture',
+        { url: `${rootPath}/static/assets/cubemaps/helipad-env-atlas.png` },
+        { type: pc.TEXTURETYPE_RGBP, mipmaps: false }
+    )
+};
+
+const gfxOptions = {
+    deviceTypes: [deviceType]
+};
+
+const device = await pc.createGraphicsDevice(canvas, gfxOptions);
+device.maxPixelRatio = Math.min(window.devicePixelRatio, 2);
+
+const createOptions = new pc.AppOptions();
+createOptions.graphicsDevice = device;
+
+createOptions.componentSystems = [pc.RenderComponentSystem, pc.CameraComponentSystem];
+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(() => {
+    app.start();
+
+    // setup skydome
+    app.scene.skyboxMip = 2;
+    app.scene.exposure = 0.7;
+    app.scene.envAtlas = assets.helipad.resource;
+
+    // Create an Entity with a camera component
+    const camera = new pc.Entity();
+    camera.addComponent('camera', {
+        toneMapping: pc.TONEMAP_ACES
+    });
+    app.root.addChild(camera);
+    camera.translate(0, 0, 10);
+
+    // create standard material that will be used on the instanced spheres
+    const material = new pc.StandardMaterial();
+    material.diffuse = new pc.Color(1, 1, 0.5);
+    material.gloss = 1;
+    material.metalness = 1;
+    material.useMetalness = true;
+    material.update();
+
+    // Create a Entity with a sphere render component and the material
+    const sphere = new pc.Entity('InstancingEntity');
+    sphere.addComponent('render', {
+        material: material,
+        type: 'sphere'
+    });
+    app.root.addChild(sphere);
+
+    // number of instances to render
+    const instanceCount = 1000;
+
+    // store matrices for individual instances into array
+    const matrices = new Float32Array(instanceCount * 16);
+    let matrixIndex = 0;
+
+    const radius = 5;
+    const pos = new pc.Vec3();
+    const rot = new pc.Quat();
+    const scl = new pc.Vec3();
+    const matrix = new pc.Mat4();
+
+    for (let i = 0; i < instanceCount; i++) {
+        // generate positions / scales and rotations
+        pos.set(
+            Math.random() * radius - radius * 0.5,
+            Math.random() * radius - radius * 0.5,
+            Math.random() * radius - radius * 0.5
+        );
+        scl.set(0.2, 0.2, 0.2);
+        rot.setFromEulerAngles(0, 0, 0);
+        matrix.setTRS(pos, rot, scl);
+
+        // copy matrix elements into array of floats
+        for (let m = 0; m < 16; m++) matrices[matrixIndex++] = matrix.data[m];
+    }
+
+    // create static vertex buffer containing the matrices
+    const vbFormat = pc.VertexFormat.getDefaultInstancingFormat(app.graphicsDevice);
+    const vertexBuffer = new pc.VertexBuffer(app.graphicsDevice, vbFormat, instanceCount, {
+        data: matrices
+    });
+
+    // initialize instancing using the vertex buffer on meshInstance of the created sphere
+    const sphereMeshInst = sphere.render.meshInstances[0];
+    sphereMeshInst.setInstancing(vertexBuffer);
+
+    // create a compute shader which will be used to update the number of instances to be rendered each frame
+    const shader = device.supportsCompute ?
+        new pc.Shader(device, {
+            name: 'ComputeShader',
+            shaderLanguage: pc.SHADERLANGUAGE_WGSL,
+            cshader: files['compute-shader.wgsl'],
+
+            // format of a uniform buffer used by the compute shader
+            computeUniformBufferFormats: {
+                ub: new pc.UniformBufferFormat(device, [
+
+                    // metadata about the mesh (how many indicies it has and similar, used to generate draw call parameters)
+                    new pc.UniformFormat('indirectMetaData', pc.UNIFORMTYPE_IVEC4),
+
+                    // time to animate number of visible instances
+                    new pc.UniformFormat('time', pc.UNIFORMTYPE_FLOAT),
+
+                    // maximum number of instances
+                    new pc.UniformFormat('maxInstanceCount', pc.UNIFORMTYPE_UINT),
+
+                    // indirect slot into storage buffer which stored draw call parameters
+                    new pc.UniformFormat('indirectSlot', pc.UNIFORMTYPE_UINT)
+                ])
+            },
+
+            // format of a bind group, providing resources for the compute shader
+            computeBindGroupFormat: new pc.BindGroupFormat(device, [
+                // a uniform buffer we provided format for
+                new pc.BindUniformBufferFormat('ub', pc.SHADERSTAGE_COMPUTE),
+
+                // the buffer with indirect draw arguments
+                new pc.BindStorageBufferFormat('indirectDrawBuffer', pc.SHADERSTAGE_COMPUTE)
+            ])
+        }) :
+        null;
+
+    // Create an instance of the compute shader, and provide it with uniform values that do not change each frame
+    const compute = new pc.Compute(device, shader, 'ComputeModifyVB');
+    compute.setParameter('maxInstanceCount', instanceCount);
+    compute.setParameter('indirectMetaData', sphereMeshInst.getIndirectMetaData());
+
+    // Set an update function on the app's update event
+    let angle = 0;
+    let time = 0;
+    app.on('update', (dt) => {
+        time += dt;
+
+        // obtain available slot in the indirect draw buffer - this needs to be done each frame
+        const indirectSlot = app.graphicsDevice.getIndirectDrawSlot();
+
+        // and assign it to the mesh instance for all cameras (null parameter)
+        sphereMeshInst.setIndirect(null, indirectSlot);
+
+        // give compute shader the indirect draw buffer - this can change between frames, so assign it each frame
+        compute.setParameter('indirectDrawBuffer', app.graphicsDevice.indirectDrawBuffer);
+
+        // update compute shader parameters
+        compute.setParameter('time', time);
+        compute.setParameter('indirectSlot', indirectSlot);
+
+        // set up the compute dispatch
+        compute.setupDispatch(1);
+
+        // dispatch the compute shader
+        device.computeDispatch([compute], 'ComputeIndirectDraw');
+
+        // orbit camera around
+        angle += dt * 0.2;
+        camera.setLocalPosition(8 * Math.sin(angle), 0, 8 * Math.cos(angle));
+        camera.lookAt(pc.Vec3.ZERO);
+    });
+});
+
+export { app };

src/platform/graphics/graphics-device.js

@@ -31,6 +31,7 @@ import { DebugGraphics } from './debug-graphics.js';
  * @import { RenderTarget } from './render-target.js'
  * @import { Shader } from './shader.js'
  * @import { Texture } from './texture.js'
+ * @import { StorageBuffer } from './storage-buffer.js';
  */
 
 const _tempSet = new Set();
@@ -125,6 +126,15 @@ class GraphicsDevice extends EventHandler {
      */
     scope;
 
+    /**
+     * The maximum number of indirect draw calls that can be used within a single frame. Used on
+     * WebGPU only. This needs to be adjusted based on the maximum number of draw calls that can
+     * be used within a single frame. Defaults to 1024.
+     *
+     * @type {number}
+     */
+    maxIndirectDrawCount = 1024;
+
     /**
      * The maximum supported texture anisotropy setting.
      *
@@ -413,6 +423,14 @@ class GraphicsDevice extends EventHandler {
      */
     capsDefines = new Map();
 
+    /**
+     * A set of maps to clear at the end of the frame.
+     *
+     * @type {Set<Map>}
+     * @ignore
+     */
+    mapsToClear = new Set();
+
     static EVENT_RESIZE = 'resizecanvas';
 
     constructor(canvas, options) {
@@ -714,6 +732,31 @@ class GraphicsDevice extends EventHandler {
         this.vertexBuffers.length = 0;
     }
 
+    /**
+     * Retrieves the available slot in the {@link indirectDrawBuffer} used for indirect rendering,
+     * which can be utilized by a {@link Compute} shader to generate indirect draw parameters and by
+     * {@link MeshInstance#setIndirect} to configure indirect draw calls.
+     *
+     * @returns {number} - The slot used for indirect rendering.
+     */
+    getIndirectDrawSlot() {
+        return 0;
+    }
+
+    /**
+     * Returns the buffer used to store arguments for indirect draw calls. The size of the buffer is
+     * controlled by the {@link maxIndirectDrawCount} property. This buffer can be passed to a
+     * {@link Compute} shader along with a slot obtained by calling {@link getIndirectDrawSlot}, in
+     * order to prepare indirect draw parameters. Also see {@link MeshInstance#setIndirect}.
+     *
+     * Only available on WebGPU, returns null on other platforms.
+     *
+     * @type {StorageBuffer|null}
+     */
+    get indirectDrawBuffer() {
+        return null;
+    }
+
     /**
      * Queries the currently set render target on the device.
      *
@@ -776,6 +819,7 @@ class GraphicsDevice extends EventHandler {
      * @param {IndexBuffer} [indexBuffer] - The index buffer to use for the draw call.
      * @param {number} [numInstances] - The number of instances to render when using instancing.
      * Defaults to 1.
+     * @param {number} [indirectSlot] - The slot of the indirect buffer to use for the draw call.
      * @param {boolean} [first] - True if this is the first draw call in a sequence of draw calls.
      * When set to true, vertex and index buffers related state is set up. Defaults to true.
      * @param {boolean} [last] - True if this is the last draw call in a sequence of draw calls.
@@ -791,7 +835,7 @@ class GraphicsDevice extends EventHandler {
      *
      * @ignore
      */
-    draw(primitive, indexBuffer, numInstances, first = true, last = true) {
+    draw(primitive, indexBuffer, numInstances, indirectSlot, first = true, last = true) {
         Debug.assert(false);
     }
 
@@ -981,6 +1025,9 @@ class GraphicsDevice extends EventHandler {
      * @ignore
      */
     frameEnd() {
+        // clear all maps scheduled for end of frame clearing
+        this.mapsToClear.forEach(map => map.clear());
+        this.mapsToClear.clear();
     }
 
     /**

src/platform/graphics/null/null-graphics-device.js

@@ -96,7 +96,7 @@ class NullGraphicsDevice extends GraphicsDevice {
         return new NullRenderTarget(renderTarget);
     }
 
-    draw(primitive, indexBuffer, numInstances, first = true, last = true) {
+    draw(primitive, indexBuffer, numInstances, indirectSlot, first = true, last = true) {
     }
 
     setShader(shader, asyncCompile = false) {

src/platform/graphics/webgl/webgl-graphics-device.js

@@ -1675,7 +1675,7 @@ class WebglGraphicsDevice extends GraphicsDevice {
         gl.bindBuffer(gl.ELEMENT_ARRAY_BUFFER, bufferId);
     }
 
-    draw(primitive, indexBuffer, numInstances, first = true, last = true) {
+    draw(primitive, indexBuffer, numInstances, indirectSlot, first = true, last = true) {
 
         const shader = this.shader;
         if (shader) {