Image filters are stencils: each output pixel is a weighted sum of its neighborhood. A 9 x 9 Gaussian blur reads 81 texels per pixel, but it is separable: a horizontal and a vertical 9-tap pass give the same result with 18. Neighbors share most taps, so each workgroup of 64 loads its row segment plus a 4-texel apron per side into workgroup memory (72 loads), and every invocation reads its taps from there; an override turns the pass vertical. The Sobel operator estimates the brightness gradient from a 3 x 3 neighborhood (weights 1, 2, 1 across, -1, 0, +1 along), and the gradient's length is the edge strength:
// image-filters.wgsl: a tiled, separable Gaussian blur and a Sobel edge filter
override horizontal = true;
@group(0) @binding(0) var src: texture_2d<f32>;
@group(0) @binding(1) var dst: texture_storage_2d<rgba8unorm, write>;
const weight = array(0.2042, 0.1802, 0.1238, 0.0663, 0.0276); // Gaussian, sigma 2
var<workgroup> run: array<vec4f, 72>; // 64 texels plus a 4-texel apron each side
@compute @workgroup_size(64) fn blur(@builtin(workgroup_id) wg: vec3u,
@builtin(local_invocation_index) l: u32) {
let size = vec2i(textureDimensions(src));
let along = select(vec2i(0, 1), vec2i(1, 0), horizontal);
let start = along * i32(wg.x) * 64 + (vec2i(1) - along) * i32(wg.y);
for (var k = i32(l); k < 72; k += 64) { // 72 loads shared by 64 invocations
run[k] = textureLoad(src, clamp(start + along * (k - 4), vec2i(0), size - 1), 0);
}
workgroupBarrier();
var sum = run[l + 4] * weight[0];
for (var t = 1u; t < 5; t++) { sum += (run[l + 4 - t] + run[l + 4 + t]) * weight[t]; }
let p = start + along * i32(l);
if (all(p < size)) { textureStore(dst, p, sum); }
}
@compute @workgroup_size(8, 8) fn sobel(@builtin(global_invocation_id) id: vec3u) {
let size = vec2i(textureDimensions(src)); let p = vec2i(id.xy);
if (any(p >= size)) { return; }
var g = vec2f(0);
for (var y = -1; y <= 1; y++) {
for (var x = -1; x <= 1; x++) { // weights 1 2 1 across the gradient
let rgb = textureLoad(src, clamp(p + vec2i(x, y), vec2i(0), size - 1), 0).rgb;
let luma = dot(rgb, vec3f(0.299, 0.587, 0.114));
g += vec2f(f32(x * (2 - y * y)), f32(y * (2 - x * x))) * luma;
}
}
textureStore(dst, p, vec4f(vec3f(1 - clamp(length(g), 0, 1)), 1)); // dark edges
}Both entry points use textureLoad() with clamped coordinates, no sampler. Blurring first keeps the edge filter from reacting to noise.
<!doctype html>
<style>
body { margin: 0; background: #f7f4ee; font: 14px system-ui, sans-serif; }
.stage { position: relative; width: 100%; max-width: 600px; }
.stage canvas { display: block; width: 100%; }
.stage canvas + canvas { position: absolute; inset: 0; pointer-events: none; }
</style>
<div class="stage">
<canvas id="view" width="600" height="300"></canvas>
<canvas id="labels" width="600" height="300"></canvas>
</div>
<script>
const canvas = document.getElementById('view');
const ink = document.getElementById('labels').getContext('2d');
function showMessage(text) { // 2D fallback when WebGPU is missing
const ctx = canvas.getContext('2d');
ctx.fillStyle = '#fbeaea'; ctx.fillRect(0, 0, canvas.width, canvas.height);
ctx.fillStyle = '#8a2b2b'; ctx.font = '18px system-ui, sans-serif'; ctx.textAlign = 'center';
ctx.fillText(text, canvas.width / 2, canvas.height / 2);
}
const W = 180, H = 250;
// image-filters.wgsl: a tiled, separable Gaussian blur and a Sobel edge filter
const filters = /* wgsl */ `
override horizontal = true;
@group(0) @binding(0) var src: texture_2d<f32>;
@group(0) @binding(1) var dst: texture_storage_2d<rgba8unorm, write>;
const weight = array(0.2042, 0.1802, 0.1238, 0.0663, 0.0276); // Gaussian, sigma 2
var<workgroup> run: array<vec4f, 72>; // 64 texels plus a 4-texel apron each side
@compute @workgroup_size(64) fn blur(@builtin(workgroup_id) wg: vec3u,
@builtin(local_invocation_index) l: u32) {
let size = vec2i(textureDimensions(src));
let along = select(vec2i(0, 1), vec2i(1, 0), horizontal);
let start = along * i32(wg.x) * 64 + (vec2i(1) - along) * i32(wg.y);
for (var k = i32(l); k < 72; k += 64) { // 72 loads shared by 64 invocations
run[k] = textureLoad(src, clamp(start + along * (k - 4), vec2i(0), size - 1), 0);
}
workgroupBarrier();
var sum = run[l + 4] * weight[0];
for (var t = 1u; t < 5; t++) { sum += (run[l + 4 - t] + run[l + 4 + t]) * weight[t]; }
let p = start + along * i32(l);
if (all(p < size)) { textureStore(dst, p, sum); }
}
@compute @workgroup_size(8, 8) fn sobel(@builtin(global_invocation_id) id: vec3u) {
let size = vec2i(textureDimensions(src)); let p = vec2i(id.xy);
if (any(p >= size)) { return; }
var g = vec2f(0);
for (var y = -1; y <= 1; y++) {
for (var x = -1; x <= 1; x++) { // weights 1 2 1 across the gradient
let rgb = textureLoad(src, clamp(p + vec2i(x, y), vec2i(0), size - 1), 0).rgb;
let luma = dot(rgb, vec3f(0.299, 0.587, 0.114));
g += vec2f(f32(x * (2 - y * y)), f32(y * (2 - x * x))) * luma;
}
}
textureStore(dst, p, vec4f(vec3f(1 - clamp(length(g), 0, 1)), 1)); // dark edges
}`;
const show = /* wgsl */ `
@group(0) @binding(0) var image: texture_2d<f32>;
@vertex fn vs(@builtin(vertex_index) v: u32) -> @builtin(position) vec4f {
let p = array(vec2f(-1, -1), vec2f(3, -1), vec2f(-1, 3))[v];
return vec4f(p, 0, 1);
}
@group(0) @binding(1) var<uniform> origin: vec2f;
@fragment fn fs(@builtin(position) pos: vec4f) -> @location(0) vec4f { return textureLoad(image, vec2u(pos.xy - origin), 0); }`;
async function main() {
const adapter = await navigator.gpu?.requestAdapter();
if (!adapter) return showMessage('WebGPU is not available in this browser');
const device = await adapter.requestDevice();
const context = canvas.getContext('webgpu');
const format = navigator.gpu.getPreferredCanvasFormat();
context.configure({ device, format });
const T = GPUTextureUsage;
// A busy cover with a little noise, so the blur has something to calm down.
const art = new OffscreenCanvas(W, H), g = art.getContext('2d');
g.fillStyle = '#2d5d2a'; g.fillRect(0, 0, W, H);
g.fillStyle = '#f5ecd7'; g.fillRect(14, 60, W - 28, 80);
g.fillStyle = '#1b1b1b'; g.font = 'bold 26px Georgia, serif'; g.textAlign = 'center'; g.fillText('Night', W / 2, 94); g.fillText('Trains', W / 2, 124);
g.strokeStyle = '#e0c060'; g.lineWidth = 3; for (let k = 0; k < 5; k++) { g.beginPath(); g.arc(W / 2, 250, 30 + k * 16, Math.PI, 0); g.stroke(); }
const noise = g.getImageData(0, 0, W, H);
for (let i = 0; i < noise.data.length; i += 4) { const n = (Math.random() - 0.5) * 40; noise.data[i] += n; noise.data[i + 1] += n; noise.data[i + 2] += n; }
g.putImageData(noise, 0, 0);
const tex = (usage) => device.createTexture({ size: [W, H], format: 'rgba8unorm', usage: usage | T.TEXTURE_BINDING });
const source = tex(T.COPY_DST | T.RENDER_ATTACHMENT), halfway = tex(T.STORAGE_BINDING), blurred = tex(T.STORAGE_BINDING), edges = tex(T.STORAGE_BINDING);
device.queue.copyExternalImageToTexture({ source: art }, { texture: source }, [W, H]);
const module = device.createShaderModule({ code: filters });
const pipe = (entryPoint, constants) => device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint, constants } });
const blurH = pipe('blur', { horizontal: 1 }), blurV = pipe('blur', { horizontal: 0 }), sobel = pipe('sobel', {}); // an override flips the blur
const bind = (p, a, b) => device.createBindGroup({ layout: p.getBindGroupLayout(0), entries: [
{ binding: 0, resource: a.createView() }, { binding: 1, resource: b.createView() }] });
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(blurH); cp.setBindGroup(0, bind(blurH, source, halfway)); cp.dispatchWorkgroups(Math.ceil(W / 64), H);
cp.setPipeline(blurV); cp.setBindGroup(0, bind(blurV, halfway, blurred)); cp.dispatchWorkgroups(Math.ceil(H / 64), W);
cp.setPipeline(sobel); cp.setBindGroup(0, bind(sobel, blurred, edges)); cp.dispatchWorkgroups(Math.ceil(W / 8), Math.ceil(H / 8));
cp.end();
const showModule = device.createShaderModule({ code: show });
const showPipeline = device.createRenderPipeline({ layout: 'auto', vertex: { module: showModule }, fragment: { module: showModule, targets: [{ format }] } });
const pass = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
pass.setPipeline(showPipeline);
[source, blurred, edges].forEach((texture, i) => {
const x = 12 + i * 196, y = 32;
const o = device.createBuffer({ size: 8, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
device.queue.writeBuffer(o, 0, new Float32Array([x, y]));
pass.setBindGroup(0, device.createBindGroup({ layout: showPipeline.getBindGroupLayout(0), entries: [
{ binding: 0, resource: texture.createView() }, { binding: 1, resource: { buffer: o } }] }));
pass.setScissorRect(x, y, W, H);
pass.draw(3);
});
pass.end();
device.queue.submit([encoder.finish()]);
ink.font = '12px system-ui, sans-serif'; ink.fillStyle = '#222';
['original (with noise)', 'blur: horizontal + vertical 9-tap', 'Sobel on the blurred image'].forEach((t, i) => ink.fillText(t, 12 + i * 196, 22));
ink.fillStyle = '#555'; ink.font = '11px system-ui, sans-serif';
ink.fillText('Each blur workgroup loads 72 texels into workgroup memory for 64 outputs; blurring first keeps the edges clean.', 12, 296);
}
main();
</script>