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_rowalignment 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§
- Dispatch
Grid - A concrete dispatch grid: the
(x, y, z)argument todispatch_workgroups. - Limits
- A conservative, portable subset of the WebGPU compute limits a planner
needs. Field names follow the WebGPU /
wgpulimit names. - Texture
RowLayout - 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_rowfield 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
valueup to the next multiple ofalign. - div_
ceil_ u32 - Integer ceil-division
ceil(a / b). Returns0whenb == 0. - div_
ceil_ u64 - Integer ceil-division for
u64. Returns0whenb == 0. - floor_
pow2 - Largest power of two
<= v(forv >= 1).floor_pow2(0) == 0. - plan_
dispatch_ 1d - Plan a 1-D dispatch covering
total_threadsitems withworkgroup_xthreads 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 noutput covered by atile.x * tile.yworkgroup, optionally with abatchZ-dimension. - plan_
workgroup_ 1d - Auto-tune a 1-D workgroup size: the largest power of two
<= preferredthat fits bothmax_workgroup_size_xandmax_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.