Ported WebGPU blur path to use fragment shaders

blur4/shaders/blend.wgsl

@@ -1,3 +1,17 @@
+struct VertexOut {
+  @builtin(position) pos: vec4f,
+  @location(0) uv: vec2f,
+}
+
+@vertex
+fn vertMain(@builtin(vertex_index) vertexIndex: u32) -> VertexOut {
+  var pos = array<vec2f, 4>(
+    vec2f(-1.0, -1.0), vec2f(1.0, -1.0), vec2f(-1.0, 1.0), vec2f(1.0, 1.0),
+  );
+  var uv = pos[vertexIndex] * vec2f(0.5, -0.5) + 0.5;
+  return VertexOut(vec4f(pos[vertexIndex], 0, 1), uv);
+}
+
 struct ImageSize {
   width : i32,
   height : i32,
@@ -7,25 +21,19 @@ struct ImageSize {
 @group(0) @binding(0) var input : ${inputTextureType};
 @group(0) @binding(1) var blurred : texture_2d<f32>;
 @group(0) @binding(2) var mask : texture_2d<f32>;
-@group(0) @binding(3) var output : texture_storage_2d<${outputFormat}, write>;
-@group(0) @binding(4) var s : sampler;
-@group(0) @binding(5) var<uniform> size : ImageSize;
+@group(0) @binding(3) var s : sampler;
+@group(0) @binding(4) var<uniform> size : ImageSize;
 
 const k00 = f32(${k00});
 const kTileSize = ${tileSize}u;
 
-@compute @workgroup_size(kTileSize, kTileSize)
-fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
-  let coord = vec2<i32>(gid.xy);
-  if (coord.x >= size.width || coord.y >= size.height) {
-    return;
-  }
-
-  let coord_norm = (vec2<f32>(coord) + vec2<f32>(0.5)) * size.texel_size;
+@fragment
+fn fragMain(@location(0) uv: vec2f) -> @location(0) vec4f {
+  let coord_norm = uv + (0.5 * size.texel_size);
   var m = textureSampleLevel(mask, s, coord_norm, 0.0).r;
   var c = ${textureSampleCall};
   var b = textureSampleLevel(blurred, s, coord_norm, 0.0);
   b = b + (k00 * m) * c;
   b = b / b.a;
-  textureStore(output, coord, mix(b, c, m));
+  return mix(b, c, m);
 }

blur4/shaders/blur.wgsl

@@ -1,3 +1,17 @@
+struct VertexOut {
+  @builtin(position) pos: vec4f,
+  @location(0) uv: vec2f,
+}
+
+@vertex
+fn vertMain(@builtin(vertex_index) vertexIndex: u32) -> VertexOut {
+  var pos = array<vec2f, 4>(
+    vec2f(-1.0, -1.0), vec2f(1.0, -1.0), vec2f(-1.0, 1.0), vec2f(1.0, 1.0),
+  );
+  var uv = pos[vertexIndex] * vec2f(0.5, -0.5) + 0.5;
+  return VertexOut(vec4f(pos[vertexIndex], 0, 1), uv);
+}
+
 struct ImageSize {
   width : i32,
   height : i32,
@@ -6,17 +20,13 @@ struct ImageSize {
 
 @group(0) @binding(0) var input : ${inputTextureType};
 @group(0) @binding(1) var mask : texture_2d<f32>;
-@group(0) @binding(2) var output : texture_storage_2d<${outputFormat}, write>;
-@group(0) @binding(3) var s : sampler;
-@group(0) @binding(4) var<uniform> size : ImageSize;
+@group(0) @binding(2) var s : sampler;
+@group(0) @binding(3) var<uniform> size : ImageSize;
 
 const kRadius = ${radius}u;
 const kTileSize = ${tileSize}u;
 
-fn blur(sample_coordinate : vec2<i32>, dir : vec2<f32>, pass_no : i32) {
-  if (sample_coordinate.x >= size.width || sample_coordinate.y >= size.height) {
-    return;
-  }
+fn blur(sample_coordinate : vec2<i32>, dir : vec2<f32>, pass_no : i32) -> vec4f {
   let sample_coordinate_norm =
       (vec2<f32>(sample_coordinate) + vec2<f32>(0.5)) * size.texel_size;
 
@@ -52,19 +62,17 @@ fn blur(sample_coordinate : vec2<i32>, dir : vec2<f32>, pass_no : i32) {
       offset = offset + step;
   }
 
-  textureStore(output, sample_coordinate, color);
+  return color;
 }
 
-@compute @workgroup_size(kTileSize, kTileSize)
-fn main_horizontal(@builtin(global_invocation_id) gid : vec3<u32>) {
-  let coord = vec2<i32>(gid.xy);
+@fragment
+fn main_horizontal(@builtin(position) pos: vec4f) -> @location(0) vec4f {
   let dir = vec2<f32>(size.texel_size.x, 0.0);
-  blur(coord, dir, 0);
+  return blur(vec2i(pos.xy), dir, 0);
 }
 
-@compute @workgroup_size(kTileSize, kTileSize)
-fn main_vertical(@builtin(global_invocation_id) gid : vec3<u32>) {
-  let coord = vec2<i32>(gid.xy);
+@fragment
+fn main_vertical(@builtin(position) pos: vec4f) -> @location(0) vec4f {
   let dir = vec2<f32>(0.0, size.texel_size.y);
-  blur(coord, dir, 1);
+  return blur(vec2i(pos.xy), dir, 1);
 }

blur4/shaders/downscale.wgsl

@@ -0,0 +1,21 @@
+struct VertexOut {
+  @builtin(position) pos: vec4f,
+  @location(0) uv: vec2f,
+}
+
+@vertex
+fn vertMain(@builtin(vertex_index) vertexIndex: u32) -> VertexOut {
+  var pos = array<vec2f, 4>(
+    vec2f(-1.0, -1.0), vec2f(1.0, -1.0), vec2f(-1.0, 1.0), vec2f(1.0, 1.0),
+  );
+  var uv = pos[vertexIndex] * vec2f(0.5, -0.5) + 0.5;
+  return VertexOut(vec4f(pos[vertexIndex], 0, 1), uv);
+}
+
+@group(0) @binding(0) var inputTexture: ${inputTextureType};
+@group(0) @binding(1) var textureSampler: sampler;
+
+@fragment
+fn fragMain(@location(0) uv: vec2f) -> @location(0) vec4f {
+    return textureSampleBaseClampToEdge(inputTexture, textureSampler, uv);
+}

blur4/shaders/render.fragment.wgsl

@@ -2,7 +2,6 @@
 @group(0) @binding(1) var textureSampler: sampler;
 
 @fragment
-fn main(@builtin(position) coord: vec4<f32>) -> @location(0) vec4<f32> {
-    let uv = coord.xy / vec2<f32>(${width}.0, ${height}.0);
+fn main(@location(0) uv: vec2f) -> @location(0) vec4f {
     return textureSample(inputTexture, textureSampler, uv);
 }

blur4/shaders/render.vertex.wgsl

@@ -1,5 +1,10 @@
+struct VertexOut {
+    @builtin(position) pos: vec4f,
+    @location(0) uv: vec2f,
+}
+
 @vertex
-fn main(@builtin(vertex_index) vertexIndex: u32) -> @builtin(position) vec4<f32> {
+fn main(@builtin(vertex_index) vertexIndex: u32) -> VertexOut {
     var pos = array<vec2<f32>, 6>(
         vec2<f32>(-1.0, -1.0),
         vec2<f32>( 1.0, -1.0),
@@ -8,5 +13,6 @@ fn main(@builtin(vertex_index) vertexIndex: u32) -> @builtin(position) vec4<f32>
         vec2<f32>( 1.0,  1.0),
         vec2<f32>(-1.0,  1.0)
     );
-    return vec4<f32>(pos[vertexIndex], 0.0, 1.0);
+    var uv = pos[vertexIndex] * vec2f(0.5, -0.5) + 0.5;
+    return VertexOut(vec4f(pos[vertexIndex], 0.0, 1.0), uv);
 }

webgpu-blur.js

@@ -70,26 +70,60 @@ export class WebGPUBlur {
     const kernel = this.calculateKernel(this.radius);
     const kernelInitializer = kernel.join(', ');
 
+    const format = this.directOutput ? navigator.gpu.getPreferredCanvasFormat() : 'rgba8unorm';
+
     const blurHorizontalShader = await this.getBlurShader(this.radius, this.tileSize, kernel.length, kernelInitializer, true);
     const blurHorizontalModule = this.device.createShaderModule({ code: blurHorizontalShader });
-    this.pipelines.horizontal = await this.device.createComputePipelineAsync({
+    this.pipelines.horizontal = await this.device.createRenderPipelineAsync({
+      label: 'blurHorizontal',
       layout: 'auto',
-      compute: { module: blurHorizontalModule, entryPoint: 'main_horizontal' },
+      vertex: {
+        module: blurHorizontalModule,
+      },
+      primitive: {
+        topology: 'triangle-strip'
+      },
+      fragment: {
+        module: blurHorizontalModule,
+        entryPoint: 'main_horizontal',
+        targets: [{ format }]
+      },
     });
 
     const blurVerticalShader = await this.getBlurShader(this.radius, this.tileSize, kernel.length, kernelInitializer, false);
     const blurVerticalModule = this.device.createShaderModule({ code: blurVerticalShader });
-    this.pipelines.vertical = await this.device.createComputePipelineAsync({
+    this.pipelines.vertical = await this.device.createRenderPipelineAsync({
+      label: 'blurVertical',
       layout: 'auto',
-      compute: { module: blurVerticalModule, entryPoint: 'main_vertical' },
+      vertex: {
+        module: blurVerticalModule,
+      },
+      primitive: {
+        topology: 'triangle-strip'
+      },
+      fragment: {
+        module: blurVerticalModule,
+        entryPoint: 'main_vertical',
+        targets: [{ format }]
+      },
     });
 
     const k00 = kernel[0] * kernel[0];
     const blendShader = await this.getBlendShader(k00, this.tileSize);
-    const blendModule = this.device.createShaderModule({ code: blendShader });
-    this.pipelines.blend = await this.device.createComputePipelineAsync({
+    const blendModule = this.device.createShaderModule({ label: 'blend', code: blendShader });
+    this.pipelines.blend = await this.device.createRenderPipelineAsync({
+      label: 'blend',
       layout: 'auto',
-      compute: { module: blendModule, entryPoint: 'main' },
+      vertex: {
+        module: blendModule,
+      },
+      primitive: {
+        topology: 'triangle-strip'
+      },
+      fragment: {
+        module: blendModule,
+        targets: [{ format }]
+      },
     });
   }
 
@@ -158,53 +192,76 @@ export class WebGPUBlur {
     device.queue.writeBuffer(blurSizeBuffer, 0, blurSizeData);
     device.queue.writeBuffer(blurSizeBuffer, 8, blurTexelSizeData);
 
-    const horizontalTexture = getOrCreateTexture(device, this.resourceCache, 'horizontal', [blurWidth, blurHeight], this.directOutput, GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.TEXTURE_BINDING);
-    const blurredTexture = getOrCreateTexture(device, this.resourceCache, 'blurred', [blurWidth, blurHeight], this.directOutput, GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.TEXTURE_BINDING);
+    const horizontalTexture = getOrCreateTexture(device, this.resourceCache, 'horizontal', [blurWidth, blurHeight], this.directOutput, GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING);
+    const blurredTexture = getOrCreateTexture(device, this.resourceCache, 'blurred', [blurWidth, blurHeight], this.directOutput, GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.TEXTURE_BINDING);
 
-    const passEncoder = commandEncoder.beginComputePass();
+    let passEncoder = commandEncoder.beginRenderPass({
+      colorAttachments: [{
+        view: horizontalTexture.createView(),
+        loadOp: 'clear',
+        storeOp: 'store',
+      }]
+    });
 
     const horizontalBindGroup = device.createBindGroup({
       layout: this.pipelines.horizontal.getBindGroupLayout(0),
       entries: [
         { binding: 0, resource: this.zeroCopy ? inputTexture : inputTexture.createView() },
         { binding: 1, resource: maskTexture.createView() },
-        { binding: 2, resource: horizontalTexture.createView() },
-        { binding: 3, resource: this.sampler },
-        { binding: 4, resource: { buffer: blurSizeBuffer } },
+        { binding: 2, resource: this.sampler },
+        { binding: 3, resource: { buffer: blurSizeBuffer } },
       ],
     });
     passEncoder.setPipeline(this.pipelines.horizontal);
     passEncoder.setBindGroup(0, horizontalBindGroup);
-    passEncoder.dispatchWorkgroups(Math.ceil(blurWidth / this.tileSize), Math.ceil(blurHeight / this.tileSize));
+    passEncoder.draw(4);
+
+    passEncoder.end();
+
+    passEncoder = commandEncoder.beginRenderPass({
+      colorAttachments: [{
+        view: blurredTexture.createView(),
+        loadOp: 'clear',
+        storeOp: 'store',
+      }]
+    });
 
     const verticalBindGroup = device.createBindGroup({
       layout: this.pipelines.vertical.getBindGroupLayout(0),
       entries: [
         { binding: 0, resource: horizontalTexture.createView() },
         { binding: 1, resource: maskTexture.createView() },
-        { binding: 2, resource: blurredTexture.createView() },
-        { binding: 3, resource: this.sampler },
-        { binding: 4, resource: { buffer: blurSizeBuffer } },
+        { binding: 2, resource: this.sampler },
+        { binding: 3, resource: { buffer: blurSizeBuffer } },
       ],
     });
     passEncoder.setPipeline(this.pipelines.vertical);
     passEncoder.setBindGroup(0, verticalBindGroup);
-    passEncoder.dispatchWorkgroups(Math.ceil(blurWidth / this.tileSize), Math.ceil(blurHeight / this.tileSize));
+    passEncoder.draw(4);
+
+    passEncoder.end();
+
+    passEncoder = commandEncoder.beginRenderPass({
+      colorAttachments: [{
+        view: outputTexture.createView(),
+        loadOp: 'clear',
+        storeOp: 'store',
+      }]
+    });
 
     const blendBindGroup = device.createBindGroup({
       layout: this.pipelines.blend.getBindGroupLayout(0),
       entries: [
         { binding: 0, resource: this.zeroCopy ? inputTexture : inputTexture.createView() },
         { binding: 1, resource: blurredTexture.createView() },
         { binding: 2, resource: maskTexture.createView() },
-        { binding: 3, resource: outputTexture.createView() },
-        { binding: 4, resource: this.sampler },
-        { binding: 5, resource: { buffer: imageSizeBuffer } },
+        { binding: 3, resource: this.sampler },
+        { binding: 4, resource: { buffer: imageSizeBuffer } },
       ],
     });
     passEncoder.setPipeline(this.pipelines.blend);
     passEncoder.setBindGroup(0, blendBindGroup);
-    passEncoder.dispatchWorkgroups(Math.ceil(width / this.tileSize), Math.ceil(height / this.tileSize));
+    passEncoder.draw(6);
 
     passEncoder.end();
   }