2d:null 2e:["$","title",null,{"children":"Blur Compute Shader with Textures | Mike Tahani | Observable"}] 33:T1a31,{ const { context, device, format } = await gpu.init(SIZE, SIZE) // input image const sampler = gpu.sampler(device, { minFilter: 'linear', magFilter: 'linear', mipmapFilter: 'linear', }) const texture = gpu.canvasTexture(device, sample.canvas) // blur const WORKGROUP_SIZE = SIZE const WORKGROUP_DISPATCH_COUNT = SIZE // Math.ceil((SIZE ** 2) / WORKGROUP_SIZE) const blurPipeline = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code: ` struct Uniforms { direction: u32, } @group(0) @binding(0) var tex_in: texture_2d; @group(0) @binding(1) var tex_out: texture_storage_2d; @group(0) @binding(2) var uniforms: Uniforms; const offsets = array(0, 1, 2, 3, 4); // k = 5 const weight = array(0.2270270270, 0.1945945946, 0.1216216216, 0.0540540541, 0.0162162162); fn blur(image_tex: texture_2d, pos: vec2u, direction: vec2u) -> vec4 { let size = textureDimensions(image_tex); var color = textureLoad(image_tex, pos, 0) * weight[0]; for (var i = 1; i < 5; i++) { let offset = vec2(offsets[i]) * direction; let o1 = pos.xy + offset; if (all(o1 < size)) { color += textureLoad(image_tex, o1, 0) * weight[i]; } let o2 = pos.xy - offset; if (all(o2 < size)) { color += textureLoad(image_tex, o2, 0) * weight[i]; } } return color; } const h = vec2u(0, 1); const v = vec2u(1, 0); @compute @workgroup_size(${WORKGROUP_SIZE}) fn cs( @builtin(workgroup_id) workgroup_id : vec3, @builtin(local_invocation_id) local_invocation_id : vec3, ) { let pos = vec2u(workgroup_id.x, local_invocation_id.x); let size = textureDimensions(tex_in, 0); if (all(pos < size)) { // let color = textureSampleLevel(tex_in, tex_sampler, vec2f(0), 0); // \`select\` is like a ternary with (false, true, condition) call sig let direction = select(h, v, bool(uniforms.direction)); let color = blur(tex_in, pos, direction); textureStore(tex_out, pos, color); } } ` }) } }) const [tex0, tex1] = util.arr(2, () => gpu.createTexture(device, { width: sample.canvas.width, height: sample.canvas.height }) ) const [buf0, buf1] = util.arr(2, b => { const buf = device.createBuffer({ label: `buf${b}`, size: 4, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST }) device.queue.writeBuffer(buf, 0, new Uint32Array([b])) return buf }) // bind groups // const computeConstants = device.createBindGroup({ // layout: blurPipeline.getBindGroupLayout(0), // entries: [ // { binding: 0, resource: sampler } // ] // }) const computeLayoutIdx = 0 const computeInitial = device.createBindGroup({ layout: blurPipeline.getBindGroupLayout(computeLayoutIdx), entries: [ { binding: 0, resource: texture.createView() }, { binding: 1, resource: tex0.createView() }, { binding: 2, resource: { buffer: buf0 } } ], }) const computeA = device.createBindGroup({ layout: blurPipeline.getBindGroupLayout(computeLayoutIdx), entries: [ { binding: 0, resource: tex0.createView() }, { binding: 1, resource: tex1.createView() }, { binding: 2, resource: { buffer: buf1 } } ], }) const computeB = device.createBindGroup({ layout: blurPipeline.getBindGroupLayout(computeLayoutIdx), entries: [ { binding: 0, resource: tex1.createView() }, { binding: 1, resource: tex0.createView() }, { binding: 2, resource: { buffer: buf0 } } ], }) // for rendering the final blurred output const module = device.createShaderModule({ code: ` struct VertexOut { @builtin(position) position: vec4f, @location(0) texcoord: vec2f, } @group(0) @binding(0) var tex_sampler: sampler; @group(0) @binding(1) var tex: texture_2d; @vertex fn vs( @builtin(vertex_index) vertexIndex : u32 ) -> VertexOut { let quad = array( vec2f(0.0, 0.0), vec2f(1.0, 0.0), vec2f(0.0, 1.0), vec2f(0.0, 1.0), vec2f(1.0, 0.0), vec2f(1.0, 1.0), ); let texcoord = quad[vertexIndex]; return VertexOut( vec4f((texcoord - 0.5) * 2, 0, 1), texcoord ); } @fragment fn fs(fsIn: VertexOut) -> @location(0) vec4f { return textureSample(tex, tex_sampler, fsIn.texcoord); } ` }) const pipeline = device.createRenderPipeline({ layout: 'auto', vertex: { module, entryPoint: 'vs', }, fragment: { module, entryPoint: 'fs', targets: [{ format }], }, }) const result = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ { binding: 0, resource: sampler }, { binding: 1, resource: tex1.createView() }, ], }) function render() { const encoder = device.createCommandEncoder(); const computePass = encoder.beginComputePass() computePass.setPipeline(blurPipeline) // computePass.setBindGroup(0, computeConstants) computePass.setBindGroup(computeLayoutIdx, computeInitial) computePass.dispatchWorkgroups(WORKGROUP_DISPATCH_COUNT) computePass.setBindGroup(computeLayoutIdx, computeA) computePass.dispatchWorkgroups(WORKGROUP_DISPATCH_COUNT) for (let i = 0; i < 4; i++) { computePass.setBindGroup(computeLayoutIdx, computeB) computePass.dispatchWorkgroups(WORKGROUP_DISPATCH_COUNT) computePass.setBindGroup(computeLayoutIdx, computeA) computePass.dispatchWorkgroups(WORKGROUP_DISPATCH_COUNT) } computePass.end() const pass = encoder.beginRenderPass({ colorAttachments: [ { view: context.getCurrentTexture().createView(), clearValue: [0, 0, 0, 1], loadOp: 'clear', storeOp: 'store', }, ], }); pass.setPipeline(pipeline); pass.setBindGroup(0, result); pass.draw(6); pass.end(); device.queue.submit([encoder.finish()]); } render() return htl.html`${[sample.canvas, context.canvas]}` }34:T678,gpu = { const init = async (width = 512, height = 512) => { const canvas = document.createElement('canvas'); canvas.width = width canvas.height = height const context = canvas.getContext('webgpu'); const adapter = await navigator.gpu.requestAdapter(); const device = await adapter.requestDevice(); const format = navigator.gpu.getPreferredCanvasFormat(); context.configure({ device, format, alphaMode: 'premultiplied', }); return { context, adapter, device, format } } const sampler = (device, options = {}) => { return device.createSampler({ magFilter: 'linear', minFilter: 'linear', ...options }) } const createTexture = ( device, { width, height, format = 'rgba8unorm', usage = GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.TEXTURE_BINDING, } ) => { return device.createTexture({ format, size: [width, height], usage }) } const canvasTexture = ( device, canvas, { format = 'rgba8unorm', flipY = true, usage = GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST | GPUTextureUsage.RENDER_ATTACHMENT } = {} ) => { const texture = device.createTexture({ format, size: [canvas.width, canvas.height], usage }) device.queue.copyExternalImageToTexture( { source: canvas, flipY }, { texture }, [canvas.width, canvas.height] ) return texture } return { init, sampler, canvasTexture, createTexture } }35:[] 31:{"id":"c3852d7a4aa39ab7","slug":"blur-compute-shader","title":"Blur Compute Shader with Textures","thumbnail":"dcab3936ff57ffe92b19c45e8c516e3c0f845e56e76fd2c2ac08deec9beb766f","access_level":"public","create_time":"$D2024-04-17T19:57:51.824Z","update_time":"$D2024-05-07T06:06:29.309Z","view_time":null,"first_public_time":"$D2024-04-23T14:52:11.692Z","has_unpublished_edits":null,"fork":false,"stars":0,"starred":false,"owner":{"type":"team","id":"f258f685ba4dff9a","login":"miketahani","name":"Mike Tahani","avatar_url":"https://avatars-next.observableusercontent.com/eaac9d07c89480b6348b8d4e15a8ada6cd440e3b8e1ba1b1a330f055aa368db8","tier":"starter_2024"},"version":544,"publish_version":null,"role":null,"authored":false,"authors":[{"type":"individual","id":"02c6b10033eb7a61","login":"miketahani","name":"Mike Tahani","avatar_url":"https://avatars-next.observableusercontent.com/8341f220c6016b9edcb3cf852d6526ce60b354ab1ec9cd907f76cfc54425f28e","tier":null}],"members":[{"type":"individual","id":"02c6b10033eb7a61","login":"miketahani","name":"Mike Tahani","avatar_url":"https://avatars-next.observableusercontent.com/8341f220c6016b9edcb3cf852d6526ce60b354ab1ec9cd907f76cfc54425f28e","tier":null,"role":"editor"},{"type":"team","id":"f258f685ba4dff9a","login":"miketahani","name":"Mike Tahani","avatar_url":"https://avatars-next.observableusercontent.com/eaac9d07c89480b6348b8d4e15a8ada6cd440e3b8e1ba1b1a330f055aa368db8","tier":"starter_2024","role":"viewer"}],"comments":[],"custom_thumbnail":true,"trashed":false,"lapsed":null,"subscribed":null,"forks":0,"collection":{"id":"b12e2b165ab97944","slug":"webgpu-basics","title":"WebGPU Basics","owner":{"id":"f258f685ba4dff9a","login":"miketahani","name":"Mike Tahani"}},"type":"notebook","body":{"title":"Blur Compute Shader with Textures","cells":[{"id":0,"value":"# Blur Compute Shader with Textures\nAdapted from the first part of this blog post: https://www.rastergrid.com/blog/2010/09/efficient-gaussian-blur-with-linear-sampling/","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined","output":"","database":"$undefined"},{"id":21,"value":"$33","pinned":true,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined","output":"$undefined","database":"$undefined"},{"id":54,"value":"SIZE = 256","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined","output":"$undefined","database":"$undefined"},{"id":41,"value":"sample.canvas","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined","output":"$undefined","database":"$undefined"},{"id":32,"value":"sample = {\n const canvas = document.createElement('canvas');\n canvas.width = SIZE;\n canvas.height = SIZE;\n const context = canvas.getContext('2d');\n\n // context.lineWidth = 3\n // context.strokeStyle = '#fff'\n \n for (let i = 0; i < 20; i++) {\n context.fillStyle = `hsl(${util.rand(360)}, 100%, 60%)`\n context.beginPath()\n const r = util.rand(SIZE / 8)\n const m = r * 1.1\n context.arc(util.rand(m, SIZE - m), util.rand(m, SIZE - m), r, 0, util.TWO_PI)\n context.fill()\n // context.stroke()\n }\n\n return context\n}","pinned":true,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined","output":"$undefined","database":"$undefined"},{"id":7,"value":"$34","pinned":true,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined","output":"$undefined","database":"$undefined"},{"id":9,"value":"util = ({\n TWO_PI: 2 * Math.PI,\n rand: (min = 1, max) => {\n if (max === undefined) {\n return Math.random() * min\n }\n return Math.random() * (max - min) + min\n },\n // random signed\n rands: () => (Math.random() - 0.5) * 2,\n arr: (size, callback) => {\n const arr = new Array(size)\n if (typeof callback !== 'function') {\n return arr.fill(callback)\n }\n return arr.fill(null).map((_, i) => callback(i))\n },\n})","pinned":true,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined","output":"$undefined","database":"$undefined"}],"files":"$W35","stdlib":"1","titling":"auto"}} 32:"$31:body"