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