diff --git a/.gitignore b/.gitignore
new file mode 100644
index 0000000..818ff30
--- /dev/null
+++ b/.gitignore
@@ -0,0 +1,2 @@
+Gemfile.lock
+_site/
\ No newline at end of file
diff --git a/DevelopmentBuild.bat b/DevelopmentBuild.bat
new file mode 100644
index 0000000..b67b3c4
--- /dev/null
+++ b/DevelopmentBuild.bat
@@ -0,0 +1 @@
+bundle install
diff --git a/DevelopmentServe.bat b/DevelopmentServe.bat
new file mode 100644
index 0000000..2f58498
--- /dev/null
+++ b/DevelopmentServe.bat
@@ -0,0 +1 @@
+jekyll serve --livereload
\ No newline at end of file
diff --git a/Gemfile b/Gemfile
new file mode 100644
index 0000000..7d234ee
--- /dev/null
+++ b/Gemfile
@@ -0,0 +1,5 @@
+source "https://rubygems.org"
+
+gem "jekyll"
+
+gem 'wdm', '>= 0.1.0' if Gem.win_platform?
\ No newline at end of file
diff --git a/LICENSE.md b/LICENSE.md
new file mode 100644
index 0000000..e3e3c6f
--- /dev/null
+++ b/LICENSE.md
@@ -0,0 +1,9 @@
+Copyright (c) 2024 Electronic Arts Inc. All rights reserved.
+
+Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:
+
+1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.
+2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.
+3. Neither the name of Electronic Arts, Inc. ("EA") nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
diff --git a/README.md b/README.md
new file mode 100644
index 0000000..9ba968b
--- /dev/null
+++ b/README.md
@@ -0,0 +1,42 @@
+
+# PB-MPM
+[](LICENSE.md)
+[](VERSION.md)
+
+This package is an accessible WebGPU implementation of Position Based MPM (PB-MPM).
+
+
+
+
+
+
+## Building
+
+PB-MPM is based on WebGPU. Currently, this restricts its use to the Chrome and Edge browsers.
+
+1. [Install Ruby](https://www.ruby-lang.org/en/documentation/installation/)
+2. Run `bundle install`
+3. Run `jekyll serve --livereload`
+4. Open a browser window at [`localhost:4000`](http://localhost:4000)
+
+## Reference
+
+> *Chris Lewin*. **[A Position Based Material Point Method](https://seed.ea.com)**. ACM SIGGRAPH 2024.
+
+## Authors
+
+
+Search for Extraordinary Experiences Division (SEED) - Electronic Arts https://seed.ea.com
+We are a cross-disciplinary team within EA Worldwide Studios.
+Our mission is to explore, build and help define the future of interactive entertainment.
+
+## Contributing
+
+Before you can contribute, EA must have a Contributor License Agreement (CLA) on file that has been signed by each contributor.
+You can sign here: http://bit.ly/electronic-arts-cla
+
+## Research Resources
+- [incremental_mpm](https://github.com/nialltl/incremental_mpm)
+
+## License
+- The source code is released under a *BSD 3-Clause License* as detailed in [LICENSE.md](LICENSE.md)
diff --git a/VERSION.md b/VERSION.md
new file mode 100644
index 0000000..8c76a01
--- /dev/null
+++ b/VERSION.md
@@ -0,0 +1,4 @@
+# Version History
+
+## Version 1.0.0 (August 2024)
+- Initial release following SIGGRAPH 2024 talk "A Position Based Material Point Method" by Christopher Lewin (SEED, Electronic Arts)
\ No newline at end of file
diff --git a/_config.yml b/_config.yml
new file mode 100644
index 0000000..10a9b68
--- /dev/null
+++ b/_config.yml
@@ -0,0 +1,3 @@
+exclude:
+ - DevelopmentBuild.bat
+ - DevelopmentServe.bat
\ No newline at end of file
diff --git a/data/SEED.jpg b/data/SEED.jpg
new file mode 100644
index 0000000..15703e7
Binary files /dev/null and b/data/SEED.jpg differ
diff --git a/data/SEED_inverted.png b/data/SEED_inverted.png
new file mode 100644
index 0000000..6ac3e06
Binary files /dev/null and b/data/SEED_inverted.png differ
diff --git a/data/blockCrusher.gif b/data/blockCrusher.gif
new file mode 100644
index 0000000..a491f2f
Binary files /dev/null and b/data/blockCrusher.gif differ
diff --git a/data/coiling.gif b/data/coiling.gif
new file mode 100644
index 0000000..5c7012d
Binary files /dev/null and b/data/coiling.gif differ
diff --git a/data/colliders.gif b/data/colliders.gif
new file mode 100644
index 0000000..b6ef069
Binary files /dev/null and b/data/colliders.gif differ
diff --git a/data/rotate.svg b/data/rotate.svg
new file mode 100644
index 0000000..bb542f9
--- /dev/null
+++ b/data/rotate.svg
@@ -0,0 +1,122 @@
+
+
+
+
diff --git a/data/splashing.gif b/data/splashing.gif
new file mode 100644
index 0000000..c720b89
Binary files /dev/null and b/data/splashing.gif differ
diff --git a/index.html b/index.html
new file mode 100644
index 0000000..510c1cb
--- /dev/null
+++ b/index.html
@@ -0,0 +1,54 @@
+
+
+
+ Home
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
Loading...
+
+
+
+
\ No newline at end of file
diff --git a/scenes/4Materials.json b/scenes/4Materials.json
new file mode 100644
index 0000000..2be6ef0
--- /dev/null
+++ b/scenes/4Materials.json
@@ -0,0 +1,77 @@
+{
+ "version": 1,
+ "resolution": [
+ 1920,
+ 953
+ ],
+ "shapes": [
+ {
+ "id": "shape1",
+ "position": {
+ "x": 291.8062500000001,
+ "y": 239
+ },
+ "halfSize": {
+ "x": "195",
+ "y": "71"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "0",
+ "emitMaterial": "0",
+ "emissionRate": "0.5",
+ "radius": 20
+ },
+ {
+ "id": "shape0",
+ "position": {
+ "x": 1231.80625,
+ "y": 751
+ },
+ "halfSize": {
+ "x": "195",
+ "y": "186"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape2",
+ "position": {
+ "x": 1668.80625,
+ "y": 702
+ },
+ "halfSize": {
+ "x": "195",
+ "y": "186"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "2",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape3",
+ "position": {
+ "x": 785.8062500000001,
+ "y": 723
+ },
+ "halfSize": {
+ "x": "195",
+ "y": "186"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ }
+ ]
+}
\ No newline at end of file
diff --git a/scenes/damBreak.json b/scenes/damBreak.json
new file mode 100644
index 0000000..1fcf95a
--- /dev/null
+++ b/scenes/damBreak.json
@@ -0,0 +1,128 @@
+{
+ "version": 1,
+ "resolution": [
+ 1530,
+ 863
+ ],
+ "shapes": [
+ {
+ "id": "shape1",
+ "position": {
+ "x": 383.90990990990986,
+ "y": 429.86149003147955
+ },
+ "halfSize": {
+ "x": "354",
+ "y": "407"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "0",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ },
+ {
+ "id": "shape2",
+ "position": {
+ "x": 962.9099099099099,
+ "y": 525.8614900314795
+ },
+ "halfSize": {
+ "x": "80",
+ "y": "80"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ },
+ {
+ "id": "shape3",
+ "position": {
+ "x": 1165.9099099099099,
+ "y": 523.8614900314795
+ },
+ "halfSize": {
+ "x": "80",
+ "y": "80"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ },
+ {
+ "id": "shape4",
+ "position": {
+ "x": 1364.9099099099099,
+ "y": 529.8614900314795
+ },
+ "halfSize": {
+ "x": "80",
+ "y": "80"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ },
+ {
+ "id": "shape5",
+ "position": {
+ "x": 1061.9099099099099,
+ "y": 348.86149003147955
+ },
+ "halfSize": {
+ "x": "80",
+ "y": "80"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ },
+ {
+ "id": "shape6",
+ "position": {
+ "x": 1253.9099099099099,
+ "y": 346.86149003147955
+ },
+ "halfSize": {
+ "x": "80",
+ "y": "80"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ },
+ {
+ "id": "shape7",
+ "position": {
+ "x": 1157.9099099099099,
+ "y": 167.86149003147955
+ },
+ "halfSize": {
+ "x": "80",
+ "y": "80"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "6.2",
+ "radius": 20.174784316989175
+ }
+ ]
+}
\ No newline at end of file
diff --git a/scenes/plasticPress.json b/scenes/plasticPress.json
new file mode 100644
index 0000000..82fd0ae
--- /dev/null
+++ b/scenes/plasticPress.json
@@ -0,0 +1,111 @@
+{
+ "version": 1,
+ "resolution": [
+ 1920,
+ 1080
+ ],
+ "shapes": [
+ {
+ "id": "shape1",
+ "position": {
+ "x": -414.2125000000001,
+ "y": 457
+ },
+ "halfSize": {
+ "x": "814",
+ "y": "575"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "1",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape0",
+ "position": {
+ "x": 1575.7875,
+ "y": 515
+ },
+ "halfSize": {
+ "x": "327",
+ "y": "584"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape2",
+ "position": {
+ "x": 1219.7875,
+ "y": 208
+ },
+ "halfSize": {
+ "x": "9",
+ "y": "62"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "1",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape3",
+ "position": {
+ "x": 1218.7875,
+ "y": -31
+ },
+ "halfSize": {
+ "x": "9",
+ "y": "62"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "1",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape4",
+ "position": {
+ "x": 1217.7875,
+ "y": 92
+ },
+ "halfSize": {
+ "x": "9",
+ "y": "35"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "1",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ },
+ {
+ "id": "shape5",
+ "position": {
+ "x": 778.7874999999999,
+ "y": 609
+ },
+ "halfSize": {
+ "x": "442",
+ "y": "425"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "1",
+ "emitMaterial": "3",
+ "emissionRate": "5.3",
+ "radius": 20
+ }
+ ]
+}
\ No newline at end of file
diff --git a/scenes/pyramid.json b/scenes/pyramid.json
new file mode 100644
index 0000000..1fcb7a4
--- /dev/null
+++ b/scenes/pyramid.json
@@ -0,0 +1,111 @@
+{
+ "version": 1,
+ "resolution": [
+ 809,
+ 818
+ ],
+ "shapes": [
+ {
+ "id": "shape6",
+ "position": {
+ "x": 429.95377929687504,
+ "y": 181.1437565582371
+ },
+ "halfSize": {
+ "x": "97",
+ "y": "88"
+ },
+ "rotation": -90.03706148292488,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "4.4",
+ "radius": 12.02773479659328
+ },
+ {
+ "id": "shape0",
+ "position": {
+ "x": 428.95377929687504,
+ "y": 680.1437565582371
+ },
+ "halfSize": {
+ "x": "97",
+ "y": "88"
+ },
+ "rotation": -90.03706148292488,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "4.4",
+ "radius": 12.02773479659328
+ },
+ {
+ "id": "shape1",
+ "position": {
+ "x": 672.953779296875,
+ "y": 689.1437565582371
+ },
+ "halfSize": {
+ "x": "97",
+ "y": "88"
+ },
+ "rotation": -90.03706148292488,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "4.4",
+ "radius": 12.02773479659328
+ },
+ {
+ "id": "shape2",
+ "position": {
+ "x": 202.95377929687504,
+ "y": 686.1437565582371
+ },
+ "halfSize": {
+ "x": "97",
+ "y": "88"
+ },
+ "rotation": -90.03706148292488,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "4.4",
+ "radius": 12.02773479659328
+ },
+ {
+ "id": "shape3",
+ "position": {
+ "x": 301.95377929687504,
+ "y": 431.1437565582371
+ },
+ "halfSize": {
+ "x": "97",
+ "y": "88"
+ },
+ "rotation": -90.03706148292488,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "4.4",
+ "radius": 12.02773479659328
+ },
+ {
+ "id": "shape4",
+ "position": {
+ "x": 576.953779296875,
+ "y": 435.1437565582371
+ },
+ "halfSize": {
+ "x": "97",
+ "y": "88"
+ },
+ "rotation": -90.03706148292488,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "1",
+ "emissionRate": "4.4",
+ "radius": 12.02773479659328
+ }
+ ]
+}
\ No newline at end of file
diff --git a/scenes/sandPile.json b/scenes/sandPile.json
new file mode 100644
index 0000000..27e0040
--- /dev/null
+++ b/scenes/sandPile.json
@@ -0,0 +1,26 @@
+{
+ "version": 1,
+ "resolution": [
+ 1920,
+ 953
+ ],
+ "shapes": [
+ {
+ "id": "shape1",
+ "position": {
+ "x": 955.3963963963963,
+ "y": 563
+ },
+ "halfSize": {
+ "x": "319",
+ "y": "283"
+ },
+ "rotation": 0,
+ "shape": "0",
+ "function": "3",
+ "emitMaterial": "2",
+ "emissionRate": "5.3",
+ "radius": 24.01200900750657
+ }
+ ]
+}
\ No newline at end of file
diff --git a/shaders/dispatch.inc.wgsl b/shaders/dispatch.inc.wgsl
new file mode 100644
index 0000000..c0ac63e
--- /dev/null
+++ b/shaders/dispatch.inc.wgsl
@@ -0,0 +1,5 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!insert DispatchSizes
diff --git a/shaders/gridToParticle.wgsl b/shaders/gridToParticle.wgsl
new file mode 100644
index 0000000..859e0e5
--- /dev/null
+++ b/shaders/gridToParticle.wgsl
@@ -0,0 +1,88 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include particle.inc
+//!include simConstants.inc
+//!include matrix.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_particleCount : array;
+@group(0) @binding(2) var g_particles : array;
+@group(0) @binding(3) var g_grid : array;
+
+@compute @workgroup_size(ParticleDispatchSize, 1, 1)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ if(id.x >= g_particleCount[0])
+ {
+ return;
+ }
+
+ var particle = g_particles[id.x];
+
+ if(particle.enabled == 0)
+ {
+ return;
+ }
+
+ var p = particle.position;
+
+ let weightInfo = quadraticWeightInit(p);
+
+ var B = ZeroMatrix;
+ var d = vec2f(0);
+ var volume = 0.0;
+ // Iterate over local 3x3 neigbourhood
+ for(var i = 0; i < 3; i++)
+ {
+ for(var j = 0; j < 3; j++)
+ {
+ // Weight corresponding to this neighbourhood cell
+ let weight = weightInfo.weights[i].x * weightInfo.weights[j].y;
+
+ // 2d index of this cell in the grid
+ let neighbourCellIndex = vec2u(vec2i(weightInfo.cellIndex) + vec2i(i,j));
+
+ // Linear index in the buffer
+ let gridVertexIdx = gridVertexIndex(neighbourCellIndex, g_simConstants.gridSize);
+
+ let weightedDisplacement = weight * vec2f(
+ decodeFixedPoint(g_grid[gridVertexIdx + 0], g_simConstants.fixedPointMultiplier),
+ decodeFixedPoint(g_grid[gridVertexIdx + 1], g_simConstants.fixedPointMultiplier)
+ );
+
+ let offset = vec2f(neighbourCellIndex) - p + 0.5;
+ B += outerProduct(weightedDisplacement, offset);
+ d += weightedDisplacement;
+
+ // This is only required if we are going to mix in the grid volume to the liquid volume
+ if(g_simConstants.useGridVolumeForLiquid != 0)
+ {
+ volume += weight * decodeFixedPoint(g_grid[gridVertexIdx + 3], g_simConstants.fixedPointMultiplier);
+ }
+ }
+ }
+
+ particle.deformationDisplacement = B * 4.0;
+ particle.displacement = d;
+
+ // Using standard MPM volume integration for liquids can lead to slow volume loss over time
+ // especially when particles are undergoing a lot of shearing motion.
+ // We can recover an objective measure of volume from the grid directly.
+ // Here we mix it in to the integrated volume, but only if the liquid is compressed.
+ // This is because the behaviour in tension of the grid volume and mpm volume is quite different.
+ // Note this runs every iteration in the PBMPM solver but we only really require it to happen occasionally
+ // because the MPM integration doesn't lose volume very fast.
+ if(g_simConstants.useGridVolumeForLiquid != 0)
+ {
+ volume = 1.0/volume;
+ if(volume < 1)
+ {
+ particle.liquidDensity = mix(particle.liquidDensity, volume, 0.1);
+ }
+ }
+
+ g_particles[id.x] = particle;
+}
\ No newline at end of file
diff --git a/shaders/gridUpdate.wgsl b/shaders/gridUpdate.wgsl
new file mode 100644
index 0000000..6a53977
--- /dev/null
+++ b/shaders/gridUpdate.wgsl
@@ -0,0 +1,94 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include simConstants.inc
+//!include particle.inc
+//!include shapes.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_grid : array;
+@group(0) @binding(2) var g_shapes : array;
+
+@compute @workgroup_size(GridDispatchSize, GridDispatchSize)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ if(any(id.xy >= g_simConstants.gridSize))
+ {
+ return;
+ }
+
+ let gridVertexAddress = gridVertexIndex(id.xy, g_simConstants.gridSize);
+
+ // Load from grid
+ var dx = decodeFixedPoint(g_grid[gridVertexAddress + 0], g_simConstants.fixedPointMultiplier);
+ var dy = decodeFixedPoint(g_grid[gridVertexAddress + 1], g_simConstants.fixedPointMultiplier);
+ var w = decodeFixedPoint(g_grid[gridVertexAddress + 2], g_simConstants.fixedPointMultiplier);
+
+ if(w < 1e-5f)
+ {
+ dx = 0;
+ dy = 0;
+ }
+
+ // Perform mass weighting to get grid displacement
+ dx = dx / w;
+ dy = dy / w;
+
+ var gridDisplacement = vec2f(dx, dy);
+
+
+ // Collision detection against collider shapes
+ for(var shapeIndex = 0u; shapeIndex < g_simConstants.shapeCount; shapeIndex++)
+ {
+ let shape = g_shapes[shapeIndex];
+
+ if(shape.functionality != ShapeFunctionCollider)
+ {
+ continue;
+ }
+
+ let gridPosition = vec2f(id.xy);
+ let displacedGridPosition = gridPosition + gridDisplacement;
+
+ let collideResult = collide(shape, displacedGridPosition);
+
+ if(collideResult.collides)
+ {
+ let gap = min(0, dot(collideResult.normal, collideResult.pointOnCollider - gridPosition));
+ let penetration = dot(collideResult.normal, gridDisplacement) - gap;
+
+ // Prevent any further penetration in radial direction
+ let radialImpulse = max(penetration, 0);
+ gridDisplacement -= radialImpulse*collideResult.normal;
+ }
+ }
+
+ // Collision detection against guardian shape
+
+ // Grid vertices near or inside the guardian region should have their displacenent values
+ // corrected in order to prevent particles moving into the guardian.
+ // We do this by finding whether a grid vertex would be inside the guardian region after displacement
+ // with the current velocity and, if it is, setting the displacement so that no further penetration can occur.
+ let gridPosition = vec2f(id.xy);
+ let displacedGridPosition = gridPosition + gridDisplacement;
+ let projectedGridPosition = projectInsideGuardian(displacedGridPosition, g_simConstants.gridSize, GuardianSize+1);
+ let projectedDifference = projectedGridPosition - displacedGridPosition;
+
+ if(projectedDifference.x != 0)
+ {
+ gridDisplacement.x = 0;
+ gridDisplacement.y = 0;
+ }
+
+ if(projectedDifference.y != 0)
+ {
+ gridDisplacement.x = 0;
+ gridDisplacement.y = 0;
+ }
+
+ // Save back to grid
+ g_grid[gridVertexAddress + 0] = encodeFixedPoint(gridDisplacement.x , g_simConstants.fixedPointMultiplier);
+ g_grid[gridVertexAddress + 1] = encodeFixedPoint(gridDisplacement.y , g_simConstants.fixedPointMultiplier);
+}
\ No newline at end of file
diff --git a/shaders/gridZero.wgsl b/shaders/gridZero.wgsl
new file mode 100644
index 0000000..62d8fd3
--- /dev/null
+++ b/shaders/gridZero.wgsl
@@ -0,0 +1,25 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include simConstants.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_grid : array;
+
+@compute @workgroup_size(GridDispatchSize, GridDispatchSize)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ if(any(id.xy >= g_simConstants.gridSize))
+ {
+ return;
+ }
+
+ let gridVertexAddress = gridVertexIndex(id.xy, g_simConstants.gridSize);
+
+ g_grid[gridVertexAddress + 0] = 0;
+ g_grid[gridVertexAddress + 1] = 0;
+ g_grid[gridVertexAddress + 2] = 0;
+ g_grid[gridVertexAddress + 3] = 0;
+}
\ No newline at end of file
diff --git a/shaders/matrix.inc.wgsl b/shaders/matrix.inc.wgsl
new file mode 100644
index 0000000..c25f290
--- /dev/null
+++ b/shaders/matrix.inc.wgsl
@@ -0,0 +1,79 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+struct SVDResult
+{
+ U: mat2x2f,
+ Sigma: vec2f,
+ Vt: mat2x2f,
+};
+
+fn svd(m: mat2x2f) -> SVDResult
+{
+ // Pedro Gimeno (https://scicomp.stackexchange.com/users/9673/pedro-gimeno),
+ // Robust algorithm for 2x2 SVD, URL (version: 2019-10-22): https://scicomp.stackexchange.com/q/14103
+ let E = (m[0][0] + m[1][1])*0.5;
+ let F = (m[0][0] - m[1][1])*0.5;
+ let G = (m[0][1] + m[1][0])*0.5;
+ let H = (m[0][1] - m[1][0])*0.5;
+
+ let Q = sqrt(E*E + H*H);
+ let R = sqrt(F*F + G*G);
+ let sx = Q + R;
+ let sy = Q - R;
+
+ let a1 = atan2(G, F);
+ let a2 = atan2(H, E);
+
+ let theta = (a2 - a1)*0.5;
+ let phi = (a2 + a1)*0.5;
+
+ let U = rot(phi);
+ let Sigma = vec2f(sx, sy);
+ let Vt = rot(theta);
+
+ return SVDResult(U, Sigma, Vt);
+}
+
+fn det(m: mat2x2f) -> f32
+{
+ return m[0][0]*m[1][1] - m[0][1]*m[1][0];
+}
+
+fn tr(m: mat2x2f) -> f32
+{
+ return m[0][0] + m[1][1];
+}
+
+fn rot(theta: f32) -> mat2x2f
+{
+ let ct = cos(theta);
+ let st = sin(theta);
+
+ return mat2x2f(ct, st, -st, ct);
+}
+
+fn inverse(m: mat2x2f) -> mat2x2f
+{
+ // This matrix is guaranteed to be numerically invertible
+ // because its singular values have been clamped in the integration stage
+ let a = m[0][0];
+ let b = m[1][0];
+ let c = m[0][1];
+ let d = m[1][1];
+ return (1.0 / det(m))*mat2x2f(d, -c, -b, a);
+}
+
+fn outerProduct(x: vec2f, y: vec2f) -> mat2x2f
+{
+ return mat2x2f(x*y.x, x*y.y);
+}
+
+fn diag(d: vec2f) -> mat2x2f
+{
+ return mat2x2f(d.x, 0, 0, d.y);
+}
+
+const Identity = mat2x2f(1,0,0,1);
+const ZeroMatrix = mat2x2f(0,0,0,0);
\ No newline at end of file
diff --git a/shaders/particle.inc.wgsl b/shaders/particle.inc.wgsl
new file mode 100644
index 0000000..7156cb2
--- /dev/null
+++ b/shaders/particle.inc.wgsl
@@ -0,0 +1,95 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+// Must be kept in sync with definition in gpu.js
+struct Particle
+{
+ position : vec2f,
+ displacement : vec2f,
+ deformationGradient : mat2x2f,
+ deformationDisplacement : mat2x2f,
+
+ liquidDensity : f32,
+ mass : f32,
+ material : f32,
+ volume: f32,
+
+ lambda: f32,
+ logJp : f32,
+ color: vec3f,
+ enabled: f32,
+};
+
+// For safety, we keep particles `guardianSize` cells away from the outside of the domain.
+// To implement this we clamp the grid values to ensure they do not contribute towards moving
+// particles outside the domain, and additionally clamp particle positions whenever they are moved.
+fn projectInsideGuardian(p : vec2f, gridSize : vec2u, guardianSize : f32) -> vec2f
+{
+ let clampMin = vec2f(guardianSize);
+ let clampMax = vec2f(gridSize) - vec2f(guardianSize,guardianSize) - vec2f(1,1);
+
+ return clamp(p, vec2f(clampMin), vec2f(clampMax));
+}
+
+fn insideGuardian(id: vec2u, gridSize: vec2u, guardianSize: u32) -> bool
+{
+ if(id.x <= guardianSize) {return false;}
+ if(id.x >= (gridSize.x-guardianSize-1)) {return false;}
+ if(id.y <= guardianSize) {return false;}
+ if(id.y >= gridSize.y-guardianSize-1) {return false;}
+
+ return true;
+}
+
+struct QuadraticWeightInfo
+{
+ weights: array,
+ cellIndex: vec2f,
+}
+
+fn pow2(x: vec2f) -> vec2f
+{
+ return x*x;
+}
+
+fn quadraticWeightInit(position: vec2f) -> QuadraticWeightInfo
+{
+ let roundDownPosition = floor(position);
+ let offset = (position - roundDownPosition) - 0.5;
+ return QuadraticWeightInfo(
+ array(
+ 0.5 * pow2(0.5 - offset),
+ 0.75 - pow2(offset),
+ 0.5 * pow2(0.5 + offset)
+ ),
+ roundDownPosition - 1,
+ );
+}
+
+fn pow3(x: vec2f) -> vec2f
+{
+ return x*x*x;
+}
+
+struct CubicWeightInfo
+{
+ weights: array,
+ cellIndex: vec2f
+};
+
+fn cubicWeightInit(position: vec2f) -> CubicWeightInfo
+{
+ let roundDownPosition = floor(position);
+ let offset = position - roundDownPosition;
+
+ return CubicWeightInfo(
+ array(
+ pow3(2.0 - (1+offset))/6.0,
+ 0.5*pow3(offset) - pow2(offset) + 2.0/3.0,
+ 0.5*pow3(1 - offset) - pow2(1 - offset) + 2.0/3.0,
+ pow3(2.0 - (2 - offset))/6.0,
+ ),
+ roundDownPosition - 1
+ );
+}
diff --git a/shaders/particleEmit.wgsl b/shaders/particleEmit.wgsl
new file mode 100644
index 0000000..e82fb15
--- /dev/null
+++ b/shaders/particleEmit.wgsl
@@ -0,0 +1,139 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include simConstants.inc
+//!include particle.inc
+//!include matrix.inc
+//!include shapes.inc
+//!include random.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_particleCount : array>;
+@group(0) @binding(2) var g_particles : array;
+@group(0) @binding(3) var g_shapes : array;
+@group(0) @binding(4) var g_freeCount : array>;
+@group(0) @binding(5) var g_freeIndices : array;
+
+fn createParticle(position: vec2f, material: f32, mass: f32, volume: f32, color: vec3f) -> Particle
+{
+ return Particle(
+ position,
+ vec2f(0,0),
+ Identity,
+ ZeroMatrix,
+ 1.0,
+ mass,
+ material,
+ volume,
+ 0.0,
+ 1.0,
+ color,
+ 1.0
+ );
+}
+
+fn addParticle(position: vec2f, material: f32, volume: f32, density: f32, jitterScale: f32)
+{
+ var particleIndex = 0u;
+ // First check the free list to see if we can reuse a particle slot
+ let freeIndexSlot = atomicSub(&g_freeCount[0], 1i) - 1i;
+ if(freeIndexSlot >= 0)
+ {
+ particleIndex = g_freeIndices[u32(freeIndexSlot)];
+ }
+ else // If free list is empty then grow the particle count
+ {
+ particleIndex = atomicAdd(&g_particleCount[0], 1);
+ }
+
+ var color = vec3f(1,1,1);
+
+ if(material == MaterialLiquid)
+ {
+ color = vec3f(0.2,0.2,1);
+ }
+ else if(material == MaterialElastic)
+ {
+ color = vec3f(0.2,1,0.2);
+ }
+ else if(material == MaterialSand)
+ {
+ color = vec3f(1,1,0.2);
+ }
+ else if(material == MaterialVisco)
+ {
+ color = vec3f(1, 0.5, 1);
+ }
+
+ let jitterX = hash(particleIndex);
+ let jitterY = hash(u32(position.x * position.y * 100));
+
+ let jitter = vec2f(-0.25, -0.25) + 0.5*vec2f(f32(jitterX % 10) / 10, f32(jitterY % 10) / 10);
+
+ var newParticle = createParticle(
+ position + jitter*jitterScale,
+ material,
+ volume*density,
+ volume,
+ color
+ );
+
+ g_particles[particleIndex] = newParticle;
+}
+
+@compute @workgroup_size(GridDispatchSize, GridDispatchSize)
+fn csMain( @builtin(global_invocation_id) id: vec3u )
+{
+ if(!insideGuardian(id.xy, g_simConstants.gridSize, GuardianSize+1))
+ {
+ return;
+ }
+
+ let gridSize = g_simConstants.gridSize;
+ let pos = vec2f(id.xy);
+
+
+
+ for(var shapeIndex = 0u; shapeIndex < g_simConstants.shapeCount; shapeIndex++)
+ {
+ let shape = g_shapes[shapeIndex];
+
+ let isEmitter = shape.functionality == ShapeFunctionEmit;
+ let isInitialEmitter = shape.functionality == ShapeFunctionInitialEmit;
+
+ if(!(isEmitter || isInitialEmitter))
+ {
+ continue;
+ }
+
+ let particleCountPerCellAxis = select(u32(g_simConstants.particlesPerCellAxis), 1, shape.emitMaterial == MaterialLiquid || shape.emitMaterial == MaterialSand);
+ let volumePerParticle = 1.0f / f32(particleCountPerCellAxis*particleCountPerCellAxis);
+
+ var c = collide(shape, pos);
+ if(c.collides)
+ {
+ let emitEvery = u32(1.0 / (shape.emissionRate * g_simConstants.deltaTime));
+
+
+ for(var i = 0u; i < particleCountPerCellAxis; i++)
+ {
+ for(var j = 0u; j < particleCountPerCellAxis; j++)
+ {
+ let hashCodeX = hash(id.x*particleCountPerCellAxis + i);
+ let hashCodeY = hash(id.y*particleCountPerCellAxis + j);
+ let hashCode = hash(hashCodeX + hashCodeY);
+
+ let emitDueToMyTurnHappening = isEmitter && 0 == ((hashCode + g_simConstants.simFrame) % emitEvery);
+ let emitDueToInitialEmission = isInitialEmitter && g_simConstants.simFrame == 0;
+
+ if(emitDueToInitialEmission || emitDueToMyTurnHappening)
+ {
+ addParticle(pos + vec2f(f32(i),f32(j))/f32(particleCountPerCellAxis), shape.emitMaterial, volumePerParticle, 1.0, 1.0/f32(particleCountPerCellAxis));
+ }
+ }
+ }
+ }
+ }
+}
\ No newline at end of file
diff --git a/shaders/particleIntegrate.wgsl b/shaders/particleIntegrate.wgsl
new file mode 100644
index 0000000..774edba
--- /dev/null
+++ b/shaders/particleIntegrate.wgsl
@@ -0,0 +1,183 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include particle.inc
+//!include simConstants.inc
+//!include matrix.inc
+//!include random.inc
+//!include shapes.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_particleCount : array;
+@group(0) @binding(2) var g_particles : array;
+@group(0) @binding(3) var g_shapes : array;
+@group(0) @binding(4) var g_freeCount : array>;
+@group(0) @binding(5) var g_freeIndices : array;
+
+@compute @workgroup_size(ParticleDispatchSize, 1, 1)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ if(id.x >= g_particleCount[0])
+ {
+ return;
+ }
+
+ var particle = g_particles[id.x];
+
+ if(particle.enabled == 0)
+ {
+ return;
+ }
+
+ if(particle.material == MaterialLiquid)
+ {
+ // The liquid material only cares about the determinant of the deformation gradient.
+ // We can use the regular MPM integration below to evolve the deformation gradient, but
+ // this approximation actually conserves more volume.
+ // This is based on det(F^n+1) = det((I + D)*F^n) = det(I+D)*det(F^n)
+ // and noticing that D is a small velocity, we can use the identity det(I + D) ≈ 1 + tr(D) to first order
+ // ending up with det(F^n+1) = (1+tr(D))*det(F^n)
+ // Then we directly set particle.liquidDensity to reflect the approximately integrated volume.
+ // The liquid material does not actually use the deformation gradient matrix.
+ particle.liquidDensity *= (tr(particle.deformationDisplacement) + 1.0);
+
+ // Safety clamp to avoid instability with very small densities.
+ particle.liquidDensity = max(particle.liquidDensity, 0.1);
+ }
+ else
+ {
+ // Integrate transform using standard MPM formula
+ particle.deformationGradient = (Identity + particle.deformationDisplacement) * particle.deformationGradient;
+ }
+
+ if(particle.material != MaterialLiquid)
+ {
+ // SVD is necessary at least for safety clamp
+ var svdResult = svd(particle.deformationGradient);
+
+ // Clamp the lower bound of scale to prevent situations where large forces are generated leading to explosions
+ svdResult.Sigma = clamp(svdResult.Sigma, vec2f(0.2), vec2f(10000.0));
+
+ // Plasticity implementations
+ if(particle.material == MaterialSand)
+ {
+ // Drucker-Prager sand based on:
+ // Gergely Klár, Theodore Gast, Andre Pradhana, Chuyuan Fu, Craig Schroeder, Chenfanfu Jiang, and Joseph Teran. 2016.
+ // Drucker-prager elastoplasticity for sand animation. ACM Trans. Graph. 35, 4, Article 103 (July 2016), 12 pages.
+ // https://doi.org/10.1145/2897824.2925906
+ let sinPhi = sin(g_simConstants.frictionAngle/180.0 * 3.14159);
+ let alpha = sqrt(2.0/3.0)*2.0*sinPhi/(3.0 - sinPhi);
+ let beta = g_simConstants.beta;
+
+ let eDiag = log(max(abs(svdResult.Sigma), vec2f(1e-6)));
+
+ let eps = diag(eDiag);
+ let trace = tr(eps) + particle.logJp;
+
+ let eHat = eps - (trace / 2) * Identity;
+ let frobNrm = length(vec2f(eHat[0][0], eHat[1][1]));
+
+ if(trace >= 0.0)
+ {
+ // In this case the motion is expansionary and we should not resist it at all.
+ // This means we should forget about any deformation that has occurred, which we do by setting Sigma to 1.
+ svdResult.Sigma = vec2f(1);
+ particle.logJp = beta * trace;
+ }
+ else
+ {
+ particle.logJp = 0;
+ let deltaGammaI = frobNrm + (g_simConstants.elasticityRatio + 1.0) * trace * alpha;
+ if(deltaGammaI > 0)
+ {
+ // Project to cone surface.
+ // This means we have to forget some deformation that the particle has undergone.
+ let h = eDiag - deltaGammaI/frobNrm * (eDiag - (trace*0.5)) ;
+ svdResult.Sigma = exp(h);
+ }
+ }
+ }
+ else if(particle.material == MaterialVisco)
+ {
+ // Very simple plasticity with volume preservation
+ let yieldSurface = exp(1-g_simConstants.plasticity);
+
+ // Record the volume before plasticity calculation
+ let J = svdResult.Sigma.x*svdResult.Sigma.y;
+
+ // Forget any deformation beyond the yield surface
+ svdResult.Sigma = clamp(svdResult.Sigma, vec2f(1.0/yieldSurface), vec2f(yieldSurface));
+
+ // Re-scale to original volume
+ let newJ = svdResult.Sigma.x*svdResult.Sigma.y;
+ svdResult.Sigma *= sqrt(J/newJ);
+ }
+
+ particle.deformationGradient = svdResult.U * diag(svdResult.Sigma) * svdResult.Vt;
+ }
+
+ // Integrate position
+ particle.position += particle.displacement;
+
+ // Mouse interaction
+ if(g_simConstants.mouseActivation > 0)
+ {
+ let offset = particle.position - g_simConstants.mousePosition;
+ let lenOffset = max(length(offset), 0.0001);
+ if(lenOffset < g_simConstants.mouseRadius)
+ {
+ let normOffset = offset/lenOffset;
+
+ if(g_simConstants.mouseFunction == MouseFunctionPush)
+ {
+ particle.displacement += normOffset * g_simConstants.mouseActivation;
+ }
+ else if(g_simConstants.mouseFunction == MouseFunctionGrab)
+ {
+ particle.displacement = 0.7*g_simConstants.mouseVelocity*g_simConstants.deltaTime;
+ }
+ }
+ }
+
+ // Gravity acceleration is normalized to the vertical size of the window.
+ particle.displacement.y -= f32(g_simConstants.gridSize.y)*g_simConstants.gravityStrength*g_simConstants.deltaTime*g_simConstants.deltaTime;
+
+ // Free count may be negative because of emission. So make sure it is at last zero before incrementing.
+ atomicMax(&g_freeCount[0], 0i);
+
+ for(var shapeIndex = 0u; shapeIndex < g_simConstants.shapeCount; shapeIndex++)
+ {
+ let shape = g_shapes[shapeIndex];
+
+ // Push particles out of colliders. Most of the work should have been done at the grid level already.
+ if(shape.functionality == ShapeFunctionCollider)
+ {
+ let collideResult = collide(shape, particle.position);
+
+ if(collideResult.collides)
+ {
+ particle.displacement -= collideResult.penetration*collideResult.normal;
+ }
+ }
+
+ // Delete particles if they are inside a drain shape
+ if(shape.functionality == ShapeFunctionDrain)
+ {
+ if(collide(shape, particle.position).collides)
+ {
+ particle.enabled = 0;
+
+ // Add index of this particle to free list
+ let freeIndex = atomicAdd(&g_freeCount[0], 1i);
+ g_freeIndices[u32(freeIndex)] = id.x;
+ }
+ }
+ }
+
+ // Ensure particles are inside the simulation limits
+ particle.position = projectInsideGuardian(particle.position, g_simConstants.gridSize, GuardianSize);
+
+ g_particles[id.x] = particle;
+}
\ No newline at end of file
diff --git a/shaders/particleRender.wgsl b/shaders/particleRender.wgsl
new file mode 100644
index 0000000..d372779
--- /dev/null
+++ b/shaders/particleRender.wgsl
@@ -0,0 +1,101 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include particle.inc
+//!include simConstants.inc
+
+// Code-generated by UniformBufferFactory
+//!insert RenderConstants
+
+// const values like RenderModeStandard etc
+// that are plugged in from js
+//!insert RenderEnums
+
+@group(0) @binding(0) var g_constants : RenderConstants;
+@group(0) @binding(1) var g_particles : array;
+
+
+struct VertexOutput {
+ @builtin(position) interpolatedVertexPosition : vec4f,
+ @location(0) particlePosition : vec2f,
+ @location(1) vertexPosition : vec2f,
+ @location(2) particleColor : vec3f,
+ @location(3) particleRadius : f32,
+};
+
+const s_quadVertices = array(
+ vec2f(-1, -1),
+ vec2f(1, -1),
+ vec2f(1, 1),
+ vec2f(-1, -1),
+ vec2f(1, 1),
+ vec2f(-1, 1)
+);
+
+const s_antialiasingWidth = 1;
+
+@vertex
+fn vertexMain(@builtin(vertex_index) vertexId: u32, @builtin(instance_index) instanceId : u32) -> VertexOutput {
+ let onePixel = 1.0 / f32(g_constants.canvasSize.x);
+
+ let particle = g_particles[instanceId];
+
+ if(particle.enabled == 0)
+ {
+ return VertexOutput(
+ vec4f(0,0,0,0),
+ vec2f(0,0),
+ vec2f(0,0),
+ vec3f(0,0,0),
+ 0.0
+ );
+ }
+
+ var particlePosition : vec2f = particle.position;
+
+ var particleMaterialRadius = g_constants.particleRadiusTimestamp.x * sqrt(particle.volume);
+
+ var quadVertexPosition = s_quadVertices[vertexId] * (1+s_antialiasingWidth*onePixel)*particleMaterialRadius;
+ var vertexPosition = quadVertexPosition + particlePosition;
+
+ var vertexPositionInRenderSpace = (vertexPosition - g_constants.viewPos) / g_constants.viewExtent;
+
+ var color = particle.color;
+
+ if(g_constants.renderMode == RenderModeStandard)
+ {
+
+ }
+ if(g_constants.renderMode == RenderModeCompression)
+ {
+ let density = particle.liquidDensity;
+ color = vec3f(-5*log(density), 5*log(density), color.z);
+ }
+ else if(g_constants.renderMode == RenderModeVelocity)
+ {
+ color = vec3f((particle.displacement*g_constants.deltaTime*1000.0 + 0.5f), color.z);
+ }
+
+ return VertexOutput(
+ vec4f(vertexPositionInRenderSpace, 0, 1),
+ particlePosition,
+ vertexPosition,
+ color,
+ particleMaterialRadius,
+ );
+}
+
+@fragment
+fn fragmentMain(input: VertexOutput) -> @location(0) vec4f {
+ let onePixel = 1.0 / f32(g_constants.canvasSize.x);
+
+ let particleRadius = input.particleRadius;
+ let fragmentPosition = input.vertexPosition;
+ let fragOffsetFromCenter = fragmentPosition - input.particlePosition;
+ let distanceFromCenter = length(fragOffsetFromCenter);
+
+ let alpha = smoothstep(1.0, 0.0, (distanceFromCenter-particleRadius)/(onePixel*s_antialiasingWidth*particleRadius) );
+
+ return vec4f(input.particleColor,alpha);
+}
\ No newline at end of file
diff --git a/shaders/particleToGrid.wgsl b/shaders/particleToGrid.wgsl
new file mode 100644
index 0000000..9e1190e
--- /dev/null
+++ b/shaders/particleToGrid.wgsl
@@ -0,0 +1,66 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include particle.inc
+//!include simConstants.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_particleCount : array;
+@group(0) @binding(2) var g_particles : array;
+@group(0) @binding(3) var g_grid : array>;
+
+@compute @workgroup_size(ParticleDispatchSize, 1, 1)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ if(id.x >= g_particleCount[0])
+ {
+ return;
+ }
+
+ let particle = g_particles[id.x];
+
+ if(particle.enabled == 0)
+ {
+ return;
+ }
+
+ var p = particle.position;
+ let d = particle.displacement;
+ let D = particle.deformationDisplacement;
+ let m = particle.mass;
+
+ let weightInfo = quadraticWeightInit(p);
+
+ // Iterate over local 3x3 neigbourhood
+ for(var i = 0; i < 3; i++)
+ {
+ for(var j = 0; j < 3; j++)
+ {
+ // Weight corresponding to this neighbourhood cell
+ let weight = weightInfo.weights[i].x * weightInfo.weights[j].y;
+
+ // 2d index of this cell in the grid
+ let neighbourCellIndex = vec2u(vec2i(weightInfo.cellIndex) + vec2i(i,j));
+
+ // Linear index in the buffer
+ let gridVertexIdx = gridVertexIndex(neighbourCellIndex, g_simConstants.gridSize);
+
+ let offset = vec2f(neighbourCellIndex) - p + 0.5;
+
+ let weightedMass = weight * m;
+ let momentum = weightedMass * (d + D * offset);
+
+ atomicAdd(&g_grid[gridVertexIdx + 0], encodeFixedPoint(momentum.x, g_simConstants.fixedPointMultiplier));
+ atomicAdd(&g_grid[gridVertexIdx + 1], encodeFixedPoint(momentum.y, g_simConstants.fixedPointMultiplier));
+ atomicAdd(&g_grid[gridVertexIdx + 2], encodeFixedPoint(weightedMass, g_simConstants.fixedPointMultiplier));
+
+ // This is only required if we are going to mix in the grid volume to the liquid volume
+ if(g_simConstants.useGridVolumeForLiquid != 0)
+ {
+ atomicAdd(&g_grid[gridVertexIdx + 3], encodeFixedPoint(particle.volume * weight, g_simConstants.fixedPointMultiplier));
+ }
+ }
+ }
+}
\ No newline at end of file
diff --git a/shaders/particleUpdatePBMPM.wgsl b/shaders/particleUpdatePBMPM.wgsl
new file mode 100644
index 0000000..e43e778
--- /dev/null
+++ b/shaders/particleUpdatePBMPM.wgsl
@@ -0,0 +1,90 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include particle.inc
+//!include simConstants.inc
+//!include matrix.inc
+
+@group(0) @binding(0) var g_simConstants : SimConstants;
+@group(0) @binding(1) var g_particleCount : array;
+@group(0) @binding(2) var g_particles : array;
+
+@compute @workgroup_size(ParticleDispatchSize, 1, 1)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ if(id.x >= g_particleCount[0])
+ {
+ return;
+ }
+
+ var particle = g_particles[id.x];
+
+ if(particle.enabled == 0)
+ {
+ return;
+ }
+
+ if(particle.material == MaterialLiquid)
+ {
+ // Simple liquid viscosity: just remove deviatoric part of the deformation displacement
+ let deviatoric = -1.0*(particle.deformationDisplacement + transpose(particle.deformationDisplacement));
+ particle.deformationDisplacement += g_simConstants.liquidViscosity*0.5*deviatoric;
+
+ // Volume preservation constraint:
+ // we want to generate hydrostatic impulses with the form alpha*I
+ // and we want the liquid volume integration (see particleIntegrate) to yield 1 = (1+tr(alpha*I + D))*det(F) at the end of the timestep.
+ // where det(F) is stored as particle.liquidDensity.
+ // Rearranging, we get the below expression that drives the deformation displacement towards preserving the volume.
+ let alpha = 0.5*(1.0/particle.liquidDensity - tr(particle.deformationDisplacement) - 1.0);
+ particle.deformationDisplacement += g_simConstants.liquidRelaxation*alpha*Identity;
+ }
+ else if(particle.material == MaterialElastic || particle.material == MaterialVisco)
+ {
+ let F = (Identity + particle.deformationDisplacement) * particle.deformationGradient;
+
+ var svdResult = svd(F);
+
+ // Closest matrix to F with det == 1
+ let df = det(F);
+ let cdf = clamp(abs(df), 0.1, 1000);
+ let Q = (1.0f/(sign(df)*sqrt(cdf)))*F;
+ // Interpolate between the two target shapes
+ let alpha = g_simConstants.elasticityRatio;
+ let tgt = alpha*(svdResult.U*svdResult.Vt) + (1.0-alpha)*Q;
+
+ let diff = (tgt*inverse(particle.deformationGradient) - Identity) - particle.deformationDisplacement;
+ particle.deformationDisplacement += g_simConstants.elasticRelaxation*diff;
+
+ }
+ else if(particle.material == MaterialSand)
+ {
+
+
+ let F = (Identity + particle.deformationDisplacement) * particle.deformationGradient;
+
+ var svdResult = svd(F);
+
+ if(particle.logJp == 0)
+ {
+ svdResult.Sigma = clamp(svdResult.Sigma, vec2f(1, 1), vec2f(1000, 1000));
+ }
+
+ // Closest matrix to F with det == 1
+ let df = det(F);
+ let cdf = clamp(abs(df), 0.1, 1);
+ let Q = (1.0f/(sign(df)*sqrt(cdf)))*F;
+ // Interpolate between the two target shapes
+ let alpha = g_simConstants.elasticityRatio;
+ let tgt = alpha*(svdResult.U*mat2x2f(svdResult.Sigma.x, 0, 0, svdResult.Sigma.y)*svdResult.Vt) + (1.0-alpha)*Q;
+
+ let diff = (tgt*inverse(particle.deformationGradient) - Identity) - particle.deformationDisplacement;
+ particle.deformationDisplacement += g_simConstants.elasticRelaxation*diff;
+
+ let deviatoric = -1.0*(particle.deformationDisplacement + transpose(particle.deformationDisplacement));
+ particle.deformationDisplacement += g_simConstants.liquidViscosity*0.5*deviatoric;
+ }
+
+ g_particles[id.x] = particle;
+}
\ No newline at end of file
diff --git a/shaders/random.inc.wgsl b/shaders/random.inc.wgsl
new file mode 100644
index 0000000..dacc11d
--- /dev/null
+++ b/shaders/random.inc.wgsl
@@ -0,0 +1,15 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+fn hash(input : u32) -> u32
+{
+ let state = input * 747796405u + 2891336453u;
+ let word = ((state >> ((state >> 28u) + 4u)) ^ state) * 277803737u;
+ return (word >> 22u) ^ word;
+}
+
+fn randomFloat(input : u32) -> f32
+{
+ return f32(hash(input) % 10000) / 9999.0;
+}
diff --git a/shaders/setIndirectArgs.wgsl b/shaders/setIndirectArgs.wgsl
new file mode 100644
index 0000000..65450f5
--- /dev/null
+++ b/shaders/setIndirectArgs.wgsl
@@ -0,0 +1,21 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+
+@group(0) @binding(0) var g_particleCount : array;
+@group(0) @binding(1) var g_simIndirectArgs : array;
+@group(0) @binding(2) var g_renderIndirectArgs : array;
+
+fn divUp(threadCount : u32, divisor : u32) -> u32
+{
+ return (threadCount + divisor - 1) / divisor;
+}
+
+@compute @workgroup_size(1)
+fn csMain( @builtin(global_invocation_id) id: vec3 )
+{
+ g_simIndirectArgs[0] = divUp(g_particleCount[0], ParticleDispatchSize);
+ g_renderIndirectArgs[1] = g_particleCount[0];
+}
\ No newline at end of file
diff --git a/shaders/shapes.inc.wgsl b/shaders/shapes.inc.wgsl
new file mode 100644
index 0000000..18a2ef4
--- /dev/null
+++ b/shaders/shapes.inc.wgsl
@@ -0,0 +1,56 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include matrix.inc
+
+// Code-generated by BufferFactory
+//!insert SimShape
+
+struct CollideResult
+{
+ collides: bool,
+ penetration: f32,
+ normal: vec2f,
+ pointOnCollider: vec2f,
+};
+
+fn collide(shape: SimShape, pos: vec2f) -> CollideResult
+{
+ if(shape.shapeType == ShapeTypeCircle)
+ {
+ let offset = shape.position - pos;
+ let offsetLen = length(offset);
+ let normal = offset * select(1.0/offsetLen, 0, offsetLen == 0);
+ return CollideResult(
+ offsetLen <= shape.radius,
+ -(offsetLen - shape.radius),
+ normal,
+ shape.position+normal*shape.radius,
+ );
+ }
+ else if(shape.shapeType == ShapeTypeBox)
+ {
+ let offset = pos - shape.position;
+ let R = rot(shape.rotation / 180.0 * 3.14159);
+ let rotOffset = R * offset;
+ let sx = sign(rotOffset.x);
+ let sy = sign(rotOffset.y);
+ let penetration = -(abs(rotOffset) - shape.halfSize);
+ let normal = transpose(R)*select(vec2f(sx,0), vec2f(0,sy), penetration.y < penetration.x);
+ let minPen = min(penetration.x, penetration.y);
+
+ let pointOnBox = shape.position + transpose(R)*clamp(rotOffset, -shape.halfSize, shape.halfSize);
+
+ return CollideResult(
+ minPen > 0,
+ minPen,
+ -normal,
+ pointOnBox
+ );
+ }
+ else
+ {
+ return CollideResult(false, 0.0, vec2f(0,0), vec2f(0,0));
+ }
+}
diff --git a/shaders/simConstants.inc.wgsl b/shaders/simConstants.inc.wgsl
new file mode 100644
index 0000000..bc71c04
--- /dev/null
+++ b/shaders/simConstants.inc.wgsl
@@ -0,0 +1,27 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+// Code-generated by UniformBufferFactory
+//!insert SimConstants
+
+// const values like SolverTypeExplicit etc
+// that are plugged in from js
+//!insert SimEnums
+
+fn gridVertexIndex(gridVertex : vec2u, gridSize : vec2u) -> u32
+{
+ // Currently using lexicographical ordering
+ // 4 components per grid vertex
+ return u32(4*(gridVertex.y * gridSize.x + gridVertex.x));
+}
+
+fn decodeFixedPoint(fixedPoint : i32, fixedPointMultiplier : u32) -> f32
+{
+ return f32(fixedPoint) / f32(fixedPointMultiplier);
+}
+
+fn encodeFixedPoint(floatingPoint : f32, fixedPointMultiplier: u32) -> i32
+{
+ return i32(floatingPoint * f32(fixedPointMultiplier));
+}
diff --git a/src/buffer_factory.js b/src/buffer_factory.js
new file mode 100644
index 0000000..34be180
--- /dev/null
+++ b/src/buffer_factory.js
@@ -0,0 +1,295 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+"use strict";
+
+export const f32 = 'f32';
+export const vec2f = 'vec2f';
+export const vec3f = 'vec3f';
+export const u32 = 'u32';
+export const vec2u = 'vec2u';
+
+export const Storage = 'storage';
+export const Uniform = 'uniform';
+
+// This class allows us to conveniently code-generate structs to be used as uniform buffers
+// and copy data to the gpu, including respecting packing rules, in order to avoid
+// having to do a lot of boilerplate for every parameter we want to pass.
+export class BufferFactory
+{
+ constructor(name, mode)
+ {
+ this.name = name;
+ this.mode = mode;
+ this.elements = []
+ this.compiled = false;
+ this.paddingCount = 0;
+ this.totalCount = 0;
+ }
+
+ // Add a parameter with the given name and type.
+ // The name is used as a key for binding data on the js side and
+ // also in the generated wgsl code.
+ add(name, type)
+ {
+ console.assert(!this.compiled);
+
+ const requiredSlotCount = getSize(type);
+ const requiredAlignment = getAlignment(type);
+
+ let alignmentAdjustment = this.totalCount % requiredAlignment;
+ while(alignmentAdjustment !== 0)
+ {
+ // Add a padding element to align up the member to the correct address
+ this.add(`padding${this.paddingCount}`, f32);
+ this.paddingCount += 1;
+ alignmentAdjustment -= 1;
+ }
+
+ this.elements.push({
+ name: name,
+ type: type,
+ value: undefined,
+ offset: this.totalCount,
+ slotCount: requiredSlotCount
+ });
+
+ this.totalCount += requiredSlotCount;
+ }
+
+ // Finalize the factory after parameters have been added
+ compile()
+ {
+ console.assert(!this.compiled);
+
+ if(this.mode == Uniform)
+ {
+ // Round up size to a multiple of 16
+ this.totalCount = Math.ceil(this.totalCount/16)*16;
+ }
+
+ this.compiled = true;
+ }
+
+ // Assemble text that can be pasted into the const buffer definition
+ getShaderText()
+ {
+ console.assert(this.compiled);
+
+ let shaderText = `struct ${this.name}\n{\n`
+ for(const elem of this.elements)
+ {
+ shaderText += `${elem.name}: ${elem.type},\n`;
+ }
+ shaderText += "};\n";
+
+ return shaderText;
+ }
+
+ getTotalSizeInWords()
+ {
+ console.assert(this.compiled);
+ return this.totalCount;
+ }
+
+ // Build a uniform buffer object containing the currently stored values
+ constructUniformBuffer(device, values)
+ {
+ console.assert(this.compiled);
+ console.assert(this.mode == Uniform);
+ console.assert(Array.isArray(values));
+
+ for(const elem of this.elements)
+ {
+ elem.value = undefined;
+ }
+
+ for(const valueObject of values)
+ {
+ for(const elem of this.elements)
+ {
+ if(elem.name in valueObject)
+ {
+ elem.value = valueObject[elem.name];
+ }
+ }
+ }
+
+ for(const elem of this.elements)
+ {
+ if(elem.value === undefined && elem.name.indexOf('padding') == -1)
+ {
+ throw `Element ${elem.name} has never had its value set.`;
+ }
+ }
+
+ const uniformBuffer = device.createBuffer({
+ label: this.name,
+ size: this.totalCount * 4,
+ usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST
+ });
+
+ const uniformValues = new Float32Array(this.totalCount);
+
+ for(const elem of this.elements)
+ {
+ if(elem.value === undefined)
+ {
+ continue;
+ }
+
+ // If we need to write an integer type then we have to trick
+ // the data into the uniform values array using this mechanism.
+ if(elem.type == u32)
+ {
+ const castArray = new Int32Array(1);
+ castArray.set([elem.value], 0);
+ const castArrayFloat = new Float32Array(castArray.buffer);
+ uniformValues.set(castArrayFloat, elem.offset);
+ }
+ else if(elem.type == vec2u)
+ {
+ const castArray = new Int32Array(2);
+ castArray.set(elem.value, 0);
+ const castArrayFloat = new Float32Array(castArray.buffer);
+ uniformValues.set(castArrayFloat, elem.offset);
+ }
+ else if(elem.type == f32)
+ {
+ // Float values can be written through an array
+ uniformValues.set([elem.value], elem.offset);
+ }
+ else
+ {
+ // Float vector values can be written directly
+ uniformValues.set(elem.value, elem.offset);
+ }
+ }
+
+ device.queue.writeBuffer(uniformBuffer, 0, uniformValues);
+
+ return uniformBuffer;
+ }
+
+ constructStorageBuffer(device, elements)
+ {
+ console.assert(this.compiled);
+ console.assert(this.mode == Storage);
+ console.assert(Array.isArray(elements));
+
+ // Note - no check that all data is properly set
+ const elementCount = elements.length;
+
+ const storageBuffer = device.createBuffer({
+ label: this.name,
+ size: this.totalCount * elementCount * 4,
+ usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
+ });
+
+ const storageValues = new Float32Array(this.totalCount*elementCount);
+
+ for(var i = 0; i < elementCount; ++i)
+ {
+ const outputOffset = i*this.totalCount;
+
+ for(const elem of this.elements)
+ {
+ elem.value = undefined;
+ }
+
+ for(const elem of this.elements)
+ {
+ if(elem.name in elements[i])
+ {
+ elem.value = elements[i][elem.name];
+ }
+ }
+
+ for(const elem of this.elements)
+ {
+ if(elem.value === undefined && elem.name.indexOf('padding') == -1)
+ {
+ throw `Element ${elem.name} has never had its value set.`;
+ }
+ }
+
+ for(const elem of this.elements)
+ {
+ if(elem.value === undefined)
+ {
+ continue;
+ }
+
+ // If we need to write an integer type then we have to trick
+ // the data into the uniform values array using this mechanism.
+ if(elem.type == u32)
+ {
+ const castArray = new Int32Array(1);
+ castArray.set([elem.value], 0);
+ const castArrayFloat = new Float32Array(castArray.buffer);
+ storageValues.set(castArrayFloat, outputOffset + elem.offset);
+ }
+ else if(elem.type == vec2u)
+ {
+ const castArray = new Int32Array(2);
+ castArray.set(elem.value, 0);
+ const castArrayFloat = new Float32Array(castArray.buffer);
+ storageValues.set(castArrayFloat, outputOffset + elem.offset);
+ }
+ else if(elem.type == f32)
+ {
+ // Float values can be written through an array
+ storageValues.set([elem.value], outputOffset + elem.offset);
+ }
+ else
+ {
+ // Float vector values can be written directly
+ storageValues.set(elem.value, outputOffset + elem.offset);
+ }
+ }
+ }
+
+ device.queue.writeBuffer(storageBuffer, 0, storageValues);
+
+ return storageBuffer;
+ }
+}
+
+// What should the size of each type be in multiples of the size of
+// an f32
+function getSize(type)
+{
+ switch(type)
+ {
+ case f32:
+ case u32:
+ return 1;
+ case vec2f:
+ case vec2u:
+ return 2;
+ case vec3f:
+ return 3;
+ default:
+ throw `Unsupported type [${type}]`;
+ }
+}
+
+// What should the alignment of each type be in multiples of
+// the size of an f32
+function getAlignment(type)
+{
+ switch(type)
+ {
+ case f32:
+ case u32:
+ return 1;
+ case vec2f:
+ case vec2u:
+ return 2;
+ case vec3f:
+ return 4;
+ default:
+ throw `Unsupported type [${type}]`;
+ }
+}
diff --git a/src/gpu.js b/src/gpu.js
new file mode 100644
index 0000000..80fef41
--- /dev/null
+++ b/src/gpu.js
@@ -0,0 +1,387 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+"use strict";
+
+import * as shader from "./shader.js"
+
+let context = {
+ pipelines: {},
+ maxParticleCount: 500000,
+ maxConstraintData: 100000,
+ // Note - not necessary to have these set dynamically
+ hashGridBucketCount: 100*1024,
+ hashGridBucketSize: 128,
+ maxNeighbourCount: 64,
+ maxTimeStampCount: 2048,
+
+ encoder: null,
+ frameTimeStampCount: 0,
+ frameTimeStampNames: {},
+
+ movingAverageTimeStamps: {},
+
+ timingStatsDirty: false,
+ particleCountDirty: false,
+
+ particleCount: 0,
+ particleFreeCount: 0,
+};
+
+export function getGpuContext() {return context;}
+
+export function divUp(threadCount, divisor)
+{
+ return Math.ceil(threadCount / divisor);
+}
+
+export function createBindGroup(name, shaderName, resources)
+{
+ let entries = [];
+ for(let i = 0; i < resources.length; ++i)
+ {
+ entries.push({binding: i, resource: {buffer: resources[i]}});
+ }
+
+ return context.device.createBindGroup({
+ label: name,
+ layout: context.pipelines[shaderName].getBindGroupLayout(0),
+ entries: entries
+ });
+}
+
+export function resetBuffers()
+{
+ // Constructs a buffer containing 4 integers of the given values
+ function construct4IntBuffer(name, usage, values)
+ {
+ context[name] = context.device.createBuffer({
+ name: name,
+ size: 16,
+ usage: usage
+ })
+
+ const valueArray = new Int32Array(4);
+ valueArray.set(values);
+ context.device.queue.writeBuffer(context[name], 0, valueArray);
+ }
+
+ // Construct various small buffers used for indirect dispatch, counting and staging
+ construct4IntBuffer('particleCountBuffer', GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, [0,0,0,0]);
+ construct4IntBuffer('particleCountStagingBuffer', GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, [0,0,0,0]);
+ construct4IntBuffer('particleRenderDispatchBuffer', GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.INDIRECT, [6,0,0,0]);
+ construct4IntBuffer('particleSimDispatchBuffer', GPUBufferUsage.STORAGE | GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_DST, [0,1,1,0]);
+ construct4IntBuffer('particleFreeCountBuffer', GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, [0,0,0,0]);
+ construct4IntBuffer('particleFreeCountStagingBuffer', GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, [0,0,0,0]);
+
+ // Construct particle buffer.
+ // Must be kept in sync with MPMParticle in particle.inc.wgsl
+ const particleFloatCount = 25;
+
+ context.particleBuffer = context.device.createBuffer({
+ label: "particles",
+ size: context.maxParticleCount * 4 * particleFloatCount,
+ usage: GPUBufferUsage.STORAGE
+ });
+
+ context.particleFreeIndicesBuffer = context.device.createBuffer({
+ label: 'freeIndices',
+ size: context.maxParticleCount * 4,
+ usage: GPUBufferUsage.STORAGE
+ });
+
+ // Data needed for Position Based Fluids impl
+ {
+ // Extra particle buffer required for ping-pong access
+ context.tempParticleBuffer = context.device.createBuffer({
+ label: 'tempParticles',
+ size: context.maxParticleCount * 4 * particleFloatCount,
+ usage: GPUBufferUsage.STORAGE
+ });
+
+ // Each bucket has one count
+ context.hashGridBucketCounts = context.device.createBuffer({
+ label: "hashGridBucketCounts",
+ size: context.hashGridBucketCount * 4,
+ usage: GPUBufferUsage.STORAGE
+ });
+
+ // Each bucket has hashGridBucketSize elements of 2 4-byte words
+ context.hashGridData = context.device.createBuffer({
+ label: "hashGridData",
+ size: context.hashGridBucketCount * context.hashGridBucketSize * 2 * 4,
+ usage: GPUBufferUsage.STORAGE
+ });
+
+ context.particleNeighbours = context.device.createBuffer({
+ label: 'particleNeighbours',
+ size: context.maxParticleCount * 4 * context.maxNeighbourCount,
+ usage: GPUBufferUsage.STORAGE
+ });
+ }
+}
+
+export function beginFrame()
+{
+ context.frameTimeStampCount = 0;
+ context.frameTimeStampNames = {};
+ context.encoder = context.device.createCommandEncoder();
+}
+
+export function endFrame()
+{
+ const canReadbackParticleCount = context.particleCountStagingBuffer.mapState === 'unmapped';
+ const canReadbackParticleFreeCount = context.particleFreeCountStagingBuffer.mapState === 'unmapped';
+ const canReadbackTimeStamps = context.canTimeStamp && context.timeStampResultBuffer.mapState === 'unmapped';
+
+ if(canReadbackParticleCount)
+ {
+ context.encoder.copyBufferToBuffer(context.particleCountBuffer, 0, context.particleCountStagingBuffer, 0, 4);
+ }
+
+ if(canReadbackParticleFreeCount)
+ {
+ context.encoder.copyBufferToBuffer(context.particleFreeCountBuffer, 0, context.particleFreeCountStagingBuffer, 0, 4);
+ }
+
+ if(canReadbackTimeStamps)
+ {
+ context.encoder.resolveQuerySet(context.timeStampQuerySet, 0, context.frameTimeStampCount, context.timeStampResolveBuffer, 0);
+
+ context.encoder.copyBufferToBuffer(context.timeStampResolveBuffer, 0, context.timeStampResultBuffer, 0, context.timeStampResultBuffer.size);
+ }
+
+ context.device.queue.submit([context.encoder.finish()]);
+
+ if(canReadbackParticleCount)
+ {
+ readbackParticleCount();
+ }
+
+ if(canReadbackParticleFreeCount)
+ {
+ readbackParticleFreeCount();
+ }
+
+ if(canReadbackTimeStamps)
+ {
+ readbackTimeStamps();
+ }
+
+ context.encoder = null;
+}
+
+// Helper function to dispatch a compute shader with the given name, resources
+// and dispatch size.
+// if GroupCount is an array then a regular dispatch is done, otherwise
+// it is bound as an indirect buffer and an indirect dispatch is done.
+export function computeDispatch(shaderName, resources, groupCount)
+{
+ // Construct array of resources in the required format for
+ // creating a bind group
+ let entries = []
+ for(let i = 0; i < resources.length; ++i)
+ {
+ entries.push({binding: i, resource: {buffer: resources[i]}});
+ }
+
+ const pipeline = shader.getComputePipeline(shaderName);
+
+
+ // Create bind group
+ const bindGroup = context.device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: entries});
+
+ //
+ const computePass = context.encoder.beginComputePass({
+ label: shaderName,
+ ...(context.canTimeStamp && {
+ timestampWrites: {
+ querySet: context.timeStampQuerySet,
+ beginningOfPassWriteIndex: context.frameTimeStampCount,
+ endOfPassWriteIndex: context.frameTimeStampCount + 1
+ }
+ })
+ });
+
+ computePass.setPipeline(pipeline);
+ computePass.setBindGroup(0, bindGroup);
+
+ if(Array.isArray(groupCount))
+ {
+ computePass.dispatchWorkgroups(groupCount[0], groupCount[1], groupCount[2]);
+ }
+ else
+ {
+ computePass.dispatchWorkgroupsIndirect(groupCount, 0);
+ }
+ computePass.end()
+
+ if(context.canTimeStamp)
+ {
+ if(shaderName in context.frameTimeStampNames)
+ {
+ context.frameTimeStampNames[shaderName].push(context.frameTimeStampCount);
+ }
+ else
+ {
+ context.frameTimeStampNames[shaderName] = [context.frameTimeStampCount];
+ }
+
+ context.frameTimeStampCount += 2;
+ }
+}
+
+// Try to initialize webgpu device and set up the basic objects.
+// This may fail on some browsers like Firefox, in which case we put a message in the dom.
+// It may also fail if the browser has blacklisted webgpu due to previously going OOM.
+// If this happens then it may be necessary to restart the whole program.
+export async function init(insertHandlers)
+{
+ context.insertHandlers = insertHandlers;
+
+ // Initialize device
+ if (!navigator.gpu) {
+ throw "WebGPU not supported on this browser.";
+ }
+
+ context.adapter = await navigator.gpu.requestAdapter();
+ if (!context.adapter) {
+ throw "No appropriate GPUAdapter found.";
+ }
+
+ context.canTimeStamp = context.adapter.features.has('timestamp-query');
+
+ if(!context.canTimeStamp)
+ {
+ console.warn('This WebGPU implementation does not support timestamp queries. Timing info will not be available.');
+ }
+
+ context.device = await context.adapter.requestDevice({
+ requiredFeatures: [
+ context.canTimeStamp ? ['timestamp-query']: []
+ ]
+ });
+
+ await shader.init(context.device, insertHandlers);
+
+ // Set back buffer pixel format
+ context.context = document.getElementById('canvas').getContext("webgpu", {alpha: true});
+ const canvasFormat = navigator.gpu.getPreferredCanvasFormat();
+ context.context.configure({
+ device: context.device,
+ format: canvasFormat,
+ alphaMode: 'premultiplied'
+ });
+
+ // Load the particle rendering module
+ const renderShaderModule = shader.getShaderModule(shader.Shaders.particleRender);
+
+ // Construct pipeline for particle rendering
+ context.pipelines['particleRender'] = context.device.createRenderPipeline({
+ label: "Render Pipeline",
+ layout: "auto",
+ vertex: {
+ module: renderShaderModule,
+ entryPoint: "vertexMain",
+ buffers: []
+ },
+ fragment: {
+ module: renderShaderModule,
+ entryPoint: "fragmentMain",
+ targets: [{
+ format: canvasFormat,
+ blend: {
+ alpha: {
+ dstFactor: 'one-minus-src-alpha',
+ srcFactor: 'src-alpha',
+ operation: 'add'
+ },
+ color: {
+ dstFactor: 'one-minus-src-alpha',
+ srcFactor: 'src-alpha',
+ operation: 'add'
+ }
+ }
+ }]
+ },
+ });
+
+ if(context.canTimeStamp)
+ {
+ context.timeStampQuerySet = context.device.createQuerySet({
+ type: 'timestamp',
+ count: context.maxTimeStampCount*2 // Begin and end
+ });
+
+ context.timeStampResolveBuffer = context.device.createBuffer({
+ size: context.timeStampQuerySet.count * 8, // Timestamps are uint64
+ usage: GPUBufferUsage.QUERY_RESOLVE | GPUBufferUsage.COPY_SRC
+ });
+
+ context.timeStampResultBuffer = context.device.createBuffer({
+ size: context.timeStampQuerySet.count * 8,
+ usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
+ });
+ }
+}
+
+// Initiate readback of the particle count staging buffer.
+function readbackParticleCount()
+{
+ context.particleCountStagingBuffer.mapAsync(GPUMapMode.READ, 0, 16).then(() => {
+ const buf = context.particleCountStagingBuffer.getMappedRange(0, 16);
+ const view = new Int32Array(buf);
+ context.particleCount = view[0];
+ context.particleCount = Math.min(context.maxParticleCount, context.particleCount);
+ context.particleCountStagingBuffer.unmap();
+ context.particleCountDirty = true;
+ });
+}
+
+function readbackParticleFreeCount()
+{
+ context.particleFreeCountStagingBuffer.mapAsync(GPUMapMode.READ, 0, 16).then(() => {
+ const buf = context.particleFreeCountStagingBuffer.getMappedRange(0, 16);
+ const view = new Int32Array(buf);
+ context.particleFreeCount = view[0];
+ context.particleFreeCountStagingBuffer.unmap();
+ context.particleCountDirty = true;
+ });
+}
+
+function readbackTimeStamps()
+{
+ context.timeStampResultBuffer.mapAsync(GPUMapMode.READ).then(() => {
+ const times = new BigInt64Array(context.timeStampResultBuffer.getMappedRange());
+
+ var movingAverageTimeStamps = {};
+
+
+ for(const name of Object.keys(context.frameTimeStampNames))
+ {
+ var total = 0;
+ for(const index of context.frameTimeStampNames[name])
+ {
+ total += Number(times[index+1] - times[index]);
+ }
+
+ movingAverageTimeStamps[name] = total;
+ }
+ context.timeStampResultBuffer.unmap();
+
+ for(const name of Object.keys(movingAverageTimeStamps))
+ {
+ if(name in context.movingAverageTimeStamps)
+ {
+ movingAverageTimeStamps[name] = 0.01*movingAverageTimeStamps[name] + 0.99*context.movingAverageTimeStamps[name]
+ }
+ }
+
+ context.movingAverageTimeStamps = movingAverageTimeStamps;
+
+ context.timingStatsDirty = true;
+ });
+}
\ No newline at end of file
diff --git a/src/main.js b/src/main.js
new file mode 100644
index 0000000..0bc485a
--- /dev/null
+++ b/src/main.js
@@ -0,0 +1,475 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts. All rights reserved.
+//-----------------------------------------------------------------------------
+
+"use strict";
+
+import * as ui from "./ui.js"
+import * as gpu from "./gpu.js"
+import * as sim from "./sim.js"
+import {getThrottlingRatio} from "./time.js"
+import * as render from "./render.js"
+
+let g_prevInputs;
+
+let g_loading = true;
+let g_uiIsHidden = true;
+let g_reset = true;
+let g_pause = false;
+
+let g_mouseRadius = 100;
+
+init();
+
+function init()
+{
+ // Build list of insert handlers.
+ // An insert handler is just a KVP that can substitute a //!insert directive
+ // in a wgsl shader.
+ let insertHandlers = {};
+
+ // Setup the UI sidebar
+ ui.init(() => {g_reset = true;}, () => {g_pause = !g_pause;});
+
+ loadSceneFromUrl('./scenes/damBreak.json');
+
+ // Setup sim.
+ // Sim will add to the list of insert handlers.
+ sim.init(insertHandlers);
+
+ // Setup rendering.
+ // Rendering will add to the list of insert handlers.
+ render.init(insertHandlers);
+
+ // Initialize gpu subsystem.
+ // This consumes the insert handlers.
+ gpu.init(insertHandlers).then(() => {
+ // Cache initial input state
+ g_prevInputs = ui.getInputs();
+
+ g_loading = false;
+ document.getElementById('ui').style = 'left:-20%';
+ document.getElementById('main').style = '';
+ document.getElementById('loadingContainer').style = 'display:none'
+
+ }).catch((e) => {
+ console.error(`Caught error while initializing WebGPU: [${e}]`)
+ g_loading = false;
+ document.getElementById('loadingText').style = 'display:none'
+ document.getElementById('errorText').innerHTML = `
+
+
+
+
+## Building
+
+PB-MPM is based on WebGPU. Currently, this restricts its use to the Chrome and Edge browsers.
+
+1. [Install Ruby](https://www.ruby-lang.org/en/documentation/installation/)
+2. Run `bundle install`
+3. Run `jekyll serve --livereload`
+4. Open a browser window at [`localhost:4000`](http://localhost:4000)
+
+## Reference
+
+> *Chris Lewin*. **[A Position Based Material Point Method](https://seed.ea.com)**. ACM SIGGRAPH 2024.
+
+## Authors
+
+
+