From ebe853c1058adeb6f2c9f91867a6ae0aeacc5f5f Mon Sep 17 00:00:00 2001
From: Jarrod Doyle <thejarroddoyle@gmail.com>
Date: Wed, 5 Jul 2023 14:01:14 +0100
Subject: [PATCH] Add brickmap/grid unpack shader

---
 assets/shaders/brickmap_upload.wgsl | 75 +++++++++++++++++++++++++++++
 1 file changed, 75 insertions(+)
 create mode 100644 assets/shaders/brickmap_upload.wgsl

diff --git a/assets/shaders/brickmap_upload.wgsl b/assets/shaders/brickmap_upload.wgsl
new file mode 100644
index 0000000..6728afe
--- /dev/null
+++ b/assets/shaders/brickmap_upload.wgsl
@@ -0,0 +1,75 @@
+@group(0) @binding(0) var<uniform> world_state: WorldState;
+@group(0) @binding(1) var<storage, read_write> brickgrid: array<atomic<u32>>;
+@group(0) @binding(2) var<storage, read_write> brickmap_cache: array<Brickmap>;
+@group(0) @binding(3) var<storage, read_write> shading_table: array<ShadingElement>;
+@group(0) @binding(4) var<storage, read> brickmap_unpack: BrickmapUnpack;
+@group(0) @binding(5) var<storage, read> brickgrid_unpack: BrickgridUnpack;
+
+struct ShadingElement {
+    albedo: u32,
+}
+
+struct Brickmap {
+    bitmask: array<u32, 16>,
+    shading_table_offset: u32,
+    lod_color: u32,
+}
+
+struct WorldState {
+    brickgrid_dims: vec3<u32>,
+    _pad: u32,
+};
+
+struct BrickmapUnpack {
+    max_count: u32,
+    count: u32,
+    _pad1: u32,
+    _pad2: u32,
+    elements: array<BrickmapUnpackElement>,
+}
+
+struct BrickmapUnpackElement {
+    cache_idx: u32,
+    brickmap: Brickmap,
+    shading_element_count: u32,
+    shading_elements: array<ShadingElement, 512>, // Always have space for a full map.
+}
+
+struct BrickgridUnpack {
+    max_count: u32,
+    count: u32,
+    _pad1: u32,
+    _pad2: u32,
+    elements: array<BrickgridUnpackElement>,
+}
+
+struct BrickgridUnpackElement {
+    grid_idx: u32,
+    grid_val: u32,
+}
+
+// Utility function. Converts a position in 3d to a 1d index.
+fn to_1d_index(p: vec3<i32>, dims: vec3<i32>) -> u32 {
+    return u32(p.x + p.y * dims.x + p.z * dims.x * dims.y);
+}
+
+@compute @workgroup_size(8,1,1)
+fn compute(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let unpack_idx = global_id.x;
+
+    // Brickgrid unpacking
+    if (unpack_idx < brickgrid_unpack.count){
+        let element = brickgrid_unpack.elements[unpack_idx];
+        brickgrid[element.grid_idx] = element.grid_val;
+    }
+
+    // Brickmap unpacking
+    if (unpack_idx < brickmap_unpack.count) {
+        let element = &brickmap_unpack.elements[unpack_idx];
+        brickmap_cache[(*element).cache_idx] = (*element).brickmap;
+        let st_offset = (*element).brickmap.shading_table_offset;
+        for (var i: u32 = 0u; i < (*element).shading_element_count; i++) {
+            shading_table[st_offset + i] = (*element).shading_elements[i];
+        }
+    }
+}
\ No newline at end of file