Skip to main content

Module planner

Module planner 

Source
Expand description

Host-side dispatch and workgroup planning for WebGPU compute.

Everything in this module is pure arithmetic over adapter limits and problem shapes — no wgpu device, queue, or adapter is required, so the logic is fully unit-testable on CPU. The planners answer questions a WebGPU backend must resolve before recording a compute pass:

  • how large a workgroup may be, given the adapter’s max_compute_invocations_per_workgroup / per-axis caps (auto-tuning);
  • how many workgroups to dispatch for a 1-D / 2-D / 3-D problem while staying below WebGPU’s 65 535-per-axis limit;
  • how a linear problem must be folded into a 2-D grid when it would otherwise exceed that per-axis cap;
  • texture upload row pitch under WebGPU’s 256-byte bytes_per_row alignment rule.

These mirror the constants in crate::shader (e.g. the reductions use a 256-thread workgroup and a 2-D dispatch decode wgid.y * grid_x + wgid.x).

Structs§

DispatchGrid
A concrete dispatch grid: the (x, y, z) argument to dispatch_workgroups.
Limits
A conservative, portable subset of the WebGPU compute limits a planner needs. Field names follow the WebGPU / wgpu limit names.
TextureRowLayout
Texture upload layout under WebGPU’s 256-byte row-alignment rule.
Workgroup1D
A planned 1-D workgroup size: a single @workgroup_size(x) value.
Workgroup2D
A planned 2-D workgroup size for tiled kernels (@workgroup_size(x, y)).

Constants§

COPY_BYTES_PER_ROW_ALIGNMENT
WebGPU’s required alignment (in bytes) for the bytes_per_row field of a texture copy (COPY_BYTES_PER_ROW_ALIGNMENT).
MAX_WORKGROUPS_PER_DIM
WebGPU’s hard cap on workgroup count per dispatch axis.

Functions§

align_up_u32
Round value up to the next multiple of align.
div_ceil_u32
Integer ceil-division ceil(a / b). Returns 0 when b == 0.
div_ceil_u64
Integer ceil-division for u64. Returns 0 when b == 0.
floor_pow2
Largest power of two <= v (for v >= 1). floor_pow2(0) == 0.
plan_dispatch_1d
Plan a 1-D dispatch covering total_threads items with workgroup_x threads each, automatically folding into a 2-D grid when the required workgroup count would exceed the per-axis limit.
plan_dispatch_2d
Plan a 2-D tiled dispatch for an m x n output covered by a tile.x * tile.y workgroup, optionally with a batch Z-dimension.
plan_workgroup_1d
Auto-tune a 1-D workgroup size: the largest power of two <= preferred that fits both max_workgroup_size_x and max_invocations_per_workgroup.
plan_workgroup_square
Auto-tune a square 2-D workgroup tile for GEMM-style kernels.
texture_row_layout
Compute the 256-aligned row pitch for a texture upload/readback.