From 6cf71c1846c70efb6f100d6e92e57ae536180dd3 Mon Sep 17 00:00:00 2001 From: James Price Date: Tue, 4 Jun 2024 16:51:51 -0400 Subject: [PATCH] temp - testing SPIR-V codegen fix --- src/nbody.ts | 134 ++++++++++++++++++++++++++++------------------- src/shaders.wgsl | 6 +++ 2 files changed, 87 insertions(+), 53 deletions(-) diff --git a/src/nbody.ts b/src/nbody.ts index 917e63c..5300947 100644 --- a/src/nbody.ts +++ b/src/nbody.ts @@ -1,5 +1,5 @@ -import { mat4, vec3 } from 'gl-matrix' -import shaders from './shaders.wgsl' +import { mat4, vec3 } from "gl-matrix"; +import shaders from "./shaders.wgsl"; // Simulation parameters. let numBodies; @@ -26,25 +26,26 @@ let positionsOut: GPUBuffer = null; let velocities: GPUBuffer = null; let renderParams: GPUBuffer = null; let computeBindGroup: GPUBindGroup = null; +let othergroup: GPUBindGroup = null; let renderBindGroup: GPUBindGroup = null; const init = async () => { // Initialize the WebGPU device. - const powerPref = document.getElementById('powerpref'); + const powerPref = document.getElementById("powerpref"); const adapter = await navigator.gpu.requestAdapter({ powerPreference: powerPref.selectedOptions[0].value, }); - device = await adapter.requestDevice() + device = await adapter.requestDevice(); queue = device.queue; // Set up the canvas context. - canvas = document.getElementById('canvas'); - canvasContext = canvas.getContext('webgpu'); -} + canvas = document.getElementById("canvas"); + canvasContext = canvas.getContext("webgpu"); +}; // Generate WGSL shader source. function getShaders() { - let preamble = '' + let preamble = ""; preamble += `const kWorkgroupSize = ${workgroupSize};\n`; preamble += `const kNumBodies = ${numBodies};\n`; return preamble + shaders; @@ -69,8 +70,7 @@ const updateRenderParams = async () => { // Generate the view projection matrix. let projectionMatrix = mat4.create(); let viewProjectionMatrix = mat4.create(); - mat4.perspectiveZO(projectionMatrix, - 1.0, canvas.width / canvas.height, 0.1, 50.0); + mat4.perspectiveZO(projectionMatrix, 1.0, canvas.width / canvas.height, 0.1, 50.0); mat4.translate(viewProjectionMatrix, viewProjectionMatrix, eyePosition); mat4.multiply(viewProjectionMatrix, projectionMatrix, viewProjectionMatrix); @@ -79,7 +79,7 @@ const updateRenderParams = async () => { let viewProjectionMatrixHost = new Float32Array(renderParamsHost); viewProjectionMatrixHost.set(viewProjectionMatrix); queue.writeBuffer(renderParams, 0, renderParamsHost); -} +}; function initPipelines() { // Reset pipelines. @@ -104,42 +104,44 @@ function initPipelines() { const positionsAttribute: GPUVertexAttribute = { shaderLocation: 0, offset: 0, - format: 'float32x4', + format: "float32x4", }; const positionsLayout: GPUVertexBufferLayout = { attributes: [positionsAttribute], arrayStride: 4 * 4, - stepMode: 'instance', + stepMode: "instance", }; renderPipeline = device.createRenderPipeline({ vertex: { module: module, - entryPoint: 'vs_main', + entryPoint: "vs_main", buffers: [positionsLayout], }, fragment: { module: module, - entryPoint: 'fs_main', - targets: [{ - format: navigator.gpu.getPreferredCanvasFormat(), - blend: { - color: { - operation: "add", - srcFactor: "one", - dstFactor: "one", - }, - alpha: { - operation: "add", - srcFactor: "one", - dstFactor: "one", + entryPoint: "fs_main", + targets: [ + { + format: navigator.gpu.getPreferredCanvasFormat(), + blend: { + color: { + operation: "add", + srcFactor: "one", + dstFactor: "one", + }, + alpha: { + operation: "add", + srcFactor: "one", + dstFactor: "one", + }, }, - } - }], + }, + ], }, primitive: { - frontFace: 'cw', - cullMode: 'none', - topology: 'triangle-list', + frontFace: "cw", + cullMode: "none", + topology: "triangle-list", }, layout: "auto", }); @@ -148,9 +150,9 @@ function initPipelines() { computePipeline = device.createComputePipeline({ compute: { module: module, - entryPoint: 'cs_main', + entryPoint: "cs_main", }, - layout: "auto" + layout: "auto", }); } @@ -159,17 +161,17 @@ function initBodies() { positionsIn = device.createBuffer({ size: numBodies * 4 * 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX, - mappedAtCreation: true + mappedAtCreation: true, }); positionsOut = device.createBuffer({ size: numBodies * 4 * 4, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX, - mappedAtCreation: false + mappedAtCreation: false, }); velocities = device.createBuffer({ size: numBodies * 4 * 4, usage: GPUBufferUsage.STORAGE, - mappedAtCreation: false + mappedAtCreation: false, }); // Generate initial positions on the surface of a sphere. @@ -177,7 +179,7 @@ function initBodies() { let positions = new Float32Array(positionsIn.getMappedRange()); for (let i = 0; i < numBodies; i++) { let longitude = 2.0 * Math.PI * Math.random(); - let latitude = Math.acos((2.0 * Math.random() - 1.0)); + let latitude = Math.acos(2.0 * Math.random() - 1.0); positions[i * 4 + 0] = kRadius * Math.sin(latitude) * Math.cos(longitude); positions[i * 4 + 1] = kRadius * Math.sin(latitude) * Math.sin(longitude); positions[i * 4 + 2] = kRadius * Math.cos(latitude); @@ -203,7 +205,7 @@ function draw() { const timeSinceLastLog = now - lastFpsUpdateTime; if (timeSinceLastLog >= kFpsUpdateInterval) { const fps = numFramesSinceFpsUpdate / (timeSinceLastLog / 1000.0); - document.getElementById("fps").innerHTML = fps.toFixed(1) + ' FPS'; + document.getElementById("fps").innerHTML = fps.toFixed(1) + " FPS"; lastFpsUpdateTime = performance.now(); numFramesSinceFpsUpdate = 0; } @@ -215,9 +217,9 @@ function draw() { // Update render parameters based on key presses. if (currentKey) { let zInc = 0.025; - if (currentKey.key == 'ArrowUp') { + if (currentKey.key == "ArrowUp") { eyePosition[2] += zInc; - } else if (currentKey.key == 'ArrowDown') { + } else if (currentKey.key == "ArrowDown") { eyePosition[2] -= zInc; } updateRenderParams(); @@ -250,6 +252,31 @@ function draw() { ], }); + let image_a = device.createTexture({ + format: "r32uint", + size: [16, 16], + usage: GPUTextureUsage.STORAGE_BINDING, + }); + let image_b = device.createTexture({ + format: "r32uint", + size: [16, 16], + usage: GPUTextureUsage.STORAGE_BINDING, + }); + + othergroup = device.createBindGroup({ + layout: computePipeline.getBindGroupLayout(1), + entries: [ + { + binding: 0, + resource: image_a.createView(), + }, + { + binding: 1, + resource: image_b.createView(), + }, + ], + }); + // Create the bind group for the compute shader. renderBindGroup = device.createBindGroup({ layout: renderPipeline.getBindGroupLayout(0), @@ -268,6 +295,7 @@ function draw() { const computePassEncoder = commandEncoder.beginComputePass(); computePassEncoder.setPipeline(computePipeline); computePassEncoder.setBindGroup(0, computeBindGroup); + computePassEncoder.setBindGroup(1, othergroup); computePassEncoder.dispatchWorkgroups(numBodies / workgroupSize); computePassEncoder.end(); @@ -282,7 +310,7 @@ function draw() { view: colorTextureView, loadOp: "clear", clearValue: { r: 0, g: 0, b: 0.1, a: 1 }, - storeOp: 'store' + storeOp: "store", }; const renderPassEncoder = commandEncoder.beginRenderPass({ colorAttachments: [colorAttachment], @@ -317,43 +345,43 @@ const reset = async () => { initPipelines(); paused = false; -} +}; function pause() { paused = !paused; - document.getElementById("pause").innerText = paused ? 'Unpause' : 'Pause'; + document.getElementById("pause").innerText = paused ? "Unpause" : "Pause"; } reset(); draw(); // Set up button onclick handlers. -document.querySelector('#reset').addEventListener('click', reset); -document.querySelector('#pause').addEventListener('click', pause); +document.querySelector("#reset").addEventListener("click", reset); +document.querySelector("#pause").addEventListener("click", pause); // Automatically reset when the number of bodies is changed. -document.querySelector('#numbodies').addEventListener('change', reset); +document.querySelector("#numbodies").addEventListener("change", reset); // Automatically reset when the power preference is changed. -document.querySelector('#powerpref').addEventListener('change', () => { +document.querySelector("#powerpref").addEventListener("change", () => { device = null; computePipeline = null; reset(); }); // Automatically rebuild the pipelines when the workgroup size is changed. -document.querySelector('#wgsize').addEventListener('change', initPipelines); +document.querySelector("#wgsize").addEventListener("change", initPipelines); // Add an event handler to update render parameters when the window is resized. -window.addEventListener('resize', updateRenderParams); +window.addEventListener("resize", updateRenderParams); // Handle key presses for user controls. -document.addEventListener('keydown', (e: KeyboardEvent) => { - if (e.key == ' ') { +document.addEventListener("keydown", (e: KeyboardEvent) => { + if (e.key == " ") { pause(); } currentKey = e; }); -document.addEventListener('keyup', (e: KeyboardEvent) => { +document.addEventListener("keyup", (e: KeyboardEvent) => { currentKey = null; }); diff --git a/src/shaders.wgsl b/src/shaders.wgsl index e6226a1..07f4331 100644 --- a/src/shaders.wgsl +++ b/src/shaders.wgsl @@ -21,6 +21,9 @@ fn computeForce(ipos : vec4, return coeff * d; } +@group(1) @binding(0) var image_a: texture_storage_2d; +@group(1) @binding(1) var image_b: texture_storage_2d; + @compute @workgroup_size(kWorkgroupSize) fn cs_main( @builtin(global_invocation_id) gid : vec3, @@ -39,6 +42,9 @@ fn cs_main( velocity = velocity + force * kDelta; velocities[idx] = velocity; + _ = image_a; + _ = image_b; + // Update position. positionsOut[idx] = pos + velocity * kDelta; }