Remove tiling and LRU caching, add circles basic shape to improve overall performances.

This commit is contained in:
2026-05-18 16:58:30 +02:00
parent dc708de354
commit 8ea6a0bf3e
11 changed files with 1252 additions and 988 deletions

View File

@@ -1,10 +1,24 @@
#ifndef API_DEFINITION
#define API_DEFINITION
#include <emscripten.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <stdarg.h>
#include <emscripten/emmalloc.h>
#define CIMGUI_DEFINE_ENUMS_AND_STRUCTS
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 < 3) return;
fprintf(stderr, "[%s - %s]: (%d) %s \tat %s:%d\r\n", tag, log_level == 0 ? "FATAL" : log_level == 1 ? "ERROR" : log_level == 2 ? "WARNING" : "INFO", log_item_id, message_or_null, filename_or_null, line_nr);
}
#define SOKOL_ASSERT(x) ((void)((x) || (log_fn("ASSERT", 1, 1, "Assertion failed", __LINE__, __FILE__, NULL),0)))
#define SOKOL_IMPL
#define SOKOL_WGPU
#define SOKOL_IMGUI_IMPL
@@ -15,6 +29,22 @@
#define STBDS_REALLOC(context,ptr,size) emmalloc_realloc(ptr, size)
#define STBDS_FREE(context,ptr) emmalloc_free(ptr)
#define COMPUTE_VIEWIDX_segments 0
#define COMPUTE_VIEWIDX_shapes 1
#define COMPUTE_VIEWIDX_circles 2
#define COMPUTE_VIEWIDX_tiles 3
#define COMPUTE_VIEWIDX_indices 4
#define COMPUTE_VIEWIDX_SDF 5
#define DISPLAY_VIEWIDX_tiles 0
#define DISPLAY_VIEWIDX_SDF 1
#define DISPLAY_VIEWIDX_Sampler 2
static float clampf(float x, float a, float b)
{
return x < a ? a : (x > b ? b : x);
}
#include "sokol_gfx.h"
#include "sokol_app.h"
#include "sokol_glue.h"
@@ -22,14 +52,9 @@
#include "sokol_imgui.h"
#include "stb_ds.h"
#include <emscripten.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <stdarg.h>
#include "shape.h"
#include "cache.h"
//#include "cache.h"
#include "pool.h"
#include "generated/compute.h"
#include "generated/display.h"

View File

@@ -37,10 +37,9 @@ uint32_t cache_allocate(scene_t* s)
return layer;
}
tile_slot_t* cache_search(scene_t* s, int lod, int tx, int ty, uint64_t frame_count)
tile_slot_t* cache_search(scene_t* s, tile_key_t* key, uint64_t frame_count)
{
tile_key_t key = { lod, tx, ty };
uint32_t index = stbds_hmgeti(s->cache.map, key);
uint32_t index = stbds_hmgeti(s->cache.map, *key);
if(index >= 0)
{
@@ -59,11 +58,50 @@ tile_slot_t* cache_search(scene_t* s, int lod, int tx, int ty, uint64_t frame_co
// 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->key.lod = key->lod; slot->key.tx = key->tx; slot->key.ty = key->ty;
slot->layer = layer;
slot->ready = false;
slot->state = TILE_STATE_DIRTY;
slot->last_used = frame_count;
stbds_hmput(s->cache.map, key, layer);
stbds_hmput(s->cache.map, *key, layer);
return slot;
}
// Works in 3 steps, first we select every required tiles and store the mising one from the cache
// Then we evict the LRU cache until we have enough space for the missing tiles
// Finally we allocate the missing tiles
static int cache_query(scene_t* s, uint32_t lod, box_t* box, tile_key_t** buffer)
{
uint64_t frame = sapp_frame_count();
tile_key_t tile_key = { lod, 0, 0 };
for(uint32_t ty = box->min_y; ty < box->max_y; ty++)
{
tile_key.ty = ty;
for(uint32_t tx = box->min_x; tx < box->max_x; tx++)
{
tile_key.tx = tx;
stbds_hmget(s->cache.map, tile_key);
}
}
}
static int cache_cull(scene_t* s, uint32_t tile_count)
{
uint32_t count = 0;
for(uint32_t i = 0; i < tile_count; i++)
{
tile_slot_t slot = s->cache.slots[i];
tile_task_t* tile = &s->cache.tiles[i];
tile->bounds.min_x = slot.key.tx * s->LODs[slot.key.lod].texel_size;
tile->bounds.min_y = slot.key.ty * s->LODs[slot.key.lod].texel_size;
tile->bounds.max_x = (slot.key.tx + 1) * s->LODs[slot.key.lod].texel_size;
tile->bounds.max_y = (slot.key.ty + 1) * s->LODs[slot.key.lod].texel_size;
tile->layer = slot.layer;
uint32_t start = count;
for(uint32_t j = 0; j < s->num_shapes; j++)
{
if(BOX_INTERSECTS(tile->bounds, s->shapes[j].aabb)) s->cache.indices[count++] = j;
}
tile->offset = start; tile->count = count - start;
}
}

File diff suppressed because it is too large Load Diff

View File

@@ -1,55 +1,59 @@
unsigned char src_shaders_display_wgsl[] = {
0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x55, 0x6e, 0x69, 0x66, 0x6f,
0x72, 0x6d, 0x73, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x70,
0x61, 0x6e, 0x3a, 0x20, 0x76, 0x65, 0x63, 0x32, 0x66, 0x2c, 0x0d, 0x0a,
0x20, 0x20, 0x20, 0x20, 0x7a, 0x6f, 0x6f, 0x6d, 0x3a, 0x20, 0x66, 0x33,
0x32, 0x2c, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 0x73, 0x74, 0x72, 0x75, 0x63,
0x74, 0x20, 0x54, 0x69, 0x6c, 0x65, 0x50, 0x61, 0x72, 0x61, 0x6d, 0x73,
0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x64, 0x3a,
0x20, 0x75, 0x33, 0x32, 0x2c, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x74,
0x69, 0x6c, 0x65, 0x5f, 0x73, 0x69, 0x7a, 0x65, 0x3a, 0x20, 0x66, 0x33,
0x32, 0x2c, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x77, 0x6f, 0x72, 0x6c,
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,
0x66, 0x6f, 0x72, 0x6d, 0x3e, 0x20, 0x75, 0x6e, 0x69, 0x66, 0x6f, 0x72,
0x6d, 0x73, 0x3a, 0x20, 0x55, 0x6e, 0x69, 0x66, 0x6f, 0x72, 0x6d, 0x73,
0x3b, 0x0d, 0x0a, 0x40, 0x67, 0x72, 0x6f, 0x75, 0x70, 0x28, 0x30, 0x29,
0x20, 0x40, 0x62, 0x69, 0x6e, 0x64, 0x69, 0x6e, 0x67, 0x28, 0x31, 0x29,
0x20, 0x76, 0x61, 0x72, 0x3c, 0x75, 0x6e, 0x69, 0x66, 0x6f, 0x72, 0x6d,
0x3e, 0x20, 0x74, 0x69, 0x6c, 0x65, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d,
0x73, 0x3a, 0x20, 0x54, 0x69, 0x6c, 0x65, 0x50, 0x61, 0x72, 0x61, 0x6d,
0x73, 0x3b, 0x0d, 0x0a, 0x0d, 0x0a, 0x40, 0x76, 0x65, 0x72, 0x74, 0x65,
0x78, 0x20, 0x66, 0x6e, 0x20, 0x76, 0x73, 0x5f, 0x6d, 0x61, 0x69, 0x6e,
0x28, 0x40, 0x62, 0x75, 0x69, 0x6c, 0x74, 0x69, 0x6e, 0x28, 0x76, 0x65,
0x72, 0x74, 0x65, 0x78, 0x5f, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x29, 0x20,
0x76, 0x65, 0x72, 0x74, 0x65, 0x78, 0x5f, 0x69, 0x6e, 0x64, 0x65, 0x78,
0x20, 0x3a, 0x20, 0x75, 0x33, 0x32, 0x29, 0x20, 0x2d, 0x3e, 0x20, 0x40,
0x62, 0x75, 0x69, 0x6c, 0x74, 0x69, 0x6e, 0x28, 0x70, 0x6f, 0x73, 0x69,
0x74, 0x69, 0x6f, 0x6e, 0x29, 0x20, 0x76, 0x65, 0x63, 0x34, 0x66, 0x20,
0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x70,
0x6f, 0x73, 0x20, 0x3d, 0x20, 0x61, 0x72, 0x72, 0x61, 0x79, 0x28, 0x0d,
0x0a, 0x20, 0x20, 0x20, 0x20, 0x76, 0x65, 0x63, 0x32, 0x28, 0x20, 0x30,
0x2e, 0x30, 0x2c, 0x20, 0x20, 0x30, 0x2e, 0x35, 0x29, 0x2c, 0x0d, 0x0a,
0x20, 0x20, 0x20, 0x20, 0x76, 0x65, 0x63, 0x32, 0x28, 0x2d, 0x30, 0x2e,
0x35, 0x2c, 0x20, 0x2d, 0x30, 0x2e, 0x35, 0x29, 0x2c, 0x0d, 0x0a, 0x20,
0x20, 0x20, 0x20, 0x76, 0x65, 0x63, 0x32, 0x28, 0x20, 0x30, 0x2e, 0x35,
0x2c, 0x20, 0x2d, 0x30, 0x2e, 0x35, 0x29, 0x0d, 0x0a, 0x20, 0x20, 0x29,
0x3b, 0x0d, 0x0a, 0x0d, 0x0a, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72,
0x6e, 0x20, 0x76, 0x65, 0x63, 0x34, 0x66, 0x28, 0x70, 0x6f, 0x73, 0x5b,
0x76, 0x65, 0x72, 0x74, 0x65, 0x78, 0x5f, 0x69, 0x6e, 0x64, 0x65, 0x78,
0x5d, 0x2c, 0x20, 0x30, 0x2c, 0x20, 0x31, 0x29, 0x3b, 0x0d, 0x0a, 0x7d,
0x72, 0x6d, 0x73, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x74,
0x69, 0x6d, 0x65, 0x3a, 0x20, 0x75, 0x33, 0x32, 0x2c, 0x0d, 0x0a, 0x20,
0x20, 0x20, 0x20, 0x66, 0x72, 0x61, 0x6d, 0x65, 0x3a, 0x20, 0x75, 0x33,
0x32, 0x2c, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6d, 0x76, 0x70, 0x3a,
0x20, 0x6d, 0x61, 0x74, 0x34, 0x78, 0x34, 0x66, 0x2c, 0x0d, 0x0a, 0x7d,
0x0d, 0x0a, 0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x54, 0x69, 0x6c,
0x65, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x77, 0x6f, 0x72,
0x6c, 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, 0x6f, 0x66, 0x66, 0x73, 0x65,
0x74, 0x3a, 0x20, 0x75, 0x33, 0x32, 0x2c, 0x0d, 0x0a, 0x20, 0x20, 0x20,
0x20, 0x63, 0x69, 0x72, 0x63, 0x6c, 0x65, 0x5f, 0x63, 0x6f, 0x75, 0x6e,
0x74, 0x3a, 0x20, 0x75, 0x33, 0x32, 0x2c, 0x0d, 0x0a, 0x20, 0x20, 0x20,
0x20, 0x63, 0x6f, 0x75, 0x6e, 0x74, 0x3a, 0x20, 0x75, 0x33, 0x32, 0x2c,
0x0d, 0x0a, 0x7d, 0x3b, 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, 0x66, 0x6f, 0x72, 0x6d, 0x3e, 0x20, 0x75, 0x6e, 0x69, 0x66, 0x6f,
0x72, 0x6d, 0x73, 0x3a, 0x20, 0x55, 0x6e, 0x69, 0x66, 0x6f, 0x72, 0x6d,
0x73, 0x3b, 0x0d, 0x0a, 0x0d, 0x0a, 0x40, 0x67, 0x72, 0x6f, 0x75, 0x70,
0x28, 0x31, 0x29, 0x20, 0x40, 0x62, 0x69, 0x6e, 0x64, 0x69, 0x6e, 0x67,
0x28, 0x30, 0x29, 0x20, 0x76, 0x61, 0x72, 0x3c, 0x73, 0x74, 0x6f, 0x72,
0x61, 0x67, 0x65, 0x3e, 0x20, 0x74, 0x69, 0x6c, 0x65, 0x73, 0x3a, 0x20,
0x61, 0x72, 0x72, 0x61, 0x79, 0x3c, 0x54, 0x69, 0x6c, 0x65, 0x3e, 0x3b,
0x0d, 0x0a, 0x40, 0x67, 0x72, 0x6f, 0x75, 0x70, 0x28, 0x31, 0x29, 0x20,
0x40, 0x62, 0x69, 0x6e, 0x64, 0x69, 0x6e, 0x67, 0x28, 0x31, 0x29, 0x20,
0x76, 0x61, 0x72, 0x20, 0x73, 0x64, 0x66, 0x5f, 0x62, 0x75, 0x66, 0x66,
0x65, 0x72, 0x20, 0x3a, 0x20, 0x74, 0x65, 0x78, 0x74, 0x75, 0x72, 0x65,
0x5f, 0x73, 0x74, 0x6f, 0x72, 0x61, 0x67, 0x65, 0x5f, 0x32, 0x64, 0x3c,
0x72, 0x31, 0x36, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2c, 0x20, 0x72, 0x65,
0x61, 0x64, 0x3e, 0x3b, 0x0d, 0x0a, 0x40, 0x67, 0x72, 0x6f, 0x75, 0x70,
0x28, 0x31, 0x29, 0x20, 0x40, 0x62, 0x69, 0x6e, 0x64, 0x69, 0x6e, 0x67,
0x28, 0x32, 0x29, 0x20, 0x76, 0x61, 0x72, 0x20, 0x73, 0x64, 0x66, 0x5f,
0x73, 0x61, 0x6d, 0x70, 0x6c, 0x65, 0x72, 0x3a, 0x20, 0x73, 0x61, 0x6d,
0x70, 0x6c, 0x65, 0x72, 0x3b, 0x0d, 0x0a, 0x0d, 0x0a, 0x40, 0x76, 0x65,
0x72, 0x74, 0x65, 0x78, 0x20, 0x66, 0x6e, 0x20, 0x76, 0x73, 0x5f, 0x6d,
0x61, 0x69, 0x6e, 0x28, 0x40, 0x6c, 0x6f, 0x63, 0x61, 0x74, 0x69, 0x6f,
0x6e, 0x28, 0x30, 0x29, 0x20, 0x70, 0x6f, 0x73, 0x69, 0x74, 0x69, 0x6f,
0x6e, 0x3a, 0x20, 0x76, 0x65, 0x63, 0x32, 0x66, 0x29, 0x20, 0x2d, 0x3e,
0x20, 0x40, 0x62, 0x75, 0x69, 0x6c, 0x74, 0x69, 0x6e, 0x28, 0x70, 0x6f,
0x73, 0x69, 0x74, 0x69, 0x6f, 0x6e, 0x29, 0x20, 0x76, 0x65, 0x63, 0x34,
0x66, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74,
0x75, 0x72, 0x6e, 0x20, 0x76, 0x65, 0x63, 0x34, 0x66, 0x28, 0x70, 0x6f,
0x73, 0x69, 0x74, 0x69, 0x6f, 0x6e, 0x2c, 0x20, 0x30, 0x2e, 0x30, 0x2c,
0x20, 0x31, 0x2e, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x75, 0x6e, 0x69, 0x66,
0x6f, 0x72, 0x6d, 0x73, 0x2e, 0x6d, 0x76, 0x70, 0x3b, 0x0d, 0x0a, 0x7d,
0x0d, 0x0a, 0x0d, 0x0a, 0x40, 0x66, 0x72, 0x61, 0x67, 0x6d, 0x65, 0x6e,
0x74, 0x20, 0x66, 0x6e, 0x20, 0x66, 0x73, 0x5f, 0x6d, 0x61, 0x69, 0x6e,
0x28, 0x29, 0x20, 0x2d, 0x3e, 0x20, 0x40, 0x6c, 0x6f, 0x63, 0x61, 0x74,
0x69, 0x6f, 0x6e, 0x28, 0x30, 0x29, 0x20, 0x76, 0x65, 0x63, 0x34, 0x66,
0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e,
0x20, 0x76, 0x65, 0x63, 0x34, 0x28, 0x31, 0x2c, 0x20, 0x30, 0x2c, 0x20,
0x30, 0x2c, 0x20, 0x31, 0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a
0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75,
0x72, 0x6e, 0x20, 0x76, 0x65, 0x63, 0x34, 0x66, 0x28, 0x31, 0x2e, 0x30,
0x2c, 0x20, 0x30, 0x2e, 0x30, 0x2c, 0x20, 0x30, 0x2e, 0x30, 0x2c, 0x20,
0x31, 0x2e, 0x30, 0x29, 0x3b, 0x0d, 0x0a, 0x7d
};
unsigned int src_shaders_display_wgsl_len = 623;
unsigned int src_shaders_display_wgsl_len = 668;

View File

@@ -1,16 +1,8 @@
#include "api.h"
#define COMPUTE_VIEWIDX_segments 0
#define COMPUTE_VIEWIDX_shapes 1
#define COMPUTE_VIEWIDX_tiles 2
#define COMPUTE_VIEWIDX_SDF 3
#define DISPLAY_VIEWIDX_tiles 0
#define DISPLAY_VIEWIDX_SDF 1
typedef struct display_uniforms {
uint64_t frame;
uint64_t time;
uint32_t frame;
uint32_t time;
float mvp[16];
} display_uniforms;
@@ -23,19 +15,10 @@ typedef struct renderer_t {
display_uniforms uniforms;
} renderer_t;
typedef struct compute_uniforms {
uint64_t frame;
uint64_t time;
float pan_x, pan_y;
float zoom, rotation;
} compute_uniforms;
typedef struct compute_t {
sg_pipeline pipeline;
sg_bindings bindings;
compute_uniforms uniforms;
sg_pass pass;
} compute_t;
typedef struct userdata_t {
@@ -59,12 +42,6 @@ static void _free(void* ptr, 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);
}
static void refresh_SDF(userdata_t* ud)
{
//Get visible tiles
@@ -77,56 +54,116 @@ static void refresh_SDF(userdata_t* ud)
ud->pan_y + view_world_h * 0.5f,
};
bool dirty = false;
if(ud->scene.segment_dirty)
{
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 });
pool_view_grow(ud->compute.bindings.views[COMPUTE_VIEWIDX_segments], sizeof(segment_t), ud->scene.shapes.num_segments, ud->scene.shapes.segments);
ud->scene.segment_dirty = false;
dirty = true;
}
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 });
pool_view_grow(ud->compute.bindings.views[COMPUTE_VIEWIDX_shapes], sizeof(shape_meta_t), ud->scene.shapes.num_shapes, ud->scene.shapes.shapes);
ud->scene.shape_dirty = false;
dirty = true;
}
if(ud->scene.circle_dirty)
{
pool_view_grow(ud->compute.bindings.views[COMPUTE_VIEWIDX_circles], sizeof(circle_t), ud->scene.circles.num_circles, ud->scene.circles.circles);
ud->scene.circle_dirty = false;
dirty = true;
}
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_bindings(&ud->compute.bindings);
if(dirty)
{
uint32_t tile_count = scene_process_tiles(&ud->scene, ud->pan_x, ud->pan_y, ud->zoom, &ud->compute.bindings);
sg_dispatch((uint32_t) 16 /* num_tiles */ * 16, 16, 1);
if(tile_count > 0)
{
pool_view_grow(ud->compute.bindings.views[COMPUTE_VIEWIDX_tiles], sizeof(tile_task_t), tile_count, ud->scene.cache.tiles);
pool_view_grow(ud->compute.bindings.views[COMPUTE_VIEWIDX_indices], sizeof(uint32_t), ud->scene.cache.num_indices, ud->scene.cache.indices);
sg_end_pass();
ud->compute.bindings.views[COMPUTE_VIEWIDX_SDF] = ud->scene.texture;
sg_begin_pass(&ud->compute.pass);
sg_apply_pipeline(ud->compute.pipeline);
sg_apply_bindings(&ud->compute.bindings);
sg_dispatch(tile_count * 16, 16, 1);
sg_end_pass();
}
dirty = false;
}
}
static sg_pass pass = (sg_pass){
.label = "Render pass",
};
static simgui_frame_desc_t imgui_frame = {0};
static void compute_mvp(userdata_t *ud)
{
const float w = sapp_widthf();
const float h = sapp_heightf();
const float z = ud->zoom;
const float px = ud->pan_x;
const float py = ud->pan_y;
float* mvp = ud->renderer.uniforms.mvp;
mvp[0] = (2.0f / w) * z;
mvp[1] = 0.0f;
mvp[2] = 0.0f;
mvp[3] = (2.0f / w) * px;
mvp[4] = 0.0f;
mvp[5] = (2.0f / h) * z;
mvp[6] = 0.0f;
mvp[7] = (2.0f / h) * py;
mvp[8] = 0.0f;
mvp[9] = 0.0f;
mvp[10] = 0.0f;
mvp[11] = 0.0f;
mvp[12] = 0.0f;
mvp[13] = 0.0f;
mvp[14] = 0.0f;
mvp[15] = 1.0f;
}
static void frame(void* _userdata)
{
userdata_t* ud = (userdata_t*) _userdata;
ud->renderer.uniforms.frame = ud->compute.uniforms.frame = sapp_frame_count();
ud->renderer.uniforms.time = ud->compute.uniforms.time += sapp_frame_duration_unfiltered();
ud->renderer.uniforms.frame = sapp_frame_count();
ud->renderer.uniforms.time += sapp_frame_duration_unfiltered();
draw_scene(ud);
refresh_SDF(ud);
sg_begin_pass(&(sg_pass){
.action = ud->renderer.clear_pass,
.swapchain = sglue_swapchain(),
});
pass.swapchain = sglue_swapchain();
sg_begin_pass(&pass);
simgui_new_frame(&(simgui_frame_desc_t){
.width = sapp_width(),
.height = sapp_height(),
.delta_time = sapp_frame_duration_unfiltered(),
.dpi_scale = sapp_dpi_scale(),
});
imgui_frame.width = sapp_width(),
imgui_frame.height = sapp_height(),
imgui_frame.delta_time = sapp_frame_duration_unfiltered(),
imgui_frame.dpi_scale = sapp_dpi_scale(),
simgui_new_frame(&imgui_frame);
igBegin("Framerate", (bool*) true, ImGuiWindowFlags_NoResize | ImGuiWindowFlags_NoCollapse | ImGuiWindowFlags_NoTitleBar);
igText("%.0f FPS", 1 / sapp_frame_duration_unfiltered());
igText("%.3fms", sapp_frame_duration_unfiltered() * 1000);
igEnd();
ud->renderer.uniforms.frame = sapp_frame_count();
ud->renderer.uniforms.time += (uint32_t) ceil(sapp_frame_duration_unfiltered() / 1000.0f);
sg_apply_pipeline(ud->renderer.pipeline);
sg_apply_uniforms(0, &SG_RANGE(ud->renderer.uniforms));
sg_apply_bindings(&ud->renderer.bindings);
sg_draw(0, 4, 1);
simgui_render();
sg_end_pass();
sg_commit();
@@ -152,7 +189,7 @@ static void init(void* _userdata)
});
ud->scene = (scene_t) {0};
if(!scene_init(&ud->scene, 128, 1024, (vec2) { 8192.0f, 8192.0f })) return;
if(!scene_init(&ud->scene, (vec2) { 8192.0f, 8192.0f })) return;
ud->compute = (compute_t) {
.pipeline = sg_make_pipeline(&(sg_pipeline_desc){
@@ -170,10 +207,17 @@ static void init(void* _userdata)
.readonly = true,
}
},
[COMPUTE_VIEWIDX_segments] = {
[COMPUTE_VIEWIDX_shapes] = {
.storage_buffer = {
.stage = SG_SHADERSTAGE_COMPUTE,
.wgsl_group1_binding_n = COMPUTE_VIEWIDX_segments,
.wgsl_group1_binding_n = COMPUTE_VIEWIDX_shapes,
.readonly = true,
}
},
[COMPUTE_VIEWIDX_circles] = {
.storage_buffer = {
.stage = SG_SHADERSTAGE_COMPUTE,
.wgsl_group1_binding_n = COMPUTE_VIEWIDX_circles,
.readonly = true,
}
},
@@ -181,15 +225,22 @@ static void init(void* _userdata)
.storage_buffer = {
.stage = SG_SHADERSTAGE_COMPUTE,
.wgsl_group1_binding_n = COMPUTE_VIEWIDX_tiles,
.readonly = false,
.readonly = true,
}
},
[COMPUTE_VIEWIDX_indices] = {
.storage_buffer = {
.stage = SG_SHADERSTAGE_COMPUTE,
.wgsl_group1_binding_n = COMPUTE_VIEWIDX_indices,
.readonly = true,
}
},
[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,
.access_format = SG_PIXELFORMAT_R16F,
.image_type = SG_IMAGETYPE_2D,
.writeonly = true,
}
}
@@ -208,7 +259,7 @@ static void init(void* _userdata)
.storage_buffer = true,
.stream_update = true,
},
.size = sizeof(segment_t) * 1024 * 1024,
.size = sizeof(segment_t) * INITIAL_SEGMENTS_SIZE,
}),
},
}),
@@ -221,7 +272,20 @@ static void init(void* _userdata)
.storage_buffer = true,
.stream_update = true,
},
.size = sizeof(shape_meta_t) * 1024 * 256,
.size = sizeof(shape_meta_t) * INITIAL_SHAPE_SIZE,
}),
},
}),
.views[COMPUTE_VIEWIDX_circles] = sg_make_view(&(sg_view_desc) {
.label = "Circles view",
.storage_buffer = {
.buffer = sg_make_buffer(&(sg_buffer_desc) {
.label = "Circles buffer",
.usage = {
.storage_buffer = true,
.stream_update = true,
},
.size = sizeof(circle_t) * INITIAL_CIRCLE_SIZE,
}),
},
}),
@@ -234,26 +298,26 @@ static void init(void* _userdata)
.storage_buffer = true,
.stream_update = true,
},
.size = sizeof(tile_info_t) * TILE_LAYERS,
.size = sizeof(tile_task_t) * 256,
}),
},
}),
.views[COMPUTE_VIEWIDX_SDF] = sg_make_view(&(sg_view_desc) {
.label = "SDF tiles view",
.texture = {
.image = sg_make_image(&(sg_image_desc) {
.height = TILE_SIZE,
.width = TILE_SIZE,
.label = "SDF tiles texture",
.num_slices = TILE_LAYERS,
.pixel_format = SG_PIXELFORMAT_R32F,
.type = SG_IMAGETYPE_ARRAY,
.usage = { .storage_image = true },
.views[COMPUTE_VIEWIDX_indices] = sg_make_view(&(sg_view_desc) {
.label = "Culling indices view",
.storage_buffer = {
.buffer = sg_make_buffer(&(sg_buffer_desc) {
.label = "Culling indices buffer",
.usage = {
.storage_buffer = true,
.stream_update = true,
},
.size = sizeof(uint32_t) * INITIAL_INDICES_SIZE,
}),
.slices = { .base = 0, .count = TILE_LAYERS },
},
}),
}
.views[COMPUTE_VIEWIDX_SDF] = (sg_view) {0},
},
.pass = { .compute = true, .label = "Compute Pass" }
};
const vec2 quad[4] = {
@@ -267,7 +331,7 @@ static void init(void* _userdata)
};
ud->renderer = (renderer_t) {
.pipeline = sg_make_pipeline(&(sg_pipeline_desc){
.pipeline = sg_make_pipeline(&(sg_pipeline_desc) {
.shader = sg_make_shader(&(sg_shader_desc) {
.vertex_func = {
.source = (const char*) src_shaders_display_wgsl,
@@ -277,6 +341,9 @@ static void init(void* _userdata)
.source = (const char*) src_shaders_display_wgsl,
.entry = "fs_main",
},
.attrs[0] = {
.base_type = SG_SHADERATTRBASETYPE_FLOAT,
},
.views = {
[DISPLAY_VIEWIDX_tiles] = {
.storage_buffer = {
@@ -284,21 +351,36 @@ static void init(void* _userdata)
.wgsl_group1_binding_n = DISPLAY_VIEWIDX_tiles,
.readonly = true,
}
}
},
[DISPLAY_VIEWIDX_SDF] = {
.storage_buffer = {
.texture = {
.image_type = SG_IMAGETYPE_ARRAY,
.stage = SG_SHADERSTAGE_FRAGMENT,
.wgsl_group1_binding_n = DISPLAY_VIEWIDX_SDF,
.readonly = true,
.sample_type = SG_IMAGESAMPLETYPE_FLOAT,
}
}
},
.samplers = {
[0] = {
.sampler_type = SG_SAMPLERTYPE_FILTERING,
.stage = SG_SHADERSTAGE_FRAGMENT,
.wgsl_group1_binding_n = DISPLAY_VIEWIDX_Sampler,
}
},
.texture_sampler_pairs = {
[0] = {
.sampler_slot = 0,
.view_slot = DISPLAY_VIEWIDX_SDF,
.stage = SG_SHADERSTAGE_FRAGMENT,
}
},
.uniform_blocks = {
[0] = {
.size = sizeof(display_uniforms),
.stage = SG_SHADERSTAGE_VERTEX,
.wgsl_group0_binding_n = 0,
}
},
},
.label = "Display Shader",
}),
@@ -309,9 +391,6 @@ static void init(void* _userdata)
},
.label = "Render pipeline",
}),
.clear_pass = (sg_pass_action) {
.colors[0] = { .clear_value = { 0.0f, 0.0f, 0.0f, 1.0f }, .load_action = SG_LOADACTION_CLEAR }
},
.bindings = {
.index_buffer = sg_make_buffer(&(sg_buffer_desc) {
.label = "Index buffer",
@@ -329,7 +408,17 @@ static void init(void* _userdata)
[DISPLAY_VIEWIDX_tiles] = ud->compute.bindings.views[COMPUTE_VIEWIDX_tiles],
[DISPLAY_VIEWIDX_SDF] = ud->compute.bindings.views[COMPUTE_VIEWIDX_SDF],
}
}
},
.uniforms = {
.frame = 0,
.time = 0,
.mvp = { 0, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, 0, 0 },
},
};
compute_mvp(ud);
pass.action = (sg_pass_action) {
.colors[0] = { .clear_value = { 0.0f, 0.0f, 0.0f, 1.0f }, .load_action = SG_LOADACTION_CLEAR }
};
add_rectangle(&ud->scene, (vec2) { 200.0f, 100.0f }, (vec2) { 100.0f, 150.0f });

36
src/pool.h Normal file
View File

@@ -0,0 +1,36 @@
#include "api.h"
#define MAX_BUFFER_SIZE 1024 * 1024 * 1024 // 1GB
static void pool_buffer_grow(sg_buffer buffer, uint32_t stripe, uint32_t count, void* data)
{
size_t size = sg_query_buffer_size(buffer);
if(size == MAX_BUFFER_SIZE) return sg_update_buffer(buffer, &(sg_range) { data, MAX_BUFFER_SIZE });
if(size < stripe * count)
{
sg_buffer_desc desc = sg_query_buffer_desc(buffer);
while(desc.size < stripe * count) desc.size = fmaxf(MAX_BUFFER_SIZE, size * 2);
sg_uninit_buffer(buffer);
sg_init_buffer(buffer, &desc);
}
sg_update_buffer(buffer, &(sg_range) { data, stripe * count });
}
static void pool_view_grow(sg_view view, uint32_t stripe, uint32_t count, void* data)
{
sg_buffer buffer = sg_query_view_buffer(view);
size_t size = sg_query_buffer_size(buffer);
if(size == MAX_BUFFER_SIZE) return sg_update_buffer(buffer, &(sg_range) { data, MAX_BUFFER_SIZE });
if(size < stripe * count)
{
sg_buffer_desc desc = sg_query_buffer_desc(buffer);
while(desc.size < stripe * count) desc.size = fmaxf(MAX_BUFFER_SIZE, size * 2);
sg_uninit_buffer(buffer);
sg_init_buffer(buffer, &desc);
}
sg_update_buffer(buffer, &(sg_range) { data, stripe * count });
}

View File

@@ -4,22 +4,30 @@ struct Segment {
p2: vec2f,
p3: vec2f,
}
struct ShapeMeta {
start_segment: u32,
struct ShapeMeta {
start_segment: u32, //Offset in the segments buffer
segment_count: u32,
bbox_min: vec2f,
bbox_max: vec2f,
}
struct Circle {
center: vec2f,
radius: f32,
}
struct Tile {
world_min: vec2f,
world_max: vec2f,
layer: u32,
offset: u32,
circle_count: u32,
count: 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> tiles: array<Tile>;
@group(1) @binding(3) var sdf_buffer: texture_storage_2d_array<r32float, write>;
@group(1) @binding(2) var<storage> circles: array<Circle>;
@group(1) @binding(3) var<storage> tiles: array<Tile>;
@group(1) @binding(4) var<storage> indices: array<u32>;
@group(1) @binding(5) var sdf_buffer: texture_storage_2d<r16float, write>;
override TILE_SIZE: u32 = 256u;
@@ -166,7 +174,9 @@ fn ray_intersections_cubic(p: vec2f, seg: Segment) -> i32 {
// ---------- Main compute shader ----------
@compute @workgroup_size(8, 8)
fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) local: vec3<u32>) {
fn main(@builtin(workgroup_id) wg: vec3<u32>, // each workgroup is a tile
@builtin(local_invocation_id) local: vec3<u32>) // each local invocation is a pixel
{
let tile_idx = wg.x / 16u;
let tile = tiles[tile_idx];
@@ -177,37 +187,49 @@ fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_id) loca
return;
}
let uv = (vec2f(px, py) + 0.5) / f32(TILE_SIZE);
let uv = (vec2f(f32(px), f32(py)) + 0.5) / f32(TILE_SIZE);
let world = mix(tile.world_min, tile.world_max, uv);
var minDist: f32 = 1e20;
var sdf: f32 = 1e20;
var winding: i32 = 0;
for (var s = 0u; s < arrayLength(&shapes); s++) {
let shape = shapes[s];
for (var s = 0u; s < tile.count; s++) {
// Use a pre-culled shape range per tile to reduce the per-pixel processing time
let index = indices[tile.offset + s];
// Quick bbox culling in world space
let bbox_min = shape.bbox_min;
let bbox_max = shape.bbox_max;
let dx = max(bbox_min.x - world.x, max(0.0, world.x - bbox_max.x));
let dy = max(bbox_min.y - world.y, max(0.0, world.y - bbox_max.y));
let dist_to_box = sqrt(dx * dx + dy * dy);
if dist_to_box >= minDist {
continue; // This shape cannot improve the distance
// If the index is lower than the circle count, it mean we are processing circles
if(index < tile.circle_count)
{
let c = circles[index];
sdf = min(sdf, length(world - c.center) - c.radius);
}
else
{
let shape = shapes[index];
// Process all segments of the shape
for (var i = shape.start_segment; i < shape.start_segment + shape.segment_count; i++) {
let seg = segments[i];
let d = distance_to_cubic(world, seg);
minDist = min(minDist, d);
/*let bbox_min = shape.bbox_min;
let bbox_max = shape.bbox_max;
let dx = max(bbox_min.x - world.x, max(0.0, world.x - bbox_max.x));
let dy = max(bbox_min.y - world.y, max(0.0, world.y - bbox_max.y));
let dist_to_box = sqrt(dx * dx + dy * dy);
if dist_to_box >= minDist {
continue;
}*/
winding += ray_intersections_cubic(world, seg);
// Process all segments of the shape
for (var i = shape.start_segment; i < shape.start_segment + shape.segment_count; i++) {
let seg = segments[i];
let d = distance_to_cubic(world, seg);
minDist = min(minDist, d);
winding += ray_intersections_cubic(world, seg);
}
let sign = select(1.0, -1.0, winding != 0i);
sdf = min(sdf, sign * minDist);
}
}
let sign = select(1.0, -1.0, winding != 0i);
let sdf = sign * minDist;
textureStore(sdf_buffer, vec2(px, py), tile.layer, vec4f(sdf, 0.0, 0.0, 1.0));
textureStore(sdf_buffer, vec2(u32(world.x), u32(world.y)), vec4f(sdf, 0.0, 0.0, 1.0));
}

View File

@@ -1,28 +1,26 @@
struct Uniforms {
pan: vec2f,
zoom: f32,
time: u32,
frame: u32,
mvp: mat4x4f,
}
struct TileParams {
lod: u32,
tile_size: f32,
struct Tile {
world_min: vec2f,
world_max: vec2f,
buffer_index: u32,
}
offset: u32,
circle_count: u32,
count: u32,
};
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<uniform> tile_params: TileParams;
@vertex fn vs_main(@builtin(vertex_index) vertex_index : u32) -> @builtin(position) vec4f {
const pos = array(
vec2( 0.0, 0.5),
vec2(-0.5, -0.5),
vec2( 0.5, -0.5)
);
@group(1) @binding(0) var<storage> tiles: array<Tile>;
@group(1) @binding(1) var sdf_buffer : texture_storage_2d<r16float, read>;
@group(1) @binding(2) var sdf_sampler: sampler;
return vec4f(pos[vertex_index], 0, 1);
@vertex fn vs_main(@location(0) position: vec2f) -> @builtin(position) vec4f {
return vec4f(position, 0.0, 1.0) * uniforms.mvp;
}
@fragment fn fs_main() -> @location(0) vec4f {
return vec4(1, 0, 0, 1);
}
return vec4f(1.0, 0.0, 0.0, 1.0);
}

View File

@@ -3,6 +3,9 @@
typedef struct uvec2 { uint32_t x, y; } uvec2;
typedef struct vec2 { float x, y; } vec2;
typedef struct ubox_t {
uint32_t min_x, min_y, max_x, max_y;
} ubox_t;
typedef struct box_t {
float min_x, min_y, max_x, max_y;
} box_t;
@@ -12,6 +15,11 @@ typedef struct box_t {
|| b.min_y > a.max_y \
|| b.max_y < a.min_y)
#define CIRCLE_INTERSECTS(box, circle) !(circle.center.x - circle.radius > box.max_x \
|| circle.center.x + circle.radius < box.min_x \
|| circle.center.y - circle.radius > box.max_y \
|| circle.center.y + circle.radius < box.min_y)
typedef struct segment_t {
vec2 p0, p1, p2, p3; // cubic Bézier: start, control1, control2, end
} segment_t;
@@ -22,167 +30,232 @@ typedef struct shape_meta_t {
box_t aabb;
} shape_meta_t;
#define TILE_SIZE 256 // pixels per tile (same for all LODs)
#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
typedef struct circle_t {
vec2 center;
float radius;
} circle_t;
typedef struct __attribute__((packed)) tile_key_t {
uint32_t lod, tx, ty;
} tile_key_t;
#define TILE_SIZE 256 // texels per tile
// Tile descriptor (separated from tile_cache_entry to be sent to the GPU)
typedef struct tile_t {
tile_key_t key;
float texel_size, tile_world;
} tile_t;
#define TILE_COUNT(size) (uint32_t)ceilf(size.x / TILE_SIZE) * (uint32_t)ceilf(size.y / TILE_SIZE)
typedef struct LOD_t {
tile_t* tiles;
#define INITIAL_SHAPE_SIZE 256
#define INITIAL_SEGMENTS_SIZE 8192
#define INITIAL_CIRCLE_SIZE 1024
#define INITIAL_INDICES_SIZE 16384
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 {
// Tile metadata for targetted updates
typedef struct tile_task_t {
box_t bounds;
uint32_t layer;
} tile_info_t;
uint32_t offset;
uint32_t circle_count;
uint32_t count;
} tile_task_t;
typedef struct scene_t {
vec2 world_size;
uint32_t num_shapes;
uint32_t max_shapes;
shape_meta_t* shapes;
struct {
uint32_t num_shapes;
uint32_t max_shapes;
shape_meta_t* shapes;
uint32_t num_segments;
uint32_t max_segments;
segment_t* segments;
//Theorically, the LOD amount and tile per LOD are computable, so it's not relevant to store them.
uint32_t num_LOD;
LOD_t* LODs;
sg_view texture_cache;
uint32_t num_segments;
uint32_t max_segments;
segment_t* segments;
} shapes;
struct {
tile_slot_t slots[TILE_LAYERS];
uint32_t lru_head, lru_tail;
uint32_t num_circles;
uint32_t max_circles;
circle_t* circles;
} circles;
struct { tile_key_t key; uint32_t value; }* map;
uint32_t free_layers[TILE_LAYERS];
uint32_t free_count;
struct {
tile_task_t* tiles; // Array of tiles metadata for the compute buffer
bool* tile_dirty; // Array of tiles state, since this info don't needs to be uploaded to the GPU, it is stored separately
uint32_t max_indices;
uint32_t num_indices;
uint32_t* indices; // Array of shapes/circles indices per tile, pre culled on the CPU side
} cache;
bool shape_dirty, segment_dirty;
sg_view texture; //Layer count depends on the world size so the texture has to be allocate on the scene side
bool shape_dirty, segment_dirty, circle_dirty;
} scene_t;
// 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*) emmalloc_malloc(init_shape_cap * sizeof(shape_meta_t));
static uint32_t scene_init(scene_t* s, vec2 world_size) {
s->shapes.num_shapes = 0;
s->shapes.max_shapes = INITIAL_SHAPE_SIZE;
s->shapes.shapes = (shape_meta_t*) emmalloc_malloc(INITIAL_SHAPE_SIZE * sizeof(shape_meta_t));
s->num_segments = 0;
s->max_segments = init_seg_cap;
s->segments = (segment_t*) emmalloc_malloc(init_seg_cap * sizeof(segment_t));
s->shapes.num_segments = 0;
s->shapes.max_segments = INITIAL_SEGMENTS_SIZE;
s->shapes.segments = (segment_t*) emmalloc_malloc(INITIAL_SEGMENTS_SIZE * sizeof(segment_t));
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->LODs = emmalloc_malloc(max_LOD * sizeof(LOD_t));
for(int i = 0; i < max_LOD; ++i)
s->circles.num_circles = 0;
s->circles.max_circles = INITIAL_SEGMENTS_SIZE;
s->circles.circles = (circle_t*) emmalloc_malloc(INITIAL_CIRCLE_SIZE * sizeof(circle_t));
const uint32_t tile_count_x = (uint32_t)ceilf(world_size.x / TILE_SIZE), tile_count_y = (uint32_t)ceilf(world_size.y / TILE_SIZE);
const uint32_t tile_count = tile_count_x * tile_count_y;
s->cache.tiles = (tile_task_t*) emmalloc_malloc(sizeof(tile_task_t) * tile_count);
s->cache.tile_dirty = (bool*) emmalloc_malloc(sizeof(bool) * tile_count);
uint32_t idx = 0;
for(uint32_t y = 0; y < tile_count_y; y++)
{
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 = emmalloc_malloc(sizeof(tile_t) * tiles_count);
s->LODs[i].range = (vec2) { }; //TODO
uint32_t x = 0; uint32_t y = 0;
for(int j = 0; j < tiles_count; j++)
for(uint32_t x = 0; x < tile_count_x; x++)
{
x++;
if(x >= s->LODs[i].dimension.x) { x = 0; y++; }
//s->LODs[i].tiles[j] = (tile_t) { .lod = i, .tile = { x, y }, . };
s->cache.tile_dirty[idx] = true;
idx++;
}
}
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->cache.max_indices = INITIAL_INDICES_SIZE;
s->cache.num_indices = 0;
s->cache.indices = (uint32_t*) emmalloc_malloc(sizeof(uint32_t) * INITIAL_INDICES_SIZE);
s->world_size = world_size;
s->texture = sg_make_view(&(sg_view_desc) {
.label = "SDF tiles view",
.texture = {
.image = sg_make_image(&(sg_image_desc) {
.width = (uint32_t) ceilf(world_size.x),
.height = (uint32_t) ceilf(world_size.y),
.label = "SDF tiles texture",
.pixel_format = SG_PIXELFORMAT_R16F,
.type = SG_IMAGETYPE_2D,
.usage = { .storage_image = true },
}),
},
});
if (!s->shapes || !s->segments || !s->LODs) return 0; // allocation failure
if (!s->shapes.shapes || !s->shapes.segments || !s->circles.circles || !s->cache.indices || !s->cache.tiles || !s->cache.tile_dirty) return 0; // allocation failure
return 1;
}
// 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) {
int new_cap = s->max_segments * 2;
segment_t* tmp = (segment_t*) realloc(s->segments, new_cap * sizeof(segment_t));
if (!tmp) return -1;
s->segments = tmp;
s->max_segments = new_cap;
#define COORD_TO_INDEX(x, y, scene) y * ceilf(s->world_size.x / TILE_SIZE) + x
static void scene_compute_culling(scene_t* s, tile_task_t* task)
{
task->count = 0;
for(uint32_t i = 0; i < s->circles.num_circles; i++)
{
if(CIRCLE_INTERSECTS(task->bounds, s->circles.circles[i]))
{
if(s->cache.num_indices >= s->cache.max_indices)
{
s->cache.max_indices *= 2;
s->cache.indices = emmalloc_realloc(s->cache.indices, s->cache.max_indices * sizeof(uint32_t));
}
s->cache.indices[s->cache.num_indices++] = i;
task->count++;
}
}
int idx = s->num_segments++;
s->segments[idx] = seg;
task->circle_count = task->count;
for(uint32_t i = 0; i < s->shapes.num_shapes; i++)
{
if(BOX_INTERSECTS(task->bounds, s->shapes.shapes[i].aabb))
{
if(s->cache.num_indices >= s->cache.max_indices)
{
s->cache.max_indices *= 2;
s->cache.indices = emmalloc_realloc(s->cache.indices, s->cache.max_indices * sizeof(uint32_t));
}
s->cache.indices[s->cache.num_indices++] = i;
task->count++;
}
}
}
// Find the proper LOD, the relevant tiles then run the culling process for the dirty tiles and upload data to the GPU for SDF rendering
static uint32_t scene_process_tiles(scene_t* s, float pan_x, float pan_y, float zoom, sg_bindings* bindings)
{
const float width = (float)sapp_width(), height = (float) sapp_height();
const float wpp = fmaxf(s->world_size.y / zoom / height, s->world_size.x / zoom / width); //World point per pixel
float view_w = width * wpp, view_h = height * wpp;
box_t frustum = {
.min_x = ceilf((pan_x - view_w * 0.5) / TILE_SIZE),
.min_y = ceilf((pan_y - view_h * 0.5) / TILE_SIZE),
.max_x = ceilf((pan_x + view_w * 0.5) / TILE_SIZE),
.max_y = ceilf((pan_y + view_h * 0.5) / TILE_SIZE)
};
uint32_t dirty_tile_count = 0, offset = 0;
for(uint32_t y = frustum.min_y; y < frustum.max_y; y++)
{
for(uint32_t x = frustum.min_x; x < frustum.max_x; x++)
{
uint32_t idx = COORD_TO_INDEX(x, y, s);
if(s->cache.tile_dirty[idx])
{
tile_task_t* task = &s->cache.tiles[dirty_tile_count++];
task->bounds.min_x = x * TILE_SIZE;
task->bounds.min_y = y * TILE_SIZE;
task->bounds.max_x = fminf((x + 1) * TILE_SIZE, s->world_size.x);
task->bounds.max_y = fminf((y + 1) * TILE_SIZE, s->world_size.y);
task->offset = offset;
scene_compute_culling(s, task);
offset += task->count;
s->cache.tile_dirty[idx] = false;
}
}
}
return dirty_tile_count;
}
// Append a segment, returns its index. (Reallocs if needed, simplified.)
static uint32_t scene_add_segment(scene_t* s, segment_t seg) {
if (s->shapes.num_segments >= s->shapes.max_segments) {
int new_cap = s->shapes.max_segments * 2;
segment_t* tmp = (segment_t*) realloc(s->shapes.segments, new_cap * sizeof(segment_t));
if (!tmp) return -1;
s->shapes.segments = tmp;
s->shapes.max_segments = new_cap;
}
int idx = s->shapes.num_segments++;
s->shapes.segments[idx] = seg;
return idx;
}
// Append a shape (meta data) and return its index.
static int scene_add_shape(scene_t* s, shape_meta_t meta) {
if (s->num_shapes >= s->max_shapes) {
int new_cap = s->max_shapes * 2;
shape_meta_t* tmp = (shape_meta_t*) realloc(s->shapes, new_cap * sizeof(shape_meta_t));
static uint32_t scene_add_shape(scene_t* s, shape_meta_t meta) {
if (s->shapes.num_shapes >= s->shapes.max_shapes) {
int new_cap = s->shapes.max_shapes * 2;
shape_meta_t* tmp = (shape_meta_t*) emmalloc_realloc(s->shapes.shapes, new_cap * sizeof(shape_meta_t));
if (!tmp) return -1;
s->shapes = tmp;
s->max_shapes = new_cap;
s->shapes.shapes = tmp;
s->shapes.max_shapes = new_cap;
}
int idx = s->num_shapes++;
s->shapes[idx] = meta;
int idx = s->shapes.num_shapes++;
s->shapes.shapes[idx] = meta;
return idx;
}
static int scene_shutdown(scene_t* s) {
emmalloc_free(s->segments);
emmalloc_free(s->shapes);
for(uint32_t i = 0; i < s->num_LOD; ++i)
emmalloc_free(s->LODs[i].tiles);
emmalloc_free(s->LODs);
static uint32_t scene_shutdown(scene_t* s) {
emmalloc_free(s->shapes.segments);
emmalloc_free(s->shapes.shapes);
emmalloc_free(s->circles.circles);
emmalloc_free(s->cache.indices);
emmalloc_free(s->cache.tiles);
emmalloc_free(s->cache.tile_dirty);
emmalloc_free(s);
return 1;
}
// Compute axis-aligned bounding box for a set of segments (used after adding a shape).
static void compute_bbox(segment_t* segs, int start, int count,
float* minx, float* miny, float* maxx, float* maxy) {
static void compute_bbox(segment_t* segs, uint32_t start, uint32_t count, float* minx, float* miny, float* maxx, float* maxy) {
float mx = 1e30f, my = 1e30f, Mx = -1e30f, My = -1e30f;
for (int i = 0; i < count; i++) {
segment_t s = segs[start + i];
@@ -200,7 +273,7 @@ static void compute_bbox(segment_t* segs, int start, int count,
// Add an axisaligned rectangle centred at `center` with given width and height.
// Returns the shape index, or -1 on error.
int add_rectangle(scene_t* s, vec2 center, vec2 size) {
uint32_t add_rectangle(scene_t* s, vec2 center, vec2 size) {
float hw = size.x * 0.5f, hh = size.y * 0.5f;
vec2 corners[4] = {
{center.x - hw, center.y - hh}, // bottomleft
@@ -228,14 +301,14 @@ 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.aabb.min_x, &meta.aabb.min_y, &meta.aabb.max_x, &meta.aabb.max_y);
meta.aabb = (box_t) { center.x - hw, center.y - hh, center.x + hw, center.y + hh };
return scene_add_shape(s, meta);
}
// Add a circle centred at `center` with given radius.
// approximated by 4 cubic Bézier segments (one per quadrant).
// Returns shape index or -1.
int add_circle(scene_t* s, vec2 center, float radius) {
uint32_t add_circle_as_shape(scene_t* s, vec2 center, float radius) {
const float k = 0.552284749831f; // magic constant for 90° arc (4/3 * (sqrt(2)-1))
float rk = radius * k;
@@ -279,6 +352,12 @@ 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.aabb.min_x, &meta.aabb.min_y, &meta.aabb.max_x, &meta.aabb.max_y);
compute_bbox(s->shapes.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);
}
uint32_t add_circle(scene_t* s, vec2 center, float radius) {
s->circles.circles[s->circles.num_circles++] = (circle_t) { center, radius };
return s->circles.num_circles - 1;
}