From c6b6ba76a592ecab1a421fdc06ef015cd2c2ac5c Mon Sep 17 00:00:00 2001 From: Jackrabbit Date: Wed, 23 Nov 2022 04:53:24 -0500 Subject: [PATCH 1/3] finished compute pipeline, but the buffer seems not working --- src/gltf_group.ts | 4 +- src/gltf_renderer.ts | 86 ++++++++++++++++++++++++++++++++++++++++++- src/main.ts | 11 +++--- src/shaders/comp.wgsl | 16 ++++++++ 4 files changed, 110 insertions(+), 7 deletions(-) create mode 100644 src/shaders/comp.wgsl diff --git a/src/gltf_group.ts b/src/gltf_group.ts index ff29b3c..612687e 100644 --- a/src/gltf_group.ts +++ b/src/gltf_group.ts @@ -45,7 +45,9 @@ export default class GLTFGroup // Calculate matrix for each node locally // since gltf alone does not give this info directly this.nodeMatrics = new Map(); - const defaultTransform : mat4 = [1,0,0,0, 0,1,0,0, 0,0,1,0, 0,0,0,1]; + //const defaultTransform : mat4 = [1,0,0,0, 0,1,0,0, 0,0,1,0, 0,0,0,1]; + + let defaultTransform = mat4.identity(mat4.create()); for(const node of this.gltf.nodes) { this.#calcNodeMatrix(node, defaultTransform); diff --git a/src/gltf_renderer.ts b/src/gltf_renderer.ts index 0bb0998..b74c3f3 100644 --- a/src/gltf_renderer.ts +++ b/src/gltf_renderer.ts @@ -1,6 +1,7 @@ import vertShaderCode from './shaders/gltf.vert.wgsl'; import fragShaderCode from './shaders/gltf.frag.wgsl'; +import compShaderCode from './shaders/comp.wgsl'; import * as GLTFSpace from 'gltf-loader-ts/lib/gltf'; import {mat4, vec3} from 'gl-matrix'; @@ -77,6 +78,17 @@ export default class GltfRenderer commandEncoder: GPUCommandEncoder; passEncoder: GPURenderPassEncoder; + //ComputePipeline + computepassEncoder: GPUComputePassEncoder; + computePipeline: GPUComputePipeline; + computePipelineLayout: GPUPipelineLayout; + computeBindGroup: GPUBindGroup; + computeBindGroupLayout: GPUBindGroupLayout; + computeBuffer: GPUBuffer; + compShaderModule : GPUShaderModule; + + + // Web stuff canvas : HTMLCanvasElement; @@ -139,6 +151,26 @@ export default class GltfRenderer this.initFrameBindGroup(); this.initNodeBindGroup(); this.initInstanceBindGroup(); + this.initComputeBindGroup(); + + //create compute pipeline here, maybe not + this.computePipelineLayout = this.device.createPipelineLayout + ({ + label: 'glTF Compute Pipeline Layout', + bindGroupLayouts: [ + this.computeBindGroupLayout, + ] + + }); + const computeModule = this.getComputeShaderModule(); + this.computePipeline = this.device.createComputePipeline({ + layout: this.computePipelineLayout, + compute: { + module: computeModule, + entryPoint: 'simulate', + }, + }); + // Pipeline Layout this.gltfPipelineLayout = this.device.createPipelineLayout @@ -240,6 +272,34 @@ export default class GltfRenderer }); } + initComputeBindGroup() { + //a 4x4 transformation matrix + const computeBufferSize = 4 * 16; + this.computeBuffer = this.device.createBuffer ({ + size: this.gltf_group.instanceCount * computeBufferSize, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + }); + this.computeBindGroupLayout = this.device.createBindGroupLayout + ({ + label: `Compute BindGroupLayout`, + entries: + [{ + binding: 0, // transformation matrix + visibility: GPUShaderStage.COMPUTE, + buffer: {type: 'storage'}, + }], + }); + this.computeBindGroup = this.device.createBindGroup + ({ + label: `Compute BindGroup`, + layout: this.computeBindGroupLayout, + entries: + [{ + binding: 0, // transformation matrix + resource: { buffer: this.computeBuffer }, + }], + }); + } async loadGPUBuffers() { // TODO:: Create instanced bind group @@ -312,6 +372,16 @@ export default class GltfRenderer } return this.fragShaderModule; } + getComputeShaderModule() { + if (!this.compShaderModule) + { + this.compShaderModule = this.device.createShaderModule({ + label: 'glTF compute shader module', + code : compShaderCode + }); + } + return this.compShaderModule; + } setupMeshNodeBindGroup(node : GLTFSpace.Node) { @@ -456,7 +526,21 @@ export default class GltfRenderer }; this.commandEncoder = this.device.createCommandEncoder(); - + let transformationMatrixData = new Float32Array (this.gltf_group.transforms[0]).buffer; + console.log("before: "); + console.log(this.gltf_group.transforms[0]); + this.device.queue.writeBuffer(this.computeBuffer, 0, transformationMatrixData); + + //compute shader first + this.computepassEncoder = this.commandEncoder.beginComputePass(); + this.computepassEncoder.setPipeline(this.computePipeline); + this.computepassEncoder.setBindGroup(0, this.computeBindGroup); + this.computepassEncoder.dispatchWorkgroups(1); + + + this.computepassEncoder.end(); + console.log("after: "); + console.log(this.gltf_group.transforms[0]); // Render pass this.passEncoder = this.commandEncoder.beginRenderPass(renderPassDesc); diff --git a/src/main.ts b/src/main.ts index 1425255..0319c26 100644 --- a/src/main.ts +++ b/src/main.ts @@ -53,11 +53,12 @@ export default class Application // Butterfly let s : number = 0.01; this.gltf_butterfly = new GLTFGroup(); - await this.gltf_butterfly.init( - 'https://raw.githubusercontent.com/Li-Jia-Jun/WebGPU-Butterfly/gltf/models/butterfly/butterfly.gltf', - 3, - ["b1", "b2", "b3"], - [[s,0,0,0, 0,s,0,0, 0,0,s,0, 3,0,0,1], [s,0,0,0, 0,s,0,0, 0,0,s,0, -3,0,0,1], [s,0,0,0, 0,s,0,0, 0,0,s,0, 0,4.5,0,1]]); + // await this.gltf_butterfly.init( + // 'https://raw.githubusercontent.com/Li-Jia-Jun/WebGPU-Butterfly/gltf/models/butterfly/butterfly.gltf', + // 3, + // ["b1", "b2", "b3"], + // [[s,0,0,0, 0,s,0,0, 0,0,s,0, 3,0,0,1], [s,0,0,0, 0,s,0,0, 0,0,s,0, -3,0,0,1], [s,0,0,0, 0,s,0,0, 0,0,s,0, 0,4.5,0,1]]); + await this.gltf_butterfly.init('https://raw.githubusercontent.com/Li-Jia-Jun/WebGPU-Butterfly/gltf/models/butterfly/butterfly.gltf',1, ["b1"],[[s,0,0,0, 0,s,0,0, 0,0,s,0, 0,4.5,0,1] ]); this.renderer_butterfly = new GltfRenderer(); await this.renderer_butterfly.init(this.adapter, this.device, this.queue, this.canvas, this.gltf_butterfly); diff --git a/src/shaders/comp.wgsl b/src/shaders/comp.wgsl new file mode 100644 index 0000000..4c134be --- /dev/null +++ b/src/shaders/comp.wgsl @@ -0,0 +1,16 @@ +// struct Transforms { +// transformationMatrix : mat4x4, +// } + + +@binding(0) @group(0) var transform : mat4x4; + + +@compute @workgroup_size(1) +fn simulate( + @builtin(global_invocation_id) GlobalInvocationID : vec3 +) { + transform[0][3] -= 100; + // Store the new particle value + //transform.transformationMatrix = ; +} \ No newline at end of file From 0ec32346872c7138c24e164e8731932f7ecf6ff4 Mon Sep 17 00:00:00 2001 From: Jackrabbit Date: Wed, 23 Nov 2022 22:43:20 -0500 Subject: [PATCH 2/3] 9z --- src/gltf_renderer.ts | 16 ++++++++++++---- src/shaders/comp.wgsl | 13 ++++++------- 2 files changed, 18 insertions(+), 11 deletions(-) diff --git a/src/gltf_renderer.ts b/src/gltf_renderer.ts index b74c3f3..b755e8a 100644 --- a/src/gltf_renderer.ts +++ b/src/gltf_renderer.ts @@ -296,7 +296,7 @@ export default class GltfRenderer entries: [{ binding: 0, // transformation matrix - resource: { buffer: this.computeBuffer }, + resource: { buffer: this.instanceBuffer }, }], }); } @@ -526,11 +526,18 @@ export default class GltfRenderer }; this.commandEncoder = this.device.createCommandEncoder(); + + let transformationMatrixData = new Float32Array (this.gltf_group.transforms[0]).buffer; - console.log("before: "); - console.log(this.gltf_group.transforms[0]); + + this.device.queue.writeBuffer(this.computeBuffer, 0, transformationMatrixData); + // console.log("before: "); + // let tmpBufferArray = new Uint8Array(this.computeBuffer.getMappedRange()); + // console.log(tmpBufferArray); + // this.instanceBuffer + //compute shader first this.computepassEncoder = this.commandEncoder.beginComputePass(); this.computepassEncoder.setPipeline(this.computePipeline); @@ -540,7 +547,8 @@ export default class GltfRenderer this.computepassEncoder.end(); console.log("after: "); - console.log(this.gltf_group.transforms[0]); + // let tmpBufferArray = new Uint8Array(this.computeBuffer.getMappedRange()); + // console.log(tmpBufferArray); // Render pass this.passEncoder = this.commandEncoder.beginRenderPass(renderPassDesc); diff --git a/src/shaders/comp.wgsl b/src/shaders/comp.wgsl index 4c134be..27f8e77 100644 --- a/src/shaders/comp.wgsl +++ b/src/shaders/comp.wgsl @@ -1,16 +1,15 @@ -// struct Transforms { -// transformationMatrix : mat4x4, -// } - -@binding(0) @group(0) var transform : mat4x4; +@binding(0) @group(0) var transform : array>; @compute @workgroup_size(1) fn simulate( @builtin(global_invocation_id) GlobalInvocationID : vec3 ) { - transform[0][3] -= 100; + var m = transform[0]; + m[0][3] += 5; + transform[0] = m; + } + // Store the new particle value //transform.transformationMatrix = ; -} \ No newline at end of file From 02eca0cd1869311c6a87b5da63453ba275df68c8 Mon Sep 17 00:00:00 2001 From: Jackrabbit Date: Sat, 26 Nov 2022 22:42:27 -0500 Subject: [PATCH 3/3] update --- src/gltf_renderer.ts | 37 ++++++++++++++++++++----------------- src/shaders/comp.wgsl | 11 ++++++++++- 2 files changed, 30 insertions(+), 18 deletions(-) diff --git a/src/gltf_renderer.ts b/src/gltf_renderer.ts index b755e8a..56805c0 100644 --- a/src/gltf_renderer.ts +++ b/src/gltf_renderer.ts @@ -7,7 +7,8 @@ import * as GLTFSpace from 'gltf-loader-ts/lib/gltf'; import {mat4, vec3} from 'gl-matrix'; import GLTFGroup from './gltf_group'; - +var frame = 0; +var tmp; // Make sure the shaders follow this mapping const ShaderLocations = { @@ -249,7 +250,7 @@ export default class GltfRenderer this.instanceBuffer = this.device.createBuffer ({ size: 16 * this.gltf_group.instanceCount * Float32Array.BYTES_PER_ELEMENT, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, }); this.instanceBindGroupLayout = this.device.createBindGroupLayout ({ @@ -275,10 +276,7 @@ export default class GltfRenderer initComputeBindGroup() { //a 4x4 transformation matrix const computeBufferSize = 4 * 16; - this.computeBuffer = this.device.createBuffer ({ - size: this.gltf_group.instanceCount * computeBufferSize, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST - }); + this.computeBuffer = this.instanceBuffer; this.computeBindGroupLayout = this.device.createBindGroupLayout ({ label: `Compute BindGroupLayout`, @@ -296,7 +294,7 @@ export default class GltfRenderer entries: [{ binding: 0, // transformation matrix - resource: { buffer: this.instanceBuffer }, + resource: { buffer: this.computeBuffer }, }], }); } @@ -527,12 +525,12 @@ export default class GltfRenderer this.commandEncoder = this.device.createCommandEncoder(); - - let transformationMatrixData = new Float32Array (this.gltf_group.transforms[0]).buffer; - - - this.device.queue.writeBuffer(this.computeBuffer, 0, transformationMatrixData); + //let transformationMatrixData = new Float32Array (this.gltf_group.transforms[0]).buffer; + //let gpuBufferArray = new Uint8Array(this.computeBuffer.getMappedRange()); + //const tablesArray = new Float32Array(this.computeBuffer.getMappedRange()); + // this.device.queue.writeBuffer(this.computeBuffer, 0, transformationMatrixData); + // console.log("before: "); // let tmpBufferArray = new Uint8Array(this.computeBuffer.getMappedRange()); // console.log(tmpBufferArray); @@ -546,15 +544,19 @@ export default class GltfRenderer this.computepassEncoder.end(); - console.log("after: "); - // let tmpBufferArray = new Uint8Array(this.computeBuffer.getMappedRange()); - // console.log(tmpBufferArray); + + if(tmp != this.computeBuffer) { + + } + + // Render pass this.passEncoder = this.commandEncoder.beginRenderPass(renderPassDesc); this.passEncoder.setBindGroup(0, this.frameBindGroup); this.passEncoder.setBindGroup(2, this.instanceBindGroup); + //+5 // Bind gltf data to render pass for (const [node, bindGroup] of this.nodeGpuData) { @@ -606,7 +608,9 @@ export default class GltfRenderer // Submit command queue this.queue.submit([this.commandEncoder.finish()]); - requestAnimationFrame(this.renderGLTF); + frame++; + requestAnimationFrame(this.renderGLTF); + } updateFrameBuffer(projMat : mat4, viewMat : mat4, pos : vec3, time : number) @@ -635,7 +639,6 @@ export default class GltfRenderer let arr = new Float32Array(instanceArrayBuffer, st, 16); arr.set(mat); } - this.device.queue.writeBuffer(this.instanceBuffer, 0, instanceArrayBuffer); } } diff --git a/src/shaders/comp.wgsl b/src/shaders/comp.wgsl index 27f8e77..f7257c3 100644 --- a/src/shaders/comp.wgsl +++ b/src/shaders/comp.wgsl @@ -7,9 +7,18 @@ fn simulate( @builtin(global_invocation_id) GlobalInvocationID : vec3 ) { var m = transform[0]; - m[0][3] += 5; + + //column major + + + + + + m[3][0] = m[3][0] + 0.001; transform[0] = m; } + + // Store the new particle value //transform.transformationMatrix = ;