GPUComputePassEncoder: dispatchWorkgroupsIndirect() method
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()
).
Syntax
dispatchWorkgroupsIndirect(indirectBuffer, indirectOffset)
Parameters
-
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
.
Return value
Validation
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.
Examples
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()]);
Specifications
Browser compatibility
|
Desktop |
Mobile |
|
Chrome |
Edge |
Firefox |
Internet Explorer |
Opera |
Safari |
WebView Android |
Chrome Android |
Firefox for Android |
Opera Android |
Safari on IOS |
Samsung Internet |
dispatchWorkgroupsIndirect |
113Currently supported on ChromeOS, macOS, and Windows only.
|
113Currently supported on ChromeOS, macOS, and Windows only.
|
previewCurrently supported on Linux and Windows only.
|
No |
99Currently supported on ChromeOS, macOS, and Windows only.
|
No |
No |
No |
No |
No |
No |
No |
See also