diff --git a/.gitignore b/.gitignore index 3209a61..7fd4781 100644 --- a/.gitignore +++ b/.gitignore @@ -79,3 +79,7 @@ examples/intel-sponza/assets/outdoor.hdr examples/**/main.exe examples/**/main examples/**/main.app + +# DXC runtime DLLs (local, for the PT golden tests on Windows) +native/shared/dxcompiler.dll +native/shared/dxil.dll diff --git a/docs/pt/PT-6-7-8-skinned-tlas-motion-oracle.md b/docs/pt/PT-6-7-8-skinned-tlas-motion-oracle.md new file mode 100644 index 0000000..18c4ed5 --- /dev/null +++ b/docs/pt/PT-6-7-8-skinned-tlas-motion-oracle.md @@ -0,0 +1,82 @@ +# PT-6/7/8 — skinned TLAS, real motion vectors, correctness oracle + +Status: **landed** (2026-07-14, 760M / DX12+DXC). + +## PT-6 — skinned meshes enter the TLAS + +Skinned characters were ghosts to the tracer: skinning lived only in +the vertex shader, so no posed geometry existed for a BLAS, and the +TLAS was scene-nodes-only. Now: + +- `cache_model_if_static` retains bind-pose CPU geometry (+ STORAGE on + the VB) for skinned meshes; `draw_model_cached_skinned` registers + each mesh as a `PtDynamicDraw` per frame. +- `rebuild_instance_data` appends a megabuffer window + instance entry + per dynamic draw; a compute pre-skin pass (`PT_SKIN_WGSL`, the same + palette blend as the raster VS) overwrites the window's + position/normal with posed WORLD-SPACE data (the palette bakes + placement → identity TLAS transform). +- Per-slot BLASes rebuild every frame FROM the megabuffer windows + (`first_vertex`/`first_index`) — intersection and hit shading read + the same bytes. The megabuffers gained `BLAS_INPUT` usage. +- The TLAS recreates when the total instance count outgrows the + capacity it was built with (`tlas_created_cap`) — wave spawns grow + the count mid-run, which the old check missed. + +Lumen's HW probe trace consumes the same instances, so skinned +characters now contribute to GI too. Verified via the debug-6 traced +view: the skinned player renders with interpolated normals — the first +time any skinned mesh has appeared in a traced image. + +## PT-7 — real skinned motion vectors + velocity-driven PT reprojection + +Skinned draws wrote EXACTLY zero velocity (no previous-frame joint +palette existed anywhere; the cached path stamped `prev_mvp` with the +current VP). Enemies ghosted under TAA/TSR and invalidated their own +PT history every frame. + +- The joint palette is double-buffered. `set_joint_matrices_scaled` + takes a pairing key (the FFI anim handle); the previous palette for + the same key stages in lockstep so both arenas share offsets. First + sighting (spawn) pairs with itself = zero velocity, correct. +- The scene VS reconstructs last frame's world position from + `joints_prev` (group 3 binding 1) and projects it through the EN-022 + velocity-reference VP — skeletal AND locomotion motion land in the + velocity MRT. +- The PT kernel binds the velocity MRT (binding 22). `compute_reproj` + follows per-pixel motion when non-zero (TAA's convention: + `prev_uv = (uv.x − vel.x, uv.y + vel.y)`) and falls back to the + camera `prev_vp` math otherwise. Moving skinned characters keep + their SVGF history instead of resetting to 1 spp (debug-20 history + heat: no rejection hole on the animating player). + +## PT-8 — the correctness oracle + +Two golden-image tests in `native/shared/tests/golden_render.rs`, +running the REAL engine headless on a ray-query device (skip cleanly +without one; on Windows `dxcompiler.dll`/`dxil.dll` must be loadable — +untracked local copies next to the crate, see .gitignore): + +- **`pt_progressive`** — converged progressive mode (300 static + frames) on a node scene. Catches transport regressions (BRDF energy, + NEE, sky handling, accumulation math) as an image diff. +- **`pt_realtime_motion`** — realtime mode while the camera orbits. + Catches reprojection/temporal regressions: a broken history (the + prev_vp-transpose class that survived three human review rounds) + floods the image with unconverged speckle, far past tolerance. + +`BLOOM_UPDATE_GOLDEN=1 cargo test golden` regenerates; the strict +outlier gate stays global so looser mean tolerances cannot hide a +broken region. + +**The oracle caught a real bug before it was even committed**: the +kernel seeded its RNG from `taa_frame_index`, which freezes when TAA +is off — the sample sequence froze and progressive accumulation +silently never converged (300 frames = the same image as 1). A player +disabling TAA in settings would have hit exactly this in-game. PT now +keeps its own rolling `pt_frame_index`. + +The bloom-reference RMSE parity run remains a manual protocol (see +PT-5 doc): number-parity on the game scene is now *closer* (skinned +meshes are in the TLAS) but cutout foliage and card-resolution albedo +still make it a human-judged comparison, not a gate. diff --git a/native/shared/src/ffi_core/models.rs b/native/shared/src/ffi_core/models.rs index feaa707..80232df 100644 --- a/native/shared/src/ffi_core/models.rs +++ b/native/shared/src/ffi_core/models.rs @@ -378,7 +378,9 @@ macro_rules! __bloom_ffi_models { eng.models.update_model_animation(handle, anim_index as usize, time as f32); if let Some(anim) = eng.models.get_animation(handle) { if !anim.joint_matrices.is_empty() { - eng.renderer.set_joint_matrices_scaled(&anim.joint_matrices, scale as f32, [px as f32, py as f32, pz as f32], rot_sin, rot_cos); + // PT-7: the anim handle keys the prev-palette + // pairing for skinned motion vectors. + eng.renderer.set_joint_matrices_scaled(handle.to_bits(), &anim.joint_matrices, scale as f32, [px as f32, py as f32, pz as f32], rot_sin, rot_cos); } } }) @@ -934,8 +936,9 @@ macro_rules! __bloom_ffi_models { eng.models.advance_and_update(handle, dt as f32); if let Some(anim) = eng.models.get_animation(handle) { if !anim.joint_matrices.is_empty() { + // PT-7: anim handle = prev-palette pairing key. eng.renderer.set_joint_matrices_scaled( - &anim.joint_matrices, scale as f32, + handle.to_bits(), &anim.joint_matrices, scale as f32, [px as f32, py as f32, pz as f32], rot_sin, rot_cos); } } diff --git a/native/shared/src/ffi_core/ragdoll_ffi.rs b/native/shared/src/ffi_core/ragdoll_ffi.rs index e59996d..2c9fc8c 100644 --- a/native/shared/src/ffi_core/ragdoll_ffi.rs +++ b/native/shared/src/ffi_core/ragdoll_ffi.rs @@ -205,8 +205,9 @@ macro_rules! __bloom_ffi_ragdoll { if let Some(a) = eng.models.get_animation(anim) { if !a.joint_matrices.is_empty() { let (s, c) = (rot.sin(), rot.cos()); + // PT-7: anim handle = prev-palette pairing key. eng.renderer.set_joint_matrices_scaled( - &a.joint_matrices, scale, pos, s, c); + anim.to_bits(), &a.joint_matrices, scale, pos, s, c); } } age as f64 diff --git a/native/shared/src/models.rs b/native/shared/src/models.rs index f287641..af25a9b 100644 --- a/native/shared/src/models.rs +++ b/native/shared/src/models.rs @@ -586,7 +586,18 @@ impl ModelManager { if let Some(ma) = self.animations.get_mut(handle) { if clip >= ma.animations.len() { return; } let m = &mut ma.mixer; - if m.started && m.cur_clip == clip && m.fade_dur <= 0.0 { + // Re-requesting the clip that is ALREADY current is a no-op by + // contract ("safe to call every frame with the clip you want"), and + // that has to hold DURING a fade too — a fade is a fade *to* + // cur_clip, so the right thing is to let it finish. + // + // This used to also require `m.fade_dur <= 0.0`. A caller doing the + // documented thing then fell through to the restart path on every + // frame of the fade: it re-seeded the fade (fade_t = 0, so the fade + // could never complete) and reset cur_time = 0. The base clip was + // pinned at t≈0 forever, from the first clip change onward — every + // animated character in a game froze the moment it stopped idling. + if m.started && m.cur_clip == clip { m.cur_speed = speed; m.cur_loop = looping; return; diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 5bcd904..7cab18d 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -749,6 +749,48 @@ struct GpuMesh { /// its real shape. `None` → opaque caster, uses the plain shadow pipeline. shadow_cutout_bg: Option, _shadow_cutoff_buf: Option, + /// PT-6 — retained CPU geometry for SKINNED meshes only: the + /// bind-pose vertices/indices are appended into the PT geometry + /// megabuffer each frame as the window the compute pre-skin pass + /// then overwrites with posed world-space data. `None` for rigid + /// meshes (they enter the TLAS as scene nodes or not at all). + cpu_vertices: Option>, + cpu_indices: Option>, + /// PT-6 — hit-shading inputs for the dynamic TLAS instance. + base_color_idx: u32, + metallic_factor: f32, + roughness_factor: f32, +} + +/// PT-6 — a dynamic instance's megabuffer window for this frame: +/// where its geometry lives (Vertex3D-slot / index units), which pose +/// to apply, and which cached mesh is its compute-skin source. +struct PtDynWindow { + v_base: u32, + v_count: u32, + i_base: u32, + i_count: u32, + joint_offset: f32, + model: [[f32; 4]; 4], + cache_handle: u64, + mesh_idx: usize, +} + +/// PT-6 — one skinned mesh drawn this frame, registered by +/// `draw_model_cached_skinned` for the dynamic TLAS path. Cleared with +/// the per-frame draw lists; consumed by `rebuild_instance_data`. +pub(crate) struct PtDynamicDraw { + cache_handle: u64, + mesh_idx: usize, + /// Base slot of the pose in the frame joint buffer (world-space + /// palette — the compute pre-skin outputs world-space vertices, + /// so the TLAS instance transform is identity). + joint_offset: f32, + /// Places the rare rigid (weightless) verts, same as the VS. + model: [[f32; 4]; 4], + /// World AABB (joint-union bound from the draw call). + wmin: [f32; 3], + wmax: [f32; 3], } struct CachedModelDraw { @@ -857,6 +899,8 @@ pub struct Renderer { // Joint matrices for GPU skinning (64 joints × 4x4 matrix) joint_buffer: wgpu::Buffer, + /// PT-7 — previous frame's palette (same slot offsets). + joint_prev_buffer: wgpu::Buffer, joint_bind_group: wgpu::BindGroup, // False until the material system's per-view bind group has been @@ -1232,6 +1276,12 @@ pub struct Renderer { /// rounded up to 1024 instances; resized only when exceeded. pub tlas: Option, pub tlas_max_instances: u32, + /// Capacity the CURRENT tlas object was created with. The dynamic + /// (skinned) instance count grows mid-run when waves spawn, so the + /// TLAS must be recreated when the total exceeds what it was built + /// for — `tlas_max_instances` alone can't tell (the instance-data + /// buffer grows it before the TLAS check runs). + tlas_created_cap: u32, /// `SceneGraph::tlas_version` the last time we rebuilt the TLAS + /// instance-data buffer. Mismatch → rebuild. Mirrors `shadow_version`'s /// cache pattern (ticket 004). @@ -1296,10 +1346,31 @@ pub struct Renderer { pt_last_tlas_version: u64, /// BLOOM_PT_DEBUG view selector, forwarded to the kernel via cfg.w. pt_debug: f32, + /// PT's own rolling frame counter (kernel RNG seed). Deliberately + /// NOT taa_frame_index: that freezes when TAA is off, which froze + /// the sample sequence and stopped accumulation from converging. + pt_frame_index: u32, /// BLOOM_PT_RESTIR=1 — PT-4 experimental ReSTIR DI (kernel ext.w). pt_restir: bool, /// PT-4 — ReSTIR reservoir ping-pong, sized with the accum pair. pt_resv_buffers: [Option; 2], + /// PT-6 — skinned meshes drawn this frame (dynamic TLAS instances). + /// Pushed by draw_model_cached_skinned, consumed by + /// rebuild_instance_data, cleared with the per-frame draw lists. + pt_dynamic_draws: Vec, + /// PT-6 — per-slot megabuffer windows for this frame's dynamic + /// instances. Feeds the compute pre-skin dispatches and the + /// per-frame BLAS builds. + pt_dyn_windows: Vec, + /// PT-6 — dynamic BLAS pool, one slot per dynamic instance. + /// Recreated when a slot's (vertex_count, index_count) changes. + pt_dyn_blas: Vec<(wgpu::Blas, u32, u32)>, + /// PT-6 — compute pre-skin pipeline (posed world-space vertices + /// written into the PT megabuffer window before the BLAS build). + pt_skin_pipeline: Option, + pt_skin_layout: Option, + /// Per-slot uniform buffers for the pre-skin params (grow-only). + pt_skin_params: Vec, /// PT-2 — concatenated Vertex3D words (f32) / indices (u32) for /// interpolated hit shading. Grow-only, rebuilt alongside the /// instance-data buffer. @@ -1655,6 +1726,17 @@ pub struct Renderer { // At `end_frame` the accumulator is flushed in one write. pub pending_skin_groups: Vec>, pub frame_joint_data: Vec<[[f32; 4]; 4]>, + /// PT-7 — the PREVIOUS frame's palette for each staged group, in + /// lockstep with `pending_skin_groups`/`frame_joint_data` (same + /// offsets). The skinned VS reconstructs last frame's world + /// position from it, which is what makes skeletal motion vectors + /// real (they were exactly zero before — no history existed). + pub pending_skin_groups_prev: Vec>, + pub frame_joint_data_prev: Vec<[[f32; 4]; 4]>, + /// Last staged palette per animation key (FFI anim handle). A few + /// dozen entries at most — anim handles are per model slot, not + /// per spawn — so no eviction is needed. + skin_prev_palettes: std::collections::HashMap>, pub model_skin_scale: f32, // Shadow mapping @@ -2338,16 +2420,31 @@ impl Renderer { // caster pipeline needs this layout at construction time. let joint_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("joint_layout"), - entries: &[wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: wgpu::ShaderStages::VERTEX, - ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Uniform, - has_dynamic_offset: false, - min_binding_size: None, + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::VERTEX, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, }, - count: None, - }], + // PT-7 — previous frame's palette (same offsets). Only + // the scene VS reads it; the shadow shader ignores the + // binding, which wgpu permits. + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::VERTEX, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], }); // Shadow map needs to be created before the lighting bind @@ -2559,13 +2656,25 @@ impl Renderer { contents: &joint_data, usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, }); + // PT-7 — previous frame's palette, parallel offsets. + let joint_prev_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("joint_prev_buffer"), + contents: &joint_data, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); let joint_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("joint_bg"), layout: &joint_layout, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: joint_buffer.as_entire_binding(), - }], + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: joint_buffer.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: joint_prev_buffer.as_entire_binding(), + }, + ], }); // Initialize with identity matrices { @@ -2580,6 +2689,7 @@ impl Renderer { identity_data[offset+60..offset+64].copy_from_slice(&one); // [3][3] } queue.write_buffer(&joint_buffer, 0, &identity_data); + queue.write_buffer(&joint_prev_buffer, 0, &identity_data); } // --- 3D Pipeline --- @@ -5239,6 +5349,16 @@ impl Renderer { has_dynamic_offset: false, min_binding_size: None, }, count: None, }, + // PT-7 — raster velocity MRT for object-motion + // reprojection (skinned characters). + wgpu::BindGroupLayoutEntry { + binding: 22, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: false }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, ]; let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("pt_layout"), @@ -6816,6 +6936,7 @@ impl Renderer { lighting_buffer, lighting_bind_group, joint_buffer, + joint_prev_buffer, joint_bind_group, material_per_view_bg_live: false, texture_bind_group_layout, @@ -6996,6 +7117,7 @@ impl Renderer { hw_rt_enabled, tlas: None, tlas_max_instances: 1024, + tlas_created_cap: 0, tlas_built_version: 0, tlas_instance_data_buffer: None, probe_place_pipeline, @@ -7020,8 +7142,15 @@ impl Renderer { pt_prev_vp: [[0.0; 4]; 4], pt_last_tlas_version: 0, pt_debug, + pt_frame_index: 0, pt_restir, pt_resv_buffers: [None, None], + pt_dynamic_draws: Vec::new(), + pt_dyn_windows: Vec::new(), + pt_dyn_blas: Vec::new(), + pt_skin_pipeline: None, + pt_skin_layout: None, + pt_skin_params: Vec::new(), pt_geo_vertex_buffer: None, pt_geo_index_buffer: None, pt_texture_arrays_enabled, @@ -7179,6 +7308,9 @@ impl Renderer { render_mode: RenderMode::ScreenSpace, pending_skin_groups: Vec::with_capacity(8), frame_joint_data: Vec::with_capacity(256), + pending_skin_groups_prev: Vec::with_capacity(8), + frame_joint_data_prev: Vec::with_capacity(256), + skin_prev_palettes: std::collections::HashMap::new(), model_skin_scale: 1.0, clear_color: wgpu::Color::BLACK, custom_pipelines: Vec::new(), @@ -8218,7 +8350,10 @@ impl Renderer { scene: &crate::scene::SceneGraph, instance_handles: &[f64], ) -> (u32, bool) { - let instance_count = instance_handles.len() as u32; + // PT-6 — dynamic (skinned) instances ride after the node + // instances in both the instance-data buffer and the TLAS. + let instance_count = + instance_handles.len() as u32 + self.pt_dynamic_draws.len() as u32; let mut resized = false; let needed_cap = instance_count.max(self.tlas_max_instances).max(64); if self.tlas_instance_data_buffer.is_none() @@ -8312,6 +8447,63 @@ impl Renderer { mat_params: [n.material.roughness, n.material.metalness, 0.0, 0.0], }); } + // PT-6 — dynamic skinned instances: append the BIND-POSE + // geometry as this frame's megabuffer window (the compute + // pre-skin overwrites position/normal with posed world-space + // data before the BLAS build) plus an instance entry per mesh. + // Hit shading multiplies the texture by inst.albedo, so the + // flat albedo is white; no card (geo window always present). + self.pt_dyn_windows.clear(); + let dyn_draws = std::mem::take(&mut self.pt_dynamic_draws); + for d in &dyn_draws { + let mesh = match self.model_gpu_cache.get(&d.cache_handle) { + Some(Some(meshes)) => match meshes.get(d.mesh_idx) { + Some(m) => m, + None => continue, + }, + _ => continue, + }; + let (cv, ci) = match (&mesh.cpu_vertices, &mesh.cpu_indices) { + (Some(v), Some(i)) if !v.is_empty() && !i.is_empty() => (v, i), + _ => continue, + }; + let vb = (geo_vertices.len() / 24) as u32; + let ib = geo_indices.len() as u32; + geo_vertices.extend_from_slice(bytemuck::cast_slice(cv)); + geo_indices.extend_from_slice(ci); + instance_data.push(InstanceGiDataCpu { + albedo: [1.0, 1.0, 1.0], + emissive_luma: 0.0, + normal_ws: [0.0, 1.0, 0.0], + _pad0: 0.0, + card_slot: [0.0, 0.0, 0.0, 0.0], + card_aabb_min: [mesh.local_min[0], mesh.local_min[1], mesh.local_min[2], 0.0], + card_aabb_max: [mesh.local_max[0], mesh.local_max[1], mesh.local_max[2], 0.0], + world_aabb_min: [d.wmin[0], d.wmin[1], d.wmin[2], 0.0], + world_aabb_max: [d.wmax[0], d.wmax[1], d.wmax[2], 0.0], + geo: [ + vb, + ib, + ci.len() as u32, + if (mesh.base_color_idx as usize) < PT_MAX_TEXTURES { + mesh.base_color_idx + } else { + 0 + }, + ], + mat_params: [mesh.roughness_factor, mesh.metallic_factor, 0.0, 0.0], + }); + self.pt_dyn_windows.push(PtDynWindow { + v_base: vb, + v_count: cv.len() as u32, + i_base: ib, + i_count: ci.len() as u32, + joint_offset: d.joint_offset, + model: d.model, + cache_handle: d.cache_handle, + mesh_idx: d.mesh_idx, + }); + } // PT-2 — upload the megabuffers (grow-only; a minimum size keeps // bind-group creation valid before any geometry is committed). { @@ -8319,11 +8511,21 @@ impl Renderer { let i_bytes = (geo_indices.len() * 4).max(16) as u64; let v_recreate = self.pt_geo_vertex_buffer.as_ref().map_or(true, |b| b.size() < v_bytes); let i_recreate = self.pt_geo_index_buffer.as_ref().map_or(true, |b| b.size() < i_bytes); + // PT-6 — on RT adapters the megabuffers double as BLAS + // geometry inputs for the dynamic skinned instances (the + // BLAS reads a window via first_vertex/first_index). + let mega_usage = if self.hw_rt_enabled { + wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::BLAS_INPUT + } else { + wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST + }; if v_recreate { self.pt_geo_vertex_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_geo_vertices"), size: v_bytes.next_power_of_two(), - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + usage: mega_usage, mapped_at_creation: false, })); } @@ -8331,7 +8533,7 @@ impl Renderer { self.pt_geo_index_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_geo_indices"), size: i_bytes.next_power_of_two(), - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + usage: mega_usage, mapped_at_creation: false, })); } @@ -8383,7 +8585,9 @@ impl Renderer { bytemuck::cast_slice(&instance_data), ); } - (instance_count, resized) + // Actual written count (dynamic draws with a missing cache + // entry were skipped, so this can be below the requested cap). + (instance_data.len() as u32, resized) } fn rebuild_acceleration_structures( @@ -8397,7 +8601,10 @@ impl Renderer { // branch on `hw_rt_enabled` for TLAS/BLAS-specific work. let pending_blas = !scene.pending_blas_builds.is_empty(); let version_changed = scene.tlas_version != self.tlas_built_version; - if !pending_blas && !version_changed { + // PT-6 — dynamic skinned instances re-pose every frame, so any + // frame that drew one rebuilds regardless of scene versioning. + let has_dyn = !self.pt_dynamic_draws.is_empty(); + if !pending_blas && !version_changed && !has_dyn { return; } @@ -8407,7 +8614,10 @@ impl Renderer { instance_handles.push(h); } } + let node_count = instance_handles.len(); let (instance_count, _resized) = self.rebuild_instance_data(scene, &instance_handles); + // Dynamic instances occupy the slots after the nodes. + let dyn_count = (instance_count as usize).saturating_sub(node_count); // Everything below this point is HW-ray-tracing-specific // (TLAS build + BLAS batch). On SW adapters we're done. @@ -8421,7 +8631,11 @@ impl Renderer { return; } - if self.tlas.is_none() || instance_count > self.tlas_max_instances { + // PT-6 — pose the dynamic windows (compute pre-skin into the + // megabuffer) and make sure each slot has a matching BLAS. + self.encode_dynamic_skin(encoder, dyn_count); + + if self.tlas.is_none() || instance_count > self.tlas_created_cap { let cap = self.tlas_max_instances.max(64); self.tlas = Some(self.device.create_tlas(&wgpu::CreateTlasDescriptor { label: Some("bloom_tlas"), @@ -8429,6 +8643,7 @@ impl Renderer { flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE, update_mode: wgpu::AccelerationStructureUpdateMode::Build, })); + self.tlas_created_cap = cap; self.probe_trace_hw_bg_cache = [None, None]; // V14 — same reason as the resize path above. self.wsrc_bake_hw_bg_cache = None; @@ -8441,7 +8656,7 @@ impl Renderer { // instance count (extras beyond stay None from the initial vec). { let tlas = self.tlas.as_mut().unwrap(); - for slot in 0..(instance_count as usize) { + for slot in 0..node_count { let n = scene.nodes.get(instance_handles[slot]).unwrap(); let blas = n.blas.as_ref().unwrap(); let t = &n.transform; @@ -8460,6 +8675,23 @@ impl Renderer { 0xff, )); } + // PT-6 — dynamic instances: IDENTITY transform (the joint + // palette bakes world placement, so the compute pre-skin + // output is already world-space). + for i in 0..dyn_count { + let slot = node_count + i; + let blas = &self.pt_dyn_blas[i].0; + tlas[slot] = Some(wgpu::TlasInstance::new( + blas, + [ + 1.0, 0.0, 0.0, 0.0, + 0.0, 1.0, 0.0, 0.0, + 0.0, 0.0, 1.0, 0.0, + ], + slot as u32, + 0xff, + )); + } // Clear any slots beyond the current instance count left // over from a previous rebuild that had more instances. for slot in (instance_count as usize)..(self.tlas_max_instances as usize) { @@ -8517,6 +8749,42 @@ impl Renderer { }); } + // PT-6 — dynamic BLASes rebuild EVERY frame from the freshly + // posed megabuffer windows (topology constant, positions move). + let dyn_size_descs: Vec = self + .pt_dyn_windows + .iter() + .take(dyn_count) + .map(|w| wgpu::BlasTriangleGeometrySizeDescriptor { + vertex_format: wgpu::VertexFormat::Float32x3, + vertex_count: w.v_count, + index_format: Some(wgpu::IndexFormat::Uint32), + index_count: Some(w.i_count), + flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE, + }) + .collect(); + if dyn_count > 0 { + let mega_vb = self.pt_geo_vertex_buffer.as_ref().unwrap(); + let mega_ib = self.pt_geo_index_buffer.as_ref().unwrap(); + for (i, w) in self.pt_dyn_windows.iter().take(dyn_count).enumerate() { + build_entries.push(wgpu::BlasBuildEntry { + blas: &self.pt_dyn_blas[i].0, + geometry: wgpu::BlasGeometries::TriangleGeometries(vec![ + wgpu::BlasTriangleGeometry { + size: &dyn_size_descs[i], + vertex_buffer: mega_vb, + first_vertex: w.v_base, + vertex_stride: std::mem::size_of::() as u64, + index_buffer: Some(mega_ib), + first_index: Some(w.i_base), + transform_buffer: None, + transform_buffer_offset: None, + }, + ]), + }); + } + } + let tlas_ref = self.tlas.as_ref().unwrap(); encoder.build_acceleration_structures( build_entries.iter(), @@ -8526,6 +8794,166 @@ impl Renderer { self.tlas_built_version = scene.tlas_version; } + /// PT-6 — encode the compute pre-skin dispatches for this frame's + /// dynamic windows (posed world-space position/normal written over + /// the bind-pose data in the megabuffer) and (re)create each slot's + /// BLAS when its geometry size changed. Runs before the combined + /// BLAS+TLAS build on the same encoder, so the build reads posed + /// vertices. + fn encode_dynamic_skin( + &mut self, + encoder: &mut wgpu::CommandEncoder, + dyn_count: usize, + ) { + if dyn_count == 0 { + return; + } + // Lazy pipeline: first dynamic draw on an RT adapter. + if self.pt_skin_pipeline.is_none() { + let shader = self.device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("pt_skin_shader"), + source: wgpu::ShaderSource::Wgsl(shaders::PT_SKIN_WGSL.into()), + }); + let layout = self.device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("pt_skin_layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 3, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + ], + }); + let pl = self.device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("pt_skin_pl"), + bind_group_layouts: &[Some(&layout)], + immediate_size: 0, + }); + self.pt_skin_pipeline = Some(self.device.create_compute_pipeline( + &wgpu::ComputePipelineDescriptor { + label: Some("pt_skin_pipeline"), + layout: Some(&pl), + module: &shader, + entry_point: Some("cs_skin"), + compilation_options: Default::default(), + cache: None, + }, + )); + self.pt_skin_layout = Some(layout); + } + // Params pool (mat4 model + vec4 = 80 bytes per slot). + while self.pt_skin_params.len() < dyn_count { + self.pt_skin_params.push(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_skin_params"), + size: 80, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + } + + // Per-slot: params upload, BLAS pool upkeep, bind group. + let mut bind_groups: Vec = Vec::with_capacity(dyn_count); + let mut dispatch_counts: Vec = Vec::with_capacity(dyn_count); + for i in 0..dyn_count { + let w = &self.pt_dyn_windows[i]; + let mut words = [0f32; 20]; + for c in 0..4 { + for r in 0..4 { + words[c * 4 + r] = w.model[c][r]; + } + } + words[16] = f32::from_bits(w.v_base); + words[17] = f32::from_bits(w.v_count); + words[18] = f32::from_bits(w.joint_offset.max(0.0) as u32); + words[19] = 0.0; + self.queue.write_buffer( + &self.pt_skin_params[i], + 0, + bytemuck::cast_slice(&words), + ); + + let need_new = match self.pt_dyn_blas.get(i) { + Some((_, vc, ic)) => *vc != w.v_count || *ic != w.i_count, + None => true, + }; + if need_new { + let blas = self.device.create_blas( + &wgpu::CreateBlasDescriptor { + label: Some("pt_dyn_blas"), + flags: wgpu::AccelerationStructureFlags::PREFER_FAST_TRACE, + update_mode: wgpu::AccelerationStructureUpdateMode::Build, + }, + wgpu::BlasGeometrySizeDescriptors::Triangles { + descriptors: vec![wgpu::BlasTriangleGeometrySizeDescriptor { + vertex_format: wgpu::VertexFormat::Float32x3, + vertex_count: w.v_count, + index_format: Some(wgpu::IndexFormat::Uint32), + index_count: Some(w.i_count), + flags: wgpu::AccelerationStructureGeometryFlags::OPAQUE, + }], + }, + ); + if i < self.pt_dyn_blas.len() { + self.pt_dyn_blas[i] = (blas, w.v_count, w.i_count); + } else { + self.pt_dyn_blas.push((blas, w.v_count, w.i_count)); + } + } + + let src_vb = match self.model_gpu_cache.get(&w.cache_handle) { + Some(Some(meshes)) => match meshes.get(w.mesh_idx) { + Some(m) => &m.vb, + None => continue, + }, + _ => continue, + }; + bind_groups.push(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_skin_bg"), + layout: self.pt_skin_layout.as_ref().unwrap(), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.pt_skin_params[i].as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: src_vb.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: self.pt_geo_vertex_buffer.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: self.joint_buffer.as_entire_binding() }, + ], + })); + dispatch_counts.push(w.v_count); + } + + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("pt_skin"), + timestamp_writes: None, + }); + pass.set_pipeline(self.pt_skin_pipeline.as_ref().unwrap()); + for (bg, vc) in bind_groups.iter().zip(dispatch_counts.iter()) { + pass.set_bind_group(0, bg, &[]); + pass.dispatch_workgroups(vc.div_ceil(64), 1, 1); + } + } + /// SSGI intensity multiplier (0 = off, 0.5 = default, 1+ = strong). /// Controls the brightness of indirect bounce light. pub fn set_ssgi_intensity(&mut self, intensity: f32) { @@ -9775,6 +10203,10 @@ impl Renderer { self.indices_3d.clear(); self.draw_calls_3d.clear(); self.model_draw_commands.clear(); + // PT-6 — normally consumed by rebuild_instance_data (mem::take); + // this clear covers frames where the accel path is skipped + // (PT and Lumen both off) so the registry can't grow unbounded. + self.pt_dynamic_draws.clear(); self.next_model_uniform_slot = 0; self.current_texture_3d = 0; self.current_uniform_idx = 0; @@ -11126,6 +11558,9 @@ impl Renderer { } pub fn set_joint_matrices(&mut self, matrices: &[[[f32; 4]; 4]]) { + // Unkeyed path (editor/tests): no palette history, so previous + // pose = current pose (zero skeletal velocity). + self.pending_skin_groups_prev.push(matrices.to_vec()); self.pending_skin_groups.push(matrices.to_vec()); } @@ -11133,7 +11568,12 @@ impl Renderer { self.model_skin_scale = scale; } - pub fn set_joint_matrices_scaled(&mut self, matrices: &[[[f32; 4]; 4]], scale: f32, position: [f32; 3], rot_sin: f32, rot_cos: f32) { + /// `key` pairs this pose with the SAME model's pose last frame + /// (PT-7 motion vectors) — pass the FFI animation handle. The + /// world placement is baked into every joint matrix, so the prev + /// palette carries last frame's position/rotation too: skeletal + /// AND locomotion motion both land in the velocity buffer. + pub fn set_joint_matrices_scaled(&mut self, key: u64, matrices: &[[[f32; 4]; 4]], scale: f32, position: [f32; 3], rot_sin: f32, rot_cos: f32) { let cos_r = rot_cos; let sin_r = rot_sin; let mut scaled = Vec::with_capacity(matrices.len()); @@ -11159,6 +11599,15 @@ impl Renderer { scaled.push(sm); } + // First sighting of a key (spawn) has no history: previous = + // current, i.e. zero velocity — correct for something that + // just appeared. + let prev = match self.skin_prev_palettes.get(&key) { + Some(p) if p.len() == scaled.len() => p.clone(), + _ => scaled.clone(), + }; + self.pending_skin_groups_prev.push(prev); + self.skin_prev_palettes.insert(key, scaled.clone()); self.pending_skin_groups.push(scaled); } @@ -11262,10 +11711,16 @@ impl Renderer { local_min = [1.0; 3]; local_max = [-1.0; 3]; } + // PT-6: skinned VBs are also the compute pre-skin pass's + // INPUT (read as a raw storage array), hence STORAGE. let vb = self.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("cached_model_vb"), contents: bytemuck::cast_slice(&mesh.vertices), - usage: wgpu::BufferUsages::VERTEX, + usage: if is_skinned { + wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::STORAGE + } else { + wgpu::BufferUsages::VERTEX + }, }); let ib = self.device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("cached_model_ib"), @@ -11326,6 +11781,14 @@ impl Renderer { _material_uniform: material_uniform, shadow_cutout_bg, _shadow_cutoff_buf: shadow_cutoff_buf, + // PT-6 — bind-pose geometry retained for the dynamic + // TLAS path (skinned only; a handful of enemy models, + // a few thousand verts each). + cpu_vertices: if is_skinned { Some(mesh.vertices.clone()) } else { None }, + cpu_indices: if is_skinned { Some(mesh.indices.clone()) } else { None }, + base_color_idx: base_color_idx as u32, + metallic_factor: mesh.metallic_factor, + roughness_factor: mesh.roughness_factor, } }).collect(); @@ -11346,10 +11809,22 @@ impl Renderer { all_data[i] = self.frame_joint_data[i]; } self.queue.write_buffer(&self.joint_buffer, 0, bytemuck::cast_slice(&all_data)); + // PT-7 — previous frame's palette, same offsets. + let prev_count = self.frame_joint_data_prev.len().min(count); + for i in 0..count { + all_data[i] = if i < prev_count { + self.frame_joint_data_prev[i] + } else { + self.frame_joint_data[i] + }; + } + self.queue.write_buffer(&self.joint_prev_buffer, 0, bytemuck::cast_slice(&all_data)); } self.frame_joint_data.clear(); + self.frame_joint_data_prev.clear(); // Any leftover staged poses are stale (no draw consumed them). self.pending_skin_groups.clear(); + self.pending_skin_groups_prev.clear(); } // ============================================================ diff --git a/native/shared/src/renderer/model_draw.rs b/native/shared/src/renderer/model_draw.rs index 875bf66..b985745 100644 --- a/native/shared/src/renderer/model_draw.rs +++ b/native/shared/src/renderer/model_draw.rs @@ -182,12 +182,17 @@ impl Renderer { self.next_model_uniform_slot += 1; self.ensure_model_uniform_slot(slot); - // Skinned draws: mvp/prev_mvp are the BARE view-projection — - // the joint matrices bake world placement, so the VS goes + // Skinned draws: mvp is the BARE view-projection — the + // joint matrices bake world placement, so the VS goes // joint-space → world → clip without a model term. + // prev_mvp is the EN-022 velocity-reference VP (previous + // frame, jitter-cancelled): with the previous palette in + // the VS this yields REAL skeletal+locomotion velocity + // (it used to be the current VP -> exactly zero velocity, + // which is why enemies ghosted under TAA/TSR). self.stage_model_uniform(slot, &Uniforms3D { mvp: vp, model: model_matrix, - prev_mvp: vp, model_tint: tint, + prev_mvp: self.velocity_ref_vp, model_tint: tint, misc: [joint_offset, 1.0, 0.0, 0.0], }); @@ -200,6 +205,23 @@ impl Renderer { joint_offset, bounds_override, }); + + // PT-6 — register the mesh for the dynamic TLAS path: a + // compute pass skins it into the PT megabuffer and a + // per-frame BLAS build makes it visible to every ray + // (traced shadows, bounce light, reflections). + if self.hw_rt_enabled { + let (wmin, wmax) = bounds_override + .unwrap_or(([0.0; 3], [-1.0; 3])); + self.pt_dynamic_draws.push(super::PtDynamicDraw { + cache_handle: handle_bits, + mesh_idx, + joint_offset, + model: model_matrix, + wmin, + wmax, + }); + } } } @@ -304,12 +326,23 @@ impl Renderer { return None; } let group = self.pending_skin_groups.remove(0); + // PT-7 — the prev-palette queue moves in lockstep so both + // arenas assign the SAME offset to this draw. A length + // mismatch (or an unpaired group from a legacy caller) falls + // back to the current pose = zero skeletal velocity. + let prev_group = if !self.pending_skin_groups_prev.is_empty() { + let p = self.pending_skin_groups_prev.remove(0); + if p.len() == group.len() { p } else { group.clone() } + } else { + group.clone() + }; let start = self.frame_joint_data.len(); // Cap at the 1024-slot buffer. Overflowing poses land at offset 0, // which at least avoids an out-of-range read — the model will look // mis-posed but not corrupt memory. if start + group.len() <= 1024 { self.frame_joint_data.extend_from_slice(&group); + self.frame_joint_data_prev.extend_from_slice(&prev_group); Some(start as f32) } else { Some(0.0) diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index 98a9aa7..eac2a9c 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -251,7 +251,12 @@ impl Renderer { amb[2] * sky_intensity, 0.0, ], - size: [trace_w, trace_h, self.taa_frame_index, self.pt_accum_count], + // size.z: PT's OWN frame counter, not taa_frame_index — + // the TAA index freezes when TAA is disabled (settings, + // headless tests), which froze the sample sequence and + // silently stopped progressive accumulation from ever + // converging (found by the pt_progressive golden). + size: [trace_w, trace_h, self.pt_frame_index, self.pt_accum_count], cfg: [ self.pt_mode as f32, max_bounces, @@ -315,6 +320,9 @@ impl Renderer { // PT-4 ReSTIR reservoirs, same ping-pong pairing. wgpu::BindGroupEntry { binding: 20, resource: self.pt_resv_buffers[i].as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 21, resource: self.pt_resv_buffers[1 - i].as_ref().unwrap().as_entire_binding() }, + // PT-7 — velocity MRT (written by hdr_scene, which + // runs before the PT node every frame). + wgpu::BindGroupEntry { binding: 22, resource: wgpu::BindingResource::TextureView(&self.velocity_rt_view) }, ]; self.pt_bg[i] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("pt_bg"), @@ -466,6 +474,7 @@ impl Renderer { // frames — the gates downstream check pt_owns_frame(). self.pt_wrote_frame = self.pt_mode >= 2 || self.pt_accum_count >= 8; self.pt_accum_count = self.pt_accum_count.saturating_add(1); + self.pt_frame_index = self.pt_frame_index.wrapping_add(1); // ---- debug 16: numeric readback of traced intersections ---- // Copies a window of the accum buffer (center of frame) into a diff --git a/native/shared/src/renderer/shaders/core.rs b/native/shared/src/renderer/shaders/core.rs index dc07476..a7d8e89 100644 --- a/native/shared/src/renderer/shaders/core.rs +++ b/native/shared/src/renderer/shaders/core.rs @@ -318,6 +318,10 @@ struct VertexOutputScene { @group(2) @binding(9) var occ_tex: texture_2d; @group(2) @binding(10) var occ_samp: sampler; @group(3) @binding(0) var joints: JointMatrices; +// PT-7 — previous frame's palette, same slot offsets: skinned verts +// reconstruct last frame's world position from it, giving skeletal +// motion a REAL velocity (it was exactly zero before). +@group(3) @binding(1) var joints_prev: JointMatrices; const PI: f32 = 3.14159265; @@ -363,6 +367,7 @@ fn vs_main_scene(in: VertexInputScene) -> VertexOutputScene { // sway here — characters aren't foliage. let total_weight = in.weights.x + in.weights.y + in.weights.z + in.weights.w; var world4: vec4; + var prev_world4: vec4; var nrm4: vec4; var tan4: vec4; let pos4l = vec4(in.position, 1.0); @@ -377,6 +382,13 @@ fn vs_main_scene(in: VertexInputScene) -> VertexOutputScene { + joints.matrices[j1] * pos4l * in.weights.y + joints.matrices[j2] * pos4l * in.weights.z + joints.matrices[j3] * pos4l * in.weights.w; + // PT-7 — where this vertex WAS: previous palette, same + // slots. Feeds the velocity MRT so TAA/TSR and the path + // tracer can reproject skeletal motion. + prev_world4 = joints_prev.matrices[j0] * pos4l * in.weights.x + + joints_prev.matrices[j1] * pos4l * in.weights.y + + joints_prev.matrices[j2] * pos4l * in.weights.z + + joints_prev.matrices[j3] * pos4l * in.weights.w; nrm4 = joints.matrices[j0] * nrm4l * in.weights.x + joints.matrices[j1] * nrm4l * in.weights.y + joints.matrices[j2] * nrm4l * in.weights.z @@ -387,6 +399,7 @@ fn vs_main_scene(in: VertexInputScene) -> VertexOutputScene { + joints.matrices[j3] * tan4l * in.weights.w; } else { world4 = u.model * pos4l; + prev_world4 = world4; nrm4 = u.model * nrm4l; tan4 = u.model * tan4l; } @@ -394,7 +407,7 @@ fn vs_main_scene(in: VertexInputScene) -> VertexOutputScene { let c = u.mvp * world4; o.clip_position = c; o.curr_clip = c; - o.prev_clip = u.prev_mvp * world4; + o.prev_clip = u.prev_mvp * prev_world4; o.world_pos = world4.xyz; o.normal = normalize(nrm4.xyz); o.color = in.color * u.model_tint; diff --git a/native/shared/src/renderer/shaders/mod.rs b/native/shared/src/renderer/shaders/mod.rs index 6e3d8d0..9a096be 100644 --- a/native/shared/src/renderer/shaders/mod.rs +++ b/native/shared/src/renderer/shaders/mod.rs @@ -14,6 +14,6 @@ pub(super) use gi::{CARD_CAPTURE_WGSL, SDF_BAKE_WGSL, SDF_CLIPMAP_BAKE_WGSL, CAR mod ssgi; pub(super) use ssgi::{PROBE_HELPERS_WGSL, SSGI_PROBE_PLACE_WGSL, SSGI_PROBE_TRACE_SW_WGSL, SSGI_PROBE_TRACE_HW_WGSL, SSGI_PROBE_TRACE_SDF_WGSL, SSGI_PROBE_TEMPORAL_WGSL, SSGI_PROBE_RESOLVE_WGSL, SSR_TEMPORAL_SHADER_WGSL, SSR_SHADER_WGSL}; mod pt; -pub(super) use pt::{PT_KERNEL_WGSL, PT_ATROUS_WGSL}; +pub(super) use pt::{PT_KERNEL_WGSL, PT_ATROUS_WGSL, PT_SKIN_WGSL}; mod post; pub(super) use post::{BLOOM_SHADER_WGSL, DOF_SHADER_WGSL, MOTION_BLUR_SHADER_WGSL, SSS_SHADER_WGSL, SCENE_COMPOSE_SHADER_WGSL, TAA_SHADER_WGSL, EXPOSURE_SHADER_WGSL, COMPOSITE_SHADER_WGSL, UPSCALE_SHADER_WGSL, RCAS_SHADER_WGSL}; diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 4bb7217..c222692 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -112,6 +112,11 @@ struct InstanceGiData { // with the accum pair: (light index, W, M, target pdf) per trace texel. @group(0) @binding(20) var resv: array>; @group(0) @binding(21) var resv_out: array>; +// PT-7 — the raster velocity MRT (uv-space delta, current − previous, +// no Y flip at write; see core.rs). Non-zero where a surface MOVED — +// reprojection follows it instead of the camera-only prev_vp math, +// exactly like TAA, so moving skinned characters keep their history. +@group(0) @binding(22) var velocity_tex: texture_2d; // Reprojection of the current surface into the previous frame's trace // grid — computed once per texel (module privates because both the @@ -122,19 +127,39 @@ var rp_fr: vec2; var rp_zl_here: f32; var rp_nearest: u32; -fn compute_reproj(p0: vec3) { +fn compute_reproj(p0: vec3, px_full: vec2, depth_cur: f32) { rp_valid = false; - let clip_prev = u.prev_vp * vec4(p0, 1.0); - if (u.size.w == 0u || clip_prev.w <= 1e-4) { return; } - let ndc_prev = clip_prev.xyz / clip_prev.w; - let uv_prev = vec2(ndc_prev.x * 0.5 + 0.5, 0.5 - ndc_prev.y * 0.5); + if (u.size.w == 0u) { return; } + // PT-7 — object motion first: the velocity buffer knows how THIS + // pixel's surface moved (including skeletal motion, which no + // camera matrix can express). TAA's convention: + // prev_uv = (uv.x - vel.x, uv.y + vel.y). Camera-only pixels + // write ~zero velocity and fall through to the prev_vp math. + var uv_prev: vec2; + var zl_here: f32; + let vel = textureLoad(velocity_tex, px_full, 0).rg; + if (abs(vel.x) + abs(vel.y) > 1e-5) { + let uv_cur = (vec2(px_full) + 0.5) + / vec2(f32(u.ext.x), f32(u.ext.y)); + uv_prev = vec2(uv_cur.x - vel.x, uv_cur.y + vel.y); + // Depth along a moving surface changes slowly frame-to-frame; + // the tap tolerance absorbs it, and a fast approach degrades + // to a disocclusion reset — the safe direction. + zl_here = lin_depth(depth_cur); + } else { + let clip_prev = u.prev_vp * vec4(p0, 1.0); + if (clip_prev.w <= 1e-4) { return; } + let ndc_prev = clip_prev.xyz / clip_prev.w; + uv_prev = vec2(ndc_prev.x * 0.5 + 0.5, 0.5 - ndc_prev.y * 0.5); + zl_here = lin_depth(ndc_prev.z); + } if (uv_prev.x < 0.0 || uv_prev.x >= 1.0 || uv_prev.y < 0.0 || uv_prev.y >= 1.0) { return; } let pos = uv_prev * vec2(f32(u.size.x), f32(u.size.y)) - 0.5; rp_base = vec2(floor(pos)); rp_fr = pos - floor(pos); - rp_zl_here = lin_depth(ndc_prev.z); + rp_zl_here = zl_here; let np = vec2( min(u32(max(rp_base.x + i32(round(rp_fr.x)), 0)), u.size.x - 1u), min(u32(max(rp_base.y + i32(round(rp_fr.y)), 0)), u.size.y - 1u), @@ -1130,7 +1155,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // Reproject this surface into the previous trace grid ONCE — the // ReSTIR temporal reuse and the SVGF colour accumulation below both // consume the result (rp_* privates). - compute_reproj(p0); + compute_reproj(p0, px_full, depth); // PT-4 experimental: ext.w routes the primary point-light NEE // through the ReSTIR reservoirs; sun NEE is untouched either way. let use_restir = u.ext.w == 1u && u.cfg.x >= 2.0; @@ -1688,3 +1713,69 @@ fn cs_final(@builtin(global_invocation_id) gid: vec3) { textureStore(out_hdr_a, px, vec4((sum / wsum) * alb, 1.0)); } "#; + +/// PT-6 — compute pre-skin: poses one skinned mesh into its PT geometry +/// megabuffer window (world space; the joint palette bakes placement). +/// The CPU wrote the bind-pose vertex data into the window this frame; +/// this pass overwrites position + normal, leaving color/uv untouched. +/// The BLAS for the dynamic instance then reads the same window +/// (first_vertex offset), so intersection and hit shading share bytes. +pub(in crate::renderer) const PT_SKIN_WGSL: &str = r#" +struct SkinParams { + // Places the rare rigid (weightless) verts, same as the raster VS. + model: mat4x4, + // x = megabuffer vertex slot base (Vertex3D units), y = vertex + // count, z = joint palette base offset, w unused. + p: vec4, +}; +struct SkinJoints { + m: array, 1024>, +}; +@group(0) @binding(0) var sp: SkinParams; +@group(0) @binding(1) var src_v: array; +@group(0) @binding(2) var dst_v: array; +@group(0) @binding(3) var joints: SkinJoints; + +@compute @workgroup_size(64, 1, 1) +fn cs_skin(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + if (i >= sp.p.y) { return; } + // Vertex3D words: pos +0, normal +3, color +6, uv +10, joints +12, + // weights +16, tangent +20 (stride 24 f32). + let s = i * 24u; + let d = (sp.p.x + i) * 24u; + let pos = vec4(src_v[s], src_v[s + 1u], src_v[s + 2u], 1.0); + let nrm = vec4(src_v[s + 3u], src_v[s + 4u], src_v[s + 5u], 0.0); + let w = vec4( + src_v[s + 16u], src_v[s + 17u], src_v[s + 18u], src_v[s + 19u], + ); + var wp: vec3; + var wn: vec3; + if (w.x + w.y + w.z + w.w > 0.01) { + // Same palette blend as the raster VS (core.rs vs_main_scene). + let j0 = u32(src_v[s + 12u]) + sp.p.z; + let j1 = u32(src_v[s + 13u]) + sp.p.z; + let j2 = u32(src_v[s + 14u]) + sp.p.z; + let j3 = u32(src_v[s + 15u]) + sp.p.z; + let m0 = joints.m[j0]; + let m1 = joints.m[j1]; + let m2 = joints.m[j2]; + let m3 = joints.m[j3]; + wp = ((m0 * pos) * w.x + (m1 * pos) * w.y + + (m2 * pos) * w.z + (m3 * pos) * w.w).xyz; + wn = ((m0 * nrm) * w.x + (m1 * nrm) * w.y + + (m2 * nrm) * w.z + (m3 * nrm) * w.w).xyz; + } else { + wp = (sp.model * pos).xyz; + wn = (sp.model * nrm).xyz; + } + let ln = length(wn); + if (ln > 1e-6) { wn = wn / ln; } + dst_v[d] = wp.x; + dst_v[d + 1u] = wp.y; + dst_v[d + 2u] = wp.z; + dst_v[d + 3u] = wn.x; + dst_v[d + 4u] = wn.y; + dst_v[d + 5u] = wn.z; +} +"#; diff --git a/native/shared/tests/golden/pt_progressive.png b/native/shared/tests/golden/pt_progressive.png new file mode 100644 index 0000000..8cf9667 Binary files /dev/null and b/native/shared/tests/golden/pt_progressive.png differ diff --git a/native/shared/tests/golden/pt_realtime_motion.png b/native/shared/tests/golden/pt_realtime_motion.png new file mode 100644 index 0000000..93ccbbd Binary files /dev/null and b/native/shared/tests/golden/pt_realtime_motion.png differ diff --git a/native/shared/tests/golden_render.rs b/native/shared/tests/golden_render.rs index 92059d9..fb1c400 100644 --- a/native/shared/tests/golden_render.rs +++ b/native/shared/tests/golden_render.rs @@ -462,3 +462,145 @@ fn golden_lit_primitives_taa() { }); compare_or_update("lit_primitives_taa", w, h, &rgba); } + +// ============================================================================ +// PT-8 — path-tracer correctness goldens. +// +// Nothing automated guarded the path tracer before these: a transposed +// reprojection matrix survived three review rounds because every check +// was a human looking at screenshots. Two scenes: +// +// - `pt_progressive`: converged progressive mode on a static node scene. +// Catches transport regressions (BRDF, NEE, sky, accumulation math) as +// an energy/structure diff. +// - `pt_realtime_motion`: realtime mode while the camera orbits. Catches +// reprojection/temporal regressions — a broken history (the prev_vp +// transpose class) floods the image with unconverged noise and blows +// straight past the tolerance. +// +// Both need a ray-query device (DX12+DXC / Vulkan RQ / Metal) and skip +// gracefully without one — same contract as the CPU-adapter skip. On +// Windows, dxcompiler.dll + dxil.dll must be loadable (cwd or PATH); +// without them DX12 is FXC-capped and the tests skip. + +fn try_engine_rt() -> Option { + let mut backend_options = wgpu::BackendOptions::default(); + backend_options.dx12.shader_compiler = wgpu::Dx12Compiler::DynamicDxc { + dxc_path: String::from("dxcompiler.dll"), + }; + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor { + backends: wgpu::Backends::all(), + backend_options, + ..wgpu::InstanceDescriptor::new_without_display_handle() + }); + let rt_mask = wgpu::Features::EXPERIMENTAL_RAY_QUERY; + // The default adapter pick may be an FXC-capped DX12 view of a GPU + // whose Vulkan view traces fine — enumerate and prefer ray query. + let adapter = pollster::block_on(instance.enumerate_adapters(wgpu::Backends::all())) + .into_iter() + .find(|a| { + a.get_info().device_type != wgpu::DeviceType::Cpu + && a.features().contains(rt_mask) + })?; + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + required_features: rt_mask, + required_limits: adapter.limits(), + experimental_features: unsafe { wgpu::ExperimentalFeatures::enabled() }, + ..Default::default() + })) + .ok()?; + let renderer = Renderer::new_headless(device, queue, W, H); + let mut eng = EngineState::new(renderer); + eng.renderer.set_taa_enabled(false); + // Auto-exposure adapts over the accumulation window; a fixed + // exposure keeps the golden a pure function of the transport. + eng.renderer.set_manual_exposure(1.0); + Some(eng) +} + +/// Shared PT test scene: floor slab + a ring of cubes as SCENE NODES so +/// each gets a BLAS and the TLAS has real occluders (traced shadows and +/// bounce light are the whole point). +fn build_pt_scene(eng: &mut EngineState) { + let scale_translate = |sx: f32, sy: f32, sz: f32, x: f32, y: f32, z: f32| -> [[f32; 4]; 4] { + let mut m = [[0.0f32; 4]; 4]; + m[0][0] = sx; m[1][1] = sy; m[2][2] = sz; m[3][3] = 1.0; + m[3][0] = x; m[3][1] = y; m[3][2] = z; + m + }; + let (floor_v, floor_i) = cube_verts(0.5, [0.55, 0.5, 0.45, 1.0]); + let floor = eng.scene.create_node(); + eng.scene.update_geometry(floor, floor_v, floor_i); + eng.scene.set_transform(floor, scale_translate(16.0, 0.2, 16.0, 0.0, -0.1, 0.0)); + let colors: [[f32; 4]; 3] = [ + [0.85, 0.2, 0.15, 1.0], + [0.2, 0.65, 0.9, 1.0], + [0.9, 0.8, 0.2, 1.0], + ]; + for i in 0..6u32 { + let t = i as f32 / 6.0 * std::f32::consts::TAU; + let (cv, ci) = cube_verts(0.5, colors[(i % 3) as usize]); + let node = eng.scene.create_node(); + eng.scene.update_geometry(node, cv, ci); + eng.scene.set_transform( + node, + scale_translate(1.0, 1.0 + (i % 2) as f32, 1.0, t.cos() * 2.4, 0.5, t.sin() * 2.4), + ); + } +} + +#[test] +fn golden_pt_progressive() { + let Some(mut eng) = try_engine_rt() else { + eprintln!("skip: no ray-query adapter (or DXC unavailable)"); + return; + }; + build_pt_scene(&mut eng); + eng.renderer.set_path_tracing(1); + // Static camera: progressive accumulates 96 samples — converged + // enough at 256x256 that the residual noise sits well under the + // tolerance while transport regressions (wrong BRDF energy, + // broken NEE, sky double-count) land far above it. + let (w, h, rgba) = render(&mut eng, 300, |eng| { + let r = &mut eng.renderer; + r.set_clear_color(0.05, 0.07, 0.1, 1.0); + r.begin_mode_3d(5.0, 4.0, 7.0, 0.0, 0.5, 0.0, 0.0, 1.0, 0.0, 50.0, 0.0); + r.add_directional_light(-0.5, -1.0, -0.3, 1.0, 0.95, 0.9, 1.2); + }); + // Accumulated stochastic content: same seed sequence every run on + // one GPU; cross-GPU fp differences get a little extra headroom. + compare_or_update_tol("pt_progressive", w, h, &rgba, 4.0); +} + +#[test] +fn golden_pt_realtime_motion() { + let Some(mut eng) = try_engine_rt() else { + eprintln!("skip: no ray-query adapter (or DXC unavailable)"); + return; + }; + build_pt_scene(&mut eng); + eng.renderer.set_path_tracing(2); + // The camera orbits ~0.5 deg/frame: every frame reprojects real + // motion through the SVGF history. A reprojection regression (the + // prev_vp-transpose class) rejects all history, the denoiser gets + // 1-spp input with zero variance signal, and the image fills with + // speckle — far past any tolerance here. + let mut frame = 0u32; + let (w, h, rgba) = render(&mut eng, 48, move |eng| { + let r = &mut eng.renderer; + let a = 0.6 + frame as f32 * 0.009; + frame += 1; + r.set_clear_color(0.05, 0.07, 0.1, 1.0); + r.begin_mode_3d( + a.cos() * 8.0, 4.0, a.sin() * 8.0, + 0.0, 0.5, 0.0, + 0.0, 1.0, 0.0, + 50.0, 0.0, + ); + r.add_directional_light(-0.5, -1.0, -0.3, 1.0, 0.95, 0.9, 1.2); + }); + // Denoised 1-spp under motion: noisier baseline than the converged + // progressive golden, hence the wider mean gate. The outlier gate + // (broken-region detector) stays at the global strict value. + compare_or_update_tol("pt_realtime_motion", w, h, &rgba, 6.0); +}