G2P2G optimization

shaders/bukkit.inc.wgsl

@@ -0,0 +1,23 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts.  All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include simConstants.inc
+
+fn bukkitAddressToIndex(address: vec2u, bukkitCountX: u32) -> u32
+{
+    return address.y*bukkitCountX + address.x;
+}
+
+fn positionToBukkitId(position: vec2f) -> vec2i
+{
+    return vec2i((position) / f32(BukkitSize));
+}
+
+struct BukkitThreadData
+{
+    rangeStart: u32,
+    rangeCount: u32,
+    bukkitX: u32,
+    bukkitY: u32,
+};

shaders/bukkitAllocate.wgsl

@@ -0,0 +1,58 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts.  All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include simConstants.inc
+//!include bukkit.inc
+
+@group(0) @binding(0) var<uniform> g_simConstants : SimConstants;
+@group(0) @binding(1) var<storage> g_bukkitCounts : array<u32>;
+@group(0) @binding(2) var<storage, read_write> g_bukkitIndirectDispatch : array<atomic<u32>>;
+@group(0) @binding(3) var<storage, read_write> g_bukkitThreadData : array<BukkitThreadData>;
+@group(0) @binding(4) var<storage, read_write> g_bukkitParticleAlloctor : array<atomic<u32>>;
+@group(0) @binding(5) var<storage, read_write> g_bukkitIndexStart : array<u32>;
+
+@compute @workgroup_size(GridDispatchSize, GridDispatchSize)
+fn csMain( @builtin(global_invocation_id) id: vec3<u32> )
+{
+    if(id.x >= g_simConstants.bukkitCountX || id.y >= g_simConstants.bukkitCountY)
+    {
+        return;
+    }
+
+    let bukkitIndex = bukkitAddressToIndex(id.xy, g_simConstants.bukkitCountX);
+
+    let bukkitCount = g_bukkitCounts[bukkitIndex];
+    let bukkitCountResidual = bukkitCount % ParticleDispatchSize;
+
+    if(bukkitCount == 0)
+    {
+        return;
+    }
+
+    let dispatchCount = divUp(bukkitCount, ParticleDispatchSize);
+
+    let dispatchStartIndex = atomicAdd(&g_bukkitIndirectDispatch[0], dispatchCount);
+    let particleStartIndex = atomicAdd(&g_bukkitParticleAlloctor[0], bukkitCount);
+
+    g_bukkitIndexStart[bukkitIndex] = particleStartIndex;
+
+    for(var i: u32 = 0; i < dispatchCount; i++)
+    {
+        // Group count is equal to ParticleDispatchSize except for the final dispatch for this
+        // bukkit in which case it's equal to the residual count
+        var groupCount : u32 = ParticleDispatchSize;
+        if(bukkitCountResidual != 0 && i == dispatchCount - 1)
+        {
+            groupCount = bukkitCountResidual;
+        }
+
+        g_bukkitThreadData[i + dispatchStartIndex] = BukkitThreadData(
+            particleStartIndex + i * ParticleDispatchSize,
+            groupCount,
+            id.x,
+            id.y
+        );
+    }
+}

shaders/bukkitCount.wgsl

@@ -0,0 +1,36 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts.  All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include simConstants.inc
+//!include particle.inc
+//!include bukkit.inc
+
+@group(0) @binding(0) var<uniform> g_simConstants : SimConstants;
+@group(0) @binding(1) var<storage> g_particleCount : array<u32>;
+@group(0) @binding(2) var<storage> g_particles : array<Particle>;
+@group(0) @binding(3) var<storage, read_write> g_bukkitCounts : array<atomic<u32>>;
+
+@compute @workgroup_size(ParticleDispatchSize)
+fn csMain( @builtin(global_invocation_id) id: vec3<u32> )
+{
+    if(id.x >= g_particleCount[0])
+    {
+        return;
+    }
+
+    let particle = g_particles[id.x];
+    let position = particle.position;
+
+    let particleBukkit = positionToBukkitId(position);
+
+    if(any(particleBukkit < vec2i(0)) || u32(particleBukkit.x) >= g_simConstants.bukkitCountX || u32(particleBukkit.y) >= g_simConstants.bukkitCountY)
+    {
+        return;
+    }
+
+    let bukkitIndex = bukkitAddressToIndex(vec2u(particleBukkit), g_simConstants.bukkitCountX);
+
+    atomicAdd(&g_bukkitCounts[bukkitIndex], 1);    
+}

shaders/bukkitInsert.wgsl

@@ -0,0 +1,42 @@
+//-----------------------------------------------------------------------------
+// Copyright (c) 2024 Electronic Arts.  All rights reserved.
+//-----------------------------------------------------------------------------
+
+//!include dispatch.inc
+//!include simConstants.inc
+//!include bukkit.inc
+//!include particle.inc
+
+@group(0) @binding(0) var<uniform> g_simConstants : SimConstants;
+@group(0) @binding(1) var<storage> g_particleCount : array<u32>;
+@group(0) @binding(2) var<storage, read_write> g_particleInsertCounters : array<atomic<u32>>;
+@group(0) @binding(3) var<storage> g_particles : array<Particle>;
+@group(0) @binding(4) var<storage, read_write> g_particleData : array<u32>;
+@group(0) @binding(5) var<storage> g_bukkitIndexStart : array<u32>;
+
+@compute @workgroup_size(ParticleDispatchSize)
+fn csMain( @builtin(global_invocation_id) id: vec3<u32> )
+{
+    if(id.x >= g_particleCount[0])
+    {
+        return;
+    }
+
+    let particle = g_particles[id.x];
+    let position = particle.position;
+
+    let particleBukkit = positionToBukkitId(position);
+
+    if(any(particleBukkit < vec2i(0)) || u32(particleBukkit.x) >= g_simConstants.bukkitCountX || u32(particleBukkit.y) >= g_simConstants.bukkitCountY)
+    {
+        return;
+    }
+
+    let bukkitIndex = bukkitAddressToIndex(vec2u(particleBukkit), g_simConstants.bukkitCountX);
+    let bukkitIndexStart = g_bukkitIndexStart[bukkitIndex];
+
+    let particleInsertCounter = atomicAdd(&g_particleInsertCounters[bukkitIndex], 1u);
+
+
+    g_particleData[particleInsertCounter + bukkitIndexStart] = id.x;
+}

shaders/dispatch.inc.wgsl

@@ -3,3 +3,8 @@
 //-----------------------------------------------------------------------------
 
 //!insert DispatchSizes
+
+fn divUp(threadCount : u32, divisor : u32) -> u32
+{
+    return (threadCount + divisor - 1) / divisor;
+}