feat: replace mesh queries with typed predicates

Amp-Thread-ID: https://ampcode.com/threads/T-019f9d91-77c1-7206-a60f-ed6554ce92ab

Co-authored-by: Heaust Azure <heaust.azure@gmail.com>

renderer/src/command_ring.rs

@@ -2,12 +2,12 @@
 use std::sync::atomic::{AtomicU32, Ordering};
 
 pub const MAGIC: u32 = u32::from_le_bytes(*b"YAWN");
-pub const VERSION: u32 = 1;
+pub const VERSION: u32 = 2;
 pub const CAPACITY: usize = 1024;
-pub const SLOT_WORDS: usize = 24;
-pub const SLOT_BYTES: usize = 96;
+pub const SLOT_WORDS: usize = 40;
+pub const SLOT_BYTES: usize = 160;
 pub const HEADER_BYTES: usize = 64;
-pub const SLOT_VERSION: u32 = 1;
+pub const SLOT_VERSION: u32 = 2;
 const STATE_OPEN: u32 = 0;
 const STATE_CORRUPT: u32 = 1;
 
@@ -121,7 +121,7 @@ mod tests {
     #[test]
     fn malformed_slot_fails_closed() {
         for (version, request, expected) in [
-            (2, 1, RingError::SlotVersion),
+            (SLOT_VERSION + 1, 1, RingError::SlotVersion),
             (SLOT_VERSION, 0, RingError::ZeroRequest),
         ] {
             let ring = CommandRing::new();

renderer/src/gltf.rs

@@ -4,8 +4,8 @@ use gltf::Gltf;
 use ultraviolet::{Mat4, Vec3};
 
 use crate::render_data::{
-    InstanceHandle, MaterialKey, MeshCreateInfo, MeshHandle, ModelTransform, PipelineKey,
-    RenderData, RenderDataError, RenderFlags,
+    InstanceHandle, InstanceType, MaterialKey, MeshCreateInfo, MeshHandle, ModelTransform,
+    PipelineKey, RenderData, RenderDataError,
 };
 
 #[derive(Clone, Copy, Debug, Default, PartialEq, Eq)]
@@ -549,8 +549,26 @@ pub fn install_imported(
             indices: &geometry.indices,
             pipeline: pipelines[usize::from(geometry.double_sided)],
             material: geometry.material,
-            flags: RenderFlags::VISIBLE,
-            default_instance_flags: RenderFlags::VISIBLE,
+            default_instance_type: InstanceType {
+                words: [
+                    1 | 4 | (geometry.double_sided as u32) * 8,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                    0,
+                ],
+            },
             default_transform: transform,
         })?;
         handles.insert(geometry.key, created.mesh);
@@ -570,10 +588,11 @@ pub fn install_imported(
             .get(&occurrence.key)
             .ok_or_else(|| ImportError::InvalidPrimitive("occurrence has no geometry".into()))?;
         if consumed.insert(occurrence.key, ()).is_some() {
+            let instance_type = stage.mesh(mesh).unwrap().default_instance_type;
             instance_handles.push(stage.create_instance(
                 mesh,
                 occurrence.transform,
-                RenderFlags::VISIBLE,
+                instance_type,
             )?);
         }
         let geometry = geometries
@@ -584,7 +603,7 @@ pub fn install_imported(
             let point = transform.transform_point3(Vec3::from(*position));
             [point.x, point.y, point.z]
         }));
-        let local = stage.mesh(mesh).unwrap().aabb;
+        let local = stage.mesh(mesh).unwrap().local_aabb;
         for x in [local.min[0], local.max[0]] {
             for y in [local.min[1], local.max[1]] {
                 for z in [local.min[2], local.max[2]] {

renderer/src/render_data/mod.rs

@@ -4,7 +4,7 @@ mod range_allocator;
 pub use handle::{InstanceHandle, MeshHandle};
 
 use std::{
-    ops::{BitAnd, BitAndAssign, BitOr, BitOrAssign, Deref},
+    ops::Deref,
     sync::atomic::{AtomicU64, Ordering},
 };
 
@@ -58,50 +58,36 @@ impl MaterialKey {
     }
 }
 
-#[repr(transparent)]
-#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)]
-pub struct RenderFlags(u32);
+#[repr(C)]
+#[derive(Clone, Copy, Debug, PartialEq, Eq, bytemuck::Pod, bytemuck::Zeroable)]
+pub struct InstanceType {
+    pub words: [u32; 16],
+}
 
-impl RenderFlags {
-    pub const NONE: Self = Self(0);
-    pub const VISIBLE: Self = Self(1);
+impl InstanceType {
+    pub const VISIBLE_MASK: u32 = 1;
+    pub const ZERO: Self = Self { words: [0; 16] };
+    pub const VISIBLE: Self = Self {
+        words: [1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],
+    };
 
-    pub const fn from_bits_retain(bits: u32) -> Self {
-        Self(bits)
+    pub const fn is_visible(self) -> bool {
+        self.words[0] & Self::VISIBLE_MASK != 0
     }
-
-    pub const fn bits(self) -> u32 {
-        self.0
-    }
-
-    pub const fn contains(self, other: Self) -> bool {
-        self.0 & other.0 == other.0
+    pub fn set_visible(&mut self, visible: bool) {
+        self.words[0] = (self.words[0] & !Self::VISIBLE_MASK) | visible as u32;
     }
 }
 
-impl BitOr for RenderFlags {
-    type Output = Self;
-    fn bitor(self, rhs: Self) -> Self {
-        Self(self.0 | rhs.0)
-    }
-}
-impl BitOrAssign for RenderFlags {
-    fn bitor_assign(&mut self, rhs: Self) {
-        self.0 |= rhs.0;
-    }
-}
-impl BitAnd for RenderFlags {
-    type Output = Self;
-    fn bitand(self, rhs: Self) -> Self {
-        Self(self.0 & rhs.0)
-    }
-}
-impl BitAndAssign for RenderFlags {
-    fn bitand_assign(&mut self, rhs: Self) {
-        self.0 &= rhs.0;
+impl Default for InstanceType {
+    fn default() -> Self {
+        Self::VISIBLE
     }
 }
 
+const _: [(); 64] = [(); std::mem::size_of::<InstanceType>()];
+const _: [(); 4] = [(); std::mem::align_of::<InstanceType>()];
+
 #[derive(Clone, Copy, Debug, PartialEq)]
 pub struct Aabb {
     pub min: [f32; 3],
@@ -124,8 +110,7 @@ pub struct MeshCreateInfo<'a> {
     pub indices: &'a [u32],
     pub pipeline: PipelineKey,
     pub material: MaterialKey,
-    pub flags: RenderFlags,
-    pub default_instance_flags: RenderFlags,
+    pub default_instance_type: InstanceType,
     pub default_transform: ModelTransform,
 }
 
@@ -141,8 +126,8 @@ pub struct MeshView {
     pub geometry: GeometryRange,
     pub pipeline: PipelineKey,
     pub material: MaterialKey,
-    pub flags: RenderFlags,
-    pub aabb: Aabb,
+    pub default_instance_type: InstanceType,
+    pub local_aabb: Aabb,
     pub default_instance: InstanceHandle,
 }
 
@@ -152,7 +137,7 @@ pub struct InstanceView {
     pub mesh: MeshHandle,
     pub model: ModelTransform,
     pub normal: NormalMatrix,
-    pub flags: RenderFlags,
+    pub instance_type: InstanceType,
     pub is_default: bool,
 }
 
@@ -311,7 +296,7 @@ struct MeshSoa {
     index_counts: Vec<u32>,
     pipeline_keys: Vec<PipelineKey>,
     material_keys: Vec<MaterialKey>,
-    flags: Vec<RenderFlags>,
+    default_instance_types: Vec<InstanceType>,
     aabb_mins: Vec<[f32; 3]>,
     aabb_maxs: Vec<[f32; 3]>,
     default_instance_slots: Vec<u32>,
@@ -329,7 +314,7 @@ struct InstanceSoa {
     normal_col_0: Vec<[f32; 3]>,
     normal_col_1: Vec<[f32; 3]>,
     normal_col_2: Vec<[f32; 3]>,
-    flags: Vec<RenderFlags>,
+    instance_types: Vec<InstanceType>,
 }
 
 pub struct RenderData {
@@ -369,9 +354,9 @@ impl ReplacementStage {
         &mut self,
         mesh: MeshHandle,
         model: ModelTransform,
-        flags: RenderFlags,
+        instance_type: InstanceType,
     ) -> Result<InstanceHandle, RenderDataError> {
-        self.data.create_instance(mesh, model, flags)
+        self.data.create_instance(mesh, model, instance_type)
     }
 }
 
@@ -595,7 +580,7 @@ impl RenderData {
             },
             info.pipeline,
             info.material,
-            info.flags,
+            info.default_instance_type,
             bounds,
             default_instance,
         );
@@ -604,7 +589,7 @@ impl RenderData {
             mesh,
             info.default_transform,
             normal,
-            info.default_instance_flags,
+            info.default_instance_type,
         );
         self.revision = next_revision;
         Ok(CreatedMesh {
@@ -617,19 +602,20 @@ impl RenderData {
         &mut self,
         mesh: MeshHandle,
         model: ModelTransform,
-        flags: RenderFlags,
+        instance_type: InstanceType,
     ) -> Result<InstanceHandle, RenderDataError> {
         if !self.meshes.slots.contains(mesh.slot(), mesh.generation()) {
             return Err(RenderDataError::InvalidMeshHandle);
         }
         let normal = normal_matrix(model)?;
-        affine_world_aabb(self.mesh(mesh).unwrap().aabb, model)?;
+        affine_world_aabb(self.mesh(mesh).unwrap().local_aabb, model)?;
         let next_revision = self.next_revision()?;
         let required = self.instances.slots.required_len_for_prepare()?;
         self.instances.reserve(required)?;
         let prepared = self.instances.slots.prepare()?;
         let handle = InstanceHandle::from_parts(prepared.slot, prepared.generation);
-        self.instances.commit(prepared, mesh, model, normal, flags);
+        self.instances
+            .commit(prepared, mesh, model, normal, instance_type);
         self.revision = next_revision;
         Ok(handle)
     }
@@ -699,10 +685,10 @@ impl RenderData {
         Ok(())
     }
 
-    pub fn set_mesh_flags(
+    pub fn set_mesh_visible(
         &mut self,
         handle: MeshHandle,
-        flags: RenderFlags,
+        visible: bool,
     ) -> Result<(), RenderDataError> {
         if !self
             .meshes
@@ -711,16 +697,24 @@ impl RenderData {
         {
             return Err(RenderDataError::InvalidMeshHandle);
         }
+        let owned: Vec<u32> = self
+            .instances
+            .slots
+            .occupied()
+            .filter_map(|(slot, _)| (self.instances.mesh_handle(slot) == handle).then_some(slot))
+            .collect();
         let next_revision = self.next_revision()?;
-        self.meshes.flags[handle.slot() as usize] = flags;
+        for slot in owned {
+            self.instances.instance_types[slot as usize].set_visible(visible);
+        }
         self.revision = next_revision;
         Ok(())
     }
 
-    pub fn set_instance_flags(
+    pub fn set_instance_type(
         &mut self,
         handle: InstanceHandle,
-        flags: RenderFlags,
+        instance_type: InstanceType,
     ) -> Result<(), RenderDataError> {
         if !self
             .instances
@@ -730,7 +724,25 @@ impl RenderData {
             return Err(RenderDataError::InvalidInstanceHandle);
         }
         let next_revision = self.next_revision()?;
-        self.instances.flags[handle.slot() as usize] = flags;
+        self.instances.instance_types[handle.slot() as usize] = instance_type;
+        self.revision = next_revision;
+        Ok(())
+    }
+
+    pub fn set_instance_visible(
+        &mut self,
+        handle: InstanceHandle,
+        visible: bool,
+    ) -> Result<(), RenderDataError> {
+        if !self
+            .instances
+            .slots
+            .contains(handle.slot(), handle.generation())
+        {
+            return Err(RenderDataError::InvalidInstanceHandle);
+        }
+        let next_revision = self.next_revision()?;
+        self.instances.instance_types[handle.slot() as usize].set_visible(visible);
         self.revision = next_revision;
         Ok(())
     }
@@ -749,7 +761,7 @@ impl RenderData {
         }
         let normal = normal_matrix(model)?;
         let mesh = self.instances.mesh_handle(handle.slot());
-        affine_world_aabb(self.mesh(mesh).unwrap().aabb, model)?;
+        affine_world_aabb(self.mesh(mesh).unwrap().local_aabb, model)?;
         let next_revision = self.next_revision()?;
         self.instances.set_transform(handle.slot(), model, normal);
         self.revision = next_revision;
@@ -890,7 +902,7 @@ impl MeshSoa {
             index_counts: Vec::new(),
             pipeline_keys: Vec::new(),
             material_keys: Vec::new(),
-            flags: Vec::new(),
+            default_instance_types: Vec::new(),
             aabb_mins: Vec::new(),
             aabb_maxs: Vec::new(),
             default_instance_slots: Vec::new(),
@@ -910,7 +922,7 @@ impl MeshSoa {
         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.flags, target, "meshes")?;
+        reserve_vec(&mut self.default_instance_types, target, "meshes")?;
         reserve_vec(&mut self.aabb_mins, target, "meshes")?;
         reserve_vec(&mut self.aabb_maxs, target, "meshes")?;
         reserve_vec(&mut self.default_instance_slots, target, "meshes")?;
@@ -925,7 +937,7 @@ impl MeshSoa {
         geometry: GeometryRange,
         pipeline: PipelineKey,
         material: MaterialKey,
-        flags: RenderFlags,
+        default_instance_type: InstanceType,
         bounds: Aabb,
         default: InstanceHandle,
     ) {
@@ -936,7 +948,7 @@ impl MeshSoa {
         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.flags, len, RenderFlags::NONE);
+        resize_column(&mut self.default_instance_types, len, InstanceType::ZERO);
         resize_column(&mut self.aabb_mins, len, [0.0; 3]);
         resize_column(&mut self.aabb_maxs, len, [0.0; 3]);
         resize_column(&mut self.default_instance_slots, len, 0);
@@ -948,7 +960,7 @@ impl MeshSoa {
         self.index_counts[index] = geometry.index_count;
         self.pipeline_keys[index] = pipeline;
         self.material_keys[index] = material;
-        self.flags[index] = flags;
+        self.default_instance_types[index] = default_instance_type;
         self.aabb_mins[index] = bounds.min;
         self.aabb_maxs[index] = bounds.max;
         self.default_instance_slots[index] = default.slot();
@@ -968,8 +980,8 @@ impl MeshSoa {
             },
             pipeline: self.pipeline_keys[index],
             material: self.material_keys[index],
-            flags: self.flags[index],
-            aabb: Aabb {
+            default_instance_type: self.default_instance_types[index],
+            local_aabb: Aabb {
                 min: self.aabb_mins[index],
                 max: self.aabb_maxs[index],
             },
@@ -994,7 +1006,7 @@ impl InstanceSoa {
             normal_col_0: Vec::new(),
             normal_col_1: Vec::new(),
             normal_col_2: Vec::new(),
-            flags: Vec::new(),
+            instance_types: Vec::new(),
         };
         soa.reserve(initial)?;
         Ok(soa)
@@ -1013,7 +1025,7 @@ impl InstanceSoa {
         reserve_vec(&mut self.normal_col_0, target, "instances")?;
         reserve_vec(&mut self.normal_col_1, target, "instances")?;
         reserve_vec(&mut self.normal_col_2, target, "instances")?;
-        reserve_vec(&mut self.flags, target, "instances")?;
+        reserve_vec(&mut self.instance_types, target, "instances")?;
         self.slots.reserve_for_len(target, "instances")?;
         Ok(())
     }
@@ -1024,7 +1036,7 @@ impl InstanceSoa {
         mesh: MeshHandle,
         model: ModelTransform,
         normal: NormalMatrix,
-        flags: RenderFlags,
+        instance_type: InstanceType,
     ) {
         let len = prepared.slot as usize + 1;
         resize_column(&mut self.mesh_slots, len, 0);
@@ -1036,11 +1048,11 @@ impl InstanceSoa {
         resize_column(&mut self.normal_col_0, len, [0.0; 3]);
         resize_column(&mut self.normal_col_1, len, [0.0; 3]);
         resize_column(&mut self.normal_col_2, len, [0.0; 3]);
-        resize_column(&mut self.flags, len, RenderFlags::NONE);
+        resize_column(&mut self.instance_types, len, InstanceType::ZERO);
         let index = prepared.slot as usize;
         self.mesh_slots[index] = mesh.slot();
         self.mesh_generations[index] = mesh.generation();
-        self.flags[index] = flags;
+        self.instance_types[index] = instance_type;
         self.set_transform(prepared.slot, model, normal);
         self.slots.commit(prepared);
     }
@@ -1077,7 +1089,7 @@ impl InstanceSoa {
                 self.normal_col_1[index],
                 self.normal_col_2[index],
             ],
-            flags: self.flags[index],
+            instance_type: self.instance_types[index],
             is_default,
         }
     }

renderer/src/render_data/tests.rs

@@ -16,8 +16,9 @@ fn info() -> MeshCreateInfo<'static> {
         indices: &INDICES,
         pipeline: PipelineKey::new(7),
         material: MaterialKey::new(11),
-        flags: RenderFlags::from_bits_retain(2),
-        default_instance_flags: RenderFlags::VISIBLE,
+        default_instance_type: InstanceType {
+            words: [3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0],
+        },
         default_transform: IDENTITY_MODEL_TRANSFORM,
     }
 }
@@ -81,29 +82,42 @@ fn world_bounds_reject_projective_and_overflowing_transforms() {
 }
 
 #[test]
-fn default_instance_is_protected_and_flags_are_separate() {
+fn default_instance_is_protected_and_preserves_its_type() {
     let mut data = data();
     let created = data.create_mesh(info()).unwrap();
     assert!(data.instance(created.default_instance).unwrap().is_default);
-    assert_eq!(data.mesh(created.mesh).unwrap().flags.bits(), 2);
+    assert_eq!(
+        data.mesh(created.mesh).unwrap().default_instance_type.words[0],
+        3
+    );
     assert_eq!(
         data.mesh(created.mesh).unwrap().material,
         MaterialKey::new(11)
     );
     assert_eq!(
-        data.instance(created.default_instance).unwrap().flags,
-        RenderFlags::VISIBLE
+        data.instance(created.default_instance)
+            .unwrap()
+            .instance_type,
+        info().default_instance_type
     );
     assert_eq!(
         data.destroy_instance(created.default_instance),
         Err(RenderDataError::CannotDestroyDefaultInstance)
     );
-    data.set_mesh_flags(created.mesh, RenderFlags::NONE)
-        .unwrap();
-    assert_eq!(data.mesh(created.mesh).unwrap().flags, RenderFlags::NONE);
+    data.set_mesh_visible(created.mesh, false).unwrap();
     assert_eq!(
-        data.instance(created.default_instance).unwrap().flags,
-        RenderFlags::VISIBLE
+        data.mesh(created.mesh).unwrap().default_instance_type,
+        info().default_instance_type
+    );
+    assert_eq!(
+        data.instance(created.default_instance)
+            .unwrap()
+            .instance_type,
+        {
+            let mut expected = info().default_instance_type;
+            expected.set_visible(false);
+            expected
+        }
     );
 }
 
@@ -112,11 +126,11 @@ fn stale_mesh_and_instance_handles_are_rejected_after_reuse() {
     let mut data = data();
     let first = data.create_mesh(info()).unwrap();
     let old_instance = data
-        .create_instance(first.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+        .create_instance(first.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     data.destroy_instance(old_instance).unwrap();
     let replacement = data
-        .create_instance(first.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+        .create_instance(first.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     assert_eq!(old_instance.slot(), replacement.slot());
     assert_ne!(old_instance.generation(), replacement.generation());
@@ -133,7 +147,7 @@ fn clear_handles_all_slot_states_retains_capacity_and_never_reuses_retired() {
     let mut data = data();
     let mesh = data.create_mesh(info()).unwrap();
     let vacant = data
-        .create_instance(mesh.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+        .create_instance(mesh.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     data.destroy_instance(vacant).unwrap();
     data.instances
@@ -188,7 +202,7 @@ fn all_storage_classes_grow_and_retired_slots_force_max_checked_append() {
             instances: 1,
         }
     );
-    data.create_instance(mesh.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+    data.create_instance(mesh.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     assert_eq!(data.capacities().instances, 2);
 
@@ -302,7 +316,7 @@ fn aabb_supports_one_point_and_multiple_points() {
     let mut data = data();
     let mesh = data.create_mesh(one).unwrap();
     assert_eq!(
-        data.mesh(mesh.mesh).unwrap().aabb,
+        data.mesh(mesh.mesh).unwrap().local_aabb,
         Aabb {
             min: point[0],
             max: point[0]
@@ -310,7 +324,7 @@ fn aabb_supports_one_point_and_multiple_points() {
     );
     let mesh = data.create_mesh(info()).unwrap();
     assert_eq!(
-        data.mesh(mesh.mesh).unwrap().aabb,
+        data.mesh(mesh.mesh).unwrap().local_aabb,
         Aabb {
             min: [-1.0, -2.0, -3.0],
             max: [4.0, 2.0, 3.0],
@@ -444,7 +458,7 @@ fn normal_matrices_and_failed_transform_operations_are_transactional() {
         assert_eq!(data.instance(mesh.default_instance).unwrap(), old);
         let count = data.instance_count();
         assert_eq!(
-            data.create_instance(mesh.mesh, invalid, RenderFlags::NONE),
+            data.create_instance(mesh.mesh, invalid, InstanceType::ZERO),
             Err(RenderDataError::InvalidTransform)
         );
         assert_eq!(data.instance_count(), count);
@@ -463,14 +477,14 @@ fn destroying_mesh_invalidates_exact_owner_instances_with_reused_generations() {
     let first = data.create_mesh(info()).unwrap();
     let second = data.create_mesh(info()).unwrap();
     let first_extra = data
-        .create_instance(first.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+        .create_instance(first.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     let second_extra = data
-        .create_instance(second.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+        .create_instance(second.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     data.destroy_instance(first_extra).unwrap();
     let reused = data
-        .create_instance(second.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::NONE)
+        .create_instance(second.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::ZERO)
         .unwrap();
     assert_eq!(first_extra.slot(), reused.slot());
     data.destroy_mesh(first.mesh).unwrap();

renderer/src/render_graph/compiler.rs

@@ -21,20 +21,13 @@ struct CullParameters {
 }
 #[derive(Deserialize)]
 #[serde(deny_unknown_fields, rename_all = "camelCase")]
-struct QueryParameters {
-    visible_predicate: TriStatePredicate,
-    visible_default: bool,
-    frustum_culled_predicate: TriStatePredicate,
-    frustum_culled_default: bool,
-}
-#[derive(Deserialize)]
-#[serde(deny_unknown_fields, rename_all = "camelCase")]
 struct PipelineParameters {
     pipeline: String,
     depth_compare: CompareFunction,
     depth_write_enabled: bool,
     clear_depth: f32,
     clear_color: [f64; 4],
+    predicate_default: bool,
 }
 #[derive(Deserialize)]
 #[serde(deny_unknown_fields)]
@@ -146,7 +139,6 @@ struct OutputKey(usize, u16);
 #[derive(Clone, Copy)]
 struct BoundInput {
     producer: OutputKey,
-    active: bool,
 }
 #[derive(Clone)]
 struct DependencyEdge {
@@ -232,15 +224,6 @@ fn validate_name_grammar(s: &str, path: impl Into<String>) -> Result<(), GraphEr
     }
 }
 
-pub fn mesh_predicate_matches(predicate: RuntimePredicate, flag: bool) -> bool {
-    match predicate {
-        RuntimePredicate::Any => true,
-        RuntimePredicate::RequiredTrue => flag,
-        RuntimePredicate::RequiredFalse => !flag,
-        RuntimePredicate::Never => false,
-    }
-}
-
 pub fn parse_and_compile(bytes: &[u8]) -> Result<CompiledGraph, GraphError> {
     if bytes.len() > MAX_JSON_BYTES {
         return Err(GraphError::new(
@@ -424,6 +407,89 @@ fn decode(node: &Node, i: usize) -> Result<NormalizedParameters, GraphError> {
             $variant
         }};
     }
+    fn literal(value: &serde_json::Value, ty: SemanticType) -> Option<TypedLiteral> {
+        let floats = |value: &serde_json::Value, n: usize| -> Option<Vec<f32>> {
+            let values = value.as_array()?;
+            if values.len() != n {
+                return None;
+            }
+            values
+                .iter()
+                .map(|value| value.as_f64().map(|value| value as f32))
+                .collect::<Option<Vec<_>>>()
+                .filter(|values| values.iter().all(|value| value.is_finite()))
+        };
+        let vector = |n| floats(value, n);
+        Some(match ty {
+            SemanticType::Bool => TypedLiteral::Bool(value.as_bool()?),
+            SemanticType::F32 => {
+                let value = value.as_f64()? as f32;
+                if !value.is_finite() {
+                    return None;
+                }
+                TypedLiteral::F32(value)
+            }
+            SemanticType::U32 => TypedLiteral::U32(value.as_u64()?.try_into().ok()?),
+            SemanticType::Vec2 => TypedLiteral::Vec2(vector(2)?.try_into().ok()?),
+            SemanticType::Vec3 => TypedLiteral::Vec3(vector(3)?.try_into().ok()?),
+            SemanticType::Vec4 => TypedLiteral::Vec4(vector(4)?.try_into().ok()?),
+            SemanticType::U32x16 => TypedLiteral::U32x16(
+                value
+                    .as_array()?
+                    .iter()
+                    .map(|v| v.as_u64()?.try_into().ok())
+                    .collect::<Option<Vec<u32>>>()?
+                    .try_into()
+                    .ok()?,
+            ),
+            SemanticType::LocalAabb => TypedLiteral::LocalAabb {
+                min: floats(value.get("min")?, 3)?.try_into().ok()?,
+                max: floats(value.get("max")?, 3)?.try_into().ok()?,
+            },
+            ty @ (SemanticType::Mat2 | SemanticType::Mat3 | SemanticType::Mat4) => {
+                let n = match ty {
+                    SemanticType::Mat2 => 2,
+                    SemanticType::Mat3 => 3,
+                    _ => 4,
+                };
+                let columns = value.as_array()?;
+                if columns.len() != n {
+                    return None;
+                }
+                let columns = columns
+                    .iter()
+                    .map(|v| floats(v, n))
+                    .collect::<Option<Vec<_>>>()?;
+                match ty {
+                    SemanticType::Mat2 => TypedLiteral::Mat2(
+                        columns
+                            .into_iter()
+                            .map(|v| v.try_into().ok())
+                            .collect::<Option<Vec<_>>>()?
+                            .try_into()
+                            .ok()?,
+                    ),
+                    SemanticType::Mat3 => TypedLiteral::Mat3(
+                        columns
+                            .into_iter()
+                            .map(|v| v.try_into().ok())
+                            .collect::<Option<Vec<_>>>()?
+                            .try_into()
+                            .ok()?,
+                    ),
+                    _ => TypedLiteral::Mat4(
+                        columns
+                            .into_iter()
+                            .map(|v| v.try_into().ok())
+                            .collect::<Option<Vec<_>>>()?
+                            .try_into()
+                            .ok()?,
+                    ),
+                }
+            }
+            SemanticType::MeshData | SemanticType::Texture => return None,
+        })
+    }
     Ok(match node.executor.key.as_str() {
         "mesh" => empty!(NormalizedParameters::Mesh),
         "frustum_cull" => {
@@ -607,37 +673,6 @@ fn decode(node: &Node, i: usize) -> Result<NormalizedParameters, GraphError> {
                 descriptor: normalize_texture(p.texture, &base)?,
             }
         }
-        "mesh_query" => {
-            let p: QueryParameters =
-                serde_json::from_value(node.parameters.clone()).map_err(invalid)?;
-            let fold = |predicate, default, linked| match (predicate, linked, default) {
-                (TriStatePredicate::Any, _, _) => RuntimePredicate::Any,
-                (TriStatePredicate::RequiredTrue, true, _) => RuntimePredicate::RequiredTrue,
-                (TriStatePredicate::RequiredFalse, true, _) => RuntimePredicate::RequiredFalse,
-                (TriStatePredicate::RequiredTrue, false, true)
-                | (TriStatePredicate::RequiredFalse, false, false) => RuntimePredicate::Any,
-                _ => RuntimePredicate::Never,
-            };
-            let mut visible = fold(
-                p.visible_predicate,
-                p.visible_default,
-                node.inputs.contains_key("isVisible"),
-            );
-            let mut culled = fold(
-                p.frustum_culled_predicate,
-                p.frustum_culled_default,
-                node.inputs.contains_key("isFrustumCulled"),
-            );
-            if visible == RuntimePredicate::Never || culled == RuntimePredicate::Never {
-                visible = RuntimePredicate::Never;
-                culled = RuntimePredicate::Never;
-            }
-            NormalizedParameters::MeshQuery {
-                visible_predicate: visible,
-                frustum_culled_predicate: culled,
-            }
-        }
-        "pipeline_registry" => empty!(NormalizedParameters::PipelineRegistry),
         "pipeline" => {
             let p: PipelineParameters =
                 serde_json::from_value(node.parameters.clone()).map_err(invalid)?;
@@ -673,8 +708,50 @@ fn decode(node: &Node, i: usize) -> Result<NormalizedParameters, GraphError> {
                 depth_write_enabled: p.depth_write_enabled,
                 clear_depth: p.clear_depth,
                 clear_color: p.clear_color,
+                predicate_default: p.predicate_default,
             }
         }
+        key if contract(key)
+            .is_some_and(|contract| contract.execution == ExecutionClass::Expression) =>
+        {
+            let contract = contract(key).unwrap();
+            let object = node.parameters.as_object().ok_or_else(|| {
+                error(
+                    "GRAPH_PARAMETERS_INVALID",
+                    "parameters must be an object",
+                    base.clone(),
+                )
+            })?;
+            if object.len() != contract.inputs.len() {
+                return Err(error(
+                    "GRAPH_PARAMETERS_INVALID",
+                    "expression defaults must exactly match inputs",
+                    base,
+                ));
+            }
+            let mut defaults = Vec::with_capacity(contract.inputs.len());
+            for input in contract.inputs {
+                let key = format!("{}Default", input.name);
+                let value = object.get(&key).ok_or_else(|| {
+                    error(
+                        "GRAPH_PARAMETERS_INVALID",
+                        "missing expression default",
+                        format!("{base}.{key}"),
+                    )
+                })?;
+                let TypeConstraint::Exact(ty) = input.accepted else {
+                    unreachable!()
+                };
+                defaults.push(literal(value, ty).ok_or_else(|| {
+                    error(
+                        "GRAPH_PARAMETERS_INVALID",
+                        "invalid typed expression default",
+                        format!("{base}.{key}"),
+                    )
+                })?);
+            }
+            NormalizedParameters::ExpressionDefaults { defaults }
+        }
         _ => unreachable!(),
     })
 }
@@ -846,13 +923,8 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
     }
     for (i, n) in graph.nodes.iter().enumerate() {
         for input in contracts[i].inputs {
-            let inactive = matches!(&params[i], NormalizedParameters::MeshQuery { visible_predicate, frustum_culled_predicate } if (input.name == "isVisible" && matches!(visible_predicate, RuntimePredicate::Any | RuntimePredicate::Never)) || (input.name == "isFrustumCulled" && matches!(frustum_culled_predicate, RuntimePredicate::Any | RuntimePredicate::Never)));
             if !n.inputs.contains_key(input.name) {
-                if input.cardinality == InputCardinality::RequiredOne
-                    || (!inactive
-                        && matches!(params[i], NormalizedParameters::MeshQuery { .. })
-                        && input.name != "mesh")
-                {
+                if input.cardinality == InputCardinality::RequiredOne {
                     return Err(error(
                         "GRAPH_SOCKET_CARDINALITY",
                         "required input is missing",
@@ -866,7 +938,6 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
     let mut bound: Vec<BTreeMap<&str, BoundInput>> = vec![BTreeMap::new(); graph.nodes.len()];
     for (i, n) in graph.nodes.iter().enumerate() {
         for input in contracts[i].inputs {
-            let inactive = matches!(&params[i], NormalizedParameters::MeshQuery { visible_predicate, frustum_culled_predicate } if (input.name == "isVisible" && matches!(visible_predicate, RuntimePredicate::Any | RuntimePredicate::Never)) || (input.name == "isFrustumCulled" && matches!(frustum_culled_predicate, RuntimePredicate::Any | RuntimePredicate::Never)));
             let Some(r) = n.inputs.get(input.name) else {
                 continue;
             };
@@ -884,97 +955,18 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
                     format!("nodes[{i}].inputs.{}", input.name),
                 ));
             }
-            if let Some(flag) = MeshFlag::ORDERED
-                .iter()
-                .find(|f| f.input_socket() == input.name)
-            {
-                if out.metadata != (OutputMetadata::BooleanFlag { flag: *flag }) {
-                    return Err(error(
-                        "GRAPH_SOCKET_TYPE_MISMATCH",
-                        "mesh flag metadata mismatch",
-                        format!("nodes[{i}].inputs.{}", input.name),
-                    ));
-                }
-            }
             bound[i].insert(
                 input.name,
                 BoundInput {
                     producer: OutputKey(pn, ordinal as u16),
-                    active: !inactive,
                 },
             );
         }
     }
-    let root = |key: OutputKey,
-                bound: &Vec<BTreeMap<&str, BoundInput>>,
-                contracts: &Vec<&Contract>|
-     -> Option<OutputKey> {
-        let mut k = key;
-        let mut seen = HashSet::new();
-        loop {
-            if !seen.insert(k.0) {
-                return None;
-            }
-            if contracts[k.0].key == "mesh" {
-                let ordinal = contracts[k.0]
-                    .outputs
-                    .iter()
-                    .position(|output| output.semantic_type == SemanticType::MeshData)?;
-                return Some(OutputKey(k.0, ordinal as u16));
-            }
-            if contracts[k.0].outputs[k.1 as usize].semantic_type == SemanticType::MeshData {
-                return Some(k);
-            }
-            k = bound[k.0]
-                .get("mesh")
-                .or_else(|| bound[k.0].get("pipelineIndices"))
-                .or_else(|| bound[k.0].get("activation"))?
-                .producer;
-        }
-    };
-    for (i, c) in contracts.iter().enumerate() {
-        if c.key == "frustum_cull"
-            && root(bound[i]["mesh"].producer, &bound, &contracts)
-                != root(bound[i]["localAabbs"].producer, &bound, &contracts)
-        {
-            return Err(error(
-                "GRAPH_SOCKET_TYPE_MISMATCH",
-                "scene roots differ",
-                format!("nodes[{i}].inputs.localAabbs"),
-            ));
-        }
-        if matches!(c.key, "mesh_query" | "pipeline_registry" | "pipeline") {
-            let scene_socket = if c.key == "pipeline_registry" {
-                "pipelineIndices"
-            } else {
-                "mesh"
-            };
-            let scene = root(bound[i][scene_socket].producer, &bound, &contracts);
-            for (s, b) in &bound[i] {
-                if b.active
-                    && matches!(
-                        *s,
-                        "isVisible"
-                            | "isFrustumCulled"
-                            | "draws"
-                            | "pipelineIndices"
-                            | "activation"
-                    )
-                    && root(b.producer, &bound, &contracts) != scene
-                {
-                    return Err(error(
-                        "GRAPH_SOCKET_TYPE_MISMATCH",
-                        "scene roots differ",
-                        format!("nodes[{i}].inputs.{s}"),
-                    ));
-                }
-            }
-        }
-    }
     let mut edges = Vec::new();
     for i in 0..graph.nodes.len() {
         for (input_ordinal, input) in contracts[i].inputs.iter().enumerate() {
-            if let Some(b) = bound[i].get(input.name).filter(|b| b.active) {
+            if let Some(b) = bound[i].get(input.name) {
                 edges.push(DependencyEdge {
                     from_node: b.producer.0,
                     from_socket: contracts[b.producer.0].outputs[b.producer.1 as usize]
@@ -1030,6 +1022,9 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
     for i in 0..graph.nodes.len() {
         if live.contains(&i) {
             for (o, out) in contracts[i].outputs.iter().enumerate() {
+                if out.semantic_type.is_virtual() {
+                    continue;
+                }
                 let key = OutputKey(i, o as u16);
                 if contracts[i].execution == ExecutionClass::Source
                     && !referenced_outputs.contains(&key)
@@ -1586,7 +1581,6 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
     for (i, o, out) in resource_meta {
         let key = OutputKey(i, o);
         let id = output_ids[&key];
-        let mesh = || output_ids[&root(key, &bound, &contracts).unwrap()];
         let plan = match out.semantic_type {
             SemanticType::Texture if matches!(params[i], NormalizedParameters::Texture { .. }) => {
                 if let NormalizedParameters::Texture {
@@ -1615,19 +1609,19 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
                 }
             }
             SemanticType::MeshData => ResourcePlan::MeshData,
-            SemanticType::LocalAabbBuffer => ResourcePlan::LocalAabbBuffer { mesh: mesh() },
-            SemanticType::BooleanFlagBuffer => {
-                if let OutputMetadata::BooleanFlag { flag } = out.metadata {
-                    ResourcePlan::BooleanFlagBuffer { mesh: mesh(), flag }
-                } else {
-                    unreachable!()
-                }
+            SemanticType::Bool
+            | SemanticType::F32
+            | SemanticType::U32
+            | SemanticType::Vec2
+            | SemanticType::Vec3
+            | SemanticType::Vec4
+            | SemanticType::Mat2
+            | SemanticType::Mat3
+            | SemanticType::Mat4
+            | SemanticType::U32x16
+            | SemanticType::LocalAabb => {
+                unreachable!("pure expression outputs are never materialized")
             }
-            SemanticType::PipelineIndexStream => ResourcePlan::PipelineIndexStream { mesh: mesh() },
-            SemanticType::PipelineActivation => ResourcePlan::PipelineActivation {
-                pipeline_indices: output_ids[&bound[i]["pipelineIndices"].producer],
-            },
-            SemanticType::DrawStream => ResourcePlan::DrawStream { mesh: mesh() },
         };
         resources.push(CompiledResource {
             original_node_index: i as u32,
@@ -1644,17 +1638,20 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
         let _ = id;
     }
     let mut executions = Vec::new();
-    let mut node_execution = HashMap::new();
     for &i in &order {
-        if contracts[i].execution == ExecutionClass::Source {
+        if matches!(
+            contracts[i].execution,
+            ExecutionClass::Source | ExecutionClass::Expression
+        ) {
             continue;
         }
-        let ordinal = executions.len() as u32;
-        node_execution.insert(i, ordinal);
         let input_resource = |s: &str| output_ids[&bound[i][s].producer];
         let mut inputs = Vec::new();
         for s in contracts[i].inputs {
-            if let Some(b) = bound[i].get(s.name).filter(|b| b.active) {
+            if s.role != InputRole::Expression {
+                let Some(b) = bound[i].get(s.name) else {
+                    continue;
+                };
                 inputs.push(CompiledSocketInput {
                     socket: s.name.into(),
                     resource: output_ids[&b.producer],
@@ -1665,6 +1662,7 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
             .outputs
             .iter()
             .enumerate()
+            .filter(|(_, s)| !s.semantic_type.is_virtual())
             .map(|(o, s)| CompiledSocketOutput {
                 socket: s.name.into(),
                 resource: output_ids[&OutputKey(i, o as u16)],
@@ -1672,58 +1670,6 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
             .collect();
         let mut accesses = Vec::new();
         let kind = match contracts[i].key {
-            "frustum_cull" => {
-                for (s, m) in [
-                    ("mesh", AccessMode::StorageRead),
-                    ("localAabbs", AccessMode::StorageRead),
-                ] {
-                    accesses.push(CompiledAccess {
-                        socket: s.into(),
-                        resource: input_resource(s),
-                        mode: m,
-                    });
-                }
-                let r = output_ids[&OutputKey(i, 0)];
-                accesses.push(CompiledAccess {
-                    socket: "isFrustumCulled".into(),
-                    resource: r,
-                    mode: AccessMode::StorageWrite {
-                        full_overwrite: true,
-                    },
-                });
-                ExecutionKind::Compute {
-                    work: ComputeWork::FrustumCull,
-                }
-            }
-            "mesh_query" => {
-                for s in ["mesh", "isVisible", "isFrustumCulled"] {
-                    if let Some(b) = bound[i].get(s).filter(|b| b.active) {
-                        accesses.push(CompiledAccess {
-                            socket: s.into(),
-                            resource: output_ids[&b.producer],
-                            mode: AccessMode::StorageRead,
-                        });
-                    }
-                }
-                accesses.push(CompiledAccess {
-                    socket: "draws".into(),
-                    resource: output_ids[&OutputKey(i, 0)],
-                    mode: AccessMode::StorageWrite {
-                        full_overwrite: true,
-                    },
-                });
-                ExecutionKind::Compute {
-                    work: ComputeWork::MeshQuery,
-                }
-            }
-            "pipeline_registry" => {
-                accesses.push(CompiledAccess {
-                    socket: "pipelineIndices".into(),
-                    resource: input_resource("pipelineIndices"),
-                    mode: AccessMode::SemanticRead,
-                });
-                ExecutionKind::CpuPreparation
-            }
             "pipeline" => {
                 let color = output_ids[&OutputKey(i, 0)];
                 let depth = output_ids[&OutputKey(i, 1)];
@@ -1747,15 +1693,11 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
                 } else {
                     NormalizedDepthLoad::Load
                 };
-                for s in ["mesh", "draws", "activation"] {
+                for s in ["mesh"] {
                     accesses.push(CompiledAccess {
                         socket: s.into(),
                         resource: input_resource(s),
-                        mode: if s == "draws" {
-                            AccessMode::IndirectRead
-                        } else {
-                            AccessMode::SemanticRead
-                        },
+                        mode: AccessMode::SemanticRead,
                     });
                 }
                 accesses.push(CompiledAccess {
@@ -1849,6 +1791,270 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
             accesses,
         });
     }
+    // Lower virtual values into a stable, device-independent expression plan.
+    let mut expression_plan = ExpressionPlan::default();
+    let mut expression_ids = HashMap::<OutputKey, ExprId>::new();
+    let mut expression_provenance = HashMap::<OutputKey, Option<u32>>::new();
+    let mut cse = HashMap::<String, ExprId>::new();
+    let mut requires_camera = false;
+    let mut mesh_root = None;
+    let mut intern = |semantic_type: SemanticType,
+                      op: ExpressionOp,
+                      origin: NodeOutputRef,
+                      mesh_provenance: Option<u32>| {
+        let key = format!("{semantic_type:?}:{op:?}:{mesh_provenance:?}");
+        if let Some(id) = cse.get(&key) {
+            return *id;
+        }
+        let id = ExprId(expression_plan.expressions.len() as u32);
+        expression_plan.expressions.push(Expression {
+            semantic_type,
+            op,
+            origin,
+            mesh_provenance,
+        });
+        cse.insert(key, id);
+        id
+    };
+    for &i in &order {
+        if contracts[i].execution != ExecutionClass::Expression {
+            continue;
+        }
+        let defaults: &[TypedLiteral] = match &params[i] {
+            NormalizedParameters::ExpressionDefaults { defaults } => defaults.as_slice(),
+            NormalizedParameters::FrustumCull { .. } => &[],
+            _ => unreachable!(),
+        };
+        let mut operands = Vec::new();
+        let mut operand_provenance = Vec::new();
+        for (ordinal, input) in contracts[i].inputs.iter().enumerate() {
+            if let Some(binding) = bound[i].get(input.name) {
+                let key = binding.producer;
+                let producer_type = contracts[key.0].outputs[key.1 as usize].semantic_type;
+                let operand_mesh = if contracts[key.0].key == "mesh" {
+                    Some(output_ids[&OutputKey(key.0, 0)])
+                } else {
+                    expression_provenance[&key]
+                };
+                let id = if producer_type.is_virtual() {
+                    if contracts[key.0].key == "mesh" {
+                        let mesh = output_ids[&OutputKey(key.0, 0)];
+                        if mesh_root.is_some_and(|root| root != mesh) {
+                            return Err(error(
+                                "GRAPH_SOCKET_TYPE_MISMATCH",
+                                "instance traversal has multiple mesh roots",
+                                format!("nodes[{i}].inputs.{}", input.name),
+                            ));
+                        }
+                        mesh_root = Some(mesh);
+                        let op = match producer_type {
+                            SemanticType::U32x16 => ExpressionOp::InstanceType { mesh },
+                            SemanticType::LocalAabb => ExpressionOp::LocalAabb { mesh },
+                            _ => unreachable!(),
+                        };
+                        intern(
+                            producer_type,
+                            op,
+                            graph.nodes[key.0]
+                                .inputs
+                                .get("")
+                                .cloned()
+                                .unwrap_or(NodeOutputRef {
+                                    node: graph.nodes[key.0].id.clone(),
+                                    socket: contracts[key.0].outputs[key.1 as usize].name.into(),
+                                }),
+                            Some(mesh),
+                        )
+                    } else {
+                        expression_ids[&key]
+                    }
+                } else {
+                    continue;
+                };
+                operands.push(id);
+                operand_provenance.push(operand_mesh);
+            } else {
+                let literal = defaults[ordinal].clone();
+                operands.push(intern(
+                    literal.semantic_type(),
+                    ExpressionOp::Literal { literal },
+                    NodeOutputRef {
+                        node: graph.nodes[i].id.clone(),
+                        socket: input.name.into(),
+                    },
+                    None,
+                ));
+                operand_provenance.push(None);
+            }
+        }
+        let mut provenances = operand_provenance.into_iter().flatten();
+        let provenance = provenances.next();
+        if provenances.any(|candidate| Some(candidate) != provenance) {
+            return Err(error(
+                "GRAPH_SOCKET_TYPE_MISMATCH",
+                "expression mixes mesh provenance",
+                format!("nodes[{i}].inputs"),
+            ));
+        }
+        for (output_ordinal, output) in contracts[i].outputs.iter().enumerate() {
+            let key = contracts[i].key;
+            let op = match key {
+                "not" => ExpressionOp::Not { value: operands[0] },
+                "and" | "or" | "xor" | "xnor" => ExpressionOp::BooleanBinary {
+                    operation: match key {
+                        "and" => BooleanBinaryOp::And,
+                        "or" => BooleanBinaryOp::Or,
+                        "xor" => BooleanBinaryOp::Xor,
+                        _ => BooleanBinaryOp::Xnor,
+                    },
+                    left: operands[0],
+                    right: operands[1],
+                },
+                "greater_than_f32" | "less_than_f32" | "equals_f32" => ExpressionOp::CompareF32 {
+                    operation: if key.starts_with("greater") {
+                        CompareOp::GreaterThan
+                    } else if key.starts_with("less") {
+                        CompareOp::LessThan
+                    } else {
+                        CompareOp::Equals
+                    },
+                    left: operands[0],
+                    right: operands[1],
+                },
+                "greater_than_u32" | "less_than_u32" | "equals_u32" => ExpressionOp::CompareU32 {
+                    operation: if key.starts_with("greater") {
+                        CompareOp::GreaterThan
+                    } else if key.starts_with("less") {
+                        CompareOp::LessThan
+                    } else {
+                        CompareOp::Equals
+                    },
+                    left: operands[0],
+                    right: operands[1],
+                },
+                k if k.starts_with("separate_vec") => ExpressionOp::VectorProject {
+                    vector: operands[0],
+                    index: output_ordinal as u8,
+                },
+                k if k.starts_with("combine_vec") => ExpressionOp::VectorConstruct {
+                    components: operands.clone(),
+                },
+                k if k.starts_with("separate_mat") => ExpressionOp::MatrixColumn {
+                    matrix: operands[0],
+                    index: output_ordinal as u8,
+                },
+                k if k.starts_with("combine_mat") => ExpressionOp::MatrixConstruct {
+                    columns: operands.clone(),
+                },
+                "separate_u32x16" => ExpressionOp::TypeWord {
+                    value: operands[0],
+                    index: output_ordinal as u8,
+                },
+                "combine_u32x16" => ExpressionOp::TypeConstruct {
+                    words: operands.clone(),
+                },
+                "separate_u32_bits" => ExpressionOp::U32Bit {
+                    value: operands[0],
+                    index: output_ordinal as u8,
+                },
+                "combine_u32_bits" => ExpressionOp::U32Construct {
+                    bits: operands.clone(),
+                },
+                "separate_local_aabb" if output_ordinal == 0 => {
+                    ExpressionOp::AabbMin { aabb: operands[0] }
+                }
+                "separate_local_aabb" => ExpressionOp::AabbMax { aabb: operands[0] },
+                "frustum_cull" => {
+                    requires_camera = true;
+                    let mesh = output_ids[&bound[i]["mesh"].producer];
+                    if mesh_root.is_some_and(|root| root != mesh) {
+                        return Err(error(
+                            "GRAPH_SOCKET_TYPE_MISMATCH",
+                            "instance traversal has multiple mesh roots",
+                            format!("nodes[{i}].inputs.mesh"),
+                        ));
+                    }
+                    mesh_root = Some(mesh);
+                    ExpressionOp::FrustumCulled {
+                        mesh,
+                        local_aabb: operands[0],
+                    }
+                }
+                _ => unreachable!(),
+            };
+            let id = intern(
+                output.semantic_type,
+                op,
+                NodeOutputRef {
+                    node: graph.nodes[i].id.clone(),
+                    socket: output.name.into(),
+                },
+                provenance,
+            );
+            expression_ids.insert(OutputKey(i, output_ordinal as u16), id);
+            expression_provenance.insert(OutputKey(i, output_ordinal as u16), provenance);
+        }
+    }
+    let mut predicates = Vec::new();
+    for (execution, compiled) in executions.iter().enumerate() {
+        let node = compiled.original_node_index as usize;
+        if contracts[node].key != "pipeline" {
+            continue;
+        }
+        let mesh = output_ids[&bound[node]["mesh"].producer];
+        if mesh_root.is_some_and(|root| root != mesh) {
+            return Err(error(
+                "GRAPH_SOCKET_TYPE_MISMATCH",
+                "instance traversal has multiple mesh roots",
+                format!("nodes[{node}].inputs.mesh"),
+            ));
+        }
+        mesh_root = Some(mesh);
+        let predicate = if let Some(binding) = bound[node].get("predicate") {
+            expression_ids[&binding.producer]
+        } else {
+            let NormalizedParameters::Pipeline {
+                predicate_default, ..
+            } = params[node]
+            else {
+                unreachable!()
+            };
+            intern(
+                SemanticType::Bool,
+                ExpressionOp::Literal {
+                    literal: TypedLiteral::Bool(predicate_default),
+                },
+                NodeOutputRef {
+                    node: graph.nodes[node].id.clone(),
+                    socket: "predicate".into(),
+                },
+                None,
+            )
+        };
+        predicates.push(PipelinePredicatePlan {
+            execution: execution as u32,
+            predicate,
+            ordinal: 0,
+        });
+    }
+    for (ordinal, predicate) in predicates.iter_mut().enumerate() {
+        predicate.ordinal = ordinal as u32;
+    }
+    if expression_plan.expressions.len() > MAX_EXPRESSIONS
+        || predicates.len() > MAX_PREDICATE_PIPELINES
+    {
+        return Err(error(
+            "GRAPH_LIMIT_EXCEEDED",
+            "instance traversal plan exceeds limits",
+            "nodes",
+        ));
+    }
+    let instance_traversal = mesh_root.map(|mesh| InstanceTraversalPlan {
+        mesh,
+        expressions: expression_plan,
+        pipelines: predicates,
+        requires_camera,
+    });
     for (ordinal, e) in executions.iter().enumerate() {
         for o in &e.outputs {
             resources[o.resource as usize].producer_execution = Some(ordinal as u32);
@@ -1917,6 +2123,7 @@ pub fn compile(graph: Graph) -> Result<CompiledGraph, GraphError> {
         culled_node_count: (graph.nodes.len() - live.len()) as u32,
         culled_resource_count: (all_outputs - output_ids.len()) as u32,
         transient_slot_count: transient,
+        instance_traversal,
     })
 }
 

renderer/src/render_graph/contracts.rs

@@ -1,27 +1,35 @@
-use super::MeshFlag;
-
-#[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize)]
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash, serde::Serialize)]
 #[serde(rename_all = "snake_case")]
 pub enum SemanticType {
     MeshData,
     Texture,
-    LocalAabbBuffer,
-    BooleanFlagBuffer,
-    PipelineIndexStream,
-    PipelineActivation,
-    DrawStream,
+    Bool,
+    F32,
+    U32,
+    Vec2,
+    Vec3,
+    Vec4,
+    Mat2,
+    Mat3,
+    Mat4,
+    U32x16,
+    LocalAabb,
+}
+
+impl SemanticType {
+    pub const fn is_virtual(self) -> bool {
+        !matches!(self, Self::MeshData | Self::Texture)
+    }
 }
 
 #[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize)]
 #[serde(rename_all = "snake_case")]
 pub enum ExecutionClass {
     Source,
-    CpuPreparation,
-    Compute,
+    Expression,
     Render,
     Frame,
 }
-
 #[derive(Clone, Copy, Debug, PartialEq, Eq)]
 pub enum FullscreenPolicy {
     Copy,
@@ -29,21 +37,18 @@ pub enum FullscreenPolicy {
     BloomExtract,
     BloomComposite,
 }
-
 #[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize)]
 #[serde(rename_all = "snake_case")]
 pub enum InputCardinality {
     RequiredOne,
     OptionalOne,
 }
-
 #[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize)]
 #[serde(tag = "kind", content = "types", rename_all = "snake_case")]
 pub enum TypeConstraint {
     Exact(SemanticType),
     OneOf(&'static [SemanticType]),
 }
-
 #[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize)]
 #[serde(tag = "kind", rename_all = "snake_case")]
 pub enum InputRole {
@@ -54,15 +59,8 @@ pub enum InputRole {
     SampledTexture,
     ColorTarget { location: u32 },
     DepthTarget,
+    Expression,
 }
-
-#[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize)]
-#[serde(tag = "kind", rename_all = "snake_case")]
-pub enum OutputMetadata {
-    None,
-    BooleanFlag { flag: MeshFlag },
-}
-
 #[derive(Clone, Copy, Debug, serde::Serialize)]
 #[serde(rename_all = "camelCase")]
 pub struct InputSocketContract {
@@ -71,15 +69,12 @@ pub struct InputSocketContract {
     pub cardinality: InputCardinality,
     pub role: InputRole,
 }
-
 #[derive(Clone, Copy, Debug, serde::Serialize)]
 #[serde(rename_all = "camelCase")]
 pub struct OutputSocketContract {
     pub name: &'static str,
     pub semantic_type: SemanticType,
-    pub metadata: OutputMetadata,
 }
-
 #[derive(Clone, Debug, serde::Serialize)]
 #[serde(rename_all = "camelCase")]
 pub struct Contract {
@@ -94,329 +89,294 @@ pub struct Contract {
 }
 
 use SemanticType::*;
-
-const fn input(
+const R: InputCardinality = InputCardinality::RequiredOne;
+const O: InputCardinality = InputCardinality::OptionalOne;
+const fn i(
     name: &'static str,
-    accepted: TypeConstraint,
+    ty: SemanticType,
     cardinality: InputCardinality,
     role: InputRole,
 ) -> InputSocketContract {
     InputSocketContract {
         name,
-        accepted,
+        accepted: TypeConstraint::Exact(ty),
         cardinality,
         role,
     }
 }
-
-const fn output(
-    name: &'static str,
-    semantic_type: SemanticType,
-    metadata: OutputMetadata,
-) -> OutputSocketContract {
+const fn o(name: &'static str, semantic_type: SemanticType) -> OutputSocketContract {
     OutputSocketContract {
         name,
         semantic_type,
-        metadata,
     }
 }
-
-const REQUIRED: InputCardinality = InputCardinality::RequiredOne;
-const OPTIONAL: InputCardinality = InputCardinality::OptionalOne;
-const NONE_IN: &[InputSocketContract] = &[];
-const NONE_OUT: &[OutputSocketContract] = &[];
-const TEXTURE_OUT: &[OutputSocketContract] = &[output("texture", Texture, OutputMetadata::None)];
-const MESH_OUT: &[OutputSocketContract] = &[
-    output("mesh", MeshData, OutputMetadata::None),
-    output("localAabbs", LocalAabbBuffer, OutputMetadata::None),
-    output(
-        "isVisible",
-        BooleanFlagBuffer,
-        OutputMetadata::BooleanFlag {
-            flag: MeshFlag::IsVisible,
-        },
-    ),
-    output("pipelineIndices", PipelineIndexStream, OutputMetadata::None),
+const NONE_I: &[InputSocketContract] = &[];
+const NONE_O: &[OutputSocketContract] = &[];
+const MESH_O: &[OutputSocketContract] = &[
+    o("mesh", MeshData),
+    o("type", U32x16),
+    o("localAabb", LocalAabb),
 ];
-const CULLED_OUT: &[OutputSocketContract] = &[output(
-    "isFrustumCulled",
-    BooleanFlagBuffer,
-    OutputMetadata::BooleanFlag {
-        flag: MeshFlag::IsFrustumCulled,
-    },
-)];
-const DRAW_OUT: &[OutputSocketContract] = &[output("draws", DrawStream, OutputMetadata::None)];
-const ACTIVATION_OUT: &[OutputSocketContract] = &[output(
-    "activation",
-    PipelineActivation,
-    OutputMetadata::None,
-)];
-const PIPELINE_OUT: &[OutputSocketContract] = &[
-    output("color", Texture, OutputMetadata::None),
-    output("depth", Texture, OutputMetadata::None),
-];
-const FULLSCREEN_COPY_OUT: &[OutputSocketContract] =
-    &[output("color", Texture, OutputMetadata::None)];
-const CULL_IN: &[InputSocketContract] = &[
-    input(
-        "mesh",
-        TypeConstraint::Exact(MeshData),
-        REQUIRED,
-        InputRole::StorageRead,
-    ),
-    input(
-        "localAabbs",
-        TypeConstraint::Exact(LocalAabbBuffer),
-        REQUIRED,
-        InputRole::StorageRead,
-    ),
-];
-const QUERY_IN: &[InputSocketContract] = &[
-    input(
-        "mesh",
-        TypeConstraint::Exact(MeshData),
-        REQUIRED,
-        InputRole::StorageRead,
-    ),
-    input(
-        "isVisible",
-        TypeConstraint::Exact(BooleanFlagBuffer),
-        OPTIONAL,
-        InputRole::StorageRead,
-    ),
-    input(
-        "isFrustumCulled",
-        TypeConstraint::Exact(BooleanFlagBuffer),
-        OPTIONAL,
-        InputRole::StorageRead,
-    ),
-];
-const REGISTRY_IN: &[InputSocketContract] = &[input(
-    "pipelineIndices",
-    TypeConstraint::Exact(PipelineIndexStream),
-    REQUIRED,
-    InputRole::SemanticRead,
-)];
-const PIPELINE_IN: &[InputSocketContract] = &[
-    input(
-        "mesh",
-        TypeConstraint::Exact(MeshData),
-        REQUIRED,
-        InputRole::SemanticRead,
-    ),
-    input(
-        "draws",
-        TypeConstraint::Exact(DrawStream),
-        REQUIRED,
-        InputRole::IndirectRead,
-    ),
-    input(
-        "activation",
-        TypeConstraint::Exact(PipelineActivation),
-        REQUIRED,
-        InputRole::SemanticRead,
-    ),
-    input(
+const TEXTURE_O: &[OutputSocketContract] = &[o("texture", Texture)];
+const PIPE_I: &[InputSocketContract] = &[
+    i("mesh", MeshData, R, InputRole::SemanticRead),
+    i("predicate", Bool, O, InputRole::Expression),
+    i(
         "colorTarget",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
+        Texture,
+        R,
         InputRole::ColorTarget { location: 0 },
     ),
-    input(
-        "depthTarget",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
-        InputRole::DepthTarget,
-    ),
+    i("depthTarget", Texture, R, InputRole::DepthTarget),
 ];
-const FULLSCREEN_COPY_IN: &[InputSocketContract] = &[
-    input(
-        "source",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
-        InputRole::SampledTexture,
-    ),
-    input(
+const PIPE_O: &[OutputSocketContract] = &[o("color", Texture), o("depth", Texture)];
+const CULL_I: &[InputSocketContract] = &[
+    i("mesh", MeshData, R, InputRole::Expression),
+    i("localAabb", LocalAabb, R, InputRole::Expression),
+];
+const CULL_O: &[OutputSocketContract] = &[o("isFrustumCulled", Bool)];
+const COPY_I: &[InputSocketContract] = &[
+    i("source", Texture, R, InputRole::SampledTexture),
+    i(
         "colorTarget",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
+        Texture,
+        R,
         InputRole::ColorTarget { location: 0 },
     ),
 ];
-const BLOOM_COMPOSITE_IN: &[InputSocketContract] = &[
-    input(
-        "source",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
-        InputRole::SampledTexture,
-    ),
-    input(
-        "bloom",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
-        InputRole::SampledTexture,
-    ),
-    input(
+const BLOOM_I: &[InputSocketContract] = &[
+    i("source", Texture, R, InputRole::SampledTexture),
+    i("bloom", Texture, R, InputRole::SampledTexture),
+    i(
         "colorTarget",
-        TypeConstraint::Exact(Texture),
-        REQUIRED,
+        Texture,
+        R,
         InputRole::ColorTarget { location: 0 },
     ),
 ];
-const FRAME_OUT_IN: &[InputSocketContract] = &[input(
-    "color",
-    TypeConstraint::Exact(Texture),
-    REQUIRED,
-    InputRole::SampledTexture,
-)];
+const COLOR_O: &[OutputSocketContract] = &[o("color", Texture)];
+const FRAME_I: &[InputSocketContract] = &[i("color", Texture, R, InputRole::SampledTexture)];
+macro_rules! ins { ($($n:literal:$t:ident),*) => { &[$(i($n,$t,O,InputRole::Expression)),*] } }
+macro_rules! outs { ($($n:literal:$t:ident),*) => { &[$(o($n,$t)),*] } }
+macro_rules! c {
+    ($k:literal,$v:expr,$e:ident,$ins:expr,$outs:expr,$obs:expr,$policy:expr) => {
+        Contract {
+            key: $k,
+            version: $v,
+            execution: ExecutionClass::$e,
+            inputs: $ins,
+            outputs: $outs,
+            inherently_observable: $obs,
+            fullscreen_policy: $policy,
+        }
+    };
+}
+macro_rules! ex {
+    ($k:literal,$ins:expr,$outs:expr) => {
+        c!($k, 1, Expression, $ins, $outs, false, None)
+    };
+}
 
 pub static CONTRACTS: &[Contract] = &[
-    Contract {
-        key: "mesh",
-        version: 1,
-        execution: ExecutionClass::Source,
-        inputs: NONE_IN,
-        outputs: MESH_OUT,
-        inherently_observable: false,
-        fullscreen_policy: None,
-    },
-    Contract {
-        key: "texture",
-        version: 1,
-        execution: ExecutionClass::Source,
-        inputs: NONE_IN,
-        outputs: TEXTURE_OUT,
-        inherently_observable: false,
-        fullscreen_policy: None,
-    },
-    Contract {
-        key: "frustum_cull",
-        version: 1,
-        execution: ExecutionClass::Compute,
-        inputs: CULL_IN,
-        outputs: CULLED_OUT,
-        inherently_observable: false,
-        fullscreen_policy: None,
-    },
-    Contract {
-        key: "mesh_query",
-        version: 1,
-        execution: ExecutionClass::Compute,
-        inputs: QUERY_IN,
-        outputs: DRAW_OUT,
-        inherently_observable: false,
-        fullscreen_policy: None,
-    },
-    Contract {
-        key: "pipeline_registry",
-        version: 1,
-        execution: ExecutionClass::CpuPreparation,
-        inputs: REGISTRY_IN,
-        outputs: ACTIVATION_OUT,
-        inherently_observable: false,
-        fullscreen_policy: None,
-    },
-    Contract {
-        key: "pipeline",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: PIPELINE_IN,
-        outputs: PIPELINE_OUT,
-        inherently_observable: false,
-        fullscreen_policy: None,
-    },
-    Contract {
-        key: "fullscreen_copy",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::Copy),
-    },
-    Contract {
-        key: "color_balance",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::HdrSameExtent),
-    },
-    Contract {
-        key: "exposure_contrast",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::HdrSameExtent),
-    },
-    Contract {
-        key: "saturation",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::HdrSameExtent),
-    },
-    Contract {
-        key: "channel_mixer",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::HdrSameExtent),
-    },
-    Contract {
-        key: "bloom_extract",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::BloomExtract),
-    },
-    Contract {
-        key: "bloom_blur",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::HdrSameExtent),
-    },
-    Contract {
-        key: "bloom_composite",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: BLOOM_COMPOSITE_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::BloomComposite),
-    },
-    Contract {
-        key: "luminance_edge",
-        version: 1,
-        execution: ExecutionClass::Render,
-        inputs: FULLSCREEN_COPY_IN,
-        outputs: FULLSCREEN_COPY_OUT,
-        inherently_observable: false,
-        fullscreen_policy: Some(FullscreenPolicy::HdrSameExtent),
-    },
-    Contract {
-        key: "frame_out",
-        version: 3,
-        execution: ExecutionClass::Frame,
-        inputs: FRAME_OUT_IN,
-        outputs: NONE_OUT,
-        inherently_observable: true,
-        fullscreen_policy: None,
-    },
+    c!("mesh", 2, Source, NONE_I, MESH_O, false, None),
+    c!("texture", 1, Source, NONE_I, TEXTURE_O, false, None),
+    c!("frustum_cull", 2, Expression, CULL_I, CULL_O, false, None),
+    c!("pipeline", 2, Render, PIPE_I, PIPE_O, false, None),
+    ex!("and", ins!("left":Bool,"right":Bool), outs!("value":Bool)),
+    ex!("or", ins!("left":Bool,"right":Bool), outs!("value":Bool)),
+    ex!("not", ins!("operand":Bool), outs!("value":Bool)),
+    ex!("xor", ins!("left":Bool,"right":Bool), outs!("value":Bool)),
+    ex!("xnor", ins!("left":Bool,"right":Bool), outs!("value":Bool)),
+    ex!(
+        "greater_than_f32",
+        ins!("left":F32,"right":F32),
+        outs!("value":Bool)
+    ),
+    ex!(
+        "less_than_f32",
+        ins!("left":F32,"right":F32),
+        outs!("value":Bool)
+    ),
+    ex!(
+        "equals_f32",
+        ins!("left":F32,"right":F32),
+        outs!("value":Bool)
+    ),
+    ex!(
+        "greater_than_u32",
+        ins!("left":U32,"right":U32),
+        outs!("value":Bool)
+    ),
+    ex!(
+        "less_than_u32",
+        ins!("left":U32,"right":U32),
+        outs!("value":Bool)
+    ),
+    ex!(
+        "equals_u32",
+        ins!("left":U32,"right":U32),
+        outs!("value":Bool)
+    ),
+    ex!("separate_vec2", ins!("vector":Vec2), outs!("x":F32,"y":F32)),
+    ex!("combine_vec2", ins!("x":F32,"y":F32), outs!("vector":Vec2)),
+    ex!(
+        "separate_vec3",
+        ins!("vector":Vec3),
+        outs!("x":F32,"y":F32,"z":F32)
+    ),
+    ex!(
+        "combine_vec3",
+        ins!("x":F32,"y":F32,"z":F32),
+        outs!("vector":Vec3)
+    ),
+    ex!(
+        "separate_vec4",
+        ins!("vector":Vec4),
+        outs!("x":F32,"y":F32,"z":F32,"w":F32)
+    ),
+    ex!(
+        "combine_vec4",
+        ins!("x":F32,"y":F32,"z":F32,"w":F32),
+        outs!("vector":Vec4)
+    ),
+    ex!(
+        "separate_mat2",
+        ins!("matrix":Mat2),
+        outs!("column0":Vec2,"column1":Vec2)
+    ),
+    ex!(
+        "combine_mat2",
+        ins!("column0":Vec2,"column1":Vec2),
+        outs!("matrix":Mat2)
+    ),
+    ex!(
+        "separate_mat3",
+        ins!("matrix":Mat3),
+        outs!("column0":Vec3,"column1":Vec3,"column2":Vec3)
+    ),
+    ex!(
+        "combine_mat3",
+        ins!("column0":Vec3,"column1":Vec3,"column2":Vec3),
+        outs!("matrix":Mat3)
+    ),
+    ex!(
+        "separate_mat4",
+        ins!("matrix":Mat4),
+        outs!("column0":Vec4,"column1":Vec4,"column2":Vec4,"column3":Vec4)
+    ),
+    ex!(
+        "combine_mat4",
+        ins!("column0":Vec4,"column1":Vec4,"column2":Vec4,"column3":Vec4),
+        outs!("matrix":Mat4)
+    ),
+    ex!(
+        "separate_u32x16",
+        ins!("value":U32x16),
+        outs!("word0":U32,"word1":U32,"word2":U32,"word3":U32,"word4":U32,"word5":U32,"word6":U32,"word7":U32,"word8":U32,"word9":U32,"word10":U32,"word11":U32,"word12":U32,"word13":U32,"word14":U32,"word15":U32)
+    ),
+    ex!(
+        "combine_u32x16",
+        ins!("word0":U32,"word1":U32,"word2":U32,"word3":U32,"word4":U32,"word5":U32,"word6":U32,"word7":U32,"word8":U32,"word9":U32,"word10":U32,"word11":U32,"word12":U32,"word13":U32,"word14":U32,"word15":U32),
+        outs!("value":U32x16)
+    ),
+    ex!(
+        "separate_u32_bits",
+        ins!("value":U32),
+        outs!("bit0":Bool,"bit1":Bool,"bit2":Bool,"bit3":Bool,"bit4":Bool,"bit5":Bool,"bit6":Bool,"bit7":Bool,"bit8":Bool,"bit9":Bool,"bit10":Bool,"bit11":Bool,"bit12":Bool,"bit13":Bool,"bit14":Bool,"bit15":Bool,"bit16":Bool,"bit17":Bool,"bit18":Bool,"bit19":Bool,"bit20":Bool,"bit21":Bool,"bit22":Bool,"bit23":Bool,"bit24":Bool,"bit25":Bool,"bit26":Bool,"bit27":Bool,"bit28":Bool,"bit29":Bool,"bit30":Bool,"bit31":Bool)
+    ),
+    ex!(
+        "combine_u32_bits",
+        ins!("bit0":Bool,"bit1":Bool,"bit2":Bool,"bit3":Bool,"bit4":Bool,"bit5":Bool,"bit6":Bool,"bit7":Bool,"bit8":Bool,"bit9":Bool,"bit10":Bool,"bit11":Bool,"bit12":Bool,"bit13":Bool,"bit14":Bool,"bit15":Bool,"bit16":Bool,"bit17":Bool,"bit18":Bool,"bit19":Bool,"bit20":Bool,"bit21":Bool,"bit22":Bool,"bit23":Bool,"bit24":Bool,"bit25":Bool,"bit26":Bool,"bit27":Bool,"bit28":Bool,"bit29":Bool,"bit30":Bool,"bit31":Bool),
+        outs!("value":U32)
+    ),
+    ex!(
+        "separate_local_aabb",
+        ins!("value":LocalAabb),
+        outs!("min":Vec3,"max":Vec3)
+    ),
+    c!(
+        "fullscreen_copy",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::Copy)
+    ),
+    c!(
+        "color_balance",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::HdrSameExtent)
+    ),
+    c!(
+        "exposure_contrast",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::HdrSameExtent)
+    ),
+    c!(
+        "saturation",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::HdrSameExtent)
+    ),
+    c!(
+        "channel_mixer",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::HdrSameExtent)
+    ),
+    c!(
+        "bloom_extract",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::BloomExtract)
+    ),
+    c!(
+        "bloom_blur",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::HdrSameExtent)
+    ),
+    c!(
+        "bloom_composite",
+        1,
+        Render,
+        BLOOM_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::BloomComposite)
+    ),
+    c!(
+        "luminance_edge",
+        1,
+        Render,
+        COPY_I,
+        COLOR_O,
+        false,
+        Some(FullscreenPolicy::HdrSameExtent)
+    ),
+    c!("frame_out", 3, Frame, FRAME_I, NONE_O, true, None),
 ];
-
 pub fn contract(key: &str) -> Option<&'static Contract> {
-    CONTRACTS.iter().find(|contract| contract.key == key)
+    CONTRACTS.iter().find(|c| c.key == key)
 }

renderer/src/render_graph/expression.rs

@@ -0,0 +1,181 @@
+//! Typed, device-independent instance predicate IR.
+
+use serde::Serialize;
+
+use super::{NodeOutputRef, SemanticType};
+
+pub const MAX_EXPRESSIONS: usize = 4096;
+pub const MAX_PREDICATE_PIPELINES: usize = 64;
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash, PartialOrd, Ord, Serialize)]
+#[serde(transparent)]
+pub struct ExprId(pub u32);
+
+#[derive(Clone, Debug, PartialEq, Serialize)]
+#[serde(tag = "type", content = "value", rename_all = "snake_case")]
+pub enum TypedLiteral {
+    Bool(bool),
+    F32(f32),
+    U32(u32),
+    Vec2([f32; 2]),
+    Vec3([f32; 3]),
+    Vec4([f32; 4]),
+    Mat2([[f32; 2]; 2]),
+    Mat3([[f32; 3]; 3]),
+    Mat4([[f32; 4]; 4]),
+    U32x16([u32; 16]),
+    LocalAabb { min: [f32; 3], max: [f32; 3] },
+}
+
+impl TypedLiteral {
+    pub fn semantic_type(&self) -> SemanticType {
+        match self {
+            Self::Bool(_) => SemanticType::Bool,
+            Self::F32(_) => SemanticType::F32,
+            Self::U32(_) => SemanticType::U32,
+            Self::Vec2(_) => SemanticType::Vec2,
+            Self::Vec3(_) => SemanticType::Vec3,
+            Self::Vec4(_) => SemanticType::Vec4,
+            Self::Mat2(_) => SemanticType::Mat2,
+            Self::Mat3(_) => SemanticType::Mat3,
+            Self::Mat4(_) => SemanticType::Mat4,
+            Self::U32x16(_) => SemanticType::U32x16,
+            Self::LocalAabb { .. } => SemanticType::LocalAabb,
+        }
+    }
+
+    pub fn is_finite(&self) -> bool {
+        let finite = |values: &[f32]| values.iter().all(|value| value.is_finite());
+        match self {
+            Self::F32(value) => value.is_finite(),
+            Self::Vec2(value) => finite(value),
+            Self::Vec3(value) => finite(value),
+            Self::Vec4(value) => finite(value),
+            Self::Mat2(value) => value.iter().all(|column| finite(column)),
+            Self::Mat3(value) => value.iter().all(|column| finite(column)),
+            Self::Mat4(value) => value.iter().all(|column| finite(column)),
+            Self::LocalAabb { min, max } => finite(min) && finite(max),
+            _ => true,
+        }
+    }
+}
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash, Serialize)]
+#[serde(rename_all = "snake_case")]
+pub enum CompareOp {
+    GreaterThan,
+    LessThan,
+    Equals,
+}
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash, Serialize)]
+#[serde(rename_all = "snake_case")]
+pub enum BooleanBinaryOp {
+    And,
+    Or,
+    Xor,
+    Xnor,
+}
+
+/// All operand IDs refer to earlier entries in [`ExpressionPlan::expressions`].
+#[derive(Clone, Debug, PartialEq, Serialize)]
+#[serde(tag = "op", rename_all = "snake_case")]
+pub enum ExpressionOp {
+    Literal {
+        literal: TypedLiteral,
+    },
+    InstanceType {
+        mesh: u32,
+    },
+    LocalAabb {
+        mesh: u32,
+    },
+    Not {
+        value: ExprId,
+    },
+    BooleanBinary {
+        operation: BooleanBinaryOp,
+        left: ExprId,
+        right: ExprId,
+    },
+    CompareF32 {
+        operation: CompareOp,
+        left: ExprId,
+        right: ExprId,
+    },
+    CompareU32 {
+        operation: CompareOp,
+        left: ExprId,
+        right: ExprId,
+    },
+    VectorProject {
+        vector: ExprId,
+        index: u8,
+    },
+    VectorConstruct {
+        components: Vec<ExprId>,
+    },
+    MatrixColumn {
+        matrix: ExprId,
+        index: u8,
+    },
+    MatrixConstruct {
+        columns: Vec<ExprId>,
+    },
+    TypeWord {
+        value: ExprId,
+        index: u8,
+    },
+    TypeConstruct {
+        words: Vec<ExprId>,
+    },
+    U32Bit {
+        value: ExprId,
+        index: u8,
+    },
+    U32Construct {
+        bits: Vec<ExprId>,
+    },
+    AabbMin {
+        aabb: ExprId,
+    },
+    AabbMax {
+        aabb: ExprId,
+    },
+    FrustumCulled {
+        mesh: u32,
+        local_aabb: ExprId,
+    },
+}
+
+#[derive(Clone, Debug, PartialEq, Serialize)]
+#[serde(rename_all = "camelCase")]
+pub struct Expression {
+    pub semantic_type: SemanticType,
+    pub op: ExpressionOp,
+    pub origin: NodeOutputRef,
+    pub mesh_provenance: Option<u32>,
+}
+
+#[derive(Clone, Debug, Default, PartialEq, Serialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ExpressionPlan {
+    pub expressions: Vec<Expression>,
+}
+
+#[derive(Clone, Debug, PartialEq, Eq, Serialize)]
+#[serde(rename_all = "camelCase")]
+pub struct PipelinePredicatePlan {
+    pub execution: u32,
+    pub predicate: ExprId,
+    pub ordinal: u32,
+}
+
+#[derive(Clone, Debug, Default, PartialEq, Serialize)]
+#[serde(rename_all = "camelCase")]
+pub struct InstanceTraversalPlan {
+    pub mesh: u32,
+    pub expressions: ExpressionPlan,
+    pub pipelines: Vec<PipelinePredicatePlan>,
+    pub requires_camera: bool,
+}

renderer/src/render_graph/mod.rs

@@ -2,13 +2,15 @@
 
 mod compiler;
 mod contracts;
+mod expression;
 mod plan;
 mod registry;
 mod runtime;
 mod schema;
 
-pub use compiler::{compile, mesh_predicate_matches, parse_and_compile};
+pub use compiler::{compile, parse_and_compile};
 pub use contracts::*;
+pub use expression::*;
 pub use plan::*;
 pub use registry::{CompiledGraphId, Registry};
 pub use runtime::*;

renderer/src/render_graph/plan.rs

@@ -16,6 +16,8 @@ pub struct CompiledGraph {
     pub culled_node_count: u32,
     pub culled_resource_count: u32,
     pub transient_slot_count: u32,
+    #[serde(skip_serializing_if = "Option::is_none")]
+    pub instance_traversal: Option<InstanceTraversalPlan>,
 }
 
 #[derive(Clone, Debug, Serialize)]
@@ -47,22 +49,6 @@ pub enum ResourcePlan {
         allocation: Option<AllocationRef>,
     },
     MeshData,
-    LocalAabbBuffer {
-        mesh: u32,
-    },
-    BooleanFlagBuffer {
-        mesh: u32,
-        flag: MeshFlag,
-    },
-    PipelineIndexStream {
-        mesh: u32,
-    },
-    PipelineActivation {
-        pipeline_indices: u32,
-    },
-    DrawStream {
-        mesh: u32,
-    },
 }
 
 #[derive(Clone, Debug, Serialize)]
@@ -95,10 +81,6 @@ pub struct CompiledSocketOutput {
 #[derive(Clone, Debug, Serialize)]
 #[serde(tag = "kind", rename_all = "snake_case")]
 pub enum ExecutionKind {
-    CpuPreparation,
-    Compute {
-        work: ComputeWork,
-    },
     Render {
         color_attachments: Vec<ColorAttachmentPlan>,
         depth_stencil: Option<DepthStencilAttachmentPlan>,
@@ -108,13 +90,6 @@ pub enum ExecutionKind {
     },
 }
 
-#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, serde::Deserialize)]
-#[serde(rename_all = "snake_case")]
-pub enum ComputeWork {
-    FrustumCull,
-    MeshQuery,
-}
-
 #[derive(Clone, Debug, Serialize)]
 #[serde(rename_all = "camelCase")]
 pub struct ColorAttachmentPlan {
@@ -196,17 +171,16 @@ pub enum NormalizedParameters {
     FrustumCull {
         camera: ActiveCamera,
     },
-    MeshQuery {
-        visible_predicate: RuntimePredicate,
-        frustum_culled_predicate: RuntimePredicate,
+    ExpressionDefaults {
+        defaults: Vec<TypedLiteral>,
     },
-    PipelineRegistry,
     Pipeline {
         pipeline: String,
         depth_compare: CompareFunction,
         depth_write_enabled: bool,
         clear_depth: f32,
         clear_color: [f64; 4],
+        predicate_default: bool,
     },
     FullscreenCopy,
     ColorBalance {

renderer/src/render_graph/runtime.rs

@@ -166,12 +166,6 @@ pub struct RuntimeAllocationClass {
     pub slots: Vec<RuntimeAllocationSlot>,
 }
 
-#[derive(Clone, Copy, Debug, PartialEq, Eq)]
-pub struct MeshQueryRuntimeKey {
-    pub visible: RuntimePredicate,
-    pub frustum_culled: RuntimePredicate,
-}
-
 #[derive(Clone, Debug, PartialEq, Eq)]
 pub struct RuntimeExecution {
     pub execution: u32,
@@ -182,13 +176,13 @@ pub struct RuntimeExecution {
 pub struct RuntimeAllocationPlan {
     pub classes: Vec<RuntimeAllocationClass>,
     pub resource_allocations: Vec<Option<AllocationRef>>,
-    pub query: MeshQueryRuntimeKey,
 }
 
-#[derive(Clone, Debug, PartialEq, Eq)]
+#[derive(Clone, Debug, PartialEq)]
 pub struct RuntimePlan {
     pub allocations: RuntimeAllocationPlan,
     pub executions: Vec<RuntimeExecution>,
+    pub instance_traversal: Option<InstanceTraversalPlan>,
     pub surface: RuntimeSurfaceContract,
 }
 
@@ -381,11 +375,7 @@ fn valid_pipeline_name(name: &str) -> bool {
 
 fn execution_supported(key: &str) -> bool {
     contract(key).is_some_and(|contract| {
-        contract.fullscreen_policy.is_some()
-            || matches!(
-                key,
-                "frustum_cull" | "mesh_query" | "pipeline_registry" | "pipeline" | "frame_out"
-            )
+        contract.fullscreen_policy.is_some() || matches!(key, "pipeline" | "frame_out")
     })
 }
 
@@ -396,19 +386,6 @@ fn resource_is_mesh(graph: &CompiledGraph, id: u32) -> bool {
     })
 }
 
-fn has_exact_producer(
-    graph: &CompiledGraph,
-    consumer: usize,
-    resource: u32,
-    executor: &str,
-    socket: &str,
-) -> bool {
-    graph.executions[..consumer].iter().any(|execution| {
-        execution.executor.key == executor
-            && matches!(execution.outputs.as_slice(), [output] if output.socket == socket && output.resource == resource)
-    })
-}
-
 fn texture_descriptor<'a>(
     graph: &'a CompiledGraph,
     id: u32,
@@ -704,181 +681,15 @@ fn validate_fullscreen_execution(
     Ok(())
 }
 
-fn validate_compute_execution(
-    graph: &CompiledGraph,
-    i: usize,
-    execution: &CompiledExecution,
-) -> Result<(), GraphError> {
-    let path = |field| format!("executions[{i}].{field}");
-    match execution.executor.key.as_str() {
-        "frustum_cull" => {
-            if !matches!(
-                execution.parameters,
-                NormalizedParameters::FrustumCull {
-                    camera: ActiveCamera::Active
-                }
-            ) {
-                return Err(invalid(
-                    "frustum cull parameters mismatch",
-                    path("parameters"),
-                ));
-            }
-            if !matches!(
-                execution.kind,
-                ExecutionKind::Compute {
-                    work: ComputeWork::FrustumCull
-                }
-            ) {
-                return Err(invalid("frustum cull work mismatch", path("kind")));
-            }
-            let [mesh, aabbs] = execution.inputs.as_slice() else {
-                return Err(invalid("frustum cull inputs mismatch", path("inputs")));
-            };
-            let [flags] = execution.outputs.as_slice() else {
-                return Err(invalid("frustum cull outputs mismatch", path("outputs")));
-            };
-            if mesh.socket != "mesh"
-                || aabbs.socket != "localAabbs"
-                || flags.socket != "isFrustumCulled"
-            {
-                return Err(invalid(
-                    "frustum cull socket order mismatch",
-                    path("inputs"),
-                ));
-            }
-            if !matches!(execution.accesses.as_slice(),
-                [CompiledAccess { socket: s0, resource: r0, mode: AccessMode::StorageRead },
-                 CompiledAccess { socket: s1, resource: r1, mode: AccessMode::StorageRead },
-                 CompiledAccess { socket: s2, resource: r2, mode: AccessMode::StorageWrite { full_overwrite: true } }]
-                if s0 == "mesh" && *r0 == mesh.resource && s1 == "localAabbs" && *r1 == aabbs.resource
-                    && s2 == "isFrustumCulled" && *r2 == flags.resource)
-            {
-                return Err(invalid("frustum cull accesses mismatch", path("accesses")));
-            }
-            if !resource_is_mesh(graph, mesh.resource)
-                || !matches!(graph.resources[aabbs.resource as usize], CompiledResource { semantic_type: SemanticType::LocalAabbBuffer, plan: ResourcePlan::LocalAabbBuffer { mesh: m }, .. } if m == mesh.resource)
-                || !matches!(graph.resources[flags.resource as usize], CompiledResource { semantic_type: SemanticType::BooleanFlagBuffer, plan: ResourcePlan::BooleanFlagBuffer { mesh: m, flag: MeshFlag::IsFrustumCulled }, .. } if m == mesh.resource)
-            {
-                return Err(invalid(
-                    "frustum cull mesh provenance mismatch",
-                    path("inputs"),
-                ));
-            }
-        }
-        "mesh_query" => {
-            let NormalizedParameters::MeshQuery {
-                visible_predicate,
-                frustum_culled_predicate,
-            } = execution.parameters
-            else {
-                return Err(invalid(
-                    "mesh query parameters mismatch",
-                    path("parameters"),
-                ));
-            };
-            if (visible_predicate == RuntimePredicate::Never)
-                != (frustum_culled_predicate == RuntimePredicate::Never)
-            {
-                return Err(invalid(
-                    "mesh query never predicates must be paired",
-                    path("parameters"),
-                ));
-            }
-            if !matches!(
-                execution.kind,
-                ExecutionKind::Compute {
-                    work: ComputeWork::MeshQuery
-                }
-            ) {
-                return Err(invalid("mesh query work mismatch", path("kind")));
-            }
-            let active = |p| {
-                matches!(
-                    p,
-                    RuntimePredicate::RequiredTrue | RuntimePredicate::RequiredFalse
-                )
-            };
-            let mut sockets = vec!["mesh"];
-            if active(visible_predicate) {
-                sockets.push("isVisible");
-            }
-            if active(frustum_culled_predicate) {
-                sockets.push("isFrustumCulled");
-            }
-            if execution.inputs.len() != sockets.len()
-                || execution
-                    .inputs
-                    .iter()
-                    .zip(&sockets)
-                    .any(|(v, s)| v.socket != *s)
-                || !matches!(execution.outputs.as_slice(), [CompiledSocketOutput { socket, .. }] if socket == "draws")
-            {
-                return Err(invalid("mesh query socket order mismatch", path("inputs")));
-            }
-            let output = execution.outputs[0].resource;
-            if execution.accesses.len() != sockets.len() + 1
-                || execution
-                    .inputs
-                    .iter()
-                    .zip(&execution.accesses)
-                    .any(|(input, access)| {
-                        access.socket != input.socket
-                            || access.resource != input.resource
-                            || !matches!(access.mode, AccessMode::StorageRead)
-                    })
-                || !matches!(&execution.accesses[sockets.len()], CompiledAccess { socket, resource, mode: AccessMode::StorageWrite { full_overwrite: true } } if socket == "draws" && *resource == output)
-            {
-                return Err(invalid("mesh query accesses mismatch", path("accesses")));
-            }
-            let mesh = execution.inputs[0].resource;
-            if !resource_is_mesh(graph, mesh) {
-                return Err(invalid(
-                    "mesh query mesh provenance mismatch",
-                    path("inputs"),
-                ));
-            }
-            for input in execution.inputs.iter().skip(1) {
-                let flag = if input.socket == "isVisible" {
-                    MeshFlag::IsVisible
-                } else {
-                    MeshFlag::IsFrustumCulled
-                };
-                if !matches!(graph.resources[input.resource as usize], CompiledResource { semantic_type: SemanticType::BooleanFlagBuffer, plan: ResourcePlan::BooleanFlagBuffer { mesh: m, flag: f }, .. } if m == mesh && f == flag)
-                {
-                    return Err(invalid(
-                        "mesh query flag provenance mismatch",
-                        path("inputs"),
-                    ));
-                }
-                if flag == MeshFlag::IsFrustumCulled
-                    && !has_exact_producer(
-                        graph,
-                        i,
-                        input.resource,
-                        "frustum_cull",
-                        "isFrustumCulled",
-                    )
-                {
-                    return Err(invalid(
-                        "mesh query frustum flag producer mismatch",
-                        path("inputs"),
-                    ));
-                }
-            }
-            if !matches!(graph.resources[output as usize], CompiledResource { semantic_type: SemanticType::DrawStream, plan: ResourcePlan::DrawStream { mesh: m }, .. } if m == mesh)
-            {
-                return Err(invalid(
-                    "mesh query output provenance mismatch",
-                    path("outputs"),
-                ));
-            }
-        }
-        _ => {}
-    }
-    Ok(())
-}
-
 fn validate_canonical_plan(graph: &CompiledGraph) -> Result<(), GraphError> {
+    for (i, resource) in graph.resources.iter().enumerate() {
+        if resource.semantic_type.is_virtual() {
+            return Err(invalid(
+                "virtual semantic type was materialized",
+                format!("resources[{i}].semanticType"),
+            ));
+        }
+    }
     for (i, execution) in graph.executions.iter().enumerate() {
         if !execution_supported(&execution.executor.key) {
             return Err(error(
@@ -939,7 +750,6 @@ fn validate_canonical_plan(graph: &CompiledGraph) -> Result<(), GraphError> {
             }
             referenced.insert(access.resource);
         }
-        validate_compute_execution(graph, i, execution)?;
         validate_fullscreen_execution(graph, i, execution, contract)?;
         for resource in referenced {
             uses.get_mut(resource as usize)
@@ -1071,6 +881,291 @@ fn validate_canonical_plan(graph: &CompiledGraph) -> Result<(), GraphError> {
     Ok(())
 }
 
+fn validate_instance_traversal(graph: &CompiledGraph) -> Result<(), GraphError> {
+    let pipeline_indices: Vec<_> = graph
+        .executions
+        .iter()
+        .enumerate()
+        .filter_map(|(i, e)| (e.executor.key == "pipeline").then_some(i as u32))
+        .collect();
+    let Some(plan) = &graph.instance_traversal else {
+        return if pipeline_indices.is_empty() {
+            Ok(())
+        } else {
+            Err(invalid(
+                "live pipelines require instance traversal",
+                "instanceTraversal",
+            ))
+        };
+    };
+    if pipeline_indices.is_empty() {
+        return Err(invalid(
+            "instance traversal has no live pipelines",
+            "instanceTraversal",
+        ));
+    }
+    if !resource_is_mesh(graph, plan.mesh) {
+        return Err(invalid(
+            "traversal mesh is invalid",
+            "instanceTraversal.mesh",
+        ));
+    }
+    let expressions = &plan.expressions.expressions;
+    if expressions.len() > MAX_EXPRESSIONS || plan.pipelines.len() > MAX_PREDICATE_PIPELINES {
+        return Err(invalid(
+            "instance traversal exceeds runtime limits",
+            "instanceTraversal",
+        ));
+    }
+    let ty = |id: ExprId| expressions.get(id.0 as usize).map(|e| e.semantic_type);
+    for (i, expression) in expressions.iter().enumerate() {
+        let path = format!("instanceTraversal.expressions.expressions[{i}]");
+        let ids: Vec<ExprId> = match &expression.op {
+            ExpressionOp::Literal { literal } => {
+                if literal.semantic_type() != expression.semantic_type || !literal.is_finite() {
+                    return Err(invalid("literal type or value is invalid", &path));
+                }
+                vec![]
+            }
+            ExpressionOp::InstanceType { mesh } => {
+                if expression.semantic_type != SemanticType::U32x16 || *mesh != plan.mesh {
+                    return Err(invalid("instance type signature is invalid", &path));
+                }
+                vec![]
+            }
+            ExpressionOp::LocalAabb { mesh } => {
+                if expression.semantic_type != SemanticType::LocalAabb || *mesh != plan.mesh {
+                    return Err(invalid("local aabb signature is invalid", &path));
+                }
+                vec![]
+            }
+            ExpressionOp::Not { value } => {
+                if expression.semantic_type != SemanticType::Bool
+                    || ty(*value) != Some(SemanticType::Bool)
+                {
+                    return Err(invalid("not signature is invalid", &path));
+                }
+                vec![*value]
+            }
+            ExpressionOp::BooleanBinary { left, right, .. } => {
+                if expression.semantic_type != SemanticType::Bool
+                    || ty(*left) != Some(SemanticType::Bool)
+                    || ty(*right) != Some(SemanticType::Bool)
+                {
+                    return Err(invalid("boolean signature is invalid", &path));
+                }
+                vec![*left, *right]
+            }
+            ExpressionOp::CompareF32 { left, right, .. } => {
+                if expression.semantic_type != SemanticType::Bool
+                    || ty(*left) != Some(SemanticType::F32)
+                    || ty(*right) != Some(SemanticType::F32)
+                {
+                    return Err(invalid("f32 comparison signature is invalid", &path));
+                }
+                vec![*left, *right]
+            }
+            ExpressionOp::CompareU32 { left, right, .. } => {
+                if expression.semantic_type != SemanticType::Bool
+                    || ty(*left) != Some(SemanticType::U32)
+                    || ty(*right) != Some(SemanticType::U32)
+                {
+                    return Err(invalid("u32 comparison signature is invalid", &path));
+                }
+                vec![*left, *right]
+            }
+            ExpressionOp::VectorProject { vector, index } => {
+                let n = match ty(*vector) {
+                    Some(SemanticType::Vec2) => 2,
+                    Some(SemanticType::Vec3) => 3,
+                    Some(SemanticType::Vec4) => 4,
+                    _ => 0,
+                };
+                if expression.semantic_type != SemanticType::F32 || usize::from(*index) >= n {
+                    return Err(invalid("vector projection signature is invalid", &path));
+                }
+                vec![*vector]
+            }
+            ExpressionOp::VectorConstruct { components } => {
+                let n = match expression.semantic_type {
+                    SemanticType::Vec2 => 2,
+                    SemanticType::Vec3 => 3,
+                    SemanticType::Vec4 => 4,
+                    _ => 0,
+                };
+                if components.len() != n
+                    || components
+                        .iter()
+                        .any(|id| ty(*id) != Some(SemanticType::F32))
+                {
+                    return Err(invalid("vector constructor signature is invalid", &path));
+                }
+                components.clone()
+            }
+            ExpressionOp::MatrixColumn { matrix, index } => {
+                let (n, out) = match ty(*matrix) {
+                    Some(SemanticType::Mat2) => (2, SemanticType::Vec2),
+                    Some(SemanticType::Mat3) => (3, SemanticType::Vec3),
+                    Some(SemanticType::Mat4) => (4, SemanticType::Vec4),
+                    _ => (0, SemanticType::Bool),
+                };
+                if usize::from(*index) >= n || expression.semantic_type != out {
+                    return Err(invalid("matrix column signature is invalid", &path));
+                }
+                vec![*matrix]
+            }
+            ExpressionOp::MatrixConstruct { columns } => {
+                let (n, col) = match expression.semantic_type {
+                    SemanticType::Mat2 => (2, SemanticType::Vec2),
+                    SemanticType::Mat3 => (3, SemanticType::Vec3),
+                    SemanticType::Mat4 => (4, SemanticType::Vec4),
+                    _ => (0, SemanticType::Bool),
+                };
+                if columns.len() != n || columns.iter().any(|id| ty(*id) != Some(col)) {
+                    return Err(invalid("matrix constructor signature is invalid", &path));
+                }
+                columns.clone()
+            }
+            ExpressionOp::TypeWord { value, index } => {
+                if expression.semantic_type != SemanticType::U32
+                    || ty(*value) != Some(SemanticType::U32x16)
+                    || *index >= 16
+                {
+                    return Err(invalid("type word signature is invalid", &path));
+                }
+                vec![*value]
+            }
+            ExpressionOp::TypeConstruct { words } => {
+                if expression.semantic_type != SemanticType::U32x16
+                    || words.len() != 16
+                    || words.iter().any(|id| ty(*id) != Some(SemanticType::U32))
+                {
+                    return Err(invalid("type constructor signature is invalid", &path));
+                }
+                words.clone()
+            }
+            ExpressionOp::U32Bit { value, index } => {
+                if expression.semantic_type != SemanticType::Bool
+                    || ty(*value) != Some(SemanticType::U32)
+                    || *index >= 32
+                {
+                    return Err(invalid("bit signature is invalid", &path));
+                }
+                vec![*value]
+            }
+            ExpressionOp::U32Construct { bits } => {
+                if expression.semantic_type != SemanticType::U32
+                    || bits.len() != 32
+                    || bits.iter().any(|id| ty(*id) != Some(SemanticType::Bool))
+                {
+                    return Err(invalid("u32 constructor signature is invalid", &path));
+                }
+                bits.clone()
+            }
+            ExpressionOp::AabbMin { aabb } | ExpressionOp::AabbMax { aabb } => {
+                if expression.semantic_type != SemanticType::Vec3
+                    || ty(*aabb) != Some(SemanticType::LocalAabb)
+                {
+                    return Err(invalid("aabb projection signature is invalid", &path));
+                }
+                vec![*aabb]
+            }
+            ExpressionOp::FrustumCulled { mesh, local_aabb } => {
+                if expression.semantic_type != SemanticType::Bool
+                    || *mesh != plan.mesh
+                    || ty(*local_aabb) != Some(SemanticType::LocalAabb)
+                    || expressions
+                        .get(local_aabb.0 as usize)
+                        .and_then(|e| e.mesh_provenance)
+                        != Some(plan.mesh)
+                {
+                    return Err(invalid("frustum signature or provenance is invalid", &path));
+                }
+                vec![*local_aabb]
+            }
+        };
+        if ids.iter().any(|id| id.0 as usize >= i) {
+            return Err(invalid("expression operands must precede consumer", &path));
+        }
+        let expected_provenance = match &expression.op {
+            ExpressionOp::Literal { .. } => None,
+            ExpressionOp::InstanceType { .. } | ExpressionOp::LocalAabb { .. } => Some(plan.mesh),
+            _ => {
+                let mut p = ids
+                    .iter()
+                    .filter_map(|id| expressions[id.0 as usize].mesh_provenance);
+                let first = p.next();
+                if p.any(|v| Some(v) != first) {
+                    return Err(invalid("expression mixes mesh provenance", &path));
+                }
+                first
+            }
+        };
+        if expression.mesh_provenance != expected_provenance {
+            return Err(invalid("expression provenance is not canonical", &path));
+        }
+    }
+    let mut seen = HashSet::new();
+    let mut reachable = vec![false; expressions.len()];
+    for (ordinal, entry) in plan.pipelines.iter().enumerate() {
+        if entry.ordinal as usize != ordinal
+            || pipeline_indices.get(ordinal) != Some(&entry.execution)
+            || !seen.insert(entry.execution)
+            || ty(entry.predicate) != Some(SemanticType::Bool)
+        {
+            return Err(invalid(
+                "pipeline predicate table is not canonical",
+                format!("instanceTraversal.pipelines[{ordinal}]"),
+            ));
+        }
+        let mut stack = vec![entry.predicate];
+        while let Some(id) = stack.pop() {
+            if reachable[id.0 as usize] {
+                continue;
+            }
+            reachable[id.0 as usize] = true;
+            stack.extend(expression_operands(&expressions[id.0 as usize].op));
+        }
+    }
+    if plan.pipelines.len() != pipeline_indices.len() {
+        return Err(invalid(
+            "pipeline predicate table is incomplete",
+            "instanceTraversal.pipelines",
+        ));
+    }
+    let requires_camera = expressions
+        .iter()
+        .enumerate()
+        .any(|(i, e)| reachable[i] && matches!(e.op, ExpressionOp::FrustumCulled { .. }));
+    if plan.requires_camera != requires_camera {
+        return Err(invalid(
+            "requires_camera is not canonical",
+            "instanceTraversal.requiresCamera",
+        ));
+    }
+    Ok(())
+}
+
+fn expression_operands(op: &ExpressionOp) -> Vec<ExprId> {
+    match op {
+        ExpressionOp::Not { value }
+        | ExpressionOp::VectorProject { vector: value, .. }
+        | ExpressionOp::MatrixColumn { matrix: value, .. }
+        | ExpressionOp::TypeWord { value, .. }
+        | ExpressionOp::U32Bit { value, .. } => vec![*value],
+        ExpressionOp::AabbMin { aabb } | ExpressionOp::AabbMax { aabb } => vec![*aabb],
+        ExpressionOp::FrustumCulled { local_aabb, .. } => vec![*local_aabb],
+        ExpressionOp::BooleanBinary { left, right, .. }
+        | ExpressionOp::CompareF32 { left, right, .. }
+        | ExpressionOp::CompareU32 { left, right, .. } => vec![*left, *right],
+        ExpressionOp::VectorConstruct { components } => components.clone(),
+        ExpressionOp::MatrixConstruct { columns } => columns.clone(),
+        ExpressionOp::TypeConstruct { words } => words.clone(),
+        ExpressionOp::U32Construct { bits } => bits.clone(),
+        _ => vec![],
+    }
+}
+
 pub fn prepare_runtime_plan(
     graph: &CompiledGraph,
     surface: RuntimeSurfaceContract,
@@ -1095,35 +1190,13 @@ pub fn prepare_runtime_plan(
         ));
     }
     let mut frame_out_index = None;
-    let mut query = None;
     let mut executions = Vec::with_capacity(graph.executions.len());
     for (i, execution) in graph.executions.iter().enumerate() {
         let path = format!("executions[{i}]");
         match execution.executor.key.as_str() {
-            "mesh_query" => {
-                let NormalizedParameters::MeshQuery {
-                    visible_predicate,
-                    frustum_culled_predicate,
-                } = &execution.parameters
-                else {
-                    return Err(invalid("mesh query parameters mismatch", &path));
-                };
-                let key = MeshQueryRuntimeKey {
-                    visible: *visible_predicate,
-                    frustum_culled: *frustum_culled_predicate,
-                };
-                if query.replace(key).is_some() {
-                    return Err(error(
-                        "GRAPH_EXECUTION_UNSUPPORTED",
-                        "multiple draw stream queries",
-                        &path,
-                    ));
-                }
-            }
-            "pipeline_registry" | "pipeline" => {}
+            "pipeline" => {}
             _ if contract(&execution.executor.key)
                 .is_some_and(|contract| contract.fullscreen_policy.is_some()) => {}
-            "frustum_cull" => {}
             "frame_out" => {
                 if frame_out_index.replace(i).is_some() {
                     return Err(error(
@@ -1153,115 +1226,7 @@ pub fn prepare_runtime_plan(
             "executions",
         )
     })?;
-    let query = query.ok_or_else(|| {
-        error(
-            "GRAPH_EXECUTION_UNSUPPORTED",
-            "one mesh query is required",
-            "executions",
-        )
-    })?;
-
-    for (i, execution) in graph.executions.iter().enumerate() {
-        if execution.executor.key != "pipeline_registry" {
-            continue;
-        }
-        if !matches!(execution.parameters, NormalizedParameters::PipelineRegistry) {
-            return Err(invalid(
-                "pipeline registry parameters mismatch",
-                format!("executions[{i}].parameters"),
-            ));
-        }
-        if !matches!(execution.kind, ExecutionKind::CpuPreparation) {
-            return Err(invalid(
-                "pipeline registry kind mismatch",
-                format!("executions[{i}].kind"),
-            ));
-        }
-        let [CompiledSocketInput {
-            socket: input_socket,
-            resource: pipeline_indices,
-        }] = execution.inputs.as_slice()
-        else {
-            return Err(invalid(
-                "pipeline registry input shape mismatch",
-                format!("executions[{i}].inputs"),
-            ));
-        };
-        if input_socket != "pipelineIndices" {
-            return Err(invalid(
-                "pipeline registry input socket mismatch",
-                format!("executions[{i}].inputs"),
-            ));
-        }
-        let [CompiledSocketOutput {
-            socket: output_socket,
-            resource: activation,
-        }] = execution.outputs.as_slice()
-        else {
-            return Err(invalid(
-                "pipeline registry output shape mismatch",
-                format!("executions[{i}].outputs"),
-            ));
-        };
-        if output_socket != "activation" {
-            return Err(invalid(
-                "pipeline registry output socket mismatch",
-                format!("executions[{i}].outputs"),
-            ));
-        }
-        if !matches!(
-            execution.accesses.as_slice(),
-            [CompiledAccess { socket, resource, mode: AccessMode::SemanticRead }]
-                if socket == "pipelineIndices" && resource == pipeline_indices
-        ) {
-            return Err(invalid(
-                "pipeline registry access mismatch",
-                format!("executions[{i}].accesses"),
-            ));
-        }
-        let indices_resource =
-            graph
-                .resources
-                .get(*pipeline_indices as usize)
-                .ok_or_else(|| {
-                    invalid(
-                        "pipeline index stream is out of bounds",
-                        format!("executions[{i}].inputs"),
-                    )
-                })?;
-        let ResourcePlan::PipelineIndexStream { mesh } = indices_resource.plan else {
-            return Err(invalid(
-                "pipeline registry input is not a pipeline index stream",
-                format!("resources[{pipeline_indices}].plan"),
-            ));
-        };
-        if indices_resource.semantic_type != SemanticType::PipelineIndexStream
-            || !graph.resources.get(mesh as usize).is_some_and(|resource| {
-                resource.semantic_type == SemanticType::MeshData
-                    && matches!(resource.plan, ResourcePlan::MeshData)
-            })
-        {
-            return Err(invalid(
-                "pipeline index stream mesh provenance is invalid",
-                format!("resources[{pipeline_indices}].plan"),
-            ));
-        }
-        let activation_resource = graph.resources.get(*activation as usize).ok_or_else(|| {
-            invalid(
-                "pipeline activation is out of bounds",
-                format!("executions[{i}].outputs"),
-            )
-        })?;
-        if activation_resource.semantic_type != SemanticType::PipelineActivation
-            || !matches!(activation_resource.plan, ResourcePlan::PipelineActivation { pipeline_indices: source } if source == *pipeline_indices)
-            || activation_resource.producer_execution != Some(i as u32)
-        {
-            return Err(invalid(
-                "pipeline activation provenance is invalid",
-                format!("resources[{activation}].plan"),
-            ));
-        }
-    }
+    validate_instance_traversal(graph)?;
 
     for (i, execution) in graph.executions.iter().enumerate() {
         if execution.executor.key != "pipeline" {
@@ -1305,9 +1270,7 @@ pub fn prepare_runtime_plan(
                 format!("executions[{i}].kind"),
             ));
         };
-        let [mesh_input, draws_input, activation_input, color_input, depth_input] =
-            execution.inputs.as_slice()
-        else {
+        let [mesh_input, color_input, depth_input] = execution.inputs.as_slice() else {
             return Err(invalid(
                 "pipeline input shape mismatch",
                 format!("executions[{i}].inputs"),
@@ -1315,11 +1278,9 @@ pub fn prepare_runtime_plan(
         };
         if [
             mesh_input.socket.as_str(),
-            draws_input.socket.as_str(),
-            activation_input.socket.as_str(),
             color_input.socket.as_str(),
             depth_input.socket.as_str(),
-        ] != ["mesh", "draws", "activation", "colorTarget", "depthTarget"]
+        ] != ["mesh", "colorTarget", "depthTarget"]
         {
             return Err(invalid(
                 "pipeline input sockets mismatch",
@@ -1394,9 +1355,7 @@ pub fn prepare_runtime_plan(
                 format!("executions[{i}].kind"),
             ));
         }
-        let [mesh_access, draws_access, activation_access, color_access, depth_access] =
-            execution.accesses.as_slice()
-        else {
+        let [mesh_access, color_access, depth_access] = execution.accesses.as_slice() else {
             return Err(invalid(
                 "pipeline access shape mismatch",
                 format!("executions[{i}].accesses"),
@@ -1405,12 +1364,6 @@ pub fn prepare_runtime_plan(
         if mesh_access.socket != "mesh"
             || mesh_access.resource != mesh_input.resource
             || !matches!(mesh_access.mode, AccessMode::SemanticRead)
-            || draws_access.socket != "draws"
-            || draws_access.resource != draws_input.resource
-            || !matches!(draws_access.mode, AccessMode::IndirectRead)
-            || activation_access.socket != "activation"
-            || activation_access.resource != activation_input.resource
-            || !matches!(activation_access.mode, AccessMode::SemanticRead)
         {
             return Err(invalid(
                 "pipeline semantic accesses mismatch",
@@ -1459,73 +1412,6 @@ pub fn prepare_runtime_plan(
                 format!("resources[{}].plan", mesh_input.resource),
             ));
         }
-        let draws_mesh = match graph.resources.get(draws_input.resource as usize) {
-            Some(CompiledResource {
-                semantic_type: SemanticType::DrawStream,
-                plan: ResourcePlan::DrawStream { mesh },
-                ..
-            }) => *mesh,
-            _ => {
-                return Err(invalid(
-                    "pipeline draw stream is invalid",
-                    format!("resources[{}].plan", draws_input.resource),
-                ))
-            }
-        };
-        let activation_resource = graph
-            .resources
-            .get(activation_input.resource as usize)
-            .ok_or_else(|| {
-                invalid(
-                    "pipeline activation is out of bounds",
-                    format!("executions[{i}].inputs"),
-                )
-            })?;
-        let indices = match activation_resource.plan {
-            ResourcePlan::PipelineActivation { pipeline_indices }
-                if activation_resource.semantic_type == SemanticType::PipelineActivation =>
-            {
-                pipeline_indices
-            }
-            _ => {
-                return Err(invalid(
-                    "pipeline activation is invalid",
-                    format!("resources[{}].plan", activation_input.resource),
-                ))
-            }
-        };
-        let activation_mesh = match graph.resources.get(indices as usize) {
-            Some(CompiledResource {
-                semantic_type: SemanticType::PipelineIndexStream,
-                plan: ResourcePlan::PipelineIndexStream { mesh },
-                ..
-            }) => *mesh,
-            _ => {
-                return Err(invalid(
-                    "pipeline activation index stream is invalid",
-                    format!("resources[{indices}].plan"),
-                ))
-            }
-        };
-        let valid_activation_producer = activation_resource
-            .producer_execution
-            .and_then(|producer| graph.executions.get(producer as usize))
-            .is_some_and(|producer| producer.executor.key == "pipeline_registry");
-        if draws_mesh != mesh_input.resource
-            || activation_mesh != mesh_input.resource
-            || !valid_activation_producer
-        {
-            return Err(invalid(
-                "pipeline mesh provenance disagrees",
-                format!("executions[{i}].inputs"),
-            ));
-        }
-        if !has_exact_producer(graph, i, draws_input.resource, "mesh_query", "draws") {
-            return Err(invalid(
-                "pipeline draw stream producer mismatch",
-                format!("executions[{i}].inputs"),
-            ));
-        }
         for (output, target) in [
             (color_output.resource, color_input.resource),
             (depth_output.resource, depth_input.resource),
@@ -1969,9 +1855,9 @@ pub fn prepare_runtime_plan(
         allocations: RuntimeAllocationPlan {
             classes,
             resource_allocations,
-            query,
         },
         executions,
+        instance_traversal: graph.instance_traversal.clone(),
         surface,
     })
 }

renderer/src/render_graph/schema.rs

@@ -106,41 +106,6 @@ pub struct TextureDescriptor {
     pub view_formats: Vec<TextureFormat>,
 }
 
-#[derive(Debug, Clone, Copy, Deserialize, Serialize, PartialEq, Eq, PartialOrd, Ord)]
-#[serde(rename_all = "snake_case")]
-pub enum TriStatePredicate {
-    Any,
-    RequiredTrue,
-    RequiredFalse,
-}
-
-#[derive(Debug, Clone, Copy, Deserialize, Serialize, PartialEq, Eq, PartialOrd, Ord)]
-#[serde(rename_all = "snake_case")]
-pub enum RuntimePredicate {
-    Any,
-    RequiredTrue,
-    RequiredFalse,
-    Never,
-}
-
-#[derive(Debug, Clone, Copy, Deserialize, Serialize, PartialEq, Eq, PartialOrd, Ord)]
-#[serde(rename_all = "camelCase")]
-pub enum MeshFlag {
-    IsVisible,
-    IsFrustumCulled,
-}
-
-impl MeshFlag {
-    pub const ORDERED: [Self; 2] = [Self::IsVisible, Self::IsFrustumCulled];
-
-    pub const fn input_socket(self) -> &'static str {
-        match self {
-            Self::IsVisible => "isVisible",
-            Self::IsFrustumCulled => "isFrustumCulled",
-        }
-    }
-}
-
 #[derive(Debug, Clone, Copy, Deserialize, Serialize, PartialEq, Eq)]
 #[serde(rename_all = "snake_case")]
 pub enum CompareFunction {

renderer/src/renderer/culling.wgsl

@@ -1,56 +0,0 @@
-struct Params { planes: array<vec4<f32>, 6>, count: u32, visible_predicate: u32, frustum_predicate: u32, _pad: u32 }
-struct Instance { model: mat4x4<f32>, n0: vec4<f32>, n1: vec4<f32>, n2: vec4<f32> }
-struct Aabb { min: vec4<f32>, max: vec4<f32> }
-struct Meta { index_count: u32, first_index: u32, base_vertex: i32, instance_index: u32 }
-struct Command { index_count: u32, instance_count: u32, first_index: u32, base_vertex: i32, first_instance: u32 }
-@group(0) @binding(0) var<uniform> params: Params;
-@group(0) @binding(1) var<storage, read> instances: array<Instance>;
-@group(0) @binding(2) var<storage, read> bounds: array<Aabb>;
-@group(0) @binding(3) var<storage, read> authored_visible: array<u32>;
-@group(0) @binding(4) var<storage, read> metadata: array<Meta>;
-@group(0) @binding(5) var<storage, read_write> frustum_flags: array<u32>;
-@group(0) @binding(6) var<storage, read_write> commands: array<Command>;
-
-@compute @workgroup_size(64)
-fn frustum_cull(@builtin(global_invocation_id) id: vec3<u32>) {
-  let i = id.x;
-  if (i >= params.count) { return; }
-  let b = bounds[i];
-  let center = (b.min.xyz + b.max.xyz) * 0.5;
-  let extent = (b.max.xyz - b.min.xyz) * 0.5;
-  let m = instances[i].model;
-  let wc = (m * vec4<f32>(center, 1.0)).xyz;
-  let ax = m[0].xyz * extent.x;
-  let ay = m[1].xyz * extent.y;
-  let az = m[2].xyz * extent.z;
-  var inside = 1u;
-  for (var p = 0u; p < 6u; p++) {
-    let plane = params.planes[p];
-    let radius = abs(dot(plane.xyz, ax)) + abs(dot(plane.xyz, ay)) + abs(dot(plane.xyz, az));
-    if (dot(plane.xyz, wc) + plane.w + radius < 0.0) { inside = 0u; }
-  }
-  frustum_flags[i] = 1u - inside;
-}
-
-fn matches(value: u32, predicate: u32) -> bool {
-  return predicate == 0u || (predicate == 1u && value != 0u) || (predicate == 2u && value == 0u);
-}
-
-@compute @workgroup_size(64)
-fn mesh_query(@builtin(global_invocation_id) id: vec3<u32>) {
-  let i = id.x;
-  if (i >= params.count) { return; }
-  let draw_meta = metadata[i];
-  var selected = true;
-  if (params.visible_predicate == 3u) {
-    selected = false;
-  } else if (params.visible_predicate != 0u) {
-    selected = matches(authored_visible[i], params.visible_predicate);
-  }
-  if (selected && params.frustum_predicate == 3u) {
-    selected = false;
-  } else if (selected && params.frustum_predicate != 0u) {
-    selected = selected && matches(frustum_flags[i], params.frustum_predicate);
-  }
-  commands[i] = Command(draw_meta.index_count, select(0u, 1u, selected), draw_meta.first_index, draw_meta.base_vertex, 0u);
-}

renderer/src/renderer/executors/pipeline.rs

@@ -30,7 +30,7 @@ fn encode_scene<'a, T: Scene>(
         pass.set_vertex_buffer(4, t.slice(..));
         pass.set_index_buffer(i.slice(..), wgpu::IndexFormat::Uint32);
         for draw in &gpu.draws {
-            if !draw.effective_visible {
+            if !draw.instance_type.is_visible() {
                 continue;
             }
             pass.set_pipeline(pipelines.get_pipeline(draw.pipeline));
@@ -54,6 +54,7 @@ pub(crate) fn encode_compiled<T: Scene>(
     gpu: &GpuSceneCache,
     pipelines: &PipelineLibrary,
     materials: &MaterialResources,
+    indirect_commands: &wgpu::Buffer,
     mut profile: Option<&mut crate::renderer::profiler::ProfileFrame>,
 ) -> Result<(), &'static str> {
     use crate::render_graph::{ExecutionKind, NormalizedColorLoad, NormalizedDepthLoad, StoreOp};
@@ -73,16 +74,8 @@ pub(crate) fn encode_compiled<T: Scene>(
             .map(|s| &s.view)
             .ok_or(" allocation out of bounds")
     };
-    for (execution_index, prepared) in active.executions.iter().enumerate() {
-        let profile_id = &active.graph.executions[execution_index].id;
+    for prepared in &active.executions {
         match prepared {
-            PreparedExecution::PipelineRegistry => {}
-            PreparedExecution::FrustumCull => {
-                gpu.encode_frustum_cull(encoder, profile.as_deref_mut(), profile_id);
-            }
-            PreparedExecution::MeshQuery => {
-                gpu.encode_mesh_query(encoder, profile.as_deref_mut(), profile_id);
-            }
             PreparedExecution::Fullscreen {
                 execution,
                 frame_out,
@@ -159,6 +152,7 @@ pub(crate) fn encode_compiled<T: Scene>(
             PreparedExecution::Pipeline {
                 execution,
                 base,
+                predicate_ordinal,
                 variant,
             } => {
                 let execution = active
@@ -236,24 +230,19 @@ pub(crate) fn encode_compiled<T: Scene>(
                     pass.set_vertex_buffer(2, u.slice(..));
                     pass.set_vertex_buffer(4, t.slice(..));
                     pass.set_index_buffer(ix.slice(..), wgpu::IndexFormat::Uint32);
-                    for draw in &gpu.draws {
-                        if draw.pipeline != *base {
-                            continue;
-                        }
-                        let slot = draw.instances.start as u64;
-                        let start = slot
-                            * std::mem::size_of::<crate::renderer::gpu_scene::GpuInstance>() as u64;
-                        pass.set_vertex_buffer(3, inst.slice(start..start + 112));
+                    pass.set_vertex_buffer(3, inst.slice(..));
+                    for (draw_index, draw) in gpu.draws.iter().enumerate() {
                         pass.set_pipeline(variant);
-                        if pipelines.requires_material(draw.pipeline) {
+                        if pipelines.requires_material(*base) {
                             pass.set_bind_group(2, materials.group(draw.material), &[]);
                         }
                         pass.draw_indexed_indirect(
-                            gpu.indirect_commands
-                                .buffer
-                                .as_ref()
-                                .ok_or("indirect command buffer missing")?,
-                            slot * 20,
+                            indirect_commands,
+                            crate::renderer::instance_traversal::command_offset(
+                                *predicate_ordinal,
+                                gpu.draws.len(),
+                                draw_index,
+                            ),
                         );
                     }
                 }

renderer/src/renderer/gpu_scene.rs

@@ -1,10 +1,11 @@
 use std::mem::size_of;
 
+use bytemuck::{Pod, Zeroable};
+
 use crate::{
-    render_data::{MaterialKey, MeshHandle, PipelineKey, RenderFlags},
+    render_data::{InstanceType, MaterialKey, MeshHandle, PipelineKey},
     renderer::scene_frame::SceneFramePlan,
 };
-use bytemuck::{Pod, Zeroable};
 
 #[repr(C)]
 #[derive(Clone, Copy, Debug, Pod, Zeroable, PartialEq)]
@@ -23,7 +24,7 @@ pub struct DrawItem {
     pub indices: std::ops::Range<u32>,
     pub base_vertex: i32,
     pub instances: std::ops::Range<u32>,
-    pub effective_visible: bool,
+    pub instance_type: InstanceType,
 }
 
 #[repr(C)]
@@ -62,162 +63,142 @@ pub struct GpuScenePlan {
     pub instances: Vec<GpuInstance>,
     pub draws: Vec<DrawItem>,
     pub local_aabbs: Vec<GpuLocalAabb>,
-    pub effective_visibility: Vec<u32>,
+    pub instance_types: Vec<[u32; 16]>,
     pub draw_metadata: Vec<DrawSlotMetadata>,
-    pub commands: Vec<DrawIndexedIndirect>,
-}
-
-fn visibility_matches(
-    predicate: crate::render_graph::TriStatePredicate,
-    mesh: RenderFlags,
-    instance: RenderFlags,
-) -> bool {
-    let effective = mesh.contains(RenderFlags::VISIBLE) && instance.contains(RenderFlags::VISIBLE);
-    match predicate {
-        crate::render_graph::TriStatePredicate::Any => true,
-        crate::render_graph::TriStatePredicate::RequiredTrue => effective,
-        crate::render_graph::TriStatePredicate::RequiredFalse => !effective,
-    }
 }
 
 impl GpuScenePlan {
     pub fn build(data: &SceneFramePlan) -> Result<Self, &'static str> {
-        self::GpuScenePlan::build_with_query(
-            data,
-            crate::render_graph::MeshQueryRuntimeKey {
-                visible: crate::render_graph::RuntimePredicate::RequiredTrue,
-                frustum_culled: crate::render_graph::RuntimePredicate::Any,
-            },
-        )
-    }
-
-    pub fn build_with_query(
-        data: &SceneFramePlan,
-        query: crate::render_graph::MeshQueryRuntimeKey,
-    ) -> Result<Self, &'static str> {
-        let _ = query; // Packing is canonical; predicates are evaluated by the GPU.
-        let mut plan = Self::default();
+        let mut p = Self::default();
         let mut meshes: Vec<_> = data.meshes.iter().collect();
-        meshes.sort_by_key(|mesh| {
+        meshes.sort_by_key(|m| {
             (
-                mesh.pipeline.get(),
-                mesh.material.get(),
-                mesh.handle.slot(),
-                mesh.handle.generation(),
+                m.pipeline.get(),
+                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()
-                .map(|&index| &data.occurrences[index])
+                .map(|&i| &data.occurrences[i])
                 .collect();
             if occurrences.is_empty() {
                 continue;
             }
-            let vertex_start = plan.positions.len();
-            let source_start = mesh.geometry.vertex_start as usize;
-            let source_end = source_start
+            let vs = mesh.geometry.vertex_start as usize;
+            let ve = vs
                 .checked_add(mesh.geometry.vertex_count as usize)
                 .ok_or("vertex range overflow")?;
-            plan.positions.extend_from_slice(
-                data.positions
-                    .get(source_start..source_end)
-                    .ok_or("invalid vertex range")?,
-            );
-            plan.normals.extend_from_slice(
-                data.normals
-                    .get(source_start..source_end)
-                    .ok_or("invalid normal range")?,
-            );
-            plan.uvs.extend_from_slice(
-                data.uvs
-                    .get(source_start..source_end)
-                    .ok_or("invalid uv range")?,
-            );
-            plan.tangents.extend_from_slice(
-                data.tangents
-                    .get(source_start..source_end)
-                    .ok_or("invalid tangent range")?,
-            );
-            let index_start =
-                u32::try_from(plan.indices.len()).map_err(|_| "index start exceeds u32")?;
-            let source_index = mesh.geometry.index_start as usize;
-            let source_index_end = source_index
+            let vertex_start = p.positions.len();
+            p.positions
+                .extend_from_slice(data.positions.get(vs..ve).ok_or("invalid vertex range")?);
+            p.normals
+                .extend_from_slice(data.normals.get(vs..ve).ok_or("invalid normal range")?);
+            p.uvs
+                .extend_from_slice(data.uvs.get(vs..ve).ok_or("invalid uv range")?);
+            p.tangents
+                .extend_from_slice(data.tangents.get(vs..ve).ok_or("invalid tangent range")?);
+            let first_index =
+                u32::try_from(p.indices.len()).map_err(|_| "index start exceeds u32")?;
+            let is = mesh.geometry.index_start as usize;
+            let ie = is
                 .checked_add(mesh.geometry.index_count as usize)
                 .ok_or("index range overflow")?;
-            plan.indices.extend_from_slice(
-                data.indices
-                    .get(source_index..source_index_end)
-                    .ok_or("invalid index range")?,
-            );
-            for instance in occurrences {
-                let instance_start = u32::try_from(plan.instances.len())
-                    .map_err(|_| "instance start exceeds u32")?;
-                let m = &instance.model;
-                let determinant = m[0][0] * (m[1][1] * m[2][2] - m[2][1] * m[1][2])
+            p.indices
+                .extend_from_slice(data.indices.get(is..ie).ok_or("invalid index range")?);
+            for occurrence in occurrences {
+                let instance_index =
+                    u32::try_from(p.instances.len()).map_err(|_| "instance start exceeds u32")?;
+                let m = &occurrence.model;
+                let det = m[0][0] * (m[1][1] * m[2][2] - m[2][1] * m[1][2])
                     - m[1][0] * (m[0][1] * m[2][2] - m[2][1] * m[0][2])
                     + m[2][0] * (m[0][1] * m[1][2] - m[1][1] * m[0][2]);
-                plan.instances.push(GpuInstance {
-                    model: instance.model,
+                p.instances.push(GpuInstance {
+                    model: occurrence.model,
                     normal_0: [
-                        instance.normal[0][0],
-                        instance.normal[0][1],
-                        instance.normal[0][2],
-                        if determinant < 0.0 { -1.0 } else { 1.0 },
+                        occurrence.normal[0][0],
+                        occurrence.normal[0][1],
+                        occurrence.normal[0][2],
+                        if det < 0. { -1. } else { 1. },
                     ],
                     normal_1: [
-                        instance.normal[1][0],
-                        instance.normal[1][1],
-                        instance.normal[1][2],
-                        0.0,
+                        occurrence.normal[1][0],
+                        occurrence.normal[1][1],
+                        occurrence.normal[1][2],
+                        0.,
                     ],
                     normal_2: [
-                        instance.normal[2][0],
-                        instance.normal[2][1],
-                        instance.normal[2][2],
-                        0.0,
+                        occurrence.normal[2][0],
+                        occurrence.normal[2][1],
+                        occurrence.normal[2][2],
+                        0.,
                     ],
                 });
+                p.local_aabbs.push(GpuLocalAabb {
+                    min: [
+                        mesh.local_aabb.min[0],
+                        mesh.local_aabb.min[1],
+                        mesh.local_aabb.min[2],
+                        0.,
+                    ],
+                    max: [
+                        mesh.local_aabb.max[0],
+                        mesh.local_aabb.max[1],
+                        mesh.local_aabb.max[2],
+                        0.,
+                    ],
+                });
+                p.instance_types.push(occurrence.instance_type.words);
                 let base_vertex =
                     i32::try_from(vertex_start).map_err(|_| "base vertex exceeds i32")?;
-                let end = index_start
-                    .checked_add(mesh.geometry.index_count)
-                    .ok_or("draw index range overflow")?;
-                let effective_visible = mesh.flags.contains(RenderFlags::VISIBLE)
-                    && instance.flags.contains(RenderFlags::VISIBLE);
-                plan.local_aabbs.push(GpuLocalAabb {
-                    min: [mesh.aabb.min[0], mesh.aabb.min[1], mesh.aabb.min[2], 0.],
-                    max: [mesh.aabb.max[0], mesh.aabb.max[1], mesh.aabb.max[2], 0.],
-                });
-                plan.effective_visibility.push(effective_visible as u32);
-                plan.draw_metadata.push(DrawSlotMetadata {
+                p.draw_metadata.push(DrawSlotMetadata {
                     index_count: mesh.geometry.index_count,
-                    first_index: index_start,
+                    first_index,
                     base_vertex,
-                    instance_index: instance_start,
+                    instance_index,
                 });
-                plan.commands.push(DrawIndexedIndirect {
-                    index_count: mesh.geometry.index_count,
-                    instance_count: 0,
-                    first_index: index_start,
-                    base_vertex,
-                    first_instance: 0,
-                });
-                plan.draws.push(DrawItem {
+                p.draws.push(DrawItem {
                     pipeline: mesh.pipeline,
                     material: mesh.material,
                     mesh: mesh.handle,
-                    indices: index_start..end,
+                    indices: first_index
+                        ..first_index
+                            .checked_add(mesh.geometry.index_count)
+                            .ok_or("draw range overflow")?,
                     base_vertex,
-                    instances: instance_start..instance_start + 1,
-                    effective_visible,
+                    instances: instance_index..instance_index + 1,
+                    instance_type: occurrence.instance_type,
                 });
             }
         }
-        Ok(plan)
+        Ok(p)
     }
 }
 
+#[derive(Default)]
+pub struct BufferSlot {
+    pub buffer: Option<wgpu::Buffer>,
+    capacity: u64,
+}
+
+#[derive(Default)]
+pub struct GpuSceneCache {
+    revision: Option<u64>,
+    pub buffer_epoch: u64,
+    pub positions: BufferSlot,
+    pub normals: BufferSlot,
+    pub uvs: BufferSlot,
+    pub tangents: BufferSlot,
+    pub indices: BufferSlot,
+    pub instances: BufferSlot,
+    pub local_aabbs: BufferSlot,
+    pub instance_types: BufferSlot,
+    pub draw_metadata: BufferSlot,
+    pub draws: Vec<DrawItem>,
+}
+
 pub fn required_buffer_capacity(
     current: u64,
     required: u64,
@@ -229,92 +210,19 @@ pub fn required_buffer_capacity(
     if required == 0 || current >= required {
         return Ok(current);
     }
-    let grown = current
+    Ok(current
         .checked_mul(2)
         .ok_or("buffer capacity overflow")?
         .max(1)
-        .max(required);
-    Ok(grown.min(maximum))
+        .max(required)
+        .min(maximum))
 }
 
-pub fn vertex_layouts() -> [wgpu::VertexBufferLayout<'static>; 5] {
-    const INSTANCE_ATTRIBUTES: [wgpu::VertexAttribute; 7] = wgpu::vertex_attr_array![3 => Float32x4, 4 => Float32x4, 5 => Float32x4, 6 => Float32x4, 7 => Float32x4, 8 => Float32x4, 9 => Float32x4];
-    [
-        wgpu::VertexBufferLayout {
-            array_stride: 12,
-            step_mode: wgpu::VertexStepMode::Vertex,
-            attributes: &wgpu::vertex_attr_array![0 => Float32x3],
-        },
-        wgpu::VertexBufferLayout {
-            array_stride: 12,
-            step_mode: wgpu::VertexStepMode::Vertex,
-            attributes: &wgpu::vertex_attr_array![1 => Float32x3],
-        },
-        wgpu::VertexBufferLayout {
-            array_stride: 8,
-            step_mode: wgpu::VertexStepMode::Vertex,
-            attributes: &wgpu::vertex_attr_array![2 => Float32x2],
-        },
-        wgpu::VertexBufferLayout {
-            array_stride: size_of::<GpuInstance>() as u64,
-            step_mode: wgpu::VertexStepMode::Instance,
-            attributes: &INSTANCE_ATTRIBUTES,
-        },
-        wgpu::VertexBufferLayout {
-            array_stride: 16,
-            step_mode: wgpu::VertexStepMode::Vertex,
-            attributes: &wgpu::vertex_attr_array![10 => Float32x4],
-        },
-    ]
-}
-
-#[derive(Default)]
-pub struct BufferSlot {
-    pub buffer: Option<wgpu::Buffer>,
-    capacity: u64,
-}
-
-#[derive(Default)]
-pub struct GpuSceneCache {
-    revision: Option<u64>,
-    pub positions: BufferSlot,
-    pub normals: BufferSlot,
-    pub uvs: BufferSlot,
-    pub tangents: BufferSlot,
-    pub indices: BufferSlot,
-    pub instances: BufferSlot,
-    pub local_aabbs: BufferSlot,
-    pub effective_visibility: BufferSlot,
-    pub draw_metadata: BufferSlot,
-    pub frustum_flags: BufferSlot,
-    pub indirect_commands: BufferSlot,
-    pub draws: Vec<DrawItem>,
-    compute: Option<CullingCompute>,
-}
-
-struct CullingCompute {
-    params: wgpu::Buffer,
-    bind_group: wgpu::BindGroup,
-    frustum_pipeline: wgpu::ComputePipeline,
-    query_pipeline: wgpu::ComputePipeline,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Pod, Zeroable)]
-struct CullingParams {
-    planes: [[f32; 4]; 6],
-    count: u32,
-    visible_predicate: u32,
-    frustum_predicate: u32,
-    _pad: u32,
-}
-
-fn predicate_code(value: crate::render_graph::RuntimePredicate) -> u32 {
-    match value {
-        crate::render_graph::RuntimePredicate::Any => 0,
-        crate::render_graph::RuntimePredicate::RequiredTrue => 1,
-        crate::render_graph::RuntimePredicate::RequiredFalse => 2,
-        crate::render_graph::RuntimePredicate::Never => 3,
+fn logical_or_zero<'a, T: Pod>(values: &'a [T], zero: &'a T) -> &'a [u8] {
+    if values.is_empty() {
+        bytemuck::bytes_of(zero)
+    } else {
+        bytemuck::cast_slice(values)
     }
 }
 
@@ -324,73 +232,40 @@ impl GpuSceneCache {
         device: &wgpu::Device,
         queue: &wgpu::Queue,
         data: &SceneFramePlan,
-    ) -> Result<(), String> {
-        self.upload_with_query(
-            device,
-            queue,
-            data,
-            crate::render_graph::MeshQueryRuntimeKey {
-                visible: crate::render_graph::RuntimePredicate::RequiredTrue,
-                frustum_culled: crate::render_graph::RuntimePredicate::Any,
-            },
-        )
-    }
-
-    pub fn upload_with_query(
-        &mut self,
-        device: &wgpu::Device,
-        queue: &wgpu::Queue,
-        data: &SceneFramePlan,
-        query: crate::render_graph::MeshQueryRuntimeKey,
     ) -> Result<(), String> {
         if self.revision == Some(data.revision) {
             return Ok(());
         }
-        let plan = GpuScenePlan::build_with_query(data, query).map_err(str::to_owned)?;
-        if plan.draws.is_empty() {
-            self.draws.clear();
-            self.revision = Some(data.revision);
-            return Ok(());
-        }
-        let maximum = device.limits().max_buffer_size;
-        fn bytes<T>(values: &[T]) -> Result<u64, String> {
-            u64::try_from(values.len())
-                .map_err(|_| "buffer length overflow".to_owned())?
+        let p = GpuScenePlan::build(data).map_err(str::to_owned)?;
+        let max = device.limits().max_buffer_size;
+        fn bytes<T>(v: &[T]) -> Result<u64, String> {
+            (v.len() as u64)
                 .checked_mul(size_of::<T>() as u64)
-                .ok_or_else(|| "buffer byte size overflow".to_owned())
+                .ok_or("buffer byte size overflow".into())
         }
         let required = [
-            bytes(&plan.positions)?,
-            bytes(&plan.normals)?,
-            bytes(&plan.uvs)?,
-            bytes(&plan.tangents)?,
-            bytes(&plan.indices)?,
-            bytes(&plan.instances)?,
-            bytes(&plan.local_aabbs)?,
-            bytes(&plan.effective_visibility)?,
-            bytes(&plan.draw_metadata)?,
-            bytes(&plan.effective_visibility)?,
-            bytes(&plan.commands)?,
+            bytes(&p.positions)?,
+            bytes(&p.normals)?,
+            bytes(&p.uvs)?,
+            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 old = [
-            self.positions.capacity,
-            self.normals.capacity,
-            self.uvs.capacity,
-            self.tangents.capacity,
-            self.indices.capacity,
-            self.instances.capacity,
-            self.local_aabbs.capacity,
-            self.effective_visibility.capacity,
-            self.draw_metadata.capacity,
-            self.frustum_flags.capacity,
-            self.indirect_commands.capacity,
+        let slots = [
+            &mut self.positions,
+            &mut self.normals,
+            &mut self.uvs,
+            &mut self.tangents,
+            &mut self.indices,
+            &mut self.instances,
+            &mut self.local_aabbs,
+            &mut self.instance_types,
+            &mut self.draw_metadata,
         ];
-        let mut capacities = [0; 11];
-        for i in 0..11 {
-            capacities[i] =
-                required_buffer_capacity(old[i], required[i], maximum).map_err(str::to_owned)?;
-        }
-        let usages = [
+        let usage = [
             wgpu::BufferUsages::VERTEX,
             wgpu::BufferUsages::VERTEX,
             wgpu::BufferUsages::VERTEX,
@@ -400,64 +275,35 @@ impl GpuSceneCache {
             wgpu::BufferUsages::STORAGE,
             wgpu::BufferUsages::STORAGE,
             wgpu::BufferUsages::STORAGE,
-            wgpu::BufferUsages::STORAGE,
-            wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::INDIRECT,
         ];
-        let labels = [
-            "scene positions",
-            "scene normals",
-            "scene uvs",
-            "scene tangents",
-            "scene indices",
-            "scene instances",
-            "scene local aabbs",
-            "scene effective visibility",
-            "scene draw metadata",
-            "scene frustum flags",
-            "scene indirect commands",
-        ];
-        let mut replacements: [Option<wgpu::Buffer>; 11] = Default::default();
-        for i in 0..11 {
-            if capacities[i] != old[i] {
-                replacements[i] = Some(device.create_buffer(&wgpu::BufferDescriptor {
-                    label: Some(labels[i]),
-                    size: capacities[i],
-                    usage: usages[i] | wgpu::BufferUsages::COPY_DST,
+        let mut replaced = false;
+        for ((slot, &need), use_) in slots.into_iter().zip(&required).zip(usage) {
+            let cap = required_buffer_capacity(slot.capacity, need, max).map_err(str::to_owned)?;
+            if cap != slot.capacity {
+                slot.buffer = Some(device.create_buffer(&wgpu::BufferDescriptor {
+                    label: Some("scene buffer"),
+                    size: cap,
+                    usage: use_ | wgpu::BufferUsages::COPY_DST,
                     mapped_at_creation: false,
                 }));
+                slot.capacity = cap;
+                replaced = true
             }
         }
-        let slots = [
-            &mut self.positions,
-            &mut self.normals,
-            &mut self.uvs,
-            &mut self.tangents,
-            &mut self.indices,
-            &mut self.instances,
-            &mut self.local_aabbs,
-            &mut self.effective_visibility,
-            &mut self.draw_metadata,
-            &mut self.frustum_flags,
-            &mut self.indirect_commands,
-        ];
-        for (i, slot) in slots.into_iter().enumerate() {
-            if let Some(buffer) = replacements[i].take() {
-                slot.buffer = Some(buffer);
-                slot.capacity = capacities[i];
-            }
-        }
-        let contents = [
-            bytemuck::cast_slice(&plan.positions),
-            bytemuck::cast_slice(&plan.normals),
-            bytemuck::cast_slice(&plan.uvs),
-            bytemuck::cast_slice(&plan.tangents),
-            bytemuck::cast_slice(&plan.indices),
-            bytemuck::cast_slice(&plan.instances),
-            bytemuck::cast_slice(&plan.local_aabbs),
-            bytemuck::cast_slice(&plan.effective_visibility),
-            bytemuck::cast_slice(&plan.draw_metadata),
-            bytemuck::cast_slice(&plan.effective_visibility),
-            bytemuck::cast_slice(&plan.commands),
+        let zero_instance = GpuInstance::zeroed();
+        let zero_aabb = GpuLocalAabb::zeroed();
+        let zero_type = [0u32; 16];
+        let zero_metadata = DrawSlotMetadata::zeroed();
+        let contents: [&[u8]; 9] = [
+            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,
@@ -467,355 +313,89 @@ impl GpuSceneCache {
             &self.indices,
             &self.instances,
             &self.local_aabbs,
-            &self.effective_visibility,
+            &self.instance_types,
             &self.draw_metadata,
-            &self.frustum_flags,
-            &self.indirect_commands,
         ];
-        for (slot, contents) in slots.into_iter().zip(contents) {
-            if !contents.is_empty() {
-                queue.write_buffer(
-                    slot.buffer.as_ref().expect("nonempty slot allocated"),
-                    0,
-                    contents,
-                );
+        for (s, c) in slots.into_iter().zip(contents) {
+            if !c.is_empty() {
+                queue.write_buffer(s.buffer.as_ref().unwrap(), 0, c)
             }
         }
-        self.draws = plan.draws;
-        self.rebuild_compute(device)?;
+        if replaced {
+            self.buffer_epoch = self.buffer_epoch.wrapping_add(1).max(1)
+        }
+        self.draws = p.draws;
         self.revision = Some(data.revision);
         Ok(())
     }
+}
 
-    fn rebuild_compute(&mut self, device: &wgpu::Device) -> Result<(), String> {
-        if self.draws.is_empty() {
-            self.compute = None;
-            return Ok(());
-        }
-        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("scene culling compute"),
-            source: wgpu::ShaderSource::Wgsl(include_str!("culling.wgsl").into()),
-        });
-        let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("scene culling layout"),
-            entries: &(0..7)
-                .map(|binding| wgpu::BindGroupLayoutEntry {
-                    binding,
-                    visibility: wgpu::ShaderStages::COMPUTE,
-                    ty: wgpu::BindingType::Buffer {
-                        ty: if binding == 0 {
-                            wgpu::BufferBindingType::Uniform
-                        } else {
-                            wgpu::BufferBindingType::Storage {
-                                read_only: binding < 5,
-                            }
-                        },
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                    count: None,
-                })
-                .collect::<Vec<_>>(),
-        });
-        let params = device.create_buffer(&wgpu::BufferDescriptor {
-            label: Some("scene culling params"),
-            size: size_of::<CullingParams>() as u64,
-            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
-            mapped_at_creation: false,
-        });
-        let buffers = [
-            &self.instances,
-            &self.local_aabbs,
-            &self.effective_visibility,
-            &self.draw_metadata,
-            &self.frustum_flags,
-            &self.indirect_commands,
-        ];
-        let mut entries = vec![wgpu::BindGroupEntry {
-            binding: 0,
-            resource: params.as_entire_binding(),
-        }];
-        for (i, slot) in buffers.iter().enumerate() {
-            entries.push(wgpu::BindGroupEntry {
-                binding: i as u32 + 1,
-                resource: slot
-                    .buffer
-                    .as_ref()
-                    .ok_or("culling buffer absent")?
-                    .as_entire_binding(),
-            });
-        }
-        let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            label: Some("scene culling group"),
-            layout: &layout,
-            entries: &entries,
-        });
-        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("scene culling pipeline layout"),
-            bind_group_layouts: &[&layout],
-            push_constant_ranges: &[],
-        });
-        let pipeline = |entry| {
-            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-                label: Some(entry),
-                layout: Some(&pipeline_layout),
-                module: &shader,
-                entry_point: Some(entry),
-                compilation_options: Default::default(),
-                cache: None,
-            })
-        };
-        self.compute = Some(CullingCompute {
-            params,
-            bind_group,
-            frustum_pipeline: pipeline("frustum_cull"),
-            query_pipeline: pipeline("mesh_query"),
-        });
-        Ok(())
-    }
-
-    pub fn write_culling_params(
-        &self,
-        queue: &wgpu::Queue,
-        planes: Option<[[f32; 4]; 6]>,
-        query: crate::render_graph::MeshQueryRuntimeKey,
-    ) {
-        if let Some(compute) = &self.compute {
-            if let Some(planes) = planes {
-                queue.write_buffer(&compute.params, 0, bytemuck::bytes_of(&planes));
-            }
-            let tail = CullingParams {
-                planes: [[0.0; 4]; 6],
-                count: self.draws.len() as u32,
-                visible_predicate: predicate_code(query.visible),
-                frustum_predicate: predicate_code(query.frustum_culled),
-                _pad: 0,
-            };
-            queue.write_buffer(
-                &compute.params,
-                std::mem::offset_of!(CullingParams, count) as u64,
-                &bytemuck::bytes_of(&tail)[std::mem::offset_of!(CullingParams, count)..],
-            );
-        }
-    }
-
-    pub(crate) fn encode_frustum_cull(
-        &self,
-        encoder: &mut wgpu::CommandEncoder,
-        profile: Option<&mut crate::renderer::profiler::ProfileFrame>,
-        profile_id: &str,
-    ) {
-        let Some(c) = &self.compute else { return };
-        let timestamps = profile.and_then(|p| p.compute_writes(profile_id));
-        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-            label: Some("frustum_cull"),
-            timestamp_writes: timestamps,
-        });
-        pass.set_pipeline(&c.frustum_pipeline);
-        pass.set_bind_group(0, &c.bind_group, &[]);
-        pass.dispatch_workgroups((self.draws.len() as u32 + 63) / 64, 1, 1);
-    }
-    pub(crate) fn encode_mesh_query(
-        &self,
-        encoder: &mut wgpu::CommandEncoder,
-        profile: Option<&mut crate::renderer::profiler::ProfileFrame>,
-        profile_id: &str,
-    ) {
-        let Some(c) = &self.compute else { return };
-        let timestamps = profile.and_then(|p| p.compute_writes(profile_id));
-        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
-            label: Some("mesh_query"),
-            timestamp_writes: timestamps,
-        });
-        pass.set_pipeline(&c.query_pipeline);
-        pass.set_bind_group(0, &c.bind_group, &[]);
-        pass.dispatch_workgroups((self.draws.len() as u32 + 63) / 64, 1, 1);
-    }
+pub fn vertex_layouts() -> [wgpu::VertexBufferLayout<'static>; 5] {
+    const IA: [wgpu::VertexAttribute; 7] = wgpu::vertex_attr_array![3=>Float32x4,4=>Float32x4,5=>Float32x4,6=>Float32x4,7=>Float32x4,8=>Float32x4,9=>Float32x4];
+    [
+        wgpu::VertexBufferLayout {
+            array_stride: 12,
+            step_mode: wgpu::VertexStepMode::Vertex,
+            attributes: &wgpu::vertex_attr_array![0=>Float32x3],
+        },
+        wgpu::VertexBufferLayout {
+            array_stride: 12,
+            step_mode: wgpu::VertexStepMode::Vertex,
+            attributes: &wgpu::vertex_attr_array![1=>Float32x3],
+        },
+        wgpu::VertexBufferLayout {
+            array_stride: 8,
+            step_mode: wgpu::VertexStepMode::Vertex,
+            attributes: &wgpu::vertex_attr_array![2=>Float32x2],
+        },
+        wgpu::VertexBufferLayout {
+            array_stride: 112,
+            step_mode: wgpu::VertexStepMode::Instance,
+            attributes: &IA,
+        },
+        wgpu::VertexBufferLayout {
+            array_stride: 16,
+            step_mode: wgpu::VertexStepMode::Vertex,
+            attributes: &wgpu::vertex_attr_array![10=>Float32x4],
+        },
+    ]
 }
 
 #[cfg(test)]
 mod tests {
     use super::*;
-
     #[test]
-    fn mesh_query_source_guards_optional_flag_buffer_reads() {
-        let source = include_str!("culling.wgsl");
-        let visible_never = source.find("if (params.visible_predicate == 3u)").unwrap();
-        let visible_guard = source
-            .find("else if (params.visible_predicate != 0u)")
-            .unwrap();
-        let visible_load = source.find("matches(authored_visible[i]").unwrap();
-        let frustum_never = source
-            .find("if (selected && params.frustum_predicate == 3u)")
-            .unwrap();
-        let frustum_guard = source
-            .find("else if (selected && params.frustum_predicate != 0u)")
-            .unwrap();
-        let frustum_load = source.find("matches(frustum_flags[i]").unwrap();
-        assert!(visible_never < visible_guard && visible_guard < visible_load);
-        assert!(frustum_never < frustum_guard && frustum_guard < frustum_load);
-        assert!(
-            visible_load < frustum_load,
-            "visible rejection must precede the frustum load"
-        );
-        assert!(
-            source.contains("predicate == 0u ||"),
-            "predicate zero is handled without a load by the guards"
-        );
-        for binding in 0..=6 {
-            assert!(source.contains(&format!("@binding({binding})")));
-        }
-    }
-
-    #[test]
-    fn effective_visibility_handles_every_mesh_instance_combination() {
-        use crate::render_graph::TriStatePredicate::{Any, RequiredFalse, RequiredTrue};
-        for (mesh, instance, effective) in [
-            (false, false, false),
-            (false, true, false),
-            (true, false, false),
-            (true, true, true),
-        ] {
-            let flags = |visible| {
-                if visible {
-                    RenderFlags::VISIBLE
-                } else {
-                    RenderFlags::NONE
-                }
-            };
-            assert!(visibility_matches(Any, flags(mesh), flags(instance)));
-            assert_eq!(
-                visibility_matches(RequiredTrue, flags(mesh), flags(instance)),
-                effective
-            );
-            assert_eq!(
-                visibility_matches(RequiredFalse, flags(mesh), flags(instance)),
-                !effective
-            );
-        }
-    }
-    use crate::render_data::{
-        MeshCreateInfo, RenderData, RenderDataConfig, IDENTITY_MODEL_TRANSFORM,
-    };
-    #[test]
-    fn instance_is_112_bytes_and_padding_is_zero() {
+    fn abi() {
         assert_eq!(size_of::<GpuInstance>(), 112);
-        let value = GpuInstance {
-            model: [[1.0; 4]; 4],
-            normal_0: [1., 2., 3., 0.],
-            normal_1: [4., 5., 6., 0.],
-            normal_2: [7., 8., 9., 0.],
-        };
-        assert_eq!(value.normal_2[3], 0.0);
-    }
-    #[test]
-    fn capacity_grows_and_checks_limit() {
-        assert_eq!(required_buffer_capacity(8, 9, 32), Ok(16));
-        assert!(required_buffer_capacity(0, 33, 32).is_err());
-    }
-    #[test]
-    fn mirrored_model_stores_negative_determinant_sign_in_padding() {
-        let mut data = RenderData::new(RenderDataConfig::default()).unwrap();
-        let mut mirrored = IDENTITY_MODEL_TRANSFORM;
-        mirrored[0][0] = -1.0;
-        data.create_mesh(MeshCreateInfo {
-            positions: &[[0., 0., 0.], [1., 0., 0.], [0., 1., 0.]],
-            normals: &[[0., 0., 1.]; 3],
-            tangents: &[[1., 0., 0., 1.]; 3],
-            uvs: &[[0., 0.]; 3],
-            indices: &[0, 1, 2],
-            pipeline: PipelineKey::new(0),
-            material: MaterialKey::DEFAULT,
-            flags: RenderFlags::VISIBLE,
-            default_instance_flags: RenderFlags::VISIBLE,
-            default_transform: mirrored,
-        })
-        .unwrap();
-        let frame = crate::renderer::scene_frame::SceneFramePlan::build(&data).unwrap();
-        let plan = GpuScenePlan::build(&frame).unwrap();
-        assert_eq!(plan.instances[0].normal_0[3], -1.0);
-    }
-    #[test]
-    fn capacity_reuses_and_layout_matches_shader_contract() {
-        assert_eq!(required_buffer_capacity(16, 12, 32), Ok(16));
-        assert_eq!(required_buffer_capacity(0, 1, 32), Ok(1));
-        let layouts = vertex_layouts();
-        assert_eq!(
-            layouts
-                .iter()
-                .map(|layout| layout.array_stride)
-                .collect::<Vec<_>>(),
-            [12, 12, 8, 112, 16]
-        );
-        assert_eq!(
-            layouts[3]
-                .attributes
-                .iter()
-                .map(|attribute| attribute.shader_location)
-                .collect::<Vec<_>>(),
-            vec![3, 4, 5, 6, 7, 8, 9]
-        );
-    }
-    #[test]
-    fn plan_is_canonical_and_predicate_independent() {
-        let mut data = RenderData::new(RenderDataConfig {
-            initial_vertices: 0,
-            initial_indices: 0,
-            initial_meshes: 0,
-            initial_instances: 0,
-            ..Default::default()
-        })
-        .unwrap();
-        let p = [[0., 0., 0.], [1., 0., 0.], [0., 1., 0.]];
-        let n = [[0., 0., 1.]; 3];
-        let u = [[0., 0.]; 3];
-        let i = [0, 1, 2];
-        let mut add = |pipeline, visible| {
-            data.create_mesh(MeshCreateInfo {
-                positions: &p,
-                normals: &n,
-                tangents: &[[1., 0., 0., 1.]; 3],
-                uvs: &u,
-                indices: &i,
-                pipeline: PipelineKey::new(pipeline),
-                material: crate::render_data::MaterialKey::DEFAULT,
-                flags: RenderFlags::VISIBLE,
-                default_instance_flags: if visible {
-                    RenderFlags::VISIBLE
-                } else {
-                    RenderFlags::NONE
-                },
-                default_transform: IDENTITY_MODEL_TRANSFORM,
-            })
-            .unwrap()
-        };
-        let high = add(9, true);
-        let _hidden = add(0, false);
-        let low = add(2, true);
-        data.create_instance(low.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::VISIBLE)
-            .unwrap();
-        let frame = crate::renderer::scene_frame::SceneFramePlan::build(&data).unwrap();
-        let plan = GpuScenePlan::build(&frame).unwrap();
-        assert_eq!(
-            plan.draws
-                .iter()
-                .map(|d| d.pipeline.get())
-                .collect::<Vec<_>>(),
-            vec![0, 2, 2, 9]
-        );
-        assert_eq!(plan.draws[0].base_vertex, 0);
-        assert!(!plan.draws[0].effective_visible);
-        assert_eq!(plan.draws[1].base_vertex, 3);
-        assert_eq!(plan.draws[1].instances, 1..2);
-        assert_eq!(plan.draws[2].instances, 2..3);
-        assert_eq!(plan.indices, [0, 1, 2, 0, 1, 2, 0, 1, 2]);
-        assert_eq!(high.mesh, plan.draws[3].mesh);
         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);
-        assert!(plan
-            .commands
-            .iter()
-            .all(|command| command.first_instance == 0));
+        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() {
+        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_traversal.rs

@@ -0,0 +1,427 @@
+//! Graph-owned instance predicate compute support.
+
+use crate::render_graph::{
+    BooleanBinaryOp, 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::BooleanBinary {
+                operation,
+                left,
+                right,
+            } => format!(
+                "({} {} {})",
+                x(*left),
+                match operation {
+                    BooleanBinaryOp::And => "&&",
+                    BooleanBinaryOp::Or => "||",
+                    BooleanBinaryOp::Xor => "!=",
+                    BooleanBinaryOp::Xnor => "==",
+                },
+                x(*right)
+            ),
+            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 } => bits
+                .iter()
+                .enumerate()
+                .map(|(bit, id)| format!("select(0u,{}u,{})", 1u32 << bit, x(*id)))
+                .collect::<Vec<_>>()
+                .join("|")
+                .pipe(|terms| format!("({terms})")),
+            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(",")))
+}
+trait Pipe: Sized {
+    fn pipe<R>(self, f: impl FnOnce(Self) -> R) -> R {
+        f(self)
+    }
+}
+impl<T> Pipe for T {}
+
+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::*;
+
+    #[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>");
+    }
+}

renderer/src/renderer/mod.rs

@@ -11,12 +11,13 @@ use crate::{
     command_ring::CommandRing,
     gltf::{install_imported, ModelBounds},
     message::{camera_drag, CameraDrag, DrainEventError, MouseMessage, ResizeMessage, WindowEvent},
-    render_data::{InstanceHandle, MeshHandle, RenderData, RenderDataConfig, RenderFlags},
+    render_data::{InstanceHandle, InstanceType, MeshHandle, RenderData, RenderDataConfig},
     renderer::scene::Scene,
 };
 
 pub mod executors;
 pub mod gpu_scene;
+pub mod instance_traversal;
 pub mod material;
 pub mod pipeline_library;
 pub mod profiler;
@@ -27,6 +28,43 @@ pub use pipeline_library::PipelineLibrary;
 
 const DEPTH_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float;
 
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+struct DeviceFeaturePlan {
+    hard: wgpu::Features,
+    initial: wgpu::Features,
+    profiling_enabled: bool,
+}
+
+fn device_feature_plan(profile: bool, supported: wgpu::Features) -> DeviceFeaturePlan {
+    let hard = wgpu::Features::INDIRECT_FIRST_INSTANCE;
+    let profiling = profiler::Profiler::requested_features(profile, supported);
+    DeviceFeaturePlan {
+        hard,
+        initial: hard | profiling,
+        profiling_enabled: !profiling.is_empty(),
+    }
+}
+
+#[cfg(test)]
+mod device_feature_tests {
+    use super::*;
+
+    #[test]
+    fn retry_plan_never_drops_hard_features() {
+        let hard_only = device_feature_plan(false, wgpu::Features::INDIRECT_FIRST_INSTANCE);
+        assert_eq!(hard_only.initial, hard_only.hard);
+        assert!(!hard_only.profiling_enabled);
+
+        let profiled = device_feature_plan(
+            true,
+            wgpu::Features::INDIRECT_FIRST_INSTANCE | wgpu::Features::TIMESTAMP_QUERY,
+        );
+        assert!(profiled.initial.contains(profiled.hard));
+        assert!(profiled.initial.contains(wgpu::Features::TIMESTAMP_QUERY));
+        assert!(profiled.profiling_enabled);
+    }
+}
+
 #[repr(C, align(16))]
 #[derive(Clone, Copy, Debug, PartialEq, bytemuck::Pod, bytemuck::Zeroable)]
 struct FullscreenUniforms {
@@ -748,12 +786,10 @@ struct GpuTextureSlot {
 }
 
 enum PreparedExecution {
-    FrustumCull,
-    MeshQuery,
-    PipelineRegistry,
     Pipeline {
         execution: usize,
         base: crate::render_data::PipelineKey,
+        predicate_ordinal: u32,
         variant: wgpu::RenderPipeline,
     },
     Fullscreen {
@@ -777,7 +813,7 @@ struct ActiveCompiledGraph {
 #[derive(Clone, Copy)]
 enum UploadGraph {
     Immediate,
-    Compiled(crate::render_graph::MeshQueryRuntimeKey),
+    Compiled(bool),
 }
 
 #[derive(Clone, Copy, Debug, PartialEq, Eq)]
@@ -832,27 +868,27 @@ fn acquisition_action(source: FrameTargetSource, error: &wgpu::SurfaceError) ->
 }
 
 fn classify_upload_graph(graph: &ActiveCompiledGraph) -> UploadGraph {
-    UploadGraph::Compiled(graph.runtime.allocations.query)
+    UploadGraph::Compiled(
+        graph
+            .runtime
+            .instance_traversal
+            .as_ref()
+            .is_some_and(|p| p.requires_camera),
+    )
 }
 
-fn upload_query_for_render(
-    pending: Option<UploadGraph>,
-    active: Option<UploadGraph>,
-) -> Option<crate::render_graph::MeshQueryRuntimeKey> {
+fn upload_query_for_render(pending: Option<UploadGraph>, active: Option<UploadGraph>) -> bool {
     match pending.or(active) {
-        Some(UploadGraph::Compiled(query)) => Some(query),
-        Some(UploadGraph::Immediate) | None => None,
+        Some(UploadGraph::Compiled(value)) => value,
+        Some(UploadGraph::Immediate) | None => false,
     }
 }
 
 fn resolve_culling_frustum(
-    query: crate::render_graph::MeshQueryRuntimeKey,
+    required: bool,
     read: impl FnOnce() -> Option<Result<[[f32; 4]; 6], crate::camera::FrustumError>>,
 ) -> Result<Option<[[f32; 4]; 6]>, crate::render_graph::GraphError> {
-    if matches!(
-        query.frustum_culled,
-        crate::render_graph::RuntimePredicate::Any | crate::render_graph::RuntimePredicate::Never
-    ) {
+    if !required {
         return Ok(None);
     }
     match read() {
@@ -871,13 +907,10 @@ fn resolve_culling_frustum(
 fn update_validate_write_scene<S: scene::Scene>(
     scene: &mut S,
     queue: &wgpu::Queue,
-    query: Option<crate::render_graph::MeshQueryRuntimeKey>,
+    requires_camera: bool,
 ) -> Result<Option<[[f32; 4]; 6]>, crate::render_graph::GraphError> {
     scene.update_cpu();
-    let planes = match query {
-        Some(query) => resolve_culling_frustum(query, || scene.frustum_planes())?,
-        None => None,
-    };
+    let planes = resolve_culling_frustum(requires_camera, || scene.frustum_planes())?;
     scene.write_uniforms(queue);
     Ok(planes)
 }
@@ -935,7 +968,7 @@ fn resolve_switch_request(
         1 => {
             let id = crate::render_graph::CompiledGraphId { slot, generation };
             // Resolve the registry entry here, before any GPU preparation or pending
-            // state mutation. Registry::get is also the Phase 4 activation gate.
+            // state mutation. Registry::get is also the compiled-graph availability gate.
             registry.get(id)?;
             Ok(ResolvedSwitchRequest::Compiled(id))
         }
@@ -980,42 +1013,29 @@ mod switch_request_tests {
 
     use super::*;
 
-    fn query(visible: crate::render_graph::RuntimePredicate) -> UploadGraph {
-        UploadGraph::Compiled(crate::render_graph::MeshQueryRuntimeKey {
-            visible,
-            frustum_culled: crate::render_graph::RuntimePredicate::Any,
-        })
+    fn graph(requires_camera: bool) -> UploadGraph {
+        UploadGraph::Compiled(requires_camera)
     }
 
     #[test]
     fn upload_selection_follows_the_graph_rendered_for_the_commit_frame() {
-        use crate::render_graph::RuntimePredicate::{Any, RequiredFalse, RequiredTrue};
-        let selected =
-            |pending, active| upload_query_for_render(pending, active).map(|query| query.visible);
+        let selected = upload_query_for_render;
+        assert!(!selected(Some(graph(false)), Some(graph(true))));
         assert_eq!(
-            selected(Some(query(RequiredFalse)), Some(query(RequiredTrue))),
-            Some(RequiredFalse)
+            selected(Some(UploadGraph::Immediate), Some(graph(true))),
+            false
         );
         assert_eq!(
-            selected(Some(UploadGraph::Immediate), Some(query(RequiredTrue))),
-            None
+            selected(Some(UploadGraph::Immediate), Some(graph(true))),
+            false
         );
-        assert_eq!(
-            selected(Some(UploadGraph::Immediate), Some(query(RequiredTrue))),
-            None
-        );
-        assert_eq!(
-            selected(None, Some(query(RequiredTrue))),
-            Some(RequiredTrue)
-        );
-        assert_eq!(selected(None, Some(UploadGraph::Immediate)), None);
-        assert_eq!(selected(None, None), None);
-        assert_eq!(selected(Some(query(Any)), None), Some(Any));
+        assert!(selected(None, Some(graph(true))));
+        assert!(!selected(None, Some(UploadGraph::Immediate)));
+        assert!(!selected(None, None));
     }
 
     #[test]
     fn frame_target_precedence_and_pending_resize_upload_are_exact() {
-        use crate::render_graph::RuntimePredicate::{RequiredFalse, RequiredTrue};
         assert_eq!(
             select_frame_target_source(true, true, true),
             FrameTargetSource::PendingSwitch
@@ -1032,14 +1052,10 @@ mod switch_request_tests {
             select_frame_target_source(false, false, false),
             FrameTargetSource::Immediate
         );
-        let pending_resize = query(RequiredFalse);
-        let active = query(RequiredTrue);
-        assert_eq!(
-            upload_query_for_render(Some(pending_resize), Some(active))
-                .unwrap()
-                .visible,
-            RequiredFalse
-        );
+        assert!(!upload_query_for_render(
+            Some(graph(false)),
+            Some(graph(true))
+        ));
     }
 
     #[test]
@@ -1063,15 +1079,10 @@ mod switch_request_tests {
     }
 
     #[test]
-    fn frustum_preflight_skips_any_and_distinguishes_missing_from_invalid() {
-        use crate::render_graph::RuntimePredicate::{Any, RequiredFalse, RequiredTrue};
-        let query = |frustum_culled| crate::render_graph::MeshQueryRuntimeKey {
-            visible: RequiredTrue,
-            frustum_culled,
-        };
+    fn frustum_preflight_uses_boolean_traversal_requirement() {
         let mut reads = 0;
         assert_eq!(
-            resolve_culling_frustum(query(Any), || {
+            resolve_culling_frustum(false, || {
                 reads += 1;
                 None
             })
@@ -1082,9 +1093,9 @@ mod switch_request_tests {
             reads, 0,
             "inactive frustum filtering must not read the camera"
         );
-        let missing = resolve_culling_frustum(query(RequiredFalse), || None).unwrap_err();
+        let missing = resolve_culling_frustum(true, || None).unwrap_err();
         assert!(missing.message.contains("no camera"));
-        let invalid = resolve_culling_frustum(query(RequiredFalse), || {
+        let invalid = resolve_culling_frustum(true, || {
             Some(Err(crate::camera::FrustumError::Degenerate { plane: 2 }))
         })
         .unwrap_err();
@@ -1265,10 +1276,11 @@ pub struct Renderer<T: scene::Scene> {
     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>,
-    gpu_error: std::sync::Arc<std::sync::atomic::AtomicBool>,
+    gpu_error: std::sync::Arc<std::sync::Mutex<Option<String>>>,
     framing_radius: f32,
     graph_registry: crate::render_graph::Registry,
     active_compiled: Option<ActiveCompiledGraph>,
@@ -1471,19 +1483,39 @@ impl<T: Scene + 'static> Renderer<T> {
                     let result = js_sys::Object::new();
                     let meshes = js_sys::Array::new();
                     for h in installed.meshes {
-                        meshes.push(&js_sys::Array::of2(
-                            &h.slot().into(),
-                            &h.generation().into(),
-                        ));
+                        let item = js_sys::Object::new();
+                        js_sys::Reflect::set(
+                            &item,
+                            &"handle".into(),
+                            &js_sys::Array::of2(&h.slot().into(), &h.generation().into()),
+                        )
+                        .unwrap();
+                        let ty = self
+                            .render_data
+                            .mesh(h)
+                            .unwrap()
+                            .default_instance_type
+                            .words;
+                        js_sys::Reflect::set(
+                            &item,
+                            &"defaultType".into(),
+                            &js_sys::Array::from_iter(ty.into_iter().map(JsValue::from)),
+                        )
+                        .unwrap();
+                        meshes.push(&item);
                     }
                     js_sys::Reflect::set(&result, &"meshes".into(), &meshes).unwrap();
                     Ok(result.into())
                 }
                 2 => {
                     self.render_data
-                        .set_mesh_flags(
+                        .set_mesh_visible(
                             MeshHandle::from_parts(words[2], words[3]),
-                            RenderFlags::from_bits_retain(words[4]),
+                            match words[4] {
+                                0 => false,
+                                1 => true,
+                                _ => return Err("INVALID_VISIBILITY"),
+                            },
                         )
                         .map_err(|e| render_data_error_code(&e))?;
                     Ok(JsValue::UNDEFINED)
@@ -1496,15 +1528,25 @@ impl<T: Scene + 'static> Renderer<T> {
                     }
                     let h = self
                         .render_data
-                        .create_instance(mesh, m, RenderFlags::from_bits_retain(words[20]))
+                        .create_instance(
+                            mesh,
+                            m,
+                            InstanceType {
+                                words: std::array::from_fn(|i| words[20 + i]),
+                            },
+                        )
                         .map_err(|e| render_data_error_code(&e))?;
                     Ok(js_sys::Array::of2(&h.slot().into(), &h.generation().into()).into())
                 }
                 4 => {
                     self.render_data
-                        .set_instance_flags(
+                        .set_instance_visible(
                             InstanceHandle::from_parts(words[2], words[3]),
-                            RenderFlags::from_bits_retain(words[4]),
+                            match words[4] {
+                                0 => false,
+                                1 => true,
+                                _ => return Err("INVALID_VISIBILITY"),
+                            },
                         )
                         .map_err(|e| render_data_error_code(&e))?;
                     Ok(JsValue::UNDEFINED)
@@ -1526,6 +1568,17 @@ impl<T: Scene + 'static> Renderer<T> {
                         .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)
+                }
                 _ => Err("UNKNOWN_OPCODE"),
             })();
             match outcome {
@@ -1765,8 +1818,7 @@ impl<T: Scene + 'static> Renderer<T> {
             let contract = crate::render_graph::contract(&execution.executor.key)
                 .ok_or_else(|| fail("executor contract missing"))?;
             match execution.executor.key.as_str() {
-                "frustum_cull" => executions.push(PreparedExecution::FrustumCull),
-                "mesh_query" => executions.push(PreparedExecution::MeshQuery),
+                "frustum_cull" => continue,
                 _ if execution.executor.key == "frame_out"
                     || contract.fullscreen_policy.is_some() =>
                 {
@@ -1924,7 +1976,6 @@ impl<T: Scene + 'static> Renderer<T> {
                         _uniform: uniform,
                     });
                 }
-                "pipeline_registry" => executions.push(PreparedExecution::PipelineRegistry),
                 "pipeline" => {
                     let ExecutionKind::Render {
                         color_attachments,
@@ -2009,6 +2060,14 @@ impl<T: Scene + 'static> Renderer<T> {
                     executions.push(PreparedExecution::Pipeline {
                         execution: index,
                         base,
+                        predicate_ordinal: runtime
+                            .instance_traversal
+                            .as_ref()
+                            .and_then(|p| {
+                                p.pipelines.iter().find(|v| v.execution as usize == index)
+                            })
+                            .map(|p| p.ordinal)
+                            .ok_or_else(|| fail("pipeline predicate missing"))?,
                         variant,
                     });
                 }
@@ -2049,7 +2108,13 @@ impl<T: Scene + 'static> Renderer<T> {
         // Candidate construction allocates GPU resources, so the live scene preflight
         // belongs here: this is the earliest boundary with both the runtime query and
         // scene access, and precedes GPU work and all pending/in-flight mutation.
-        resolve_culling_frustum(runtime.allocations.query, || self.scene.frustum_planes())?;
+        resolve_culling_frustum(
+            runtime
+                .instance_traversal
+                .as_ref()
+                .is_some_and(|p| p.requires_camera),
+            || self.scene.frustum_planes(),
+        )?;
         let restart_graph = graph.clone();
         self.next_preparation_token = self.next_preparation_token.wrapping_add(1).max(1);
         let token = self.next_preparation_token;
@@ -2153,9 +2218,13 @@ impl<T: Scene + 'static> Renderer<T> {
             .await
             .unwrap();
 
-        let optional_features = profiler::Profiler::requested_features(profile, adapter.features());
+        let feature_plan = device_feature_plan(profile, adapter.features());
+        assert!(
+            adapter.features().contains(feature_plan.hard),
+            "WebGPU adapter lacks required indirect-first-instance support"
+        );
         let descriptor = wgpu::DeviceDescriptor {
-            required_features: optional_features,
+            required_features: feature_plan.initial,
             required_limits: wgpu::Limits::default(),
             label: None,
             memory_hints: wgpu::MemoryHints::default(),
@@ -2164,7 +2233,7 @@ impl<T: Scene + 'static> Renderer<T> {
 
         let (device, queue) = match adapter.request_device(&descriptor).await {
             Ok(result) => result,
-            Err(error) if !optional_features.is_empty() => {
+            Err(error) if feature_plan.profiling_enabled => {
                 log::warn!("timestamp-enabled device request failed, retrying baseline: {error}");
                 adapter = instance
                     .request_adapter(&wgpu::RequestAdapterOptions {
@@ -2174,8 +2243,12 @@ impl<T: Scene + 'static> Renderer<T> {
                     })
                     .await
                     .expect("surface-compatible adapter required for baseline device");
+                assert!(
+                    adapter.features().contains(feature_plan.hard),
+                    "reacquired WebGPU adapter lacks required indirect-first-instance support"
+                );
                 let baseline = wgpu::DeviceDescriptor {
-                    required_features: wgpu::Features::empty(),
+                    required_features: feature_plan.hard,
                     required_limits: wgpu::Limits::default(),
                     label: None,
                     memory_hints: wgpu::MemoryHints::default(),
@@ -2185,15 +2258,23 @@ impl<T: Scene + 'static> Renderer<T> {
             }
             Err(error) => panic!("baseline WebGPU device request failed: {error}"),
         };
+        assert!(
+            device.features().contains(feature_plan.hard),
+            "WebGPU device lacks required indirect-first-instance support"
+        );
         info!("Adapter info: {:?}", adapter.get_info());
         info!("Adapter features: {:?}", adapter.features());
         info!("Adapter limits: {:?}", adapter.limits());
         let profiler = profiler::Profiler::new(profile, &device, &queue).await;
-        let gpu_error = std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false));
+        let gpu_error = std::sync::Arc::new(std::sync::Mutex::new(None));
         let error_flag = gpu_error.clone();
         device.on_uncaptured_error(Box::new(move |error| {
-            error_flag.store(true, std::sync::atomic::Ordering::Relaxed);
-            log::error!("Uncaptured GPU error: {error}");
+            let message = error.to_string();
+            let mut first = error_flag.lock().unwrap();
+            if first.is_none() {
+                *first = Some(message.clone());
+            }
+            log::error!("Uncaptured GPU error: {message}");
         }));
 
         let surface_caps = surface.get_capabilities(&adapter);
@@ -2244,6 +2325,7 @@ impl<T: Scene + 'static> Renderer<T> {
             snapshot_init_sent: false,
             scene_frame: Default::default(),
             gpu_scene: Default::default(),
+            instance_traversal: None,
             materials,
             command_ring: None,
             pending_replies: Vec::new(),
@@ -2266,12 +2348,10 @@ impl<T: Scene + 'static> Renderer<T> {
             return;
         }
         self.drain_preparation_completions();
-        if self
-            .gpu_error
-            .swap(false, std::sync::atomic::Ordering::AcqRel)
-        {
+        let gpu_error = self.gpu_error.lock().unwrap().take();
+        if let Some(error) = gpu_error {
             self.halted = true;
-            self.post_fatal("GPU_VALIDATION_FAILED", "uncaptured WebGPU error");
+            self.post_fatal("GPU_VALIDATION_FAILED", &error);
             return;
         }
         if !self.drain_commands() {
@@ -2295,8 +2375,16 @@ impl<T: Scene + 'static> Renderer<T> {
                 &"controlPtr".into(),
                 &self.snapshot.control_ptr().into(),
             );
-            let _ = js_sys::Reflect::set(&message, &"controlVersion".into(), &1.into());
-            let _ = js_sys::Reflect::set(&message, &"schemaVersion".into(), &1.into());
+            let _ = js_sys::Reflect::set(
+                &message,
+                &"controlVersion".into(),
+                &crate::shared_snapshot::CONTROL_VERSION.into(),
+            );
+            let _ = js_sys::Reflect::set(
+                &message,
+                &"schemaVersion".into(),
+                &crate::shared_snapshot::SCHEMA.into(),
+            );
             let _ = global.post_message(&message);
             self.snapshot_init_sent = true;
         }
@@ -2351,26 +2439,14 @@ impl<T: Scene + 'static> Renderer<T> {
                 return;
             }
         };
-        let upload = if let Some(query) = query {
-            self.gpu_scene.upload_with_query(
-                &self.context.device,
-                &self.context.queue,
-                frame_plan,
-                query,
-            )
-        } else {
-            self.gpu_scene
-                .upload(&self.context.device, &self.context.queue, frame_plan)
-        };
+        let upload = self
+            .gpu_scene
+            .upload(&self.context.device, &self.context.queue, frame_plan);
         if let Err(error) = upload {
             log::error!("GPU scene upload failed: {error}");
             self.post_fatal("GPU_UPLOAD_FAILED", &error);
             return;
         }
-        if let Some(query) = query {
-            self.gpu_scene
-                .write_culling_params(&self.context.queue, planes, query);
-        }
 
         // Candidate publication is transactional: configure its complete contract at
         // the last possible point before acquisition, but retain the known-good
@@ -2463,16 +2539,53 @@ impl<T: Scene + 'static> Renderer<T> {
             ),
         });
         let encode_result = if let Some(active) = rendering_compiled {
-            executors::encode_compiled(
-                &mut encoder,
-                &texture_view,
-                active,
-                &self.scene,
-                &self.gpu_scene,
-                &self.resources,
-                &self.materials,
-                profile_frame.as_mut(),
-            )
+            (|| -> 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,
+                    active,
+                    &self.scene,
+                    &self.gpu_scene,
+                    &self.resources,
+                    &self.materials,
+                    &self.instance_traversal.as_ref().unwrap().commands,
+                    profile_frame.as_mut(),
+                )
+            })()
         } else {
             executors::encode_immediate(
                 &mut encoder,
@@ -2620,12 +2733,7 @@ impl<T: Scene + 'static> Renderer<T> {
                     .unwrap_or(0)
                     .into(),
             ),
-            (
-                "gpuError",
-                self.gpu_error
-                    .load(std::sync::atomic::Ordering::Relaxed)
-                    .into(),
-            ),
+            ("gpuError", self.gpu_error.lock().unwrap().is_some().into()),
         ] {
             let _ = js_sys::Reflect::set(&telemetry, &key.into(), &value);
         }

renderer/src/renderer/pipeline_library.rs

@@ -120,7 +120,7 @@ pub struct PipelineLibrary {
     default_layout: Option<PipelineLayoutKey>,
     material_layout: Option<PipelineLayoutKey>,
     next_layout: u64,
-    pipeline_registry: HashMap<String, (PipelineKey, RenderPipelineKey)>,
+    named_bases: HashMap<String, (PipelineKey, RenderPipelineKey)>,
     descriptor_cache: HashMap<RenderPipelineKey, PipelineKey>,
 }
 
@@ -134,7 +134,7 @@ impl PipelineLibrary {
             default_layout: None,
             material_layout: None,
             next_layout: 0,
-            pipeline_registry: HashMap::new(),
+            named_bases: HashMap::new(),
             descriptor_cache: HashMap::new(),
         }
     }
@@ -328,7 +328,7 @@ impl PipelineLibrary {
     ) -> Result<PipelineKey, String> {
         let spec = self.compatibility_spec(name, layouts, shader, format);
         let descriptor = spec.key();
-        if let Some((_, existing)) = self.pipeline_registry.get(name) {
+        if let Some((_, existing)) = self.named_bases.get(name) {
             return Err(if existing == &descriptor {
                 format!("Pipeline '{name}' already exists")
             } else {
@@ -336,13 +336,12 @@ impl PipelineLibrary {
             });
         }
         let key = self.get_or_create_from_spec(device, &spec, Some(name));
-        self.pipeline_registry
-            .insert(name.to_owned(), (key, descriptor));
+        self.named_bases.insert(name.to_owned(), (key, descriptor));
         Ok(key)
     }
 
     pub fn find_pipeline(&self, name: &str) -> Option<PipelineKey> {
-        self.pipeline_registry.get(name).map(|v| v.0)
+        self.named_bases.get(name).map(|v| v.0)
     }
     pub fn get_or_create_pipeline(
         &mut self,
@@ -353,7 +352,7 @@ impl PipelineLibrary {
         format: wgpu::TextureFormat,
     ) -> PipelineKey {
         let wanted = self.compatibility_spec(name, layouts, shader, format).key();
-        if let Some((key, existing)) = self.pipeline_registry.get(name) {
+        if let Some((key, existing)) = self.named_bases.get(name) {
             assert_eq!(
                 existing, &wanted,
                 "Pipeline '{name}' requested with a different descriptor"

renderer/src/renderer/scene_frame.rs

@@ -3,8 +3,8 @@ use std::collections::HashMap;
 use thiserror::Error;
 
 use crate::render_data::{
-    affine_world_aabb, Aabb, GeometryRange, InstanceHandle, MaterialKey, MeshHandle,
-    ModelTransform, NormalMatrix, PipelineKey, RenderData, RenderFlags,
+    affine_world_aabb, Aabb, GeometryRange, InstanceHandle, InstanceType, MaterialKey, MeshHandle,
+    ModelTransform, NormalMatrix, PipelineKey, RenderData,
 };
 
 #[derive(Clone, Debug)]
@@ -13,8 +13,8 @@ pub struct SceneFrameMesh {
     pub geometry: GeometryRange,
     pub pipeline: PipelineKey,
     pub material: MaterialKey,
-    pub flags: RenderFlags,
-    pub aabb: Aabb,
+    pub instance_type: InstanceType,
+    pub local_aabb: Aabb,
     pub default_instance: InstanceHandle,
     pub occurrence_range: std::ops::Range<usize>,
 }
@@ -26,7 +26,7 @@ pub struct SceneFrameOccurrence {
     pub mesh_index: usize,
     pub model: ModelTransform,
     pub normal: NormalMatrix,
-    pub flags: RenderFlags,
+    pub instance_type: InstanceType,
     pub is_default: bool,
     pub world_aabb: Aabb,
 }
@@ -84,9 +84,9 @@ impl SceneFramePlan {
                 mesh_index,
                 model: occurrence.model,
                 normal: occurrence.normal,
-                flags: occurrence.flags,
+                instance_type: occurrence.instance_type,
                 is_default: handle == mesh.default_instance,
-                world_aabb: affine_world_aabb(mesh.aabb, occurrence.model)
+                world_aabb: affine_world_aabb(mesh.local_aabb, occurrence.model)
                     .map_err(|_| SceneFrameError::InvalidWorldBounds)?,
             });
         }
@@ -117,8 +117,8 @@ impl SceneFramePlan {
                 geometry: mesh.geometry,
                 pipeline: mesh.pipeline,
                 material: mesh.material,
-                flags: mesh.flags,
-                aabb: mesh.aabb,
+                instance_type: mesh.default_instance_type,
+                local_aabb: mesh.local_aabb,
                 default_instance: mesh.default_instance,
                 occurrence_range: offsets[dense]..offsets[dense + 1],
             })
@@ -170,12 +170,11 @@ mod tests {
             indices: &[0, 1, 2],
             pipeline: PipelineKey::new(0),
             material: crate::render_data::MaterialKey::DEFAULT,
-            flags: if visible {
-                RenderFlags::VISIBLE
+            default_instance_type: if visible {
+                InstanceType::VISIBLE
             } else {
-                RenderFlags::NONE
+                InstanceType::ZERO
             },
-            default_instance_flags: RenderFlags::VISIBLE,
             default_transform: IDENTITY_MODEL_TRANSFORM,
         })
         .unwrap()
@@ -190,8 +189,7 @@ mod tests {
         let created = mesh(&mut data, true);
         let second = cache.get_or_build(&data).unwrap() as *const _;
         assert_ne!(first, second);
-        data.set_mesh_flags(created.mesh, RenderFlags::NONE)
-            .unwrap();
+        data.set_mesh_visible(created.mesh, false).unwrap();
         let third = cache.get_or_build(&data).unwrap() as *const _;
         assert_ne!(second, third);
         let mut moved = IDENTITY_MODEL_TRANSFORM;
@@ -213,7 +211,7 @@ mod tests {
         translated[3][0] = 5.;
         translated[3][1] = -2.;
         let extra = data
-            .create_instance(hidden.mesh, translated, RenderFlags::NONE)
+            .create_instance(hidden.mesh, translated, InstanceType::ZERO)
             .unwrap();
         let plan = SceneFramePlan::build(&data).unwrap();
         assert_eq!((plan.meshes.len(), plan.occurrences.len()), (2, 3));
@@ -247,13 +245,13 @@ mod tests {
         let doomed = mesh(&mut data, true);
         let c = mesh(&mut data, true);
         let a_extra = data
-            .create_instance(a.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::VISIBLE)
+            .create_instance(a.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::VISIBLE)
             .unwrap();
         let doomed_extra = data
-            .create_instance(doomed.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::VISIBLE)
+            .create_instance(doomed.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::VISIBLE)
             .unwrap();
         let c_extra = data
-            .create_instance(c.mesh, IDENTITY_MODEL_TRANSFORM, RenderFlags::VISIBLE)
+            .create_instance(c.mesh, IDENTITY_MODEL_TRANSFORM, InstanceType::VISIBLE)
             .unwrap();
         data.destroy_instance(a_extra).unwrap();
         data.destroy_mesh(doomed.mesh).unwrap();
@@ -262,7 +260,7 @@ mod tests {
             .create_instance(
                 replacement.mesh,
                 IDENTITY_MODEL_TRANSFORM,
-                RenderFlags::VISIBLE,
+                InstanceType::VISIBLE,
             )
             .unwrap();
         assert_eq!(replacement.mesh.slot(), doomed.mesh.slot());

renderer/src/shared_snapshot.rs

@@ -1,12 +1,12 @@
 //! Triple-buffered, immutable packed scene snapshot shared with JavaScript.
 use std::sync::atomic::{AtomicU32, Ordering};
 
-use crate::{render_data::RenderFlags, renderer::scene_frame::SceneFramePlan};
+use crate::renderer::scene_frame::SceneFramePlan;
 
 pub const MAGIC: u32 = u32::from_le_bytes(*b"YSNP");
-pub const BLOB_MAGIC: u32 = u32::from_le_bytes(*b"RDS1");
+pub const BLOB_MAGIC: u32 = u32::from_le_bytes(*b"RDS2");
 pub const CONTROL_VERSION: u32 = 1;
-pub const SCHEMA: u32 = 1;
+pub const SCHEMA: u32 = 2;
 pub const SLOT_COUNT: usize = 3;
 pub const INIT: u32 = 0;
 pub const OPEN: u32 = 1;
@@ -23,9 +23,9 @@ pub const ERROR_PUBLICATION: u32 = 4;
 const CONTROL_BYTES: u32 = 256;
 const SLOT_BYTES: u32 = 64;
 const SNAPSHOT_HEADER_BYTES: usize = 64;
-const DATA_OFFSET: usize = 512;
+const DATA_OFFSET: usize = 448;
 const DESCRIPTOR_BYTES: usize = 32;
-const STREAMS: usize = 14;
+const STREAMS: usize = 12;
 const SCHEMA_FLAGS: u32 = 3; // dense arrays | affine transforms
 
 #[repr(C, align(64))]
@@ -247,41 +247,40 @@ fn wasm_pages(minimum_end: usize) -> Result<u32, u32> {
 fn pack(data: &SceneFramePlan, epoch: u32) -> Result<Vec<u8>, u32> {
     let meshes = &data.meshes;
     let instances = &data.occurrences;
-    let strides = [4usize, 4, 4, 12, 12, 4, 4, 4, 4, 4, 64, 12, 12, 4];
-    let components = [1u32, 1, 1, 3, 3, 1, 1, 1, 1, 1, 16, 3, 3, 1];
-    let scalar = [1u32, 1, 1, 2, 2, 1, 1, 1, 1, 1, 2, 2, 2, 1];
-    let counts = [meshes.len(); 5]
+    let strides = [4usize, 4, 12, 12, 4, 4, 4, 4, 64, 12, 12, 64];
+    let components = [1u32, 1, 3, 3, 1, 1, 1, 1, 16, 3, 3, 16];
+    let scalar = [1u32, 1, 2, 2, 1, 1, 1, 1, 2, 2, 2, 1];
+    let counts = [meshes.len(); 4]
         .into_iter()
-        .chain([instances.len(); 9])
+        .chain([instances.len(); 8])
         .collect::<Vec<_>>();
     let mut offsets = [0usize; STREAMS];
     let mut cursor = DATA_OFFSET;
     for i in 0..STREAMS {
         offsets[i] = cursor;
-        let bytes = strides[i].checked_mul(counts[i]).ok_or(ERROR_OVERFLOW)?;
-        cursor = align16(cursor.checked_add(bytes).ok_or(ERROR_OVERFLOW)?)?;
+        cursor = align16(
+            cursor
+                .checked_add(strides[i].checked_mul(counts[i]).ok_or(ERROR_OVERFLOW)?)
+                .ok_or(ERROR_OVERFLOW)?,
+        )?;
     }
-    let total = u32::try_from(cursor).map_err(|_| ERROR_OVERFLOW)?;
-    let mesh_count = u32::try_from(meshes.len()).map_err(|_| ERROR_OVERFLOW)?;
-    let instance_count = u32::try_from(instances.len()).map_err(|_| ERROR_OVERFLOW)?;
     let mut out = vec![0u8; cursor];
-    let put32 = |out: &mut [u8], at: usize, value: u32| {
-        out[at..at + 4].copy_from_slice(&value.to_le_bytes())
-    };
+    let put32 =
+        |out: &mut [u8], at: usize, v: u32| out[at..at + 4].copy_from_slice(&v.to_le_bytes());
     let revision = data.revision;
-    for (i, value) in [
+    for (i, v) in [
         BLOB_MAGIC,
         SCHEMA,
-        SNAPSHOT_HEADER_BYTES as u32,
-        total,
+        64,
+        cursor as u32,
         epoch,
         revision as u32,
         (revision >> 32) as u32,
         STREAMS as u32,
-        SNAPSHOT_HEADER_BYTES as u32,
-        DESCRIPTOR_BYTES as u32,
-        mesh_count,
-        instance_count,
+        64,
+        32,
+        meshes.len() as u32,
+        instances.len() as u32,
         0x0102_0304,
         SCHEMA_FLAGS,
         0,
@@ -290,11 +289,11 @@ fn pack(data: &SceneFramePlan, epoch: u32) -> Result<Vec<u8>, u32> {
     .into_iter()
     .enumerate()
     {
-        put32(&mut out, i * 4, value);
+        put32(&mut out, i * 4, v);
     }
     for i in 0..STREAMS {
-        let at = SNAPSHOT_HEADER_BYTES + i * DESCRIPTOR_BYTES;
-        for (j, value) in [
+        let at = 64 + i * 32;
+        for (j, v) in [
             i as u32 + 1,
             scalar[i],
             offsets[i] as u32,
@@ -307,75 +306,61 @@ fn pack(data: &SceneFramePlan, epoch: u32) -> Result<Vec<u8>, u32> {
         .into_iter()
         .enumerate()
         {
-            put32(&mut out, at + j * 4, value);
+            put32(&mut out, at + j * 4, v);
         }
     }
-    for (dense, mesh) in meshes.iter().enumerate() {
-        for (i, value) in [
-            mesh.handle.slot(),
-            mesh.handle.generation(),
-            mesh.flags.bits(),
-        ]
-        .into_iter()
-        .enumerate()
-        {
-            put32(&mut out, offsets[i] + dense * 4, value);
-        }
+    for (d, m) in meshes.iter().enumerate() {
+        put32(&mut out, offsets[0] + d * 4, m.handle.slot());
+        put32(&mut out, offsets[1] + d * 4, m.handle.generation());
         for i in 0..3 {
             put32(
                 &mut out,
-                offsets[3] + dense * 12 + i * 4,
-                mesh.aabb.min[i].to_bits(),
+                offsets[2] + d * 12 + i * 4,
+                m.local_aabb.min[i].to_bits(),
             );
             put32(
                 &mut out,
-                offsets[4] + dense * 12 + i * 4,
-                mesh.aabb.max[i].to_bits(),
+                offsets[3] + d * 12 + i * 4,
+                m.local_aabb.max[i].to_bits(),
             );
         }
     }
-    for (dense, instance) in instances.iter().enumerate() {
-        let mesh = data
-            .meshes
-            .get(instance.mesh_index)
-            .ok_or(ERROR_INVARIANT)?;
-        for (i, value) in [
-            instance.handle.slot(),
-            instance.handle.generation(),
-            instance.mesh.slot(),
-            instance.mesh.generation(),
-            instance.flags.bits(),
+    for (d, x) in instances.iter().enumerate() {
+        for (i, v) in [
+            x.handle.slot(),
+            x.handle.generation(),
+            x.mesh.slot(),
+            x.mesh.generation(),
         ]
         .into_iter()
         .enumerate()
         {
-            put32(&mut out, offsets[5 + i] + dense * 4, value);
+            put32(&mut out, offsets[4 + i] + d * 4, v);
         }
         for i in 0..16 {
             put32(
                 &mut out,
-                offsets[10] + dense * 64 + i * 4,
-                instance.model[i / 4][i % 4].to_bits(),
+                offsets[8] + d * 64 + i * 4,
+                x.model[i / 4][i % 4].to_bits(),
+            );
+            put32(
+                &mut out,
+                offsets[11] + d * 64 + i * 4,
+                x.instance_type.words[i],
             );
         }
         for i in 0..3 {
             put32(
                 &mut out,
-                offsets[11] + dense * 12 + i * 4,
-                instance.world_aabb.min[i].to_bits(),
+                offsets[9] + d * 12 + i * 4,
+                x.world_aabb.min[i].to_bits(),
             );
             put32(
                 &mut out,
-                offsets[12] + dense * 12 + i * 4,
-                instance.world_aabb.max[i].to_bits(),
+                offsets[10] + d * 12 + i * 4,
+                x.world_aabb.max[i].to_bits(),
             );
         }
-        put32(
-            &mut out,
-            offsets[13] + dense * 4,
-            (mesh.flags.contains(RenderFlags::VISIBLE)
-                && instance.flags.contains(RenderFlags::VISIBLE)) as u32,
-        );
     }
     Ok(out)
 }
@@ -393,199 +378,14 @@ const _: [(); 256] = [(); std::mem::size_of::<SnapshotControl>()];
 #[cfg(test)]
 mod tests {
     use super::*;
-    use crate::render_data::{
-        MeshCreateInfo, PipelineKey, RenderData, RenderDataConfig, IDENTITY_MODEL_TRANSFORM,
-    };
-
     #[test]
-    fn exact_control_layout_and_initial_values() {
+    fn exact_schema_two_layout() {
+        assert_eq!(BLOB_MAGIC, u32::from_le_bytes(*b"RDS2"));
+        assert_eq!(SCHEMA, 2);
+        assert_eq!(STREAMS, 12);
+        assert_eq!(DATA_OFFSET, 448);
         assert_eq!(std::mem::size_of::<SnapshotControl>(), 256);
-        assert_eq!(std::mem::size_of::<SnapshotDescriptor>(), 64);
         let snapshot = SharedSnapshot::new();
-        let values: Vec<_> = snapshot
-            .control
-            .header
-            .iter()
-            .map(|v| v.load(Ordering::Relaxed))
-            .collect();
-        assert_eq!(&values[..7], &[MAGIC, 1, 256, 3, 64, 1, INIT]);
-        assert_eq!(values[9], u32::MAX);
-    }
-
-    #[test]
-    fn claiming_never_overwrites_reading_and_prefers_free() {
-        let snapshot = SharedSnapshot::new();
-        snapshot.control.slots[0].0[0].store(READING, Ordering::Relaxed);
-        assert_eq!(snapshot.claim_slot(), Some(1));
-        assert_eq!(
-            snapshot.control.slots[0].0[0].load(Ordering::Relaxed),
-            READING
-        );
-        assert_eq!(
-            snapshot.control.slots[1].0[0].load(Ordering::Relaxed),
-            WRITING
-        );
-    }
-
-    #[test]
-    fn producer_blob_abi_matches_schema_exactly() {
-        let mut data = RenderData::new(RenderDataConfig::default()).unwrap();
-        let create = |data: &mut RenderData, flags, x: f32| {
-            data.create_mesh(MeshCreateInfo {
-                positions: &[[x, 0., 0.], [x + 2., 0., 0.], [x, 3., 0.]],
-                normals: &[[0., 0., 1.]; 3],
-                tangents: &[[1., 0., 0., 1.]; 3],
-                uvs: &[[0., 0.]; 3],
-                indices: &[0, 1, 2],
-                pipeline: PipelineKey::new(7),
-                material: crate::render_data::MaterialKey::DEFAULT,
-                flags,
-                default_instance_flags: RenderFlags::VISIBLE,
-                default_transform: IDENTITY_MODEL_TRANSFORM,
-            })
-            .unwrap()
-        };
-        let visible = create(&mut data, RenderFlags::VISIBLE, -1.0);
-        let hidden = create(&mut data, RenderFlags::NONE, 10.0);
-        let mut model = IDENTITY_MODEL_TRANSFORM;
-        model[0][0] = 2.0;
-        model[1][1] = 0.5;
-        model[3][0] = 4.0;
-        model[3][1] = -2.0;
-        let extra = data
-            .create_instance(hidden.mesh, model, RenderFlags::NONE)
-            .unwrap();
-        let plan = SceneFramePlan::build(&data).unwrap();
-        let epoch = 0x1234_5678;
-        let blob = pack(&plan, epoch).unwrap();
-        let word = |at: usize| u32::from_le_bytes(blob[at..at + 4].try_into().unwrap());
-
-        assert_eq!(word(0), BLOB_MAGIC);
-        assert_eq!(word(4), SCHEMA);
-        assert_eq!(word(8), SNAPSHOT_HEADER_BYTES as u32);
-        assert_eq!(word(12), blob.len() as u32);
-        assert_eq!(word(16), epoch);
-        assert_eq!(word(20), plan.revision as u32);
-        assert_eq!(word(24), (plan.revision >> 32) as u32);
-        assert_eq!(word(28), STREAMS as u32);
-        assert_eq!(word(32), SNAPSHOT_HEADER_BYTES as u32);
-        assert_eq!(word(36), DESCRIPTOR_BYTES as u32);
-        assert_eq!(word(40), 2);
-        assert_eq!(word(44), 3);
-        assert_eq!(word(48), 0x0102_0304);
-        assert_eq!(word(52), SCHEMA_FLAGS);
-
-        let strides = [4, 4, 4, 12, 12, 4, 4, 4, 4, 4, 64, 12, 12, 4];
-        let scalars = [1, 1, 1, 2, 2, 1, 1, 1, 1, 1, 2, 2, 2, 1];
-        let components = [1, 1, 1, 3, 3, 1, 1, 1, 1, 1, 16, 3, 3, 1];
-        let counts = [2usize; 5]
-            .into_iter()
-            .chain([3usize; 9])
-            .collect::<Vec<_>>();
-        let mut offsets = Vec::new();
-        let mut cursor = DATA_OFFSET;
-        for i in 0..STREAMS {
-            let at = SNAPSHOT_HEADER_BYTES + i * DESCRIPTOR_BYTES;
-            offsets.push(cursor);
-            assert_eq!(
-                [
-                    word(at),
-                    word(at + 4),
-                    word(at + 8),
-                    word(at + 12),
-                    word(at + 16),
-                    word(at + 20),
-                    word(at + 24),
-                    word(at + 28)
-                ],
-                [
-                    i as u32 + 1,
-                    scalars[i],
-                    cursor as u32,
-                    counts[i] as u32,
-                    components[i],
-                    strides[i] as u32,
-                    4,
-                    0
-                ]
-            );
-            assert_eq!(cursor % 16, 0);
-            cursor = (cursor + strides[i] as usize * counts[i] + 15) & !15;
-        }
-        assert_eq!(cursor, blob.len());
-
-        for (dense, mesh) in plan.meshes.iter().enumerate() {
-            assert_eq!(word(offsets[0] + dense * 4), mesh.handle.slot());
-            assert_eq!(word(offsets[1] + dense * 4), mesh.handle.generation());
-            assert_eq!(word(offsets[2] + dense * 4), mesh.flags.bits());
-            for axis in 0..3 {
-                assert_eq!(
-                    word(offsets[3] + dense * 12 + axis * 4),
-                    mesh.aabb.min[axis].to_bits()
-                );
-                assert_eq!(
-                    word(offsets[4] + dense * 12 + axis * 4),
-                    mesh.aabb.max[axis].to_bits()
-                );
-            }
-        }
-        for (dense, occurrence) in plan.occurrences.iter().enumerate() {
-            assert_eq!(
-                [
-                    word(offsets[5] + dense * 4),
-                    word(offsets[6] + dense * 4),
-                    word(offsets[7] + dense * 4),
-                    word(offsets[8] + dense * 4),
-                    word(offsets[9] + dense * 4)
-                ],
-                [
-                    occurrence.handle.slot(),
-                    occurrence.handle.generation(),
-                    occurrence.mesh.slot(),
-                    occurrence.mesh.generation(),
-                    occurrence.flags.bits()
-                ]
-            );
-            for i in 0..16 {
-                assert_eq!(
-                    word(offsets[10] + dense * 64 + i * 4),
-                    occurrence.model[i / 4][i % 4].to_bits()
-                );
-            }
-            for axis in 0..3 {
-                assert_eq!(
-                    word(offsets[11] + dense * 12 + axis * 4),
-                    occurrence.world_aabb.min[axis].to_bits()
-                );
-                assert_eq!(
-                    word(offsets[12] + dense * 12 + axis * 4),
-                    occurrence.world_aabb.max[axis].to_bits()
-                );
-            }
-            let mesh_visible = plan.meshes[occurrence.mesh_index]
-                .flags
-                .contains(RenderFlags::VISIBLE);
-            assert_eq!(
-                word(offsets[13] + dense * 4),
-                (mesh_visible && occurrence.flags.contains(RenderFlags::VISIBLE)) as u32
-            );
-        }
-        assert!(plan
-            .meshes
-            .windows(2)
-            .all(|w| w[0].handle.slot() < w[1].handle.slot()));
-        assert!(plan
-            .occurrences
-            .windows(2)
-            .all(|w| w[0].handle.slot() < w[1].handle.slot()));
-        assert_eq!(
-            plan.occurrences
-                .iter()
-                .find(|o| o.handle == extra)
-                .unwrap()
-                .model,
-            model
-        );
-        assert!(plan.occurrences.iter().any(|o| o.mesh == visible.mesh));
+        assert_eq!(snapshot.control.header[5].load(Ordering::Relaxed), 2);
     }
 }