vgpu4 symbolsView source ↗

Symbols in this topic

Compute

Compute pipeline created by compute(gpu). It uses the same WGSL reflection and set() ownership rules as render draws, then dispatch(x, y?, z?) — or dispatch({ indirect }) for GPU-driven counts — encodes and submits one compute pass.

Import

TypeScript
1
import type { Compute, ComputeOptions, DispatchOptions, StorageAccess, StorageBuffer, StorageOptions } from "vgpu";

Signature

TypeScript
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
interface ComputeOptions {
  readonly label?: string;
  readonly set?: Record<string, unknown>;
  readonly constants?: Readonly<Record<string, number | boolean>>;
  readonly entry?: string;
}
 
interface DispatchOptions {
  readonly indirect: StorageBuffer | { readonly buffer: StorageBuffer; readonly offset?: number };
}
 
interface Compute {
  set(values: Record<string, unknown>): this;
  dispatch(x: number, y?: number, z?: number): void;
  dispatch(opts: DispatchOptions): void;
}
 
type StorageAccess = "read" | "read-write";
 
interface StorageOptions {
  readonly access?: StorageAccess;
  readonly indirect?: boolean;
}
 
interface StorageBuffer {
  readonly size: number;
  readonly access: StorageAccess;
  read(): Promise<ArrayBuffer>;
  write(data: BufferSource): void;
}

Parameters

ParamTypeRequiredDefaultNotes
compute.sourcestring | ShaderSourceWGSL string or ShaderSource. Must include at least one @compute entry point.
compute.optsComputeOptions{}Initial compute options.
opts.labelstring"compute"Used in shader reflection, GPU labels, and error where fields.
opts.setRecord<string, unknown>undefinedInitial .set() call.
opts.constantsReadonly<Record<string, number | boolean>>WGSL defaultsConstructor-only values for WGSL override constants, applied to the compute stage. Use them to tune workgroup size per device or workload at pipeline creation: @workgroup_size(WG) with override WG: u32. Keying (@id(N) → decimal string of N) and number/boolean conversion match DrawOptions.constants.
opts.entrystringfirst @compute entry pointConstructor-only entry point selection when one WGSL module packs several @compute kernels sharing structs and bindings (e.g. emit/simulate/compact). The name must exist in the shader with the @compute stage. Binding visibility, bind group layouts, and the storage-aliasing preflight follow the selected entry.
compute.set.valuesRecord<string, unknown>Binding values by WGSL variable name. JS values are packed; buffers/resources are bound by identity.
compute.dispatch.xnumberWorkgroup count X passed to dispatchWorkgroups.
compute.dispatch.ynumber1Workgroup count Y.
compute.dispatch.znumber1Workgroup count Z.
compute.dispatch.opts.indirectStorageBuffer | { buffer, offset? }✔ in the overloadGPU-driven dispatch via dispatchWorkgroupsIndirect: the GPU reads [x, y, z] workgroup counts (3 tightly packed u32, 12 bytes) from the buffer at the byte offset (default 0). Use it when an earlier pass decides how much work exists — variable particle populations, stream compaction. Requires a buffer created with storage(gpu, bytes, { indirect: true }); offset must be a multiple of 4 and offset + 12 <= size. Cannot be combined with explicit counts.
storage.bytesnumberByte size for a main API (vgpu) storage buffer.
storage.accessStorageAccess | StorageOptions"read-write"Access string, or a StorageOptions bag with access and indirect. Stored on the resource facade and used by binding normalization.
storage.access.indirectbooleanfalseAppends the "indirect" buffer usage so the buffer can supply GPU-read draw/dispatch arguments.
storage.write.dataBufferSourceArrayBuffer or ArrayBufferView; writes at offset 0 in the public main API (vgpu) type.

Returns: compute(gpu) returns Compute; set() returns the same Compute; dispatch() returns void after submitting; storage(gpu) returns a main API (vgpu) StorageBuffer; StorageBuffer.read() resolves an ArrayBuffer copy.

Throws: VGPU-RING1-UNSUPPORTED when the shader has no @compute entry point; VGPU-INDIRECT-INVALID at dispatch time for a malformed indirect (neither a StorageBuffer nor { buffer, offset? }), a buffer created without the indirect flag (use storage(gpu, bytes, { indirect: true })), an offset that is not a non-negative integer multiple of 4, counts that do not fit the buffer (offset + 12 > size), or indirect combined with explicit workgroup counts in the same call; VGPU-CONSTANTS-INVALID for a malformed constants option (non-object value, a key that matches no override in the shader — the message lists the available overrides — or a value that is neither a finite number nor a boolean), and for an override declared without a default that constants does not provide; VGPU-ENTRY-INVALID for a non-string entry, a name that matches no entry point in the shader, or a name whose entry point is not @compute — the message lists the shader's available entry points with their stages; VGPU-R1-STORAGE-ALIASING when the same storage buffer is bound more than once and at least one reflected binding is writable; VGPU-R1-BINDING-NEVER-SET, VGPU-R1-OWNERSHIP-FLIP, and VGPU-R1-BINDING-INCOMPATIBLE-RESOURCE for binding errors; VGPU-SHADER-SOURCE-INVALID for malformed ShaderSource; TypeError if StorageBuffer.write() receives a non-buffer source.

Examples

TypeScript
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
import { init, compute, storage } from "vgpu/mock";
 
const gpu = await init();
const bytes = 4 * 16;
const src = storage(gpu, bytes, "read");
const dst = storage(gpu, bytes, "read-write");
src.write(new Float32Array(16));
 
const sim = compute(gpu, `
  @group(0) @binding(0) var<storage, read> src: array<vec4f>;
  @group(0) @binding(1) var<storage, read_write> dst: array<vec4f>;
  @compute @workgroup_size(1)
  fn cs_main(@builtin(global_invocation_id) id: vec3u) {
    dst[id.x] = src[id.x] + vec4f(1.0, 0.0, 0.0, 0.0);
  }
`, { label: "sim", set: { src, dst } });
 
sim.dispatch(4);
TypeScript
1
2
3
4
5
6
7
8
9
10
11
12
13
14
import { init, compute, pingPongStorage } from "vgpu/mock";
 
const gpu = await init();
const particles = pingPongStorage(gpu, 1024);
const step = compute(gpu, `
  @group(0) @binding(0) var<storage, read> src: array<u32>;
  @group(0) @binding(1) var<storage, read_write> dst: array<u32>;
  @compute @workgroup_size(64)
  fn cs_main(@builtin(global_invocation_id) id: vec3u) { dst[id.x] = src[id.x]; }
`);
 
step.set({ src: particles.read, dst: particles.write });
step.dispatch(Math.ceil(256 / 64));
particles.swap();
TypeScript
1
2
3
4
5
6
7
8
9
10
11
12
13
import { init, compute, storage } from "vgpu/mock";
 
const gpu = await init();
const wg = 64; // tune per device or workload without editing WGSL
const data = storage(gpu, 4 * 256);
const scale = compute(gpu, `
  override WG: u32 = 64;
  @group(0) @binding(0) var<storage, read_write> data: array<f32>;
  @compute @workgroup_size(WG)
  fn cs_main(@builtin(global_invocation_id) id: vec3u) { data[id.x] = data[id.x] * 2.0; }
`, { constants: { WG: wg }, set: { data } });
 
scale.dispatch(Math.ceil(256 / wg));

One JS constant drives both the pipeline's workgroup size and the dispatch math, so retuning wg cannot desynchronize them.

TypeScript
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
import { init, compute, storage } from "vgpu/mock";
 
const gpu = await init();
const alive = storage(gpu, 4, "read");             // live-particle count, e.g. from emission/compaction
const args = storage(gpu, 12, { indirect: true }); // [x, y, z] workgroup counts
const particles = storage(gpu, 4 * 1024);
 
const prepare = compute(gpu, `
  @group(0) @binding(0) var<storage, read> alive: u32;
  @group(0) @binding(1) var<storage, read_write> args: array<u32, 3>;
  @compute @workgroup_size(1) fn cs_main() {
    args[0] = (alive + 63u) / 64u; args[1] = 1u; args[2] = 1u; // one workgroup per 64 live particles
  }
`, { set: { alive, args } });
 
const step = compute(gpu, `
  @group(0) @binding(0) var<storage, read_write> particles: array<f32>;
  @compute @workgroup_size(64)
  fn cs_main(@builtin(global_invocation_id) id: vec3u) { particles[id.x] = particles[id.x] + 0.016; }
`, { set: { particles } });
 
prepare.dispatch(1);               // GPU computes how much work exists
step.dispatch({ indirect: args }); // GPU reads the counts; JS never sees them

GPU-driven dispatch: the first pass writes the workgroup counts from GPU-side state, so the population can vary every frame without a readback stall.

Notes

  • Use explicit dispatch(x, y, z) when the CPU already knows stable workgroup counts. Use dispatch({ indirect }) when a preceding GPU pass decides the count (compaction, particles), so no CPU readback is needed.
  • Declare storage read for source-only buffers and read-write for state that a kernel updates. For iterative simulation, bind pingPongStorage(gpu, ...) read/write pairs and swap after each step instead of aliasing one writable buffer.
  • StorageBuffer.read() is for tests, snapshots, and diagnostics; avoid awaiting it in a hot loop unless CPU synchronization is intentional.
  • Use pingPongStorage(gpu, bytes) when a compute step reads previous state and writes next state; binding the same writable storage identity twice is rejected before dispatch.
  • Bindings use compute visibility only when statically reachable from the selected compute entry point; unused declarations stay in the layout with visibility 0.
  • constants maps to GPUProgrammableStage.constants of the compute stage; the pipeline is created inside compute(gpu), so recreate the compute to change them.
  • Dispatch counts are forwarded to WebGPU; validate domain-specific bounds in your app.
  • Write indirect counts from another compute pass (bind the same buffer as storage) or from JS via write(). The same option shape drives GPU-driven draws via DrawCallOptions.indirect.
  • storage(gpu) creates storage buffers with copy_src and copy_dst, so they can be read back and rewritten from JS.
  • See also: compute, Draw.set, SharedUniforms, Target, StorageBuffer from vgpu/core.