GPU-compute building extrusion — design
The chosen design for FLIGHT-SIM roadmap §3b: turn building extrusion from a CPU geometry generator into a CPU data marshaller feeding a GPU compute pass. The CPU packs raw footprints into tight storage buffers; extrude.wgsl builds the mesh, one invocation per building. The landed CPU extrude_prism in facett-map3d is untouched and stays the reference oracle + the adapter-free fallback (L2 additive, no regression); the shared robust ear-clip (§3a) is reused, never forked (L5).
The pieces (src/extrude_gpu.rs + src/extrude.wgsl)
-
Footprint— one building to extrude: a design-plane(x, z)ring, a roofheight, a basemin_height(OSMbuilding:min_height,0.0= ground, matching the CPU path), abuilding_typestyling tag. -
BuildingMeta— the GPU-side per-building record, byte-for-byte the WGSLstruct BuildingMeta:vertex_start/vertex_count(slice into the flat contour buffer),cap_start/cap_count(slice into the flat cap-index buffer),out_vertex_base/out_index_base(prefix-summed output offsets),height/min_height,building_type.#[repr(C)], 40 bytes, 4-byte-clean →bytemuck::Pod(the derive is gated behindwgpu, so the marshaller compiles + runs GPU-less). -
BuildingMarshal::from_footprints— the whole CPU-side job: open + CCW-normalise each ring (facett_map3d::open_cap_ring, the exact contour the caps index), precompute its cap triangulation (facett_map3d::cap_triangulation, the shared §3a ear-clip — concave/H-shaped footprints handled on the CPU), and stamp the prefix-summedBuildingMeta. No geometry is generated here. Degenerate (< 3-vertex / collinear) footprints are dropped. -
extrude.wgsl·cs_extrude—@workgroup_size(64), one invocation per building: emit the top ring (y = height,+Y), the bottom ring (y = min_height,−Y), andnwall quads (each 4 own verts carrying one flat outward normal), plus the cap- wall indices —
6nvertices,4n−4triangles, at the marshaller's precomputed bases. The output vertex/normal/index buffers are exactly what an indirect draw (draw_indexed_indirect) consumes.
- wall indices —
-
ExtrudeGpu(behindwgpu) — builds the pipeline, dispatches over aBuildingMarshal, reads the geometry back as aMesh3.
The vertex contract (identical to extrude_prism)
cs_extrude is a twin of the CPU extrude_prism risen case — same order, same math: top ring, then bottom ring, then per-edge wall quads with the centroid-away outward normal. So the GPU output is not merely similar to the CPU mesh; it is vertex-for- vertex the same geometry, which is what makes an exact oracle-parity proof possible.
The hybrid cap tessellation (the key nuance)
Robust concave tessellation is not attempted from scratch in the compute shader. The CPU cap_triangulation (the §3a ear-clip) precomputes the cap triangle indices — it handles concave + H-shaped footprints correctly — and the GPU only extrudes them. This fixes concave-footprint breakage and keeps the parallel per-building GPU work simple, while the walls (the bulk of the geometry) are fully GPU-generated.
Testability without a GPU (session constraint)
The RED-when-broken correctness gate runs GPU-less: marshal.cpu_extrude() — the CPU statement of the exact op cs_extrude runs — is proven vertex-for-vertex equal to extrude_prism_mesh for every footprint (vertex count, triangle count, positions/normals bit-equal, indices equal after rebasing, per-building AABB = footprint × [min_h, h]). A wrong offset / centroid / height / cap winding diverges it, on any box. The adapter-needing ExtrudeGpu dispatch test self-skips with a recorded sentinel when no adapter is present (the scatter_gpu / cdlod pattern), and when it runs it proves the GPU output equals that same CPU reference (position tol 1e-4, exact indices).
Alternatives rejected
Full in-shader concave tessellation (fragile, serial per polygon), a geometry-shader prism extrusion (deprecated on modern targets, no compute reuse), and keeping the CPU generator as the only path (the very thing §3b offloads). See the roadmap §3b spike notes; the CPU generator remains as the fallback + oracle, not the hot path.