GPUComputePassEncoder: dispatchWorkgroupsIndirect() method
Limited availability
This feature is not Baseline because it does not work in some of the most widely-used browsers.
Experimental: This is an experimental technology
Check the Browser compatibility table carefully before using this in production.
Secure context: This feature is available only in secure contexts (HTTPS), in some or all supporting browsers.
Note: This feature is available in Web Workers.
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 forGPUComputePassEncoder.dispatchWorkgroups()
. So for example:jsconst uint32 = new Uint32Array(3); uint32[0] = 25; // The X value uint32[1] = 1; // The Y value uint32[2] = 1; // The Z value // Write values into a GPUBuffer 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
None (Undefined
).
Validation
The following criteria must be met when calling dispatchWorkgroupsIndirect()
, otherwise a GPUValidationError
is generated and the GPUComputePassEncoder
becomes invalid:
indirectBuffer
'sGPUBuffer.usage
contains theGPUBufferUsage.INDIRECT
flag.indirectOffset
+ the total size specified by theX
,Y
, andZ
dimensions is less than or equal to theindirectBuffer
'sGPUBuffer.size
.indirectOffset
is a multiple of 4.
Examples
// Set global buffer size
const BUFFER_SIZE = 1000;
// Compute shader; note workgroup size of 64
const shader = `
@group(0) @binding(0)
var<storage, read_write> output: array<f32>;
@compute @workgroup_size(64)
...
`;
// ...
// Create GPUCommandEncoder to encode commands to issue to the GPU
const commandEncoder = device.createCommandEncoder();
// Initiate render pass
const passEncoder = commandEncoder.beginComputePass();
// Issue commands
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const uint32 = new Uint32Array(3);
// Note workgroupCountX is set based on the global buffer size and the shader workgroup count.
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);
// End the render pass
passEncoder.end();
// Copy output buffer to staging buffer
commandEncoder.copyBufferToBuffer(
output,
0, // Source offset
stagingBuffer,
0, // Destination offset
BUFFER_SIZE,
);
// End frame by passing array of command buffers to command queue for execution
device.queue.submit([commandEncoder.finish()]);
// ...
Specifications
Specification |
---|
WebGPU # dom-gpucomputepassencoder-dispatchworkgroupsindirect |
Browser compatibility
BCD tables only load in the browser
See also
- The WebGPU API