compute render pipeline

compute.js

@@ -0,0 +1,145 @@
+import { getPath } from "./common.js";
+
+class Compute {
+  constructor(device, canvasFormat, wgsl)
+  {
+    const label = "compute";
+    this.label = label;
+
+    //////////////////////////////////////////////////////////////////////
+    // shader module
+    //////////////////////////////////////////////////////////////////////
+
+    const shaderModule = device.createShaderModule({
+      label: `${label} shader`,
+      code: wgsl,
+    });
+
+    //////////////////////////////////////////////////////////////////////
+    // bind group layout
+    //////////////////////////////////////////////////////////////////////
+
+    this.bindGroupLayout = {
+      compute: device.createBindGroupLayout({
+        label: `${label} bind group layout`,
+        entries: [
+          {
+            binding: 0,
+            visibility: GPUShaderStage.COMPUTE,
+            storageTexture: {
+              access: "write-only",
+              format: canvasFormat,
+              viewDimension: "2d",
+            }
+          }
+        ]
+      }),
+    };
+
+    //////////////////////////////////////////////////////////////////////
+    // pipeline
+    //////////////////////////////////////////////////////////////////////
+
+    const pipelineLayout = device.createPipelineLayout({
+      label: `${label} pipeline layout`,
+      bindGroupLayouts: [
+        this.bindGroupLayout.compute,
+      ],
+    });
+
+    this.computePipeline = device.createComputePipeline({
+      label: `${label} compute pipeline`,
+      layout: pipelineLayout,
+      compute: {
+        module: shaderModule,
+      },
+    });
+
+    this.textures = undefined;
+  }
+
+  recreateTextures(device, canvasTexture)
+  {
+    const label = this.label;
+
+    if (this.textures !== undefined && this.textures[0].width == canvasTexture.width && this.textures[0].height == canvasTexture.height) {
+      return;
+    }
+    if (this.textures !== undefined) {
+      for (let i = 0; i < 2; i++) {
+        this.textures[i].destroy();
+      }
+    }
+
+    //////////////////////////////////////////////////////////////////////
+    // textures
+    //////////////////////////////////////////////////////////////////////
+
+    this.textures = [];
+    for (let i = 0; i < 2; i++) {
+      const texture = device.createTexture({
+        label: `${label} texture ${i}`,
+        format: canvasTexture.format,
+        usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_SRC,
+        size: [canvasTexture.width, canvasTexture.height],
+      });
+      this.textures.push(texture);
+    }
+
+    //////////////////////////////////////////////////////////////////////
+    // bind group
+    //////////////////////////////////////////////////////////////////////
+
+    this.bindGroups = []
+    for (let i = 0; i < 2; i++) {
+      const bindGroup = device.createBindGroup({
+        label: `${label} bind group ${i}`,
+        layout: this.bindGroupLayout.compute,
+        entries: [{
+          binding: 0,
+          resource: this.textures[i].createView(),
+        }],
+      });
+      this.bindGroups.push(bindGroup);
+    }
+  }
+
+  render(device, frameNumber, canvasTexture)
+  {
+    this.recreateTextures(device, canvasTexture);
+
+    const label = this.label;
+    const encoder = device.createCommandEncoder({ label: `${label} command encoder` });
+    const computePass = encoder.beginComputePass({ label: `${label} compute pass` });
+
+    computePass.setPipeline(this.computePipeline);
+    computePass.setBindGroup(0, this.bindGroups[frameNumber]);
+    computePass.dispatchWorkgroups(canvasTexture.width / 16, canvasTexture.height / 16);
+    computePass.end();
+
+    encoder.copyTextureToTexture({
+      texture: this.textures[frameNumber],
+    }, {
+      texture: canvasTexture
+    }, {
+      width: canvasTexture.width,
+      height: canvasTexture.height,
+      depthOrArrayLayers: 1,
+    });
+
+    const commandBuffer = encoder.finish();
+    device.queue.submit([commandBuffer]);
+  }
+}
+
+async function loadCompute(device, canvasFormat)
+{
+  const canvasStorage = `@group(0) @binding(0) var out: texture_storage_2d<${canvasFormat}, write>;\n`
+  const wgsl = canvasStorage + await getPath("compute.wgsl");
+
+  const compute = new Compute(device, canvasFormat, wgsl);
+
+  return compute;
+}
+
+export { loadCompute };

compute.wgsl

@@ -1,56 +1,17 @@
-struct Configuration {
-  gridSize: vec2f,
-  aspect: vec2f,
-  matrix: mat4x4f
-};
-
-@group(0) @binding(0) var<uniform> config: Configuration;
-@group(0) @binding(1) var<storage> stateIn: array<u32>;
-@group(0) @binding(2) var<storage, read_write> stateOut: array<u32>;
-
-fn cellIndex(cell: vec2i) -> u32
-{
-  let x = (cell.x % i32(config.gridSize.x));
-  let y = (cell.y % i32(config.gridSize.y));
-  return u32(y * i32(config.gridSize.x) + x);
-}
-
-fn cellActive(cell: vec2i) -> u32
-{
-  return stateIn[cellIndex(cell)];
-}
-
-const directions: array<vec2i, 8> = array(vec2i(-1, -1),
-                                          vec2i(-1, 0),
-                                          vec2i(-1, 1),
-                                          vec2i(0, -1),
-                                          //0, 0)
-                                          vec2i(0, 1),
-                                          vec2i(1, -1),
-                                          vec2i(1, 0),
-                                          vec2i(1, 1));
 @compute
-@workgroup_size(16)
+@workgroup_size(16, 16)
-fn computeMain(@builtin(global_invocation_id) globalInvocationId: vec3u)
+fn computeMain(@builtin(global_invocation_id) id: vec3u)
 {
-  let cell: vec2i = vec2i(i32(globalInvocationId.x) % i32(config.gridSize.x),
+  let outSize = textureDimensions(out);
-                          i32(globalInvocationId.x) / i32(config.gridSize.y));
+  if (id.x >= outSize.x || id.y >= outSize.y) {
-
+    return;
-  var activeNeighbors: u32 = 0;
-  for (var i: u32 = 0; i < 8; i++) {
-    activeNeighbors += cellActive(directions[i] + cell);
   }
 
-  let index = cellIndex(cell);
+  let coordinate = vec2f(f32(id.x) + 0.5, f32(outSize.y - id.y) - 0.5);
-  switch activeNeighbors {
+  let outSizef = vec2f(outSize);
-    case 2: {
+  let uv = (coordinate * 2.0 - outSizef) / outSizef.y;
-      stateOut[index] = stateIn[index];
+
-    }
+  let color = vec3f(uv.xy, 0);
-    case 3: {
+
-      stateOut[index] = 1;
+  textureStore(out, id.xy, vec4f(color, 1.0));
-    }
-    default: {
-      stateOut[index] = 0;
-    }
-  }
 }

font.js

@@ -228,7 +228,6 @@ class FontRenderer {
       ],
     });
 
-    console.log(canvasFormat);
     this.renderPipelines = {}
     for (let i = 1; i <= 4; i += 3) {
       this.renderPipelines[i] = device.createRenderPipeline({

game-of-life.wgsl

@@ -0,0 +1,56 @@
+struct Configuration {
+  gridSize: vec2f,
+  aspect: vec2f,
+  matrix: mat4x4f
+};
+
+@group(0) @binding(0) var<uniform> config: Configuration;
+@group(0) @binding(1) var<storage> stateIn: array<u32>;
+@group(0) @binding(2) var<storage, read_write> stateOut: array<u32>;
+
+fn cellIndex(cell: vec2i) -> u32
+{
+  let x = (cell.x % i32(config.gridSize.x));
+  let y = (cell.y % i32(config.gridSize.y));
+  return u32(y * i32(config.gridSize.x) + x);
+}
+
+fn cellActive(cell: vec2i) -> u32
+{
+  return stateIn[cellIndex(cell)];
+}
+
+const directions: array<vec2i, 8> = array(vec2i(-1, -1),
+                                          vec2i(-1, 0),
+                                          vec2i(-1, 1),
+                                          vec2i(0, -1),
+                                          //0, 0)
+                                          vec2i(0, 1),
+                                          vec2i(1, -1),
+                                          vec2i(1, 0),
+                                          vec2i(1, 1));
+@compute
+@workgroup_size(16)
+fn computeMain(@builtin(global_invocation_id) globalInvocationId: vec3u)
+{
+  let cell: vec2i = vec2i(i32(globalInvocationId.x) % i32(config.gridSize.x),
+                          i32(globalInvocationId.x) / i32(config.gridSize.y));
+
+  var activeNeighbors: u32 = 0;
+  for (var i: u32 = 0; i < 8; i++) {
+    activeNeighbors += cellActive(directions[i] + cell);
+  }
+
+  let index = cellIndex(cell);
+  switch activeNeighbors {
+    case 2: {
+      stateOut[index] = stateIn[index];
+    }
+    case 3: {
+      stateOut[index] = 1;
+    }
+    default: {
+      stateOut[index] = 0;
+    }
+  }
+}

index.js

@@ -4,6 +4,7 @@ import { loadLight } from "./light.js";
 import { loadFlat } from "./flat.js";
 import { loadSnake } from "./snake.js";
 import { loadFont } from "./font.js";
+import { loadCompute } from "./compute.js";
 
 if (!navigator.gpu) {
   throw new Error("WebGPU not supported on this browser.");
@@ -14,19 +15,24 @@ if (!adapter) {
   throw new Error("No WebGPU adapter");
 }
 
-const device = await adapter.requestDevice();
+const hasBgraStorage = adapter.features.has("bgra8unorm-storage");
+const device = await adapter.requestDevice({
+  requiredFeatures: hasBgraStorage ? ["bgra8unorm-storage"] : []
+});
 
 const canvas = document.querySelector("canvas");
 canvas.width = canvas.clientWidth
 canvas.height = canvas.clientHeight
 
 const context = canvas.getContext("webgpu");
-const canvasFormat = navigator.gpu.getPreferredCanvasFormat();
+const canvasFormat = hasBgraStorage ? navigator.gpu.getPreferredCanvasFormat() : "rgba8unorm";
 context.configure({
   device: device,
   format: canvasFormat,
   alphaMode: 'premultiplied',
+  usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT
 });
+console.log("canvasFormat", canvasFormat);
 
 const encoder = device.createCommandEncoder();
 
@@ -151,7 +157,8 @@ module.instance.exports.camera_init(cameraStateAddress);
 //const lightRenderer = await loadLight(device, canvasFormat, viewUniformBuffer);
 //const flatRenderer = await loadFlat(device, canvasFormat, viewUniformBuffer, module);
 //const snakeRenderer = await loadSnake(device, canvasFormat, viewUniformBuffer, module);
-const fontRenderer = await loadFont(device, canvasFormat, viewUniformBuffer, module);
+//const fontRenderer = await loadFont(device, canvasFormat, viewUniformBuffer, module);
+const compute = await loadCompute(device, canvasFormat, module);
 
 const KEY = {
   A: 65,
@@ -333,8 +340,9 @@ function render2()
     canvas.width = canvas.clientWidth * window.devicePixelRatio;
     canvas.height = canvas.clientHeight * window.devicePixelRatio;
 
-    const configuration = fontRenderer.update(device, frameNumber);
+    //const configuration = fontRenderer.update(device, frameNumber);
-    const sampleCount = configuration.multisampleCount;
+    //const sampleCount = configuration.multisampleCount;
+    const sampleCount = 1;
 
     //const configuration = updateSliders(flatSliders);
     //const configuration = updateSliders(snakeSliders);
@@ -380,7 +388,7 @@ function render2()
     //snakeRenderer.update(device, frameNumber, mousePosition, configuration);
     //snakeRenderer.render(renderPass, frameNumber);
 
-    fontRenderer.render(renderPass, frameNumber, sampleCount);
+    //fontRenderer.render(renderPass, frameNumber, sampleCount);
 
     renderPass.end();
 
@@ -397,8 +405,10 @@ function render2()
       });
     }
 
-    const commandBuffer = encoder.finish();
+    //const commandBuffer = encoder.finish();
-    device.queue.submit([commandBuffer]);
+    //device.queue.submit([commandBuffer]);
+
+    compute.render(device, frameNumber, canvasTexture);
 
     frameNumber = (frameNumber + 1) % 2;