Your browser sucks, JavaScript disabled! Gunnar's console needs JavaScript and WebAssembly. Git over SSH on port 2222 does not.

nordisk/facett

facett-map3d-flight/src/extrude.wgsl

// GPU-compute building extrusion (FLIGHT-SIM #3b / roadmap §3b) — the compute twin of
// `facett-map3d`'s CPU `extrude_prism`. One invocation per building: it reads a
// `BuildingMeta` record + a slice of the flat contour buffer + the CPU-precomputed cap
// triangle indices, and writes the prism geometry (top cap + bottom cap + side walls)
// into the shared output vertex/normal/index buffers at the marshaller's precomputed
// bases. The Rust side reads the result back and proves it is the SAME geometry as the
// CPU `extrude_prism_mesh` oracle (the scatter_lut.wgsl / cdlod_morph.wgsl parity
// pattern). Layout mirrors `BuildingMeta` / the buffers in `extrude_gpu.rs`.
//
// Geometry, vertex-for-vertex identical to `extrude_prism` (risen case):
//   vb+0        .. vb+n     : top ring   (y = height,     normal +Y)
//   vb+n        .. vb+2n    : bottom ring (y = min_height, normal -Y)
//   vb+2n + 4*i .. +4       : wall quad i (4 own verts, one flat outward normal)
// with 4n-4 triangles: (n-2) top cap + (n-2) bottom cap (reversed) + 2n wall.

struct Params {
    count : u32,   // number of buildings
    _pad0 : u32,
    _pad1 : u32,
    _pad2 : u32,
};

struct BuildingMeta {
    vertex_start    : u32,  // offset into `contour` (open, CCW ring vertices)
    vertex_count    : u32,  // n — ring vertex count
    cap_start       : u32,  // starting triangle index into `caps` (3 u32 per tri)
    cap_count       : u32,  // number of cap triangles (n-2 for a clean ring)
    out_vertex_base : u32,  // base into out_pos / out_nrm  (vb)
    out_index_base  : u32,  // base into out_idx (u32 count, = first tri * 3)
    height          : f32,
    min_height      : f32,
    building_type   : u32,
    _pad            : u32,
};

@group(0) @binding(0) var<uniform>              params  : Params;
@group(0) @binding(1) var<storage, read>        contour : array<vec2<f32>>;
@group(0) @binding(2) var<storage, read>        caps    : array<u32>;
@group(0) @binding(3) var<storage, read>        buildings : array<BuildingMeta>;
@group(0) @binding(4) var<storage, read_write>  out_pos : array<vec4<f32>>;
@group(0) @binding(5) var<storage, read_write>  out_nrm : array<vec4<f32>>;
@group(0) @binding(6) var<storage, read_write>  out_idx : array<u32>;

@compute @workgroup_size(64)
fn cs_extrude(@builtin(global_invocation_id) gid : vec3<u32>) {
    let b = gid.x;
    if (b >= params.count) {
        return;
    }
    let m  = buildings[b];
    let n  = m.vertex_count;
    if (n < 3u) {
        return; // degenerate footprint — no geometry
    }
    let vs = m.vertex_start;
    let vb = m.out_vertex_base;
    let h  = m.height;
    let mh = m.min_height;

    // Footprint centroid (for the outward wall-normal disambiguation) — mean of the
    // ring vertices, exactly as `extrude_prism`.
    var c = vec2<f32>(0.0, 0.0);
    for (var i = 0u; i < n; i = i + 1u) {
        c = c + contour[vs + i];
    }
    c = c / f32(n);

    // ── top ring (y = height, +Y) and bottom ring (y = min_height, -Y) ──
    for (var i = 0u; i < n; i = i + 1u) {
        let p = contour[vs + i];
        out_pos[vb + i]     = vec4<f32>(p.x, h,  p.y, 1.0);
        out_nrm[vb + i]     = vec4<f32>(0.0, 1.0, 0.0, 0.0);
        out_pos[vb + n + i] = vec4<f32>(p.x, mh, p.y, 1.0);
        out_nrm[vb + n + i] = vec4<f32>(0.0, -1.0, 0.0, 0.0);
    }

    // ── side walls: one quad per rim edge, four OWN verts carrying one flat normal ──
    let wall_base = vb + 2u * n;
    for (var i = 0u; i < n; i = i + 1u) {
        let p0 = contour[vs + i];
        let p1 = contour[vs + ((i + 1u) % n)];
        let e  = p1 - p0;
        let elen = max(length(e), 1e-6);
        var nrm = vec2<f32>(e.y / elen, -e.x / elen); // perpendicular
        let mid = (p0 + p1) * 0.5 - c;                // edge midpoint, centroid-relative
        if (dot(nrm, mid) < 0.0) {
            nrm = -nrm; // flip so it points outward (centroid-away)
        }
        let wn = vec4<f32>(nrm.x, 0.0, nrm.y, 0.0);
        let q = wall_base + 4u * i;
        out_pos[q + 0u] = vec4<f32>(p0.x, h,  p0.y, 1.0); out_nrm[q + 0u] = wn; // i0t
        out_pos[q + 1u] = vec4<f32>(p0.x, mh, p0.y, 1.0); out_nrm[q + 1u] = wn; // i0b
        out_pos[q + 2u] = vec4<f32>(p1.x, h,  p1.y, 1.0); out_nrm[q + 2u] = wn; // i1t
        out_pos[q + 3u] = vec4<f32>(p1.x, mh, p1.y, 1.0); out_nrm[q + 3u] = wn; // i1b
    }

    // ── indices ──
    var w = m.out_index_base;
    // top cap — ear-clip tris, as-is (CCW from +Y)
    for (var t = 0u; t < m.cap_count; t = t + 1u) {
        let ci = (m.cap_start + t) * 3u;
        out_idx[w + 0u] = vb + caps[ci + 0u];
        out_idx[w + 1u] = vb + caps[ci + 1u];
        out_idx[w + 2u] = vb + caps[ci + 2u];
        w = w + 3u;
    }
    // bottom cap — same tris, reversed winding (points -Y)
    for (var t = 0u; t < m.cap_count; t = t + 1u) {
        let ci = (m.cap_start + t) * 3u;
        out_idx[w + 0u] = vb + n + caps[ci + 0u];
        out_idx[w + 1u] = vb + n + caps[ci + 2u];
        out_idx[w + 2u] = vb + n + caps[ci + 1u];
        w = w + 3u;
    }
    // walls — two tris per quad, outward-facing CCW (matches `extrude_prism`)
    for (var i = 0u; i < n; i = i + 1u) {
        let q = wall_base + 4u * i;
        out_idx[w + 0u] = q + 0u; out_idx[w + 1u] = q + 2u; out_idx[w + 2u] = q + 1u;
        out_idx[w + 3u] = q + 2u; out_idx[w + 4u] = q + 3u; out_idx[w + 5u] = q + 1u;
        w = w + 6u;
    }
}