diff --git a/.gitignore b/.gitignore
index db1fb845..a633d4ca 100644
--- a/.gitignore
+++ b/.gitignore
@@ -405,3 +405,7 @@ res_path.hpp
 
 build_emscripten
 build_analysis
+.aider*
+.env
+
+build_xcode
diff --git a/.vscode/settings.json b/.vscode/settings.json
index 095119ad..00276e15 100644
--- a/.vscode/settings.json
+++ b/.vscode/settings.json
@@ -89,7 +89,10 @@
         "print": "cpp",
         "filesystem": "cpp",
         "forward_list": "cpp",
-        "ranges": "cpp"
+        "ranges": "cpp",
+        "cfenv": "cpp",
+        "regex": "cpp",
+        "valarray": "cpp"
     },
     "files.exclude": {
         "build": true,
diff --git a/OpenGL/res/shaders/particles.wgsl b/OpenGL/res/shaders/particles.wgsl
new file mode 100644
index 00000000..65890a6e
--- /dev/null
+++ b/OpenGL/res/shaders/particles.wgsl
@@ -0,0 +1,46 @@
+struct Particle {
+    @location(1) position: vec2<f32>,  // World space position
+    @location(2) velocity: vec2<f32>,
+    @location(3) color: vec4<f32>,
+};
+
+struct VertexInput {
+    @location(0) position: vec2<f32>,  // Local vertex position (relative to particle center)
+};
+
+// Separate matrices for clearer transform chain
+struct VertexUniforms {
+    world_to_clip: mat4x4<f32>,    // View-Projection matrix only
+};
+
+@group(0) @binding(0)
+var<uniform> vertexUniforms: VertexUniforms;
+
+struct VertexOutput {
+    @builtin(position) Position: vec4<f32>,
+    @location(0) color: vec4<f32>,
+};
+
+struct FragmentOutput {
+    @location(0) color: vec4<f32>,
+};
+
+@vertex
+fn vertex_main(vertex: VertexInput, particle: Particle) -> VertexOutput {
+    var output: VertexOutput;
+    
+    // Transform the local vertex position to world space relative to particle position
+    let world_pos = vertex.position + particle.position;
+    
+    // Transform from world space to clip space using world_to_clip matrix
+    output.Position = vertexUniforms.world_to_clip * vec4<f32>(world_pos, 0.0, 1.0);
+    output.color = particle.color;
+    return output;
+}
+
+@fragment
+fn fragment_main(input: VertexOutput) -> FragmentOutput {
+    var output: FragmentOutput;
+    output.color = input.color;
+    return output;
+}
diff --git a/OpenGL/res/shaders/particlesCompute.wgsl b/OpenGL/res/shaders/particlesCompute.wgsl
new file mode 100644
index 00000000..e44742d5
--- /dev/null
+++ b/OpenGL/res/shaders/particlesCompute.wgsl
@@ -0,0 +1,209 @@
+struct WorldInfo {
+    deltaTime : f32,
+    mousePos  : vec2<f32>,
+};
+
+struct Particle {
+    position : vec2<f32>,  // World space position
+    velocity : vec2<f32>,
+    color    : vec4<f32>,
+};
+
+struct Segment {
+    start : vec2<f32>,
+    end   : vec2<f32>,
+};
+
+struct BvhNode {
+    leftTypeCount : u32,    // leftType :1, leftCount :31
+    leftOffset    : u32,
+
+    rightTypeCount : u32,   // rightType :1, rightCount :31
+    rightOffset    : u32,
+
+    leftBBoxMin  : vec2<f32>,
+    leftBBoxMax  : vec2<f32>,
+
+    rightBBoxMin : vec2<f32>,
+    rightBBoxMax : vec2<f32>,
+};
+
+struct AABB {
+    min : vec2<f32>,
+    max : vec2<f32>,
+};
+
+// Helper functions to unpack BVH Node data
+fn getLeftType(node : BvhNode) -> u32 {
+    return node.leftTypeCount >> 31u;
+}
+
+fn getLeftCount(node : BvhNode) -> u32 {
+    return node.leftTypeCount & 0x7FFFFFFFu;
+}
+
+fn getRightType(node : BvhNode) -> u32 {
+    return node.rightTypeCount >> 31u;
+}
+
+fn getRightCount(node : BvhNode) -> u32 {
+    return node.rightTypeCount & 0x7FFFFFFFu;
+}
+
+fn isLeafNode(node : BvhNode, isLeft : bool) -> bool {
+    if isLeft {
+        return getLeftType(node) == 1u;
+    } else {
+        return getRightType(node) == 1u;
+    }
+}
+
+fn getChildOffset(node : BvhNode, isLeft : bool) -> u32 {
+    if isLeft {
+        return node.leftOffset;
+    } else {
+        return node.rightOffset;
+    }
+}
+
+fn getChildCount(node : BvhNode, isLeft : bool) -> u32 {
+    if isLeft {
+        return getLeftCount(node);
+    } else {
+        return getRightCount(node);
+    }
+}
+
+fn getChildBBox(node : BvhNode, isLeft : bool) -> AABB {
+    if isLeft {
+        return AABB(node.leftBBoxMin, node.leftBBoxMax);
+    } else {
+        return AABB(node.rightBBoxMin, node.rightBBoxMax);
+    }
+}
+
+// Utility functions
+fn cross2D(a : vec2<f32>, b : vec2<f32>) -> f32 {
+    return a.x * b.y - a.y * b.x;
+}
+
+fn dot2D(a : vec2<f32>, b : vec2<f32>) -> f32 {
+    return a.x * b.x + a.y * b.y;
+}
+
+struct Ray {
+    origin    : vec2<f32>,
+    direction : vec2<f32>,
+};
+
+struct SegmentIntersection {
+    hit      : bool,
+    position : vec2<f32>,
+    normal   : vec2<f32>,
+    t        : f32,
+};
+
+// Buffer bindings
+@group(0) @binding(0) var<storage, read_write> particleBuffer : array<Particle>;
+@group(0) @binding(1) var<uniform> world : WorldInfo;
+@group(0) @binding(2) var<storage, read> segments : array<Segment>;
+@group(0) @binding(3) var<storage, read> bvhNodes : array<BvhNode>;
+
+// Constants
+const G : f32 = 30.0; 
+const MIN_DISTANCE_SQUARED : f32 = 1.0; // Prevent division by zero
+
+// Compute shader entry point
+@compute @workgroup_size(256)
+fn compute_main(@builtin(global_invocation_id) id : vec3<u32>) {
+    let index = id.x;
+
+    let particle = &particleBuffer[index];
+
+    // Calculate direction to mouse
+    let toMouse = world.mousePos - particle.position;
+    let distanceSquared = max(dot2D(toMouse, toMouse), MIN_DISTANCE_SQUARED);
+
+    // Calculate gravitational force (F = G * m1 * m2 / r^2)
+    // Since mass is uniform we can simplify
+    let force = normalize(toMouse) * G / distanceSquared;
+
+    // Update velocity (a = F/m, simplified since mass = 1)
+    particleBuffer[index].velocity += force * world.deltaTime;
+
+    // New position based on velocity
+    let newPosition = particle.position + particleBuffer[index].velocity * world.deltaTime;
+
+    // Check for collision with walls
+    let intersection = findWallCollision(particle.position, newPosition);
+    
+    if (intersection.hit) {
+        // Bounce coefficient (1.0 = perfect bounce, 0.0 = full stop)
+        let bounce = 0.8;
+        
+        // Calculate reflection vector
+        let v = particleBuffer[index].velocity;
+        let n = intersection.normal;
+        let reflected = v - 2.0 * dot2D(v, n) * n;
+        
+        // Update velocity with bounce effect
+        let newVelocity = reflected * bounce;
+        particleBuffer[index].velocity = newVelocity;
+        
+        // Place particle at intersection point
+        particleBuffer[index].position = intersection.position + newVelocity * world.deltaTime;
+    } else {
+        // No collision, update particle position normally
+        particleBuffer[index].position = newPosition;
+    }
+}
+
+fn findWallCollision(particlePosition : vec2<f32>, newPosition : vec2<f32>) -> SegmentIntersection {
+    var closest : SegmentIntersection;
+    closest.hit = false;
+    closest.t = 999999.0;
+
+    let particlePath = Segment(particlePosition, newPosition);
+    
+    for (var i: u32 = 0; i < arrayLength(&segments); i++) {
+        let wall = segments[i];
+        let intersection = segmentIntersection(wall, particlePath);
+        
+        if (intersection.hit && intersection.t < closest.t) {
+            closest = intersection;
+        }
+    }
+    
+    return closest;
+}
+
+fn segmentIntersection(s1 : Segment, s2 : Segment) -> SegmentIntersection {
+    var result : SegmentIntersection;
+    result.hit = false;
+    
+    let p = s1.start;
+    let r = s1.end - s1.start;
+    let q = s2.start;
+    let s = s2.end - s2.start;
+
+    let r_cross_s = cross2D(r, s);
+    let q_p = q - p;
+
+    if (abs(r_cross_s) < 1e-8) {
+        return result; // Lines are parallel
+    }
+
+    let t = cross2D(q_p, s) / r_cross_s;
+    let u = cross2D(q_p, r) / r_cross_s;
+
+    if (t >= 0.0 && t <= 1.0 && u >= 0.0 && u <= 1.0) {
+        result.hit = true;
+        result.t = t;
+        result.position = p + t * r;
+        result.normal = normalize(vec2<f32>(-r.y, r.x)); // Perpendicular to wall
+        return result;
+    }
+
+    return result;
+}
+
diff --git a/OpenGL/res/shaders/stars.wgsl b/OpenGL/res/shaders/stars.wgsl
index 8409099c..3e6d5d7f 100644
--- a/OpenGL/res/shaders/stars.wgsl
+++ b/OpenGL/res/shaders/stars.wgsl
@@ -116,7 +116,7 @@ fn starColor(seed: f32) -> vec3<f32> {
 fn main_1() {
     var uv_2: vec2<f32>;
     var aspectRatio: f32;
-    var numStars: i32 = 502i;
+    var numStars: i32 = 2i;
     var finalColor: vec3<f32> = vec3(0f);
     var i: i32 = 0i;
     var seed_2: f32;
diff --git a/OpenGL/src/Application.cpp b/OpenGL/src/Application.cpp
index 5b879af1..6d40c9f6 100644
--- a/OpenGL/src/Application.cpp
+++ b/OpenGL/src/Application.cpp
@@ -39,6 +39,7 @@
 #include "rendering/Texture.h"
 #include "rendering/CommandEncoder.h"
 #include "rendering/RenderPass.h"
+#include "rendering/ComputePass.h"
 #include "AudioEngine.h"
 
 #include "glm/glm.hpp"
@@ -79,18 +80,6 @@ void mainLoop(Application& application, Renderer& renderer) {
       World::settingTimeSpeed = false;
    }
 
-   World::UpdateObjects();
-
-   if (!World::ticksPaused()) {
-      if (World::shouldTick) {
-         World::TickObjects();
-         Input::lastTick   = Input::currentTime;
-         World::shouldTick = false;
-      } else if (Input::lastTick + (1.0 / TICKS_PER_SECOND) <= Input::currentTime) {
-         World::TickObjects();
-         Input::lastTick = Input::lastTick + (1.0 / TICKS_PER_SECOND);
-      }
-   }
 
    auto nextTexture = application.GetNextSurfaceTextureView();
    if (nextTexture) {
@@ -98,33 +87,63 @@ void mainLoop(Application& application, Renderer& renderer) {
       {
          // Create a command encoder for the draw call
          CommandEncoder encoder(device);
+         application.encoder = &encoder.get();
 
-         // Create the render pass
-         RenderPass renderPass(encoder, targetView);
-         renderer.renderPass = renderPass.get();
-
-         // Start the Dear ImGui frame
-         ImGui_ImplWGPU_NewFrame();
-         ImGui_ImplGlfw_NewFrame();
-         ImGui::NewFrame();
+         // CPU game logic
+         {
+            World::UpdateObjects();
+
+            if (!World::ticksPaused()) {
+               if (World::shouldTick) {
+                  World::TickObjects();
+                  Input::lastTick   = Input::currentTime;
+                  World::shouldTick = false;
+               } else if (Input::lastTick + (1.0 / TICKS_PER_SECOND) <= Input::currentTime) {
+                  World::TickObjects();
+                  Input::lastTick = Input::lastTick + (1.0 / TICKS_PER_SECOND);
+               }
+            }
+         }
 
-         World::RenderObjects(renderer, renderPass);
-         renderer.DrawDebug(renderPass);
+         // Pre-compute pass
+         { World::PreComputeObjects(); }
 
-         // Performance info
+         // Compute pass
          {
-            ImGui::PushFont(application.pixelify);
-            ImGui::Begin("Performance Info");
-            ImGui::Text("%.3f ms/frame (%.1f FPS)", 1000.0f / application.getImGuiIO().Framerate,
-                        application.getImGuiIO().Framerate);
-            ImGui::End();
-            ImGui::PopFont();
+            ComputePass computePass(encoder, targetView);
+            World::ComputeObjects(renderer, computePass);
          }
 
-         ImGui::Render();
-         ImGui_ImplWGPU_RenderDrawData(ImGui::GetDrawData(), renderPass.get());
+         // Render pass
+         {
+            RenderPass renderPass(encoder, targetView);
+
+            // Start the Dear ImGui frame
+            ImGui_ImplWGPU_NewFrame();
+            ImGui_ImplGlfw_NewFrame();
+            ImGui::NewFrame();
+
+            World::RenderObjects(renderer, renderPass);
+            renderer.DrawDebug(renderPass);
+
+            // Performance info
+            {
+               ImGui::PushFont(application.pixelify);
+               ImGui::Begin("Performance Info");
+               ImGui::Text("%.3f ms/frame (%.1f FPS)", 1000.0f / application.getImGuiIO().Framerate,
+                           application.getImGuiIO().Framerate);
+               ImGui::End();
+               ImGui::PopFont();
+            }
+
+            ImGui::Render();
+            ImGui_ImplWGPU_RenderDrawData(ImGui::GetDrawData(), renderPass.get());
+
+            // The render pass will be ended and submitted in its destructor
+         }
 
-         // The render pass and command encoder will be ended and submitted in their destructors
+         // The command encoder will be ended and submitted in their destructors
+         application.encoder = nullptr;
       }
       renderer.FinishFrame();
       targetView.release();
@@ -285,7 +304,7 @@ auto getUncapturedErrorCallbackHandle(wgpu::Device& device) {
 #ifdef _MSC_VER
       __debugbreak();
 #else
-      assert(false);
+      std::terminate();
 #endif
    });
 }
@@ -505,14 +524,18 @@ int main(void) {
    Application& application = Application::get();
    Renderer     renderer    = Renderer();
 
-   World::LoadMap("SpaceShip.txt");
-   World::gameobjects.push_back(std::make_unique<Fog>());
-
-   audio().Song.play();
+   {
+      CommandEncoder encoder(application.getDevice());
+      application.encoder = &encoder.get();
+      World::LoadMap("SpaceShip.txt");
+      World::gameobjects.push_back(std::make_unique<Fog>());
+      application.encoder = nullptr;
+   }
 
    Input::currentTime       = glfwGetTime();
    Input::realTimeLastFrame = Input::currentTime;
    Input::lastTick          = Input::currentTime;
+
    // audio().Song.play();
 
    // Not Emscripten-friendly
diff --git a/OpenGL/src/Application.h b/OpenGL/src/Application.h
index 735d8a01..f5719485 100644
--- a/OpenGL/src/Application.h
+++ b/OpenGL/src/Application.h
@@ -46,6 +46,8 @@ class Application {
    int currentDecals = 0;
    int maxDecals     = 50;
 
+   wgpu::CommandEncoder* encoder;
+
 private:
    GLFWwindow*                          window;
    wgpu::Instance                       instance;
diff --git a/OpenGL/src/World.cpp b/OpenGL/src/World.cpp
index ef53da15..436b68b6 100644
--- a/OpenGL/src/World.cpp
+++ b/OpenGL/src/World.cpp
@@ -8,6 +8,7 @@
 #include "game_objects/Player.h"
 #include "game_objects/Background.h"
 #include "game_objects/Camera.h"
+#include "game_objects/Particles.h"
 #include "game_objects/Tile.h"
 #include "game_objects/enemies/Bomber.h"
 #include "game_objects/enemies/Turret.h"
@@ -53,6 +54,7 @@ void World::LoadMap(const std::filesystem::path& map_path) {
             if (c == 'p') { // player
                gameobjects.push_back(std::make_shared<Player>("Coolbox", (float)x, (float)y));
                gameobjects.push_back(std::make_shared<Tile>("Floor", (float)x, (float)y));
+               gameobjects.push_back(std::make_shared<Particles>("Floor", DrawPriority::Character, glm::vec2(x, y)));
             }
             if (c == 'f') { // floor
                gameobjects.push_back(std::make_shared<Tile>("Floor", (float)x, (float)y));
@@ -135,6 +137,24 @@ void World::RenderObjects(Renderer& renderer, RenderPass& renderPass) {
    }
 }
 
+void World::ComputeObjects(Renderer& renderer, ComputePass& computePass) {
+   auto objects = get_gameobjects();
+   sortGameObjectsByPriority(objects);
+
+   for (auto& gameobject : objects) {
+      gameobject->compute(renderer, computePass);
+   }
+}
+
+void World::PreComputeObjects() {
+   auto objects = get_gameobjects();
+   sortGameObjectsByPriority(objects);
+
+   for (auto& gameobject : objects) {
+      gameobject->pre_compute();
+   }
+}
+
 bool World::ticksPaused() {
    auto player = getFirst<Player>();
    return player->pauseTicks();
diff --git a/OpenGL/src/World.h b/OpenGL/src/World.h
index 8d380ca9..b4a527a6 100644
--- a/OpenGL/src/World.h
+++ b/OpenGL/src/World.h
@@ -18,17 +18,26 @@ class World {
 
    static bool ticksPaused();
 
+   static void add_gameobject_recursive(GameObject* parent, GameObject* obj, std::vector<GameObject*>& allGameObjects) {
+      if (!obj)
+         return;
+
+      obj->parent = parent; // Set parent (null for top-level objects)
+      allGameObjects.push_back(obj);
+
+      auto children = obj->children();
+      for (auto& child : children) {
+         if (child) {
+            add_gameobject_recursive(obj, child, allGameObjects);
+         }
+      }
+   }
+
    static std::vector<GameObject*> get_gameobjects() {
-      // return a vector of all gameobjects and their gameobjects.children
+      // return a vector of all gameobjects and their gameobjects.children recursively
       std::vector<GameObject*> allGameObjects;
       for (auto& gameobject : gameobjects) {
-         allGameObjects.push_back(gameobject.get());
-         auto children = gameobject->children();
-         for (auto& child : children) {
-            if (child) {
-               allGameObjects.push_back(child);
-            }
-         }
+         add_gameobject_recursive(nullptr, gameobject.get(), allGameObjects);
       }
       return allGameObjects;
    }
@@ -78,5 +87,7 @@ class World {
    static void UpdateObjects();
    static void TickObjects();
    static void RenderObjects(Renderer& renderer, RenderPass& renderPass);
+   static void ComputeObjects(Renderer& renderer, ComputePass& computePass);
+   static void PreComputeObjects();
    static bool shouldTick;
 };
diff --git a/OpenGL/src/game_objects/Fog.cpp b/OpenGL/src/game_objects/Fog.cpp
index 5bd6e7c8..d91ca561 100644
--- a/OpenGL/src/game_objects/Fog.cpp
+++ b/OpenGL/src/game_objects/Fog.cpp
@@ -3,7 +3,8 @@
 #include "Player.h"
 #include <vector>
 #include "clipper2/clipper.h"
-#include "../GeometryUtils.h"
+#include "../geometry/GeometryUtils.h"
+#include "../geometry/SceneGeometry.h"
 #include "earcut.hpp"
 
 using namespace Clipper2Lib;
@@ -11,68 +12,27 @@ using namespace GeometryUtils;
 
 Fog::Fog()
    : GameObject("Fog of War", DrawPriority::Fog, {0, 0})
-   , vertexUniform(
-        UniformBuffer<FogVertexUniform>({FogVertexUniform(CalculateMVP(position, rotation, scale))},
-                                        wgpu::bothBufferUsages(wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Uniform)))
+   , vertexUniform(UniformBuffer<FogVertexUniform>(
+        {FogVertexUniform(MVP())}, wgpu::bothBufferUsages(wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Uniform)))
    , fragmentUniformWalls(
         UniformBufferView<FogFragmentUniform>::create(FogFragmentUniform(mainFogColor, tintFogColor, {0, 0})))
    , fragmentUniformOther(
         UniformBufferView<FogFragmentUniform>::create(FogFragmentUniform(mainFogColor, mainFogColor, {0, 0}))) {}
 
 void Fog::render(Renderer& renderer, RenderPass& renderPass) {
-   bool showWalls = true;
-   vertexUniform.upload({FogVertexUniform(CalculateMVP(position, rotation, scale))});
-
-   // Collect all tile bounds
-   std::vector<std::vector<glm::vec2>> allBounds;
-   auto                                tiles = World::getAll<Tile>(); // Simplified retrieval of all tiles
-   for (auto tile : tiles) {
-      allBounds.push_back(tile->getBounds());
-   }
+   vertexUniform.upload({FogVertexUniform(MVP())});
 
    // Get the player
    auto player = World::getFirst<Player>(); // Simplified retrieval of the first player
    fragmentUniformWalls.Update(FogFragmentUniform(mainFogColor, tintFogColor, player->position));
    fragmentUniformOther.Update(FogFragmentUniform(mainFogColor, mainFogColor, player->position));
 
-   // Compute the union of all tile bounds
-   PolyTreeD combined;
-   findPolygonUnion(allBounds, combined);
-   auto flattened = FlattenPolyPathD(combined);
-
-   // Prepare the hull for clipping
-   ClipperD clipper;
-   PathsD   hullPaths;
-   for (auto& child : combined) {
-      hullPaths.push_back(child->Polygon());
-   }
-   clipper.AddSubject(hullPaths);
-
-   // Compute the visibility polygon
-   auto visibility = ComputeVisibilityPolygon(player->position, flattened);
-
-   // Compute the areas occluded
-   clipper.AddClip({visibility});
-   if (showWalls) {
-      clipper.AddClip({flattened});
-   }
-   // Compute the difference to get invisibility regions
-   PolyTreeD invisibilityPaths;
-   clipper.Execute(ClipType::Difference, FillRule::NonZero, invisibilityPaths);
+   auto walls      = SceneGeometry::computeWallPaths();
+   auto visibility = SceneGeometry::computeVisibility(walls, player->position);
 
    // Render the invisibility regions
-   renderPolyTree(renderer, renderPass, invisibilityPaths, fragmentUniformOther);
-
-   if (showWalls) {
-      // Tint all the walls that are not visible
-      ClipperD tint;
-      tint.AddSubject({flattened});
-      tint.AddClip({visibility});
-      PolyTreeD tintPaths;
-      tint.Execute(ClipType::Difference, FillRule::NonZero, tintPaths);
-
-      renderPolyTree(renderer, renderPass, tintPaths, fragmentUniformWalls);
-   }
+   renderPolyTree(renderer, renderPass, *visibility.invisibilityPaths, fragmentUniformOther);
+   renderPolyTree(renderer, renderPass, *walls.wallPaths, fragmentUniformWalls);
 }
 
 void Fog::update() {}
diff --git a/OpenGL/src/game_objects/GameObject.cpp b/OpenGL/src/game_objects/GameObject.cpp
index 3e2ddb39..296df73e 100644
--- a/OpenGL/src/game_objects/GameObject.cpp
+++ b/OpenGL/src/game_objects/GameObject.cpp
@@ -19,3 +19,29 @@ void GameObject::tickUpdate() {}
 void GameObject::render(Renderer& renderer, RenderPass& renderPass) {
    std::cout << "Rendering GameObject (you should not see this lol) " << name << std::endl;
 }
+
+void GameObject::pre_compute() {}
+void GameObject::compute(Renderer& renderer, ComputePass& computePass) {}
+
+glm::mat4 GameObject::getLocalTransform() const {
+   return CalculateModel(position, rotation, scale);
+}
+
+glm::mat4 GameObject::MVP() const {
+   glm::mat4 localTransform = getLocalTransform();
+
+   if (parent) {
+      // Get parent's transform and combine with local transform
+      glm::mat4 parentMVP = parent->MVP();
+      return parentMVP * localTransform;
+   }
+
+   // No parent, just apply view and projection to local transform
+   glm::mat4 view       = CalculateView();
+   glm::mat4 projection = CalculateProjection();
+   return projection * view * localTransform;
+}
+
+glm::mat4 GameObject::VP() const {
+   return CalculateProjection() * CalculateView();
+}
diff --git a/OpenGL/src/game_objects/GameObject.h b/OpenGL/src/game_objects/GameObject.h
index 657db17a..731f3363 100644
--- a/OpenGL/src/game_objects/GameObject.h
+++ b/OpenGL/src/game_objects/GameObject.h
@@ -10,6 +10,7 @@
 
 #include "../rendering/Renderer.h"
 #include "../rendering/RenderPass.h"
+#include "../rendering/ComputePass.h"
 
 enum class DrawPriority {
    Background,
@@ -33,6 +34,8 @@ class GameObject {
    virtual std::vector<GameObject*> children() { return {}; }
 
    virtual void render(Renderer& renderer, RenderPass& renderPass);
+   virtual void pre_compute();
+   virtual void compute(Renderer& renderer, ComputePass& computePass);
    virtual void update();
    virtual void tickUpdate();
 
@@ -41,6 +44,16 @@ class GameObject {
    glm::vec2    position;
    float        rotation = 0;
    float        scale    = 1.0f;
+   GameObject*  parent   = nullptr;
+
+   // Get this object's local transform matrix
+   glm::mat4 getLocalTransform() const;
+
+   // Get the Model-View-Projection matrix for this object
+   glm::mat4 MVP() const;
+
+   // Get the View-Projection matrix for this object
+   glm::mat4 VP() const;
 
    // Add coroutine
    void addCoroutine(Generator coroutine) { coroutines.emplace_back(std::move(coroutine)); }
diff --git a/OpenGL/src/game_objects/Particles.cpp b/OpenGL/src/game_objects/Particles.cpp
new file mode 100644
index 00000000..419ccbd3
--- /dev/null
+++ b/OpenGL/src/game_objects/Particles.cpp
@@ -0,0 +1,86 @@
+#include "Particles.h"
+#include "../Input.h"
+#include "../geometry/SceneGeometry.h"
+
+#include <random>
+
+Particles::Particles(const std::string& name, DrawPriority drawPriority, glm::vec2 position)
+   : GameObject(name, drawPriority, position)
+   , particles(std::vector<Particle>{
+        Particle{position + glm::vec2(0.0f, 0.0f), glm::vec2(1.0f, 0.0f), glm::vec4(1.0f, 1.0f, 0.0f, 1.0f)},
+        Particle{position + glm::vec2(0.0f, 0.0f), glm::vec2(0.0f, 1.0f), glm::vec4(0.0f, 1.0f, 0.0f, 1.0f)}
+}),
+   particleBuffer(
+      std::make_shared<Buffer<Particle>>(particles,
+                                         wgpu::bothBufferUsages(wgpu::BufferUsage::Vertex, wgpu::BufferUsage::CopyDst,
+                                                                wgpu::BufferUsage::CopySrc, wgpu::BufferUsage::Storage),
+                                         "Particles")),
+   pointBuffer(Buffer<ParticleVertex>::create(
+      {
+         ParticleVertex{glm::vec2(-0.5f, -0.5f) * (1.0f / 16.0f)}, // 0
+         ParticleVertex{glm::vec2(0.5f, -0.5f) * (1.0f / 16.0f)},  // 1
+         ParticleVertex{glm::vec2(0.5f, 0.5f) * (1.0f / 16.0f)},   // 2
+         ParticleVertex{glm::vec2(-0.5f, 0.5f) * (1.0f / 16.0f)},  // 3
+      },
+      wgpu::bothBufferUsages(wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Vertex))),
+   indexBuffer(IndexBuffer::create(
+      {
+         0, 1, 2, // Triangle #0 connects points #0, #1 and #2
+         0, 2, 3  // Triangle #1 connects points #0, #2 and #3
+      },
+      wgpu::bothBufferUsages(wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Index))),
+   segmentBuffer(Buffer<Segment>(
+      {}, wgpu::bothBufferUsages(wgpu::BufferUsage::CopySrc, wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Storage),
+      "segments")),
+   bvhBuffer(Buffer<BvhNode>(
+      {}, wgpu::bothBufferUsages(wgpu::BufferUsage::CopySrc, wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Storage),
+      "bvh")),
+   vertexUniform(UniformBufferView<ParticleVertexUniform>::create(ParticleVertexUniform{VP()})),
+   worldInfo(UniformBufferView<ParticleWorldInfo>::create(ParticleWorldInfo(0.0f, glm::vec2(0.0f)))) {}
+
+void Particles::render(Renderer& renderer, RenderPass& renderPass) {
+   if (particles.empty())
+      return;
+
+   // Update VP matrix
+   this->vertexUniform.Update(ParticleVertexUniform{VP()});
+
+   // Create bind group and draw
+   BindGroup bindGroup = ParticleLayout::ToBindGroup(renderer.device, vertexUniform);
+   renderPass.DrawInstanced(renderer.particles, *indexBuffer, bindGroup, {(uint32_t)vertexUniform.getOffset()},
+                            particles.size(), *pointBuffer, *particleBuffer);
+}
+
+void Particles::pre_compute() {
+   auto walls = SceneGeometry::computeWallPaths();
+   bvhBuffer.upload(walls.bvh.nodes);
+   segmentBuffer.upload(walls.bvh.segments);
+}
+
+void Particles::compute(Renderer& renderer, ComputePass& computePass) {
+   worldInfo.Update(ParticleWorldInfo(Input::deltaTime, Renderer::MousePos()));
+   BindGroup bindGroup =
+      ParticleComputeLayout::ToBindGroup(renderer.device, std::forward_as_tuple(*particleBuffer, 0), worldInfo,
+                                         std::forward_as_tuple(segmentBuffer, 0), std::forward_as_tuple(bvhBuffer, 0));
+   computePass.dispatch(renderer.particlesCompute, bindGroup, {(uint32_t)worldInfo.getOffset()}, particles.size());
+}
+
+void Particles::update() {
+   // add a particle if the p key is pressed
+   if (Input::keys_pressed[GLFW_KEY_P]) {
+      std::random_device rd;
+      std::mt19937       gen(rd()); // Mersenne Twister generator
+
+      // Define distribution from 0 to 1
+      std::uniform_real_distribution<> vel_dist(-0.5, 0.5);
+      std::uniform_real_distribution<> color_dist(0.0, 1.0);
+
+      auto random_vel = glm::vec2(vel_dist(gen), vel_dist(gen)) * 2.0f;
+      addParticle(position, random_vel, glm::vec4(color_dist(gen), color_dist(gen), color_dist(gen), 1.0f));
+   }
+}
+
+void Particles::addParticle(const glm::vec2& pos, const glm::vec2& vel, const glm::vec4& color) {
+   particles.push_back({pos, vel, color});
+   particleViews.push_back(particleBuffer->Add(particles.back()));
+}
diff --git a/OpenGL/src/game_objects/Particles.h b/OpenGL/src/game_objects/Particles.h
new file mode 100644
index 00000000..9a8e708e
--- /dev/null
+++ b/OpenGL/src/game_objects/Particles.h
@@ -0,0 +1,30 @@
+#pragma once
+#include "GameObject.h"
+#include "../geometry/BVH.h"
+#include "../rendering/Renderer.h"
+#include "../rendering/Texture.h"
+#include <glm/glm.hpp>
+
+class Particles : public GameObject {
+public:
+   Particles(const std::string& name, DrawPriority drawPriority, glm::vec2 position);
+   virtual void render(Renderer& renderer, RenderPass& renderPass) override;
+   virtual void update() override;
+   virtual void pre_compute() override;
+   virtual void compute(Renderer& renderer, ComputePass& computePass) override;
+   void         addParticle(const glm::vec2& pos, const glm::vec2& vel, const glm::vec4& color);
+
+private:
+   std::vector<Particle>                    particles;
+   std::vector<BufferView<Particle, false>> particleViews;
+   std::shared_ptr<Buffer<Particle>>        particleBuffer;
+   std::shared_ptr<Buffer<ParticleVertex>>  pointBuffer;
+   std::shared_ptr<IndexBuffer>             indexBuffer;
+   Buffer<Segment>                          segmentBuffer;
+   Buffer<BvhNode>                          bvhBuffer;
+   UniformBufferView<ParticleVertexUniform> vertexUniform;
+   UniformBufferView<ParticleWorldInfo>     worldInfo;
+
+private:
+protected:
+};
diff --git a/OpenGL/src/game_objects/SquareObject.cpp b/OpenGL/src/game_objects/SquareObject.cpp
index accc6d68..be1909b3 100644
--- a/OpenGL/src/game_objects/SquareObject.cpp
+++ b/OpenGL/src/game_objects/SquareObject.cpp
@@ -23,14 +23,13 @@ SquareObject::SquareObject(const std::string& name, DrawPriority drawPriority, i
          0, 2, 3  // Triangle #1 connects points #0, #2 and #3
       },
       wgpu::bothBufferUsages(wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Index))),
-   vertexUniform(BufferView<SquareObjectVertexUniform>::create(
-      SquareObjectVertexUniform{CalculateMVP(position, rotation, scale)})),
-   fragmentUniform(
-      BufferView<SquareObjectFragmentUniform>::create(SquareObjectFragmentUniform{glm::vec4(0.0f, 0.0f, 0.0f, 0.0f)})),
+   vertexUniform(UniformBufferView<SquareObjectVertexUniform>::create(SquareObjectVertexUniform{MVP()})),
+   fragmentUniform(UniformBufferView<SquareObjectFragmentUniform>::create(
+      SquareObjectFragmentUniform{glm::vec4(0.0f, 0.0f, 0.0f, 0.0f)})),
    texture(Texture::create(texturePath)) {}
 
 void SquareObject::render(Renderer& renderer, RenderPass& renderPass) {
-   this->vertexUniform.Update(SquareObjectVertexUniform{CalculateMVP(position, rotation, scale)});
+   this->vertexUniform.Update(SquareObjectVertexUniform{MVP()});
    this->fragmentUniform.Update(SquareObjectFragmentUniform{tintColor});
    BindGroup bindGroup =
       SquareObjectLayout::ToBindGroup(renderer.device, vertexUniform, fragmentUniform, texture.get(), renderer.sampler);
diff --git a/OpenGL/src/game_objects/enemies/Turret.cpp b/OpenGL/src/game_objects/enemies/Turret.cpp
index f75739eb..3884d399 100644
--- a/OpenGL/src/game_objects/enemies/Turret.cpp
+++ b/OpenGL/src/game_objects/enemies/Turret.cpp
@@ -5,7 +5,7 @@ Turret::Turret(const std::string& name, float x, float y)
    : Character(name, x, y, "turret_base.png") {
    drawPriority = DrawPriority::Character;
    health       = 1;
-   turretHead   = std::make_unique<TurretHead>("Turret Head", x, y);
+   turretHead   = std::make_unique<TurretHead>("Turret Head", 0, 0);
 }
 
 void Turret::update() {
@@ -14,7 +14,6 @@ void Turret::update() {
    if (health <= 0) {
       ShouldDestroy = true;
    }
-   turretHead->position = position;
    if (aimDirection_x != 0) {
       turretHead->rotation = zeno(turretHead->rotation, aimDirection_x > 0 ? 0 : 180, 0.05);
    } else if (aimDirection_y != 0) {
diff --git a/OpenGL/src/geometry/BVH.cpp b/OpenGL/src/geometry/BVH.cpp
new file mode 100644
index 00000000..6444811a
--- /dev/null
+++ b/OpenGL/src/geometry/BVH.cpp
@@ -0,0 +1,421 @@
+#include "BVH.h"
+#include <algorithm>
+#include <cmath>
+
+const float  EPSILON       = 1e-8f;
+const size_t MAX_LEAF_SIZE = 1; // Adjust as needed
+
+float cross(const glm::vec2& v, const glm::vec2& w) {
+   return v.x * w.y - v.y * w.x;
+}
+
+BVH BVH::build(std::vector<Segment> segments) {
+   BVH bvh;
+   if (segments.empty()) {
+      bvh.segments = std::move(segments);
+      return bvh;
+   }
+   bvh.segments = std::move(segments);
+   // Build the tree recursively starting with all segments
+   bvh.build_recursive(0, bvh.segments.size());
+   return bvh;
+}
+
+glm::vec4 BVH::getBoundingBoxOfNode(const BvhNode& node) const {
+   float min_x = node.leftBBox.x;
+   float min_y = node.leftBBox.y;
+   float max_x = node.leftBBox.z;
+   float max_y = node.leftBBox.w;
+
+   // Check if right child has a valid bounding box
+   if (node.rightBBox != glm::vec4(0.0f)) {
+      min_x = std::min(min_x, node.rightBBox.x);
+      min_y = std::min(min_y, node.rightBBox.y);
+      max_x = std::max(max_x, node.rightBBox.z);
+      max_y = std::max(max_y, node.rightBBox.w);
+   }
+
+   return glm::vec4(min_x, min_y, max_x, max_y);
+}
+
+size_t BVH::build_recursive(size_t segment_start, size_t segment_end) {
+   // Calculate AABB for current node
+   AABB aabb = compute_aabb(segment_start, segment_end);
+
+   // Create leaf node if we have few enough segments
+   if (segment_end - segment_start <= MAX_LEAF_SIZE) {
+      BvhNode node;
+
+      // Leaf node - set left child as line bucket
+      uint32_t leftType  = 1; // 1 indicates a line bucket
+      uint32_t leftCount = static_cast<uint32_t>(segment_end - segment_start);
+      node.setLeft(leftType, leftCount);
+      node.leftOffset = static_cast<uint32_t>(segment_start); // Index into segments array
+
+      // Set leftBBox
+      node.leftBBox = glm::vec4(aabb.min, aabb.max);
+
+      // No right child for leaf node
+      node.setRight(0, 0); // Right child is unused
+      node.rightOffset = 0;
+      node.rightBBox   = glm::vec4(0.0f); // Empty bounding box
+
+      size_t node_idx = nodes.size();
+      nodes.push_back(node);
+      return node_idx;
+   }
+
+   // Partition segments
+   size_t mid = partition_segments(segment_start, segment_end);
+
+   // Recursively build left and right children
+   size_t leftChildIndex  = build_recursive(segment_start, mid);
+   size_t rightChildIndex = build_recursive(mid, segment_end);
+
+   // Compute bounding boxes of children
+   glm::vec4 leftChildBBox  = getBoundingBoxOfNode(nodes[leftChildIndex]);
+   glm::vec4 rightChildBBox = getBoundingBoxOfNode(nodes[rightChildIndex]);
+
+   // Create internal node
+   BvhNode node;
+
+   // Left child
+   node.setLeft(0, 0); // 0 indicates a node
+   node.leftOffset = static_cast<uint32_t>(leftChildIndex);
+   node.leftBBox   = leftChildBBox;
+
+   // Right child
+   node.setRight(0, 0); // 0 indicates a node
+   node.rightOffset = static_cast<uint32_t>(rightChildIndex);
+   node.rightBBox   = rightChildBBox;
+
+   size_t node_idx = nodes.size();
+   nodes.push_back(node);
+   return node_idx;
+}
+
+AABB BVH::compute_aabb(size_t segment_start, size_t segment_end) const {
+   float min_x = std::numeric_limits<float>::infinity();
+   float min_y = std::numeric_limits<float>::infinity();
+   float max_x = -std::numeric_limits<float>::infinity();
+   float max_y = -std::numeric_limits<float>::infinity();
+
+   for (size_t i = segment_start; i < segment_end; ++i) {
+      const Segment& segment = segments[i];
+      min_x                  = std::min({min_x, segment.start.x, segment.end.x});
+      min_y                  = std::min({min_y, segment.start.y, segment.end.y});
+      max_x                  = std::max({max_x, segment.start.x, segment.end.x});
+      max_y                  = std::max({max_y, segment.start.y, segment.end.y});
+   }
+
+   return AABB{
+      glm::vec2{min_x, min_y},
+      glm::vec2{max_x, max_y}
+   };
+}
+
+size_t BVH::partition_segments(size_t segment_start, size_t segment_end) {
+   AABB  aabb   = compute_aabb(segment_start, segment_end);
+   float width  = aabb.max.x - aabb.min.x;
+   float height = aabb.max.y - aabb.min.y;
+
+   // Choose axis with the largest spread
+   bool axis_vertical = height > width;
+
+   // Lambda to compute the centroid of a segment
+   auto centroid = [axis_vertical](const Segment& segment) {
+      if (axis_vertical) {
+         return (segment.start.y + segment.end.y) / 2.0f;
+      } else {
+         return (segment.start.x + segment.end.x) / 2.0f;
+      }
+   };
+
+   // Get mutable iterator range for the segments
+   auto segments_begin  = segments.begin() + segment_start;
+   auto segments_end_it = segments.begin() + segment_end;
+
+   // Sort the segments in place based on their centroids
+   std::sort(segments_begin, segments_end_it,
+             [centroid](const Segment& a, const Segment& b) { return centroid(a) < centroid(b); });
+
+   // Return the midpoint index for splitting
+   return (segment_start + segment_end) / 2;
+}
+
+std::optional<std::pair<glm::vec2, const Segment*>> BVH::ray_intersect(const Ray& ray) const {
+   if (nodes.empty()) {
+      return std::nullopt;
+   }
+   auto result = ray_intersect_node(ray, nodes.size() - 1, std::nullopt); // Start from root node
+   if (result.has_value()) {
+      auto [p, t, segment_ptr] = result.value();
+      return std::make_optional(std::make_pair(p, segment_ptr));
+   } else {
+      return std::nullopt;
+   }
+}
+
+std::optional<std::tuple<glm::vec2, float, const Segment*>>
+BVH::ray_intersect_node(const Ray& ray, size_t node_idx,
+                        std::optional<std::tuple<glm::vec2, float, const Segment*>> closest) const {
+   const BvhNode& node = nodes[node_idx];
+
+   // Check intersection with left child bounding box
+   AABB leftAABB{glm::vec2(node.leftBBox.x, node.leftBBox.y), glm::vec2(node.leftBBox.z, node.leftBBox.w)};
+   auto left_aabb_hit = intersect_ray_aabb(ray, leftAABB);
+
+   if (left_aabb_hit.has_value()) {
+      uint32_t leftType = node.getLeftType();
+      if (leftType == 0) {
+         // Left child is a node
+         size_t childIdx = node.leftOffset;
+         closest         = ray_intersect_node(ray, childIdx, closest);
+      } else {
+         // Left child is a line bucket (leaf node)
+         uint32_t count  = node.getLeftCount();
+         uint32_t offset = node.leftOffset;
+         for (uint32_t i = 0; i < count; ++i) {
+            const Segment& segment      = segments[offset + i];
+            auto           intersection = intersect_ray_segment(ray, segment);
+            if (intersection.has_value()) {
+               glm::vec2 p = intersection->first;
+               float     t = intersection->second;
+               if (t >= 0.0f) {
+                  if (!closest.has_value() || t < std::get<1>(closest.value())) {
+                     closest = std::make_optional(std::make_tuple(p, t, &segment));
+                  }
+               }
+            }
+         }
+      }
+   }
+
+   // Check intersection with right child bounding box
+   if (node.rightBBox != glm::vec4(0.0f)) { // Check if right child exists
+      AABB rightAABB{glm::vec2(node.rightBBox.x, node.rightBBox.y), glm::vec2(node.rightBBox.z, node.rightBBox.w)};
+      auto right_aabb_hit = intersect_ray_aabb(ray, rightAABB);
+
+      if (right_aabb_hit.has_value()) {
+         uint32_t rightType = node.getRightType();
+         if (rightType == 0) {
+            // Right child is a node
+            size_t childIdx = node.rightOffset;
+            closest         = ray_intersect_node(ray, childIdx, closest);
+         } else {
+            // Right child is a line bucket (leaf node)
+            uint32_t count  = node.getRightCount();
+            uint32_t offset = node.rightOffset;
+            for (uint32_t i = 0; i < count; ++i) {
+               const Segment& segment      = segments[offset + i];
+               auto           intersection = intersect_ray_segment(ray, segment);
+               if (intersection.has_value()) {
+                  glm::vec2 p = intersection->first;
+                  float     t = intersection->second;
+                  if (t >= 0.0f) {
+                     if (!closest.has_value() || t < std::get<1>(closest.value())) {
+                        closest = std::make_optional(std::make_tuple(p, t, &segment));
+                     }
+                  }
+               }
+            }
+         }
+      }
+   }
+
+   return closest;
+}
+
+std::optional<std::pair<glm::vec2, const Segment*>> BVH::segment_intersect(const Segment& segment) const {
+   if (nodes.empty()) {
+      return std::nullopt;
+   }
+   auto result = segment_intersect_node(segment, nodes.size() - 1, std::nullopt); // Start from root node
+   if (result.has_value()) {
+      auto [p, t, segment_ptr] = result.value();
+      return std::make_optional(std::make_pair(p, segment_ptr));
+   } else {
+      return std::nullopt;
+   }
+}
+
+std::optional<std::tuple<glm::vec2, float, const Segment*>>
+BVH::segment_intersect_node(const Segment& segment, size_t node_idx,
+                            std::optional<std::tuple<glm::vec2, float, const Segment*>> closest) const {
+   const BvhNode& node = nodes[node_idx];
+
+   // Check intersection with left child bounding box
+   AABB leftAABB{glm::vec2(node.leftBBox.x, node.leftBBox.y), glm::vec2(node.leftBBox.z, node.leftBBox.w)};
+   if (intersect_segment_aabb(segment, leftAABB)) {
+      uint32_t leftType = node.getLeftType();
+      if (leftType == 0) {
+         // Left child is a node
+         size_t childIdx = node.leftOffset;
+         closest         = segment_intersect_node(segment, childIdx, closest);
+      } else {
+         // Left child is a line bucket (leaf node)
+         uint32_t count  = node.getLeftCount();
+         uint32_t offset = node.leftOffset;
+         for (uint32_t i = 0; i < count; ++i) {
+            const Segment& other_segment = segments[offset + i];
+            auto           intersection  = intersect_segment_segment(segment, other_segment);
+            if (intersection.has_value()) {
+               glm::vec2 p = intersection->first;
+               float     t = intersection->second;
+               if (t >= 0.0f && t <= 1.0f) {
+                  if (!closest.has_value() || t < std::get<1>(closest.value())) {
+                     closest = std::make_optional(std::make_tuple(p, t, &other_segment));
+                  }
+               }
+            }
+         }
+      }
+   }
+
+   // Check intersection with right child bounding box
+   if (node.rightBBox != glm::vec4(0.0f)) { // Check if right child exists
+      AABB rightAABB{glm::vec2(node.rightBBox.x, node.rightBBox.y), glm::vec2(node.rightBBox.z, node.rightBBox.w)};
+      if (intersect_segment_aabb(segment, rightAABB)) {
+         uint32_t rightType = node.getRightType();
+         if (rightType == 0) {
+            // Right child is a node
+            size_t childIdx = node.rightOffset;
+            closest         = segment_intersect_node(segment, childIdx, closest);
+         } else {
+            // Right child is a line bucket (leaf node)
+            uint32_t count  = node.getRightCount();
+            uint32_t offset = node.rightOffset;
+            for (uint32_t i = 0; i < count; ++i) {
+               const Segment& other_segment = segments[offset + i];
+               auto           intersection  = intersect_segment_segment(segment, other_segment);
+               if (intersection.has_value()) {
+                  glm::vec2 p = intersection->first;
+                  float     t = intersection->second;
+                  if (t >= 0.0f && t <= 1.0f) {
+                     if (!closest.has_value() || t < std::get<1>(closest.value())) {
+                        closest = std::make_optional(std::make_tuple(p, t, &other_segment));
+                     }
+                  }
+               }
+            }
+         }
+      }
+   }
+
+   return closest;
+}
+
+// Utility functions for intersection tests
+
+std::optional<std::pair<glm::vec2, float>> intersect_ray_segment(const Ray& ray, const Segment& segment) {
+   const glm::vec2& p = ray.origin;
+   const glm::vec2& r = ray.direction;
+   const glm::vec2& q = segment.start;
+   glm::vec2        s = segment.end - segment.start;
+
+   float r_cross_s = cross(r, s);
+   if (std::abs(r_cross_s) < EPSILON) {
+      // Lines are parallel
+      return std::nullopt;
+   }
+
+   glm::vec2 q_minus_p = q - p;
+
+   float t = cross(q_minus_p, s) / r_cross_s;
+   float u = cross(q_minus_p, r) / r_cross_s;
+
+   if (t >= 0.0f && u >= 0.0f && u <= 1.0f) {
+      // Intersection occurs at p + t * r
+      glm::vec2 intersection_point = p + t * r;
+      return std::make_optional(std::make_pair(intersection_point, t));
+   } else {
+      return std::nullopt;
+   }
+}
+
+std::optional<std::pair<glm::vec2, float>> intersect_segment_segment(const Segment& s1, const Segment& s2) {
+   const glm::vec2& p = s1.start;
+   glm::vec2        r = s1.end - s1.start;
+   const glm::vec2& q = s2.start;
+   glm::vec2        s = s2.end - s2.start;
+
+   float r_cross_s = cross(r, s);
+   if (std::abs(r_cross_s) < EPSILON) {
+      // Lines are parallel
+      return std::nullopt;
+   }
+
+   glm::vec2 q_minus_p = q - p;
+
+   float t = cross(q_minus_p, s) / r_cross_s;
+   float u = cross(q_minus_p, r) / r_cross_s;
+
+   if (t >= 0.0f && t <= 1.0f && u >= 0.0f && u <= 1.0f) {
+      // Intersection occurs at p + t * r
+      glm::vec2 intersection_point = p + t * r;
+      return std::make_optional(std::make_pair(intersection_point, t));
+   } else {
+      return std::nullopt;
+   }
+}
+
+std::optional<std::pair<float, float>> intersect_ray_aabb(const Ray& ray, const AABB& aabb) {
+   float tmin = (aabb.min.x - ray.origin.x) / ray.direction.x;
+   float tmax = (aabb.max.x - ray.origin.x) / ray.direction.x;
+
+   if (tmin > tmax)
+      std::swap(tmin, tmax);
+
+   float tymin = (aabb.min.y - ray.origin.y) / ray.direction.y;
+   float tymax = (aabb.max.y - ray.origin.y) / ray.direction.y;
+
+   if (tymin > tymax)
+      std::swap(tymin, tymax);
+
+   if ((tmin > tymax) || (tymin > tmax))
+      return std::nullopt;
+
+   if (tymin > tmin)
+      tmin = tymin;
+
+   if (tymax < tmax)
+      tmax = tymax;
+
+   if (tmax < 0.0f)
+      // Intersection is behind the ray origin
+      return std::nullopt;
+
+   return std::make_optional(std::make_pair(tmin, tmax));
+}
+
+bool intersect_segment_aabb(const Segment& segment, const AABB& aabb) {
+   float t0 = 0.0f;
+   float t1 = 1.0f;
+   float dx = segment.end.x - segment.start.x;
+   float dy = segment.end.y - segment.start.y;
+
+   float p[4] = {-dx, dx, -dy, dy};
+   float q[4] = {segment.start.x - aabb.min.x, aabb.max.x - segment.start.x, segment.start.y - aabb.min.y,
+                 aabb.max.y - segment.start.y};
+
+   for (int i = 0; i < 4; ++i) {
+      if (std::abs(p[i]) < EPSILON) {
+         if (q[i] < 0.0f)
+            return false;
+      } else {
+         float r = q[i] / p[i];
+         if (p[i] < 0.0f) {
+            if (r > t1)
+               return false;
+            else if (r > t0)
+               t0 = r;
+         } else {
+            if (r < t0)
+               return false;
+            else if (r < t1)
+               t1 = r;
+         }
+      }
+   }
+   return true;
+}
diff --git a/OpenGL/src/geometry/BVH.h b/OpenGL/src/geometry/BVH.h
new file mode 100644
index 00000000..055229f8
--- /dev/null
+++ b/OpenGL/src/geometry/BVH.h
@@ -0,0 +1,94 @@
+#pragma once
+
+#include <glm/glm.hpp>
+#include <glm/vec2.hpp>
+#include <limits>
+#include <vector>
+#include <optional>
+#include <cstdint>
+
+struct Ray {
+   glm::vec2 origin;
+   glm::vec2 direction;
+};
+
+struct AABB {
+   glm::vec2 min;
+   glm::vec2 max;
+};
+
+struct Segment {
+   glm::vec2 start;
+   glm::vec2 end;
+};
+
+struct BvhNode {
+   uint32_t leftTypeCount; // leftType :1, leftCount :31
+   uint32_t leftOffset;
+
+   uint32_t rightTypeCount; // rightType :1, rightCount :31
+   uint32_t rightOffset;
+
+   glm::vec4 leftBBox; // (min.x, min.y, max.x, max.y)
+   glm::vec4 rightBBox;
+
+   // Helper methods for left child
+   uint32_t getLeftType() const { return leftTypeCount >> 31; }
+   uint32_t getLeftCount() const { return leftTypeCount & 0x7FFFFFFF; }
+   void     setLeft(uint32_t type, uint32_t count) { leftTypeCount = (type << 31) | (count & 0x7FFFFFFF); }
+
+   // Helper methods for right child
+   uint32_t getRightType() const { return rightTypeCount >> 31; }
+   uint32_t getRightCount() const { return rightTypeCount & 0x7FFFFFFF; }
+   void     setRight(uint32_t type, uint32_t count) { rightTypeCount = (type << 31) | (count & 0x7FFFFFFF); }
+};
+
+struct BVH {
+   std::vector<BvhNode> nodes;
+   std::vector<Segment> segments;
+
+   static BVH build(std::vector<Segment> segments);
+
+   size_t build_recursive(size_t segment_start, size_t segment_end);
+
+   AABB compute_aabb(size_t segment_start, size_t segment_end) const;
+
+   size_t partition_segments(size_t segment_start, size_t segment_end);
+
+   glm::vec4 getBoundingBoxOfNode(const BvhNode& node) const;
+
+   /// Find the closest intersection point between a ray and any segment in the BVH.
+   std::optional<std::pair<glm::vec2, const Segment*>> ray_intersect(const Ray& ray) const;
+
+   std::optional<std::tuple<glm::vec2, float, const Segment*>>
+   ray_intersect_node(const Ray& ray, size_t node_idx,
+                      std::optional<std::tuple<glm::vec2, float, const Segment*>> closest) const;
+
+   /// Find the closest intersection point between a segment and any segment in the BVH.
+   std::optional<std::pair<glm::vec2, const Segment*>> segment_intersect(const Segment& segment) const;
+
+   std::optional<std::tuple<glm::vec2, float, const Segment*>>
+   segment_intersect_node(const Segment& segment, size_t node_idx,
+                          std::optional<std::tuple<glm::vec2, float, const Segment*>> closest) const;
+};
+
+// Utility functions for intersection tests
+
+float cross(const glm::vec2& v, const glm::vec2& w);
+
+extern const float EPSILON;
+
+/// Compute the intersection point between a ray and a segment.
+/// Returns (intersection_point, t) where t is the parameter along the ray.
+std::optional<std::pair<glm::vec2, float>> intersect_ray_segment(const Ray& ray, const Segment& segment);
+
+/// Compute the intersection point between two segments.
+/// Returns (intersection_point, t) where t is the parameter along the first segment.
+std::optional<std::pair<glm::vec2, float>> intersect_segment_segment(const Segment& s1, const Segment& s2);
+
+/// Check if a ray intersects an AABB.
+/// Returns Some((tmin, tmax)) if there is an intersection, None otherwise.
+std::optional<std::pair<float, float>> intersect_ray_aabb(const Ray& ray, const AABB& aabb);
+
+/// Check if a segment intersects an AABB.
+bool intersect_segment_aabb(const Segment& segment, const AABB& aabb);
diff --git a/OpenGL/src/GeometryUtils.cpp b/OpenGL/src/geometry/GeometryUtils.cpp
similarity index 73%
rename from OpenGL/src/GeometryUtils.cpp
rename to OpenGL/src/geometry/GeometryUtils.cpp
index c68f649a..e875eaa1 100644
--- a/OpenGL/src/GeometryUtils.cpp
+++ b/OpenGL/src/geometry/GeometryUtils.cpp
@@ -5,6 +5,7 @@
 #include <cmath>
 #include <cstdlib>
 #include "earcut.hpp"
+#include "rendering/Renderer.h"
 
 namespace GeometryUtils {
 
@@ -13,6 +14,11 @@ float length2(const glm::vec2& a, const glm::vec2& b) {
    return glm::dot(temp, temp);
 }
 
+float length2(const glm::vec2& a) {
+   return glm::dot(a, a);
+}
+
+
 bool findPolygonUnion(const std::vector<std::vector<glm::vec2>>& polygons, PolyTreeD& output) {
    ClipperD clipper;
    for (const auto& polygon : polygons) {
@@ -101,28 +107,6 @@ float distancePointToLineSegment(const glm::vec2& point, const glm::vec2& lineSt
    return glm::distance(point, projection);
 }
 
-std::optional<glm::vec2> RayIntersect(const glm::vec2& ray_origin, double dx, double dy,
-                                      const std::vector<std::array<glm::vec2, 2>>& obstructionLines) {
-   float                    closest_distance = std::numeric_limits<float>::max();
-   std::optional<glm::vec2> closest_intersection;
-   for (const auto& line : obstructionLines) {
-      const auto& pa = line[0];
-      const auto& pb = line[1];
-
-      auto intersection_opt = RaySegmentIntersect(ray_origin, dx, dy, {pa.x, pa.y}, {pb.x, pb.y});
-      if (intersection_opt) {
-         auto current_distance = length2(*intersection_opt, ray_origin);
-         if (current_distance > 0.01) {
-            if (current_distance < closest_distance) {
-               closest_distance     = current_distance;
-               closest_intersection = intersection_opt;
-            }
-         }
-      }
-   }
-   return closest_intersection;
-}
-
 // Function to compute intersection between two line segments
 std::optional<glm::vec2> LineSegmentIntersect(const glm::vec2& line1_start, const glm::vec2& line1_end,
                                               const glm::vec2& line2_start, const glm::vec2& line2_end) {
@@ -158,35 +142,29 @@ std::optional<glm::vec2> LineSegmentIntersect(const glm::vec2& line1_start, cons
 }
 
 
-std::optional<glm::vec2> LineIntersect(const glm::vec2& line1_start, const glm::vec2& line1_end,
-                                       const std::vector<std::array<glm::vec2, 2>>& obstructionLines) {
-   float                    closest_distance = std::numeric_limits<float>::max();
-   std::optional<glm::vec2> closest_intersection;
-   for (const auto& line : obstructionLines) {
-      const auto& pa = line[0];
-      const auto& pb = line[1];
-
-      auto intersection_opt = LineSegmentIntersect(line1_start, line1_end, {pa.x, pa.y}, {pb.x, pb.y});
-      if (intersection_opt) {
-         auto current_distance = length2(*intersection_opt, line1_start);
-         if (current_distance > 0.01) {
-            if (current_distance < closest_distance) {
-               closest_distance     = current_distance;
-               closest_intersection = intersection_opt;
-            }
-         }
-      }
+bool isPointObstructed(const glm::vec2& position, const glm::vec2& point, const BVH& bvh) {
+   // shrink the segment a bit so it doesn't intersect with itself
+   auto segment_direction = glm::normalize(point - position);
+   auto segment_start     = position + 0.01f * segment_direction;
+   auto segment_end       = point - 0.01f * segment_direction;
+   auto intersection_opt  = bvh.segment_intersect(Segment{segment_start, segment_end});
+
+   if (intersection_opt) {
+      auto [intersection, segment] = intersection_opt.value();
+      return length2(intersection, position) < length2(point, position);
    }
-   return closest_intersection;
+   return false;
 }
 
-bool isPointObstructed(const glm::vec2& position, const glm::vec2& point,
-                       const std::vector<std::array<glm::vec2, 2>>& obstructionLines) {
-   auto intersection_opt = LineIntersect(position, point, obstructionLines);
+std::optional<glm::vec2> continue_ray(const glm::vec2& origin, const glm::vec2& direction, const BVH& bvh) {
+   // advance the starting point of the ray a tiny bit so it doesn't intersect with the ray's own origin
+   auto advanced_origin = origin + 0.01f * direction;
+
+   auto intersection_opt = bvh.ray_intersect(Ray{advanced_origin, direction});
    if (intersection_opt) {
-      return length2(intersection_opt.value(), position) < length2(point, position);
+      return intersection_opt.value().first;
    }
-   return false;
+   return std::nullopt;
 }
 
 bool adjacentInVectorCircular(size_t a, size_t b, size_t size) {
@@ -207,12 +185,9 @@ Side pointSide(const glm::vec2& point, const glm::vec2& linePoint, const glm::ve
    }
 }
 
-PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacles) {
-   bool cull_frontfaces = false;
-
+PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacles, const BVH& bvh) {
    enum class PointType { Start, End, Middle };
 
-   // Struct to hold a point along with its angle and type
    struct TaggedPoint {
       PointD    point;
       double    angle;
@@ -245,24 +220,16 @@ PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacle
 
          if (prevSide == Side::RIGHT && nextSide == Side::RIGHT) {
             tagged_points.push_back(TaggedPoint{pt, angle, PointType::Start});
-            if (cull_frontfaces) {
-               obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
-            }
+            obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
          } else if (prevSide == Side::LEFT && nextSide == Side::LEFT) {
             tagged_points.push_back(TaggedPoint{pt, angle, PointType::End});
-            if (!cull_frontfaces) {
-               obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
-            }
+            obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
          } else if (prevSide == Side::LEFT && nextSide == Side::RIGHT) {
-            if (cull_frontfaces) {
-               tagged_points.push_back(TaggedPoint{pt, angle, PointType::Middle});
-               obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
-            }
+            tagged_points.push_back(TaggedPoint{pt, angle, PointType::Middle});
+            obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
          } else {
-            if (!cull_frontfaces) {
-               tagged_points.push_back(TaggedPoint{pt, angle, PointType::Middle});
-               obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
-            }
+            tagged_points.push_back(TaggedPoint{pt, angle, PointType::Middle});
+            obstructionLines.push_back(std::array<glm::vec2, 2>{vertex, next});
          }
       }
 
@@ -287,7 +254,7 @@ PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacle
          }
       }
 
-      if (!isPointObstructed(position, {point.point.x, point.point.y}, obstructionLines)) {
+      if (!isPointObstructed(position, {point.point.x, point.point.y}, bvh)) {
          filtered_points.push_back(pointCopy);
       }
    }
@@ -304,8 +271,9 @@ PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacle
       std::optional<glm::vec2> extendedPoint;
       if (point.end != PointType::Middle) {
          glm::vec2 direction = glm::normalize(vertex - position);
-         extendedPoint       = RayIntersect(vertex, direction.x, direction.y, obstructionLines);
-         if (extendedPoint && length2(*extendedPoint, vertex) < 0.1) {
+         extendedPoint       = continue_ray(vertex, direction, bvh);
+
+         if (extendedPoint.has_value() && length2(*extendedPoint, vertex) < 0.1) {
             std::cout << "vertex super close to extended: " << length2(*extendedPoint, vertex) << std::endl;
          }
       }
diff --git a/OpenGL/src/GeometryUtils.h b/OpenGL/src/geometry/GeometryUtils.h
similarity index 94%
rename from OpenGL/src/GeometryUtils.h
rename to OpenGL/src/geometry/GeometryUtils.h
index 10faf5c1..21e1b8d5 100644
--- a/OpenGL/src/GeometryUtils.h
+++ b/OpenGL/src/geometry/GeometryUtils.h
@@ -6,6 +6,7 @@
 #include "clipper2/clipper.h"
 #include <glm/glm.hpp>
 #include "earcut.hpp"
+#include "geometry/BVH.h"
 
 namespace GeometryUtils {
 using namespace Clipper2Lib;
@@ -34,7 +35,7 @@ PathsD FlattenPolyPathD(const PolyPathD& polyPath);
  * @param obstacles The obstacles represented as PathsD (vector of paths).
  * @return PathD The visibility polygon as a vector of points.
  */
-PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacles);
+PathD ComputeVisibilityPolygon(const glm::vec2& position, const PathsD& obstacles, const BVH& bvh);
 
 /**
  * @brief Computes the intersection point between a ray and a line segment.
@@ -61,6 +62,8 @@ std::optional<glm::vec2> RaySegmentIntersect(const glm::vec2& ray_origin, double
 std::optional<glm::vec2> LineSegmentIntersect(const glm::vec2& line1_start, const glm::vec2& line1_end,
                                               const glm::vec2& line2_start, const glm::vec2& line2_end);
 
+float length2(const glm::vec2& a, const glm::vec2& b);
+float length2(const glm::vec2& a);
 } // namespace GeometryUtils
 
 
diff --git a/OpenGL/src/geometry/SceneGeometry.cpp b/OpenGL/src/geometry/SceneGeometry.cpp
new file mode 100644
index 00000000..ac80a9c4
--- /dev/null
+++ b/OpenGL/src/geometry/SceneGeometry.cpp
@@ -0,0 +1,65 @@
+#include "SceneGeometry.h"
+#include "GeometryUtils.h"
+#include "World.h"
+#include "game_objects/Tile.h"
+
+using namespace Clipper2Lib;
+using namespace GeometryUtils;
+
+SceneGeometry::WallResult SceneGeometry::computeWallPaths() {
+   std::vector<std::vector<glm::vec2>> allBounds;
+   auto                                tiles = World::getAll<Tile>(); // Simplified retrieval of all tiles
+   for (auto tile : tiles) {
+      allBounds.push_back(tile->getBounds());
+   }
+
+   WallResult result;
+   result.allBounds = allBounds;
+   result.wallPaths = std::make_unique<PolyTreeD>();
+
+   // Compute the union of all tile bounds
+   findPolygonUnion(allBounds, *result.wallPaths);
+   result.flattened = FlattenPolyPathD(*result.wallPaths);
+
+   // Build BVH from flattened paths
+   std::vector<Segment> segments;
+   for (const auto& path : result.flattened) {
+      for (size_t i = 0; i < path.size(); i++) {
+         const auto& p1 = path[i];
+         const auto& p2 = path[(i + 1) % path.size()];
+         segments.push_back(Segment{glm::vec2(p1.x, p1.y), glm::vec2(p2.x, p2.y)});
+      }
+   }
+   result.bvh = BVH::build(segments);
+
+   // Debug visualization of the BVH structure
+
+   return result;
+}
+
+
+SceneGeometry::VisibilityResult SceneGeometry::computeVisibility(SceneGeometry::WallResult& wallResult,
+                                                                 const glm::vec2&           playerPosition) {
+   SceneGeometry::VisibilityResult result{Clipper2Lib::PathD(), std::make_unique<PolyTreeD>()};
+
+   // Compute the visibility polygon
+   result.visibility = ComputeVisibilityPolygon(playerPosition, wallResult.flattened, wallResult.bvh);
+
+   // Prepare the hull for clipping
+   PathsD    hullPaths;
+   PolyTreeD combined;
+   findPolygonUnion(wallResult.allBounds, combined);
+   for (auto& child : combined) {
+      hullPaths.push_back(child->Polygon());
+   }
+
+   // Compute invisibility paths
+   result.invisibilityPaths = std::make_unique<PolyTreeD>();
+   ClipperD clipper;
+   clipper.AddSubject(hullPaths);
+   clipper.AddClip({result.visibility});
+   clipper.AddClip({wallResult.flattened});
+   clipper.Execute(ClipType::Difference, FillRule::NonZero, *result.invisibilityPaths);
+
+   return result;
+}
diff --git a/OpenGL/src/geometry/SceneGeometry.h b/OpenGL/src/geometry/SceneGeometry.h
new file mode 100644
index 00000000..85a108f9
--- /dev/null
+++ b/OpenGL/src/geometry/SceneGeometry.h
@@ -0,0 +1,26 @@
+#pragma once
+#include <vector>
+#include <glm/glm.hpp>
+#include "clipper2/clipper.h"
+#include "BVH.h"
+
+class SceneGeometry {
+public:
+   struct WallResult {
+      std::vector<std::vector<glm::vec2>>     allBounds;
+      Clipper2Lib::PathsD                     flattened;
+      std::unique_ptr<Clipper2Lib::PolyTreeD> wallPaths;
+      BVH                                     bvh;
+   };
+
+   struct VisibilityResult {
+      Clipper2Lib::PathD                      visibility;
+      std::unique_ptr<Clipper2Lib::PolyTreeD> invisibilityPaths;
+   };
+
+   static WallResult computeWallPaths();
+
+   static VisibilityResult computeVisibility(SceneGeometry::WallResult& wallResult, const glm::vec2& playerPosition);
+
+private:
+};
diff --git a/OpenGL/src/rendering/BindGroupLayout.h b/OpenGL/src/rendering/BindGroupLayout.h
index 00fe2869..ef57d3b1 100644
--- a/OpenGL/src/rendering/BindGroupLayout.h
+++ b/OpenGL/src/rendering/BindGroupLayout.h
@@ -19,7 +19,7 @@ template <typename Binding>
 struct GetToBind;
 
 // Helper struct for buffer bindings (e.g., Uniform or Storage buffers)
-template <typename T, wgpu::ShaderStage Visibility, wgpu::BufferBindingType BufferType, bool DynamicOffset = false>
+template <typename T, wgpu::ShaderStage Visibility, wgpu::BufferBindingType BufferType, bool DynamicOffset>
 struct BufferBinding {
    using Type                                             = T;
    static constexpr wgpu::ShaderStage       visibility    = Visibility;
@@ -74,7 +74,7 @@ concept BindingC = requires { typename ToBind<T>; };
 // Helper for caching
 // ------------------------------------------------------
 template <std::size_t N>
-using Ids = std::array<std::array<int32_t, 2>, N>;
+using Ids = std::array<std::array<int32_t, 3>, N>;
 template <std::size_t N>
 struct IdsHash {
    std::size_t operator()(const Ids<N>& ids) const {
@@ -145,19 +145,22 @@ struct BindGroupLayout {
    }
 
    template <size_t I, typename Binding, typename Resource>
-   static std::array<int32_t, 2> getId(Resource& resource) {
+   static std::array<int32_t, 3> getId(Resource& resource) {
       if constexpr (Binding::bindingType == BindingType::Buffer) {
          if constexpr (!Binding::dynamicOffset) {
             // Assuming ToBind<T> is Buffer
-            return std::array<int32_t, 2>{std::get<0>(resource).summed_id(), (int32_t)std::get<1>(resource)};
+            return std::array<int32_t, 3>{std::get<0>(resource).summed_id(), 
+                                        static_cast<int32_t>(std::get<1>(resource)),
+                                        static_cast<int32_t>(std::get<0>(resource).count())};
          } else if constexpr (Binding::dynamicOffset) {
-            return std::array<int32_t, 2>{resource.getBuffer()->summed_id(), -1};
+            return std::array<int32_t, 3>{resource.getBuffer()->summed_id(), -1, -1};
          }
       } else if constexpr (Binding::bindingType == BindingType::Sampler) {
-         return std::array<int32_t, 2>{resource.id, -1};
+         return std::array<int32_t, 3>{resource.id, -1, -1};
       } else if constexpr (Binding::bindingType == BindingType::Texture) {
-         return std::array<int32_t, 2>{resource->id, -1};
+         return std::array<int32_t, 3>{resource->id, -1, -1};
       }
+      return std::array<int32_t, 3>{-1, -1, -1}; // Default case
    }
 
    template <size_t I, typename Binding, typename Resource>
@@ -169,7 +172,7 @@ struct BindGroupLayout {
             // Assuming ToBind<T> is Buffer
             entry.buffer = std::get<0>(resource).get();
             entry.offset = std::get<1>(resource);
-            entry.size   = sizeof(typename Binding::Type);
+            entry.size   = sizeof(typename Binding::Type) * std::get<0>(resource).count();
          } else if constexpr (Binding::dynamicOffset) {
             entry.buffer = resource.getBuffer()->get();
             entry.offset = 0;
@@ -238,5 +241,3 @@ struct BindGroupLayouts {
                         std::forward<Tuple>(tuple));
    }
 };
-
-
diff --git a/OpenGL/src/rendering/Buffer.h b/OpenGL/src/rendering/Buffer.h
index ba221df1..c8f528c3 100644
--- a/OpenGL/src/rendering/Buffer.h
+++ b/OpenGL/src/rendering/Buffer.h
@@ -46,13 +46,15 @@ class Buffer : public std::enable_shared_from_this<Buffer<T, Uniform>> {
 public:
    // Friend declaration to allow BufferView access to private members
    friend class BufferView<T, Uniform>;
-   int32_t id;
-   int32_t generation; // Generation counter for BufferView invalidation
+   int32_t     id;
+   int32_t     generation; // Generation counter for BufferView invalidation
+   std::string name;
 
    // Constructor: Creates a buffer with specified usage and data
-   Buffer(const std::vector<T>& data, wgpu::BufferUsage usage)
+   Buffer(const std::vector<T>& data, wgpu::BufferUsage usage, std::string name = "Unnamed Buffer")
       : id(Id::get())
       , generation(0)
+      , name(name)
       , device_(Application::get().getDevice())
       , queue_(Application::get().getQueue())
       , usage_(usage) {
@@ -64,6 +66,7 @@ class Buffer : public std::enable_shared_from_this<Buffer<T, Uniform>> {
       bufferDesc.usage                  = usage_;
       bufferDesc.mappedAtCreation       = false;
       bufferDesc.size                   = capacityBytes();
+      bufferDesc.label                  = name.c_str();
 
       // Create the buffer
       buffer_ = device_.createBuffer(bufferDesc);
@@ -92,7 +95,11 @@ class Buffer : public std::enable_shared_from_this<Buffer<T, Uniform>> {
 
    // Method to upload data to the buffer
    void upload(const std::vector<T>& data) {
-      assert(data.size() <= count_ && "Data size exceeds buffer capacity.");
+      if (data.size() > capacity_) {
+         expandBuffer(data.size() * elementStride());
+      }
+
+      count_ = data.size();
 
       if constexpr (Uniform) {
          // Create a temporary buffer with padding
@@ -130,6 +137,7 @@ class Buffer : public std::enable_shared_from_this<Buffer<T, Uniform>> {
 
    // Add method: Allocates a new index and returns a BufferView
    BufferView<T, Uniform> Add(const T& data) {
+
       size_t allocatedIndex;
 
       // Reuse a deleted index if available
@@ -188,16 +196,18 @@ class Buffer : public std::enable_shared_from_this<Buffer<T, Uniform>> {
    }
 
    // Method to resize the buffer
-   void expandBuffer() {
+   void expandBuffer() { expandBuffer(capacityBytes() * 2); }
+
+   void expandBuffer(size_t newSize) {
       std::cout << "Expanding buffer" << std::endl;
-      size_t newSize = capacityBytes() * 2;
-      newSize        = std::max(newSize, elementStride()); // Ensure the buffer is at least 256 bytes
+      newSize = std::max(newSize, elementStride()); // Ensure the buffer is at least the length of one element
 
       // Create a new buffer with the new size
       wgpu::BufferDescriptor newBufferDesc = {};
       newBufferDesc.size                   = newSize;
       newBufferDesc.usage                  = usage_ | wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc;
       newBufferDesc.mappedAtCreation       = false;
+      newBufferDesc.label                  = (name + " (gen " + std::to_string(generation) + ")").c_str();
 
       wgpu::Buffer newBuffer = device_.createBuffer(newBufferDesc);
       if (!newBuffer) {
@@ -210,24 +220,27 @@ class Buffer : public std::enable_shared_from_this<Buffer<T, Uniform>> {
       // Instead, maybe we should store the encoder in Application and reuse it within a frame, or accumulate a queue
       // of things to add and add them all once we call `flush`
       // (which would also minimize wasted copies when resizing multiple times within a frame)
-      wgpu::CommandEncoder encoder = device_.createCommandEncoder();
+      wgpu::CommandEncoder* encoder = Application::get().encoder;
+      if (!encoder) {
+         std::cerr << "No encoder found when trying to resize buffer " << name << std::endl;
+         assert(false);
+      }
 
       // Copy existing data from old buffer to new buffer
-      encoder.copyBufferToBuffer(buffer_, 0, newBuffer, 0, capacityBytes());
+      auto bytes_to_copy = sizeBytes();
+      std::cout << "Copying " << bytes_to_copy << " bytes from old buffer to new buffer" << std::endl;
+      encoder->copyBufferToBuffer(buffer_, 0, newBuffer, 0, bytes_to_copy);
 
-      // Finish encoding and submit the commands
-      wgpu::CommandBuffer commands = encoder.finish();
-      queue_.submit(1, &commands);
-      encoder.release();
-      commands.release();
       generation++;
 
       // Queue the old buffer for destruction
       DeadBuffers::buffers.push_back(buffer_);
       buffer_ = newBuffer;
 
-      // Update buffer capaticy now that it's been resized
+      // Update buffer capacity now that it's been resized
       capacity_ = newSize / elementStride();
+
+      std::cout << "Buffer " << name << " resized - new capacity: " << capacityBytes() << " bytes" << std::endl;
    }
 
    // Method to free an index (called by BufferView destructor)
@@ -257,26 +270,54 @@ using IndexBuffer = Buffer<uint16_t, false>;
 template <typename T>
 using UniformBuffer = Buffer<T, true>;
 
-
 template <typename T, bool Uniform = true>
 class BufferView {
 public:
    // Constructor: Acquires an index from the Buffer
    BufferView(std::shared_ptr<Buffer<T, Uniform>> buffer, size_t index)
-      : buffer_(buffer)
-      , index_(index) {}
-
-   // Destructor: Frees the index back to the Buffer
-   ~BufferView() { buffer_->freeIndex(index_); }
+      : buffer_(std::move(buffer))
+      , index_(index)
+      , valid_(true) {}
+
+   // Destructor: Frees the index back to the Buffer if valid
+   ~BufferView() {
+      if (valid_ && buffer_) {
+         buffer_->freeIndex(index_);
+      }
+   }
 
    // Deleted copy constructor and assignment operator
-   BufferView(const BufferView&)             = delete;
-   BufferView& operator=(const BufferView&)  = delete;
-   BufferView(BufferView&& other)            = delete;
-   BufferView& operator=(BufferView&& other) = delete;
+   BufferView(const BufferView&)            = delete;
+   BufferView& operator=(const BufferView&) = delete;
+
+   // Move constructor
+   BufferView(BufferView&& other)
+      : buffer_(std::move(other.buffer_))
+      , index_(other.index_)
+      , valid_(other.valid_) {
+      other.valid_ = false; // Invalidate the moved-from object
+   }
+
+   // Move assignment
+   BufferView& operator=(BufferView&& other) {
+      if (this != &other) {
+         if (valid_ && buffer_) {
+            buffer_->freeIndex(index_); // Free the current index if valid
+         }
+         buffer_      = std::move(other.buffer_);
+         index_       = other.index_;
+         valid_       = other.valid_;
+         other.valid_ = false; // Invalidate the moved-from object
+      }
+      return *this;
+   }
 
    // Update method to modify the data at this index
-   void Update(const T& data) { buffer_->updateBuffer(data, index_); }
+   void Update(const T& data) {
+      if (valid_ && buffer_) {
+         buffer_->updateBuffer(data, index_);
+      }
+   }
 
    // Getter for the index
    std::shared_ptr<Buffer<T, Uniform>> getBuffer() const { return buffer_; }
@@ -286,15 +327,16 @@ class BufferView {
    static BufferView<T, Uniform> create(const T& data) {
       static std::shared_ptr<Buffer<T, Uniform>> buffer = std::make_shared<Buffer<T, Uniform>>(
          std::vector<T>{},
-         wgpu::bothBufferUsages(wgpu::bothBufferUsages(wgpu::BufferUsage::CopySrc, wgpu::BufferUsage::CopyDst),
-                                wgpu::BufferUsage::Uniform));
+         wgpu::bothBufferUsages(wgpu::BufferUsage::CopySrc, wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Uniform));
       return buffer->Add(data);
    }
 
 private:
    std::shared_ptr<Buffer<T, Uniform>> buffer_;
    size_t                              index_;
+   bool                                valid_; // Added to manage ownership state
 };
 
+
 template <typename T>
 using UniformBufferView = BufferView<T, true>;
diff --git a/OpenGL/src/rendering/ComputePass.cpp b/OpenGL/src/rendering/ComputePass.cpp
new file mode 100644
index 00000000..7bc75479
--- /dev/null
+++ b/OpenGL/src/rendering/ComputePass.cpp
@@ -0,0 +1,20 @@
+#include "ComputePass.h"
+#include "CommandEncoder.h"
+#include "Application.h"
+
+
+ComputePass::ComputePass(CommandEncoder& encoder, wgpu::TextureView& targetView) {
+   auto& application = Application::get();
+
+   wgpu::ComputePassDescriptor computePassDesc;
+   computePass_ = encoder.get().beginComputePass(computePassDesc);
+}
+
+ComputePass::~ComputePass() {
+   computePass_.end();
+   computePass_.release();
+}
+
+wgpu::ComputePassEncoder& ComputePass::get() {
+   return computePass_;
+}
diff --git a/OpenGL/src/rendering/ComputePass.h b/OpenGL/src/rendering/ComputePass.h
new file mode 100644
index 00000000..a96b842d
--- /dev/null
+++ b/OpenGL/src/rendering/ComputePass.h
@@ -0,0 +1,57 @@
+#pragma once
+
+#include <webgpu/webgpu.hpp>
+#include "Buffer.h"
+#include "BindGroup.h"
+#include "ComputePipeline.h"
+
+class CommandEncoder;
+
+class ComputePass {
+public:
+   explicit ComputePass(CommandEncoder& encoder, wgpu::TextureView& targetView);
+   ~ComputePass();
+
+   wgpu::ComputePassEncoder& get();
+
+   template <typename T>
+   void setPipeline(const T& pass) {
+      if (last_set_compute_pipeline != pass.id) {
+         computePass_.setPipeline(pass.GetPipeline());
+         last_set_compute_pipeline = pass.id;
+      }
+   }
+
+   void setBindGroup(uint32_t group, BindGroup bindGroup, std::vector<uint32_t> offset) {
+      // TODO: this is overly conservative because we could be checking the that the bind group / offset is the same as
+      // the last time we set that index
+      if ((last_set_bind_group != (int32_t)group) ||           // Check index
+          (bindGroup.id != (int32_t)last_set_bind_group_id) || // Check bind group id
+          offset != last_set_bind_group_offset                 // check offset
+      ) {
+         computePass_.setBindGroup(group, bindGroup.get(), offset.size(), offset.data());
+         last_set_bind_group        = group;
+         last_set_bind_group_offset = offset;
+         last_set_bind_group_id     = bindGroup.id;
+      }
+   }
+
+   template <typename BGLs>
+   void dispatch(ComputePipeline<BGLs>& pipeline, BindGroup bindGroup, std::vector<uint32_t> offsets, uint32_t x,
+                 uint32_t y = 1, uint32_t z = 1) {
+      setPipeline(pipeline);
+      setBindGroup(0, bindGroup, offsets);
+      computePass_.dispatchWorkgroups(x, y, z);
+   }
+
+
+private:
+   wgpu::ComputePassEncoder computePass_;
+
+   int32_t               last_set_compute_pipeline = -1;
+   int32_t               last_set_bind_group       = -1;
+   int32_t               last_set_bind_group_id    = -1;
+   std::vector<uint32_t> last_set_bind_group_offset;
+   int32_t               last_set_vertex_buffer = -1;
+   int32_t               last_set_index_buffer  = -1;
+};
diff --git a/OpenGL/src/rendering/ComputePipeline.h b/OpenGL/src/rendering/ComputePipeline.h
new file mode 100644
index 00000000..570f7133
--- /dev/null
+++ b/OpenGL/src/rendering/ComputePipeline.h
@@ -0,0 +1,69 @@
+#pragma once
+
+#include <webgpu/webgpu.hpp>
+#include <vector>
+#include <string>
+#include "Shader.h"
+#include "glm/glm.hpp"
+#include "VertexBufferLayout.h"
+#include "BindGroupLayout.h"
+#include "DataFormats.h"
+#include "Id.h"
+
+static int32_t created_render_pipelines = 0;
+
+template <typename BGLs>
+class ComputePipeline {
+public:
+   ComputePipeline(std::filesystem::path shaderPath)
+      : id(Id::get())
+      , device(Application::get().getDevice())
+      , bindGroupLayouts(BGLs::CreateLayouts(device)) {
+      std::string label = shaderPath.stem().string();
+      Shader      shader(device, shaderPath);
+
+      wgpu::PipelineLayoutDescriptor layoutDesc{};
+      layoutDesc.bindGroupLayoutCount = bindGroupLayouts.size();
+      layoutDesc.bindGroupLayouts     = (WGPUBindGroupLayout*)bindGroupLayouts.data();
+      layoutDesc.label                = label.c_str();
+      wgpu::PipelineLayout layout     = device.createPipelineLayout(layoutDesc);
+
+      // Define pipeline descriptor
+      wgpu::ComputePipelineDescriptor computePipelineDesc = wgpu::Default;
+      computePipelineDesc.compute.constantCount           = 0;
+      computePipelineDesc.compute.constants               = nullptr;
+      computePipelineDesc.compute.entryPoint              = "compute_main";
+      computePipelineDesc.compute.module                  = shader.GetShaderModule();
+      computePipelineDesc.layout                          = layout;
+
+      // Create compute pipeline
+      pipeline = device.createComputePipeline(computePipelineDesc);
+   };
+
+   ~ComputePipeline() {
+      if (pipeline) {
+         pipeline.release();
+      }
+   };
+
+   // Deleted copy constructor and assignment operator
+   ComputePipeline(const ComputePipeline&)             = delete;
+   ComputePipeline& operator=(const ComputePipeline&)  = delete;
+   ComputePipeline(ComputePipeline&& other)            = delete;
+   ComputePipeline& operator=(ComputePipeline&& other) = delete;
+
+   template <typename... Ts>
+   auto BindGroups(Ts&&... ts) {
+      return BGLs::BindGroups(device, std::forward<Ts>(ts)...);
+   }
+
+   wgpu::ComputePipeline              GetPipeline() const { return pipeline; }
+   std::vector<wgpu::BindGroupLayout> GetBindGroupLayouts() const { return bindGroupLayouts; }
+
+   const int32_t id;
+
+private:
+   wgpu::Device                       device;
+   wgpu::ComputePipeline              pipeline;
+   std::vector<wgpu::BindGroupLayout> bindGroupLayouts;
+};
diff --git a/OpenGL/src/rendering/DataFormats.h b/OpenGL/src/rendering/DataFormats.h
index 9b00e225..5ccdd50d 100644
--- a/OpenGL/src/rendering/DataFormats.h
+++ b/OpenGL/src/rendering/DataFormats.h
@@ -3,6 +3,7 @@
 #include <webgpu/webgpu.hpp>
 #include "VertexBufferLayout.h"
 #include "BindGroupLayout.h"
+#include "../geometry/BVH.h"
 
 #define GLM_ENABLE_EXPERIMENTAL
 #include "glm/gtx/hash.hpp"
@@ -23,9 +24,10 @@ struct StarUniforms {
 };
 
 using StarUniformBinding =
-   BufferBinding<StarUniforms,                                                       // Type of the buffer
+   BufferBinding<StarUniforms,                                                                   // Type of the buffer
                  wgpu::bothShaderStages(wgpu::ShaderStage::Vertex, wgpu::ShaderStage::Fragment), // Shader visibility
-                 wgpu::BufferBindingType::Uniform                                    // Buffer binding type
+                 wgpu::BufferBindingType::Uniform,                                               // Buffer binding type
+                 false // Buffer is not dynamically offset
                  >;
 // ============================================================
 
@@ -138,10 +140,52 @@ using FogLayout = BindGroupLayout<FogVertexUniformBinding, FogFragmentUniformBin
 using FogVertex = glm::vec2;
 // ============================================================
 
-// Cinders
+// Particles
 // ============================================================
-//struct Particle {
-//   glm::vec2 position;
-//   glm::vec2 velocity;
-//   glm::vec4 color;
-//};  
\ No newline at end of file
+struct Particle {
+   glm::vec2 position;
+   glm::vec2 velocity;
+   glm::vec4 color;
+
+   typedef InstanceBufferLayout<glm::vec2> Layout;
+};
+
+struct ParticleVertex {
+   glm::vec2 position;
+
+   bool operator==(const ParticleVertex& other) const { return position == other.position; }
+};
+
+struct ParticleVertexUniform {
+   glm::mat4 u_MVP;
+
+   bool operator==(const ParticleVertexUniform& other) const { return u_MVP == other.u_MVP; }
+};
+
+namespace std {
+template <>
+struct hash<ParticleVertex> {
+   size_t operator()(const ParticleVertex& v) const {
+      size_t h1 = std::hash<glm::vec2>{}(v.position);
+      return h1;
+   }
+};
+} // namespace std
+
+struct ParticleWorldInfo {
+   float     deltaTime; // at byte offset 0
+   float     _pad0;
+   glm::vec2 mousePos; // at byte offset 8
+
+   ParticleWorldInfo(float deltaTime, glm::vec2 mousePos)
+      : deltaTime(deltaTime)
+      , mousePos(mousePos) {}
+};
+
+using ParticleLayout = BindGroupLayout<
+   BufferBinding<ParticleVertexUniform, wgpu::ShaderStage::Vertex, wgpu::BufferBindingType::Uniform, true>>;
+using ParticleComputeLayout =
+   BindGroupLayout<BufferBinding<Particle, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage, false>,
+                   BufferBinding<ParticleWorldInfo, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Uniform, true>,
+                   BufferBinding<Segment, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::ReadOnlyStorage, false>,
+                   BufferBinding<BvhNode, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::ReadOnlyStorage, false>>;
diff --git a/OpenGL/src/rendering/RenderPass.h b/OpenGL/src/rendering/RenderPass.h
index fb73c532..7fbe8e51 100644
--- a/OpenGL/src/rendering/RenderPass.h
+++ b/OpenGL/src/rendering/RenderPass.h
@@ -39,10 +39,11 @@ class RenderPass {
    }
 
    template <typename T>
-   void setVertexBuffer(const Buffer<T>& buffer) {
-      if (last_set_vertex_buffer != (int32_t)buffer.summed_id()) {
-         renderPass_.setVertexBuffer(0, buffer.get(), 0, buffer.sizeBytes());
-         last_set_vertex_buffer = buffer.summed_id();
+   void setVertexBuffer(Buffer<T>& buffer, uint32_t index) {
+      if (last_set_vertex_buffer != (int32_t)buffer.summed_id() || last_set_vertex_buffer_index != (int32_t)index) {
+         renderPass_.setVertexBuffer(index, buffer.get(), 0, buffer.sizeBytes());
+         last_set_vertex_buffer       = buffer.summed_id();
+         last_set_vertex_buffer_index = index;
       }
    }
 
@@ -53,14 +54,24 @@ class RenderPass {
       }
    }
 
-   template <typename Pipeline, typename Vertex>
-   void Draw(const Pipeline& pipeline, const Buffer<Vertex>& pointBuffer, const IndexBuffer& indexBuffer,
-             BindGroup bindGroup, std::vector<uint32_t> offset) {
+   template <typename Pipeline, typename... Vertices>
+   void DrawInstanced(const Pipeline& pipeline, const IndexBuffer& indexBuffer, BindGroup bindGroup,
+                      std::vector<uint32_t> offset, uint32_t instanceCount, Buffer<Vertices>&... bufs) {
       setPipeline(pipeline);
       setBindGroup(0, bindGroup, offset);
-      setVertexBuffer(pointBuffer);
+
+      size_t bufferIndex = 0;
+      // Apply setBuffer to each vertex buffer
+      (setVertexBuffer(bufs, bufferIndex++), ...);
+
       setIndexBuffer(indexBuffer);
-      renderPass_.drawIndexed(indexBuffer.count(), 1, 0, 0, 0);
+      renderPass_.drawIndexed(indexBuffer.count(), instanceCount, 0, 0, 0);
+   }
+
+   template <typename Pipeline, typename Vertex>
+   void Draw(const Pipeline& pipeline, Buffer<Vertex>& pointBuffer, const IndexBuffer& indexBuffer, BindGroup bindGroup,
+             std::vector<uint32_t> offset) {
+      DrawInstanced(pipeline, indexBuffer, bindGroup, offset, 1, pointBuffer);
    }
 
 private:
@@ -70,6 +81,7 @@ class RenderPass {
    int32_t               last_set_bind_group      = -1;
    int32_t               last_set_bind_group_id   = -1;
    std::vector<uint32_t> last_set_bind_group_offset;
-   int32_t               last_set_vertex_buffer = -1;
-   int32_t               last_set_index_buffer  = -1;
+   int32_t               last_set_vertex_buffer       = -1;
+   int32_t               last_set_vertex_buffer_index = -1;
+   int32_t               last_set_index_buffer        = -1;
 };
diff --git a/OpenGL/src/rendering/RenderPipeline.h b/OpenGL/src/rendering/RenderPipeline.h
index e2890861..6ef72141 100644
--- a/OpenGL/src/rendering/RenderPipeline.h
+++ b/OpenGL/src/rendering/RenderPipeline.h
@@ -8,25 +8,26 @@
 #include "VertexBufferLayout.h"
 #include "BindGroupLayout.h"
 #include "DataFormats.h"
+#include "Id.h"
 
-static int32_t created_render_pipelines = 0;
-
-template <typename BGLs, typename VBL>
+template <typename BGLs, typename VBLs>
 class RenderPipeline {
 public:
    RenderPipeline(std::filesystem::path   shaderPath,
                   wgpu::PrimitiveTopology topology = wgpu::PrimitiveTopology::TriangleList)
-      : id(created_render_pipelines++)
+      : id(Id::get())
       , device(Application::get().getDevice())
       , bindGroupLayouts(BGLs::CreateLayouts(device)) {
       std::string label = shaderPath.stem().string();
       Shader      shader(device, shaderPath);
 
       // Create vertex buffer layouts
-      auto vertexInfo = VBL::CreateLayout();
+      auto vertexInfos = VBLs::CreateLayouts();
 
       std::vector<wgpu::VertexBufferLayout> wgpuVertexLayouts;
-      wgpuVertexLayouts.push_back(vertexInfo.layout);
+      for (const auto& info : vertexInfos) {
+         wgpuVertexLayouts.push_back(info.layout);
+      }
 
       wgpu::PipelineLayoutDescriptor layoutDesc{};
       layoutDesc.bindGroupLayoutCount = bindGroupLayouts.size();
@@ -75,7 +76,7 @@ class RenderPipeline {
 
       // Define pipeline descriptor
       wgpu::RenderPipelineDescriptor pipelineDesc     = {};
-      pipelineDesc.label                              = "Render Pipeline";
+      pipelineDesc.label                              = label.c_str();
       pipelineDesc.layout                             = layout;
       pipelineDesc.vertex                             = vertexState;
       pipelineDesc.fragment                           = &fragmentState;
diff --git a/OpenGL/src/rendering/Renderer.cpp b/OpenGL/src/rendering/Renderer.cpp
index 738b46e7..465eb410 100644
--- a/OpenGL/src/rendering/Renderer.cpp
+++ b/OpenGL/src/rendering/Renderer.cpp
@@ -5,12 +5,19 @@
 #include "../game_objects/Camera.h"
 
 Renderer::Renderer()
-   : stars(RenderPipeline<BindGroupLayouts<BindGroupLayout<StarUniformBinding>>, VertexBufferLayout<glm::vec2>>(
-        "stars.wgsl"))
-   , squareObject(RenderPipeline<BindGroupLayouts<SquareObjectLayout>, VertexBufferLayout<glm::vec2, glm::vec2>>(
-        "square_object.wgsl"))
-   , line(RenderPipeline<BindGroupLayouts<LineLayout>, VertexBufferLayout<LineVertex>>("line.wgsl"))
-   , fog(RenderPipeline<BindGroupLayouts<FogLayout>, VertexBufferLayout<FogVertex>>("fog.wgsl"))
+   : stars(RenderPipeline<BindGroupLayouts<BindGroupLayout<StarUniformBinding>>,
+                          VertexBufferLayouts<VertexBufferLayout<glm::vec2>>>("stars.wgsl"))
+   , squareObject(RenderPipeline<BindGroupLayouts<SquareObjectLayout>,
+                                 VertexBufferLayouts<VertexBufferLayout<glm::vec2, glm::vec2>>>("square_object.wgsl"))
+   , line(
+        RenderPipeline<BindGroupLayouts<LineLayout>, VertexBufferLayouts<VertexBufferLayout<LineVertex>>>("line.wgsl"))
+   , fog(RenderPipeline<BindGroupLayouts<FogLayout>, VertexBufferLayouts<VertexBufferLayout<FogVertex>>>("fog.wgsl"))
+   , particles(
+        RenderPipeline<
+           BindGroupLayouts<ParticleLayout>,
+           VertexBufferLayouts<VertexBufferLayout<glm::vec2>, InstanceBufferLayout<glm::vec2, glm::vec2, glm::vec4>>>(
+           "particles.wgsl"))
+   , particlesCompute(ComputePipeline<BindGroupLayouts<ParticleComputeLayout>>("particlesCompute.wgsl"))
    , device(Application::get().getDevice())
    , linePoints(
         std::vector<LineVertex>{
@@ -27,35 +34,39 @@ Renderer::Renderer()
       },
       wgpu::bothBufferUsages(wgpu::BufferUsage::CopyDst, wgpu::BufferUsage::Index)) {}
 
-glm::mat4 CalculateMVP(const glm::vec2& objectPosition, float objectRotationDegrees, float objectScale) {
-   glm::ivec2 windowSize = Application::get().windowSize();
-
-   // Calculate aspect ratio
-   float aspectRatio = static_cast<float>(windowSize.x) / static_cast<float>(windowSize.y);
-
-   // Create orthographic projection matrix
-   float     orthoWidth = Camera::scale * aspectRatio;
-   float     left       = -orthoWidth / 2.0f;
-   float     right      = orthoWidth / 2.0f;
-   float     bottom     = -Camera::scale / 2.0f;
-   float     top        = Camera::scale / 2.0f;
-   float     near       = -1.0f;
-   float     far        = 1.0f;
-   glm::mat4 projection = glm::ortho(left, right, bottom, top, near, far);
-
-   // Create view matrix
-   glm::mat4 view = glm::translate(glm::mat4(1.0f), glm::vec3(-Camera::position, 0.0f));
-
-   // Create model matrix
+glm::mat4 CalculateModel(const glm::vec2& objectPosition, float objectRotationDegrees, float objectScale) {
    glm::mat4 model           = glm::translate(glm::mat4(1.0f), glm::vec3(objectPosition, 0.0f));
    float     rotationRadians = glm::radians(objectRotationDegrees);
    model                     = glm::rotate(model, rotationRadians, glm::vec3(0.0f, 0.0f, 1.0f));
    model                     = glm::scale(model, glm::vec3(objectScale, objectScale, 1.0f));
+   return model;
+}
+
+glm::mat4 CalculateView() {
+   return glm::translate(glm::mat4(1.0f), glm::vec3(-Camera::position, 0.0f));
+}
+
+glm::mat4 CalculateProjection() {
+   glm::ivec2 windowSize  = Application::get().windowSize();
+   float      aspectRatio = static_cast<float>(windowSize.x) / static_cast<float>(windowSize.y);
 
-   // Combine matrices to form MVP
-   glm::mat4 mvp = projection * view * model;
+   float orthoWidth = Camera::scale * aspectRatio;
+   float left       = -orthoWidth / 2.0f;
+   float right      = orthoWidth / 2.0f;
+   float bottom     = -Camera::scale / 2.0f;
+   float top        = Camera::scale / 2.0f;
+   float near       = -1.0f;
+   float far        = 1.0f;
+
+   return glm::ortho(left, right, bottom, top, near, far);
+}
+
+glm::mat4 CalculateMVP(const glm::vec2& objectPosition, float objectRotationDegrees, float objectScale) {
+   glm::mat4 projection = CalculateProjection();
+   glm::mat4 view       = CalculateView();
+   glm::mat4 model      = CalculateModel(objectPosition, objectRotationDegrees, objectScale);
 
-   return mvp;
+   return projection * view * model;
 }
 
 void Renderer::DrawLine(Line line, RenderPass& renderPass) {
diff --git a/OpenGL/src/rendering/Renderer.h b/OpenGL/src/rendering/Renderer.h
index 24b7a7ac..970b8085 100644
--- a/OpenGL/src/rendering/Renderer.h
+++ b/OpenGL/src/rendering/Renderer.h
@@ -2,6 +2,7 @@
 
 #include <webgpu/webgpu.hpp>
 
+#include "ComputePipeline.h"
 #include "RenderPipeline.h"
 #include "RenderPass.h"
 #include "Texture.h"
@@ -23,15 +24,22 @@ class Renderer {
 public:
    Renderer();
 
-   RenderPipeline<BindGroupLayouts<BindGroupLayout<StarUniformBinding>>, VertexBufferLayout<glm::vec2>> stars;
-   RenderPipeline<BindGroupLayouts<SquareObjectLayout>, VertexBufferLayout<glm::vec2, glm::vec2>>       squareObject;
-   RenderPipeline<BindGroupLayouts<LineLayout>, VertexBufferLayout<LineVertex>>                         line;
-   RenderPipeline<BindGroupLayouts<FogLayout>, VertexBufferLayout<FogVertex>>                           fog;
+   RenderPipeline<BindGroupLayouts<BindGroupLayout<StarUniformBinding>>,
+                  VertexBufferLayouts<VertexBufferLayout<glm::vec2>>>
+      stars;
+   RenderPipeline<BindGroupLayouts<SquareObjectLayout>, VertexBufferLayouts<VertexBufferLayout<glm::vec2, glm::vec2>>>
+                                                                                                     squareObject;
+   RenderPipeline<BindGroupLayouts<LineLayout>, VertexBufferLayouts<VertexBufferLayout<LineVertex>>> line;
+   RenderPipeline<BindGroupLayouts<FogLayout>, VertexBufferLayouts<VertexBufferLayout<FogVertex>>>   fog;
+   RenderPipeline<
+      BindGroupLayouts<ParticleLayout>,
+      VertexBufferLayouts<VertexBufferLayout<glm::vec2>, InstanceBufferLayout<glm::vec2, glm::vec2, glm::vec4>>>
+                                                            particles;
+   ComputePipeline<BindGroupLayouts<ParticleComputeLayout>> particlesCompute;
 
    TextureSampler sampler;
 
-   wgpu::RenderPassEncoder renderPass;
-   wgpu::Device            device;
+   wgpu::Device device;
 
    static glm::vec2 MousePos();
    static glm::vec2 ScreenToWorldPosition(const glm::vec2& screenPos);
@@ -51,4 +59,7 @@ class Renderer {
 };
 
 
+glm::mat4 CalculateModel(const glm::vec2& objectPosition, float objectRotationDegrees, float objectScale);
+glm::mat4 CalculateView();
+glm::mat4 CalculateProjection();
 glm::mat4 CalculateMVP(const glm::vec2& objectPosition, float objectRotationDegrees, float objectScale);
diff --git a/OpenGL/src/rendering/Shader.cpp b/OpenGL/src/rendering/Shader.cpp
index 2738a6f4..65bc0c04 100644
--- a/OpenGL/src/rendering/Shader.cpp
+++ b/OpenGL/src/rendering/Shader.cpp
@@ -4,7 +4,8 @@
 Shader::Shader(wgpu::Device device, const std::filesystem::path& filePath)
    : filePath(filePath) {
    // Read the shader source code from the file
-   std::string src = ReadFile(Application::get().res_path / "shaders" / filePath);
+   auto fullPath = Application::get().res_path / "shaders" / filePath;
+   std::string src = ReadFile(fullPath);
    if (src.empty()) {
       std::cerr << "Shader source is empty. Failed to load shader from: " << filePath << std::endl;
       return;
@@ -59,21 +60,48 @@ Shader& Shader::operator=(Shader&& other) noexcept {
    return *this;
 }
 
+std::string Shader::ProcessIncludes(const std::string& source, const std::filesystem::path& shaderPath) const {
+    std::istringstream stream(source);
+    std::ostringstream result;
+    std::string line;
+
+    while (std::getline(stream, line)) {
+        if (line.find("#include") == 0) {
+            size_t firstQuote = line.find('<');
+            size_t lastQuote = line.find('>');
+            if (firstQuote != std::string::npos && lastQuote != std::string::npos) {
+                std::string filename = line.substr(firstQuote + 1, lastQuote - firstQuote - 1);
+                auto includePath = shaderPath.parent_path() / filename;
+                
+                // Read the included file
+                std::string includeContent = ReadFile(includePath);
+                if (!includeContent.empty()) {
+                    result << includeContent << '\n';
+                }
+            }
+        } else {
+            result << line << '\n';
+        }
+    }
+    return result.str();
+}
+
 std::string Shader::ReadFile(const std::filesystem::path& path) const {
-   if (!std::filesystem::exists(path)) {
-      std::cerr << "Shader file does not exist: " << path << std::endl;
-      return "";
-   }
+    if (!std::filesystem::exists(path)) {
+        std::cerr << "Shader file does not exist: " << path << std::endl;
+        return "";
+    }
 
-   std::ifstream fileStream(path, std::ios::in | std::ios::binary);
-   if (!fileStream) {
-      std::cerr << "Failed to open shader file: " << path << std::endl;
-      return "";
-   }
+    std::ifstream fileStream(path, std::ios::in | std::ios::binary);
+    if (!fileStream) {
+        std::cerr << "Failed to open shader file: " << path << std::endl;
+        return "";
+    }
 
-   std::ostringstream contents;
-   contents << fileStream.rdbuf();
-   fileStream.close();
+    std::ostringstream contents;
+    contents << fileStream.rdbuf();
+    fileStream.close();
 
-   return contents.str();
+    // Process includes in the file content
+    return ProcessIncludes(contents.str(), path);
 }
diff --git a/OpenGL/src/rendering/Shader.h b/OpenGL/src/rendering/Shader.h
index 89b26f51..aabd56ac 100644
--- a/OpenGL/src/rendering/Shader.h
+++ b/OpenGL/src/rendering/Shader.h
@@ -29,4 +29,5 @@ class Shader {
 
    // Helper function to read file contents
    std::string ReadFile(const std::filesystem::path& path) const;
+   std::string ProcessIncludes(const std::string& source, const std::filesystem::path& shaderPath) const;
 };
diff --git a/OpenGL/src/rendering/VertexBufferLayout.h b/OpenGL/src/rendering/VertexBufferLayout.h
index 87a3e935..b835f134 100644
--- a/OpenGL/src/rendering/VertexBufferLayout.h
+++ b/OpenGL/src/rendering/VertexBufferLayout.h
@@ -3,6 +3,7 @@
 #include <webgpu/webgpu.hpp>
 #include <vector>
 #include <string>
+#include <tuple>
 #include "glm/glm.hpp"
 
 // Make a type like std::tuple<wgpu::VertexBufferLayout, std::vector<wgpu::VertexAttribute>>, but delete its copy and
@@ -21,6 +22,28 @@ struct VertexBufferInfo {
    VertexBufferInfo& operator=(VertexBufferInfo&& other) noexcept = default;
 };
 
+// Compose multiple vertex buffer layouts
+template <typename... Layouts>
+struct VertexBufferLayouts {
+   static std::vector<VertexBufferInfo> CreateLayouts() {
+      std::vector<VertexBufferInfo> layouts;
+      layouts.reserve(sizeof...(Layouts));
+      (layouts.push_back(Layouts::CreateLayout()), ...);
+
+      // Each VBLayoutEntry must have a unique location
+      // But each VBLayout as returned by CreateLayout starts at 0
+      // So we need to renumber the locations
+      std::size_t location = 0;
+      for (auto& layout : layouts) {
+         for (auto& attribute : *layout.attributes) {
+            attribute.shaderLocation = location++;
+         }
+      }
+
+      return layouts;
+   }
+};
+
 
 // template function that converts a type to a wgpu::VertexFormat
 template <typename T>
@@ -46,6 +69,7 @@ constexpr std::size_t type_size() {
    return sizeof(T);
 }
 
+// Single vertex buffer layout
 template <typename... Types>
 struct VertexBufferLayout {
    static std::vector<wgpu::VertexAttribute> Attributes() {
@@ -80,3 +104,13 @@ struct VertexBufferLayout {
       return attribute;
    }
 };
+
+
+template <typename... Types>
+struct InstanceBufferLayout {
+   static VertexBufferInfo CreateLayout() {
+      auto info            = VertexBufferLayout<Types...>::CreateLayout();
+      info.layout.stepMode = wgpu::VertexStepMode::Instance;
+      return info;
+   }
+};
diff --git a/OpenGL/src/rendering/webgpu-utils.h b/OpenGL/src/rendering/webgpu-utils.h
index cc6ad2e8..efe86e49 100644
--- a/OpenGL/src/rendering/webgpu-utils.h
+++ b/OpenGL/src/rendering/webgpu-utils.h
@@ -11,4 +11,11 @@ constexpr wgpu::ShaderStage bothShaderStages(wgpu::ShaderStage lhs, wgpu::Shader
 constexpr wgpu::BufferUsage bothBufferUsages(wgpu::BufferUsage lhs, wgpu::BufferUsage rhs) {
    return wgpu::BufferUsage(WGPUBufferUsage(lhs.m_raw | rhs.m_raw));
 }
+constexpr wgpu::BufferUsage bothBufferUsages(wgpu::BufferUsage lhs, wgpu::BufferUsage rhs, wgpu::BufferUsage rhs2) {
+   return wgpu::BufferUsage(WGPUBufferUsage(lhs.m_raw | rhs.m_raw | rhs2.m_raw));
+}
+constexpr wgpu::BufferUsage bothBufferUsages(wgpu::BufferUsage lhs, wgpu::BufferUsage rhs, wgpu::BufferUsage rhs2,
+                                             wgpu::BufferUsage rhs3) {
+   return wgpu::BufferUsage(WGPUBufferUsage(lhs.m_raw | rhs.m_raw | rhs2.m_raw | rhs3.m_raw));
+}
 } // namespace wgpu
diff --git a/readme.md b/readme.md
index a494aa3b..0f57da15 100644
--- a/readme.md
+++ b/readme.md
@@ -47,3 +47,7 @@ python -m http.server
 ```
 
 Then check out http://localhost:8000/OpenGL/Spec-Hops.html 
+
+## Thanks
+
+1. [This page from Alister Chowdhury](https://alister-chowdhury.github.io/posts/20230620-raytracing-in-2d/) which Andre wish he'd found earlier.