@group(0) @binding(0) var world_state: WorldState; @group(0) @binding(1) var brickgrid: array>; @group(0) @binding(2) var brickmap_cache: array; @group(0) @binding(3) var shading_table: array; @group(0) @binding(4) var brickmap_unpack: BrickmapUnpack; @group(0) @binding(5) var brickgrid_unpack: BrickgridUnpack; struct ShadingElement { albedo: u32, } struct Brickmap { bitmask: array, shading_table_offset: u32, lod_color: u32, } struct WorldState { brickgrid_dims: vec3, _pad: u32, }; struct BrickmapUnpack { max_count: u32, count: u32, _pad1: u32, _pad2: u32, elements: array, } struct BrickmapUnpackElement { cache_idx: u32, brickmap: Brickmap, shading_element_count: u32, shading_elements: array, // Always have space for a full map. } struct BrickgridUnpack { max_count: u32, count: u32, _pad1: u32, _pad2: u32, elements: array, } 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, dims: vec3) -> 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) { 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]; } } }