Add cache handling for tiles

fetch_libs.sh

@@ -59,14 +59,14 @@ else
     echo "  > cglm already present"
 fi
 
-# 5. sokol_gp.h
+# 5. stb_ds.h
-if [ ! -f "$LIB_DIR/sokol/sokol_gp.h" ]; then
+if [ ! -f "$LIB_DIR/util/stb_ds.h" ]; then
-    echo "  > Fetching cglm..."
+    echo "  > Fetching STB..."
-    git clone --depth 1 https://github.com/edubart/sokol_gp.git "$LIB_DIR/sokol_gp_tmp"
+    git clone --depth 1 https://github.com/nothings/stb.git "$LIB_DIR/stb_tmp"
-    cp -r "$LIB_DIR/sokol_gp_tmp/sokol_gp.h" "$LIB_DIR/sokol/sokol_gp.h"
+    cp -r "$LIB_DIR/stb_tmp/stb_ds.h" "$LIB_DIR/util/stb_ds.h"
-    rm -rf "$LIB_DIR/sokol_gp_tmp"
+    rm -rf "$LIB_DIR/stb_tmp"
 else
-    echo "  > sokol_gp.h already present"
+    echo "  > stb_ds.h already present"
 fi
 
 echo "=== Done ==="

makefile

@@ -31,8 +31,8 @@ EMCC_FLAGS = --use-port=emdawnwebgpu \
              -sALLOW_MEMORY_GROWTH \
              -msimd128 \
              -sFILESYSTEM=0 \
-             -flto \
+			 -sMALLOC=emmalloc \
-			 -Rpass=loop-vectorize
+             -flto
 
 # Shell template
 SHELL_FILE = shell.html
@@ -46,6 +46,7 @@ $(TARGET): $(SHADER_HEADERS) $(C_SOURCES) $(IMGUI_SOURCES) $(CGLM_SOURCES) $(SHE
 		-o $(TARGET) \
 		$(EMCC_FLAGS) \
 		-O3 \
+		--closure 1 \
 		-I$(LIB_DIR)/sokol \
 		-I$(LIB_DIR)/imgui \
 		-I$(LIB_DIR)/imgui/imgui \

src/api.h

@@ -1,17 +1,26 @@
 #ifndef API_DEFINITION
 #define API_DEFINITION
 
+#include <emscripten/emmalloc.h>
+
 #define CIMGUI_DEFINE_ENUMS_AND_STRUCTS
 
 #define SOKOL_IMPL
 #define SOKOL_WGPU
 #define SOKOL_IMGUI_IMPL
 
+#define STB_DS_IMPLEMENTATION
+#define STBDS_NO_SHORT_NAMES
+
+#define STBDS_REALLOC(context,ptr,size) emmalloc_realloc(ptr, size)
+#define STBDS_FREE(context,ptr)         emmalloc_free(ptr)
+
 #include "sokol_gfx.h"
 #include "sokol_app.h"
 #include "sokol_glue.h"
 #include "cimgui.h"
 #include "sokol_imgui.h"
+#include "stb_ds.h"
 
 #include <emscripten.h>
 #include <stdio.h>
@@ -20,6 +29,7 @@
 #include <stdarg.h>
 
 #include "shape.h"
+#include "cache.h"
 #include "generated/compute.h"
 #include "generated/display.h"
 

src/cache.h

@@ -0,0 +1,69 @@
+#include "api.h"
+
+tile_slot_t* cache_evict(scene_t* s)
+{
+    tile_slot_t* best = NULL;
+    uint64_t oldest = UINT64_MAX;
+
+    for(uint32_t i = 0; i < TILE_LAYERS; i++)
+    {
+        tile_slot_t* slot = &s->cache.slots[i];
+        if(slot->key.lod == UINT32_MAX) continue; // LOD == UINT32_MAX means the slot is free.
+
+        //Found a better candidate
+        if(slot->last_used < oldest)
+        {
+            oldest = slot->last_used;
+            best = slot;
+        }
+    }
+    
+    if(best == NULL)
+        return best;
+
+    stbds_hmdel(s->cache.map, best->key);
+
+    s->cache.free_layers[s->cache.free_count] = best->layer;
+    s->cache.free_count++;
+    s->cache.slots[best->layer].key.lod = UINT32_MAX; //Mark the slot as free for the evict scan using lod == UINT32_MAX
+
+    return best;
+}
+uint32_t cache_allocate(scene_t* s)
+{
+    assert(s->cache.free_count > 0);
+    s->cache.free_count--;
+    uint32_t layer = s->cache.free_layers[s->cache.free_count];
+
+    return layer;
+}
+tile_slot_t* cache_search(scene_t* s, int lod, int tx, int ty, uint64_t frame_count)
+{
+    tile_key_t key = { lod, tx, ty };
+    uint32_t index = stbds_hmgeti(s->cache.map, key);
+
+    if(index >= 0)
+    {
+        uint32_t layer = s->cache.map[index].value;
+        tile_slot_t* slot = &s->cache.slots[layer];
+        slot->last_used = frame_count;
+        return slot;
+    }
+    
+    // Evict the least recently used (LRU) slot
+    if(stbds_hmlen(s->cache.map) >= TILE_LAYERS)
+    {
+        tile_slot_t* evict = cache_evict(s);
+    }
+
+    // Allocate a free layer
+    uint32_t layer = cache_allocate(s);
+    tile_slot_t* slot = &s->cache.slots[layer];
+    slot->key.lod = lod; slot->key.tx = tx; slot->key.ty = ty;
+    slot->layer = layer;
+    slot->ready = false;
+    slot->last_used = frame_count;
+
+    stbds_hmput(s->cache.map, key, layer);
+    return slot;
+}

src/generated/display.h

@@ -12,6 +12,8 @@ unsigned char src_shaders_display_wgsl[] = {
   0x64, 0x5f, 0x6d, 0x69, 0x6e, 0x3a, 0x20, 0x76, 0x65, 0x63, 0x32, 0x66,
   0x2c, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x77, 0x6f, 0x72, 0x6c, 0x64,
   0x5f, 0x6d, 0x61, 0x78, 0x3a, 0x20, 0x76, 0x65, 0x63, 0x32, 0x66, 0x2c,
+  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x62, 0x75, 0x66, 0x66, 0x65, 0x72,
+  0x5f, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x3a, 0x20, 0x75, 0x33, 0x32, 0x2c,
   0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 0x0d, 0x0a, 0x40, 0x67, 0x72, 0x6f, 0x75,
   0x70, 0x28, 0x30, 0x29, 0x20, 0x40, 0x62, 0x69, 0x6e, 0x64, 0x69, 0x6e,
   0x67, 0x28, 0x30, 0x29, 0x20, 0x76, 0x61, 0x72, 0x3c, 0x75, 0x6e, 0x69,
@@ -50,4 +52,4 @@ unsigned char src_shaders_display_wgsl[] = {
   0x20, 0x76, 0x65, 0x63, 0x34, 0x28, 0x31, 0x2c, 0x20, 0x30, 0x2c, 0x20,
   0x30, 0x2c, 0x20, 0x31, 0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a
 };
-unsigned int src_shaders_display_wgsl_len = 599;
+unsigned int src_shaders_display_wgsl_len = 623;

src/main.c

@@ -2,13 +2,15 @@
 
 #define COMPUTE_VIEWIDX_segments 0
 #define COMPUTE_VIEWIDX_shapes 1
-#define COMPUTE_VIEWIDX_SDF 2
+#define COMPUTE_VIEWIDX_tiles 2
+#define COMPUTE_VIEWIDX_SDF 3
 
-#define DISPLAY_VIEWIDX_SDF 0
+#define DISPLAY_VIEWIDX_tiles 0
+#define DISPLAY_VIEWIDX_SDF 1
 
 typedef struct display_uniforms {
-    uint32_t frame;
+    uint64_t frame;
-    uint32_t time;
+    uint64_t time;
 
     float mvp[16];
 } display_uniforms;
@@ -22,8 +24,8 @@ typedef struct renderer_t {
 } renderer_t;
 
 typedef struct compute_uniforms {
-    uint32_t frame;
+    uint64_t frame;
-    uint32_t time;
+    uint64_t time;
 
     float pan_x, pan_y;
     float zoom, rotation;
@@ -33,7 +35,7 @@ typedef struct compute_t {
     sg_pipeline pipeline;
     sg_bindings bindings;
 
-    compute_uniforms uniforms
+    compute_uniforms uniforms;
 } compute_t;
 
 typedef struct userdata_t {
@@ -46,16 +48,24 @@ typedef struct userdata_t {
     float zoom, rotation;
 } userdata_t;
 
+static void* _malloc(size_t size, void* userdata)
+{
+    (void) userdata;
+    return emmalloc_malloc(size);
+}
+static void _free(void* ptr, void* userdata)
+{
+    (void) userdata;
+    emmalloc_free(ptr);
+}
+
 static void log_fn(const char* tag, uint32_t log_level, uint32_t log_item_id, const char* message_or_null, uint32_t line_nr, const char* filename_or_null, void* user_data)
 {
     if(log_level < 2) return;
     fprintf(stderr, "[%s - %s]: (%d) %s \nat %s:%d", tag, log_level == 0 ? "FATAL" : log_level == 1 ? "ERROR" : "WARNING", log_item_id, message_or_null, filename_or_null, line_nr);
 }
 
-//#define _SG_LOG_ITEMS \
+static void refresh_SDF(userdata_t* ud)
-    _SG_LOGITEM_XMACRO(COMPUTE_INVALID_BUFFER, "invalid buffer");
-
-static void draw_scene(userdata_t* ud)
 {
     //Get visible tiles
     float view_world_w = sapp_widthf()  * ud->zoom;
@@ -67,21 +77,26 @@ static void draw_scene(userdata_t* ud)
         ud->pan_y + view_world_h * 0.5f,
     };
 
-    sg_buffer seg_buffer = sg_query_view_buffer(ud->compute.bindings.views[0]);
+    if(ud->scene.segment_dirty)
-    //_SG_VALIDATE(seg_buffer.id == SG_INVALID_ID, COMPUTE_INVALID_BUFFER);
+    {
-    sg_update_buffer(seg_buffer, &(sg_range) { .ptr = ud->scene.segments, .size = sizeof(segment_t) * ud->scene.num_segments });
+        sg_buffer seg_buffer = sg_query_view_buffer(ud->compute.bindings.views[COMPUTE_VIEWIDX_segments]);
-
+        sg_update_buffer(seg_buffer, &(sg_range) { .ptr = ud->scene.segments, .size = sizeof(segment_t) * ud->scene.num_segments });
-    sg_buffer shape_buffer = sg_query_view_buffer(ud->compute.bindings.views[1]);
+        ud->scene.segment_dirty = false;
-    //_SG_VALIDATE(shape_buffer.id == SG_INVALID_ID, COMPUTE_INVALID_BUFFER);
+    }
-    sg_update_buffer(shape_buffer, &(sg_range) { .ptr = ud->scene.shapes, .size = sizeof(shape_meta_t) * ud->scene.num_shapes });
+    
+    if(ud->scene.shape_dirty)
+    {
+        sg_buffer shape_buffer = sg_query_view_buffer(ud->compute.bindings.views[COMPUTE_VIEWIDX_shapes]);
+        sg_update_buffer(shape_buffer, &(sg_range) { .ptr = ud->scene.shapes, .size = sizeof(shape_meta_t) * ud->scene.num_shapes });
+        ud->scene.shape_dirty = false;
+    }
 
     sg_begin_pass(&(sg_pass){ .compute = true, .label = "Compute Pass" });
     sg_apply_pipeline(ud->compute.pipeline);
     sg_apply_uniforms(0, &SG_RANGE(ud->renderer.uniforms));
-    sg_apply_uniforms(1, &SG_RANGE(ud->scene.LODs));
     sg_apply_bindings(&ud->compute.bindings);
 
-    sg_dispatch(TILE_SIZE / 8, TILE_SIZE / 8, 1);
+    sg_dispatch((uint32_t) 16 /* num_tiles */ * 16, 16, 1);
 
     sg_end_pass();
 }
@@ -124,8 +139,17 @@ static void init(void* _userdata)
     sg_setup(&(sg_desc){
         .environment = sglue_environment(),
         .logger.func = log_fn,
+        .allocator = {
+            .alloc_fn = _malloc,
+            .free_fn = _free,
+        },
+    });
+    simgui_setup(&(simgui_desc_t){
+        .allocator = {
+            .alloc_fn = _malloc,
+            .free_fn = _free,
+        },
     });
-    simgui_setup(&(simgui_desc_t){0});
 
     ud->scene = (scene_t) {0};
     if(!scene_init(&ud->scene, 128, 1024, (vec2) { 8192.0f, 8192.0f })) return;
@@ -139,40 +163,37 @@ static void init(void* _userdata)
                     .entry = "main",
                 },
                 .views = {
-                    [0] = {
+                    [COMPUTE_VIEWIDX_segments] = {
                         .storage_buffer = {
                             .stage = SG_SHADERSTAGE_COMPUTE,
-                            .wgsl_group1_binding_n = 0,
+                            .wgsl_group1_binding_n = COMPUTE_VIEWIDX_segments,
                             .readonly = true,
                         }
                     },
-                    [1] = {
+                    [COMPUTE_VIEWIDX_segments] = {
                         .storage_buffer = {
                             .stage = SG_SHADERSTAGE_COMPUTE,
-                            .wgsl_group1_binding_n = 1,
+                            .wgsl_group1_binding_n = COMPUTE_VIEWIDX_segments,
                             .readonly = true,
                         }
                     },
-                    [2] = {
+                    [COMPUTE_VIEWIDX_tiles] = {
                         .storage_buffer = {
                             .stage = SG_SHADERSTAGE_COMPUTE,
-                            .wgsl_group1_binding_n = 2,
+                            .wgsl_group1_binding_n = COMPUTE_VIEWIDX_tiles,
                             .readonly = false,
                         }
+                    },
+                    [COMPUTE_VIEWIDX_SDF] = {
+                        .storage_image = {
+                            .stage = SG_SHADERSTAGE_COMPUTE,
+                            .wgsl_group1_binding_n = COMPUTE_VIEWIDX_SDF,
+                            .access_format = SG_PIXELFORMAT_R32F,
+                            .image_type = SG_IMAGETYPE_ARRAY,
+                            .writeonly = true,
+                        }
                     }
                 },
-                .uniform_blocks = {
-                    [0] = {
-                        .size = sizeof(compute_uniforms),
-                        .stage = SG_SHADERSTAGE_COMPUTE,
-                        .wgsl_group0_binding_n = 0,
-                    },
-                    [1] = {
-                        .size = sizeof(tile_ID),
-                        .stage = SG_SHADERSTAGE_COMPUTE,
-                        .wgsl_group0_binding_n = 1,
-                    },
-                },
                 .label = "SDF Compute Shader",
             }),
             .label = "SDF Compute Pipeline",
@@ -187,7 +208,7 @@ static void init(void* _userdata)
                             .storage_buffer = true,
                             .stream_update = true,
                         },
-                        .size = sizeof(segment_t),
+                        .size = sizeof(segment_t) * 1024 * 1024,
                     }),
                 },
             }),
@@ -200,13 +221,26 @@ static void init(void* _userdata)
                             .storage_buffer = true,
                             .stream_update = true,
                         },
-                        .size = sizeof(shape_meta_t),
+                        .size = sizeof(shape_meta_t) * 1024 * 256,
+                    }),
+                },
+            }),
+            .views[COMPUTE_VIEWIDX_tiles] = sg_make_view(&(sg_view_desc) {
+                .label = "Tiles rebuild view",
+                .storage_buffer = {
+                    .buffer = sg_make_buffer(&(sg_buffer_desc) {
+                        .label = "Tiles rebuild buffer",
+                        .usage = {
+                            .storage_buffer = true,
+                            .stream_update = true,
+                        },
+                        .size = sizeof(tile_info_t) * TILE_LAYERS,
                     }),
                 },
             }),
             .views[COMPUTE_VIEWIDX_SDF] = sg_make_view(&(sg_view_desc) {
                 .label = "SDF tiles view",
-                .storage_image = {
+                .texture = {
                     .image = sg_make_image(&(sg_image_desc) {
                         .height = TILE_SIZE,
                         .width = TILE_SIZE,
@@ -216,6 +250,7 @@ static void init(void* _userdata)
                         .type = SG_IMAGETYPE_ARRAY,
                         .usage = { .storage_image = true },
                     }),
+                    .slices = { .base = 0, .count = TILE_LAYERS },
                 },
             }),
         }
@@ -243,10 +278,17 @@ static void init(void* _userdata)
                     .entry = "fs_main",
                 },
                 .views = {
+                    [DISPLAY_VIEWIDX_tiles] = {
+                        .storage_buffer = {
+                            .stage = SG_SHADERSTAGE_FRAGMENT,
+                            .wgsl_group1_binding_n = DISPLAY_VIEWIDX_tiles,
+                            .readonly = true,
+                        }
+                    }
                     [DISPLAY_VIEWIDX_SDF] = {
                         .storage_buffer = {
                             .stage = SG_SHADERSTAGE_FRAGMENT,
-                            .wgsl_group1_binding_n = 1,
+                            .wgsl_group1_binding_n = DISPLAY_VIEWIDX_SDF,
                             .readonly = true,
                         }
                     }
@@ -256,12 +298,7 @@ static void init(void* _userdata)
                         .size = sizeof(display_uniforms),
                         .stage = SG_SHADERSTAGE_VERTEX,
                         .wgsl_group0_binding_n = 0,
-                    },
+                    }
-                    [1] = {
-                        .size = sizeof(tile_ID),
-                        .stage = SG_SHADERSTAGE_VERTEX,
-                        .wgsl_group0_binding_n = 1,
-                    },
                 },
                 .label = "Display Shader",
             }),
@@ -289,14 +326,14 @@ static void init(void* _userdata)
                 .data = SG_RANGE(quad),
             }),
             .views = {
-                [DISPLAY_VIEWIDX_SDF] = ud->compute.bindings.views[COMPUTE_VIEWIDX_SDF]
+                [DISPLAY_VIEWIDX_tiles] = ud->compute.bindings.views[COMPUTE_VIEWIDX_tiles],
+                [DISPLAY_VIEWIDX_SDF] = ud->compute.bindings.views[COMPUTE_VIEWIDX_SDF],
             }
         }
     };
 
     add_rectangle(&ud->scene, (vec2) { 200.0f, 100.0f }, (vec2) { 100.0f, 150.0f });
     add_circle(&ud->scene, (vec2) { 400.0f, 300.0f }, 125.0f);
-    scene_add_tile(&ud->scene, (tile_ID) { .lod = 0, .tile = TILE_SIZE, .bounds = { 0.0f, 0.0f, 1024.0f, 1024.0f } });
 }
 
 static void cleanup(void* _userdata)
@@ -304,7 +341,7 @@ static void cleanup(void* _userdata)
     userdata_t* ud = (userdata_t*) _userdata;
 
     scene_shutdown(&ud->scene);
-    free(ud);
+    emmalloc_free(ud);
 
     simgui_shutdown();
     sg_shutdown();
@@ -344,7 +381,7 @@ static void event(const sapp_event* event, void* _userdata)
 
 sapp_desc sokol_main(int argc, char* argv[])
 {
-    userdata_t* userdata = (userdata_t*) malloc(sizeof(userdata_t));
+    userdata_t* userdata = (userdata_t*) emmalloc_malloc(sizeof(userdata_t));
 
     (void)argc;
     (void)argv;
@@ -355,6 +392,10 @@ sapp_desc sokol_main(int argc, char* argv[])
         .cleanup_userdata_cb = cleanup,
         .event_userdata_cb = event,
         .window_title = "Sokol",
+        .allocator = {
+            .alloc_fn = _malloc,
+            .free_fn = _free,
+        },
         .logger.func = log_fn,
     };
 }

src/shaders/compute.wgsl

@@ -1,4 +1,3 @@
-// ---------- Bindings (exactly as you defined) ----------
 struct Segment {
     p0: vec2f,
     p1: vec2f,
@@ -11,25 +10,18 @@ struct ShapeMeta {
     bbox_min: vec2f,
     bbox_max: vec2f,
 }
+struct Tile {
+    world_min: vec2f,
+    world_max: vec2f,
+    layer: u32,
+};
 
 @group(1) @binding(0) var<storage> segments: array<Segment>;
 @group(1) @binding(1) var<storage> shapes: array<ShapeMeta>;
-@group(1) @binding(2) var<storage, read_write> sdf_buffer: array<f32>;
+@group(1) @binding(2) var<storage> tiles: array<Tile>;
-
+@group(1) @binding(3) var sdf_buffer: texture_storage_2d_array<r32float, write>;
-struct Uniforms {
-    pan: vec2f,
-    zoom: f32,
-};
-struct TileParams {
-    lod: u32,
-    tile_size: f32,
-    world_min: vec2f,
-    world_max: vec2f,
-};
-
-@group(0) @binding(0) var<uniform> uniforms: Uniforms;
-@group(0) @binding(1) var<uniform> tile_params: TileParams;
 
+override TILE_SIZE: u32 = 256u;
 
 // ---------- Cubic Bézier helpers ----------
 fn eval_cubic(t: f32, p0: vec2f, p1: vec2f, p2: vec2f, p3: vec2f) -> vec2f {
@@ -174,15 +166,19 @@ fn ray_intersections_cubic(p: vec2f, seg: Segment) -> i32 {
 
 // ---------- Main compute shader ----------
 @compute @workgroup_size(8, 8)
-fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
-    let tile_dim = vec2f(tile_params.tile_size);
+    let tile_idx = wg.x / 16u;
-    if id.x >= u32(tile_dim.x) || id.y >= u32(tile_dim.y) {
+    let tile = tiles[tile_idx];
+
+    let px = (wg.x % 16u) * 8u + local.x;
+    let py = wg.y * 8u + local.y;
+
+    if (px >= TILE_SIZE || py >= TILE_SIZE) {
         return;
     }
 
-    // Pixel center UV, then map to world space
+    let uv = (vec2f(px, py) + 0.5) / f32(TILE_SIZE);
-    let uv = (vec2f(id.xy) + 0.5) / tile_dim;
+    let world = mix(tile.world_min, tile.world_max, uv);
-    let world = mix(tile_params.world_min, tile_params.world_max, uv);
 
     var minDist: f32 = 1e20;
     var winding: i32 = 0;
@@ -213,5 +209,5 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>) {
     let sign = select(1.0, -1.0, winding != 0i);
     let sdf = sign * minDist;
 
-    sdf_buffer[id.y * u32(tile_params.tile_size) + id.x] = sdf;
+    textureStore(sdf_buffer, vec2(px, py), tile.layer, vec4f(sdf, 0.0, 0.0, 1.0));
 }

src/shaders/display.wgsl

@@ -7,6 +7,7 @@ struct TileParams {
     tile_size: f32,
     world_min: vec2f,
     world_max: vec2f,
+    buffer_index: u32,
 }
 
 @group(0) @binding(0) var<uniform> uniforms: Uniforms;

src/shape.h

@@ -7,42 +7,56 @@ typedef struct box_t {
     float min_x, min_y, max_x, max_y;
 } box_t;
 
+#define BOX_INTERSECTS(a, b) !(b.min_x > a.max_x \
+        || b.max_x < a.min_x \
+        || b.min_y > a.max_y \
+        || b.max_y < a.min_y) 
+
 typedef struct segment_t {
     vec2 p0, p1, p2, p3;    // cubic Bézier: start, control1, control2, end
 } segment_t;
 
 typedef struct shape_meta_t {
-    int   start_segment;      // index into global segments array
+    uint32_t start_segment;      // index into global segments array
-    int   segment_count;      // number of segments forming this closed loop
+    uint32_t segment_count;      // number of segments forming this closed loop
-    float bbox_min_x, bbox_min_y, bbox_max_x, bbox_max_y;
+    box_t    aabb;
 } shape_meta_t;
 
-#define TILE_SIZE 128   // pixels per tile (same for all LODs)
+#define TILE_SIZE 256   // pixels per tile (same for all LODs)
-#define TILE_LAYERS 256 // 16 MB buffer
+#define TILE_LAYERS 128 // texture_2d_array length with TILE_SIZE dimensions in r32float pixel format (TILE_SIZE² * 4 * TILE_LAYERS bytes) = 32 MB buffer
 
-#define _INTERSECTS(a, b) !(b.min_x > a.max_x \
+typedef struct __attribute__((packed)) tile_key_t {
-        || b.max_x < a.min_x \
+    uint32_t lod, tx, ty;
-        || b.min_y > a.max_y \
+} tile_key_t;
-        || b.max_y < a.min_y) 
 
 // Tile descriptor (separated from tile_cache_entry to be sent to the GPU)
-typedef struct tile_ID {
+typedef struct tile_t {
-    uint32_t lod;
+    tile_key_t key;
-    uvec2    tile;
+    float texel_size, tile_world;
-
+} tile_t;
-    // Cached tile bounds
-    box_t bounds;
-} tile_ID;
 
 typedef struct LOD_t {
-    tile_ID* tiles;
+    tile_t* tiles;
 
     uvec2 dimension; //Tile amount per dimension
     vec2 range; //Zoom min-max range
 } LOD_t;
 
+typedef struct tile_slot_t {
+    uint64_t last_used; //Last frame number for LURA
+    tile_key_t key;
+    uint32_t layer;
+    bool ready;
+} tile_slot_t;
+
+// Used for storage buffer/uniform upload
+typedef struct tile_info_t {
+    box_t bounds;
+    uint32_t layer;
+} tile_info_t;
+
 typedef struct scene_t {
-    vec2 world_size;
+    vec2          world_size;
 
     uint32_t      num_shapes;
     uint32_t      max_shapes;
@@ -53,35 +67,44 @@ typedef struct scene_t {
     segment_t*    segments;
 
     //Theorically, the LOD amount and tile per LOD are computable, so it's not relevant to store them.
-    uint32_t max_LOD;
+    uint32_t      num_LOD;
-    LOD_t* LODs;
+    LOD_t*        LODs;
+    
+    sg_view       texture_cache;
 
-    sg_view texture_cache;
+    struct {
+        tile_slot_t slots[TILE_LAYERS];
+        uint32_t    lru_head, lru_tail;
+
+        struct { tile_key_t key; uint32_t value; }* map;
+
+        uint32_t free_layers[TILE_LAYERS];
+        uint32_t free_count;
+    } cache;
+
+    bool shape_dirty, segment_dirty;
 } scene_t;
 
-typedef int (*tile_callback)(scene_t* s, tile_ID* tile);
-
 // Initialise a scene_t with a given capacity for shapes and segments.
 static int scene_init(scene_t* s, int init_shape_cap, int init_seg_cap, vec2 world_size) {
     s->num_shapes = 0;
     s->max_shapes = init_shape_cap;
-    s->shapes = (shape_meta_t*) malloc(init_shape_cap * sizeof(shape_meta_t));
+    s->shapes = (shape_meta_t*) emmalloc_malloc(init_shape_cap * sizeof(shape_meta_t));
 
     s->num_segments = 0;
     s->max_segments = init_seg_cap;
-    s->segments = (segment_t*) malloc(init_seg_cap * sizeof(segment_t));
+    s->segments = (segment_t*) emmalloc_malloc(init_seg_cap * sizeof(segment_t));
 
-    const uint32_t max_LOD = max(ceilf(log2(world_size.x / TILE_SIZE)), ceilf(log2(world_size.y / TILE_SIZE)));
+    const uint32_t max_LOD = fmax(ceilf(log2(world_size.x / TILE_SIZE)), ceilf(log2(world_size.y / TILE_SIZE)));
-
+    s->num_LOD = max_LOD;
-    s->max_LOD = max_LOD;
+    s->LODs = emmalloc_malloc(max_LOD * sizeof(LOD_t));
-    s->LODs = malloc(max_LOD * sizeof(LOD_t));
     for(int i = 0; i < max_LOD; ++i)
     {
         const float factor = 1 << i;
         s->LODs[i].dimension = (uvec2) { (uint32_t) ceilf(world_size.x * factor / TILE_SIZE), (uint32_t) ceilf(world_size.y * factor / TILE_SIZE) };
 
         const uint32_t tiles_count = s->LODs[i].dimension.x * s->LODs[i].dimension.y;
-        s->LODs[i].tiles = malloc(sizeof(tile_ID) * tiles_count);
+        s->LODs[i].tiles = emmalloc_malloc(sizeof(tile_t) * tiles_count);
 
         s->LODs[i].range = (vec2) {  }; //TODO
 
@@ -91,29 +114,28 @@ static int scene_init(scene_t* s, int init_shape_cap, int init_seg_cap, vec2 wor
             x++;
             if(x >= s->LODs[i].dimension.x) { x = 0; y++; }
 
-            s->LODs[i].tiles[j] = (tile_ID) { .bounds = { x * factor, y * factor, (x + 1) * factor, (y + 1) * factor }, .lod = i, .tile = { x, y } };
+            //s->LODs[i].tiles[j] = (tile_t) { .lod = i, .tile = { x, y }, . };
         }
     }
 
+    s->cache.map = NULL;
+    s->cache.free_count = TILE_LAYERS;
+    for(uint32_t i = 0; i < TILE_LAYERS; i++)
+    {
+        s->cache.free_layers[i] = i;
+
+        s->cache.slots[i].key = (tile_key_t) { UINT32_MAX, 0, 0 };
+        s->cache.slots[i].last_used = UINT32_MAX;
+        s->cache.slots[i].layer = UINT32_MAX;
+        s->cache.slots[i].ready = false;
+    }
+
     s->world_size = world_size;
 
     if (!s->shapes || !s->segments || !s->LODs) return 0;   // allocation failure
     return 1;
 }
 
-//Viewport bounds in *world* space
-static void scene_for_each_tiles(scene_t* s, box_t viewport, tile_callback callback)
-{
-    const uint32_t lod = 0;
-    const uint32_t count = s->LODs[lod].dimension.x * s->LODs[lod].dimension.y;
-    for(uint32_t i = 0; i < count; ++i)
-    {
-        tile_ID tile = s->LODs[lod].tiles[i];
-        if(_INTERSECTS(tile.bounds, viewport))
-            callback(s, &tile);
-    }
-}
-
 // Append a segment, returns its index. (Reallocs if needed, simplified.)
 static int scene_add_segment(scene_t* s, segment_t seg) {
     if (s->num_segments >= s->max_segments) {
@@ -145,15 +167,16 @@ static int scene_add_shape(scene_t* s, shape_meta_t meta) {
 }
 
 static int scene_shutdown(scene_t* s) {
-    free(s->segments);
+    emmalloc_free(s->segments);
-    free(s->shapes);
+    emmalloc_free(s->shapes);
 
-    for(uint32_t i = 0; i < s->max_LOD; ++i)
+    for(uint32_t i = 0; i < s->num_LOD; ++i)
-        free(s->LODs[i].tiles);
+        emmalloc_free(s->LODs[i].tiles);
 
-    free(s->LODs);
+    emmalloc_free(s->LODs);
+    
 
-    free(s);
+    emmalloc_free(s);
     return 1;
 }
 
@@ -205,7 +228,7 @@ int add_rectangle(scene_t* s, vec2 center, vec2 size) {
     shape_meta_t meta;
     meta.start_segment = start;
     meta.segment_count = 4;
-    compute_bbox(s->segments, start, 4, &meta.bbox_min_x, &meta.bbox_min_y, &meta.bbox_max_x, &meta.bbox_max_y);
+    compute_bbox(s->segments, start, 4, &meta.aabb.min_x, &meta.aabb.min_y, &meta.aabb.max_x, &meta.aabb.max_y);
     return scene_add_shape(s, meta);
 }
 
@@ -256,6 +279,6 @@ int add_circle(scene_t* s, vec2 center, float radius) {
     shape_meta_t meta;
     meta.start_segment = start;
     meta.segment_count = 4;
-    compute_bbox(s->segments, start, 4, &meta.bbox_min_x, &meta.bbox_min_y, &meta.bbox_max_x, &meta.bbox_max_y);
+    compute_bbox(s->segments, start, 4, &meta.aabb.min_x, &meta.aabb.min_y, &meta.aabb.max_x, &meta.aabb.max_y);
     return scene_add_shape(s, meta);
 }