The dispatchWorkgroupsIndirect() method of the GPUComputePassEncoder interface dispatches a grid of workgroups, defined by the parameters of a GPUBuffer, to perform the work being done by the current GPUComputePipeline (i.e. set via GPUComputePassEncoder.setPipeline()).
dispatchWorkgroupsIndirect(indirectBuffer, indirectOffset)
indirectBuffer -
A GPUBuffer containing the X, Y, and Z dimensions of the grid of workgroups to dispatch. The buffer must contain a tightly packed block of three 32-bit unsigned integer values representing the dimensions (12 bytes total), given in the same order as the arguments for GPUComputePassEncoder.dispatchWorkgroups(). So for example:
const uint32 = new Uint32Array(3);
uint32[0] = 25;
uint32[1] = 1;
uint32[2] = 1;
device.queue.writeBuffer(buffer, 0, uint32, 0, uint32.length);
indirectOffset -
The offset, in bytes, into indirectBuffer where the dimension data begins.
Note: The X, Y, and Z dimension values passed to GPUComputePassEncoder.dispatchWorkgroups() and dispatchWorkgroupsIndirect() are the number of workgroups to dispatch for each dimension, not the number of shader invocations to perform across each dimension. This matches the behavior of modern native GPU APIs, but differs from the behavior of OpenCL. This means that if a GPUShaderModule defines an entry point with @workgroup_size(4, 4), and work is dispatched to it with the call dispatchWorkgroupsIndirect(indirectBuffer); with indirectBuffer specifying X and Y dimensions of 8 and 8, the entry point will be invoked 1024 times total — Dispatching a 4 x 4 workgroup 8 times along both the X and Y axes. 4 * 4 * 8 * 8 = 1024.
The following criteria must be met when calling dispatchWorkgroupsIndirect(), otherwise a GPUValidationError is generated and the GPUComputePassEncoder becomes invalid:
-
indirectBuffer's GPUBuffer.usage contains the GPUBufferUsage.INDIRECT flag. -
indirectOffset + the total size specified by the X, Y, and Z dimensions is less than or equal to the indirectBuffer's GPUBuffer.size. -
indirectOffset is a multiple of 4.
const BUFFER_SIZE = 1000;
const shader = `
@group(0) @binding(0)
var<storage, read_write> output: array<f32>;
@compute @workgroup_size(64)
...
`;
const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const uint32 = new Uint32Array(3);
uint32[0] = Math.ceil(BUFFER_SIZE / 64);
uint32[1] = 1;
uint32[2] = 1;
const workgroupDimensions = device.createBuffer({
size: 12,
usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.INDIRECT,
});
device.queue.writeBuffer(workgroupDimensions, 0, uint32, 0, uint32.length);
passEncoder.dispatchWorkgroupsIndirect(workgroupDimensions, 0);
passEncoder.end();
commandEncoder.copyBufferToBuffer(
output,
0,
stagingBuffer,
0,
BUFFER_SIZE,
);
device.queue.submit([commandEncoder.finish()]);