computeShader = createShaderModule( device, ` override tileSize = 512u; override tileBuffer = 1u; override workgroupSize = 128; struct UniformStruct { tilesize: vec2, step: vec2, buffer: i32, pixelSize: f32, normalization: f32 } @binding(0) @group(0) var uniforms: UniformStruct; @binding(1) @group(0) var terrainData: array; @binding(2) @group(0) var outputData: array; fn unbufferedIndex(ij: vec2) -> u32 { return (u32(ij.x) % tileSize) + u32(ij.y) * tileSize; } fn bufferedIndex(ij: vec2) -> u32 { let w = tileSize + 2u * tileBuffer; let ijbuf = vec2(ij + i32(tileBuffer)); return (ijbuf.x % w) + (ijbuf.y) * w; } @compute @workgroup_size(workgroupSize) fn main( @builtin(global_invocation_id) coord: vec3u) { var ij = vec2(i32(coord.x), i32(coord.x)); var dijz: vec3; if (uniforms.step.y == 0) { ij.x = select(0i, i32(tileSize - 1), uniforms.step.x < 0i); } else if (uniforms.step.x == 0) { ij.y = select(0i, i32(tileSize - 1), uniforms.step.y < 0i); } var hull: array, 64>; hull[0] = vec3(vec2(ij - uniforms.step), terrainData[bufferedIndex(ij - uniforms.step)]); var hullPtr = 0u; for (var i = 0u; i < tileSize; i = i + 1u) { let uidx = unbufferedIndex(ij); let bidx = bufferedIndex(ij); let z = terrainData[bidx]; let ijz = vec3(vec2(ij), z); dijz = hull[hullPtr] - ijz; var s0 = dijz.z * dijz.z / dot(dijz, dijz); s0 = select(-s0, s0, dijz.z > 0.0); while(hullPtr > 0) { dijz = hull[hullPtr - 1] - ijz; var s1 = dijz.z * dijz.z / dot(dijz, dijz); s1 = select(-s1, s1, dijz.z > 0.0); if (s0 > s1) { break; } s0 = s1; hullPtr -= 1u; } dijz = hull[hullPtr] - ijz; dijz = vec3(dijz.xy, dijz.z / uniforms.pixelSize); outputData[uidx] = outputData[uidx] + uniforms.normalization * exp(-dijz.z / length(dijz)); // Fail silently but gracefully if we overflow the stack hullPtr = min(hullPtr + 1u, 63u); hull[hullPtr] = ijz; ij = ij + uniforms.step; } }` )

Purpose-built for displays of data