facett-core/src/render/wgsl.rs
//! **THE shader TEXT — one place, and reachable without a GPU feature.**
//!
//! Every `.wgsl` file facett ships is `include_str!`ed exactly once, here. The
//! pipeline modules that build `wgpu::ShaderModule`s out of them (`render::gpu::*`,
//! all behind feature `wgpu`) re-export these constants under their old paths, so
//! nothing outside this file changed name.
//!
//! # Why the text is NOT inside `render::gpu`
//!
//! It used to be, and that made the shader-CONTRACT guards unreachable. `render::gpu`
//! is `#[cfg(feature = "wgpu")]`, so a plain `cargo test -p facett-core` compiled the
//! whole tree out and printed a green `… filtered out` — including
//! [`tests::cull_shader_keeps_its_per_way_contract`], the guard on
//! `cull.wgsl`, which is the shader **every road on the map goes through**.
//! MEASURED 2026-08-22 at facett `e6084d7d`: `cargo test -p facett-core` listed 430
//! tests and `--features wgpu` listed 525, so 95 of this crate's tests (18.1 %) never
//! ran on a default build, and 56 of the 76 inside `render/gpu` need no device at all.
//!
//! A shader source string is text. It needs `include_str!` and nothing else — no
//! `wgpu`, no adapter, no `bytemuck`. Putting it behind the GPU feature bought
//! nothing and cost the contract guards their only chance to run. They live here now
//! and run on **every** build of this crate.
//!
//! This is the LAW-5 shape: one writer, not two copies watched by a third test.
/// **THE shared WGSL prelude** — `px_to_ndc`, `coverage_from_sd`, `coverage_inside`
/// and `resolve_line_width_px`, the geometry maths every facett line/SDF shader
/// needs. WGSL has no `#include`, so it is prepended as text by [`wgsl`]. See
/// `gpu/common.wgsl` for why each function is bit-identical to the copies it replaced.
pub const COMMON_WGSL: &str = include_str!("gpu/common.wgsl");
/// Compose a shader: [`COMMON_WGSL`] followed by `body`.
///
/// **The one composer.** Every `create_shader_module` in facett that wants the shared
/// maths goes through here — facett-core's SDF/line/draw pipelines and facett-map's
/// pretty-line pipeline alike — so the prelude cannot be attached two different ways.
#[must_use]
pub fn wgsl(body: &str) -> String {
let mut s = String::with_capacity(COMMON_WGSL.len() + body.len() + 1);
s.push_str(COMMON_WGSL);
s.push('\n');
s.push_str(body);
s
}
/// The GPU frustum-cull + LOD + stream-compaction compute shader — the country-scale
/// chunked pipeline. **The single copy**: facett-map's `GpuMapRenderer` builds its
/// compute pipeline from this constant (it used to hold a byte-identical twin at
/// `facett-map/src/gpu/cull.wgsl`, watched by a guard test; the twin is gone, so
/// there is nothing left to drift). The chunking partition in `render::gpu::buffer`
/// keeps this per-way shader unchanged.
pub const CULL_WGSL: &str = include_str!("gpu/cull.wgsl");
/// The line-draw vertex/fragment shader for OSM ways — **the single copy**, consumed
/// by facett-map's `GpuMapRenderer` (its byte-identical twin at
/// `facett-map/src/gpu/draw.wgsl` is gone). Needs the [`COMMON_WGSL`] prelude:
/// compose it with [`wgsl`].
pub const DRAW_WGSL: &str = include_str!("gpu/draw.wgsl");
/// The SDF quad shader BODY, byte-identical to the CPU coverage math. Needs the
/// [`COMMON_WGSL`] prelude — compose with [`wgsl`].
pub const SDF_WGSL: &str = include_str!("gpu/sdf.wgsl");
/// The thick-AA line shader BODY. Needs the [`COMMON_WGSL`] prelude — compose with
/// [`wgsl`].
pub const LINE_WGSL: &str = include_str!("gpu/line.wgsl");
/// The linked-list OIT shader (gather + resolve). Composed behind [`COMMON_WGSL`] by
/// [`wgsl`], from which it takes `px_to_ndc` — so its screen mapping is the same one
/// the line lanes use.
pub const OIT_WGSL: &str = include_str!("gpu/oit.wgsl");
/// The id-writing vertex/fragment shader. Needs the [`COMMON_WGSL`] prelude (it takes
/// `px_to_ndc` from there rather than re-spelling it) — compose with [`wgsl`].
pub const PICK_WGSL: &str = include_str!("gpu/pick.wgsl");
/// The reusable blur shader (separable Gaussian; `blur_vs`/`blur_fs`). The fragment
/// reads the per-axis uniform for the texel step + the CPU-computed weights.
pub const BLUR_WGSL: &str = include_str!("gpu/blur.wgsl");
/// The Scene3D shader — flat quads, screen-width line ribbons and SDF glyph quads on
/// world planes, one view-projection, depth-tested (`render::gpu::scene3d`).
pub const SCENE3D_WGSL: &str = include_str!("gpu/scene3d.wgsl");
/// The colour-grade shader (`grade_vs`/`grade_fs`).
pub const COLORGRADE_WGSL: &str = include_str!("gpu/colorgrade.wgsl");
/// The MSDF glyph shader (`msdf_vs`/`msdf_fs`).
pub const MSDF_WGSL: &str = include_str!("gpu/msdf.wgsl");
/// The TAA resolve shader. Not prelude-composed: it is a pure screen-space pass and
/// uses none of `common.wgsl`'s geometry maths.
pub const TAA_WGSL: &str = include_str!("gpu/taa.wgsl");
/// The compute shader (`cs_step`) + the additive point render (`pt_vs`/`pt_fs`).
pub const PARTICLES_WGSL: &str = include_str!("gpu/particles.wgsl");
/// The label collision compute passes (`cs_clear_frame` / `cs_clear_round` /
/// `cs_claim` / `cs_name_bid` / `cs_emit`).
pub const LABEL_COLLIDE_WGSL: &str = include_str!("gpu/label_collide.wgsl");
/// The instanced glyph draw whose instance count comes from the emit pass.
pub const LABEL_DRAW_WGSL: &str = include_str!("gpu/label_draw.wgsl");
/// The instanced emissive node-disc / edge-filament shader behind the shared graph
/// cloud.
pub const GRAPHCLOUD_WGSL: &str = include_str!("gpu/graphcloud.wgsl");
#[cfg(test)]
mod tests {
use super::*;
/// The cull shader's per-way contract (6 u32 words/vertex) is unchanged.
///
/// This replaces `shaders_are_byte_identical_to_facett_map`, which compared this
/// constant against `include_str!("../../../../facett-map/src/gpu/cull.wgsl")`.
/// That guard is deleted rather than kept: facett-map no longer HAS its own copy
/// (LAW #5 — fixed by construction, not by a watcher), and the cross-crate
/// `include_str!` it needed reached outside this crate's package directory, so it
/// only ever resolved inside the monorepo checkout.
///
/// The compaction is an INDEX LIST (GG audit row 21): `main` writes
/// `vert_start + v` into `compact`, never a vertex word, so it has no stride and
/// needs no `VPV` override — one pipeline compacts the 24 B `Vertex` streams and the
/// 48 B `LineVertex` streams alike. A stride returning here would mean the copy is
/// moving vertices again, and with it the second whole copy of the geometry in
/// VRAM that `facett-map/tests/vram_ratio.rs` holds at ≤ 1.2×.
#[test]
fn cull_shader_keeps_its_per_way_contract() {
// CODE lines only, here and below: the shader's own header explains at length
// what `VPV` and `atomicAdd` used to do and must keep being allowed to say so.
// A guard that cannot tell the prose from the program is a guard that fires on
// its own documentation — which is exactly what this one did on its first run.
let code: String =
CULL_WGSL.lines().filter(|l| !l.trim_start().starts_with("//")).collect::<Vec<_>>().join("\n");
assert!(code.contains("compact[dst + v] = m.vert_start + v;"), "the compaction must write INDICES");
assert!(!code.contains("VPV"), "a words-per-vertex stride means the cull is copying vertices again");
assert!(code.contains("const HEADER_WORDS: u32 = 5u;"), "the control block leads with DrawIndexedIndirectArgs");
assert!(code.contains("@workgroup_size(64)"));
// ── THE ORDERED COMPACTION (facett-map `tests/cull_order.rs` holds the
// behaviour; this holds the SHAPE the pipelines are built against).
//
// `atomicAdd` on the draw count is what made the GPU lane's paint order the
// workgroup scheduler's: MEASURED 24 000 of 24 000 vertices out of bake
// order and 21 888 of 24 000 moving between two dispatches of the same
// buffers (oden / RTX 4090, 2026-08-22). A revert to it would be a one-line
// edit here and a flickering map there, so the string is pinned.
for entry in ["fn count_pass(", "fn scan_pass(", "fn main("] {
assert!(CULL_WGSL.contains(entry), "the ordered compaction's `{entry}` entry point is gone");
}
assert!(
!code.contains("atomicAdd"),
"the cull must not claim its destination with `atomicAdd` — that is arrival order, \
not paint order, and it is not even stable between two runs of the same dispatch"
);
assert!(code.contains("atomicStore"), "the scan pass must still WRITE the count it computed");
// FOUR storage bindings, no more: a compute stage is guaranteed only that many,
// and wgpu refuses the layout outright at five (MEASURED: "limit is 4, count
// was 6"). The scan scratch therefore rides inside the indirect buffer.
assert!(!CULL_WGSL.contains("binding(5)"), "binding 5 would take the cull over the 4-storage-buffer floor");
}
/// The shared prelude really is prepended, and the bodies really do delegate: a
/// shader body that declared its own `px_to_ndc` would collide at compile time,
/// and one that never calls it would have silently kept a private copy.
#[test]
fn every_shared_shader_takes_its_maths_from_the_one_prelude() {
for (name, body) in
[("sdf", SDF_WGSL), ("line", LINE_WGSL), ("draw", DRAW_WGSL), ("oit", OIT_WGSL)]
{
assert!(
!body.contains("fn px_to_ndc"),
"{name}.wgsl re-declares px_to_ndc instead of using the prelude"
);
assert!(
!body.contains("fn coverage_from_sd"),
"{name}.wgsl re-declares coverage_from_sd instead of using the prelude"
);
assert!(body.contains("px_to_ndc("), "{name}.wgsl calls the shared px_to_ndc");
let composed = wgsl(body);
assert!(composed.starts_with(COMMON_WGSL), "{name}.wgsl is composed prelude-first");
assert!(composed.contains("fn px_to_ndc"), "the composed {name} module defines px_to_ndc once");
assert_eq!(
composed.matches("fn px_to_ndc").count(),
1,
"exactly ONE definition of px_to_ndc reaches the {name} shader module"
);
}
// The prelude carries the maths the "two scales of land and lines" bug lived
// in, so facett-map's line lane can share it rather than re-spell it.
assert!(COMMON_WGSL.contains("fn resolve_line_width_px"));
assert!(COMMON_WGSL.contains("fn coverage_inside"));
}
/// **Every `.wgsl` file in the tree is included exactly once, and from here.**
///
/// The reason the two guards above sat unreachable for so long is that the text
/// they watch lived in a feature-gated module, so nobody noticed they were being
/// filtered out. This one keeps that from happening again by CONSTRUCTION rather
/// than by discipline: it reads the `gpu/` directory off disk, and fails both ways
/// —
///
/// * a `.wgsl` file that no constant in this module includes is a shader no guard
/// here can ever see;
/// * an `include_str!("….wgsl")` anywhere ELSE under `src/` is a second copy of
/// the text, which is the exact twin this module exists to prevent.
///
/// RED (both proven 2026-08-22): drop `GRAPHCLOUD_WGSL` → "graphcloud.wgsl is not
/// included"; put `include_str!("gpu/cull.wgsl")` back into `render/gpu/mod.rs` →
/// "a second include_str".
#[test]
fn every_wgsl_file_is_included_here_and_nowhere_else() {
let src = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("src");
let me = src.join("render/wgsl.rs");
let my_text = std::fs::read_to_string(&me).expect("this file is readable");
let shaders: Vec<String> = std::fs::read_dir(src.join("render/gpu"))
.expect("render/gpu exists")
.filter_map(|e| {
let n = e.ok()?.file_name().to_string_lossy().into_owned();
n.ends_with(".wgsl").then_some(n)
})
.collect();
assert!(shaders.len() >= 14, "the shader set shrank unexpectedly: {}", shaders.len());
for s in &shaders {
let want = format!("include_str!(\"gpu/{s}\")");
assert!(
my_text.contains(&want),
"{s} is not included by render/wgsl.rs — no shader-text guard can reach it \
(expected the literal `{want}`)"
);
}
// …and nowhere else. Walk every .rs under src/ except this file.
let mut stack = vec![src.clone()];
let mut strays: Vec<String> = Vec::new();
while let Some(dir) = stack.pop() {
for e in std::fs::read_dir(&dir).expect("readable dir").flatten() {
let p = e.path();
if p.is_dir() {
stack.push(p);
} else if p.extension().is_some_and(|x| x == "rs") && p != me {
let t = std::fs::read_to_string(&p).unwrap_or_default();
for line in t.lines() {
// Only real code — a doc comment may quote the old shape.
let l = line.trim_start();
if l.starts_with("//") {
continue;
}
if l.contains("include_str!(") && l.contains(".wgsl") {
strays.push(format!(
"{}: {}",
p.strip_prefix(&src).unwrap_or(&p).display(),
l.trim()
));
}
}
}
}
}
assert!(
strays.is_empty(),
"a second include_str of shader text — the twin render::wgsl exists to prevent:\n {}",
strays.join("\n ")
);
}
}