Skip to content

[Web] Pre-allocate TypedArray views for pod args in WebGPU dispatch #18961

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
gnguralnick wants to merge 3 commits into
base: main
Choose a base branch
from
Open

[Web] Pre-allocate TypedArray views for pod args in WebGPU dispatch #18961

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
0e6a873
[Web] Pre-allocate TypedArray views for pod args in WebGPU dispatch
Mar 31, 2026
ccb9b77
Address review: use BYTES_PER_ELEMENT, hoist podArgBytes
Apr 1, 2026
7d850cf
Address review: rename maxPodArgs, restore packDimX comment
Apr 2, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
29 changes: 16 additions & 13 deletions web/src/webgpu.ts
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -697,6 +697,14 @@ export class WebGPUContext {
bindGroupLayouts: [bindGroupLayout]
});

// Pre-allocate typed array views for pod args (reused across dispatches)
const numPodSlots = podArgIndices.length + 1; // +1 for packDimX
const podArgBytes = numPodSlots * Int32Array.BYTES_PER_ELEMENT;
const podArgsArrayBuffer = new ArrayBuffer(podArgBytes);
const i32ViewCached = new Int32Array(podArgsArrayBuffer);
const u32ViewCached = new Uint32Array(podArgsArrayBuffer);
const f32ViewCached = new Float32Array(podArgsArrayBuffer);

// Function to create the pipeline.
const createShaderFunc = (pipeline: GPUComputePipeline): Function => {
const submitShader = (...args: Array<GPUPointer | number>): void => {
Expand Down Expand Up @@ -756,35 +764,30 @@ export class WebGPUContext {
});
}

const sizeOfI32 = 4;
const bufBytes = (podArgIndices.length + 1) * sizeOfI32;
const podArgBuffer = this.getUniformFromPool(bufBytes);
const i32View = new Int32Array(podArgIndices.length + 1);
const u32View = new Uint32Array(i32View.buffer);
const f32View = new Float32Array(i32View.buffer);
const podArgBuffer = this.getUniformFromPool(podArgBytes);

for (let i = 0; i < podArgIndices.length; ++i) {
const value = args[podArgIndices[i]];
const dtype = finfo.arg_types[podArgIndices[i]];
if (dtype.startsWith("int")) {
i32View[i] = value;
i32ViewCached[i] = value;
} else if (dtype.startsWith("uint")) {
u32View[i] = value;
u32ViewCached[i] = value;
} else if (dtype.startsWith("float")) {
f32View[i] = value;
f32ViewCached[i] = value;
Comment on lines +773 to +777
Copy link
Copy Markdown
Contributor

@gemini-code-assist gemini-code-assist Bot Mar 31, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

The dtype.startsWith string operations are executed for every POD argument on every dispatch. Since the argument types are fixed for each shader, consider pre-calculating an array of type indicators (e.g., an enum or numeric constants) in the createShadeInternal scope. This would allow replacing the string operations with a faster numeric check in the submitShader loop, which is beneficial for workloads with many small dispatches.

Copy link
Copy Markdown
Contributor Author

@gnguralnick gnguralnick Apr 1, 2026 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not sure if this needs to be in scope here though it does seem like a good suggestion

} else {
throw Error("Unknown pod dtype " + dtype);
}
}
// always pass in dim z launching grid size in
u32View[podArgIndices.length] = packDimX;
this.device.queue.writeBuffer(podArgBuffer, 0, i32View.buffer);
// Pass the original grid X dimension so the shader can recover blockIdx.x from the z-split
u32ViewCached[podArgIndices.length] = packDimX;
this.device.queue.writeBuffer(podArgBuffer, 0, podArgsArrayBuffer);

bindGroupEntries.push({
binding: bufferArgIndices.length,
resource: {
buffer: podArgBuffer,
size: i32View.buffer.byteLength
size: podArgsArrayBuffer.byteLength
}
});

Expand Down
Loading