// 5x5 convolution with a user-supplied kernel.
//
// Parameters: 25 floats stored row-major (top-left to bottom-right).
// Caller is responsible for kernel normalisation. All four channels are
// convolved together, which is the correct linear operation on
// premultiplied-alpha data; a kernel summing to zero therefore also
// zeroes coverage.
@compute @workgroup_size(WORKGROUP_X, WORKGROUP_Y)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let dims = vec2<u32>(uniforms.output_dimensions);
if gid.x >= dims.x || gid.y >= dims.y {
return;
}
let center = map_to_input(gid.xy);
var acc = vec4<f32>(0.0);
var idx: u32 = 0u;
for (var dy: i32 = -2; dy <= 2; dy = dy + 1) {
for (var dx: i32 = -2; dx <= 2; dx = dx + 1) {
acc += load_input(center + vec2<i32>(dx, dy)) * param(idx);
idx = idx + 1u;
}
}
textureStore(output_texture, vec2<i32>(gid.xy), acc);
}