Can WebGPU really compete with server-side inference for production ML?

For models under ~500M parameters with proper quantization, absolutely. We measured 23ms inference on a quantized DistilBERT with WebGPU compute shaders on a mid-range laptop GPU, which is competitive with a cold-start serverless function when you factor in network latency.

What browsers support WebGPU compute shaders right now?

As of early 2026, Chrome 121+, Edge 121+, and Firefox 130+ ship with full WebGPU compute support. Safari has partial support behind a flag. For production, you should feature-detect with navigator.gpu and fall back to ONNX Runtime Web with WebGL.

How does WGSL compare to GLSL for compute workloads?

WGSL is more explicit about memory layout and has better safety guarantees. You lose some of the syntactic sugar from GLSL compute shaders, but you gain predictable behavior across vendors. The type system catches buffer layout mismatches at pipeline creation time rather than at runtime.

WebGPU Compute Shaders for ML Inference: Client-Side Model Execution

Why Not Just Use ONNX Runtime Web?

ONNX Runtime Web is a solid tool for production use. However, it is an abstraction over WebGPU (or WebGL, or WASM), and abstractions have costs. Profiling an ONNX Runtime Web session running a DistilBERT model reveals that roughly 18% of the total inference time was spent on operator dispatch overhead, memory allocation bookkeeping, and graph optimization passes that had already been applied to the model.

Writing compute shaders directly gives you control over:

Exact workgroup dimensions and dispatch geometry
Memory layout in storage and uniform buffers
Synchronization between passes
Quantized arithmetic without runtime dequantization

The tradeoff is that you're writing GPU code. There's no free lunch.

Setting Up the WebGPU Pipeline

First, the boilerplate. The adapter/device negotiation code is standard boilerplate; the interesting parts follow.

const adapter = await navigator.gpu.requestAdapter({
 powerPreference: 'high-performance'
});

const device = await adapter.requestDevice({
 requiredFeatures: ['timestamp-query'],
 requiredLimits: {
 maxStorageBufferBindingSize: 256 * 1024 * 1024,
 maxBufferSize: 256 * 1024 * 1024,
 maxComputeWorkgroupSizeX: 256,
 maxComputeInvocationsPerWorkgroup: 256,
 }
});

That maxStorageBufferBindingSize request is important, the default is 128MB on many implementations, and a quantized 125M parameter model can easily exceed that when you include KV cache buffers.

The Core: Matrix Multiplication in WGSL

The backbone of transformer inference is matrix multiplication. Here's a WGSL compute shader that does tiled matmul with shared memory, which is the minimum viable approach for acceptable performance:

@group(0) @binding(0) var<storage, read> a: array<f32>;
@group(0) @binding(1) var<storage, read> b: array<f32>;
@group(0) @binding(2) var<storage, read_write> result: array<f32>;

struct Dimensions {
 M: u32,
 N: u32,
 K: u32,
}
@group(0) @binding(3) var<uniform> dims: Dimensions;

const TILE_SIZE: u32 = 16u;

var<workgroup> tileA: array<array<f32, 16>, 16>;
var<workgroup> tileB: array<array<f32, 16>, 16>;

@compute @workgroup_size(16, 16, 1)
fn main(
 @builtin(global_invocation_id) gid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>,
) {
 let row = gid.x;
 let col = gid.y;
 var sum: f32 = 0.0;

 let numTiles = (dims.K + TILE_SIZE - 1u) / TILE_SIZE;

 for (var t: u32 = 0u; t < numTiles; t = t + 1u) {
 let tiledCol = t * TILE_SIZE + lid.y;
 let tiledRow = t * TILE_SIZE + lid.x;

 if (row < dims.M && tiledCol < dims.K) {
 tileA[lid.x][lid.y] = a[row * dims.K + tiledCol];
 } else {
 tileA[lid.x][lid.y] = 0.0;
 }

 if (tiledRow < dims.K && col < dims.N) {
 tileB[lid.x][lid.y] = b[tiledRow * dims.N + col];
 } else {
 tileB[lid.x][lid.y] = 0.0;
 }

 workgroupBarrier();

 for (var k: u32 = 0u; k < TILE_SIZE; k = k + 1u) {
 sum = sum + tileA[lid.x][k] * tileB[k][lid.y];
 }

 workgroupBarrier();
 }

 if (row < dims.M && col < dims.N) {
 result[row * dims.N + col] = sum;
 }
}

This is a standard tiled approach. The key detail most tutorials miss: those two workgroupBarrier() calls are non-negotiable. The first ensures the tile is fully loaded before computation; the second ensures computation is complete before the tile memory is reused. Remove either one and you'll get intermittent wrong results that are absolute hell to debug.

Going Quantized: INT4 Matmul

Running FP32 matmul works, but it's wasteful for inference. Most production models use INT4 or INT8 quantization. Here's where it gets interesting, WGSL doesn't have native int4 types, so you pack two int4 values into a single u8 (well, u32, since WGSL's smallest type is 32-bit).

fn unpack_int4(packed: u32, index: u32) -> f32 {
 let shift = (index % 8u) * 4u;
 let nibble = (packed >> shift) & 0xFu;
 // Sign-extend from 4-bit
 let signed = i32(nibble) - select(0, 16, nibble >= 8u);
 return f32(signed);
}

@compute @workgroup_size(256, 1, 1)
fn quantized_matmul(
 @builtin(global_invocation_id) gid: vec3<u32>,
) {
 let row = gid.x / dims.N;
 let col = gid.x % dims.N;

 if (row >= dims.M) { return; }

 var sum: f32 = 0.0;
 let groupSize: u32 = 32u;

 for (var k: u32 = 0u; k < dims.K; k = k + 1u) {
 let group_idx = k / groupSize;
 let scale = scales[row * (dims.K / groupSize) + group_idx];
 let zero = zeros[row * (dims.K / groupSize) + group_idx];

 let packed_idx = (row * dims.K + k) / 8u;
 let w = unpack_int4(weights_packed[packed_idx], k);
 let dequant = (w - zero) * scale;

 sum = sum + input[col * dims.K + k] * dequant;
 }

 result[row * dims.N + col] = sum;
}

The unpack_int4 function is doing the heavy lifting. We store 8 int4 values per u32, unpack on the fly, and apply group-wise dequantization. The group size of 32 matches what GPTQ and AWQ use, so you can load quantized models directly.

The Attention Mechanism

Self-attention is the other performance-critical piece. A common implementation error is as follows: they implement attention as three separate dispatches (QK^T, softmax, attnV). But you can fuse the softmax into the same pass as the Q*K^T multiply if you use the online softmax trick:

@compute @workgroup_size(256, 1, 1)
fn fused_attention(
 @builtin(global_invocation_id) gid: vec3<u32>,
 @builtin(local_invocation_id) lid: vec3<u32>,
) {
 let head = gid.x / seq_len;
 let query_pos = gid.x % seq_len;

 if (head >= num_heads) { return; }

 var max_score: f32 = -1e10;
 var sum_exp: f32 = 0.0;
 var acc: array<f32, 64>; // head_dim

 // Online softmax: single pass over K
 for (var k: u32 = 0u; k < seq_len; k = k + 1u) {
 var score: f32 = 0.0;
 for (var d: u32 = 0u; d < head_dim; d = d + 1u) {
 score += q[head * seq_len * head_dim + query_pos * head_dim + d]
 * keys[head * seq_len * head_dim + k * head_dim + d];
 }
 score = score / sqrt(f32(head_dim));

 let new_max = max(max_score, score);
 let correction = exp(max_score - new_max);
 sum_exp = sum_exp * correction + exp(score - new_max);

 for (var d: u32 = 0u; d < head_dim; d = d + 1u) {
 acc[d] = acc[d] * correction
 + exp(score - new_max)
 * values[head * seq_len * head_dim + k * head_dim + d];
 }
 max_score = new_max;
 }

 // Normalize
 for (var d: u32 = 0u; d < head_dim; d = d + 1u) {
 output[head * seq_len * head_dim + query_pos * head_dim + d]
 = acc[d] / sum_exp;
 }
}

The online softmax avoids a second pass over the attention scores to normalize. On a 512-token sequence with 12 heads, this fusion saved ~4ms per layer on an RTX 3060, which adds up fast across 6 or 12 layers.

The Dispatch Orchestration

Tying it all together in JavaScript. This is where you coordinate the multi-pass inference:

class WebGPUInferenceEngine {
 constructor(device, model) {
 this.device = device;
 this.pipelines = {};
 this.buffers = {};
 this.model = model;
 }

 async buildPipelines() {
 this.pipelines.embed = await this.createComputePipeline(embeddingShader);
 this.pipelines.layerNorm = await this.createComputePipeline(layerNormShader);
 this.pipelines.attention = await this.createComputePipeline(fusedAttentionShader);
 this.pipelines.ffn = await this.createComputePipeline(quantizedMatmulShader);
 this.pipelines.softmax = await this.createComputePipeline(finalSoftmaxShader);
 }

 runInference(inputTokens) {
 const encoder = this.device.createCommandEncoder();

 // Embedding lookup
 this.dispatchEmbedding(encoder, inputTokens);

 // Transformer layers
 for (let layer = 0; layer < this.model.numLayers; layer++) {
 this.dispatchLayerNorm(encoder, layer, 'pre_attn');
 this.dispatchAttention(encoder, layer);
 this.dispatchLayerNorm(encoder, layer, 'pre_ffn');
 this.dispatchFFN(encoder, layer);
 }

 // Final norm + logits
 this.dispatchLayerNorm(encoder, -1, 'final');
 this.dispatchLogits(encoder);

 this.device.queue.submit([encoder.finish()]);
 }
}

All the dispatches go into a single command encoder. The GPU driver handles dependency tracking between dispatches based on buffer usage, this is one area where WebGPU actually shines compared to doing this in CUDA, where you'd manually manage streams.

Benchmarks: WebGPU vs WebGL vs Server

I ran these on a 2024 MacBook Pro (M3 Pro) and a desktop with an RTX 3060, using a quantized DistilBERT (66M params, INT4):

Method	M3 Pro	RTX 3060	Notes
WebGPU Compute (ours)	19ms	23ms	Custom WGSL shaders
ONNX Runtime Web (WebGPU)	31ms	38ms	Same model, default config
ONNX Runtime Web (WebGL)	67ms	54ms	WebGL2 backend
Server (Lambda + ONNX)	45ms*	45ms*	*Includes ~30ms network
Server (Lambda cold start)	2.8s	2.8s	Warm-up penalty

The M3 Pro numbers are noteworthy. Apple's WebGPU implementation is genuinely excellent, their Metal backend for WebGPU is mature and well-optimized. The RTX 3060 numbers are slightly higher because Chrome's Dawn backend adds a thin Vulkan translation layer.

The real win isn't the raw latency, it's the elimination of cold starts and network variability. A P99 latency of 25ms in the browser beats a P99 of 3 seconds from a serverless function that might need to cold-start.

Memory Layout Pitfalls

A common pitfall worth highlighting: WGSL's storage buffer alignment rules are strict and poorly documented. Every struct member must be aligned to a multiple of its size, and the struct itself must be aligned to the largest member alignment. For a mixed f32/u32 struct:

struct LayerParams {
 weight_offset: u32, // offset 0, align 4
 bias_offset: u32, // offset 4, align 4
 scale: f32, // offset 8, align 4
 num_elements: u32, // offset 12, align 4
 // Total: 16 bytes, align 4, this is fine
}

// But if you add a vec4:
struct LayerParamsV2 {
 weight_offset: u32, // offset 0
 // 12 bytes padding!
 params: vec4<f32>, // offset 16, align 16
 num_elements: u32, // offset 32
 // 12 bytes padding!
 // Total: 48 bytes
}

That padding is silent and brutal. Your JavaScript-side buffer writes will be wrong if you don't account for it. Use offset_of annotations or just avoid vec types in uniform structs.

When NOT to Use This Approach

There are important limitations to consider. This approach is not suitable when:

Your model is larger than ~500M parameters. GPU memory in browser contexts is limited, and you're sharing it with the renderer.
You need deterministic results across devices. Floating-point behavior varies between GPU vendors in ways that matter for some applications.
Your users are on integrated graphics. Intel UHD 630 can run these shaders, but at 10x the latency. Feature-detect and have a fallback.
You need to keep model weights private. Everything shipped to the browser is inspectable. If your model is proprietary, this isn't for you.

What's Next

An active area of exploration is WebGPU implementation of flash attention for longer context windows. The challenge is that WGSL doesn't give you the fine-grained shared memory control that CUDA does, so the tiling strategy needs to be different. This will be worth examining once meaningful benchmarks become available.

The broader point is this: the industry spent a decade assuming that ML inference belongs on servers. For a growing category of models and use cases, the client GPU is not just adequate, it's superior. Zero latency variance, zero cold starts, zero server costs. The tooling is finally catching up to make this practical.

WebGPU Compute Shaders for ML Inference: Client-Side Model Execution