Rebuild core around render graph AST and shared memory

Amp-Thread-ID: https://ampcode.com/threads/T-01a01380-b478-77d0-84a0-102880a5c5ae
Co-authored-by: Heaust Azure <heaust.azure@gmail.com>

renderer/src/app_setup.rs

@@ -1,193 +1,68 @@
 #![cfg(target_arch = "wasm32")]
 
-use std::sync::mpsc::{self, Sender};
-use wasm_bindgen::closure::Closure;
+use std::cell::RefCell;
+use std::sync::mpsc;
+
 use wasm_bindgen::prelude::*;
-use wasm_bindgen::JsCast;
+use wasm_bindgen_futures::spawn_local;
 
 use crate::command_ring::CommandRing;
-use crate::message::WindowEvent;
-use crate::platform::web;
-use crate::platform::web::worker::MainWorker;
-use wasm_bindgen_futures::spawn_local;
-use web_sys::AddEventListenerOptions;
+use crate::message::{MouseMessage, ResizeMessage, WheelMessage, WindowEvent};
+use crate::platform::web::worker;
 
-pub struct EventListeners {
-    _resize_listener: Closure<dyn FnMut()>,
-    _pointer_listener: Closure<dyn FnMut(web_sys::PointerEvent)>,
-    _click_listener: Closure<dyn FnMut(web_sys::MouseEvent)>,
-    _wheel_listener: Closure<dyn FnMut(web_sys::WheelEvent)>,
-    _contextmenu_listener: Closure<dyn FnMut(web_sys::MouseEvent)>,
+thread_local! {
+    static WORKER_EVENTS: RefCell<Option<mpsc::Sender<WindowEvent>>> = const { RefCell::new(None) };
 }
 
-/// Setup default window event listeners that forward events to the worker thread
-pub fn setup_event_listeners(
-    worker_chan: &Sender<WindowEvent>,
-    canvas: &web_sys::HtmlCanvasElement,
-) -> Result<EventListeners, JsValue> {
-    let window = web_sys::window().unwrap();
-    let resize_worker_chan = worker_chan.clone();
-    let resize_canvas = canvas.clone();
-
-    let resize_listener: Closure<dyn FnMut()> = Closure::new(move || {
-        use crate::message::ResizeMessage;
-
-        let window = web_sys::window().unwrap();
-        let width = f64::from(resize_canvas.client_width().max(1));
-        let height = f64::from(resize_canvas.client_height().max(1));
-
-        let _ = resize_worker_chan.send(WindowEvent::Resize(ResizeMessage {
-            width,
-            height,
-            scale_factor: window.device_pixel_ratio(),
-        }));
-    });
-
-    window.add_event_listener_with_callback("resize", resize_listener.as_ref().unchecked_ref())?;
-
-    let pointer_worker_chan = worker_chan.clone();
-    let pointer_canvas = canvas.clone();
-    let pointer_listener: Closure<dyn FnMut(web_sys::PointerEvent)> =
-        Closure::new(move |event: web_sys::PointerEvent| {
-            use crate::message::{camera_drag, MouseMessage};
-
-            if event.pointer_type() != "mouse" {
-                return;
+/// Deliver a low-frequency browser event to the worker-owned renderer channel.
+#[wasm_bindgen]
+pub fn worker_window_event(kind: u32, values: js_sys::Float64Array) {
+    let values = values.to_vec();
+    let value = |index: usize| values.get(index).copied().unwrap_or_default();
+    let event = match kind {
+        0 => WindowEvent::Resize(ResizeMessage {
+            width: value(0),
+            height: value(1),
+            scale_factor: value(2),
+        }),
+        1 | 2 => {
+            let message = MouseMessage {
+                scale_factor: value(0),
+                buttons: value(1) as u16,
+                movement_x: value(2),
+                movement_y: value(3),
+                offset_x: value(4),
+                offset_y: value(5),
+                viewport_height: value(6),
+            };
+            if kind == 1 {
+                WindowEvent::PointerMove(message)
+            } else {
+                WindowEvent::PointerClick(message)
             }
-            match event.type_().as_str() {
-                "pointerdown" if matches!(event.button(), 1 | 2) => {
-                    event.prevent_default();
-                    let _ = pointer_canvas.set_pointer_capture(event.pointer_id());
-                }
-                "pointermove"
-                    if pointer_canvas.has_pointer_capture(event.pointer_id())
-                        && camera_drag(event.buttons()).is_some() =>
-                {
-                    event.prevent_default();
-                    let message = MouseMessage::from_pointer_evt(
-                        &event,
-                        f64::from(pointer_canvas.client_height().max(1)),
-                    );
-                    let _ = pointer_worker_chan.send(WindowEvent::PointerMove(message));
-                }
-                "pointerup" | "pointercancel"
-                    if pointer_canvas.has_pointer_capture(event.pointer_id()) =>
-                {
-                    let _ = pointer_canvas.release_pointer_capture(event.pointer_id());
-                }
-                _ => {}
-            }
-        });
-
-    for event_name in ["pointerdown", "pointermove", "pointerup", "pointercancel"] {
-        canvas.add_event_listener_with_callback(
-            event_name,
-            pointer_listener.as_ref().unchecked_ref(),
-        )?;
-    }
-
-    let click_worker_chan = worker_chan.clone();
-    let click_canvas = canvas.clone();
-    let click_listener: Closure<dyn FnMut(web_sys::MouseEvent)> =
-        Closure::new(move |event: web_sys::MouseEvent| {
-            use crate::message::MouseMessage;
-            if event.button() != 0 {
-                return;
-            }
-            let message =
-                MouseMessage::from_evt(&event, f64::from(click_canvas.client_height().max(1)));
-            let _ = click_worker_chan.send(WindowEvent::PointerClick(message));
-        });
-    canvas.add_event_listener_with_callback("click", click_listener.as_ref().unchecked_ref())?;
-
-    let wheel_worker_chan = worker_chan.clone();
-    let wheel_canvas = canvas.clone();
-    let wheel_listener: Closure<dyn FnMut(web_sys::WheelEvent)> =
-        Closure::new(move |event: web_sys::WheelEvent| {
-            use crate::message::WheelMessage;
-
-            event.prevent_default();
-            if let Some(message) =
-                WheelMessage::from_evt(&event, f64::from(wheel_canvas.client_height().max(1)))
-            {
-                let _ = wheel_worker_chan.send(WindowEvent::PointerWheel(message));
-            }
-        });
-
-    let wheel_options = {
-        let options = AddEventListenerOptions::new();
-        options.set_passive(false);
-        options
+        }
+        3 => WindowEvent::PointerWheel(WheelMessage {
+            delta_y_pixels: value(0) as f32,
+        }),
+        _ => return,
     };
-
-    canvas.add_event_listener_with_callback_and_add_event_listener_options(
-        "wheel",
-        wheel_listener.as_ref().unchecked_ref(),
-        &wheel_options,
-    )?;
-
-    let contextmenu_listener: Closure<dyn FnMut(web_sys::MouseEvent)> =
-        Closure::new(move |event: web_sys::MouseEvent| event.prevent_default());
-    canvas.add_event_listener_with_callback(
-        "contextmenu",
-        contextmenu_listener.as_ref().unchecked_ref(),
-    )?;
-
-    Ok(EventListeners {
-        _resize_listener: resize_listener,
-        _pointer_listener: pointer_listener,
-        _click_listener: click_listener,
-        _wheel_listener: wheel_listener,
-        _contextmenu_listener: contextmenu_listener,
-    })
+    WORKER_EVENTS.with(|sender| {
+        if let Some(sender) = sender.borrow().as_ref() {
+            let _ = sender.send(event);
+        }
+    });
 }
 
-/// Runtime resources required to keep a WASM application running.
-pub struct WebAppRuntime {
-    worker: MainWorker,
-    _event_listeners: EventListeners,
-    ring: Box<CommandRing>,
-}
-
-impl WebAppRuntime {
-    /// Initialize the web worker, canvas ownership, and event listeners.
-    pub fn new<T: crate::renderer::scene::Scene + 'static>(
-        worker_name: &str,
-        canvas_selector: &str,
-        profile: bool,
-    ) -> Result<Self, JsValue> {
-        let (sender, receiver) = mpsc::channel::<WindowEvent>();
-
-        let canvas = web::get_canvas_element(canvas_selector);
-        let window = web_sys::window().unwrap();
-        let dpr = window.device_pixel_ratio();
-        canvas.set_width((canvas.client_width() as f64 * dpr).round() as u32);
-        canvas.set_height((canvas.client_height() as f64 * dpr).round() as u32);
-        let ring = CommandRing::new();
-        let ring_ptr = ring.ptr();
-        let worker = MainWorker::spawn(worker_name, 1, ring_ptr, move || {
-            spawn_local(async move {
-                let ring = unsafe { &*(ring_ptr as *const CommandRing) };
-                MainWorker::run_render_loop::<T>(receiver, ring, profile).await;
-            });
-        })?;
-
-        worker.transfer_ownership(&canvas);
-
-        let event_listeners = setup_event_listeners(&sender, &canvas)?;
-
-        Ok(Self {
-            worker,
-            _event_listeners: event_listeners,
-            ring,
-        })
-    }
-
-    /// Access the spawned worker reference.
-    pub fn worker(&self) -> &MainWorker {
-        &self.worker
-    }
-    pub fn ring_ptr(&self) -> u32 {
-        self.ring.ptr()
-    }
+/// Start the typed renderer and return its SAB command-ring pointer.
+pub fn worker_entrypoint<T: crate::renderer::scene::Scene + 'static>(profile: bool) -> u32 {
+    let (sender, events) = mpsc::channel();
+    // The render worker owns this allocation for its entire lifetime. Publishing a
+    // stable address lets every connected thread use the same shared command ring.
+    let ring: &'static CommandRing = Box::leak(CommandRing::new());
+    let ring_ptr = ring.ptr();
+    WORKER_EVENTS.with(|worker_events| *worker_events.borrow_mut() = Some(sender));
+    spawn_local(async move {
+        worker::run_render_loop::<T>(events, ring, profile).await;
+    });
+    ring_ptr
 }

renderer/src/bin/app.rs

@@ -1,3 +0,0 @@
-fn main() {
-    println!("hello world!");
-}

renderer/src/example.wgsl

@@ -1,35 +0,0 @@
-struct UniformData {
-    mouse_move: vec2<f32>,
-    mouse_click: vec2<f32>,
-    resolution: vec2<f32>,
-    time: f32,
-}
-
-@group(0) @binding(0) var<uniform> uni: UniformData;
-@group(1) @binding(0) var<uniform> view_proj: mat4x4<f32>;
-
-struct VertexInput {
-    @location(0) pos: vec3<f32>,
-    @location(1) color: vec3<f32>
-}
-
-struct VertexOutput {
-    @builtin(position) pos: vec4<f32>,
-    @location(1) color: vec3<f32>
-}
-
-@vertex
-fn v_main(in: VertexInput) -> VertexOutput {
-    var out: VertexOutput;
-    out.pos = vec4<f32>(in.pos, 1.0);
-    let fluc = sin(modf(uni.time).fract * 3.141592) * 0.3 + 0.7;
-    out.color = in.color * fluc;
-    return out;
-}
-
-@fragment
-fn f_main(in: VertexOutput) -> @location(0) vec4<f32> {
-    let x = select(0.0, 0.3, distance(in.pos.xy, uni.mouse_move) < 25.0);
-    let y = select(0.0, 0.3, distance(in.pos.xy, uni.mouse_click) < 25.0);
-    return vec4f(in.color + x - y, 1.0);
-}

renderer/src/gltf.rs

@@ -5,7 +5,7 @@ use ultraviolet::{Mat4, Vec3};
 
 use crate::render_data::{
     InstanceHandle, InstanceType, MaterialKey, MeshCreateInfo, MeshHandle, ModelTransform,
-    PipelineKey, RenderData, RenderDataError,
+    RenderData, RenderDataError,
 };
 
 #[derive(Clone, Copy, Debug, Default, PartialEq, Eq)]
@@ -527,7 +527,6 @@ fn decode_gltf_model(mut model: Gltf) -> Result<ImportedScene, ImportError> {
 pub fn install_imported(
     target: &mut RenderData,
     imported: &ImportedScene,
-    pipelines: [PipelineKey; 2],
 ) -> Result<InstalledScene, ImportError> {
     let mut stage = target.replacement_stage()?;
     let mut handles = HashMap::new();
@@ -547,7 +546,6 @@ pub fn install_imported(
             tangents: &geometry.tangents,
             uvs: &geometry.uvs,
             indices: &geometry.indices,
-            pipeline: pipelines[usize::from(geometry.double_sided)],
             material: geometry.material,
             default_instance_type: InstanceType {
                 words: [

renderer/src/gltf.wgsl

@@ -1,52 +0,0 @@
-struct UniformData { mouse_move: vec2<f32>, mouse_click: vec2<f32>, resolution: vec2<f32>, time: f32, _padding0: f32, camera_position: vec4<f32> }
-struct MaterialData { base_color_factor: vec4<f32>, emissive_factor: vec4<f32>, surface_factors: vec4<f32>, alpha_optics: vec4<f32>, flags: vec4<u32>, uv_sets: vec4<u32>, debug_extras: vec4<u32> }
-@group(0) @binding(0) var<uniform> uni: UniformData;
-@group(1) @binding(0) var<uniform> view_proj: mat4x4<f32>;
-@group(2) @binding(0) var<uniform> material: MaterialData;
-@group(2) @binding(1) var base_tex: texture_2d<f32>;
-@group(2) @binding(2) var mr_tex: texture_2d<f32>;
-@group(2) @binding(3) var normal_tex: texture_2d<f32>;
-@group(2) @binding(4) var occlusion_tex: texture_2d<f32>;
-@group(2) @binding(5) var emissive_tex: texture_2d<f32>;
-@group(2) @binding(6) var base_sampler: sampler;
-@group(2) @binding(7) var mr_sampler: sampler;
-@group(2) @binding(8) var normal_sampler: sampler;
-@group(2) @binding(9) var occlusion_sampler: sampler;
-@group(2) @binding(10) var emissive_sampler: sampler;
-
-struct VertexInput { @location(0) pos: vec3<f32>, @location(1) normal: vec3<f32>, @location(2) uv: vec2<f32>, @location(3) model_col0: vec4<f32>, @location(4) model_col1: vec4<f32>, @location(5) model_col2: vec4<f32>, @location(6) model_col3: vec4<f32>, @location(7) normal_col0: vec4<f32>, @location(8) normal_col1: vec4<f32>, @location(9) normal_col2: vec4<f32>, @location(10) tangent: vec4<f32> }
-struct VertexOutput { @builtin(position) clip_position: vec4<f32>, @location(0) world_pos: vec3<f32>, @location(1) normal: vec3<f32>, @location(2) tangent: vec3<f32>, @location(3) bitangent: vec3<f32>, @location(4) uv: vec2<f32>, @location(5) @interpolate(flat) determinant_sign: f32 }
-fn safe_normalize(v: vec3<f32>, fallback: vec3<f32>) -> vec3<f32> { let l2 = dot(v, v); return select(fallback, v * inverseSqrt(l2), l2 > 1e-12 && l2 < 1e30); }
-@vertex fn vs_main(in: VertexInput) -> VertexOutput {
-    var out: VertexOutput; let model = mat4x4<f32>(in.model_col0, in.model_col1, in.model_col2, in.model_col3);
-    let linear = mat3x3<f32>(in.model_col0.xyz, in.model_col1.xyz, in.model_col2.xyz); let nm = mat3x3<f32>(in.normal_col0.xyz, in.normal_col1.xyz, in.normal_col2.xyz);
-    let world = model * vec4<f32>(in.pos, 1.0); let n = safe_normalize(nm * in.normal, vec3<f32>(0,1,0)); let raw_t = linear * in.tangent.xyz;
-    var t = raw_t - n * dot(n, raw_t); if dot(t,t) < 1e-8 { t = cross(select(vec3<f32>(0,1,0), vec3<f32>(1,0,0), abs(n.x) < 0.9), n); } t = safe_normalize(t, vec3<f32>(1,0,0));
-    out.clip_position = view_proj * world; out.world_pos = world.xyz; out.normal = n; out.tangent = t; out.bitangent = safe_normalize(cross(n,t), vec3<f32>(0,0,1)) * in.tangent.w * in.normal_col0.w; out.uv = in.uv; out.determinant_sign = in.normal_col0.w; return out;
-}
-struct Closure { base: vec4<f32>, mr: vec2<f32>, normal_map: vec3<f32>, ao: f32, emissive: vec3<f32> }
-fn sample_closure(uv: vec2<f32>) -> Closure {
-    let bits = material.flags.x; var c: Closure;
-    c.base = material.base_color_factor * select(vec4<f32>(1), textureSample(base_tex, base_sampler, uv), (bits & 1u) != 0u);
-    let mr = select(vec4<f32>(1), textureSample(mr_tex, mr_sampler, uv), (bits & 2u) != 0u); c.mr = vec2<f32>(clamp(material.surface_factors.x * mr.b,0,1), clamp(material.surface_factors.y * mr.g,0.045,1));
-    c.normal_map = select(vec3<f32>(0.5,0.5,1), textureSample(normal_tex, normal_sampler, uv).xyz, (bits & 4u) != 0u);
-    let occ = select(1.0, textureSample(occlusion_tex, occlusion_sampler, uv).r, (bits & 8u) != 0u); c.ao = mix(1.0, occ, material.surface_factors.w);
-    c.emissive = material.emissive_factor.rgb * select(vec3<f32>(1), textureSample(emissive_tex, emissive_sampler, uv).rgb, (bits & 16u) != 0u); return c;
-}
-fn schlick(f0: vec3<f32>, v_h: f32) -> vec3<f32> { return f0 + (vec3<f32>(1)-f0) * pow(1.0-clamp(v_h,0,1),5.0); }
-fn ggx_d(n_h_input: f32, a: f32) -> f32 { let n_h=clamp(n_h_input,0.0,1.0); let a2=a*a; let nh2=n_h*n_h; let q=(1.0-nh2)+a2*nh2; return a2/(3.14159265*q*q); }
-fn smith_v(n_v: f32, n_l: f32, a: f32) -> f32 { let a2=a*a; let gv=n_l*sqrt(max(n_v*n_v*(1.0-a2)+a2,0)); let gl=n_v*sqrt(max(n_l*n_l*(1.0-a2)+a2,0)); return 0.5/max(gv+gl,1e-6); }
-@fragment fn fs_main(in: VertexOutput, @builtin(front_facing) front: bool) -> @location(0) vec4<f32> {
-    let c=sample_closure(in.uv); if material.alpha_optics.x == 1.0 && c.base.a < material.alpha_optics.y { discard; }
-    let physical_front=front == (in.determinant_sign > 0); let orientation=select(-1.0,1.0,physical_front || material.flags.y == 0u); let map=c.normal_map*2.0-1.0;
-    let n=safe_normalize(mat3x3<f32>(in.tangent,in.bitangent,in.normal)*safe_normalize(vec3<f32>(map.xy*material.surface_factors.z,map.z),vec3<f32>(0,0,1)),in.normal)*orientation;
-    let v=safe_normalize(uni.camera_position.xyz-in.world_pos,n); let l=safe_normalize(vec3<f32>(0.35,1,0.45),vec3<f32>(0,1,0)); let h=safe_normalize(v+l,n);
-    let nv=max(dot(n,v),0); let nl=max(dot(n,l),0); let nh=dot(n,h); let vh=max(dot(v,h),0); let a=c.mr.y*c.mr.y;
-    let f0=mix(vec3<f32>(material.alpha_optics.w),c.base.rgb,c.mr.x); let direct_f=schlick(f0,vh); let env_f=schlick(f0,nv); let spec=direct_f*ggx_d(nh,a)*smith_v(nv,nl,a); let diffuse=(vec3<f32>(1)-direct_f)*(1.0-c.mr.x)*c.base.rgb/3.14159265;
-    let sun=(diffuse+spec)*nl*vec3<f32>(3.0,2.85,2.65);
-    let up=clamp(n.y*0.5+0.5,0,1); let sky=mix(vec3<f32>(0.055,0.045,0.035),vec3<f32>(0.24,0.36,0.58),up); let env_diff=(vec3<f32>(1)-env_f)*(1.0-c.mr.x)*c.base.rgb*sky;
-    let reflection=reflect(-v,n); let horizon=clamp(reflection.y*0.5+0.5,0,1); let env_spec=env_f*mix(vec3<f32>(0.04,0.035,0.03),vec3<f32>(0.28,0.42,0.7),horizon)*(1.0-0.65*c.mr.y);
-    var color=sun+(env_diff+env_spec)*c.ao+c.emissive;
-    if material.debug_extras.y == 1u { color=n*0.5+0.5; } else if material.debug_extras.y == 2u { color=vec3<f32>(c.mr.x,c.mr.y,c.ao); } else if material.debug_extras.y == 3u { color=f0; }
-    return vec4<f32>(color,1.0); // BLEND remains intentionally opaque.
-}

renderer/src/lib.rs

@@ -8,6 +8,7 @@ pub mod render_data;
 pub mod render_graph;
 pub mod renderer;
 pub mod shared_snapshot;
+pub mod shared_soa;
 
 #[cfg(target_arch = "wasm32")]
 thread_local! { static PAYLOADS: std::cell::RefCell<std::collections::HashMap<u32, Vec<u8>>> = Default::default(); }
@@ -43,27 +44,3 @@ pub(crate) fn take_payload(id: u32) -> Option<Vec<u8>> {
 pub(crate) fn take_payload(_id: u32) -> Option<Vec<u8>> {
     None
 }
-
-/// Worker entrypoint helper - executes the closure it is spawned with
-/// Applications should export this with #[wasm_bindgen]
-pub fn worker_entrypoint_impl(ptr: u32) {
-    let work = unsafe { Box::from_raw(ptr as *mut Box<dyn FnOnce()>) };
-    (*work)();
-}
-
-/// Macro to export the worker_entrypoint function in application crates
-///
-/// Usage:
-/// ```rust
-/// use renderer::export_worker_entrypoint;
-/// export_worker_entrypoint!();
-/// ```
-#[macro_export]
-macro_rules! export_worker_entrypoint {
-    () => {
-        #[wasm_bindgen::prelude::wasm_bindgen]
-        pub fn worker_entrypoint(ptr: u32) {
-            $crate::worker_entrypoint_impl(ptr);
-        }
-    };
-}

renderer/src/platform/mod.rs

@@ -1,5 +1,2 @@
 #[cfg(target_arch = "wasm32")]
 pub mod web;
-
-#[cfg(not(target_arch = "wasm32"))]
-mod native;

renderer/src/platform/native/mod.rs

@@ -1 +0,0 @@
-

renderer/src/platform/web/mod.rs

@@ -1,10 +1 @@
-use wasm_bindgen::JsCast;
-
 pub mod worker;
-
-pub fn get_canvas_element(selectors: &str) -> web_sys::HtmlCanvasElement {
-    let window = web_sys::window().unwrap();
-    let document = window.document().unwrap();
-    let element = document.query_selector(selectors).unwrap().unwrap();
-    element.dyn_into::<web_sys::HtmlCanvasElement>().unwrap()
-}

renderer/src/platform/web/worker/mainWorker.js

@@ -1,8 +1,14 @@
-// Generic worker that imports the app's WASM module relative to the generated pkg folder.
-// Works for any application because the relative depth from this file to pkg is stable.
-import initWasm, { clear_payloads, discard_payload, stage_payload, worker_entrypoint } from "/level-editor/pkg/level_editor.js";
+// The level editor's render worker owns its only WASM and WebGPU runtime.
+import initWasm, {
+  clear_payloads,
+  discard_payload,
+  stage_payload,
+  worker_main,
+  worker_memory,
+  worker_window_event,
+} from "/level-editor/pkg/level_editor.js";
 
-export function listenerReady() {
+function listenerReady() {
   if (state !== "waiting-listener") return;
   state = "replaying";
   for (const queued of pending.splice(0)) route(queued);
@@ -24,22 +30,24 @@ addEventListener("message", async (event) => {
   }
   if (state !== "uninitialized") return;
   state = "initializing";
-  const { wasmModule, workerId, memory, entryPtr } = message;
+  const { canvas, profile } = message;
 
-  console.log(
-    "worker: initializing with WASM module",
-    wasmModule,
-    "id:",
-    workerId,
-  );
-
-  // Initialize WASM with the shared module and memory forwarded from the main thread.
+  // The renderer worker exclusively owns the one WASM instance. Other threads
+  // receive only its shared memory and mutate the published SAB layouts.
   try {
-    api = await initWasm({ module_or_path: wasmModule, memory });
-    state = "waiting-listener";
-    worker_entrypoint(entryPtr);
+    api = await initWasm();
   } catch (error) {
-    fatal("WORKER_INIT_FAILED", String(error));
+    fatal("WORKER_INIT_FAILED", error?.stack || String(error));
+    return;
+  }
+  state = "waiting-listener";
+  pending.push({ type: "canvas", canvas });
+  try {
+    const ringPtr = worker_main(profile);
+    postMessage({ type: "bootstrap", memory: worker_memory(), ringPtr });
+    setTimeout(listenerReady, 0);
+  } catch (error) {
+    fatal("WORKER_ENTRY_FAILED", error?.stack || String(error));
   }
 });
 
@@ -51,6 +59,8 @@ function route(message) {
     postMessage({ type: "payload-ready", id: message.id });
   } else if (message?.type === "payload-release") {
     discard_payload(message.id);
+  } else if (message?.type === "window-event") {
+    worker_window_event(message.kind, message.values);
   }
 }
 

renderer/src/platform/web/worker/mod.rs

@@ -2,130 +2,25 @@ use crate::command_ring::CommandRing;
 use crate::message::WindowEvent;
 use log::info;
 use std::sync::mpsc::Receiver;
-use std::{cell::RefCell, fmt::Debug, ops::Deref, rc::Rc};
+use std::{cell::RefCell, rc::Rc};
 use wasm_bindgen::{prelude::*, JsValue};
 use wasm_bindgen_futures::JsFuture;
 use web_sys::MessageEvent;
 
-/// Binds JS.
-#[wasm_bindgen(module = "/src/platform/web/worker/workerGen.js")]
-extern "C" {
-    /// Spawn new worker in JS side in order to make bundler know about dependency.
-    #[wasm_bindgen(js_name = "createWorker")]
-    fn create_worker(kind: &str, name: &str) -> web_sys::Worker;
-}
+pub async fn run_render_loop<T: crate::renderer::scene::Scene + 'static>(
+    events_chan: Receiver<WindowEvent>,
+    ring: &'static CommandRing,
+    profile: bool,
+) {
+    use crate::renderer::Renderer;
 
-/// Binds JS.
-/// This makes wasm-bindgen bring `mainWorker.js` to the `pkg` directory.
-/// So that bundler can bundle it together.
-#[wasm_bindgen(module = "/src/platform/web/worker/mainWorker.js")]
-extern "C" {
-    #[wasm_bindgen(js_name = "listenerReady")]
-    fn listener_ready();
-}
+    let canvas = wait_for_canvas_transfer().await;
 
-pub struct MainWorker {
-    handle: web_sys::Worker,
-    name: String,
-    _callback: Closure<dyn FnMut(web_sys::Event)>,
-}
-
-impl Drop for MainWorker {
-    /// Terminates web worker *immediately*.
-    fn drop(&mut self) {
-        self.handle.terminate();
-        info!("Worker({}) was terminated", &self.name);
-    }
-}
-
-impl Debug for MainWorker {
-    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
-        f.debug_struct("MainWorker")
-            .field("handle", &self.handle)
-            .field("name", &self.name)
-            .finish()
-    }
-}
-
-impl MainWorker {
-    /// Spawns main worker from the window context.
-    pub fn spawn(
-        name: &str,
-        id: usize,
-        ring_ptr: u32,
-        f: impl FnOnce() + Send + 'static,
-    ) -> Result<Self, JsValue> {
-        // Creates a new worker.
-        let handle = create_worker("main", name);
-
-        // Double-boxing because `dyn FnOnce` is unsized and so `Box<dyn FnOnce()>` has
-        // an undefined layout (although I think in practice its a pointer and a length?).
-        let ptr = Box::into_raw(Box::new(Box::new(f) as Box<dyn FnOnce()>));
-
-        // Sets default callback.
-        let callback = Closure::new(|_ev| {});
-        handle.set_onmessage(Some(callback.as_ref().unchecked_ref()));
-
-        let msg = js_sys::Object::new();
-        for (key, value) in [
-            ("type", JsValue::from("init")),
-            ("wasmModule", wasm_bindgen::module()),
-            ("workerId", id.into()),
-            ("memory", wasm_bindgen::memory()),
-            ("entryPtr", (ptr as u32).into()),
-            ("ringPtr", ring_ptr.into()),
-        ] {
-            js_sys::Reflect::set(&msg, &key.into(), &value)?;
-        }
-
-        info!("posting message");
-        handle.post_message(&msg)?;
-
-        Ok(Self {
-            handle,
-            name: name.to_owned(),
-            _callback: callback,
-        })
-    }
-
-    pub fn transfer_ownership(&self, canvas: &web_sys::HtmlCanvasElement) {
-        let offscreen_canvas = canvas.transfer_control_to_offscreen().unwrap();
-        let transfer_list = js_sys::Array::new();
-        transfer_list.push(&offscreen_canvas);
-        let msg = js_sys::Object::new();
-        js_sys::Reflect::set(&msg, &"type".into(), &"canvas".into()).unwrap();
-        js_sys::Reflect::set(&msg, &"canvas".into(), &offscreen_canvas).unwrap();
-
-        info!("posting canvas (is_undefined: {})", canvas.is_undefined());
-        self.handle
-            .post_message_with_transfer(&msg, &transfer_list)
-            .unwrap();
-    }
-
-    pub async fn run_render_loop<T: crate::renderer::scene::Scene + 'static>(
-        events_chan: Receiver<WindowEvent>,
-        ring: &'static CommandRing,
-        profile: bool,
-    ) {
-        use crate::renderer::Renderer;
-
-        let canvas = wait_for_canvas_transfer().await;
-
-        let renderer = Rc::new(RefCell::new(
-            Renderer::<T>::new(canvas, events_chan, profile).await,
-        ));
-        renderer.borrow_mut().command_ring = Some(ring);
-        Renderer::run_render_loop(renderer);
-    }
-}
-
-impl Deref for MainWorker {
-    type Target = web_sys::Worker;
-
-    #[inline]
-    fn deref(&self) -> &Self::Target {
-        &self.handle
-    }
+    let renderer = Rc::new(RefCell::new(
+        Renderer::<T>::new(canvas, events_chan, profile).await,
+    ));
+    renderer.borrow_mut().command_ring = Some(ring);
+    Renderer::run_render_loop(renderer);
 }
 
 pub async fn wait_for_canvas_transfer() -> web_sys::OffscreenCanvas {
@@ -149,8 +44,6 @@ pub async fn wait_for_canvas_transfer() -> web_sys::OffscreenCanvas {
             .add_event_listener_with_callback("renderer-canvas", handler.as_ref().unchecked_ref())
             .unwrap();
         handler.forget();
-
-        listener_ready();
     });
 
     let canvas: web_sys::OffscreenCanvas = JsFuture::from(promise)

renderer/src/platform/web/worker/workerGen.js

@@ -1,13 +0,0 @@
-export function createWorker(kind, name) {
-  switch (kind) {
-    case 'main':
-      const main = new Worker(new URL('./mainWorker.js', import.meta.url), {
-        type: 'module',
-        /* @vite-ignore */ name, // vite doesn't allow non static value here.
-      });
-      return main;
-    default:
-      console.log("unsurpported type of worker: ", kind);
-      return undefined;
-  }
-}

renderer/src/render_data/mod.rs

@@ -26,20 +26,6 @@ pub const IDENTITY_MODEL_TRANSFORM: ModelTransform = [
 pub const IDENTITY_NORMAL_MATRIX: NormalMatrix =
     [[1.0, 0.0, 0.0], [0.0, 1.0, 0.0], [0.0, 0.0, 1.0]];
 
-#[repr(transparent)]
-#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)]
-pub struct PipelineKey(u32);
-
-impl PipelineKey {
-    pub const fn new(value: u32) -> Self {
-        Self(value)
-    }
-
-    pub const fn get(self) -> u32 {
-        self.0
-    }
-}
-
 /// Stable CPU-side identity for a device-independent material.
 #[repr(transparent)]
 #[derive(Clone, Copy, Debug, Default, PartialEq, Eq, Hash)]
@@ -91,7 +77,6 @@ pub struct MeshCreateInfo<'a> {
     pub tangents: &'a [[f32; 4]],
     pub uvs: &'a [[f32; 2]],
     pub indices: &'a [u32],
-    pub pipeline: PipelineKey,
     pub material: MaterialKey,
     pub default_instance_type: InstanceType,
     pub default_transform: ModelTransform,
@@ -107,7 +92,6 @@ pub struct CreatedMesh {
 pub struct MeshView {
     pub handle: MeshHandle,
     pub geometry: GeometryRange,
-    pub pipeline: PipelineKey,
     pub material: MaterialKey,
     pub default_instance_type: InstanceType,
     pub local_aabb: Aabb,
@@ -277,7 +261,6 @@ struct MeshSoa {
     vertex_counts: Vec<u32>,
     index_starts: Vec<u32>,
     index_counts: Vec<u32>,
-    pipeline_keys: Vec<PipelineKey>,
     material_keys: Vec<MaterialKey>,
     default_instance_types: Vec<InstanceType>,
     aabb_mins: Vec<[f32; 3]>,
@@ -561,7 +544,6 @@ impl RenderData {
                 index_start: index_range.start,
                 index_count,
             },
-            info.pipeline,
             info.material,
             info.default_instance_type,
             bounds,
@@ -839,7 +821,6 @@ impl MeshSoa {
             vertex_counts: Vec::new(),
             index_starts: Vec::new(),
             index_counts: Vec::new(),
-            pipeline_keys: Vec::new(),
             material_keys: Vec::new(),
             default_instance_types: Vec::new(),
             aabb_mins: Vec::new(),
@@ -859,7 +840,6 @@ impl MeshSoa {
         reserve_vec(&mut self.vertex_counts, target, "meshes")?;
         reserve_vec(&mut self.index_starts, target, "meshes")?;
         reserve_vec(&mut self.index_counts, target, "meshes")?;
-        reserve_vec(&mut self.pipeline_keys, target, "meshes")?;
         reserve_vec(&mut self.material_keys, target, "meshes")?;
         reserve_vec(&mut self.default_instance_types, target, "meshes")?;
         reserve_vec(&mut self.aabb_mins, target, "meshes")?;
@@ -874,7 +854,6 @@ impl MeshSoa {
         &mut self,
         prepared: PreparedSlot,
         geometry: GeometryRange,
-        pipeline: PipelineKey,
         material: MaterialKey,
         default_instance_type: InstanceType,
         bounds: Aabb,
@@ -885,7 +864,6 @@ impl MeshSoa {
         resize_column(&mut self.vertex_counts, len, 0);
         resize_column(&mut self.index_starts, len, 0);
         resize_column(&mut self.index_counts, len, 0);
-        resize_column(&mut self.pipeline_keys, len, PipelineKey::new(0));
         resize_column(&mut self.material_keys, len, MaterialKey::DEFAULT);
         resize_column(&mut self.default_instance_types, len, InstanceType::ZERO);
         resize_column(&mut self.aabb_mins, len, [0.0; 3]);
@@ -897,7 +875,6 @@ impl MeshSoa {
         self.vertex_counts[index] = geometry.vertex_count;
         self.index_starts[index] = geometry.index_start;
         self.index_counts[index] = geometry.index_count;
-        self.pipeline_keys[index] = pipeline;
         self.material_keys[index] = material;
         self.default_instance_types[index] = default_instance_type;
         self.aabb_mins[index] = bounds.min;
@@ -917,7 +894,6 @@ impl MeshSoa {
                 index_start: self.index_starts[index],
                 index_count: self.index_counts[index],
             },
-            pipeline: self.pipeline_keys[index],
             material: self.material_keys[index],
             default_instance_type: self.default_instance_types[index],
             local_aabb: Aabb {

renderer/src/render_data/tests.rs

@@ -14,7 +14,6 @@ fn info() -> MeshCreateInfo<'static> {
         tangents: &TANGENTS,
         uvs: &UVS,
         indices: &INDICES,
-        pipeline: PipelineKey::new(7),
         material: MaterialKey::new(11),
         default_instance_type: InstanceType {
             words: [3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],

renderer/src/render_graph/ast.rs

@@ -0,0 +1,535 @@
+//! Canonical S-expression wire AST.
+//!
+//! Nodes are definitions and `(ref "node" "socket")` forms are references, so one
+//! output may feed any number of consumers without expanding the source expression.
+
+use std::collections::{BTreeMap, HashSet};
+
+use serde_json::Value;
+
+use super::{
+    ComputePipelineDeclaration, ExecutorRef, Graph, GraphError, Node, NodeOutputRef, NodeState,
+    PipelineDeclarations, RenderPipelineDeclaration, MAX_AST_BYTES,
+};
+
+const AST_VERSION: u32 = 1;
+
+#[derive(Clone, Debug, PartialEq)]
+enum SExpr {
+    List(Vec<SExpr>),
+    Atom(String),
+    String(String),
+}
+
+struct Parser<'a> {
+    source: &'a str,
+    offset: usize,
+}
+
+impl Parser<'_> {
+    fn skip_trivia(&mut self) {
+        loop {
+            while self
+                .source
+                .as_bytes()
+                .get(self.offset)
+                .is_some_and(u8::is_ascii_whitespace)
+            {
+                self.offset += 1;
+            }
+            if self.source.as_bytes().get(self.offset) != Some(&b';') {
+                return;
+            }
+            while self
+                .source
+                .as_bytes()
+                .get(self.offset)
+                .is_some_and(|byte| *byte != b'\n')
+            {
+                self.offset += 1;
+            }
+        }
+    }
+
+    fn expression(&mut self) -> Result<SExpr, GraphError> {
+        self.skip_trivia();
+        match self.source.as_bytes().get(self.offset).copied() {
+            Some(b'(') => self.list(),
+            Some(b'"') => self.string(),
+            Some(b')') | None => Err(invalid("expected expression", self.offset)),
+            Some(_) => self.atom(),
+        }
+    }
+
+    fn list(&mut self) -> Result<SExpr, GraphError> {
+        self.offset += 1;
+        let mut values = Vec::new();
+        loop {
+            self.skip_trivia();
+            match self.source.as_bytes().get(self.offset).copied() {
+                Some(b')') => {
+                    self.offset += 1;
+                    return Ok(SExpr::List(values));
+                }
+                None => return Err(invalid("unterminated list", self.offset)),
+                _ => values.push(self.expression()?),
+            }
+        }
+    }
+
+    fn string(&mut self) -> Result<SExpr, GraphError> {
+        let start = self.offset;
+        self.offset += 1;
+        let mut escaped = false;
+        while let Some(byte) = self.source.as_bytes().get(self.offset).copied() {
+            self.offset += 1;
+            if escaped {
+                escaped = false;
+            } else if byte == b'\\' {
+                escaped = true;
+            } else if byte == b'"' {
+                let encoded = &self.source[start..self.offset];
+                let value = serde_json::from_str(encoded)
+                    .map_err(|_| invalid("invalid string literal", start))?;
+                return Ok(SExpr::String(value));
+            }
+        }
+        Err(invalid("unterminated string", start))
+    }
+
+    fn atom(&mut self) -> Result<SExpr, GraphError> {
+        let start = self.offset;
+        while self.source.as_bytes().get(self.offset).is_some_and(|byte| {
+            !byte.is_ascii_whitespace() && !matches!(*byte, b'(' | b')' | b'"' | b';')
+        }) {
+            self.offset += 1;
+        }
+        if start == self.offset {
+            Err(invalid("invalid token", start))
+        } else {
+            Ok(SExpr::Atom(self.source[start..self.offset].to_owned()))
+        }
+    }
+}
+
+fn invalid(message: impl Into<String>, offset: usize) -> GraphError {
+    let message = message.into();
+    GraphError {
+        code: "GRAPH_AST_INVALID",
+        message: message.clone(),
+        details: serde_json::json!({"message":message,"offset":offset}),
+    }
+}
+
+fn list(value: &SExpr) -> Result<&[SExpr], GraphError> {
+    match value {
+        SExpr::List(values) => Ok(values),
+        _ => Err(invalid("expected list", 0)),
+    }
+}
+
+fn atom(value: &SExpr) -> Result<&str, GraphError> {
+    match value {
+        SExpr::Atom(value) => Ok(value),
+        _ => Err(invalid("expected symbol", 0)),
+    }
+}
+
+fn string(value: &SExpr) -> Result<String, GraphError> {
+    match value {
+        SExpr::String(value) => Ok(value.clone()),
+        _ => Err(invalid("expected string", 0)),
+    }
+}
+
+fn u32_value(value: &SExpr) -> Result<u32, GraphError> {
+    atom(value)?.parse().map_err(|_| invalid("expected u32", 0))
+}
+
+fn named_fields<'a>(values: &'a [SExpr]) -> Result<BTreeMap<&'a str, &'a [SExpr]>, GraphError> {
+    let mut fields = BTreeMap::new();
+    for value in values {
+        let field = list(value)?;
+        let Some(name) = field.first() else {
+            return Err(invalid("empty field", 0));
+        };
+        let name = atom(name)?;
+        if fields.insert(name, &field[1..]).is_some() {
+            return Err(invalid(format!("duplicate field '{name}'"), 0));
+        }
+    }
+    Ok(fields)
+}
+
+fn exact_field<'a>(
+    fields: &BTreeMap<&str, &'a [SExpr]>,
+    name: &str,
+    length: usize,
+) -> Result<&'a [SExpr], GraphError> {
+    let values = fields
+        .get(name)
+        .copied()
+        .ok_or_else(|| invalid(format!("missing field '{name}'"), 0))?;
+    if values.len() != length {
+        return Err(invalid(format!("field '{name}' has invalid arity"), 0));
+    }
+    Ok(values)
+}
+
+fn json_value(value: &SExpr) -> Result<Value, GraphError> {
+    match value {
+        SExpr::String(value) => Ok(Value::String(value.clone())),
+        SExpr::Atom(value) if value == "true" => Ok(Value::Bool(true)),
+        SExpr::Atom(value) if value == "false" => Ok(Value::Bool(false)),
+        SExpr::Atom(value) if value == "null" => Ok(Value::Null),
+        SExpr::Atom(value) => serde_json::from_str(value)
+            .map_err(|_| invalid("value atom must be a finite JSON number", 0)),
+        SExpr::List(values)
+            if values.first().and_then(|value| atom(value).ok()) == Some("array") =>
+        {
+            values[1..]
+                .iter()
+                .map(json_value)
+                .collect::<Result<Vec<_>, _>>()
+                .map(Value::Array)
+        }
+        SExpr::List(values)
+            if values.first().and_then(|value| atom(value).ok()) == Some("object") =>
+        {
+            let mut object = serde_json::Map::new();
+            for field in &values[1..] {
+                let field = list(field)?;
+                if field.len() != 3 || atom(&field[0])? != "field" {
+                    return Err(invalid("object entries must be (field string value)", 0));
+                }
+                let key = string(&field[1])?;
+                if object.insert(key.clone(), json_value(&field[2])?).is_some() {
+                    return Err(invalid(format!("duplicate object field '{key}'"), 0));
+                }
+            }
+            Ok(Value::Object(object))
+        }
+        _ => Err(invalid("invalid data value", 0)),
+    }
+}
+
+fn node(value: &SExpr) -> Result<Node, GraphError> {
+    let values = list(value)?;
+    if values.len() < 4 || atom(&values[0])? != "node" {
+        return Err(invalid("invalid node definition", 0));
+    }
+    let id = string(&values[1])?;
+    let state = match atom(&values[2])? {
+        "enabled" => NodeState::Enabled,
+        "muted" => NodeState::Muted,
+        _ => return Err(invalid("node state must be enabled or muted", 0)),
+    };
+    let fields = named_fields(&values[3..])?;
+    if fields.len() != 3 {
+        return Err(invalid("node requires executor, params, and inputs", 0));
+    }
+    let executor = exact_field(&fields, "executor", 2)?;
+    let parameters = json_value(&exact_field(&fields, "params", 1)?[0])?;
+    let input_forms = fields
+        .get("inputs")
+        .copied()
+        .ok_or_else(|| invalid("missing field 'inputs'", 0))?;
+    let mut inputs = BTreeMap::new();
+    for input in input_forms {
+        let input = list(input)?;
+        if input.len() < 2 || atom(&input[0])? != "input" {
+            return Err(invalid("invalid input definition", 0));
+        }
+        let name = string(&input[1])?;
+        let mut references = Vec::new();
+        for reference in &input[2..] {
+            let reference = list(reference)?;
+            if reference.len() != 3 || atom(&reference[0])? != "ref" {
+                return Err(invalid("invalid DAG reference", 0));
+            }
+            references.push(NodeOutputRef {
+                node: string(&reference[1])?,
+                socket: string(&reference[2])?,
+            });
+        }
+        if inputs.insert(name.clone(), references).is_some() {
+            return Err(invalid(format!("duplicate input '{name}'"), 0));
+        }
+    }
+    Ok(Node {
+        id,
+        state,
+        executor: ExecutorRef {
+            key: string(&executor[0])?,
+            version: u32_value(&executor[1])?,
+        },
+        parameters,
+        inputs,
+    })
+}
+
+/// Parses the only render-graph wire format accepted by Yawn core.
+pub fn parse(bytes: &[u8]) -> Result<Graph, GraphError> {
+    if bytes.len() > MAX_AST_BYTES {
+        return Err(GraphError::new(
+            "GRAPH_PAYLOAD_TOO_LARGE",
+            "graph AST exceeds 1 MiB",
+        ));
+    }
+    let source = std::str::from_utf8(bytes)
+        .map_err(|_| GraphError::new("GRAPH_ENCODING_INVALID", "graph AST is not UTF-8"))?;
+    let mut parser = Parser { source, offset: 0 };
+    let root = parser.expression()?;
+    parser.skip_trivia();
+    if parser.offset != source.len() {
+        return Err(invalid("trailing expression", parser.offset));
+    }
+    let root = list(&root)?;
+    if root.len() < 2 || atom(&root[0])? != "yawn-graph" {
+        return Err(invalid("root must be yawn-graph", 0));
+    }
+    if u32_value(&root[1])? != AST_VERSION {
+        return Err(GraphError::new(
+            "GRAPH_SCHEMA_UNSUPPORTED",
+            "render graph AST version must be 1",
+        ));
+    }
+    let fields = named_fields(&root[2..])?;
+    if fields.len() != 4 {
+        return Err(invalid(
+            "graph requires id, revision, pipelines, and nodes",
+            0,
+        ));
+    }
+    let graph_id = string(&exact_field(&fields, "id", 1)?[0])?;
+    let revision = u32_value(&exact_field(&fields, "revision", 1)?[0])?;
+    let pipelines_value = json_value(&exact_field(&fields, "pipelines", 1)?[0])?;
+    let pipelines: PipelineDeclarations = serde_json::from_value(pipelines_value)
+        .map_err(|error| invalid(format!("invalid pipeline declarations: {error}"), 0))?;
+    let nodes = fields
+        .get("nodes")
+        .copied()
+        .ok_or_else(|| invalid("missing field 'nodes'", 0))?
+        .iter()
+        .map(node)
+        .collect::<Result<Vec<_>, _>>()?;
+    Ok(Graph {
+        schema_version: 3,
+        graph_id,
+        revision,
+        pipelines,
+        nodes,
+    })
+}
+
+fn push_string(out: &mut String, value: &str) {
+    out.push_str(&serde_json::to_string(value).expect("strings always serialize"));
+}
+
+fn push_json(out: &mut String, value: &Value) {
+    match value {
+        Value::Null => out.push_str("null"),
+        Value::Bool(value) => out.push_str(if *value { "true" } else { "false" }),
+        Value::Number(value) => out.push_str(&value.to_string()),
+        Value::String(value) => push_string(out, value),
+        Value::Array(values) => {
+            out.push_str("(array");
+            for value in values {
+                out.push(' ');
+                push_json(out, value);
+            }
+            out.push(')');
+        }
+        Value::Object(values) => {
+            out.push_str("(object");
+            let mut fields: Vec<_> = values.iter().collect();
+            fields.sort_by(|left, right| left.0.cmp(right.0));
+            for (name, value) in fields {
+                out.push_str(" (field ");
+                push_string(out, name);
+                out.push(' ');
+                push_json(out, value);
+                out.push(')');
+            }
+            out.push(')');
+        }
+    }
+}
+
+/// Serializes an internal graph for fixtures and cross-language conformance tests.
+pub fn serialize(graph: &Graph) -> String {
+    let mut out = format!("(yawn-graph {AST_VERSION}\n  (id ");
+    push_string(&mut out, &graph.graph_id);
+    out.push_str(&format!(
+        ")\n  (revision {})\n  (pipelines ",
+        graph.revision
+    ));
+    push_json(
+        &mut out,
+        &serde_json::to_value(&graph.pipelines).expect("pipeline declarations serialize"),
+    );
+    out.push_str(")\n  (nodes");
+    for node in &graph.nodes {
+        out.push_str("\n    (node ");
+        push_string(&mut out, &node.id);
+        out.push(' ');
+        out.push_str(match node.state {
+            NodeState::Enabled => "enabled",
+            NodeState::Muted => "muted",
+        });
+        out.push_str("\n      (executor ");
+        push_string(&mut out, &node.executor.key);
+        out.push_str(&format!(" {})\n      (params ", node.executor.version));
+        push_json(&mut out, &node.parameters);
+        out.push_str(")\n      (inputs");
+        for (name, references) in &node.inputs {
+            out.push_str("\n        (input ");
+            push_string(&mut out, name);
+            for reference in references {
+                out.push_str(" (ref ");
+                push_string(&mut out, &reference.node);
+                out.push(' ');
+                push_string(&mut out, &reference.socket);
+                out.push(')');
+            }
+            out.push(')');
+        }
+        out.push_str(")\n    )");
+    }
+    out.push_str("))\n");
+    out
+}
+
+pub(crate) fn validate_pipeline_declarations(graph: &Graph) -> Result<(), GraphError> {
+    let mut names = HashSet::new();
+    let mut shader_bytes = 0usize;
+    for RenderPipelineDeclaration {
+        name,
+        shader,
+        vertex_entry,
+        fragment_entry,
+        ..
+    } in &graph.pipelines.render
+    {
+        if !names.insert(name) {
+            return Err(GraphError::new(
+                "GRAPH_DUPLICATE_ID",
+                format!("duplicate authored pipeline '{name}'"),
+            ));
+        }
+        for identifier in [name, vertex_entry, fragment_entry] {
+            if !super::identifier(identifier) || identifier.len() > 64 {
+                return Err(GraphError::new(
+                    "GRAPH_INVALID_ID",
+                    "invalid authored render pipeline identifier",
+                ));
+            }
+        }
+        if !super::contract(name).is_some_and(|contract| {
+            contract.is_raster_draw() || contract.fullscreen_policy.is_some() || name == "frame_out"
+        }) {
+            return Err(GraphError::new(
+                "GRAPH_EXECUTION_UNSUPPORTED",
+                format!("authored render pipeline '{name}' has no render executor"),
+            ));
+        }
+        shader_bytes = shader_bytes.saturating_add(shader.len());
+    }
+    for ComputePipelineDeclaration {
+        name,
+        shader,
+        entry,
+        dispatch,
+    } in &graph.pipelines.compute
+    {
+        if !names.insert(name) {
+            return Err(GraphError::new(
+                "GRAPH_DUPLICATE_ID",
+                format!("duplicate authored pipeline '{name}'"),
+            ));
+        }
+        for identifier in [name, entry] {
+            if !super::identifier(identifier) || identifier.len() > 64 {
+                return Err(GraphError::new(
+                    "GRAPH_INVALID_ID",
+                    "invalid authored compute pipeline identifier",
+                ));
+            }
+        }
+        if dispatch.contains(&0) {
+            return Err(GraphError::new(
+                "GRAPH_PARAMETERS_INVALID",
+                "compute dispatch dimensions must be nonzero",
+            ));
+        }
+        shader_bytes = shader_bytes.saturating_add(shader.len());
+    }
+    if shader_bytes > MAX_AST_BYTES / 2 {
+        return Err(GraphError::new(
+            "GRAPH_LIMIT_EXCEEDED",
+            "authored shader source exceeds 512 KiB",
+        ));
+    }
+    Ok(())
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn canonical_round_trip_preserves_shared_dag_references() {
+        let graph = Graph {
+            schema_version: 3,
+            graph_id: "dag".into(),
+            revision: 7,
+            pipelines: PipelineDeclarations::default(),
+            nodes: vec![Node {
+                id: "consumer".into(),
+                state: NodeState::Enabled,
+                executor: ExecutorRef {
+                    key: "and".into(),
+                    version: 2,
+                },
+                parameters: serde_json::json!({}),
+                inputs: BTreeMap::from([(
+                    "inputs".into(),
+                    vec![
+                        NodeOutputRef {
+                            node: "shared".into(),
+                            socket: "value".into(),
+                        },
+                        NodeOutputRef {
+                            node: "shared".into(),
+                            socket: "value".into(),
+                        },
+                    ],
+                )]),
+            }],
+        };
+        let encoded = serialize(&graph);
+        let decoded = parse(encoded.as_bytes()).unwrap();
+        assert_eq!(decoded.graph_id, "dag");
+        assert_eq!(decoded.nodes[0].inputs["inputs"].len(), 2);
+        assert_eq!(serialize(&decoded), encoded);
+    }
+
+    #[test]
+    fn rejects_duplicate_fields_and_trailing_expressions() {
+        let duplicate = b"(yawn-graph 1 (id \"x\") (id \"y\") (revision 1) (pipelines (object (field \"render\" (array)) (field \"compute\" (array)))) (nodes))";
+        assert_eq!(parse(duplicate).unwrap_err().code, "GRAPH_AST_INVALID");
+
+        let trailing = b"(yawn-graph 1 (id \"x\") (revision 1) (pipelines (object (field \"render\" (array)) (field \"compute\" (array)))) (nodes)) true";
+        assert_eq!(parse(trailing).unwrap_err().code, "GRAPH_AST_INVALID");
+    }
+
+    #[test]
+    fn rejects_json_and_unknown_top_level_fields() {
+        assert_eq!(
+            parse(br#"{"graphId":"old"}"#).unwrap_err().code,
+            "GRAPH_AST_INVALID"
+        );
+        let source = b"(yawn-graph 1 (id \"x\") (revision 1) (pipelines (object (field \"render\" (array)) (field \"compute\" (array)))) (nodes) (legacy true))";
+        assert_eq!(parse(source).unwrap_err().code, "GRAPH_AST_INVALID");
+    }
+}

renderer/src/render_graph/compiler.rs

@@ -306,25 +306,7 @@ fn validate_name_grammar(s: &str, path: impl Into<String>) -> Result<(), GraphEr
 }
 
 pub fn parse_and_compile(bytes: &[u8]) -> Result<CompiledGraph, GraphError> {
-    if bytes.len() > MAX_JSON_BYTES {
-        return Err(GraphError::new(
-            "GRAPH_PAYLOAD_TOO_LARGE",
-            "graph payload exceeds 1 MiB",
-        ));
-    }
-    let text = std::str::from_utf8(bytes)
-        .map_err(|_| GraphError::new("GRAPH_ENCODING_INVALID", "graph payload is not UTF-8"))?;
-    let probe: serde_json::Value = serde_json::from_str(text)
-        .map_err(|e| GraphError::new("GRAPH_JSON_INVALID", e.to_string()))?;
-    if probe.get("schemaVersion").and_then(|v| v.as_u64()) != Some(3) {
-        return Err(GraphError::new(
-            "GRAPH_SCHEMA_UNSUPPORTED",
-            "schemaVersion must be 3",
-        ));
-    }
-    let graph = serde_json::from_str(text)
-        .map_err(|e| GraphError::new("GRAPH_JSON_INVALID", e.to_string()))?;
-    compile(graph)
+    compile(super::ast::parse(bytes)?)
 }
 
 fn gcd(mut a: u32, mut b: u32) -> u32 {
@@ -834,6 +816,7 @@ fn accepts(c: TypeConstraint, ty: SemanticType) -> bool {
 }
 
 pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
+    super::ast::validate_pipeline_declarations(&graph)?;
     if graph.nodes.len() > MAX_EXECUTIONS {
         return Err(error(
             "GRAPH_LIMIT_EXCEEDED",
@@ -3008,6 +2991,7 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
         graph_id: graph.graph_id,
         revision: graph.revision,
         node_count: graph.nodes.len() as u32,
+        pipelines: graph.pipelines,
         resources,
         executions,
         render_passes,

renderer/src/render_graph/mod.rs

@@ -1,5 +1,6 @@
-//! Device-free render graph compiler and compiled graph registry.
+//! Device-free render graph AST, compiler, and compiled graph registry.
 
+mod ast;
 mod compiler;
 pub(crate) use compiler::execution_attachments;
 mod contracts;
@@ -9,6 +10,7 @@ mod registry;
 mod runtime;
 mod schema;
 
+pub use ast::{parse as parse_ast, serialize as serialize_ast};
 pub use compiler::{compile, parse_and_compile};
 pub use contracts::*;
 pub use expression::*;
@@ -17,7 +19,7 @@ pub use registry::{CompiledGraphId, Registry};
 pub use runtime::*;
 pub use schema::*;
 
-pub const MAX_JSON_BYTES: usize = 1024 * 1024;
+pub const MAX_AST_BYTES: usize = 1024 * 1024;
 pub const MAX_EXECUTIONS: usize = 1024;
 
 #[derive(Debug, Clone, serde::Serialize, PartialEq, Eq)]

renderer/src/render_graph/plan.rs

@@ -9,6 +9,7 @@ pub struct CompiledGraph {
     pub graph_id: String,
     pub revision: u32,
     pub node_count: u32,
+    pub pipelines: PipelineDeclarations,
     pub resources: Vec<CompiledResource>,
     pub executions: Vec<CompiledExecution>,
     pub render_passes: Vec<PhysicalRenderPass>,
@@ -485,6 +486,6 @@ pub enum TextureUsage {
 
 impl CompiledGraph {
     pub fn summary(&self, id: [u32; 2]) -> serde_json::Value {
-        serde_json::json!({"compiledId":id,"graphId":self.graph_id,"revision":self.revision,"schemaVersion":self.schema_version,"nodeCount":self.node_count,"executionCount":self.executions.len(),"physicalPassCount":self.render_passes.len(),"resourceCount":self.resources.len(),"culledNodeCount":self.culled_node_count,"culledResourceCount":self.culled_resource_count,"transientSlotCount":self.transient_slot_count})
+        serde_json::json!({"compiledId":id,"graphId":self.graph_id,"revision":self.revision,"schemaVersion":self.schema_version,"nodeCount":self.node_count,"executionCount":self.executions.len(),"computePassCount":self.pipelines.compute.len(),"physicalPassCount":self.render_passes.len(),"resourceCount":self.resources.len(),"culledNodeCount":self.culled_node_count,"culledResourceCount":self.culled_resource_count,"transientSlotCount":self.transient_slot_count})
     }
 }

renderer/src/render_graph/runtime.rs

@@ -1656,6 +1656,21 @@ pub fn prepare_runtime_plan(
     limits: Option<&wgpu::Limits>,
 ) -> Result<RuntimePlan, GraphError> {
     validate_canonical_plan(graph)?;
+    if let Some(limits) = limits {
+        for (index, compute) in graph.pipelines.compute.iter().enumerate() {
+            if compute
+                .dispatch
+                .iter()
+                .any(|dimension| *dimension > limits.max_compute_workgroups_per_dimension)
+            {
+                return Err(error(
+                    "GRAPH_RESOURCE_LIMIT",
+                    "compute dispatch exceeds adapter limits",
+                    format!("pipelines.compute[{index}].dispatch"),
+                ));
+            }
+        }
+    }
     if surface.width == 0 || surface.height == 0 {
         return Err(error(
             "GRAPH_SURFACE_INCOMPATIBLE",

renderer/src/render_graph/schema.rs

@@ -2,16 +2,18 @@ use std::collections::BTreeMap;
 
 use serde::{Deserialize, Serialize};
 
-#[derive(Debug, Clone, Deserialize)]
+#[derive(Debug, Clone, Deserialize, Serialize)]
 #[serde(deny_unknown_fields, rename_all = "camelCase")]
 pub struct Graph {
     pub schema_version: u32,
     pub graph_id: String,
     pub revision: u32,
+    #[serde(default)]
+    pub pipelines: PipelineDeclarations,
     pub nodes: Vec<Node>,
 }
 
-#[derive(Debug, Clone, Deserialize)]
+#[derive(Debug, Clone, Deserialize, Serialize)]
 #[serde(deny_unknown_fields)]
 pub struct Node {
     pub id: String,
@@ -21,6 +23,41 @@ pub struct Node {
     pub inputs: BTreeMap<String, Vec<NodeOutputRef>>,
 }
 
+/// GPU programs shipped with a graph AST and prepared with the graph loadout.
+#[derive(Debug, Clone, Default, Deserialize, Serialize, PartialEq, Eq)]
+#[serde(deny_unknown_fields, rename_all = "camelCase")]
+pub struct PipelineDeclarations {
+    #[serde(default)]
+    pub render: Vec<RenderPipelineDeclaration>,
+    #[serde(default)]
+    pub compute: Vec<ComputePipelineDeclaration>,
+}
+
+/// A scene render pipeline using Yawn's fixed mesh/instance SOA vertex layout.
+#[derive(Debug, Clone, Deserialize, Serialize, PartialEq, Eq)]
+#[serde(deny_unknown_fields, rename_all = "camelCase")]
+pub struct RenderPipelineDeclaration {
+    pub name: String,
+    pub shader: String,
+    pub vertex_entry: String,
+    pub fragment_entry: String,
+    #[serde(default)]
+    pub double_sided: bool,
+}
+
+/// A binding-free compute pass dispatched before the graph's render passes.
+///
+/// Bindings are deliberately not implicit: shared SOA bindings will be added as an
+/// explicit AST resource contract rather than inferred from shader source.
+#[derive(Debug, Clone, Deserialize, Serialize, PartialEq, Eq)]
+#[serde(deny_unknown_fields, rename_all = "camelCase")]
+pub struct ComputePipelineDeclaration {
+    pub name: String,
+    pub shader: String,
+    pub entry: String,
+    pub dispatch: [u32; 3],
+}
+
 #[derive(Debug, Clone, Copy, Deserialize, Serialize, PartialEq, Eq)]
 #[serde(rename_all = "snake_case")]
 pub enum NodeState {

renderer/src/render_graph/tests.rs

@@ -6,6 +6,11 @@ fn compile_value(value: Value) -> Result<CompiledGraph, GraphError> {
     compile(serde_json::from_value(value).unwrap())
 }
 
+pub(crate) fn ast_bytes(value: &Value) -> Vec<u8> {
+    let graph: Graph = serde_json::from_value(value.clone()).unwrap();
+    super::ast::serialize(&graph).into_bytes()
+}
+
 fn input(node: &str, socket: &str) -> Value {
     json!([{ "node": node, "socket": socket }])
 }

renderer/src/renderer/executors/pipeline.rs

@@ -5,44 +5,6 @@ use crate::renderer::{
 
 use super::super::scene::Scene;
 
-fn encode_scene<'a, T: Scene>(
-    pass: &mut wgpu::RenderPass<'a>,
-    scene: &'a T,
-    gpu: &'a GpuSceneCache,
-    pipelines: &'a PipelineLibrary,
-    materials: &'a MaterialResources,
-) {
-    for (i, bind_group) in scene.bind_groups().iter().enumerate() {
-        pass.set_bind_group(i as u32, bind_group, &[]);
-    }
-    if let (Some(p), Some(n), Some(u), Some(t), Some(i), Some(inst)) = (
-        &gpu.positions.buffer,
-        &gpu.normals.buffer,
-        &gpu.uvs.buffer,
-        &gpu.tangents.buffer,
-        &gpu.indices.buffer,
-        &gpu.instances.buffer,
-    ) {
-        pass.set_vertex_buffer(0, p.slice(..));
-        pass.set_vertex_buffer(1, n.slice(..));
-        pass.set_vertex_buffer(2, u.slice(..));
-        pass.set_vertex_buffer(3, inst.slice(..));
-        pass.set_vertex_buffer(4, t.slice(..));
-        pass.set_index_buffer(i.slice(..), wgpu::IndexFormat::Uint32);
-        for draw in &gpu.draws {
-            pass.set_pipeline(pipelines.get_pipeline(draw.pipeline));
-            if pipelines.requires_material(draw.pipeline) {
-                pass.set_bind_group(2, materials.group(draw.material), &[]);
-            }
-            pass.draw_indexed(
-                draw.indices.clone(),
-                draw.base_vertex,
-                draw.instances.clone(),
-            );
-        }
-    }
-}
-
 pub(crate) fn encode_compiled<T: Scene>(
     encoder: &mut wgpu::CommandEncoder,
     surface: &wgpu::TextureView,
@@ -51,10 +13,24 @@ pub(crate) fn encode_compiled<T: Scene>(
     gpu: &GpuSceneCache,
     pipelines: &PipelineLibrary,
     materials: &MaterialResources,
-    indirect_commands: &wgpu::Buffer,
+    planes: Option<&[[f32; 4]; 6]>,
     mut profile: Option<&mut crate::renderer::profiler::ProfileFrame>,
 ) -> Result<(), &'static str> {
     use crate::render_graph::{NormalizedColorLoad, NormalizedDepthLoad, StoreOp};
+    for compute in &active.compute {
+        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some(&compute.name),
+            timestamp_writes: profile
+                .as_deref_mut()
+                .and_then(|profile| profile.compute_writes(&compute.name)),
+        });
+        pass.set_pipeline(&compute.pipeline);
+        pass.dispatch_workgroups(
+            compute.dispatch[0],
+            compute.dispatch[1],
+            compute.dispatch[2],
+        );
+    }
     let view = |resource: u32| -> Result<&wgpu::TextureView, &'static str> {
         let a = active
             .runtime
@@ -193,10 +169,15 @@ pub(crate) fn encode_compiled<T: Scene>(
                 }
                 PreparedExecution::Pipeline {
                     base,
-                    predicate_ordinal,
+                    predicate,
                     variant,
                     ..
                 } => {
+                    let traversal = active
+                        .runtime
+                        .instance_traversal
+                        .as_ref()
+                        .ok_or("compiled graph instance traversal missing")?;
                     for (i, group) in scene.bind_groups().iter().enumerate() {
                         pass.set_bind_group(i as u32, group, &[]);
                     }
@@ -215,17 +196,30 @@ pub(crate) fn encode_compiled<T: Scene>(
                         pass.set_index_buffer(ix.slice(..), wgpu::IndexFormat::Uint32);
                         pass.set_vertex_buffer(3, inst.slice(..));
                         for (draw_index, draw) in gpu.draws.iter().enumerate() {
+                            if !crate::renderer::instance_filter::evaluate(
+                                traversal,
+                                *predicate,
+                                gpu.instance_records
+                                    .get(draw_index)
+                                    .ok_or("instance record missing")?,
+                                gpu.local_aabb_records
+                                    .get(draw_index)
+                                    .ok_or("local aabb record missing")?,
+                                *gpu.instance_type_records
+                                    .get(draw_index)
+                                    .ok_or("instance type record missing")?,
+                                planes,
+                            )? {
+                                continue;
+                            }
                             pass.set_pipeline(variant);
                             if pipelines.requires_material(*base) {
                                 pass.set_bind_group(2, materials.group(draw.material), &[]);
                             }
-                            pass.draw_indexed_indirect(
-                                indirect_commands,
-                                crate::renderer::instance_traversal::command_offset(
-                                    *predicate_ordinal,
-                                    gpu.draws.len(),
-                                    draw_index,
-                                ),
+                            pass.draw_indexed(
+                                draw.indices.clone(),
+                                draw.base_vertex,
+                                draw.instances.clone(),
                             );
                         }
                     }
@@ -236,18 +230,14 @@ pub(crate) fn encode_compiled<T: Scene>(
     Ok(())
 }
 
-pub(crate) fn encode_immediate<T: Scene>(
+pub(crate) fn encode_immediate(
     encoder: &mut wgpu::CommandEncoder,
     color: &wgpu::TextureView,
     depth: &wgpu::TextureView,
-    scene: &T,
-    gpu: &GpuSceneCache,
-    pipelines: &PipelineLibrary,
-    materials: &MaterialResources,
     profile: Option<&mut crate::renderer::profiler::ProfileFrame>,
 ) {
-    let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
-        label: Some("Immediate pipeline pass"),
+    let _pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+        label: Some("No active render graph"),
         color_attachments: &[Some(wgpu::RenderPassColorAttachment {
             depth_slice: None,
             view: color,
@@ -271,7 +261,6 @@ pub(crate) fn encode_immediate<T: Scene>(
             stencil_ops: None,
         }),
         occlusion_query_set: None,
-        timestamp_writes: profile.and_then(|p| p.render_writes("immediate.pipeline")),
+        timestamp_writes: profile.and_then(|p| p.render_writes("no-active-graph")),
     });
-    encode_scene(&mut pass, scene, gpu, pipelines, materials);
 }

renderer/src/renderer/fullscreen_copy.wgsl

@@ -1,76 +0,0 @@
-@group(0) @binding(0) var source_texture: texture_2d<f32>;
-@group(0) @binding(1) var second_texture: texture_2d<f32>;
-@group(0) @binding(2) var linear_clamp: sampler;
-struct Parameters { values: array<vec4<f32>, 8> }
-@group(0) @binding(3) var<uniform> parameters: Parameters;
-
-struct VertexOut { @builtin(position) position: vec4<f32>, @location(0) uv: vec2<f32> }
-@vertex fn vs_main(@builtin(vertex_index) index: u32) -> VertexOut {
-    let positions = array(vec2(-1.0, -1.0), vec2(3.0, -1.0), vec2(-1.0, 3.0));
-    let p = positions[index];
-    var out: VertexOut; out.position = vec4(p, 0.0, 1.0); out.uv = p * vec2(0.5, -0.5) + vec2(0.5); return out;
-}
-fn sample_source(uv: vec2<f32>) -> vec4<f32> { return textureSampleLevel(source_texture, linear_clamp, uv, 0.0); }
-@fragment fn fs_copy(in: VertexOut) -> @location(0) vec4<f32> { return sample_source(in.uv); }
-
-fn aces(x: vec3<f32>) -> vec3<f32> {
-    return clamp((x * (2.51 * x + vec3(0.03))) / (x * (2.43 * x + vec3(0.59)) + vec3(0.14)), vec3(0.0), vec3(1.0));
-}
-fn linear_to_srgb(x: vec3<f32>) -> vec3<f32> {
-    let safe = clamp(x, vec3(0.0), vec3(1.0));
-    let low = safe * 12.92; let high = 1.055 * pow(safe, vec3(1.0 / 2.4)) - vec3(0.055);
-    return select(high, low, safe <= vec3(0.0031308));
-}
-struct FrameCoordinates { uv: vec2<f32>, contained: bool }
-fn frame_coordinates(position: vec2<f32>, surface: vec2<f32>, source: vec2<f32>, mode: f32) -> FrameCoordinates {
-    if mode < 0.5 { return FrameCoordinates(position / surface, true); }
-    let surface_aspect = surface.x / surface.y; let source_aspect = source.x / source.y; var size = surface;
-    if (mode < 1.5 && source_aspect > surface_aspect) || (mode > 1.5 && source_aspect < surface_aspect) { size.y = surface.x / source_aspect; } else { size.x = surface.y * source_aspect; }
-    let origin = (surface - size) * 0.5;
-    return FrameCoordinates((position - origin) / size, mode > 1.5 || (all(position >= origin) && all(position < origin + size)));
-}
-@fragment fn fs_frame_out(in: VertexOut) -> @location(0) vec4<f32> {
-    let surface=parameters.values[1].yz; let source=vec2<f32>(textureDimensions(source_texture));
-    let coordinates=frame_coordinates(in.position.xy,surface,source,parameters.values[1].x);
-    if !coordinates.contained {
-        var bg=parameters.values[2]; if parameters.values[0].w > 0.5 { bg=vec4(linear_to_srgb(bg.rgb),clamp(bg.a,0.0,1.0)); } return bg;
-    }
-    let sampled=sample_source(coordinates.uv); var rgb: vec3<f32>;
-    if parameters.values[0].x > 0.5 { rgb=max(sampled.rgb*exp2(parameters.values[0].z),vec3(0.0)); if parameters.values[0].y > 1.5 { rgb=aces(rgb); } else if parameters.values[0].y > 0.5 { rgb=rgb/(vec3(1.0)+rgb); } else { rgb=clamp(rgb,vec3(0.0),vec3(1.0)); } } else { rgb=clamp(sampled.rgb,vec3(0.0),vec3(1.0)); }
-    if parameters.values[0].w > 0.5 { rgb=linear_to_srgb(rgb); }
-    return vec4(clamp(rgb,vec3(0.0),vec3(1.0)),clamp(sampled.a,0.0,1.0));
-}
-fn grading_result(source: vec4<f32>, graded: vec3<f32>, factor: f32) -> vec4<f32> { return vec4(mix(source.rgb, graded, vec3(factor)), source.a); }
-@fragment fn fs_color_balance(in: VertexOut) -> @location(0) vec4<f32> {
-    let c=sample_source(in.uv); var graded: vec3<f32>;
-    if parameters.values[0].x < 0.5 {
-        let lift=parameters.values[2].xyz+vec3(parameters.values[0].z); let lifted=(c.rgb-vec3(1.0))*(vec3(2.0)-lift)+vec3(1.0);
-        let gain=parameters.values[4].xyz*parameters.values[1].x; let gained=max(lifted*gain,vec3(0.0));
-        let gamma=max(parameters.values[3].xyz*parameters.values[0].w,vec3(0.000001)); graded=pow(gained,vec3(1.0)/gamma);
-    } else {
-        let slope=parameters.values[7].xyz*parameters.values[1].w; let offset=vec3(parameters.values[1].y)+(parameters.values[5].xyz-vec3(1.0));
-        let power=max(parameters.values[6].xyz*parameters.values[1].z,vec3(0.000001)); graded=pow(max(c.rgb*slope+offset,vec3(0.0)),power);
-    }
-    return grading_result(c,graded,parameters.values[0].y);
-}
-@fragment fn fs_exposure_contrast(in: VertexOut) -> @location(0) vec4<f32> {
-    let c=sample_source(in.uv); let exposed=c.rgb*exp2(parameters.values[0].x); let pivot=parameters.values[0].z;
-    let graded=sign(exposed)*vec3(pivot)*pow(abs(exposed)/vec3(pivot),vec3(parameters.values[0].y)); return grading_result(c,graded,parameters.values[0].w);
-}
-@fragment fn fs_saturation(in: VertexOut) -> @location(0) vec4<f32> { let c=sample_source(in.uv); let l=dot(c.rgb,vec3(0.2126,0.7152,0.0722)); return grading_result(c,mix(vec3(l),c.rgb,vec3(parameters.values[0].x)),parameters.values[0].y); }
-@fragment fn fs_channel_mixer(in: VertexOut) -> @location(0) vec4<f32> { let c=sample_source(in.uv); let graded=vec3(dot(c.rgb,parameters.values[0].xyz),dot(c.rgb,parameters.values[1].xyz),dot(c.rgb,parameters.values[2].xyz)); return grading_result(c,graded,parameters.values[0].w); }
-@fragment fn fs_bloom_extract(in: VertexOut) -> @location(0) vec4<f32> {
-    let c=sample_source(in.uv); let brightness=max(c.r,max(c.g,c.b)); let knee=max(parameters.values[0].y,0.00001); let soft=clamp((brightness-parameters.values[0].x+knee)/(2.0*knee),0.0,1.0); let contribution=max(brightness-parameters.values[0].x,0.0)+soft*soft*knee; return vec4(c.rgb*contribution/max(brightness,0.00001),1.0);
-}
-@fragment fn fs_bloom_blur(in: VertexOut) -> @location(0) vec4<f32> {
-    let size=vec2<f32>(textureDimensions(source_texture)); let step=parameters.values[0].xy*parameters.values[0].z/size;
-    var c=sample_source(in.uv)*0.227027; c+=sample_source(in.uv+step*1.384615)*0.316216; c+=sample_source(in.uv-step*1.384615)*0.316216; c+=sample_source(in.uv+step*3.230769)*0.070270; c+=sample_source(in.uv-step*3.230769)*0.070270; return c;
-}
-@fragment fn fs_bloom_composite(in: VertexOut) -> @location(0) vec4<f32> { let c=sample_source(in.uv); return vec4(c.rgb+textureSampleLevel(second_texture,linear_clamp,in.uv,0.0).rgb*parameters.values[0].x,c.a); }
-fn luminance(c: vec3<f32>) -> f32 { return dot(c,vec3(0.2126,0.7152,0.0722)); }
-@fragment fn fs_luminance_edge(in: VertexOut) -> @location(0) vec4<f32> {
-    let d=1.0/vec2<f32>(textureDimensions(source_texture)); var gx=0.0; var gy=0.0;
-    gx += -luminance(sample_source(in.uv+d*vec2(-1.0,-1.0)).rgb)+luminance(sample_source(in.uv+d*vec2(1.0,-1.0)).rgb); gx += -2.0*luminance(sample_source(in.uv+d*vec2(-1.0,0.0)).rgb)+2.0*luminance(sample_source(in.uv+d*vec2(1.0,0.0)).rgb); gx += -luminance(sample_source(in.uv+d*vec2(-1.0,1.0)).rgb)+luminance(sample_source(in.uv+d*vec2(1.0,1.0)).rgb);
-    gy += -luminance(sample_source(in.uv+d*vec2(-1.0,-1.0)).rgb)-2.0*luminance(sample_source(in.uv+d*vec2(0.0,-1.0)).rgb)-luminance(sample_source(in.uv+d*vec2(1.0,-1.0)).rgb); gy += luminance(sample_source(in.uv+d*vec2(-1.0,1.0)).rgb)+2.0*luminance(sample_source(in.uv+d*vec2(0.0,1.0)).rgb)+luminance(sample_source(in.uv+d*vec2(1.0,1.0)).rgb);
-    let edge=clamp(length(vec2(gx,gy))*parameters.values[0].x,0.0,1.0); return vec4(vec3(edge),1.0);
-}

renderer/src/renderer/gpu_scene.rs

@@ -3,7 +3,7 @@ use std::mem::size_of;
 use bytemuck::{Pod, Zeroable};
 
 use crate::{
-    render_data::{MaterialKey, MeshHandle, PipelineKey},
+    render_data::{MaterialKey, MeshHandle},
     renderer::scene_frame::SceneFramePlan,
 };
 
@@ -18,7 +18,6 @@ pub struct GpuInstance {
 
 #[derive(Clone, Debug, PartialEq, Eq)]
 pub struct DrawItem {
-    pub pipeline: PipelineKey,
     pub material: MaterialKey,
     pub mesh: MeshHandle,
     pub indices: std::ops::Range<u32>,
@@ -33,25 +32,6 @@ pub struct GpuLocalAabb {
     pub max: [f32; 4],
 }
 
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Pod, Zeroable, PartialEq, Eq)]
-pub struct DrawSlotMetadata {
-    pub index_count: u32,
-    pub first_index: u32,
-    pub base_vertex: i32,
-    pub instance_index: u32,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug, Pod, Zeroable, PartialEq, Eq)]
-pub struct DrawIndexedIndirect {
-    pub index_count: u32,
-    pub instance_count: u32,
-    pub first_index: u32,
-    pub base_vertex: i32,
-    pub first_instance: u32,
-}
-
 #[derive(Default)]
 pub struct GpuScenePlan {
     pub positions: Vec<[f32; 3]>,
@@ -63,21 +43,13 @@ pub struct GpuScenePlan {
     pub draws: Vec<DrawItem>,
     pub local_aabbs: Vec<GpuLocalAabb>,
     pub instance_types: Vec<[u32; 16]>,
-    pub draw_metadata: Vec<DrawSlotMetadata>,
 }
 
 impl GpuScenePlan {
     pub fn build(data: &SceneFramePlan) -> Result<Self, &'static str> {
         let mut p = Self::default();
         let mut meshes: Vec<_> = data.meshes.iter().collect();
-        meshes.sort_by_key(|m| {
-            (
-                m.pipeline.get(),
-                m.material.get(),
-                m.handle.slot(),
-                m.handle.generation(),
-            )
-        });
+        meshes.sort_by_key(|m| (m.material.get(), m.handle.slot(), m.handle.generation()));
         for mesh in meshes {
             let occurrences: Vec<_> = data.mesh_occurrence_indices[mesh.occurrence_range.clone()]
                 .iter()
@@ -152,14 +124,7 @@ impl GpuScenePlan {
                 p.instance_types.push(occurrence.instance_type.words);
                 let base_vertex =
                     i32::try_from(vertex_start).map_err(|_| "base vertex exceeds i32")?;
-                p.draw_metadata.push(DrawSlotMetadata {
-                    index_count: mesh.geometry.index_count,
-                    first_index,
-                    base_vertex,
-                    instance_index,
-                });
                 p.draws.push(DrawItem {
-                    pipeline: mesh.pipeline,
                     material: mesh.material,
                     mesh: mesh.handle,
                     indices: first_index
@@ -191,9 +156,9 @@ pub struct GpuSceneCache {
     pub tangents: BufferSlot,
     pub indices: BufferSlot,
     pub instances: BufferSlot,
-    pub local_aabbs: BufferSlot,
-    pub instance_types: BufferSlot,
-    pub draw_metadata: BufferSlot,
+    pub instance_records: Vec<GpuInstance>,
+    pub local_aabb_records: Vec<GpuLocalAabb>,
+    pub instance_type_records: Vec<[u32; 16]>,
     pub draws: Vec<DrawItem>,
 }
 
@@ -248,9 +213,6 @@ impl GpuSceneCache {
             bytes(&p.tangents)?,
             bytes(&p.indices)?,
             bytes(&p.instances)?.max(size_of::<GpuInstance>() as u64),
-            bytes(&p.local_aabbs)?.max(size_of::<GpuLocalAabb>() as u64),
-            bytes(&p.instance_types)?.max(size_of::<[u32; 16]>() as u64),
-            bytes(&p.draw_metadata)?.max(size_of::<DrawSlotMetadata>() as u64),
         ];
         let slots = [
             &mut self.positions,
@@ -259,9 +221,6 @@ impl GpuSceneCache {
             &mut self.tangents,
             &mut self.indices,
             &mut self.instances,
-            &mut self.local_aabbs,
-            &mut self.instance_types,
-            &mut self.draw_metadata,
         ];
         let usage = [
             wgpu::BufferUsages::VERTEX,
@@ -269,10 +228,7 @@ impl GpuSceneCache {
             wgpu::BufferUsages::VERTEX,
             wgpu::BufferUsages::VERTEX,
             wgpu::BufferUsages::INDEX,
-            wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::STORAGE,
-            wgpu::BufferUsages::STORAGE,
-            wgpu::BufferUsages::STORAGE,
-            wgpu::BufferUsages::STORAGE,
+            wgpu::BufferUsages::VERTEX,
         ];
         let mut replaced = false;
         for ((slot, &need), use_) in slots.into_iter().zip(&required).zip(usage) {
@@ -289,19 +245,13 @@ impl GpuSceneCache {
             }
         }
         let zero_instance = GpuInstance::zeroed();
-        let zero_aabb = GpuLocalAabb::zeroed();
-        let zero_type = [0u32; 16];
-        let zero_metadata = DrawSlotMetadata::zeroed();
-        let contents: [&[u8]; 9] = [
+        let contents: [&[u8]; 6] = [
             bytemuck::cast_slice(&p.positions),
             bytemuck::cast_slice(&p.normals),
             bytemuck::cast_slice(&p.uvs),
             bytemuck::cast_slice(&p.tangents),
             bytemuck::cast_slice(&p.indices),
             logical_or_zero(&p.instances, &zero_instance),
-            logical_or_zero(&p.local_aabbs, &zero_aabb),
-            logical_or_zero(&p.instance_types, &zero_type),
-            logical_or_zero(&p.draw_metadata, &zero_metadata),
         ];
         let slots = [
             &self.positions,
@@ -310,9 +260,6 @@ impl GpuSceneCache {
             &self.tangents,
             &self.indices,
             &self.instances,
-            &self.local_aabbs,
-            &self.instance_types,
-            &self.draw_metadata,
         ];
         for (s, c) in slots.into_iter().zip(contents) {
             if !c.is_empty() {
@@ -322,6 +269,9 @@ impl GpuSceneCache {
         if replaced {
             self.buffer_epoch = self.buffer_epoch.wrapping_add(1).max(1)
         }
+        self.instance_records = p.instances;
+        self.local_aabb_records = p.local_aabbs;
+        self.instance_type_records = p.instance_types;
         self.draws = p.draws;
         self.revision = Some(data.revision);
         Ok(())
@@ -367,33 +317,14 @@ mod tests {
         assert_eq!(size_of::<GpuInstance>(), 112);
         assert_eq!(size_of::<GpuLocalAabb>(), 32);
         assert_eq!(size_of::<[u32; 16]>(), 64);
-        assert_eq!(size_of::<DrawSlotMetadata>(), 16);
-        assert_eq!(size_of::<DrawIndexedIndirect>(), 20)
-    }
-    #[test]
-    fn command_offset() {
-        let n = 7u64;
-        assert_eq!((3 * n + 2) * 20, 460)
     }
 
     #[test]
-    fn empty_traversal_records_have_exact_zero_floors() {
+    fn empty_instance_buffer_has_an_exact_zero_floor() {
         assert_eq!(
             logical_or_zero::<GpuInstance>(&[], &GpuInstance::zeroed()),
             [0; 112]
         );
-        assert_eq!(
-            logical_or_zero::<GpuLocalAabb>(&[], &GpuLocalAabb::zeroed()),
-            [0; 32]
-        );
-        assert_eq!(logical_or_zero::<[u32; 16]>(&[], &[0; 16]), [0; 64]);
-        assert_eq!(
-            logical_or_zero::<DrawSlotMetadata>(&[], &DrawSlotMetadata::zeroed()),
-            [0; 16]
-        );
         assert_eq!(required_buffer_capacity(0, 112, 1024), Ok(112));
-        assert_eq!(required_buffer_capacity(0, 32, 1024), Ok(32));
-        assert_eq!(required_buffer_capacity(0, 64, 1024), Ok(64));
-        assert_eq!(required_buffer_capacity(0, 16, 1024), Ok(16));
     }
 }

renderer/src/renderer/instance_filter.rs

@@ -0,0 +1,327 @@
+//! CPU evaluation of graph-owned instance predicates.
+//!
+//! Shader source belongs to graph packages, so core evaluates its small typed
+//! predicate IR directly instead of manufacturing a hidden compute shader.
+
+use crate::render_graph::{
+    BooleanOp, CompareOp, ExprId, ExpressionOp, InstanceTraversalPlan, TypedLiteral,
+};
+
+use super::gpu_scene::{GpuInstance, GpuLocalAabb};
+
+#[derive(Clone, Debug)]
+enum Value {
+    Bool(bool),
+    F32(f32),
+    U32(u32),
+    Vector(Vec<f32>),
+    Matrix(Vec<Vec<f32>>),
+    Type([u32; 16]),
+    Aabb { min: [f32; 3], max: [f32; 3] },
+}
+
+fn literal(value: &TypedLiteral) -> Value {
+    match value {
+        TypedLiteral::Bool(value) => Value::Bool(*value),
+        TypedLiteral::F32(value) => Value::F32(*value),
+        TypedLiteral::U32(value) => Value::U32(*value),
+        TypedLiteral::Vec2(value) => Value::Vector(value.to_vec()),
+        TypedLiteral::Vec3(value) => Value::Vector(value.to_vec()),
+        TypedLiteral::Vec4(value) => Value::Vector(value.to_vec()),
+        TypedLiteral::Mat2(value) => {
+            Value::Matrix(value.iter().map(|column| column.to_vec()).collect())
+        }
+        TypedLiteral::Mat3(value) => {
+            Value::Matrix(value.iter().map(|column| column.to_vec()).collect())
+        }
+        TypedLiteral::Mat4(value) => {
+            Value::Matrix(value.iter().map(|column| column.to_vec()).collect())
+        }
+        TypedLiteral::U32x16(value) => Value::Type(*value),
+        TypedLiteral::LocalAabb { min, max } => Value::Aabb {
+            min: *min,
+            max: *max,
+        },
+    }
+}
+
+fn value<'a>(values: &'a [Value], id: ExprId) -> Result<&'a Value, &'static str> {
+    values.get(id.0 as usize).ok_or("predicate operand missing")
+}
+
+fn boolean(values: &[Value], id: ExprId) -> Result<bool, &'static str> {
+    match value(values, id)? {
+        Value::Bool(value) => Ok(*value),
+        _ => Err("predicate operand is not bool"),
+    }
+}
+
+fn f32_value(values: &[Value], id: ExprId) -> Result<f32, &'static str> {
+    match value(values, id)? {
+        Value::F32(value) => Ok(*value),
+        _ => Err("predicate operand is not f32"),
+    }
+}
+
+fn u32_value(values: &[Value], id: ExprId) -> Result<u32, &'static str> {
+    match value(values, id)? {
+        Value::U32(value) => Ok(*value),
+        _ => Err("predicate operand is not u32"),
+    }
+}
+
+fn compare<T: PartialEq + PartialOrd>(operation: CompareOp, left: T, right: T) -> bool {
+    match operation {
+        CompareOp::GreaterThan => left > right,
+        CompareOp::LessThan => left < right,
+        CompareOp::Equals => left == right,
+    }
+}
+
+fn transformed(model: &[[f32; 4]; 4], point: [f32; 3]) -> [f32; 4] {
+    std::array::from_fn(|row| {
+        model[0][row] * point[0]
+            + model[1][row] * point[1]
+            + model[2][row] * point[2]
+            + model[3][row]
+    })
+}
+
+fn frustum_culled(
+    bounds: ([f32; 3], [f32; 3]),
+    model: &[[f32; 4]; 4],
+    planes: &[[f32; 4]; 6],
+) -> bool {
+    planes.iter().any(|plane| {
+        (0..8).all(|corner| {
+            let local = std::array::from_fn(|axis| {
+                if corner & (1 << axis) == 0 {
+                    bounds.0[axis]
+                } else {
+                    bounds.1[axis]
+                }
+            });
+            let world = transformed(model, local);
+            plane
+                .iter()
+                .zip(world)
+                .map(|(left, right)| left * right)
+                .sum::<f32>()
+                < 0.0
+        })
+    })
+}
+
+/// Evaluates one compiled raster predicate for one dense scene occurrence.
+pub fn evaluate(
+    plan: &InstanceTraversalPlan,
+    predicate: ExprId,
+    instance: &GpuInstance,
+    local_aabb: &GpuLocalAabb,
+    instance_type: [u32; 16],
+    planes: Option<&[[f32; 4]; 6]>,
+) -> Result<bool, &'static str> {
+    let mut values = Vec::with_capacity(plan.expressions.expressions.len());
+    for expression in &plan.expressions.expressions {
+        let result = match &expression.op {
+            ExpressionOp::Literal { literal: item } => literal(item),
+            ExpressionOp::InstanceType { .. } => Value::Type(instance_type),
+            ExpressionOp::LocalAabb { .. } => Value::Aabb {
+                min: local_aabb.min[..3].try_into().unwrap(),
+                max: local_aabb.max[..3].try_into().unwrap(),
+            },
+            ExpressionOp::Not { value: operand } => Value::Bool(!boolean(&values, *operand)?),
+            ExpressionOp::Boolean {
+                operation,
+                operands,
+            } => {
+                let operands = operands
+                    .iter()
+                    .map(|operand| boolean(&values, *operand))
+                    .collect::<Result<Vec<_>, _>>()?;
+                Value::Bool(match operation {
+                    BooleanOp::And => operands.into_iter().all(|item| item),
+                    BooleanOp::Or => operands.into_iter().any(|item| item),
+                    BooleanOp::Xor => operands.into_iter().fold(false, |left, right| left ^ right),
+                    BooleanOp::Xnor => {
+                        !operands.into_iter().fold(false, |left, right| left ^ right)
+                    }
+                })
+            }
+            ExpressionOp::CompareF32 {
+                operation,
+                left,
+                right,
+            } => Value::Bool(compare(
+                *operation,
+                f32_value(&values, *left)?,
+                f32_value(&values, *right)?,
+            )),
+            ExpressionOp::CompareU32 {
+                operation,
+                left,
+                right,
+            } => Value::Bool(compare(
+                *operation,
+                u32_value(&values, *left)?,
+                u32_value(&values, *right)?,
+            )),
+            ExpressionOp::VectorProject { vector, index } => match value(&values, *vector)? {
+                Value::Vector(vector) => Value::F32(
+                    *vector
+                        .get(*index as usize)
+                        .ok_or("vector predicate index out of bounds")?,
+                ),
+                _ => return Err("predicate operand is not vector"),
+            },
+            ExpressionOp::VectorConstruct { components } => Value::Vector(
+                components
+                    .iter()
+                    .map(|component| f32_value(&values, *component))
+                    .collect::<Result<_, _>>()?,
+            ),
+            ExpressionOp::MatrixColumn { matrix, index } => match value(&values, *matrix)? {
+                Value::Matrix(matrix) => Value::Vector(
+                    matrix
+                        .get(*index as usize)
+                        .ok_or("matrix predicate index out of bounds")?
+                        .clone(),
+                ),
+                _ => return Err("predicate operand is not matrix"),
+            },
+            ExpressionOp::MatrixConstruct { columns } => Value::Matrix(
+                columns
+                    .iter()
+                    .map(|column| match value(&values, *column)? {
+                        Value::Vector(column) => Ok(column.clone()),
+                        _ => Err("matrix column is not vector"),
+                    })
+                    .collect::<Result<_, _>>()?,
+            ),
+            ExpressionOp::TypeWord {
+                value: operand,
+                index,
+            } => match value(&values, *operand)? {
+                Value::Type(words) => Value::U32(words[*index as usize]),
+                _ => return Err("predicate operand is not u32x16"),
+            },
+            ExpressionOp::TypeConstruct { words } => {
+                if words.len() != 16 {
+                    return Err("type predicate requires 16 words");
+                }
+                let mut result = [0; 16];
+                for (index, word) in words.iter().enumerate() {
+                    result[index] = u32_value(&values, *word)?;
+                }
+                Value::Type(result)
+            }
+            ExpressionOp::U32Bit {
+                value: operand,
+                index,
+            } => Value::Bool(u32_value(&values, *operand)? & (1 << index) != 0),
+            ExpressionOp::U32Construct { bits } => {
+                if bits.len() > 32 {
+                    return Err("u32 predicate has too many bits");
+                }
+                let mut result = 0;
+                for (index, bit) in bits.iter().enumerate() {
+                    result |= u32::from(boolean(&values, *bit)?) << index;
+                }
+                Value::U32(result)
+            }
+            ExpressionOp::AabbMin { aabb } => match value(&values, *aabb)? {
+                Value::Aabb { min, .. } => Value::Vector(min.to_vec()),
+                _ => return Err("predicate operand is not aabb"),
+            },
+            ExpressionOp::AabbMax { aabb } => match value(&values, *aabb)? {
+                Value::Aabb { max, .. } => Value::Vector(max.to_vec()),
+                _ => return Err("predicate operand is not aabb"),
+            },
+            ExpressionOp::FrustumCulled { local_aabb, .. } => {
+                let bounds = match value(&values, *local_aabb)? {
+                    Value::Aabb { min, max } => (*min, *max),
+                    _ => return Err("predicate operand is not aabb"),
+                };
+                let planes = planes.ok_or("camera frustum missing")?;
+                Value::Bool(frustum_culled(bounds, &instance.model, planes))
+            }
+        };
+        values.push(result);
+    }
+    boolean(&values, predicate)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::render_graph::{
+        Expression, ExpressionPlan, NodeOutputRef, PipelinePredicatePlan, SemanticType,
+    };
+
+    fn origin() -> NodeOutputRef {
+        NodeOutputRef {
+            node: "test".into(),
+            socket: "value".into(),
+        }
+    }
+
+    #[test]
+    fn evaluates_type_bits_without_shader_source() {
+        let plan = InstanceTraversalPlan {
+            mesh: 0,
+            expressions: ExpressionPlan {
+                expressions: vec![
+                    Expression {
+                        semantic_type: SemanticType::U32x16,
+                        op: ExpressionOp::InstanceType { mesh: 0 },
+                        origin: origin(),
+                        mesh_provenance: Some(0),
+                    },
+                    Expression {
+                        semantic_type: SemanticType::U32,
+                        op: ExpressionOp::TypeWord {
+                            value: ExprId(0),
+                            index: 0,
+                        },
+                        origin: origin(),
+                        mesh_provenance: Some(0),
+                    },
+                    Expression {
+                        semantic_type: SemanticType::Bool,
+                        op: ExpressionOp::U32Bit {
+                            value: ExprId(1),
+                            index: 3,
+                        },
+                        origin: origin(),
+                        mesh_provenance: Some(0),
+                    },
+                ],
+            },
+            pipelines: vec![PipelinePredicatePlan {
+                execution: 0,
+                predicate: ExprId(2),
+                ordinal: 0,
+            }],
+            requires_camera: false,
+        };
+        let mut words = [0; 16];
+        words[0] = 8;
+        assert!(evaluate(
+            &plan,
+            ExprId(2),
+            &GpuInstance {
+                model: crate::render_data::IDENTITY_MODEL_TRANSFORM,
+                normal_0: [0.; 4],
+                normal_1: [0.; 4],
+                normal_2: [0.; 4],
+            },
+            &GpuLocalAabb {
+                min: [-1., -1., -1., 0.],
+                max: [1., 1., 1., 0.],
+            },
+            words,
+            None,
+        )
+        .unwrap());
+    }
+}

renderer/src/renderer/instance_traversal.rs

@@ -1,531 +0,0 @@
-//! Graph-owned instance predicate compute support.
-
-use crate::render_graph::{
-    BooleanOp, CompareOp, ExpressionOp, InstanceTraversalPlan, SemanticType, TypedLiteral,
-};
-
-use super::gpu_scene::{DrawIndexedIndirect, GpuSceneCache};
-
-fn f(value: f32) -> Result<String, String> {
-    if !value.is_finite() {
-        return Err("non-finite expression literal".into());
-    }
-    Ok(format!("{value:?}"))
-}
-
-/// Generates the dense, single-invocation traversal body. Expressions are emitted in
-/// `ExprId` order, so shared IR nodes are evaluated exactly once.
-pub fn generate_wgsl(plan: &InstanceTraversalPlan) -> Result<String, String> {
-    let mut s = String::from("struct Params { planes: array<vec4<f32>,6>, instance_count:u32, pipeline_count:u32, _pad:vec2<u32> };\nstruct Inst{model:mat4x4<f32>,n0:vec4<f32>,n1:vec4<f32>,n2:vec4<f32>}; struct Aabb{min:vec4<f32>,max:vec4<f32>}; struct Type16{words:array<u32,16>}; struct Meta{index_count:u32,first_index:u32,base_vertex:i32,instance_index:u32}; struct Cmd{index_count:u32,instance_count:u32,first_index:u32,base_vertex:i32,first_instance:u32};\n@group(0) @binding(0)var<uniform>p:Params; @group(0) @binding(1)var<storage,read>instances:array<Inst>; @group(0) @binding(2)var<storage,read>aabbs:array<Aabb>; @group(0) @binding(3)var<storage,read>types:array<Type16>; @group(0) @binding(4)var<storage,read>metadata:array<Meta>; @group(0) @binding(5)var<storage,read_write>commands:array<Cmd>;\nstruct LocalAabb{min:vec3<f32>,max:vec3<f32>}; fn culled(i:u32,a:LocalAabb)->bool{var outside=false;for(var q=0u;q<6u;q++){var all=true;for(var c=0u;c<8u;c++){let v=vec3<f32>(select(a.min.x,a.max.x,(c&1u)!=0u),select(a.min.y,a.max.y,(c&2u)!=0u),select(a.min.z,a.max.z,(c&4u)!=0u));all=all&&(dot(p.planes[q],instances[i].model*vec4<f32>(v,1.0))<0.0);}outside=outside||all;}return outside;} @compute @workgroup_size(64) fn main(@builtin(global_invocation_id)gid:vec3<u32>){let i=gid.x;if(i>=p.instance_count){return;}\n");
-    for (i, e) in plan.expressions.expressions.iter().enumerate() {
-        let x = |id: crate::render_graph::ExprId| format!("e{}", id.0);
-        let rhs = match &e.op {
-            ExpressionOp::Literal { literal } => match literal {
-                TypedLiteral::Bool(v) => v.to_string(),
-                TypedLiteral::F32(v) => f(*v)?,
-                TypedLiteral::U32(v) => format!("{v}u"),
-                TypedLiteral::Vec2(v) => format!("vec2<f32>({},{})", f(v[0])?, f(v[1])?),
-                TypedLiteral::Vec3(v) => {
-                    format!("vec3<f32>({},{},{})", f(v[0])?, f(v[1])?, f(v[2])?)
-                }
-                TypedLiteral::Vec4(v) => format!(
-                    "vec4<f32>({},{},{},{})",
-                    f(v[0])?,
-                    f(v[1])?,
-                    f(v[2])?,
-                    f(v[3])?
-                ),
-                TypedLiteral::U32x16(v) => format!(
-                    "Type16(array<u32,16>({}))",
-                    v.iter()
-                        .map(|x| format!("{x}u"))
-                        .collect::<Vec<_>>()
-                        .join(",")
-                ),
-                TypedLiteral::LocalAabb { min, max } => format!(
-                    "LocalAabb(vec3<f32>({},{},{}),vec3<f32>({},{},{}))",
-                    f(min[0])?,
-                    f(min[1])?,
-                    f(min[2])?,
-                    f(max[0])?,
-                    f(max[1])?,
-                    f(max[2])?
-                ),
-                TypedLiteral::Mat2(v) => matrix_literal("mat2x2<f32>", v)?,
-                TypedLiteral::Mat3(v) => matrix_literal("mat3x3<f32>", v)?,
-                TypedLiteral::Mat4(v) => matrix_literal("mat4x4<f32>", v)?,
-            },
-            ExpressionOp::InstanceType { .. } => "types[i]".into(),
-            ExpressionOp::LocalAabb { .. } => "LocalAabb(aabbs[i].min.xyz,aabbs[i].max.xyz)".into(),
-            ExpressionOp::Not { value } => format!("!{}", x(*value)),
-            ExpressionOp::Boolean {
-                operation,
-                operands,
-            } => {
-                let identity = matches!(operation, BooleanOp::And | BooleanOp::Xnor);
-                let operator = match operation {
-                    BooleanOp::And => "&&",
-                    BooleanOp::Or => "||",
-                    BooleanOp::Xor | BooleanOp::Xnor => "!=",
-                };
-                let folded = operands
-                    .iter()
-                    .map(|operand| x(*operand))
-                    .reduce(|left, right| format!("({left} {operator} {right})"))
-                    .unwrap_or_else(|| identity.to_string());
-                if matches!(operation, BooleanOp::Xnor) && operands.len() > 1 {
-                    format!("!{folded}")
-                } else if matches!(operation, BooleanOp::Xnor) && !operands.is_empty() {
-                    format!("!({folded})")
-                } else {
-                    folded
-                }
-            }
-            ExpressionOp::CompareF32 {
-                operation,
-                left,
-                right,
-            }
-            | ExpressionOp::CompareU32 {
-                operation,
-                left,
-                right,
-            } => format!(
-                "({} {} {})",
-                x(*left),
-                match operation {
-                    CompareOp::GreaterThan => ">",
-                    CompareOp::LessThan => "<",
-                    CompareOp::Equals => "==",
-                },
-                x(*right)
-            ),
-            ExpressionOp::VectorProject { vector, index } => {
-                let limit =
-                    vector_width(&plan.expressions.expressions[vector.0 as usize].semantic_type)
-                        .ok_or("vector projection source is not a vector")?;
-                fixed_index(*index, limit, "vector projection")?;
-                format!("{}[{}]", x(*vector), index)
-            }
-            ExpressionOp::VectorConstruct { components } => format!(
-                "{}({})",
-                wgsl_type(&e.semantic_type)?,
-                components
-                    .iter()
-                    .map(|id| x(*id))
-                    .collect::<Vec<_>>()
-                    .join(",")
-            ),
-            ExpressionOp::MatrixColumn { matrix, index } => {
-                let limit =
-                    matrix_width(&plan.expressions.expressions[matrix.0 as usize].semantic_type)
-                        .ok_or("matrix projection source is not a matrix")?;
-                fixed_index(*index, limit, "matrix projection")?;
-                format!("{}[{}]", x(*matrix), index)
-            }
-            ExpressionOp::MatrixConstruct { columns } => format!(
-                "{}({})",
-                wgsl_type(&e.semantic_type)?,
-                columns
-                    .iter()
-                    .map(|id| x(*id))
-                    .collect::<Vec<_>>()
-                    .join(",")
-            ),
-            ExpressionOp::TypeWord { value, index } => {
-                fixed_index(*index, 16, "u32x16 projection")?;
-                format!("{}.words[{}]", x(*value), index)
-            }
-            ExpressionOp::TypeConstruct { words } => format!(
-                "Type16(array<u32,16>({}))",
-                words.iter().map(|id| x(*id)).collect::<Vec<_>>().join(",")
-            ),
-            ExpressionOp::U32Bit { value, index } => {
-                fixed_index(*index, 32, "u32 bit projection")?;
-                format!("(({} & (1u<<{}u))!=0u)", x(*value), index)
-            }
-            ExpressionOp::U32Construct { bits } => format!(
-                "({})",
-                bits.iter()
-                    .enumerate()
-                    .map(|(bit, id)| format!("select(0u,{}u,{})", 1u32 << bit, x(*id)))
-                    .collect::<Vec<_>>()
-                    .join("|")
-            ),
-            ExpressionOp::AabbMin { aabb } => format!("{}.min", x(*aabb)),
-            ExpressionOp::AabbMax { aabb } => format!("{}.max", x(*aabb)),
-            ExpressionOp::FrustumCulled { local_aabb, .. } => {
-                format!("culled(i,{})", x(*local_aabb))
-            }
-        };
-        s.push_str(&format!("let e{i}={rhs};\n"));
-    }
-    for entry in &plan.pipelines {
-        s.push_str(&format!("{{let m=metadata[i];commands[{}u*p.instance_count+i]=Cmd(m.index_count,select(0u,1u,e{}),m.first_index,m.base_vertex,m.instance_index);}}\n",entry.ordinal,entry.predicate.0));
-    }
-    s.push('}');
-    if s.len() >= 1024 * 1024 {
-        return Err("generated traversal WGSL exceeds 1 MiB".into());
-    }
-    Ok(s)
-}
-
-fn fixed_index(index: u8, limit: u8, kind: &str) -> Result<(), String> {
-    (index < limit)
-        .then_some(())
-        .ok_or_else(|| format!("invalid fixed {kind} index"))
-}
-fn vector_width(ty: &SemanticType) -> Option<u8> {
-    match ty {
-        SemanticType::Vec2 => Some(2),
-        SemanticType::Vec3 => Some(3),
-        SemanticType::Vec4 => Some(4),
-        _ => None,
-    }
-}
-fn matrix_width(ty: &SemanticType) -> Option<u8> {
-    match ty {
-        SemanticType::Mat2 => Some(2),
-        SemanticType::Mat3 => Some(3),
-        SemanticType::Mat4 => Some(4),
-        _ => None,
-    }
-}
-fn wgsl_type(ty: &SemanticType) -> Result<&'static str, String> {
-    match ty {
-        SemanticType::Vec2 => Ok("vec2<f32>"),
-        SemanticType::Vec3 => Ok("vec3<f32>"),
-        SemanticType::Vec4 => Ok("vec4<f32>"),
-        SemanticType::Mat2 => Ok("mat2x2<f32>"),
-        SemanticType::Mat3 => Ok("mat3x3<f32>"),
-        SemanticType::Mat4 => Ok("mat4x4<f32>"),
-        _ => Err("invalid combine result type".into()),
-    }
-}
-fn matrix_literal<const N: usize>(name: &str, columns: &[[f32; N]; N]) -> Result<String, String> {
-    let values = columns
-        .iter()
-        .flatten()
-        .map(|v| f(*v))
-        .collect::<Result<Vec<_>, _>>()?;
-    Ok(format!("{name}({})", values.join(",")))
-}
-pub fn dispatch_count(instances: u32, pipelines: u32) -> u32 {
-    if pipelines == 0 {
-        0
-    } else {
-        instances.max(1).div_ceil(64)
-    }
-}
-
-pub fn command_offset(predicate_ordinal: u32, instance_count: usize, draw_index: usize) -> u64 {
-    (u64::from(predicate_ordinal) * instance_count as u64 + draw_index as u64)
-        * std::mem::size_of::<DrawIndexedIndirect>() as u64
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
-struct Params {
-    planes: [[f32; 4]; 6],
-    instance_count: u32,
-    pipeline_count: u32,
-    pad: [u32; 2],
-}
-
-pub struct TraversalGpu {
-    graph: crate::render_graph::CompiledGraphId,
-    plan: InstanceTraversalPlan,
-    scene_epoch: u64,
-    draw_count: usize,
-    pipeline: wgpu::ComputePipeline,
-    params: wgpu::Buffer,
-    bind_group: wgpu::BindGroup,
-    pub commands: wgpu::Buffer,
-}
-
-impl TraversalGpu {
-    pub fn matches(
-        &self,
-        graph: crate::render_graph::CompiledGraphId,
-        plan: &InstanceTraversalPlan,
-        scene_epoch: u64,
-        draw_count: usize,
-    ) -> bool {
-        self.graph == graph
-            && self.plan == *plan
-            && self.scene_epoch == scene_epoch
-            && self.draw_count == draw_count
-    }
-    pub fn create(
-        device: &wgpu::Device,
-        graph: crate::render_graph::CompiledGraphId,
-        plan: &InstanceTraversalPlan,
-        gpu: &GpuSceneCache,
-    ) -> Result<Self, String> {
-        let source = generate_wgsl(plan)?;
-        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("instance traversal"),
-            source: wgpu::ShaderSource::Wgsl(source.into()),
-        });
-        let entries = [
-            (0, wgpu::BufferBindingType::Uniform),
-            (1, wgpu::BufferBindingType::Storage { read_only: true }),
-            (2, wgpu::BufferBindingType::Storage { read_only: true }),
-            (3, wgpu::BufferBindingType::Storage { read_only: true }),
-            (4, wgpu::BufferBindingType::Storage { read_only: true }),
-            (5, wgpu::BufferBindingType::Storage { read_only: false }),
-        ]
-        .map(|(binding, ty)| wgpu::BindGroupLayoutEntry {
-            binding,
-            visibility: wgpu::ShaderStages::COMPUTE,
-            ty: wgpu::BindingType::Buffer {
-                ty,
-                has_dynamic_offset: false,
-                min_binding_size: None,
-            },
-            count: None,
-        });
-        let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("instance traversal"),
-            entries: &entries,
-        });
-        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("instance traversal"),
-            bind_group_layouts: &[&layout],
-            push_constant_ranges: &[],
-        });
-        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-            label: Some("instance traversal"),
-            layout: Some(&pipeline_layout),
-            module: &module,
-            entry_point: Some("main"),
-            compilation_options: Default::default(),
-            cache: None,
-        });
-        let params = device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("instance traversal params"),
-            size: std::mem::size_of::<Params>() as u64,
-            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
-            mapped_at_creation: false,
-        });
-        let count = (gpu.draws.len() as u64)
-            .checked_mul(plan.pipelines.len() as u64)
-            .and_then(|n| n.checked_mul(std::mem::size_of::<DrawIndexedIndirect>() as u64))
-            .ok_or("indirect command size overflow")?
-            .max(std::mem::size_of::<DrawIndexedIndirect>() as u64);
-        if count > device.limits().max_buffer_size {
-            return Err("indirect command buffer exceeds device limit".into());
-        }
-        let commands = device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("instance traversal commands"),
-            size: count,
-            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::INDIRECT,
-            mapped_at_creation: false,
-        });
-        fn required(slot: &super::gpu_scene::BufferSlot) -> Result<&wgpu::Buffer, String> {
-            slot.buffer
-                .as_ref()
-                .ok_or_else(|| "instance traversal scene buffer missing".to_owned())
-        }
-        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            label: Some("instance traversal"),
-            layout: &layout,
-            entries: &[
-                wgpu::BindGroupEntry {
-                    binding: 0,
-                    resource: params.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 1,
-                    resource: required(&gpu.instances)?.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 2,
-                    resource: required(&gpu.local_aabbs)?.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 3,
-                    resource: required(&gpu.instance_types)?.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 4,
-                    resource: required(&gpu.draw_metadata)?.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 5,
-                    resource: commands.as_entire_binding(),
-                },
-            ],
-        });
-        Ok(Self {
-            graph,
-            plan: plan.clone(),
-            scene_epoch: gpu.buffer_epoch,
-            draw_count: gpu.draws.len(),
-            pipeline,
-            params,
-            bind_group,
-            commands,
-        })
-    }
-    pub(crate) fn encode(
-        &self,
-        encoder: &mut wgpu::CommandEncoder,
-        queue: &wgpu::Queue,
-        planes: Option<[[f32; 4]; 6]>,
-        instances: u32,
-        mut profile: Option<&mut super::profiler::ProfileFrame>,
-    ) {
-        queue.write_buffer(
-            &self.params,
-            0,
-            bytemuck::bytes_of(&Params {
-                planes: planes.unwrap_or([[0.; 4]; 6]),
-                instance_count: instances,
-                pipeline_count: self.plan.pipelines.len() as u32,
-                pad: [0; 2],
-            }),
-        );
-        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-            label: Some("instance traversal"),
-            timestamp_writes: profile
-                .as_deref_mut()
-                .and_then(|p| p.compute_writes("instance_traversal")),
-        });
-        pass.set_pipeline(&self.pipeline);
-        pass.set_bind_group(0, &self.bind_group, &[]);
-        pass.dispatch_workgroups(
-            dispatch_count(instances, self.plan.pipelines.len() as u32),
-            1,
-            1,
-        );
-    }
-}
-
-#[cfg(test)]
-mod tests {
-    use super::*;
-    use crate::render_graph::{Expression, ExpressionPlan, NodeOutputRef};
-
-    fn expression(semantic_type: SemanticType, op: ExpressionOp) -> Expression {
-        Expression {
-            semantic_type,
-            op,
-            origin: NodeOutputRef {
-                node: "test".into(),
-                socket: "value".into(),
-            },
-            mesh_provenance: None,
-        }
-    }
-
-    #[test]
-    fn pipeline_major_offsets_and_single_dispatch_are_deterministic() {
-        assert_eq!(command_offset(0, 7, 6), 120);
-        assert_eq!(command_offset(1, 7, 0), 140);
-        assert_eq!(command_offset(3, 7, 2), 460);
-        assert_eq!(dispatch_count(1, 4), 1);
-        assert_eq!(dispatch_count(64, 4), 1);
-        assert_eq!(dispatch_count(65, 4), 2);
-        assert_eq!(dispatch_count(0, 4), 1);
-    }
-
-    #[test]
-    fn lowering_helpers_cover_matrices_and_reject_dynamic_indexes() {
-        assert_eq!(
-            matrix_literal("mat2x2<f32>", &[[1.0, 2.0], [3.0, 4.0]]).unwrap(),
-            "mat2x2<f32>(1.0,2.0,3.0,4.0)"
-        );
-        assert!(fixed_index(3, 3, "vector projection").is_err());
-        assert_eq!(wgsl_type(&SemanticType::Mat4).unwrap(), "mat4x4<f32>");
-    }
-
-    #[test]
-    fn variadic_boolean_wgsl_uses_identities_and_ordered_parity() {
-        let expressions = vec![
-            expression(
-                SemanticType::Bool,
-                ExpressionOp::Boolean {
-                    operation: BooleanOp::And,
-                    operands: vec![],
-                },
-            ),
-            expression(
-                SemanticType::Bool,
-                ExpressionOp::Boolean {
-                    operation: BooleanOp::Xor,
-                    operands: vec![],
-                },
-            ),
-            expression(
-                SemanticType::Bool,
-                ExpressionOp::Boolean {
-                    operation: BooleanOp::Xnor,
-                    operands: vec![crate::render_graph::ExprId(0)],
-                },
-            ),
-            expression(
-                SemanticType::Bool,
-                ExpressionOp::Boolean {
-                    operation: BooleanOp::Xnor,
-                    operands: vec![
-                        crate::render_graph::ExprId(0),
-                        crate::render_graph::ExprId(1),
-                        crate::render_graph::ExprId(2),
-                    ],
-                },
-            ),
-        ];
-        let wgsl = generate_wgsl(&InstanceTraversalPlan {
-            mesh: 0,
-            expressions: ExpressionPlan { expressions },
-            pipelines: vec![],
-            requires_camera: false,
-        })
-        .unwrap();
-        assert!(wgsl.contains("let e0=true;"));
-        assert!(wgsl.contains("let e1=false;"));
-        assert!(wgsl.contains("let e2=!(e0);"));
-        assert!(wgsl.contains("let e3=!((e0 != e1) != e2);"));
-    }
-
-    #[test]
-    fn u32_construct_wgsl_is_parenthesized() {
-        let plan = InstanceTraversalPlan {
-            mesh: 0,
-            expressions: ExpressionPlan {
-                expressions: vec![
-                    expression(
-                        SemanticType::Bool,
-                        ExpressionOp::Literal {
-                            literal: TypedLiteral::Bool(true),
-                        },
-                    ),
-                    expression(
-                        SemanticType::Bool,
-                        ExpressionOp::Literal {
-                            literal: TypedLiteral::Bool(false),
-                        },
-                    ),
-                    expression(
-                        SemanticType::U32,
-                        ExpressionOp::U32Construct {
-                            bits: vec![
-                                crate::render_graph::ExprId(0),
-                                crate::render_graph::ExprId(1),
-                            ],
-                        },
-                    ),
-                ],
-            },
-            pipelines: vec![],
-            requires_camera: false,
-        };
-
-        let wgsl = generate_wgsl(&plan).unwrap();
-        assert_eq!(
-            wgsl.lines().find(|line| line.starts_with("let e2=")),
-            Some("let e2=(select(0u,1u,e0)|select(0u,2u,e1));")
-        );
-    }
-}

renderer/src/renderer/mod.rs

@@ -17,13 +17,14 @@ use crate::{
 
 pub mod executors;
 pub mod gpu_scene;
-pub mod instance_traversal;
+pub mod instance_filter;
 pub mod material;
 pub mod pipeline_library;
 pub mod profiler;
 pub mod scene;
 pub mod scene_frame;
 
+use pipeline_library::PipelineKey;
 pub use pipeline_library::PipelineLibrary;
 
 const DEPTH_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float;
@@ -273,22 +274,6 @@ fn pack_frame_out_uniforms(
     Some(FullscreenUniforms { values })
 }
 
-fn resolve_fullscreen_entry(key: &str) -> Option<&'static str> {
-    match key {
-        "fullscreen_copy" => Some("fs_copy"),
-        "frame_out" => Some("fs_frame_out"),
-        "color_balance" => Some("fs_color_balance"),
-        "exposure_contrast" => Some("fs_exposure_contrast"),
-        "saturation" => Some("fs_saturation"),
-        "channel_mixer" => Some("fs_channel_mixer"),
-        "bloom_extract" => Some("fs_bloom_extract"),
-        "bloom_blur" => Some("fs_bloom_blur"),
-        "bloom_composite" => Some("fs_bloom_composite"),
-        "luminance_edge" => Some("fs_luminance_edge"),
-        _ => None,
-    }
-}
-
 #[cfg(test)]
 mod fullscreen_tests {
     use super::*;
@@ -322,46 +307,6 @@ mod fullscreen_tests {
         );
     }
 
-    #[test]
-    fn fullscreen_entries_are_explicit() {
-        assert_eq!(resolve_fullscreen_entry("fullscreen_copy"), Some("fs_copy"));
-        assert_eq!(resolve_fullscreen_entry("frame_out"), Some("fs_frame_out"));
-        assert_eq!(resolve_fullscreen_entry("tone_map"), None);
-        assert_eq!(
-            resolve_fullscreen_entry("color_balance"),
-            Some("fs_color_balance")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("exposure_contrast"),
-            Some("fs_exposure_contrast")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("saturation"),
-            Some("fs_saturation")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("channel_mixer"),
-            Some("fs_channel_mixer")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("bloom_extract"),
-            Some("fs_bloom_extract")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("bloom_blur"),
-            Some("fs_bloom_blur")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("bloom_composite"),
-            Some("fs_bloom_composite")
-        );
-        assert_eq!(
-            resolve_fullscreen_entry("luminance_edge"),
-            Some("fs_luminance_edge")
-        );
-        assert_eq!(resolve_fullscreen_entry("unknown"), None);
-    }
-
     fn surface(format: wgpu::TextureFormat) -> RuntimeSurfaceContract {
         RuntimeSurfaceContract {
             format,
@@ -800,7 +745,7 @@ mod fullscreen_tests {
                 .iter()
                 .all(|v| v.is_finite()));
         }
-        // Both WGSL balance branches are neutral on nonnegative RGB with neutral controls.
+        // Both balance branches are neutral on nonnegative RGB with neutral controls.
         let lgg = c; // lift=0/color=1, gamma=1/color=1, gain=1/color=1
         let ops = c; // offset=0/color=1, power=1/color=1, slope=1/color=1
         assert_eq!(lgg, c);
@@ -815,8 +760,8 @@ struct GpuTextureSlot {
 
 enum PreparedExecution {
     Pipeline {
-        base: crate::render_data::PipelineKey,
-        predicate_ordinal: u32,
+        base: PipelineKey,
+        predicate: crate::render_graph::ExprId,
         variant: wgpu::RenderPipeline,
     },
     Fullscreen {
@@ -826,11 +771,18 @@ enum PreparedExecution {
     },
 }
 
+struct PreparedCompute {
+    name: String,
+    pipeline: wgpu::ComputePipeline,
+    dispatch: [u32; 3],
+}
+
 struct ActiveCompiledGraph {
     id: crate::render_graph::CompiledGraphId,
     graph: crate::render_graph::CompiledGraph,
     runtime: crate::render_graph::RuntimePlan,
     textures: Vec<Vec<GpuTextureSlot>>,
+    compute: Vec<PreparedCompute>,
     executions: Vec<PreparedExecution>,
     _fullscreen_layout: wgpu::BindGroupLayout,
 }
@@ -1033,7 +985,7 @@ mod switch_request_tests {
         let mut graph = crate::render_graph::tests::full_cull_graph();
         graph["graphId"] = serde_json::json!(graph_id);
         graph["revision"] = serde_json::json!(revision);
-        serde_json::to_vec(&graph).unwrap()
+        crate::render_graph::tests::ast_bytes(&graph)
     }
 
     use super::*;
@@ -1133,7 +1085,7 @@ mod switch_request_tests {
         let mut registry = crate::render_graph::Registry::default();
         let mut graph = crate::render_graph::tests::full_cull_graph();
         graph["graphId"] = serde_json::json!("switch");
-        let bytes = serde_json::to_vec(&graph).unwrap();
+        let bytes = crate::render_graph::tests::ast_bytes(&graph);
         let (id, _) = registry.compile(&bytes).unwrap();
         let active = "existing_graph";
         let pending: Option<&str> = None;
@@ -1151,11 +1103,10 @@ mod switch_request_tests {
         assert_eq!(active, "existing_graph");
         assert_eq!(pending, None);
 
-        let invalid_replacement =
-            br#"{"schemaVersion":3,"graphId":"switch","revision":2,"nodes":[],"unexpected":true}"#;
+        let invalid_replacement = b"(yawn-graph 1 (id \"switch\") (revision 2) (pipelines (object (field \"render\" (array)) (field \"compute\" (array)))) (nodes) (unexpected true))";
         assert_eq!(
             registry.compile(invalid_replacement).unwrap_err().code,
-            "GRAPH_JSON_INVALID"
+            "GRAPH_AST_INVALID"
         );
         let stored = registry.get(id).unwrap();
         assert_eq!(stored.revision, 1);
@@ -1296,11 +1247,12 @@ pub struct Renderer<T: scene::Scene> {
     resources: PipelineLibrary,
     scene: T,
     render_data: RenderData,
+    shared_soa: crate::shared_soa::SharedSoaRegistry,
+    shared_soa_init_sent: bool,
     snapshot: crate::shared_snapshot::SharedSnapshot,
     snapshot_init_sent: bool,
     scene_frame: scene_frame::SceneFrameCache,
     gpu_scene: gpu_scene::GpuSceneCache,
-    instance_traversal: Option<instance_traversal::TraversalGpu>,
     materials: material::MaterialResources,
     pub(crate) command_ring: Option<&'static CommandRing>,
     pending_replies: Vec<JsValue>,
@@ -1445,20 +1397,22 @@ impl<T: Scene + 'static> Renderer<T> {
             }
             let outcome: Result<JsValue, &'static str> = (|| match opcode {
                 1 => {
-                    if words[3] > 1 {
+                    if words[4] > 1 {
                         return Err("INVALID_FRAMING");
                     }
-                    let bytes = crate::take_payload(words[2]).ok_or("PAYLOAD_MISSING")?;
+                    let bytes = self
+                        .shared_soa
+                        .read_fixed_bytes(words[2], words[3])
+                        .map_err(|_| "SHARED_UPLOAD_INVALID")?;
                     let imported =
                         crate::gltf::decode_gltf_owned(bytes).map_err(|_| "GLB_INVALID")?;
-                    let pipelines = Self::ensure_gltf_pipelines(&mut self.resources, &self.context);
                     // Build a complete GPU candidate first. Neither the live scene nor
                     // its material epoch changes if image decode/resource creation fails.
                     let prepared_materials = self
                         .materials
                         .prepare(&self.context.device, &self.context.queue, &imported)
                         .map_err(|_| "MATERIAL_INVALID")?;
-                    let installed = install_imported(&mut self.render_data, &imported, pipelines)
+                    let installed = install_imported(&mut self.render_data, &imported)
                         .map_err(|_| "INSTALL_FAILED")?;
                     // RenderData replacement and material publication are adjacent in
                     // this synchronous command, preventing a frame with mixed assets.
@@ -1487,7 +1441,7 @@ impl<T: Scene + 'static> Renderer<T> {
                             (radius * 0.001).max(0.1),
                             (radius * 6.0).max(1.1),
                         );
-                        if words[3] == 1 {
+                        if words[4] == 1 {
                             self.scene.set_camera_look_at(
                                 center + ultraviolet::Vec3::new(0.0, radius * 0.05, 0.0),
                                 center + ultraviolet::Vec3::new(radius, 0.0, 0.0),
@@ -1558,33 +1512,21 @@ impl<T: Scene + 'static> Renderer<T> {
                         .map_err(|e| render_data_error_code(&e))?;
                     Ok(js_sys::Array::of2(&h.slot().into(), &h.generation().into()).into())
                 }
-                5 => {
-                    let h = InstanceHandle::from_parts(words[2], words[3]);
-                    let mut m = [[0.; 4]; 4];
-                    for i in 0..16 {
-                        m[i / 4][i % 4] = f32::from_bits(words[4 + i]);
-                    }
-                    self.render_data
-                        .set_instance_transform(h, m)
-                        .map_err(|e| render_data_error_code(&e))?;
-                    Ok(JsValue::UNDEFINED)
-                }
                 6 => {
                     self.render_data
                         .destroy_instance(InstanceHandle::from_parts(words[2], words[3]))
                         .map_err(|e| render_data_error_code(&e))?;
                     Ok(JsValue::UNDEFINED)
                 }
-                10 => {
-                    self.render_data
-                        .set_instance_type(
-                            InstanceHandle::from_parts(words[2], words[3]),
-                            InstanceType {
-                                words: std::array::from_fn(|i| words[4 + i]),
-                            },
-                        )
-                        .map_err(|e| render_data_error_code(&e))?;
-                    Ok(JsValue::UNDEFINED)
+                11 => {
+                    let bytes = crate::take_payload(words[2]).ok_or("PAYLOAD_MISSING")?;
+                    let descriptor = self
+                        .shared_soa
+                        .allocate_json(&bytes, self.render_data.capacities())
+                        .map_err(|_| "SOA_LAYOUT_INVALID")?;
+                    let json =
+                        serde_json::to_string(&descriptor).map_err(|_| "SOA_LAYOUT_INVALID")?;
+                    Ok(js_sys::JSON::parse(&json).map_err(|_| "SOA_LAYOUT_INVALID")?)
                 }
                 _ => Err("UNKNOWN_OPCODE"),
             })();
@@ -1629,28 +1571,6 @@ impl<T: Scene + 'static> Renderer<T> {
         self.context.depth_view = view;
     }
 
-    fn ensure_gltf_pipelines(
-        resources: &mut PipelineLibrary,
-        context: &RendererContext,
-    ) -> [crate::render_data::PipelineKey; 2] {
-        let layout = gpu_scene::vertex_layouts();
-        let culled = resources.get_or_create_pipeline(
-            &context.device,
-            "gltf_standard",
-            &layout,
-            include_str!("../gltf.wgsl"),
-            context.initial_surface_config.format,
-        );
-        let double_sided = resources.get_or_create_pipeline(
-            &context.device,
-            "gltf_standard_double_sided",
-            &layout,
-            include_str!("../gltf.wgsl"),
-            context.initial_surface_config.format,
-        );
-        [culled, double_sided]
-    }
-
     fn plan_compiled(
         &self,
         graph: &crate::render_graph::CompiledGraph,
@@ -1679,6 +1599,31 @@ impl<T: Scene + 'static> Renderer<T> {
     ) -> Result<ActiveCompiledGraph, crate::render_graph::GraphError> {
         use crate::render_graph::*;
         let fail = |message| GraphError::new("GRAPH_RUNTIME_PLAN_INVALID", message);
+        let vertex_layouts = gpu_scene::vertex_layouts();
+        let render_declarations: std::collections::HashMap<_, _> = graph
+            .pipelines
+            .render
+            .iter()
+            .map(|declaration| (declaration.name.as_str(), declaration))
+            .collect();
+        let authored_pipelines: std::collections::HashMap<_, _> = graph
+            .pipelines
+            .render
+            .iter()
+            .filter(|declaration| {
+                crate::render_graph::contract(&declaration.name)
+                    .is_some_and(crate::render_graph::Contract::is_raster_draw)
+            })
+            .map(|declaration| {
+                let key = self.resources.get_or_create_authored_pipeline(
+                    &self.context.device,
+                    declaration,
+                    &vertex_layouts,
+                    self.context.initial_surface_config.format,
+                );
+                (declaration.name.clone(), key)
+            })
+            .collect();
         let resolved_pipelines = graph
             .executions
             .iter()
@@ -1687,8 +1632,9 @@ impl<T: Scene + 'static> Renderer<T> {
                 let NormalizedParameters::Raster { .. } = &execution.parameters else {
                     return Ok(None);
                 };
-                self.resources
-                    .find_pipeline(&execution.executor.key)
+                authored_pipelines
+                    .get(&execution.executor.key)
+                    .copied()
                     .map(Some)
                     .ok_or_else(|| {
                         GraphError::at(
@@ -1765,6 +1711,44 @@ impl<T: Scene + 'static> Renderer<T> {
                 ));
             }
         }
+        let compute_layout =
+            self.context
+                .device
+                .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                    label: Some("graph compute pipeline layout"),
+                    bind_group_layouts: &[],
+                    push_constant_ranges: &[],
+                });
+        let compute = graph
+            .pipelines
+            .compute
+            .iter()
+            .map(|declaration| {
+                let shader =
+                    self.context
+                        .device
+                        .create_shader_module(wgpu::ShaderModuleDescriptor {
+                            label: Some(&declaration.name),
+                            source: wgpu::ShaderSource::Wgsl(declaration.shader.as_str().into()),
+                        });
+                let pipeline =
+                    self.context
+                        .device
+                        .create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                            label: Some(&declaration.name),
+                            layout: Some(&compute_layout),
+                            module: &shader,
+                            entry_point: Some(&declaration.entry),
+                            compilation_options: Default::default(),
+                            cache: None,
+                        });
+                PreparedCompute {
+                    name: declaration.name.clone(),
+                    pipeline,
+                    dispatch: declaration.dispatch,
+                }
+            })
+            .collect();
         let mut textures = Vec::with_capacity(runtime.allocations.classes.len());
         for class in &runtime.allocations.classes {
             let mut gpu_class = Vec::with_capacity(class.slots.len());
@@ -1861,13 +1845,6 @@ impl<T: Scene + 'static> Renderer<T> {
                     bind_group_layouts: &[&fullscreen_layout],
                     push_constant_ranges: &[],
                 });
-        let shader = self
-            .context
-            .device
-            .create_shader_module(wgpu::ShaderModuleDescriptor {
-                label: Some(" fullscreen"),
-                source: wgpu::ShaderSource::Wgsl(include_str!("fullscreen_copy.wgsl").into()),
-            });
         let sampler = self
             .context
             .device
@@ -1974,15 +1951,26 @@ impl<T: Scene + 'static> Renderer<T> {
                             .descriptor
                             .format
                     };
-                    let entry = resolve_fullscreen_entry(&execution.executor.key)
-                        .ok_or_else(|| fail("fullscreen executor mismatch"))?;
+                    let declaration = render_declarations
+                        .get(execution.executor.key.as_str())
+                        .copied()
+                        .ok_or_else(|| fail("fullscreen pipeline declaration missing"))?;
+                    let shader =
+                        self.context
+                            .device
+                            .create_shader_module(wgpu::ShaderModuleDescriptor {
+                                label: Some(&declaration.name),
+                                source: wgpu::ShaderSource::Wgsl(
+                                    declaration.shader.as_str().into(),
+                                ),
+                            });
                     let pipeline = self.context.device.create_render_pipeline(
                         &wgpu::RenderPipelineDescriptor {
                             label: Some(" post pipeline"),
                             layout: Some(&pipeline_layout),
                             vertex: wgpu::VertexState {
                                 module: &shader,
-                                entry_point: Some("vs_main"),
+                                entry_point: Some(&declaration.vertex_entry),
                                 buffers: &[],
                                 compilation_options: Default::default(),
                             },
@@ -1991,7 +1979,7 @@ impl<T: Scene + 'static> Renderer<T> {
                             multisample: Default::default(),
                             fragment: Some(wgpu::FragmentState {
                                 module: &shader,
-                                entry_point: Some(entry),
+                                entry_point: Some(&declaration.fragment_entry),
                                 targets: &[Some(wgpu::ColorTargetState {
                                     format: target_format,
                                     blend: None,
@@ -2137,13 +2125,13 @@ impl<T: Scene + 'static> Renderer<T> {
                         .map_err(|e| GraphError::new("GRAPH_RUNTIME_PLAN_INVALID", e))?;
                     executions.push(PreparedExecution::Pipeline {
                         base,
-                        predicate_ordinal: runtime
+                        predicate: runtime
                             .instance_traversal
                             .as_ref()
                             .and_then(|p| {
                                 p.pipelines.iter().find(|v| v.execution as usize == index)
                             })
-                            .map(|p| p.ordinal)
+                            .map(|p| p.predicate)
                             .ok_or_else(|| fail("pipeline predicate missing"))?,
                         variant,
                     });
@@ -2156,6 +2144,7 @@ impl<T: Scene + 'static> Renderer<T> {
             graph,
             runtime,
             textures,
+            compute,
             executions,
             _fullscreen_layout: fullscreen_layout,
         })
@@ -2387,9 +2376,10 @@ impl<T: Scene + 'static> Renderer<T> {
         let mut render_data =
             RenderData::new(RenderDataConfig::default()).expect("valid render data config");
         let scene = T::setup(&context, &mut resources, &mut render_data);
+        let shared_soa = crate::shared_soa::SharedSoaRegistry::new(render_data.capacities())
+            .expect("default shared SOA layouts are valid");
         let materials = material::MaterialResources::new(&context.device, &context.queue);
         resources.set_material_bind_group_layout(&materials.layout);
-        Self::ensure_gltf_pipelines(&mut resources, &context);
 
         Self {
             events_chan,
@@ -2397,11 +2387,12 @@ impl<T: Scene + 'static> Renderer<T> {
             scene,
             resources,
             render_data,
+            shared_soa,
+            shared_soa_init_sent: false,
             snapshot: crate::shared_snapshot::SharedSnapshot::new(),
             snapshot_init_sent: false,
             scene_frame: Default::default(),
             gpu_scene: Default::default(),
-            instance_traversal: None,
             materials,
             command_ring: None,
             pending_replies: Vec::new(),
@@ -2433,6 +2424,45 @@ impl<T: Scene + 'static> Renderer<T> {
         if !self.drain_commands() {
             return;
         }
+        let soa_layout_changed = match self
+            .shared_soa
+            .sync_capacities(self.render_data.capacities())
+        {
+            Ok(changed) => changed,
+            Err(error) => {
+                self.post_fatal("SOA_ALLOCATION_FAILED", &error.to_string());
+                return;
+            }
+        };
+        self.shared_soa
+            .synchronize_render_data(&mut self.render_data);
+        if !self.shared_soa_init_sent || soa_layout_changed {
+            let message = js_sys::Object::new();
+            let message_type = if self.shared_soa_init_sent {
+                "soa-layout"
+            } else {
+                "soa-init"
+            };
+            let _ = js_sys::Reflect::set(&message, &"type".into(), &message_type.into());
+            match self.shared_soa.descriptors().and_then(|descriptors| {
+                serde_json::to_string(&descriptors)
+                    .map_err(|_| crate::shared_soa::SharedSoaError::SizeOverflow)
+            }) {
+                Ok(json) => {
+                    if let Ok(descriptors) = js_sys::JSON::parse(&json) {
+                        let _ = js_sys::Reflect::set(&message, &"arrays".into(), &descriptors);
+                        let global =
+                            js_sys::global().unchecked_into::<DedicatedWorkerGlobalScope>();
+                        let _ = global.post_message(&message);
+                        self.shared_soa_init_sent = true;
+                    }
+                }
+                Err(error) => {
+                    self.post_fatal("SOA_ALLOCATION_FAILED", &error.to_string());
+                    return;
+                }
+            }
+        }
         let frame_plan = match self.scene_frame.get_or_build(&self.render_data) {
             Ok(plan) => plan,
             Err(error) => {
@@ -2616,40 +2646,6 @@ impl<T: Scene + 'static> Renderer<T> {
         });
         let encode_result = if let Some(active) = rendering_compiled {
             (|| -> Result<(), &'static str> {
-                let plan = active
-                    .runtime
-                    .instance_traversal
-                    .as_ref()
-                    .ok_or("compiled graph instance traversal missing")?;
-                let rebuild = self.instance_traversal.as_ref().is_none_or(|traversal| {
-                    !traversal.matches(
-                        active.id,
-                        plan,
-                        self.gpu_scene.buffer_epoch,
-                        self.gpu_scene.draws.len(),
-                    )
-                });
-                if rebuild {
-                    self.instance_traversal = Some(
-                        instance_traversal::TraversalGpu::create(
-                            &self.context.device,
-                            active.id,
-                            plan,
-                            &self.gpu_scene,
-                        )
-                        .map_err(|error| {
-                            log::error!("instance traversal preparation failed: {error}");
-                            "instance traversal preparation failed"
-                        })?,
-                    );
-                }
-                self.instance_traversal.as_ref().unwrap().encode(
-                    &mut encoder,
-                    &self.context.queue,
-                    planes,
-                    self.gpu_scene.draws.len() as u32,
-                    profile_frame.as_mut(),
-                );
                 executors::encode_compiled(
                     &mut encoder,
                     &texture_view,
@@ -2658,7 +2654,7 @@ impl<T: Scene + 'static> Renderer<T> {
                     &self.gpu_scene,
                     &self.resources,
                     &self.materials,
-                    &self.instance_traversal.as_ref().unwrap().commands,
+                    planes.as_ref(),
                     profile_frame.as_mut(),
                 )
             })()
@@ -2667,10 +2663,6 @@ impl<T: Scene + 'static> Renderer<T> {
                 &mut encoder,
                 &texture_view,
                 &self.context.depth_view,
-                &self.scene,
-                &self.gpu_scene,
-                &self.resources,
-                &self.materials,
                 profile_frame.as_mut(),
             );
             Ok(())

renderer/src/renderer/pipeline_library.rs

@@ -1,9 +1,21 @@
 use std::{collections::HashMap, num::NonZeroU32};
 
-use crate::render_data::PipelineKey;
-
 use super::DEPTH_FORMAT;
 
+#[repr(transparent)]
+#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)]
+pub(crate) struct PipelineKey(u32);
+
+impl PipelineKey {
+    pub const fn new(value: u32) -> Self {
+        Self(value)
+    }
+
+    pub const fn get(self) -> u32 {
+        self.0
+    }
+}
+
 /// Identity of a set of bind-group layouts registered with this library.
 #[derive(Clone, Copy, Debug, Hash, PartialEq, Eq)]
 pub struct PipelineLayoutKey(u64);
@@ -127,7 +139,6 @@ pub struct PipelineLibrary {
     default_layout: Option<PipelineLayoutKey>,
     material_layout: Option<PipelineLayoutKey>,
     next_layout: u64,
-    named_bases: HashMap<String, (PipelineKey, RenderPipelineKey)>,
     descriptor_cache: HashMap<RenderPipelineKey, PipelineKey>,
 }
 
@@ -141,7 +152,6 @@ impl PipelineLibrary {
             default_layout: None,
             material_layout: None,
             next_layout: 0,
-            named_bases: HashMap::new(),
             descriptor_cache: HashMap::new(),
         }
     }
@@ -181,11 +191,6 @@ impl PipelineLibrary {
         shader: &str,
         format: wgpu::TextureFormat,
     ) -> RenderPipelineSpec {
-        let (vertex_entry, fragment_entry) = if name == "triangle_colored" {
-            ("v_main", "f_main")
-        } else {
-            ("vs_main", "fs_main")
-        };
         let stage = |entry: &str| OwnedProgrammableStage {
             shader_source: shader.to_owned(),
             entry_point: entry.to_owned(),
@@ -198,7 +203,7 @@ impl PipelineLibrary {
             } else {
                 self.default_layout
             },
-            vertex: stage(vertex_entry),
+            vertex: stage("vs_main"),
             vertex_layouts: layouts
                 .iter()
                 .map(|layout| OwnedVertexBufferLayout {
@@ -207,7 +212,7 @@ impl PipelineLibrary {
                     attributes: layout.attributes.to_vec(),
                 })
                 .collect(),
-            fragment: Some(stage(fragment_entry)),
+            fragment: Some(stage("fs_main")),
             primitive: wgpu::PrimitiveState {
                 cull_mode: (name != "gltf_standard_double_sided").then_some(wgpu::Face::Back),
                 ..Default::default()
@@ -230,7 +235,7 @@ impl PipelineLibrary {
     }
 
     /// Creates or reuses a pipeline solely by its owned descriptor identity.
-    pub fn get_or_create_from_spec(
+    pub(crate) fn get_or_create_from_spec(
         &mut self,
         device: &wgpu::Device,
         spec: &RenderPipelineSpec,
@@ -325,60 +330,31 @@ impl PipelineLibrary {
         pipeline_key
     }
 
-    pub fn create_pipeline(
+    /// Creates a graph-owned scene pipeline from its authored declaration.
+    pub(crate) fn get_or_create_authored_pipeline(
         &mut self,
         device: &wgpu::Device,
-        name: &str,
+        declaration: &crate::render_graph::RenderPipelineDeclaration,
         layouts: &[wgpu::VertexBufferLayout],
-        shader: &str,
-        format: wgpu::TextureFormat,
-    ) -> Result<PipelineKey, String> {
-        let spec = self.compatibility_spec(name, layouts, shader, format);
-        let descriptor = spec.key();
-        if let Some((_, existing)) = self.named_bases.get(name) {
-            return Err(if existing == &descriptor {
-                format!("Pipeline '{name}' already exists")
-            } else {
-                format!("Pipeline '{name}' already exists with a different descriptor")
-            });
-        }
-        let key = self.get_or_create_from_spec(device, &spec, Some(name));
-        self.named_bases.insert(name.to_owned(), (key, descriptor));
-        Ok(key)
-    }
-
-    pub fn find_pipeline(&self, name: &str) -> Option<PipelineKey> {
-        self.named_bases.get(name).map(|v| v.0)
-    }
-    pub fn get_or_create_pipeline(
-        &mut self,
-        device: &wgpu::Device,
-        name: &str,
-        layouts: &[wgpu::VertexBufferLayout],
-        shader: &str,
         format: wgpu::TextureFormat,
     ) -> PipelineKey {
-        let wanted = self.compatibility_spec(name, layouts, shader, format).key();
-        if let Some((key, existing)) = self.named_bases.get(name) {
-            assert_eq!(
-                existing, &wanted,
-                "Pipeline '{name}' requested with a different descriptor"
-            );
-            return *key;
-        }
-        self.create_pipeline(device, name, layouts, shader, format)
-            .unwrap_or_else(|e| panic!("Failed to create pipeline '{name}': {e}"))
-    }
-    pub fn get_pipeline(&self, key: PipelineKey) -> &wgpu::RenderPipeline {
-        &self.pipelines[key.get() as usize]
+        let mut spec =
+            self.compatibility_spec(&declaration.name, layouts, &declaration.shader, format);
+        spec.vertex.entry_point = declaration.vertex_entry.clone();
+        spec.fragment
+            .as_mut()
+            .expect("scene pipelines have fragment stages")
+            .entry_point = declaration.fragment_entry.clone();
+        spec.primitive.cull_mode = (!declaration.double_sided).then_some(wgpu::Face::Back);
+        self.get_or_create_from_spec(device, &spec, Some(&declaration.name))
     }
 
-    pub fn requires_material(&self, key: PipelineKey) -> bool {
+    pub(crate) fn requires_material(&self, key: PipelineKey) -> bool {
         self.material_layout.is_some()
             && self.specs[key.get() as usize].layout == self.material_layout
     }
 
-    pub fn create_target_variant(
+    pub(crate) fn create_target_variant(
         &self,
         device: &wgpu::Device,
         base: PipelineKey,

renderer/src/renderer/profiler.rs

@@ -13,6 +13,7 @@ const SLOT_COUNT: usize = 4;
 const QUERY_COUNT: u32 = (MAX_PROFILE_PASSES * 2) as u32;
 #[cfg(any(target_arch = "wasm32", test))]
 const USED_RESOLVE_SIZE: u64 = QUERY_COUNT as u64 * 8;
+#[cfg(any(target_arch = "wasm32", test))]
 const RESOLVE_SIZE: u64 = USED_RESOLVE_SIZE.next_multiple_of(wgpu::QUERY_RESOLVE_BUFFER_ALIGNMENT);
 
 #[derive(Clone, Copy, Debug, PartialEq, Eq)]

renderer/src/renderer/scene_frame.rs

@@ -4,14 +4,13 @@ use thiserror::Error;
 
 use crate::render_data::{
     affine_world_aabb, Aabb, GeometryRange, InstanceHandle, InstanceType, MaterialKey, MeshHandle,
-    ModelTransform, NormalMatrix, PipelineKey, RenderData,
+    ModelTransform, NormalMatrix, RenderData,
 };
 
 #[derive(Clone, Debug)]
 pub struct SceneFrameMesh {
     pub handle: MeshHandle,
     pub geometry: GeometryRange,
-    pub pipeline: PipelineKey,
     pub material: MaterialKey,
     pub instance_type: InstanceType,
     pub local_aabb: Aabb,
@@ -115,7 +114,6 @@ impl SceneFramePlan {
             .map(|(dense, (handle, mesh))| SceneFrameMesh {
                 handle,
                 geometry: mesh.geometry,
-                pipeline: mesh.pipeline,
                 material: mesh.material,
                 instance_type: mesh.default_instance_type,
                 local_aabb: mesh.local_aabb,
@@ -168,7 +166,6 @@ mod tests {
             tangents: &[[1., 0., 0., 1.]; 3],
             uvs: &[[0., 0.]; 3],
             indices: &[0, 1, 2],
-            pipeline: PipelineKey::new(0),
             material: crate::render_data::MaterialKey::DEFAULT,
             default_instance_type: instance_type,
             default_transform: IDENTITY_MODEL_TRANSFORM,

renderer/src/shared_soa.rs

@@ -0,0 +1,826 @@
+//! Extensible, SIMD-aligned SOA columns in shared WebAssembly memory.
+
+use std::{
+    collections::BTreeMap,
+    sync::atomic::{AtomicU32, Ordering},
+};
+
+use serde::{Deserialize, Serialize};
+use thiserror::Error;
+
+use crate::render_data::{InstanceHandle, RenderData, RenderDataCapacities};
+
+pub const MAGIC: u32 = u32::from_le_bytes(*b"YSOA");
+pub const VERSION: u32 = 1;
+pub const HEADER_WORDS: usize = 16;
+pub const DATA_OFFSET: u32 = 64;
+
+#[repr(C, align(64))]
+pub struct SharedBlock([AtomicU32; 16]);
+
+impl SharedBlock {
+    fn zeroed() -> Self {
+        Self(std::array::from_fn(|_| AtomicU32::new(0)))
+    }
+}
+
+#[derive(Clone, Copy, Debug, Deserialize, Serialize, PartialEq, Eq)]
+#[serde(rename_all = "snake_case")]
+pub enum ScalarType {
+    U32,
+    I32,
+    F32,
+}
+
+impl ScalarType {
+    fn tag(self) -> u32 {
+        match self {
+            Self::U32 => 1,
+            Self::I32 => 2,
+            Self::F32 => 3,
+        }
+    }
+}
+
+#[derive(Clone, Copy, Debug, Deserialize, Serialize, PartialEq, Eq)]
+#[serde(rename_all = "snake_case")]
+pub enum ArrayDomain {
+    Mesh,
+    Instance,
+    Fixed,
+}
+
+impl ArrayDomain {
+    fn tag(self) -> u32 {
+        match self {
+            Self::Mesh => 1,
+            Self::Instance => 2,
+            Self::Fixed => 3,
+        }
+    }
+
+    fn capacity(self, capacities: RenderDataCapacities, fixed: Option<u32>) -> Option<u32> {
+        match self {
+            Self::Mesh => Some(capacities.meshes),
+            Self::Instance => Some(capacities.instances),
+            Self::Fixed => fixed,
+        }
+    }
+}
+
+#[derive(Clone, Debug, Deserialize, PartialEq, Eq)]
+#[serde(deny_unknown_fields, rename_all = "camelCase")]
+pub struct ArrayRequest {
+    pub name: String,
+    pub domain: ArrayDomain,
+    pub scalar: ScalarType,
+    pub lanes: u32,
+    pub stride: Option<u32>,
+    pub length: Option<u32>,
+}
+
+#[derive(Clone, Debug, Serialize, PartialEq, Eq)]
+#[serde(rename_all = "camelCase")]
+pub struct ArrayDescriptor {
+    pub id: u32,
+    pub name: String,
+    pub domain: ArrayDomain,
+    pub scalar: ScalarType,
+    pub lanes: u32,
+    pub stride: u32,
+    pub length: u32,
+    pub capacity: u32,
+    pub control_ptr: u32,
+    pub data_offset: u32,
+    pub byte_length: u32,
+    pub layout_epoch: u32,
+    pub writable: bool,
+    #[serde(skip_serializing_if = "Option::is_none")]
+    pub generation_guard: Option<ArrayDomain>,
+}
+
+#[derive(Debug, Error, PartialEq, Eq)]
+pub enum SharedSoaError {
+    #[error("shared SOA request is not valid JSON: {0}")]
+    InvalidJson(String),
+    #[error("shared SOA array name is invalid")]
+    InvalidName,
+    #[error("shared SOA lanes must be in 1..=64")]
+    InvalidLanes,
+    #[error("shared SOA stride must fit all lanes and be a multiple of 16 bytes")]
+    InvalidStride,
+    #[error("fixed shared SOA arrays require a nonzero length")]
+    InvalidLength,
+    #[error("shared SOA array already exists with a different layout")]
+    LayoutConflict,
+    #[error("shared SOA allocation size overflow")]
+    SizeOverflow,
+    #[error("shared SOA allocation failed")]
+    AllocationFailed,
+    #[error("shared SOA array is unknown")]
+    UnknownArray,
+    #[error("shared SOA byte uploads require a packed fixed array")]
+    NotPackedFixed,
+    #[error("shared SOA byte range exceeds the array length")]
+    ByteRange,
+    #[error("shared SOA array is currently being written")]
+    Busy,
+}
+
+struct SharedArray {
+    id: u32,
+    request: ArrayRequest,
+    stride_words: u32,
+    length: u32,
+    capacity: u32,
+    layout_epoch: u32,
+    storage: Box<[SharedBlock]>,
+    retired: Vec<Box<[SharedBlock]>>,
+    consumed_sequence: u32,
+    consumed_slot_sequences: Vec<u32>,
+    generation_guard: Option<ArrayDomain>,
+    writable: bool,
+}
+
+fn generation_guard(name: &str) -> Option<ArrayDomain> {
+    match name {
+        "instance.transform" | "instance.type" => Some(ArrayDomain::Instance),
+        _ => None,
+    }
+}
+
+fn writable(name: &str) -> bool {
+    !matches!(name, "instance.generation" | "mesh.generation")
+}
+
+impl SharedArray {
+    fn new(
+        id: u32,
+        request: ArrayRequest,
+        stride_words: u32,
+        capacity: u32,
+    ) -> Result<Self, SharedSoaError> {
+        let generation_guard = generation_guard(&request.name);
+        let writable = writable(&request.name);
+        let mut array = Self {
+            id,
+            request,
+            stride_words,
+            length: capacity,
+            capacity,
+            layout_epoch: 1,
+            storage: allocate(capacity, stride_words)?,
+            retired: Vec::new(),
+            consumed_sequence: 0,
+            consumed_slot_sequences: vec![0; capacity as usize],
+            generation_guard,
+            writable,
+        };
+        array.initialize_header();
+        Ok(array)
+    }
+
+    fn initialize_header(&mut self) {
+        for (index, value) in [
+            MAGIC,
+            VERSION,
+            self.id,
+            self.request.scalar.tag(),
+            self.request.lanes,
+            self.stride_words,
+            self.length,
+            self.capacity,
+            self.request.domain.tag(),
+            0,
+            self.layout_epoch,
+            0,
+            0,
+            0,
+            0,
+            0,
+        ]
+        .into_iter()
+        .enumerate()
+        {
+            self.word(index).store(value, Ordering::Relaxed);
+        }
+        self.consumed_sequence = 0;
+    }
+
+    fn word(&self, index: usize) -> &AtomicU32 {
+        let block = index / 16;
+        let lane = index % 16;
+        &self.storage[block].0[lane]
+    }
+
+    fn data_word(&self, slot: u32, lane: u32) -> &AtomicU32 {
+        let index = HEADER_WORDS + (slot * self.stride_words + lane) as usize;
+        self.word(index)
+    }
+
+    fn descriptor(&self) -> Result<ArrayDescriptor, SharedSoaError> {
+        let control_ptr = pointer(&self.storage)?;
+        let byte_length = self
+            .capacity
+            .checked_mul(self.stride_words)
+            .and_then(|words| words.checked_mul(4))
+            .ok_or(SharedSoaError::SizeOverflow)?;
+        Ok(ArrayDescriptor {
+            id: self.id,
+            name: self.request.name.clone(),
+            domain: self.request.domain,
+            scalar: self.request.scalar,
+            lanes: self.request.lanes,
+            stride: self.stride_words * 4,
+            length: self.length,
+            capacity: self.capacity,
+            control_ptr,
+            data_offset: DATA_OFFSET,
+            byte_length,
+            layout_epoch: self.layout_epoch,
+            writable: self.writable,
+            generation_guard: self.generation_guard,
+        })
+    }
+
+    fn resize(&mut self, capacity: u32) -> Result<bool, SharedSoaError> {
+        if capacity <= self.capacity {
+            self.length = capacity;
+            self.word(6).store(capacity, Ordering::Release);
+            return Ok(false);
+        }
+        let replacement = allocate(capacity, self.stride_words)?;
+        let old_words = HEADER_WORDS
+            + self
+                .capacity
+                .checked_mul(self.stride_words)
+                .ok_or(SharedSoaError::SizeOverflow)? as usize;
+        for index in HEADER_WORDS..old_words {
+            let block = index / 16;
+            let lane = index % 16;
+            replacement[block].0[lane]
+                .store(self.word(index).load(Ordering::Acquire), Ordering::Relaxed);
+        }
+        let old = std::mem::replace(&mut self.storage, replacement);
+        self.retired.push(old);
+        self.capacity = capacity;
+        self.length = capacity;
+        self.consumed_slot_sequences.resize(capacity as usize, 0);
+        self.layout_epoch = self
+            .layout_epoch
+            .checked_add(1)
+            .ok_or(SharedSoaError::SizeOverflow)?;
+        self.initialize_header();
+        self.word(10).store(self.layout_epoch, Ordering::Relaxed);
+        Ok(true)
+    }
+
+    fn try_lock(&self) -> Option<u32> {
+        let sequence = self.word(9).load(Ordering::Acquire);
+        if sequence & 1 != 0 {
+            return None;
+        }
+        self.word(9)
+            .compare_exchange(
+                sequence,
+                sequence.wrapping_add(1),
+                Ordering::AcqRel,
+                Ordering::Acquire,
+            )
+            .ok()
+            .map(|_| sequence)
+    }
+
+    fn unlock(&mut self, sequence: u32) {
+        let next = sequence.wrapping_add(2) & !1;
+        self.consumed_sequence = next;
+        self.word(9).store(next, Ordering::Release);
+    }
+
+    fn changed(&self) -> bool {
+        let sequence = self.word(9).load(Ordering::Acquire);
+        sequence & 1 == 0 && sequence != self.consumed_sequence
+    }
+}
+
+fn allocate(capacity: u32, stride_words: u32) -> Result<Box<[SharedBlock]>, SharedSoaError> {
+    let words = capacity
+        .checked_mul(stride_words)
+        .and_then(|value| value.checked_add(HEADER_WORDS as u32))
+        .ok_or(SharedSoaError::SizeOverflow)?;
+    let blocks = words.checked_add(15).ok_or(SharedSoaError::SizeOverflow)? / 16;
+    let blocks = usize::try_from(blocks).map_err(|_| SharedSoaError::SizeOverflow)?;
+    let mut storage = Vec::new();
+    storage
+        .try_reserve_exact(blocks)
+        .map_err(|_| SharedSoaError::AllocationFailed)?;
+    storage.resize_with(blocks, SharedBlock::zeroed);
+    Ok(storage.into_boxed_slice())
+}
+
+#[cfg(target_arch = "wasm32")]
+fn pointer(storage: &[SharedBlock]) -> Result<u32, SharedSoaError> {
+    u32::try_from(storage.as_ptr() as usize).map_err(|_| SharedSoaError::SizeOverflow)
+}
+
+#[cfg(not(target_arch = "wasm32"))]
+fn pointer(_storage: &[SharedBlock]) -> Result<u32, SharedSoaError> {
+    Ok(0)
+}
+
+fn valid_name(value: &str) -> bool {
+    let mut chars = value.chars();
+    chars
+        .next()
+        .is_some_and(|character| character.is_ascii_alphabetic())
+        && chars.all(|character| {
+            character.is_ascii_alphanumeric() || matches!(character, '_' | '.' | '-')
+        })
+        && value.len() <= 64
+}
+
+/// Owns stable shared columns. Replaced allocations are retained so stale external
+/// views can never alias newly allocated Rust objects.
+pub struct SharedSoaRegistry {
+    arrays: BTreeMap<String, SharedArray>,
+    next_id: u32,
+    published_instance_generations: BTreeMap<u32, u32>,
+    published_mesh_generations: BTreeMap<u32, u32>,
+}
+
+impl SharedSoaRegistry {
+    pub fn new(capacities: RenderDataCapacities) -> Result<Self, SharedSoaError> {
+        let mut registry = Self {
+            arrays: BTreeMap::new(),
+            next_id: 1,
+            published_instance_generations: BTreeMap::new(),
+            published_mesh_generations: BTreeMap::new(),
+        };
+        for request in [
+            ArrayRequest {
+                name: "instance.transform".into(),
+                domain: ArrayDomain::Instance,
+                scalar: ScalarType::F32,
+                lanes: 16,
+                stride: Some(80),
+                length: None,
+            },
+            ArrayRequest {
+                name: "instance.type".into(),
+                domain: ArrayDomain::Instance,
+                scalar: ScalarType::U32,
+                lanes: 16,
+                stride: Some(80),
+                length: None,
+            },
+            ArrayRequest {
+                name: "instance.generation".into(),
+                domain: ArrayDomain::Instance,
+                scalar: ScalarType::U32,
+                lanes: 1,
+                stride: Some(16),
+                length: None,
+            },
+            ArrayRequest {
+                name: "mesh.generation".into(),
+                domain: ArrayDomain::Mesh,
+                scalar: ScalarType::U32,
+                lanes: 1,
+                stride: Some(16),
+                length: None,
+            },
+        ] {
+            registry.allocate(request, capacities)?;
+        }
+        Ok(registry)
+    }
+
+    pub fn allocate_json(
+        &mut self,
+        bytes: &[u8],
+        capacities: RenderDataCapacities,
+    ) -> Result<ArrayDescriptor, SharedSoaError> {
+        let request = serde_json::from_slice(bytes)
+            .map_err(|error| SharedSoaError::InvalidJson(error.to_string()))?;
+        self.allocate(request, capacities)
+    }
+
+    pub fn allocate(
+        &mut self,
+        request: ArrayRequest,
+        capacities: RenderDataCapacities,
+    ) -> Result<ArrayDescriptor, SharedSoaError> {
+        if !valid_name(&request.name) {
+            return Err(SharedSoaError::InvalidName);
+        }
+        if !(1..=64).contains(&request.lanes) {
+            return Err(SharedSoaError::InvalidLanes);
+        }
+        let physical_lanes = request
+            .lanes
+            .checked_add(if generation_guard(&request.name).is_some() {
+                2
+            } else {
+                0
+            })
+            .ok_or(SharedSoaError::SizeOverflow)?;
+        let minimum_stride = physical_lanes
+            .checked_mul(4)
+            .ok_or(SharedSoaError::SizeOverflow)?;
+        let stride = request
+            .stride
+            .unwrap_or_else(|| (minimum_stride + 15) & !15);
+        if stride < minimum_stride || stride % 16 != 0 {
+            return Err(SharedSoaError::InvalidStride);
+        }
+        let capacity = request
+            .domain
+            .capacity(capacities, request.length)
+            .filter(|capacity| *capacity > 0)
+            .ok_or(SharedSoaError::InvalidLength)?;
+        if let Some(existing) = self.arrays.get_mut(&request.name) {
+            if existing.request.domain != request.domain
+                || existing.request.scalar != request.scalar
+                || existing.request.lanes != request.lanes
+                || existing.stride_words != stride / 4
+            {
+                return Err(SharedSoaError::LayoutConflict);
+            }
+            if request.domain == ArrayDomain::Fixed {
+                existing.resize(capacity)?;
+            }
+            return existing.descriptor();
+        }
+        let id = self.next_id;
+        self.next_id = self
+            .next_id
+            .checked_add(1)
+            .ok_or(SharedSoaError::SizeOverflow)?;
+        let name = request.name.clone();
+        let array = SharedArray::new(id, request, stride / 4, capacity)?;
+        let descriptor = array.descriptor()?;
+        self.arrays.insert(name, array);
+        Ok(descriptor)
+    }
+
+    pub fn descriptors(&self) -> Result<Vec<ArrayDescriptor>, SharedSoaError> {
+        self.arrays.values().map(SharedArray::descriptor).collect()
+    }
+
+    /// Copies a stable byte snapshot from a packed fixed array. Writers publish a
+    /// complete upload with the same sequence lock used by all shared SOA columns.
+    pub fn read_fixed_bytes(
+        &mut self,
+        id: u32,
+        byte_length: u32,
+    ) -> Result<Vec<u8>, SharedSoaError> {
+        let array = self
+            .arrays
+            .values_mut()
+            .find(|array| array.id == id)
+            .ok_or(SharedSoaError::UnknownArray)?;
+        if array.request.domain != ArrayDomain::Fixed
+            || array.request.scalar != ScalarType::U32
+            || array.stride_words != array.request.lanes
+        {
+            return Err(SharedSoaError::NotPackedFixed);
+        }
+        let available = array
+            .length
+            .checked_mul(array.request.lanes)
+            .and_then(|words| words.checked_mul(4))
+            .ok_or(SharedSoaError::SizeOverflow)?;
+        if byte_length == 0 || byte_length > available {
+            return Err(SharedSoaError::ByteRange);
+        }
+        let mut bytes = Vec::new();
+        bytes
+            .try_reserve_exact(byte_length as usize)
+            .map_err(|_| SharedSoaError::AllocationFailed)?;
+        let sequence = array.try_lock().ok_or(SharedSoaError::Busy)?;
+        let word_length = byte_length.div_ceil(4);
+        for index in 0..word_length {
+            let slot = index / array.request.lanes;
+            let lane = index % array.request.lanes;
+            bytes.extend_from_slice(
+                &array
+                    .data_word(slot, lane)
+                    .load(Ordering::Acquire)
+                    .to_le_bytes(),
+            );
+        }
+        bytes.truncate(byte_length as usize);
+        array.unlock(sequence);
+        Ok(bytes)
+    }
+
+    /// Reallocates matching-domain columns before a frame. Old blocks stay pinned.
+    pub fn sync_capacities(
+        &mut self,
+        capacities: RenderDataCapacities,
+    ) -> Result<bool, SharedSoaError> {
+        let mut changed = false;
+        for array in self.arrays.values_mut() {
+            if array.request.domain == ArrayDomain::Fixed {
+                continue;
+            }
+            let capacity = array
+                .request
+                .domain
+                .capacity(capacities, None)
+                .expect("non-fixed domains always resolve");
+            changed |= array.resize(capacity)?;
+        }
+        Ok(changed)
+    }
+
+    fn take_instance_words(
+        &mut self,
+        name: &str,
+        handles: &[InstanceHandle],
+    ) -> Option<Vec<(InstanceHandle, Vec<u32>)>> {
+        let array = self.arrays.get_mut(name)?;
+        if !array.changed() {
+            return None;
+        }
+        let sequence = array.try_lock()?;
+        let mut values = Vec::new();
+        for handle in handles.iter().copied() {
+            let slot = handle.slot() as usize;
+            let mutation_sequence = array
+                .data_word(handle.slot(), array.request.lanes + 1)
+                .load(Ordering::Acquire);
+            if array.consumed_slot_sequences[slot] == mutation_sequence {
+                continue;
+            }
+            array.consumed_slot_sequences[slot] = mutation_sequence;
+            let expected_generation = array
+                .data_word(handle.slot(), array.request.lanes)
+                .load(Ordering::Acquire);
+            if expected_generation == handle.generation() {
+                let words = (0..array.request.lanes)
+                    .map(|lane| array.data_word(handle.slot(), lane).load(Ordering::Acquire))
+                    .collect();
+                values.push((handle, words));
+            }
+        }
+        array.unlock(sequence);
+        Some(values)
+    }
+
+    /// Publishes newly live slots, then applies committed mutable columns. Existing
+    /// slots are never republished, so a concurrent shared-memory write cannot be
+    /// overwritten by an unrelated render-data revision.
+    pub fn synchronize_render_data(&mut self, data: &mut RenderData) {
+        self.publish_instance_handles(data);
+        self.publish_mesh_handles(data);
+        let handles: Vec<_> = data.instances().map(|(handle, _)| handle).collect();
+        if let Some(transforms) = self.take_instance_words("instance.transform", &handles) {
+            for (handle, words) in transforms {
+                let mut transform = [[0.0; 4]; 4];
+                for (index, word) in words.into_iter().enumerate() {
+                    transform[index / 4][index % 4] = f32::from_bits(word);
+                }
+                let _ = data.set_instance_transform(handle, transform);
+            }
+        }
+        if let Some(types) = self.take_instance_words("instance.type", &handles) {
+            for (handle, words) in types {
+                let Ok(words) = <Vec<u32> as TryInto<[u32; 16]>>::try_into(words) else {
+                    continue;
+                };
+                let _ = data.set_instance_type(handle, crate::render_data::InstanceType { words });
+            }
+        }
+    }
+
+    fn publish_instance_handles(&mut self, data: &RenderData) {
+        let instances: Vec<_> = data.instances().map(|(_, value)| value).collect();
+        let current: BTreeMap<_, _> = instances
+            .iter()
+            .map(|instance| (instance.handle.slot(), instance.handle.generation()))
+            .collect();
+        let changed: Vec<_> = instances
+            .iter()
+            .filter(|instance| {
+                self.published_instance_generations
+                    .get(&instance.handle.slot())
+                    != Some(&instance.handle.generation())
+            })
+            .collect();
+        let removed: Vec<_> = self
+            .published_instance_generations
+            .keys()
+            .filter(|slot| !current.contains_key(slot))
+            .copied()
+            .collect();
+        if changed.is_empty() && removed.is_empty() {
+            return;
+        }
+        for name in ["instance.transform", "instance.type"] {
+            let Some(array) = self.arrays.get_mut(name) else {
+                return;
+            };
+            let Some(sequence) = array.try_lock() else {
+                return;
+            };
+            for instance in &changed {
+                match name {
+                    "instance.transform" => {
+                        for (lane, value) in instance.model.iter().flatten().enumerate() {
+                            array
+                                .data_word(instance.handle.slot(), lane as u32)
+                                .store(value.to_bits(), Ordering::Relaxed);
+                        }
+                    }
+                    "instance.type" => {
+                        for (lane, value) in instance.instance_type.words.iter().enumerate() {
+                            array
+                                .data_word(instance.handle.slot(), lane as u32)
+                                .store(*value, Ordering::Relaxed);
+                        }
+                    }
+                    _ => unreachable!(),
+                }
+                array
+                    .data_word(instance.handle.slot(), array.request.lanes)
+                    .store(instance.handle.generation(), Ordering::Relaxed);
+                let mutation_sequence = array
+                    .data_word(instance.handle.slot(), array.request.lanes + 1)
+                    .load(Ordering::Relaxed);
+                array.consumed_slot_sequences[instance.handle.slot() as usize] = mutation_sequence;
+            }
+            array.unlock(sequence);
+        }
+        let Some(generations) = self.arrays.get_mut("instance.generation") else {
+            return;
+        };
+        let Some(sequence) = generations.try_lock() else {
+            return;
+        };
+        for slot in removed {
+            generations.data_word(slot, 0).store(0, Ordering::Relaxed);
+        }
+        for instance in changed {
+            generations
+                .data_word(instance.handle.slot(), 0)
+                .store(instance.handle.generation(), Ordering::Relaxed);
+        }
+        generations.unlock(sequence);
+        self.published_instance_generations = current;
+    }
+
+    fn publish_mesh_handles(&mut self, data: &RenderData) {
+        let meshes: Vec<_> = data.meshes().map(|(_, value)| value).collect();
+        let current: BTreeMap<_, _> = meshes
+            .iter()
+            .map(|mesh| (mesh.handle.slot(), mesh.handle.generation()))
+            .collect();
+        if current == self.published_mesh_generations {
+            return;
+        }
+        let Some(generations) = self.arrays.get_mut("mesh.generation") else {
+            return;
+        };
+        let Some(sequence) = generations.try_lock() else {
+            return;
+        };
+        for slot in self
+            .published_mesh_generations
+            .keys()
+            .filter(|slot| !current.contains_key(slot))
+        {
+            generations.data_word(*slot, 0).store(0, Ordering::Relaxed);
+        }
+        for mesh in meshes.iter().filter(|mesh| {
+            self.published_mesh_generations.get(&mesh.handle.slot())
+                != Some(&mesh.handle.generation())
+        }) {
+            generations
+                .data_word(mesh.handle.slot(), 0)
+                .store(mesh.handle.generation(), Ordering::Relaxed);
+        }
+        generations.unlock(sequence);
+        self.published_mesh_generations = current;
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn capacities(instances: u32) -> RenderDataCapacities {
+        RenderDataCapacities {
+            vertices: 0,
+            indices: 0,
+            meshes: 8,
+            instances,
+        }
+    }
+
+    #[test]
+    fn custom_layouts_are_aligned_idempotent_and_conflict_checked() {
+        let mut registry = SharedSoaRegistry::new(capacities(8)).unwrap();
+        let request = ArrayRequest {
+            name: "instance.velocity".into(),
+            domain: ArrayDomain::Instance,
+            scalar: ScalarType::F32,
+            lanes: 3,
+            stride: None,
+            length: None,
+        };
+        let first = registry.allocate(request.clone(), capacities(8)).unwrap();
+        let second = registry.allocate(request, capacities(8)).unwrap();
+        assert_eq!(first, second);
+        assert_eq!((first.stride, first.capacity), (16, 8));
+        let conflict = ArrayRequest {
+            lanes: 4,
+            ..serde_json::from_str::<ArrayRequest>(
+                r#"{"name":"instance.velocity","domain":"instance","scalar":"f32","lanes":3,"stride":null,"length":null}"#,
+            )
+            .unwrap()
+        };
+        assert_eq!(
+            registry.allocate(conflict, capacities(8)),
+            Err(SharedSoaError::LayoutConflict)
+        );
+    }
+
+    #[test]
+    fn domain_growth_replaces_layout_and_retains_old_storage() {
+        let mut registry = SharedSoaRegistry::new(capacities(8)).unwrap();
+        let old = registry.arrays["instance.transform"].storage.as_ptr();
+        assert!(registry.sync_capacities(capacities(16)).unwrap());
+        let array = &registry.arrays["instance.transform"];
+        assert_ne!(old, array.storage.as_ptr());
+        assert_eq!(array.retired.len(), 1);
+        assert_eq!(array.descriptor().unwrap().capacity, 16);
+    }
+
+    #[test]
+    fn fixed_arrays_grow_and_publish_stable_byte_uploads() {
+        let mut registry = SharedSoaRegistry::new(capacities(8)).unwrap();
+        let request = |length| ArrayRequest {
+            name: "upload.gltf".into(),
+            domain: ArrayDomain::Fixed,
+            scalar: ScalarType::U32,
+            lanes: 4,
+            stride: Some(16),
+            length: Some(length),
+        };
+        let first = registry.allocate(request(1), capacities(8)).unwrap();
+        let second = registry.allocate(request(2), capacities(8)).unwrap();
+        assert_eq!(first.id, second.id);
+        assert_eq!((second.length, second.capacity), (2, 2));
+
+        let array = registry.arrays.get_mut("upload.gltf").unwrap();
+        let sequence = array.try_lock().unwrap();
+        array
+            .data_word(0, 0)
+            .store(u32::from_le_bytes(*b"glTF"), Ordering::Relaxed);
+        array
+            .data_word(0, 1)
+            .store(u32::from_le_bytes([2, 0, 0, 0]), Ordering::Relaxed);
+        array.unlock(sequence);
+        assert_eq!(
+            registry.read_fixed_bytes(first.id, 8).unwrap(),
+            b"glTF\x02\0\0\0"
+        );
+    }
+
+    #[test]
+    fn guarded_columns_ignore_stale_slot_writers() {
+        let mut registry = SharedSoaRegistry::new(capacities(8)).unwrap();
+        let handle = InstanceHandle::from_parts(2, 7);
+        let array = registry.arrays.get_mut("instance.transform").unwrap();
+        let descriptor = array.descriptor().unwrap();
+        assert_eq!(
+            (descriptor.stride, descriptor.generation_guard),
+            (80, Some(ArrayDomain::Instance))
+        );
+
+        array
+            .data_word(2, 0)
+            .store(1.0f32.to_bits(), Ordering::Relaxed);
+        array.data_word(2, 16).store(6, Ordering::Relaxed);
+        array.data_word(2, 17).store(1, Ordering::Relaxed);
+        array.word(9).store(2, Ordering::Release);
+        assert!(registry
+            .take_instance_words("instance.transform", &[handle])
+            .unwrap()
+            .is_empty());
+
+        let array = registry.arrays.get_mut("instance.transform").unwrap();
+        array
+            .data_word(2, 0)
+            .store(2.0f32.to_bits(), Ordering::Relaxed);
+        array.data_word(2, 16).store(7, Ordering::Relaxed);
+        array.data_word(2, 17).store(2, Ordering::Relaxed);
+        array.word(9).store(6, Ordering::Release);
+        let values = registry
+            .take_instance_words("instance.transform", &[handle])
+            .unwrap();
+        assert_eq!(values.len(), 1);
+        assert_eq!(f32::from_bits(values[0].1[0]), 2.0);
+    }
+}