Webgpu Syntax Compute Pipeline

WebGPU Compute Pipeline

Create a compute pipeline, encode a compute pass, and dispatch workgroups with correct workgroup-count math. WebGPU 1.0-stable (Chrome 113+, Safari 26+, Firefox 141+).

Quick Reference

GPUComputePipelineDescriptor (device.createComputePipeline(descriptor))

createComputePipelineAsync(descriptor) returns Promise<GPUComputePipeline> and compiles off the content timeline. createComputePipeline(descriptor) returns a GPUComputePipeline synchronously.

GPUComputePassEncoder methods (from encoder.beginComputePass({ label?, timestampWrites? }))

Decision Tree

Direct vs indirect dispatch

Is the workgroup count known on the CPU when you encode the command?
├── YES (count is a JS number or a known constant)
│      ALWAYS use dispatchWorkgroups(x, y, z).
└── NO (count is produced by an earlier GPU pass, e.g. live particle count)
       ALWAYS use dispatchWorkgroupsIndirect(buffer, offset).
       The buffer holds 3 u32: workgroupCountX, Y, Z (12 bytes).
       Avoids a CPU readback round-trip.

Sizing the dispatch grid

You have N data elements and a WGSL @workgroup_size(WG).
├── dispatchWorkgroups arguments are WORKGROUP counts, NEVER element counts.
├── workgroupCount = Math.ceil(N / WG)
├── total invocations launched = workgroupCount * WG  (>= N, with a remainder)
└── The shader MUST guard: if (global_invocation_id.x >= N) { return; }
        because the last workgroup launches WG invocations even when N is not
        a multiple of WG.

Core Patterns

Pattern 1: ALWAYS pass workgroup counts to dispatchWorkgroups, NEVER invocation counts

dispatchWorkgroups(x, y, z) arguments are the number of WORKGROUPS, not the number of shader invocations. Total invocations = x * y * z * (@workgroup_size product).

const WG = 64;                          // matches WGSL @workgroup_size(64)
const elementCount = 100_000;
const workgroupCount = Math.ceil(elementCount / WG);   // 1563, NOT 100000
pass.dispatchWorkgroups(workgroupCount);

NEVER write dispatchWorkgroups(elementCount). That launches elementCount workgroups, each running 64 invocations, which is 64x too much work and reads far out of bounds.

Pattern 2: ALWAYS compute the grid with ceil and guard the remainder in the shader

Math.ceil(N / WG) rounds up so every element is covered. The final workgroup then launches invocations past N. The WGSL shader MUST bounds-check.

@group(0) @binding(0) var<storage, read_write> data: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= arrayLength(&data)) { return; }   // guard the ceil remainder
  data[i] = data[i] * 2.0;
}

NEVER omit the index guard. Without it the extra invocations write out of bounds.

Pattern 3: ALWAYS check maxComputeWorkgroupsPerDimension before a large dispatch

Each of x, y, z MUST NOT exceed device.limits.maxComputeWorkgroupsPerDimension (spec default 65535). A 1D dispatch over a large array can exceed it.

const max = device.limits.maxComputeWorkgroupsPerDimension;
let count = Math.ceil(elementCount / WG);
if (count > max) {
  // spread the count across the X and Y dimensions
  const x = max;
  const y = Math.ceil(count / max);
  pass.dispatchWorkgroups(x, y);          // shader must linearize x + y * x
} else {
  pass.dispatchWorkgroups(count);
}

Pattern 4: ALWAYS set pipeline and bind groups before dispatching

The encoder validates that a pipeline is set and that every bind group required by the pipeline layout is bound at dispatch time.

const pass = encoder.beginComputePass({ label: "double-pass" });
pass.setPipeline(computePipeline);        // first
pass.setBindGroup(0, bindGroup);          // then every required group
pass.dispatchWorkgroups(workgroupCount);  // then dispatch
pass.end();                               // then end, before encoder.finish()

Pattern 5: ALWAYS give the indirect buffer INDIRECT usage and a 12-byte payload

dispatchWorkgroupsIndirect reads exactly 3 consecutive u32 values (workgroupCountX, workgroupCountY, workgroupCountZ = 12 bytes) starting at indirectOffset. The buffer MUST be created with GPUBufferUsage.INDIRECT.

const indirectBuffer = device.createBuffer({
  label: "dispatch-args",
  size: 12,                                       // 3 * u32
  usage: GPUBufferUsage.INDIRECT | GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
});
// a prior compute pass writes [countX, countY, countZ] into this buffer
pass.dispatchWorkgroupsIndirect(indirectBuffer, 0);   // offset multiple of 4

Pattern 6: ALWAYS use createComputePipelineAsync for heavy compilation at load time

createComputePipelineAsync compiles off the content timeline so the frame loop does not stall. Use it during loading. See webgpu-core-pipeline-architecture.

const pipeline = await device.createComputePipelineAsync({
  label: "physics-step",
  layout: "auto",
  compute: { module, entryPoint: "main" },
});

Common Anti-Patterns

1. Passing invocation counts to dispatchWorkgroups. Writing dispatchWorkgroups(elementCount) instead of dispatchWorkgroups(ceil(elementCount / WG)) launches WG times too many workgroups. WHY it fails: total invocations = workgroupCount * @workgroup_size, so the GPU runs elementCount * WG invocations.

2. No remainder guard in the shader. Math.ceil(N / WG) makes the last workgroup launch invocations past N. WHY it fails: without if (gid.x >= N) { return; } those invocations write out of bounds and corrupt data or trigger a robustness clamp.

3. Wrong indirect buffer stride. Using a 16-byte or 20-byte payload for dispatchWorkgroupsIndirect. WHY it fails: the GPU reads exactly 12 bytes (3 u32); a draw-indirect layout (16 or 20 bytes) puts the wrong values into workgroupCountX/Y/Z.

Critical Warnings

• NEVER pass shader-invocation counts to dispatchWorkgroups. The arguments are WORKGROUP counts. Total invocations = workgroupCount * @workgroup_size.
• NEVER omit the if (index >= dataSize) { return; } guard in the WGSL shader when the data size is not a multiple of @workgroup_size.
• NEVER let any of x, y, z exceed device.limits.maxComputeWorkgroupsPerDimension. It is a validation error.
• NEVER use a 16-byte or 20-byte indirect payload for dispatchWorkgroupsIndirect. It reads exactly 3 u32 (12 bytes).
• NEVER call dispatchWorkgroupsIndirect with an indirectOffset that is not a multiple of 4, or with a buffer that lacks GPUBufferUsage.INDIRECT.
• NEVER confuse the WGSL @workgroup_size with the dispatch count. They are separate. @workgroup_size is fixed in the shader; the dispatch count is the number of workgroups. See webgpu-wgsl-compute-shaders.
• NEVER call encoder.finish() while a compute pass is still open. Call pass.end() first.

Reference Files

• references/methods.md : createComputePipeline, GPUComputePipelineDescriptor, and every GPUComputePassEncoder method with full signatures and constraints.
• references/examples.md : verified end-to-end compute pipeline, indirect dispatch, and workgroup-count computation code.
• references/anti-patterns.md : compute-pipeline mistakes with WHY-it-fails explanations.

Related Skills

• webgpu-core-pipeline-architecture : pipeline layouts, "auto" vs explicit, async compilation.
• webgpu-syntax-bind-groups : bind group layouts and resource binding for compute.
• webgpu-wgsl-compute-shaders : @compute, @workgroup_size, compute builtins.
• webgpu-impl-compute-usecases : particle systems, reductions, image processing.