From 2147ae1c8268836b0946a894e2978206d7a2d6f9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:20:15 +0200 Subject: [PATCH 1/7] refactor(renderer): extract GTAO compute into hiz.rs (record_gtao) Cluster 9 prep: the GTAO dispatch (temporal EMA ping-pong, Halton-5 rotation) leaves end_frame_with_scene; p22/p32 derive inside the method. Goldens pixel-identical. --- native/shared/src/renderer/hiz.rs | 96 ++++++++++++++++++++++++++++++- native/shared/src/renderer/mod.rs | 77 ++----------------------- 2 files changed, 99 insertions(+), 74 deletions(-) diff --git a/native/shared/src/renderer/hiz.rs b/native/shared/src/renderer/hiz.rs index d47815f..5db7447 100644 --- a/native/shared/src/renderer/hiz.rs +++ b/native/shared/src/renderer/hiz.rs @@ -5,7 +5,8 @@ //! policy); pipelines and the mip chain stay fields on [`Renderer`]. use super::formats::HIZ_MIP_COUNT; -use super::{HizDownsampleParams, HizLinearizeParams, SsaoBlurParams}; +use super::formats::halton; +use super::{HizDownsampleParams, HizLinearizeParams, SsaoBlurParams, SsaoParams}; use super::Renderer; impl Renderer { @@ -169,3 +170,96 @@ impl Renderer { } } } + +impl Renderer { + /// GTAO compute dispatch (half-res, Hi-Z-accelerated, temporal EMA + /// ping-pong). Caller guards on `ssao_enabled` and passes the + /// projection terms. Split from end_frame_with_scene. + #[allow(clippy::too_many_arguments)] + pub(super) fn record_gtao( + &mut self, + encoder: &mut wgpu::CommandEncoder, + profiler: &mut crate::profiler::Profiler, + half_w: u32, + half_h: u32, + p00: f32, + p11: f32, + p20: f32, + p21: f32, + ) { + let p22 = self.current_proj_matrix[2][2]; + let p32 = self.current_proj_matrix[3][2]; + // --- SSAO (compute GTAO, samples Hi-Z pyramid) -------------- + let ld = self.lighting_uniforms.light_dir; + let v = &self.current_view_matrix; + let light_dir_vs = [ + v[0][0]*ld[0] + v[1][0]*ld[1] + v[2][0]*ld[2], + v[0][1]*ld[0] + v[1][1]*ld[1] + v[2][1]*ld[2], + v[0][2]*ld[0] + v[1][2]*ld[1] + v[2][2]*ld[2], + 0.0, + ]; + // Temporal accumulation: ping-pong history textures. + // `write_idx` is the current-frame output; `read_idx` the + // previous frame's result. First 4 frames force alpha=1 + // so the initial clear never contaminates the signal. + let write_idx = self.ssao_history_idx; + let read_idx = 1 - write_idx; + let frame_phase = self.ssao_history_frame % 4; + let force_refresh = if self.ssao_history_frame < 4 { 1u32 } else { 0u32 }; + // 4-frame EMA: alpha = 1/4 = 0.25 gives equal weight to + // each of the 4 phases at steady state. + let alpha = 0.25_f32; + // Halton-5 rotation: uncorrelated with TAA's base-2/3 jitter + // so the two noise patterns don't resonate. + let halton5 = halton(self.ssao_history_frame + 1, 5); + let sp = SsaoParams { + params: [ + 1.0 / half_w as f32, + 1.0 / half_h as f32, + self.ssao_radius, + self.ssao_strength, + ], + proj_row01: [p00, p11, p20, p21], + proj_z: [p22, p32, 1.0 / p00, 1.0 / p11], + light_dir_vs, + size: [half_w, half_h, frame_phase, force_refresh], + temporal: [alpha, halton5, 0.0, 0.0], + }; + self.queue.write_buffer(&self.ssao_uniform_buffer, 0, bytemuck::bytes_of(&sp)); + + if self.ssao_bg_cache[write_idx].is_none() { + self.ssao_bg_cache[write_idx] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("ssao_bg"), + layout: &self.ssao_layout, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.ssao_uniform_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&self.ssao_rt_view) }, + wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.hiz_sampler) }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hiz_views[0]) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.hiz_views[1]) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.hiz_views[2]) }, + wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.hiz_views[3]) }, + wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&self.hiz_views[4]) }, + wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::TextureView(&self.velocity_rt_view) }, + wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&self.ssao_history_views[read_idx]) }, + wgpu::BindGroupEntry { binding: 10, resource: wgpu::BindingResource::Sampler(&self.composite_sampler) }, + wgpu::BindGroupEntry { binding: 11, resource: wgpu::BindingResource::TextureView(&self.ssao_history_views[write_idx]) }, + ], + })); + } + let bg = self.ssao_bg_cache[write_idx].as_ref().unwrap(); + + let ssao_ts = profiler.compute_pass_timestamp_writes("ssao_pass"); + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("ssao_pass"), + timestamp_writes: ssao_ts, + }); + pass.set_pipeline(&self.ssao_pipeline); + pass.set_bind_group(0, bg, &[]); + pass.dispatch_workgroups((half_w + 7) / 8, (half_h + 7) / 8, 1); + + // Flip ping-pong indices for the next frame. + self.ssao_history_idx = read_idx; + self.ssao_history_frame = self.ssao_history_frame.wrapping_add(1); + } +} diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 3678aab..86914b7 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -204,7 +204,7 @@ pub(super) struct HizDownsampleParams { #[repr(C)] #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] -struct SsaoParams { +pub(super) struct SsaoParams { /// xy = inv_size (1/half_w, 1/half_h), z = radius (world units), /// w = strength params: [f32; 4], @@ -9518,78 +9518,9 @@ impl Renderer { } } - // --- SSAO (compute GTAO, samples Hi-Z pyramid) -------------- - let ld = self.lighting_uniforms.light_dir; - let v = &self.current_view_matrix; - let light_dir_vs = [ - v[0][0]*ld[0] + v[1][0]*ld[1] + v[2][0]*ld[2], - v[0][1]*ld[0] + v[1][1]*ld[1] + v[2][1]*ld[2], - v[0][2]*ld[0] + v[1][2]*ld[1] + v[2][2]*ld[2], - 0.0, - ]; - // Temporal accumulation: ping-pong history textures. - // `write_idx` is the current-frame output; `read_idx` the - // previous frame's result. First 4 frames force alpha=1 - // so the initial clear never contaminates the signal. - let write_idx = self.ssao_history_idx; - let read_idx = 1 - write_idx; - let frame_phase = self.ssao_history_frame % 4; - let force_refresh = if self.ssao_history_frame < 4 { 1u32 } else { 0u32 }; - // 4-frame EMA: alpha = 1/4 = 0.25 gives equal weight to - // each of the 4 phases at steady state. - let alpha = 0.25_f32; - // Halton-5 rotation: uncorrelated with TAA's base-2/3 jitter - // so the two noise patterns don't resonate. - let halton5 = halton(self.ssao_history_frame + 1, 5); - let sp = SsaoParams { - params: [ - 1.0 / half_w as f32, - 1.0 / half_h as f32, - self.ssao_radius, - self.ssao_strength, - ], - proj_row01: [p00, p11, p20, p21], - proj_z: [p22, p32, 1.0 / p00, 1.0 / p11], - light_dir_vs, - size: [half_w, half_h, frame_phase, force_refresh], - temporal: [alpha, halton5, 0.0, 0.0], - }; - self.queue.write_buffer(&self.ssao_uniform_buffer, 0, bytemuck::bytes_of(&sp)); - - if self.ssao_bg_cache[write_idx].is_none() { - self.ssao_bg_cache[write_idx] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("ssao_bg"), - layout: &self.ssao_layout, - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.ssao_uniform_buffer.as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&self.ssao_rt_view) }, - wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.hiz_sampler) }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hiz_views[0]) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.hiz_views[1]) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.hiz_views[2]) }, - wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.hiz_views[3]) }, - wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&self.hiz_views[4]) }, - wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::TextureView(&self.velocity_rt_view) }, - wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&self.ssao_history_views[read_idx]) }, - wgpu::BindGroupEntry { binding: 10, resource: wgpu::BindingResource::Sampler(&self.composite_sampler) }, - wgpu::BindGroupEntry { binding: 11, resource: wgpu::BindingResource::TextureView(&self.ssao_history_views[write_idx]) }, - ], - })); - } - let bg = self.ssao_bg_cache[write_idx].as_ref().unwrap(); - - let ssao_ts = profiler.compute_pass_timestamp_writes("ssao_pass"); - let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { - label: Some("ssao_pass"), - timestamp_writes: ssao_ts, - }); - pass.set_pipeline(&self.ssao_pipeline); - pass.set_bind_group(0, bg, &[]); - pass.dispatch_workgroups((half_w + 7) / 8, (half_h + 7) / 8, 1); - - // Flip ping-pong indices for the next frame. - self.ssao_history_idx = read_idx; - self.ssao_history_frame = self.ssao_history_frame.wrapping_add(1); + // GTAO compute (samples the Hi-Z pyramid) — see + // record_gtao in hiz.rs. + self.record_gtao(&mut encoder, profiler, half_w, half_h, p00, p11, p20, p21); } // GTAO bilateral blur (or disabled-clear) — see hiz.rs. From 8d5d6f3dd5b798bcc7b53dbf195866998684c17c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:21:29 +0200 Subject: [PATCH 2/7] refactor(renderer): extract the translucent pass into scene_pass.rs Cluster 10 prep: the Phase-4b translucent/refractive/additive dispatch (back-to-front sort, scene-color snapshot for reads_scene materials, impulse-field update) becomes record_translucent_pass. Goldens pixel-identical. --- native/shared/src/renderer/mod.rs | 171 +--------------------- native/shared/src/renderer/scene_pass.rs | 178 +++++++++++++++++++++++ 2 files changed, 183 insertions(+), 166 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 86914b7..0f5816e 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -9259,171 +9259,12 @@ impl Renderer { // material pass on the inner graph) — see // record_hdr_scene_pass in scene_pass.rs. self.record_hdr_scene_pass(&mut encoder, profiler, scene); - // ============================================================ - // Phase 4b — translucent / refractive / additive material pass - // ============================================================ - // - // Runs after opaque materials, before post-FX. Loads hdr_rt so - // opaque output survives; alpha-blends into it. Depth is - // bound as read-only so translucent draws participate in the - // depth test without writing. - // - // If any submitted translucent material declared - // `reads_scene = true`, we first snapshot hdr_rt into a - // swapchain-sized transient and bind that as group 4 - // scene_color_tex for the dispatch. Free after the pass so - // the transient pool reuses on the next frame. - if !self.material_system.translucent_commands.is_empty() { - // Back-to-front by view depth — required for correct alpha - // compositing; submission order is only kept between - // equal-depth draws (stable sort). - self.material_system.sort_translucent(); - profiler.begin("translucent_pass"); - let swap_w = self.surface_config.width; - let swap_h = self.surface_config.height; - self.transient_pool.begin_frame(swap_w, swap_h); - - // Phase 7 — run the impulse decay + splat compute BEFORE - // we build scene_inputs so the front view reflects this - // frame's submissions. - self.impulse_field.update(&self.device, &self.queue, &mut encoder); - - // Does any queued translucent material need the scene - // colour snapshot? - let needs_scene = self.material_system.translucent_commands - .iter() - .any(|c| self.material_system.pipelines - .get(c.material as usize - 1) - .and_then(|p| p.as_ref()) - .map(|p| p.reads_scene) - .unwrap_or(false)); - - let scene_color_tid = if needs_scene { - let desc = transient::TransientDesc::new( - formats::HDR_FORMAT, - wgpu::TextureUsages::COPY_DST | wgpu::TextureUsages::TEXTURE_BINDING, - transient::SizePolicy::Swapchain, - ); - Some(self.transient_pool.acquire(&self.device, desc)) - } else { - None - }; - - // Phase 4c — depth snapshot. wgpu forbids sampling a - // texture that is also a depth-stencil attachment of the - // same pass, so we copy the opaque depth buffer into a - // transient before beginning the translucent pass and - // bind the transient at group 4 binding 2. Acquired - // whenever any translucent material reads_scene (same - // gate as colour) — cheap enough that it's not worth a - // separate `reads_depth` flag yet. - let scene_depth_tid = if needs_scene { - let desc = transient::TransientDesc::new( - formats::DEPTH_FORMAT, - wgpu::TextureUsages::COPY_DST | wgpu::TextureUsages::TEXTURE_BINDING, - transient::SizePolicy::Swapchain, - ); - Some(self.transient_pool.acquire(&self.device, desc)) - } else { - None - }; - - // Snapshot hdr_rt + live depth -> transients. - if let (Some(ctid), Some(dtid)) = (scene_color_tid, scene_depth_tid) { - let color_tex = self.transient_pool.texture(ctid).expect("fresh color transient"); - encoder.copy_texture_to_texture( - wgpu::TexelCopyTextureInfo { - texture: &self.hdr_rt_texture, - mip_level: 0, - origin: wgpu::Origin3d::ZERO, - aspect: wgpu::TextureAspect::All, - }, - wgpu::TexelCopyTextureInfo { - texture: color_tex, - mip_level: 0, - origin: wgpu::Origin3d::ZERO, - aspect: wgpu::TextureAspect::All, - }, - wgpu::Extent3d { width: swap_w, height: swap_h, depth_or_array_layers: 1 }, - ); - let depth_tex = self.transient_pool.texture(dtid).expect("fresh depth transient"); - encoder.copy_texture_to_texture( - wgpu::TexelCopyTextureInfo { - texture: &self.depth_texture, - mip_level: 0, - origin: wgpu::Origin3d::ZERO, - aspect: wgpu::TextureAspect::DepthOnly, - }, - wgpu::TexelCopyTextureInfo { - texture: depth_tex, - mip_level: 0, - origin: wgpu::Origin3d::ZERO, - aspect: wgpu::TextureAspect::DepthOnly, - }, - wgpu::Extent3d { width: swap_w, height: swap_h, depth_or_array_layers: 1 }, - ); - let color_view = self.transient_pool.view(ctid).unwrap(); - let depth_view = self.transient_pool.view(dtid).unwrap(); - let imp_view = self.impulse_field.front_view(); - let imp_samp = self.impulse_field.sampler(); - self.material_system.update_scene_inputs( - &self.device, color_view, Some(depth_view), - Some((imp_view, imp_samp)), - ); - } else { - // No refractive/depth-reading materials this frame — - // still need a valid bind group. None → internal stubs. - self.material_system.update_scene_inputs( - &self.device, &self.hdr_rt_view, None, None, - ); - } - - { - let t_ts = profiler.pass_timestamp_writes("translucent_pass"); - let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { - label: Some("bloom_translucent_pass"), - color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &self.hdr_rt_view, - resolve_target: None, - depth_slice: None, - ops: wgpu::Operations { - load: wgpu::LoadOp::Load, - store: wgpu::StoreOp::Store, - }, - })], - depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment { - view: &self.depth_view, - depth_ops: Some(wgpu::Operations { - load: wgpu::LoadOp::Load, - // Translucents don't write depth — keep - // the opaque pass's depth pristine so - // downstream post-FX (SSR/SSGI) still - // sees the opaque geometry. - store: wgpu::StoreOp::Store, - }), - stencil_ops: None, - }), - timestamp_writes: t_ts, - occlusion_query_set: None, - multiview_mask: None, - }); - let cache = &self.model_gpu_cache; - self.material_system.dispatch_translucent(&mut pass, |handle, idx| { - if let Some(Some(meshes)) = cache.get(&handle) { - if idx < meshes.len() { - let mesh = &meshes[idx]; - return Some((&mesh.vb, &mesh.ib, mesh.index_count)); - } - } - None - }); - } + let surf_w = self.surface_config.width; + let surf_h = self.surface_config.height; - if let Some(tid) = scene_color_tid { - self.transient_pool.release(tid); - } - profiler.end("translucent_pass"); - } + // Translucent / refractive / additive material pass — + // see record_translucent_pass in scene_pass.rs. + self.record_translucent_pass(&mut encoder, profiler); // ============================================================ // SSAO: half-res GTAO sampling a hierarchical linear-depth @@ -9431,8 +9272,6 @@ impl Renderer { // dispatch the GTAO compute pass. // ============================================================ profiler.begin("post_fx"); - let surf_w = self.surface_config.width; - let surf_h = self.surface_config.height; if self.ssao_enabled { let p = &self.current_proj_matrix; let p00 = p[0][0]; diff --git a/native/shared/src/renderer/scene_pass.rs b/native/shared/src/renderer/scene_pass.rs index f290cec..a191ae2 100644 --- a/native/shared/src/renderer/scene_pass.rs +++ b/native/shared/src/renderer/scene_pass.rs @@ -285,3 +285,181 @@ impl Renderer { } } + +impl Renderer { + /// Translucent / refractive / additive material pass: after opaque, + /// before post-FX; loads hdr_rt, depth read-only, back-to-front + /// sorted; snapshots scene color for reads_scene materials. Split + /// from end_frame_with_scene. + pub(super) fn record_translucent_pass( + &mut self, + encoder: &mut wgpu::CommandEncoder, + profiler: &mut crate::profiler::Profiler, + ) { + // ============================================================ + // Phase 4b — translucent / refractive / additive material pass + // ============================================================ + // + // Runs after opaque materials, before post-FX. Loads hdr_rt so + // opaque output survives; alpha-blends into it. Depth is + // bound as read-only so translucent draws participate in the + // depth test without writing. + // + // If any submitted translucent material declared + // `reads_scene = true`, we first snapshot hdr_rt into a + // swapchain-sized transient and bind that as group 4 + // scene_color_tex for the dispatch. Free after the pass so + // the transient pool reuses on the next frame. + if !self.material_system.translucent_commands.is_empty() { + // Back-to-front by view depth — required for correct alpha + // compositing; submission order is only kept between + // equal-depth draws (stable sort). + self.material_system.sort_translucent(); + profiler.begin("translucent_pass"); + let swap_w = self.surface_config.width; + let swap_h = self.surface_config.height; + self.transient_pool.begin_frame(swap_w, swap_h); + + // Phase 7 — run the impulse decay + splat compute BEFORE + // we build scene_inputs so the front view reflects this + // frame's submissions. + self.impulse_field.update(&self.device, &self.queue, &mut *encoder); + + // Does any queued translucent material need the scene + // colour snapshot? + let needs_scene = self.material_system.translucent_commands + .iter() + .any(|c| self.material_system.pipelines + .get(c.material as usize - 1) + .and_then(|p| p.as_ref()) + .map(|p| p.reads_scene) + .unwrap_or(false)); + + let scene_color_tid = if needs_scene { + let desc = transient::TransientDesc::new( + formats::HDR_FORMAT, + wgpu::TextureUsages::COPY_DST | wgpu::TextureUsages::TEXTURE_BINDING, + transient::SizePolicy::Swapchain, + ); + Some(self.transient_pool.acquire(&self.device, desc)) + } else { + None + }; + + // Phase 4c — depth snapshot. wgpu forbids sampling a + // texture that is also a depth-stencil attachment of the + // same pass, so we copy the opaque depth buffer into a + // transient before beginning the translucent pass and + // bind the transient at group 4 binding 2. Acquired + // whenever any translucent material reads_scene (same + // gate as colour) — cheap enough that it's not worth a + // separate `reads_depth` flag yet. + let scene_depth_tid = if needs_scene { + let desc = transient::TransientDesc::new( + formats::DEPTH_FORMAT, + wgpu::TextureUsages::COPY_DST | wgpu::TextureUsages::TEXTURE_BINDING, + transient::SizePolicy::Swapchain, + ); + Some(self.transient_pool.acquire(&self.device, desc)) + } else { + None + }; + + // Snapshot hdr_rt + live depth -> transients. + if let (Some(ctid), Some(dtid)) = (scene_color_tid, scene_depth_tid) { + let color_tex = self.transient_pool.texture(ctid).expect("fresh color transient"); + encoder.copy_texture_to_texture( + wgpu::TexelCopyTextureInfo { + texture: &self.hdr_rt_texture, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + wgpu::TexelCopyTextureInfo { + texture: color_tex, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + wgpu::Extent3d { width: swap_w, height: swap_h, depth_or_array_layers: 1 }, + ); + let depth_tex = self.transient_pool.texture(dtid).expect("fresh depth transient"); + encoder.copy_texture_to_texture( + wgpu::TexelCopyTextureInfo { + texture: &self.depth_texture, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::DepthOnly, + }, + wgpu::TexelCopyTextureInfo { + texture: depth_tex, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::DepthOnly, + }, + wgpu::Extent3d { width: swap_w, height: swap_h, depth_or_array_layers: 1 }, + ); + let color_view = self.transient_pool.view(ctid).unwrap(); + let depth_view = self.transient_pool.view(dtid).unwrap(); + let imp_view = self.impulse_field.front_view(); + let imp_samp = self.impulse_field.sampler(); + self.material_system.update_scene_inputs( + &self.device, color_view, Some(depth_view), + Some((imp_view, imp_samp)), + ); + } else { + // No refractive/depth-reading materials this frame — + // still need a valid bind group. None → internal stubs. + self.material_system.update_scene_inputs( + &self.device, &self.hdr_rt_view, None, None, + ); + } + + { + let t_ts = profiler.pass_timestamp_writes("translucent_pass"); + let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("bloom_translucent_pass"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: &self.hdr_rt_view, + resolve_target: None, + depth_slice: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Load, + store: wgpu::StoreOp::Store, + }, + })], + depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment { + view: &self.depth_view, + depth_ops: Some(wgpu::Operations { + load: wgpu::LoadOp::Load, + // Translucents don't write depth — keep + // the opaque pass's depth pristine so + // downstream post-FX (SSR/SSGI) still + // sees the opaque geometry. + store: wgpu::StoreOp::Store, + }), + stencil_ops: None, + }), + timestamp_writes: t_ts, + occlusion_query_set: None, + multiview_mask: None, + }); + let cache = &self.model_gpu_cache; + self.material_system.dispatch_translucent(&mut pass, |handle, idx| { + if let Some(Some(meshes)) = cache.get(&handle) { + if idx < meshes.len() { + let mesh = &meshes[idx]; + return Some((&mesh.vb, &mesh.ib, mesh.index_count)); + } + } + None + }); + } + + if let Some(tid) = scene_color_tid { + self.transient_pool.release(tid); + } + profiler.end("translucent_pass"); + } + } +} From cd83049634b2c25a6c335d841d9ef8006a676137 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:22:48 +0200 Subject: [PATCH 3/7] refactor(renderer): extract auto-exposure into postfx_chain.rs Cluster 11 prep: the measure+adapt pass becomes record_auto_exposure; its luminance source derives from composite_source_view() internally (it measures whatever the composite will read), and the caller keeps the src/dst ping-pong indices because the composite binds the same dst view. Goldens pixel-identical. --- native/shared/src/renderer/mod.rs | 51 ++--------------- native/shared/src/renderer/postfx_chain.rs | 64 ++++++++++++++++++++++ 2 files changed, 68 insertions(+), 47 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 0f5816e..c87472e 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -9395,7 +9395,6 @@ impl Renderer { // record_postfx_tail in postfx_chain.rs. self.record_postfx_tail(&mut encoder, profiler); - let composite_src_view = self.composite_source_view(); // ============================================================ // Auto-exposure update pass (runs only when auto_exposure is @@ -9404,53 +9403,11 @@ impl Renderer { // ============================================================ let exposure_src_idx = self.exposure_current_idx; let exposure_dst_idx = 1 - self.exposure_current_idx; - if self.auto_exposure { - let ep = ExposureParams { - params: [ - self.auto_exposure_key, - self.auto_exposure_rate, - // Wide clamp — without SSGI, Sponza's shadowed - // corridors have ~7× less average luma than its - // sunlit courtyard, so exposure needs to span - // the same range to keep perceived brightness - // stable across rotations. - 0.1, - 10.0, - ], - }; - self.queue.write_buffer(&self.exposure_uniform_buffer, 0, bytemuck::bytes_of(&ep)); + // Measurement + adaptation — see record_auto_exposure in + // postfx_chain.rs. Composite reads exposure_views[dst]. + self.record_auto_exposure(&mut encoder, exposure_src_idx, exposure_dst_idx); - let bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("exposure_bg"), - layout: &self.exposure_layout, - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.exposure_uniform_buffer.as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(composite_src_view) }, - wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.composite_sampler) }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.exposure_views[exposure_src_idx]) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(&self.composite_sampler) }, - ], - }); - let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { - label: Some("exposure_pass"), - color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &self.exposure_views[exposure_dst_idx], - resolve_target: None, - depth_slice: None, - ops: wgpu::Operations { - load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), - store: wgpu::StoreOp::Store, - }, - })], - depth_stencil_attachment: None, - timestamp_writes: None, - occlusion_query_set: None, - multiview_mask: None, - }); - pass.set_pipeline(&self.exposure_pipeline); - pass.set_bind_group(0, &bg, &[]); - pass.draw(0..3, 0..1); - } + let composite_src_view = self.composite_source_view(); // composite_uniform_buffer carries per-frame composite state. // x = tonemap kind (0 ACES / 1 AgX) diff --git a/native/shared/src/renderer/postfx_chain.rs b/native/shared/src/renderer/postfx_chain.rs index eac7864..4a9d946 100644 --- a/native/shared/src/renderer/postfx_chain.rs +++ b/native/shared/src/renderer/postfx_chain.rs @@ -625,3 +625,67 @@ impl Renderer { } } } + +impl Renderer { + /// Auto-exposure measure + adapt pass into the dst slot of the + /// ping-pong exposure texture. No-op when auto_exposure is off (the + /// composite keeps reading the stale texture, which manual_exposure + /// bypasses). The caller owns the src/dst indices because the + /// composite binds the same dst view. + pub(super) fn record_auto_exposure( + &mut self, + encoder: &mut wgpu::CommandEncoder, + exposure_src_idx: usize, + exposure_dst_idx: usize, + ) { + // The luminance source is whatever the composite will read. + let composite_src_view = self.composite_source_view(); + if self.auto_exposure { + let ep = ExposureParams { + params: [ + self.auto_exposure_key, + self.auto_exposure_rate, + // Wide clamp — without SSGI, Sponza's shadowed + // corridors have ~7× less average luma than its + // sunlit courtyard, so exposure needs to span + // the same range to keep perceived brightness + // stable across rotations. + 0.1, + 10.0, + ], + }; + self.queue.write_buffer(&self.exposure_uniform_buffer, 0, bytemuck::bytes_of(&ep)); + + let bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("exposure_bg"), + layout: &self.exposure_layout, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.exposure_uniform_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(composite_src_view) }, + wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.composite_sampler) }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.exposure_views[exposure_src_idx]) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(&self.composite_sampler) }, + ], + }); + let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("exposure_pass"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: &self.exposure_views[exposure_dst_idx], + resolve_target: None, + depth_slice: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), + store: wgpu::StoreOp::Store, + }, + })], + depth_stencil_attachment: None, + timestamp_writes: None, + occlusion_query_set: None, + multiview_mask: None, + }); + pass.set_pipeline(&self.exposure_pipeline); + pass.set_bind_group(0, &bg, &[]); + pass.draw(0..3, 0..1); + } + } +} From 8493193e922bada190c4f652f9996b68439bfe33 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:31:38 +0200 Subject: [PATCH 4/7] fix(renderer): TAA poisoned sky chroma via degenerate far-plane reprojection MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit With zero sky velocity, TAA reprojected sky pixels POSITIONALLY: the far-plane world reconstruction divides by a near-zero w, landing prev_uv on arbitrary scene geometry. The luma-only YCoCg history clamp (correct for sparkle) then locks the wrong CHROMA in forever — skies tinted toward whatever bright object the broken UV hit (uniform green with a green sphere in frame; reddish with a red cube). Every TAA-on game shipped with this. Sky is at infinity, so far-depth pixels now reproject the view DIRECTION (w=0 through prev_vp): exact under camera rotation, translation-invariant by definition, with an in.uv fallback for degenerate w. Found by the new TAA-on golden (lit_primitives_taa), which now pins both the fix and the TAA branch of the post-FX cascade. TAA-off goldens unaffected (the change is inside the far-depth else-branch). --- native/shared/examples/probe_compose.rs | 31 ++++++++++++++++++ native/shared/src/renderer/shaders/post.rs | 15 +++++++++ .../tests/golden/lit_primitives_taa.png | Bin 0 -> 27953 bytes native/shared/tests/golden_render.rs | 26 +++++++++++++++ 4 files changed, 72 insertions(+) create mode 100644 native/shared/examples/probe_compose.rs create mode 100644 native/shared/tests/golden/lit_primitives_taa.png diff --git a/native/shared/examples/probe_compose.rs b/native/shared/examples/probe_compose.rs new file mode 100644 index 0000000..eba063e --- /dev/null +++ b/native/shared/examples/probe_compose.rs @@ -0,0 +1,31 @@ +use bloom_shared::engine::EngineState; +use bloom_shared::renderer::Renderer; +fn main() { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor { + backends: wgpu::Backends::all(), + ..wgpu::InstanceDescriptor::new_without_display_handle() + }); + let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions::default())).unwrap(); + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + required_limits: adapter.limits(), ..Default::default() + })).unwrap(); + let renderer = Renderer::new_headless(device, queue, 256, 256); + let mut eng = EngineState::new(renderer); + eng.renderer.set_taa_enabled(true); + + let n: u32 = std::env::args().nth(1).map(|v| v.parse().unwrap()).unwrap_or(1); for i in 0..n { + eng.begin_frame(); + eng.renderer.set_clear_color(13.0, 18.0, 26.0, 255.0); + eng.renderer.begin_mode_3d(4.0, 3.0, 6.0, 0.0, 0.5, 0.0, 0.0, 1.0, 0.0, 45.0, 0.0); + eng.renderer.add_directional_light(-0.5, -1.0, -0.3, 1.0, 0.95, 0.9, 1.2); + eng.renderer.add_point_light(2.0, 2.0, 2.0, 10.0, 0.2, 0.4, 1.0, 2.0); + eng.renderer.draw_plane(0.0, 0.0, 0.0, 10.0, 10.0, 120.0, 120.0, 125.0, 255.0); + eng.renderer.draw_cube(-1.2, 0.5, 0.0, 1.0, 1.0, 1.0, 230.0, 41.0, 55.0, 255.0); + if i + 1 == n { eng.renderer.screenshot_requested = true; } + eng.end_frame(); + } + let (w, h, mut data) = eng.renderer.screenshot_data.take().unwrap(); + for px in data.chunks_exact_mut(4) { px.swap(0, 2); } + image::save_buffer("/tmp/probe_compose.png", &data, w, h, image::ColorType::Rgba8).unwrap(); + println!("sky pixel: {:?}", &data[..4]); +} diff --git a/native/shared/src/renderer/shaders/post.rs b/native/shared/src/renderer/shaders/post.rs index 0072b9f..a8c5733 100644 --- a/native/shared/src/renderer/shaders/post.rs +++ b/native/shared/src/renderer/shaders/post.rs @@ -792,6 +792,21 @@ fn fs_main(in: VsOut) -> @location(0) vec4 { var prev_uv: vec2; if (vel_len > 0.00001) { prev_uv = vec2(in.uv.x - vel.x, in.uv.y + vel.y); + } else if (depth >= 0.9999) { + // Sky / far plane: the positional reconstruction divides by a + // near-zero w and reprojects sky pixels onto arbitrary scene + // points — the luma-only history clamp then locks that wrong + // chroma in forever (uniform green/red sky tint). The sky is at + // infinity, so reproject the view DIRECTION instead: exact under + // camera rotation, translation-invariant by definition. + let dir = world_h.xyz; // w ~ 0 at the far plane: xyz IS the direction + let prev_clip = u.prev_vp * vec4(dir, 0.0); + if (prev_clip.w > 0.00001) { + let prev_ndc = prev_clip.xyz / prev_clip.w; + prev_uv = vec2(prev_ndc.x * 0.5 + 0.5, 1.0 - (prev_ndc.y * 0.5 + 0.5)); + } else { + prev_uv = in.uv; + } } else { let prev_clip = u.prev_vp * vec4(world, 1.0); let prev_ndc = prev_clip.xyz / prev_clip.w; diff --git a/native/shared/tests/golden/lit_primitives_taa.png b/native/shared/tests/golden/lit_primitives_taa.png new file mode 100644 index 0000000000000000000000000000000000000000..cfa8302740d36599a225b9a5f97a57efebfad958 GIT binary patch literal 27953 zcmeFZ|6j}d|3ChCJhs+WOKVX}$z!$hEn3BdXsu`|70x-4WPMR0lzdr7wAEr$F@;K< zh2z^fNaRb=(W;5G%y&_~4a+&4m^p~l`}W9rKHsnR=ko`=e|cYy%Sm6f`?zJgXSR!3V%frDg{2otT~#5AgP~2f&zXJqkpSO`#1~tv=|3A z^b}sz9$6EbP_?r2+`oTP3RHV$pHzAKxb*a{KfnCv!&|4VxO?)Fp}b*h-Cw^S>-Tr5 zB|bF#zYqVnj^MzWyFcn?cIkp7pHD8uZTU#)vHs#nxEyna$dMM^bh%k1WDHEAe?X_@ zAC#R6U>bfOApcPHB|F%!JdvXcrkSQi{QdMTPtugZL`13*#Ps&|Q@8EnZ(9q`3_aJV z-wZ#G6F)ywvJy#Ma3Xi2&ZV*A>FUiFpTBPSxV~qae=~o4$nY&*Vy)}P-wqkABc{%} z?^d!f%qsTVu(!{Di||(@x~O^0qt|<`E42?--;qgtj!4Cn~x%Neh}6)p+lM|h@57WjcS+9@B6m;q;%M$eCzrz zfx*#h?q)j?FO|^TA6I?$fw%Udau5>cxO|wYm#URGh@LXbEZ-nx4l)l@Rz1aXBKNXY=1elg*ijU;jh<3w5akE7>ihFi{Np)GtAoq#j&unVV_ ze|2&mVZ0jrq-E5Mr1*OyPH+4(y62D8`Np^Q9w%3(uH>An{p(5L;7ckgqr1f4ip(En zS>Laiaqu3>*tY$I8DA3oaf@a~{9Jn(#%(XWX}d|wZO!mzF?3rQQ626>!=^CD%5r`+ zGCGC33p1IAejRtDwf^e!$}Eh!HvrTAIIwvzl45SIYi(&+oVv2T?C|WQU+cVw*dj;q zu(@dyN5s*o-1zVdZ1P@4>oC2y$&XA$anYJz51lHZqBrQQXu%oURMwGezIKS_#{ucl zn12R-iuRqn@Z=6g-Wll@4#acfT2gtk?S8lU$8a{iv>Qi^d<>(KDWvuBOu9fO8^13u z@UWuR_D?n(@OIek4%z~I1R}Yd;f=`Vg%Hp+YKz^>sQR zB*SLIl(Os*BwSTq_m{uf?qmHUeFkha**qzH>b(0{)z4B*wzDaxwEJEye82I>s6Y4y z|GeHPZ+Jz#_OF5I8UCi#as9uRC5OHnv*naIoacUO)ddR@f@w`l5tE-b8#AUR?cHvS z-*L&D5%p~=B>TWdNXaF0`mT6OGKY9?c#=_ZtiRa?j%{!625((u?|B>DWm(aqKO96= zU#w@Q-5p54)b%1mLuuv&tRNA-7MynXN5t?%-Pi3cLO4}lb_{+&|3KTYx304HzL%#( z`LX`z#BWp~T1%0tICHb6VQ1Z64c3i64!nKPH%{Jtr5-o*dgp_z1FU=tc<$@z7TN{qf~SB{=3?1 zN(V4XBHzzE*Gyl>Vn`VJDW&#W37ob)+{S!bwUtDg9+sN>`IDFMz%gm9fNAjGiW*nj zQxo`paA^(0;Kv?mYC`<1LxlP_kBb@_l1J{b6{e-jer%dFUDvo}+D~a0wj$g1Su=E; z>bSX?#@FYj)YPbXh(51a#(}emrD*b7BnTHr=nyt7hy*dEQq+D*FU-*u78>yK_iTIg z9LnWxSFB{{;+ucjIPS4w_sG_^kyX`J%*jhOY^$B4P?)E*5 z8-#x48k17S>A863#mJsVTVIlyF)X$f_Az^_OPeI{A%v%OG7$4xZ$?fH3%gHFV+3r)A!Kx<@LYs=h%+uIw!k5og7Kt^v(m@_Dqj; z#MLdAFFv@X-4w_7vnF}ts&Cwqe|Pa-+bRllq1sIy1(+6B1eSu zytKoxkj(4aoYFe&(!&JtBW~YsDpqZ_x3?a7JuiP_{fGSNIeTan+QEJI znli?ntuEDZqkP-H73>TOdbeRmw#Di%L*8B~d%j5cZr+Z~)bo!rM_PEzDa`$KSW#`| z%ENY9Z_igIXC&a8EWSG3VqPTWjMLO8>=mdKMMeA$i5P(lqf=WY49fUZZ=XJW6|i>F z1OM7RaijUaQd1!#;MS6DbnZrvBBf?i$@b4%H(HCPNF_=9w6D(Zj#T_y6ZK+`V^v(m z@Yi8UHLBaA*4Mt=vf|X7Amql~ZBQhA=2&rc?wkN>dR~q>gPG{QW0NLO{GoXynLnD< zveW!Y)`)Yb;ws!LCos5$ty$h_g!w#y!;F8?^dJ{q&G4CzEz ztScHGHMREJt)dNCXRo}oFb}4(6$_VylyB`?JHBwk({mejroGmGl%-VQR?Oe?j_$5b zcQ0$bGkg|8MKb)+mX29=h)A2^jVr2N%Nat^{N-_8m2nWX2JZyUGg(HrweR}lfq&b0 zSXU*-ED3J-yy*0Z;Mr=4()~cikzGef88{ztRq~(J(U5e1n+1m%LM&I~gLPe6 zYVoS&3-p`EYktUahl-fkKeLI28a4~+BXg8o7sn2PH!^RHzw1NaU+?Vv+b8>4ys5nP z<5;}LPdvxqm>p=Kw*pv^1bgLccFtQpb866x?X^f|#k>h#g;)4_Uwfvt{yE>i_rmI| zF`IivjieW-w?$>`c3r=H(rKDwg-Dvr2b6%^9&uyErQVID0|^9*num04_*0%;`w7`& zOReSRNLtfAU(7PvXFD7spD1e#dGe{k_uNc->DTQ?hjntQkLNvqEj53#e$Rp(h?7&N zXvHV_G8#r-nDgr|W+%wr5>&|!t?n7?;Bsq0+QhdD?+WtyrC(d|a;|P(s`FL9YzzP= z5iu_*GiOkS7S1zQqSAnyuaYi~c31D)!qw?#YxsY^-9L0bXHAE`XVF_jRKu$Fe>`t@ zTu6HQVSK3SEf%t*hlvI&jRL4 zFGz@Y{2$1bm+N+&tSCR}KS${I&&Q2*t9zIATu4(yWhsJh9;V$AHl2-}9eLmNl+)-k z`q{K^H}|&Z4y&=;y0&d+`mfzD`e$CdU^AbD$X47KNyk-z8zy?UkJ~U^AN}#1VeNd9 zrZseHjPJRShT$Pb4~KYH;L=5C7x7UQ!pf^GqDc zRZY#vtJMZ56jLv5IlAHidFAx({0aRtA~rW2Q=c6-^ zX6qyr@kLU1rD)Tpm)h>IGbQAv(QLEV57YQdEcF&}y|L7>1D1;9VzDzI;e$Wiva@wF zJVuSqIWhvt)|8#k8n&wL+3>^nN)HnS<^2mU=e2GB9CyWYmdYU)QMdN8N1JZl*kOCf z&7m=TMe3WA@-cIVhwl1^OuaAq%i@!*`{5T-a-x;eH=}!Ic*){jNQ>Li`QCT=DLwWn z(N;gzY_|(I7*O3a?sYVuKit@)u>U1dPBZeG>=68xnZt5B#_={yfY)-%z0*FPHSC-A zVy<=ebXDJq`cxqZDoo>Qx3dHiz*%>`mXdy$SXf&8y34^BKVf zeMUXzH!)K?=}PUYjigRQZhzgdP&+W~)ubnV)3pOGSSm>*s$h2d<*T<$`?RONeVqTc z+~#41ZH_0Ww4O~4f2TR}sPR~XKGb`=?YW)%?roo!vw8XG;!~4urNw8r9y&24*tEU) zYv?NR`#8}!rs?-M!4v!Ip9G#dn8av5{^)b()Ht_k zDZd95SIZD|h>hIrKymFv&#w1ZE`Hd+oA5_d#V-x9HRr7+(|w%RtlwI=^h(yT?ym#y z-)6r1NMVX%aoUAwMwGCruR-uK8#kshb%D$x-;3v8ynXus7zAP;TuR4=WS0PgXnVC{ z|JlUlg%6G=uSm!($^I?aCuCRQlmvI%>L|-AKauV~-Xwj{Xz}39FJ18)KgoZ|`06XW zh3}j+Iw2mF-y22+!Z6v0IVh23n*xSvhj}lXO_nGN5$;ZdW^svcnR{sU!-`X}WhY*1 zpIxu`?c@RhZ(uzJe2{O4NK4OdUwYoqm0J1e^b3b?GX}U9)t4wkR`$e`hmYQ8SIs{V zd_14U5x5{BWC7;7CubdUOtCRG%Qlv|{p97m6007CVnruqvT9oM*^+Z(#&&(OTB6=` zC3ZykhqTwss>(onp!_7q-S!RZxB6VM>!9%Km!+yM_nn;iGMn`u5~rFe0umQWkhsf0 z;W$$4Y3M~pjY8u${_SV11m)g>XIcW4r|`>lhQC#qWxk^Wj- zn6mVttHst$Zf=SFzi3Ph5hIGq*8EPd`1@nsdFOZKjuFp1CrZp7sQH%FYFfCG&lO7e z7>y>QDi9?D52Ec~$UczvL%^bpJ&OaMQ#xzLhmV}rSUJ#cEABr3g?lcdJ%!_G)LR=u zA{uh#v5k_mM z&%dMU%Qo)0lIABbx19gt{Uv%?S=jsp-s>sr^sLEG#N*E?sNA(@;~gze56w!GMUj?P z2t|*>8a4z?aT;?>anvl&>nUT(F{DFM8)oRWjRFbEQO&Fuk{SBVDLcnR=b$LXU9oQ1 z)EPU=^Ss@30ex3aQ-j#Zf)y|3c{~^wIw7U?$(4ouE}t4{D~H&1D6LYf>%uS^N6E)z zW_yH+HVz+`yYHF*>+Itz;(i)?Kk1pnI(_%`uFGw!o!{T`{FoN|_a3u^$v3wxJUeN} z&#%TG-zwFTgJeoF&(vG6x^J)Nv+I{vsDB&Mv!JQfSrBQq>75fj^@rHL3oEZXbrO1D zW`fDX&@f{6fcqe1IN5gM)z3bLhCVjlb#5EyTxeu?qR;@EHaB>v5Q zBe0MTp}^+OJVZFE?C@EN#4S=EEpK+4y!qPdYSRSISr1gjyc6WDocxDghfw7lj!=M^ z!kC`BK1VS`&R40USnRSrF7uBc!L8T{bUyO~d$dbnmku+I42?N;Pwr*khDQ9=^X&fD zX(`cbPFfzf^u`IyMXeA(ZE||dRIK^ zL5$Yn*xV+=@_>e$3FT){B|cQBf2(>o^znxgk1KtyOjb~1-`});c5d2RH%`lM>D4i8 zvi@{fLV4lkg38|FQEQ%G$7tLEm;XS_`2+7O|07ba_l)oTy78lEN&wk!)5y;W*Xv&% zOKU5<&z)m?`%Y>5&*7E6E$x$I^H;iH^uE?8+!~N&8aDHhNBJ zfonCH+0l1}fpI6gnS0a?4UV~0wCMHB%{7m&FPsoyw`_W%nPckUSrJFY6yhiZ6a(@K z4s1ektw(W4K%1bXasF(V>s^Q$;EJv~g&8-V7)BNr=;_ZX% z8(ts7S7H)KPHN8DHFLjqIbO_KH1O%e)E7gL;k7v=dWT32p?CF#H)_GyhoH(MGhj^oQQM$rvxYhQjOhO7(<(#Dx94|$st&uh0GUP2lT*p8W%HX^2tCE* z!4vMH_V>^tp1oDA8I`0hfW`06|x5T>s_eLmp%#rr~Z%7fCwiM-do zmulub`Eh8#7^zhUtJ;b}F17zf%{(~9{?oho7X9kgOFy15o%!2gt=SR4o?NMX-$}DP zbqjx2)A5r-X`L*pVxDCYsQ$&{cQJ(`5{Gwi=G59H!*3l}=*za6CCbP{85A`yiiONk zTzvX!%lK$}r;5k%e_8A_JPDQ0j!V=lE)?tD_B&mhA(nI@8$?noUWX{xikFzvz7(QU zN@jxHMZ~F;j4>eWrDTK|#2-33ASlxH71lD<>*)KRUi;K{-E;^)dHz+5oc`sfosIfT zp#ahM5;FDn9WYEdThbHHMi=us{u7^=A%Duq-kl>Url_sAr+1x_tw`otg6enVP5m$Y zfoGKL<;mEc?fO5<Te(O-xJm4@3InEb*X>irl?H)?ViL>@8-sIJbHd8_RY>S z_Yh@}P}syg;X6~s7)SGrf^YWw8I{253#kSjmv*O7=W5|O;U9(cot@=ar z?6)}{sKH;r;Iv+f`j9DYw0h5CrRDs4jt2=?lp{0_RZ1QMxJXow?+}S#D6=Cbj#*0g z9ocP;v)JU)n)z}j{{N}AYcRApKs{okx-z40&XS#rR@H7vUk>$l#%teodf&-U(ujJC zT&Qa1b+Ckv+F=IokiKK`aT*_5h4|jD={Xa_XolXJok2VKO&Ei93qa9kNP1eW;ngxV zLb;3ImXT*4_&9p(d`f~P19*^x8&zr_B*_e)Ai%a>odJ<0&)&hu7i>)Q+)=wDlx~$~ z9$wwJWg^6J_*GfhG8#`9<1QuJFfg+q{$Pz6ti&`jPMoNB6iY()K^bIymi+6r8QFLL zes5@Ys9N}U+E-h^yk7P z3^x;{D*5u>_?E>>Om1D1ve)@OY90|9^T?rqzoEwMlN&rNjZ`3??}e` z16x}t8%E5|Z8w4b74c)-HA_*P%wkt9*kQ}-z!kT0;!0Utc0L^6(be0PjT92EBgQ>|93-&{ zpHRa_CtW4eT)tPIe^VP_${H#5qa|}N*Tc*YLoD?CP?b|?>?F$JWiryj1z^Iq@p!t+ zhy-QaSU18(KuuUa-oxMnD2Q-SWR_h1^nVoZ^Kyo#@u<1!1+drUECpa|aY|}UU8!+? z_WPh!4;Mll_U;GR|5pKqhW7VAp6GMi{{Qk!}u)azUQ%a-9+8o#=T9J`^ zd(`&Z-oGur%wrXlML##&`8kR?FqSPO33#7;qU7e@coi4tN>0J4Ap#7Qx24{S6x}Rg$?Az!5Ueup}5KZ!xr7K6hnFM8wOB-+mC;$Kb)sXl0^6vXFtYQ0yFxAF@6YUS>5}|3nzmL6YytK zsZuU3tE==XDYL(Z7^XOIOUk|~#*w}8*dMSQJY+&aiEz}DM=Ht}hpSOry0GC?C5|o6 zV=_>t498sfF*_-a-fW^oSj|^YL@97qff#c%AMlJQ6E>1YiRDl=MbyLzSodsR+2_^= zE8n#2oxEbidnRl-&^&{M)a%a|bs_bL-ha1d>7g5s)(ttUF19?-PdRD~;{lIE zxfD|vv|`M>bZvf~8UDsTdG-h|PoJM|?={_?QMvI$e&l_|N%G!6^gs?7$|02_moP^2 zbEvAzlu-v8w}e%*ZD&iX1j8#tfy1M7tRSRN8#o$fP>vNPwR->12{j>(=F|kY2*vcC z6XsnT`XpE0EQxqBb?aocV2Sinu)Ub{QqAT-^PD>Iya`a_DX0&hwHNOUbC3{`OI@ECqNVO)bND>D19#Km^c!J z&fA6YcpaF zk>V82+<_M%$wC{-!T!%ld;-5CChAF%tFt-$U{XO97&-Lm$5pM%Cn zF_5KS;!dt5chvr*R2!ZQ_E%Kp#7l9w5--oGsGSHg^_705p}58}xr!0_;{wCpn%gr! zp4GSh6U)0W*)V!|kaRR2Ko;TwgXM?VuF-=B$mv+|R5GzcNh4CCOu5BFNlSqUh<@+k zhJXPh#YAZ+i7Ht<$$80jQ=WZnMcR$IDdBsIx+DDZYX$X)!5jA5kSlHkOTw{}Fq7qs zPx*-C%j1&+`k3r?8uU~)W$NcwHjcaWA)zd~nOCIDiMB+_LJp^#K=s~^#P(wJ{#(>Y z&P0*gFW-u(YnF)FbO1tSZ6)W8hS$;8MmJ1hd=Z;1lkVUq+Z9%S2k$1O)h`z9&Hu~1$j z)%vW7o6kLNZ#-aZnK0D+8Q*SC?37iRBACcf3l@ z7>e?ANwEx4seK`z2}ILIRK-A;q6HD8l7}-Qp3mry>+e(_6AXpz0BnDUB@D#6REE%Q zNvwFHK>=3HCJk=uschoGMp_x{Y_!>}r?T{%dmm7NJ?-E>r>$*_e^71vY#-x~rAXr&=Vt{9JWC{rM0^EKH{ z#*qR)$>;mFc&x808DRmD#*Kb-Dm0Bw2Vi`^V%bmxu4_l5pS{U1k*f>!D@OQzmO*oG zc-P(d`1#czm3b?8`RAsr4U=udlu{K8nT)}tvZV;)oR_;*OUGe{f)w7l6kZCXgfGtge(a+=kx%Vd@U8A2TYRKc_s^ZTu(NMjLbg@qN$D;`!PSV#= z4nf6%1Sz=j#$xz9&V@Ao7nFx+Rje@rSIJ$%pRob1@Z*{ajQvYP8 zpG0hDe9P$>73*tkq9Tca8UQQe{oq+}XUV&lssR|`-<2l&=v=^x(VE->MlXBdJGtlm z&Azdz=6~@Nw{CA`9H0sUgt8zihaCg~f>3EPgpY%&Y3q;_pq(8p{D&RxUMb*YrSRss ztlT^2;Zww|Llg-1!+h)K)Z64TzV$>^F}t|Y^=`m^k!qyzfS0aBjKcH?9N8SdQN}I} z$$`;QA4$Z3%8l}4I2lEnidA7MDLsycC=L<^hTO`m@BCy-rk)zrQ<){tD89w=lVp$2Qt(i7d z4!>R`V2&@&T+2lsxgYxJ%V;X0*j+d>Dvib$;xaO>@qyl-BLf_B=}M0T%Qgs%B3k(w zT6uLy`Dpi*v$q)zJ~&wD#?U3Ia-}J<-vZp%PS?2vyspi6y}o=$;61;l$v1A_6vlE* zCr82mvWQY-gOnzF<7&Uj=|l*J{SfpYCs5?KZ)7U_BFuvxW(XB_e1jB*?@! zAuNFWz)O^%>b`P^qjMuiF+NO8Iutn+k;ym$0l=|9!XXPWK9#J`T!`g*2|HSt9ScN` zZ8G*s4_jW!;2%XwT%t|*t)$Z-2x%z%I&Ku5ba4M1Tr~#Co%Qpm$|VC8Z-+n0iM%0< z#q+bvGH)vEwe@^`85#QK+6*sGWVd|~(AF)rm7ni3!KNOk=nloyOezfFv`}IE01@SC zEQ#@xm9e-fMr-ocwX@=?xfq$s5cn-&b~tShUEet)yZC&waqGTS;e`Sraokksgj1C= z233YA#qMlXu@|+PD{@Hnh;t0>Xz@N>J?Glh7T)Qdg{e*Mskd&VR+mZ6p!*P3%zhI0 zM&rcBJ8Hk;09zl$A@kVH4`)%?psS6k6vWsEy;IMM%$ty1v7?O@9x&13J{Ex}X{dx# zKS`W4ML*L8#dAZIHZn(Ds{~O?KHIYxNC%Xd)Y)UgZi9E>qe~hLd^?k;Hzhp}F;#HL zVc~X5!^iJltA69VXz45kPT=$49yl^30=1owKsrRR_$+=4#Z`Fs>*Ad%5$K|Ts->Uutf%;Nib znhw}_mS#HZ>!=8Qm?C>5MEY4x24lbiTE*GZ;%Fq<=gB^JEFQ;4Bq_Ip8al|9(BSbi zH+h19(K zEvTm3?K#!xea@yO=Zvz@tBEsfcA_%(47<>5hz~{>bcbN^fH39zCRI^;CXgatgg^cz zO5$=$K@D>OZf80GgJ-lf{r3}~U588tT&^DvRd2bb!>W((hS1im;4L1pGMjVT!! z>$&USY^Wz{?ZA3HK%tKWssr$WCLk_8v*c|RNQWWMn@Hcxn^cA_Vzjbol=X&pX17N+ zwT@UHUwi=Pa0rOB!Pp#BfB{URFbRk`Ft7r%(_D(o%dVrzo5zZC=)CIg)pU?fBm>h) zq~<^5P#m{4Zec#fCN=U`hyqCybNY8RdbYD{aW&W*3$>Ow0~P zK_nr&yYZe;5I{>IdhDfe2{T<|abuReEfE431^(WFtMXA)LQr^$qB|~EQ&VPNnCZ;Z zHN#|J@=iDahJcgw#_UlL{1t>}d-Z2wcorfWl zim+oD3H@&ROb22+5^=xQ6ks-(0890mitfO^PKCV?6m_rdrwnr9yq=3>(-0FOq5Uaf z;v4HBL#8LFTA_Z$X-!4#H8f6jfLGKQt}-tXpPE2lAhJyG$Sp+C5Iig%XUlBN=F(Dq$f8`!p@V;Z$hAtyjgcb8FwDUN zP!S?L!%G(*i$DT+c~*c%G1hB+6y1r)@8Y?NHcdm$XP9XWtY_$Fx~Z8R40Lhu#U%4N zifBC(sLc1(6B4+dPF2G!8V0&VwY9y*0h*{ab3-G!e|wZJo+_YfhB0*^5`I(M(bbB( zN`f9h77(lvghmLJj{zQavsEEfk_hyWjVPB)(bmx73cyO|D*1?=2q%$dO4Qia5YcJE zexukEoc^$kOdpb~8AeQ^=<3s8E2Pq7FQ{7PY}LqwSS~CI?t_o~0CD349p1q^@Eb|& zglbIn?vTjreYVFwUt0R@QSqOZ>2|Bz*<=b#Newc1*wFVh*OD9%uC5ta8$MzF@HvIP z3vvrrO>JG9N}_v92^Vxu@D^U1Q4PkC0R09E5zm=d(sJF@A}O6h_Ff>8c_~mTI201bbqV5qD#zt$9FqkNJ0jQPMw?mirK4*R zW)WASF)cMUDTo`w5MvD8(93{ZG36-C_vP?mo;7>A5d+HMqsX>m$8z-G>F{E;H@Sl0 zR<0sKGsEY57zlh|F`zFE|0moBmfCNm;Q_hqs9XBPeZP6doNrEBUY?#DQ2tMtNtN>W zIU|mb0jG*f1Y2k{4TFgUW)Z|&5d*MG2~9*Qg5rbeV=IDYTq=0UgK3nB1XAfNAub^Y zF=q1@91!}-MjQ;A#eD$zBLe{5kBns;Q84idfL?whnYt2GYw1o5(b6;|)_-oP_kc^k znSMzy@Dk!`1~;WKxjH(Q{v*2aGe@;X(XqGq3J-d|G)`H{8RxM%B>pXOX?~!rG)lI zwZNhMO*u3ey)JFjn%J5e;M)9`nZzY!DePe)gADwGVA9@ZvbeK{-~9dj`68!FSH{_V z4gL^u$#Hyez0au0-HfKQf_Gsc+}X}oC4=OaXi!acJj>ZmndsKMu9JS`oWpthGnu;D z(wi5yapnnl-nyB7<(d0%{_Y$lnHnN=VGPQ(A_Q&>Diy;Ms80+Vc9n}*a`9L$Nx!g3 zDFN)D(p)`kEVk|25>n2_z+K|*l8}GElZ5&@2&!*@@5HqrAjm~3e}Fbi_X)PvB1q@} zgo+MBG!5Iu&iXvKCcx^Mz`XLqflJdCXj3)?uUftTAO4%3z@cka1k|nN*NKs`Wvu-} zA$lZa3ngmwpg?4+{eWHo42QZp5*OE;u09^IvSw;VMS8;aK#&1*4rcO}djd>Ak(#?=-k4S01v5{Jd{nF3HKLlr+nHVrMe zL%~Bc!y@yjd=ZKvQ1ncFQ26;^!1gjFVlW_em4G;Aw!rSZkqpd;7I&5(o%HbR5g2!v8Z&Rgh#S*;SDkU%H}t5~e){9m0P6j1 zfmEZX4J5uwwKP}sbI(UNuxuw9)~ z?q~!c>ji@hhYGhz!P6^%!wK_mRv8<*k$lNF@nd~C1J{5#}*f95u zX8w3!rjoa3KPz=KWrJFMaLaADZP+OS4*-1AkjlQ6b+eH{kHX(r0oREzB=$CL*A@$2N`JeV11O&x+Op>Da@L!uBekX{apUC?0JD>h6f5=tDO|Y65^wY7jPrA`U^f} zvgN3|e4H#n;3jbz)G`!Te5E1A3c3VkFm*{JEIki(@L)*11Qc?IEa-h0C5}(jhK0aE zil~an$fe#ntrYs~NQu=yt-C(92r=-pyAxHOsQlmkl8My31gQx_L+aDEys)5=h04Q=vnd?EN%g1r&&f8mv&N#obmuvqJ?cGpda(EC0aF3UNr4inczOHOs zxPkZNT-Nl>Ngp@#>EFaUx8@|<;0y^|E)+UA1L*VvSfIEGX_0V{YE)V8}ff~{l#xfDEmB5X;ro>*X^YqJ$REVl~o9Lm(o(J(FIS^=-k)yJ3Vwx30vE6X8{z?LQWzu9BKe$uKP_Y6}QHB7^|9Jn@S7I%Wmh&LO_TD*~`|0qgkRc zf!BtkICyoS8A>!Y9qm`y@F7vPv3w3Dl3o>`(VEE##Z{aecRE4*`FE#NJ|JU?Us5YEgp z;Q9K6bP=K%28C-gJg;F5e3Ppy+$Nr(X$b0y=!yW622eb!UdZWp6M$x6O2yAoxjGj* zru&(%47r}np}-WvKOZmdANsI0Au%e&wB2j$l=9mnha9hZ@a)wGua?3oL!_vbzFhIS zMBYZo9&f4FY2TflPMV_Nh$FoV3sO3UIeMt(;9!dja0ViP8uQ6kgwQ2g9~!9%ug>FA zIMPP%1wsmm$x_y7@hV0W#4?9%2^^1knyPSF#_W)Ya^^Qv*2U5=a=R7y0(F%rjSkXW z>r>Y423pD)V}Kfwk_SB!~s^lnAOOpI;wIS&9J} zg1vQ|SrL&dThyFCg*wZOZhLG+1b5Z-luhPLA=TAUOQIrR*2aQY#hGn)P}DP|Ub;Ib zPVH9?M23t(QygdLVe}uBs9A^upQS9YIqelig`8vnwTAv4LTUmR!mbHtTLO;&12;z* zLMp4y`RC`vc||M~$jzv9f3OmWfCMo0%LgyGUIT7PiDoHN! zRl^CaaLaP)WN}4~95NMrUYrARq`0=pZxkExAmO<-O)NH@dCmrNpo5o)tD&yiV6vSu z57|yW<^XKS0ricGCRImS|9OW^M!a!|O6V5YezLVrgoPTzh;W1Qp9(?^FoaZU&P2f5 zL=serYnHfLZa*8K1bk1acqf(M(KJFVv}r)O0dwye1-99q7&{Cj8>Zy7XMhF#d2%rr z0bvr9-ZIpFmZM(BdKkD+l$r~VxR5XzAAyk>)f|Kh(v@4{km1j!~sI+*|0`;C%4BB#2SqP~A9d(7aL=BPij7dG3PAM3DPLpkq)assQJ)g-$D`CvSiJ z%aJK+tAqkC3aRv~BAF?vt%O-8_=i%C42EJ6DgR%~7ko+hs}ig^8A1m0iI0GDzz%p= z@L=DFw5$s8Bd_6zwde6A@(Di+RpTru;%2X~qF9iVr2Ou5-Q(9n{ z!xqInDTHS(82;ucOgF;LXyAvaxl{K~Dh&${PWA3ietuHOgr^0*AEz!5wgsyc!P&CO z@^Nr7#21LAXe>F9x6EoB^NV70#(4yFuvSA4=@7Q418S#tTsD; zmrZ{V?0($?s3^O~@=A?v7M9dY=aOtXzzz=YbLc7~@}U0o92WFj}AsB#>nC z;;I<%6?TD*z}p+(TB<21yD^;p^^M`tadKsYI7%S>1?mvE^<2+VXc#O6cF_%LtDa0e z^uRKO0h-0bryxajffz1cMZ}u;+4qQHdX~bKi8m8B2%8};%?}m4#~D)8Gu6pYvjZ-*fJ!p+^hJWNwvI8Ad?H68HY40AawiZ~!v648J8%uJUZ<*;2?a zL;;@EP8#9lo?O~^aMprUYs;xok23~r`Xo8O@eXdG(P6oK5Z}X5!xSpyJm3upK3a3s z&<-5HPzpE2=Pajl#zfJ#?O8t$?A69onCL_hZ7{S-y#8iGS*ofS==>CbR&e8Wp$Nnd z%j5(Vm@w+Z?idQWS^+j34Px5}`_P-gUlji%S%{>P5nEaUC;y-Phj|c~Vg@xww&zY2dWpED;|l63#Q`$$D!uTfhX22A0HN zO!xuf5V<=8>InCH@h}R60KJxOFiG-tusp&AxN~C3K=4J-rUhBe1zBBkR+xCIrt4|i zCH<*mw_{#?s?!Mf$MdO<3@S-41?z?ckJRSRl7O4`O>0Vq4(fdOz#%s*oc^$`G?W}U zxx3MITFv+(vUVv`*9smE!7pj>>o@b@Xp15Cr%?GZNnkY4kT?X&E$}sj6q$Jm7B(1x zuu3qTAQ-yYM8ix9OuRuzBre$XMh`tq43300H854V8OHx`;(69W+0e^Vv(46#c4cp?h?;CINxC&k9hnuHvuiP14 z;Drz%03?3!6^&WMulG~J>jN%>=`)BB0Gr?Q6=20~tUPLZy0-ViqC3>~NcZ+P?>cis8S&GS*yuR~6Df350W6r1FC{g-d-h_|1{ zjJ@i^$8%8{_cUB{EH?pILN0Ke3hh$E1MM(M!KpciTPKd6b0N;NPWhFq4m`4A%EI-p zLUk`_Kmox@%LgY8&eVudOmNNKpKQTp5h3M3`ikqq{EBO}h#?dBmuRe~9GIG`3^ihd zVeJmUhYTV+B&#jL^puJBbD-%U)sSG_13&BYn7XKISL0Xwv9EIKuzaW^W3;l7DdGzCy{Ax^uNNgC1t=Fe=OqH}XH_~b6$1UwgLQF}D{MD(&s=6( z$;+B%t+=f{ki4cDShcy)qX~q^Wr!nH2_Ks5e)hnL!1_T&BxsbDh`XT4sR-TcTgTjvr{v-8RVE@`Yc&Q|jVmXj?f<-fZ_#$;UE1Qd1KQCRM92)veO)3MwlFA^OZebHQ%nf-Y`%uX zu;UQIT!almFB|w@+0m{Dzz&Tin|OFF97J?&B~+{bZouTGmcYiKV!DHrT;dKcF?b|Q z&%JaK!|s`=ds*(f$?vBI-#1fCkCxAX6e{RIby2S`W=QFE38wIg=m1&^1vupH?f zlmf4Fr;69+hDC4h>Y2Xfp7G=Bl3xowFS}k@ky(2(!PTIB<-qdNB_iNLnMQ%9;sU|u z4ewu2Or$UvHl^J_mb)Ur&*R~0!xYd&s+kBz;?Mlwn+StVKxZ(A2AhJ%#sG?lrU`g2 z7%bcxSA_s{n%rXHnP1@G`Ll@&ACp_i*x49b(NI=Wo90vV+w;DAGa7#WGNy|+Ep)%@ zD}W-IFih`-14D)Vz^y>tQ23q{0Yu#bW1C6PE|8o?KaGhEGh|$%k0=!f zc(5)8g&*GQ@zeCluZNzT_M__Nl6`lFB`1^4f85BT+UMlN06(!3!FUR*bK>w;nb(0+ zM_gCsht@ZVE39XTy{bm8e-$0Kac@mh{Kjam&KU_|6|=~v5rlXb41{Z)$k4SC$=`rqtU*KZzua@RZMKQ!l`*KSE@rkq|jTSm2HGkwOOCJ&;2ZvZ@l zH~cWgz?1??FicR(tKpB`~9z3fHe!^ymFr5qVh$lf3RAJwOhR?4_JHGPt^p(}^lxepslC=|V zT`Q_tciB^0*^)9A=!4ILgLud-_mw?`ZmZ8{d950LzHI5c-{n5f?WR4?fh{;bcXDWw zGLStXOc;Dk1)gTkfTy9ka03othQe3M;LN+Bi!ZT})M*Fr0-p4Ok^$3ZT2ctooh%0abqQL-eu>|k`kppjt!hrU&(3{R* zAj@sPTTy;x`-#n4NU96#My}pD{Vpm&V7y@%k4Il#-|>F_Vl;K>{5++#O@K#`&Tyi0!PEep zv=4*fA_7nwzej&5v1Gz%C&Hib3_z<8aRYk_igAa2$V_+0=6V1nV9Z_fC}q`0Po%9m z8;B>Jcz!5m%#Z?NA_5d+DnIOI>9W-DjH*SdeMvXYV)9+5z)wXfd@47J%9XT3zoj_S z&WKB=mV&5(w&xINo_$qxyTe0|z&F8u=pPUM6Ik%@RJVs0bEr8o?e(1KiSu9(3esP? z`LB+v4yRv5U-ouTci_2hUt80aUJ4Bl4o&hEbM4ibAAO`wxppIEE z=s+B>6NhN}D8|wGW(12N_>4ImP!k|KP+>LzIsxovwtzMv#7Br>djbx@X0?Qwa4fwH z+Or^;P@)Ehh{9~C%9-ukv<+b}^yvMY2AV&~HOOG8ajgP>HWRo>1oYL3R5O&p3h=m| zWN05FgeO8Hi!{S{;4BbM+&H=cB}2_Ivjt=zDZUbw6B6rx99nJVV95`I>H>rr(n@X* zZ+n2TfyU&43WRF9da2s`D!_jfTuluPCJ>=v36o8J#n2WfDa>J@?2}JvcOZ~gJ2-!0 zfLw0@w+3f~U|=gTKNZsn9OeN~VPM=Su0c)X0(FJrKp6gjQ;jk>C&DO)IFxt0!^oNw z!JwGBIS8>@0Zj%|j&7hCJSCr@BVCdi3$HOV0iN5h%yb8`p8@D#2jfNXtAO+Tdnf?# zu;8UIMhj~P{6TU9VhHs}hyg_M8|?Sr>p?-W0PsM|-S>i z1C-nX;Ab_kRG5-;Rpow;fLSx(spU{IU=J1A8sb+IZ>R(L0E^)2X2Q$n);7UWz@@-9 zVOnr7f>DZ8Tf+jiI5^YBg@YJEgn%uK$itl4;6lEyqnKC^kQJaJ0AfU)fhYva#9!ab zghP)5VW@NI?u*5-GFUF6FSqK1VcCQY$+@Y zWFt&D5zQnP8RpQrF!cq-!XO$A(mQx5@n&dz=^)+V^@$WOg@vMnUxwj3ke}pP;QWB1 zuajph0LAQ#Bl+;APM+R`SEyZzsUZ+xg{{P+ab$r6q9v*n-Y16x1s&ofJud*B z#u!}4gko^3_r2#c2j|yC(1n6%HI0H*1rJunUN~C9U;t2Yapr?8@W;Tug%c#&r4SwB z2}G%!gZB^9G?c+~;^+HbaPna1_Q7NnL?_;p3{%7K26iYBoFOn>xe8MPG7}5x#;d9! zjlgfxKcGYOzzhZW7DCECOhS;mUg%7Ij}*9rQV7*=@Jt+$KmZW`tE%&XiSkbK_&YN& zf}jjgI@r1}3>J;L&M@qj+Oi$6pm-=fb6(i4EDVT}?CzB_>~5LF>pQ4a*r5T|y3TAZ z+}aJao+npFIaC)L_m4p-=F9C zKHulBg;Jg_{vE*z80p$VC|Ivoz`S7d+)W$|N9IqILr|wheWgJXE&;Q&8sez)CedE8 z1Z*$Y;wgplpt{4mc1J~XK$a&9I!a|fs`3wD>7)x9oi@<+@xh1*MkZc|5uWBR{@=OG z=2G@`PLWW3GXPZ9oIXH^c`}IDQZ^I9ejXi+Ud z%cwb}{zj6qcVT@<24N6<@#_mB#k`!sLOG*0AJkl~#iOP#G-RRl7SDQB?<-YURlPNSkgM&!!Vl625ma5U_=C4*I@Rc4dedJIAc-` zgz`3&LMC|`kXC|uAvKL*ftnVHvxbjjC^7ZYVTB-Lemyx(VioIjPP|HJF&9#W(5A!^ z0(w-&;4T>kz!2(e>h*}saixO59LLlxz|NA(13grUi=RvE-)B>Z330KAf+n`+QI^9^ zeKv%dW)LM08CtqaK^T`7P~QI<+3JhPtCNNmG+{oS7&(rDzb40m-cEYO64deLWgyrA z*#YC4lwmcGq6xjntepK)QZ@2(#Vhz~80d$Kp? zkZMAJ+s_RgAq}hV7l;&;D_Pn|oFLd(_b!7JDSQWD2^SdAjq^&D74}xfIJ%6iDq6Jg zpLKR5=rs<9cFR*+E)vBcFmD|Hwm;Y3Cqm0(#7bhN-$l156_vtuv4)4z&yJ|yl*0Ri z0j3Z3dvqt<9I2%pv8q4pe_5Bt4nVaV`LRxig2>V>CJcb0IBtOwTIL?ImeV7>S7AtA7ZvEkXDjWhs1}G;~;Oq5@D-j!l77OnJnE8i-F@}u}Oqr zv?@sKP!P2pFW{~aUMIM57rWT$jN%;R%TmUaN@ZPyHCWEWoo zE95SSTv2$xkZo)}!`E*vg||fo>?}5JVEzfaCqBQ%2nvM|wSp8BsH0-^osnI0m;tsS zqWI;YCCN8urkB)9dsTfv8I;^56mWI=Xl{dJVms%%wXV|(gwu}E!W|kr(TFjEQq&s*(BuJh|-f;l$Lo}(8 z>QV{Q%fGcbfVRObu#W4hH92hY0?GH%HEa^S+H3|7xqd@A@M_Na8+f}oe{2*L2}-bp zoYLe$)c4d-@_q}s0UA4vjX?*@7-<_u+Rz)DLjf7Q2DywM4qyiMKQN2clC%Qb!WE^E z(O|+vu*n9>ZMxn;=xzFo_=Eu>WTI>UO+(hOe#}_4GxHrz7yG(S-|MN+6<-ETF>oXouZ@>n4@?|Mj?d>7}1?+!fD3E~N5Urokme+AkhIO-KtBq6( z2^bP=>sy;8lQW*QOL5tn*9$Bn?(45y3g=9^KvIr?bb^F_*BHvYaDGR}kIlyA6Y4PG zrt&$}G}7M%^_5v*3&Z*g@>Y(vm3&Zn3+5SFAtA`R$EZ0t#8QDNTN~s?s@*v~#bm3q zA8elcLsm}x+>XVXFdGnL$U6a$;1G**V3^ymx?y#uGCF5N(fbQgBF6nV6ztvlF6V4W zs$NQpSma?W`4IZJ?jVTe;posU$`~kaKsH1Uv7!e9;3JJnCNf5~T=X86M|wF8r8I&k zFW^UET(fNsicdi%b*Bl7fDuMPlVc?A;cjWjG|YlOzkMvsq3B&UhewESIY$ZX_s#e60zyvVf%(Ato%> z*CVc}$Q~$=yaRgqkKk5kR_G4aMCGdWmlIyzgC7g^M~N&3(E?S9*Oh_okW)Z1y>IX4 zoIo@wmYq2%e(g|_i2nKH3@r@YSKoH3Z4I$dU8`3lm~tx2clSM*kr?c0QBaZZM~^|U zlTx7juclwZC+{ddOSlR6($rvT<~dsXI*Cx}FXGs(>!?x)zcY`hC5B!WiA%6!L9iwg zG(e4Dp5`I|!rr}?ne3LqO`=^2y4qAFXAOH?*S*>^_KRWy%9?3$n@YA7yPQ<-P`j)x zg!NJe0epj4Q|JT8MKiRv{$4!&fel40Z$N~qIGtr5mfI(`(x$I2M&3&nR0td~J0SfP zg_h~%;`sSg!FK67BYPx~|C9#^6_1Y?aUmXfuh7n=IdouM8 zrN(-vQITM>pqB)!Wmtd8Oj_7Jc)ff+xDVi7i^x!LA0v-&Y=kK62Cb94Xe=^ByJX}o z6VriyF<{CM>4!kS&mkLJMr}{CTl%ni*s!HGpH&@T&@2d6Nb7N1D*c7 z)UCg1C4jA;B87{nXa36Y2xE~m3gG_oqPn9&K`8LyqA-GL>pI?EaqektD9E)dfoc1(DJRi0?q06$z1Si*Jvt!0xoIN7Ph$CLV9~* zZ%8jqP!zv;sk89`r2wr+5hG-tUWx=#cWU&%z-5eJQxja$2DdMozQc!r$zo-C`v2qr zjx)~8vur(C08(n@@$_;{+xIC<+@NiZX~&~;e~78iF3>KQmJpryd_euu^W_#lF*6$&V50a8BkuMyhOD-5OIjX_0fOwrXGI`^J74 zQ?QU;IT^)^VFm=`K4lN97Np^R102-#)O_TE$x!3Vsg>@`e}^}CLiPm~ppN*TmP5q~ zCuT#--AXh))txEfI;5SeBJ-zMMCbKsAuJ2YpZZsOYP@5v9+g^J_^O-!$H z>n~9UNeYadK}c7YFb`Fl%r#k^sbcn|?&L&#U-Y_tj6^U?vMJHjHF;xhn1LA55|eWK zcz3)DVWj$$I?Ur8+vpjO&QU(Xoqr=Mkg^x14^->VCnDp8pa7114lR&Y8T*c^PMEYa zr9>1_V@@57Ug=eEt}JpIIDm_3f*(6I7vWhs48psseU-Ai3AbyMCTsS7a+*m6h-luLoG3AME2SCRKfl@@>o3yH zTKOkf-&o|W)U_JW_B1PG6g+n-4gcOV$=YIH<~P{oD$Z=Re35xw_|Y_ourQ}PUezI? zNV@E%lGmzrPk%Ud&8_d`tn5Qn_iq-{K6{9JIg%Q40GpA?S;{!My6#aM`)6W)eS*Or z^j*oE5<8rx%Cgkd%M6K6!^=3O;v(mDW&R_y3+G{o_P6v0;pOg*r$p?rEDUmXyNY&V zl}}g~^&eO<*&hc6CWe>03CJ4b^I5HZ$V9s8`PDw~#7i%d8YbxLdQrRL(SKkp)bG&B zKce1I2K_XC*4b1yj)rp;*2VrBZUKa%#+MT~3Av*yyl2|RU(55g6CikxB86}}>Dl4X zElihUT)(>M)KKbFvP9>}s3LeZE4X@bR4x3L89}xftX|P97k2q-O-cCcDvqT42i!X z7W`<=!H@j{^(Gw$VDJ#^hKZlnf5Qi4H8plTn*KDj+=YG^hbZDtI#Iv97@;6JeZVa_ z_kM4GI9p(t(rxJ+R5;O36nqkCzf$RJrdK}UeGmDMyu}%AjYXuNBcv8y)Gx}25VeNp zXXxYXzH|_ywP-;=Xb+nYWT)0^c|zn2(fu`aXLccNokmj0P+ z^ruThY|dXO9k2~#3y6vY8MP>&_OQk{;8N;UK!CoQH8}0ORW}Zx1F^o$bBS47`cg4r z(sDY9hWJ;1{E-6svn)2+s(%EL+YS^LfS7g zzt73s3ZuU}R#&<*C-#d4hN+*R6-XSpcCsnm)Eg7Cd>@Ul_8FNX|!~ z%U^nivpmFqaDj%q#3+C)f6S*;E3dobTbTOe)ACbpmlKW0UdWA>8KQ7WOq1#bZjg)f zCXOt6FCGA?ntUDGBoWHD_BEwl4AXhzY6Hy>yTZaP@uCT4SWu}5tk@1Qz?n0-aa_5F z?}0Ul#kZC^e@8eAbBHX&Z)4(Sk`o?MF-5jO0UKC=CedJE%QDotl`N12LJJudCHoj%R+tf8L>2W$AF)@c{-)+`y&k zJc(7W4t>8@cqGhQYE}g*!~$Rssg?SbM%D z!QhjWaKdq{7}F0s=jWQS2&5mtr}%H{Y=fybs$;Z10#91`?+^&el_QIwNlX*QaxY&& z8tfDPi7j$z{c?SbIndugfaUtA7-5?;Jy9NOVe0-yZDs$1nQkCZok-b37=)%1naa?v zd}NoTkeY1~YSrF_hhyUZjj_h`!>Bt_b{=cL3zdJKZ44Yi*C(G9L?>jI!?m|@4`60q zFC~J+BlzruQkhGKWfwq)YRZkJTc|ZXd*kY`9b3Y3P+<)t1QkrL2p@5IbSu|YheWms zODJQY-&OV#;s~{WFF;PZNv$gQ-%jtvNS6a zy%+$nrI^Cek%jGu{Lcjo^9TDMN% literal 0 HcmV?d00001 diff --git a/native/shared/tests/golden_render.rs b/native/shared/tests/golden_render.rs index 0c70d8c..acdda08 100644 --- a/native/shared/tests/golden_render.rs +++ b/native/shared/tests/golden_render.rs @@ -345,3 +345,29 @@ fn cooked_bc7_texture_matches_raw() { "cooked render diverges from raw render: max channel diff {max_diff}" ); } + +#[test] +fn golden_lit_primitives_taa() { + let Some(mut eng) = try_engine() else { + eprintln!("skip: no GPU adapter"); + return; + }; + // Same scene as lit_primitives_3d but with TAA ON: pins the TAA + // branch of the post-FX cascade (reprojection, neighborhood clamp, + // Catmull-Rom upscale path) that the TAA-off goldens never touch. + // The Halton jitter sequence is indexed by frame number, so a fixed + // frame count renders deterministically. + eng.renderer.set_taa_enabled(true); + let (w, h, rgba) = render(&mut eng, 10, |eng| { + let r = &mut eng.renderer; + r.set_clear_color(13.0, 18.0, 26.0, 255.0); + r.begin_mode_3d(4.0, 3.0, 6.0, 0.0, 0.5, 0.0, 0.0, 1.0, 0.0, 45.0, 0.0); + r.add_directional_light(-0.5, -1.0, -0.3, 1.0, 0.95, 0.9, 1.2); + r.add_point_light(2.0, 2.0, 2.0, 10.0, 0.2, 0.4, 1.0, 2.0); + r.draw_plane(0.0, 0.0, 0.0, 10.0, 10.0, 120.0, 120.0, 125.0, 255.0); + r.draw_cube(-1.2, 0.5, 0.0, 1.0, 1.0, 1.0, 230.0, 41.0, 55.0, 255.0); + r.draw_sphere(1.2, 0.75, 0.5, 0.75, 0.0, 228.0, 48.0, 255.0); + r.draw_cube(0.0, 1.6, -1.0, 0.8, 0.8, 0.8, 253.0, 249.0, 0.0, 255.0); + }); + compare_or_update("lit_primitives_taa", w, h, &rgba); +} From 8c988f6247c38b754ed897a28f9141cf52805f0a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:32:16 +0200 Subject: [PATCH 5/7] =?UTF-8?q?feat(renderer):=20the=20frame=20runs=20on?= =?UTF-8?q?=20the=20render=20graph=20=E2=80=94=20RFC=200001=20Phase=202b?= =?UTF-8?q?=20complete?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Every render pass between geometry upload and the terminal composite now executes as a PassNode in one frame Graph: shadow → hdr_scene → translucent → hiz_build → occlusion_capture → gtao → ssao_blur → ssr_march → ssr_temporal → ssgi → bloom → compose → postfx_tail → auto_exposure. Reads/writes declare the real data dependencies (Shadow/HdrColor/MRT/Depth named outputs; Transient tokens for the intermediates the enum doesn't name); each node additionally carries a with_after pin to its predecessor so the schedule reproduces the hand-tuned order exactly. Loosening those pins so the scheduler can interleave independent passes is the documented next refinement — do it dependency-by-dependency with the goldens watching. The context owns &mut Renderer (the cluster-1 pattern at full scale): closures borrow nothing at build time and call the record_* methods. Feature toggles are checked inside closures/methods, never by omitting nodes (with_after on a missing node is a schedule error by design). Immediate-mode uploads moved ahead of the graph build — queue write-buffer ordering is submission-scoped, so within this one encoder the move is semantically identical. All 6 goldens (incl. the new TAA-on scene) pixel-identical. --- native/shared/src/renderer/mod.rs | 341 +++++++++++++++++------------- tools/file-lines-baseline.json | 2 +- 2 files changed, 194 insertions(+), 149 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index c87472e..d4e306d 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -9228,10 +9228,6 @@ impl Renderer { profiler.end("card_light"); } - // Cascaded shadow maps (with the ticket-004 cache-hit skip) — - // see record_shadow_pass in shadow_pass.rs. - self.record_shadow_pass(&mut encoder, profiler, scene); - // Upload immediate-mode 2D data profiler.begin("upload_geometry"); let has_2d = !self.vertices_2d.is_empty(); @@ -9254,158 +9250,207 @@ impl Renderer { } profiler.end("upload_geometry"); - // HDR scene pass (sky-view LUT refresh, sky + 3D batch + - // scene-graph render into the HDR MRTs, then the opaque - // material pass on the inner graph) — see - // record_hdr_scene_pass in scene_pass.rs. - self.record_hdr_scene_pass(&mut encoder, profiler, scene); let surf_w = self.surface_config.width; let surf_h = self.surface_config.height; - - // Translucent / refractive / additive material pass — - // see record_translucent_pass in scene_pass.rs. - self.record_translucent_pass(&mut encoder, profiler); + let exposure_src_idx = self.exposure_current_idx; + let exposure_dst_idx = 1 - self.exposure_current_idx; // ============================================================ - // SSAO: half-res GTAO sampling a hierarchical linear-depth - // pyramid. Build hiz (linearize + 4 min-downsamples), then - // dispatch the GTAO compute pass. + // Frame render graph (RFC 0001 Phase 2b — complete). + // + // Every render pass between geometry upload and the terminal + // composite runs as a PassNode. Reads/writes document the real + // data dependencies; in addition, each node carries a with_after + // pin to its predecessor so the schedule reproduces the + // hand-tuned order exactly. Loosening those pins (to let the + // scheduler interleave independent passes) is the documented + // next refinement — do it dependency-by-dependency with the + // golden tests watching. + // + // The context owns &mut Renderer, so node closures borrow + // nothing at build time and can call the record_* methods. + // Feature toggles (ssao/ssr/ssgi/bloom) are checked inside the + // closures (or inside the methods), never by omitting nodes — + // with_after on a missing node is a schedule error. // ============================================================ - profiler.begin("post_fx"); - if self.ssao_enabled { - let p = &self.current_proj_matrix; - let p00 = p[0][0]; - let p11 = p[1][1]; - let p20 = p[2][0]; - let p21 = p[2][1]; - let p22 = p[2][2]; - let p32 = p[3][2]; - let half_w = (surf_w / 2).max(1); - let half_h = (surf_h / 2).max(1); - - // Hi-Z build + occlusion capture run on the render graph - // (Phase 2b, cluster 1). Unlike the older material-pass nodes - // that capture individual field refs, these use the - // ctx-owns-renderer pattern: the context carries &mut Renderer - // and closures borrow nothing at build time — the shape the - // rest of end_frame_with_scene migrates onto. - { - use graph::{Graph, PassInput, PassNode, PassOutput}; - // Transient id 0 = the linearized Hi-Z pyramid for this - // frame (graph-internal ordering token; the textures - // themselves are persistent renderer fields). - const HIZ_PYRAMID: u32 = 0; - - struct HizCtx<'a> { - r: &'a mut Renderer, - encoder: &'a mut wgpu::CommandEncoder, - profiler: &'a mut crate::profiler::Profiler, - half: (u32, u32), - p22: f32, - p32: f32, - } - - let mut g: Graph = Graph::new(); - g.push( - PassNode::new( - "hiz_build", - Box::new(|ctx: &mut HizCtx| { - let (hw, hh) = ctx.half; - ctx.r.record_hiz_chain(ctx.encoder, ctx.profiler, hw, hh, ctx.p22, ctx.p32); - }), - ) - .with_reads(&[PassInput::SceneDepth]) - .with_writes(&[PassOutput::Transient(HIZ_PYRAMID)]), - ); - // Max-reduce the linearized depth into the 64x64 occlusion - // grid and queue its readback; scene.prepare consumes it - // next frame (one-frame latency, no stall). - g.push( - PassNode::new( - "occlusion_capture", - Box::new(|ctx: &mut HizCtx| { - let vp = ctx.r.vp_matrix(); - let (hw, hh) = ctx.half; - // Split borrows: occlusion is a sibling field - // of device/queue; record() also needs the - // hiz view. - let occlusion = &mut ctx.r.occlusion as *mut OcclusionCuller; - unsafe { - (*occlusion).record( - &ctx.r.device, - &ctx.r.queue, - ctx.encoder, - &ctx.r.hiz_views[0], - (hw, hh), - vp, - ); - } - }), - ) - .with_reads(&[PassInput::Transient(HIZ_PYRAMID)]), - ); - let mut ctx = HizCtx { - r: self, - encoder: &mut encoder, - profiler, - half: (half_w, half_h), - p22, - p32, - }; - if let Err(e) = g.execute(&mut ctx) { - eprintln!("[graph] hiz/occlusion cluster failed: {:?}", e); - } + { + use graph::{Graph, PassInput, PassNode, PassOutput}; + // Transient ordering tokens for resources the enum doesn't + // name. The textures themselves are persistent renderer + // fields; these ids only express producer→consumer edges. + const HIZ_PYRAMID: u32 = 0; + const SSAO_TEX: u32 = 1; + const SSR_TEX: u32 = 2; + const SSGI_TEX: u32 = 3; + const BLOOM_CHAIN: u32 = 4; + const COMPOSED: u32 = 5; + const LDR_FINAL: u32 = 6; + + struct FrameCtx2<'a> { + r: &'a mut Renderer, + encoder: &'a mut wgpu::CommandEncoder, + profiler: &'a mut crate::profiler::Profiler, + scene: &'a mut crate::scene::SceneGraph, + surf: (u32, u32), + exposure_idx: (usize, usize), } - // GTAO compute (samples the Hi-Z pyramid) — see - // record_gtao in hiz.rs. - self.record_gtao(&mut encoder, profiler, half_w, half_h, p00, p11, p20, p21); - } - - // GTAO bilateral blur (or disabled-clear) — see hiz.rs. - self.record_ssao_blur(&mut encoder, surf_w, surf_h); - - // SSR ray march — see record_ssr_march in ssr_pass.rs. - self.record_ssr_march(&mut encoder, profiler); - - // SSR temporal denoiser — see record_ssr_temporal in ssr_pass.rs. - self.record_ssr_temporal(&mut encoder); - - // The compose pass reads denoised SSR from the current history - // texture when ssr_enabled; otherwise the raw ssr_rt (which was - // cleared to transparent above) so it contributes nothing. - // Lumen-style screen-probe SSGI (place/trace/temporal/resolve) - // or disabled-clear — see record_ssgi_passes in ssgi_pass.rs. - self.record_ssgi_passes(&mut encoder, profiler, surf_w, surf_h); - - - // The resolve pass writes directly into `ssgi_rt_view`, so - // downstream composite + TAA reads are unchanged from the - // legacy path. - // Bloom chain (Karis-thresholded downsample + additive upsample) - // — see record_bloom_chain in postfx_chain.rs. - self.record_bloom_chain(&mut encoder, profiler, surf_w, surf_h); - - - // Scene compose (HDR + SSR + SSGI*albedo + bloom + fog + shafts - // -> composed_rt) — see record_scene_compose in postfx_chain.rs. - self.record_scene_compose(&mut encoder); - // Post-FX tail: upscale/TAA/DoF/motion-blur/SSS/CAS, each - // reading the previous enabled stage — see - // record_postfx_tail in postfx_chain.rs. - self.record_postfx_tail(&mut encoder, profiler); - + let mut g: Graph = Graph::new(); + g.push( + PassNode::new("shadow", Box::new(|c: &mut FrameCtx2| { + c.r.record_shadow_pass(c.encoder, c.profiler, c.scene); + })) + .with_writes(&[PassOutput::Shadow(0), PassOutput::Shadow(1), PassOutput::Shadow(2)]), + ); + g.push( + PassNode::new("hdr_scene", Box::new(|c: &mut FrameCtx2| { + c.r.record_hdr_scene_pass(c.encoder, c.profiler, c.scene); + })) + .with_reads(&[PassInput::Shadow(0), PassInput::Shadow(1), PassInput::Shadow(2)]) + .with_writes(&[ + PassOutput::HdrColor, + PassOutput::MaterialRt, + PassOutput::VelocityRt, + PassOutput::AlbedoRt, + PassOutput::Depth, + ]) + .with_after(&["shadow"]), + ); + g.push( + PassNode::new("translucent", Box::new(|c: &mut FrameCtx2| { + c.r.record_translucent_pass(c.encoder, c.profiler); + })) + // Reads the opaque HDR + depth and alpha-blends back into + // HdrColor; the pin (not a second HdrColor write) keeps a + // single declared writer per resource. + .with_after(&["hdr_scene"]), + ); + g.push( + PassNode::new("hiz_build", Box::new(|c: &mut FrameCtx2| { + if !c.r.ssao_enabled { + return; + } + let (hw, hh) = ((c.surf.0 / 2).max(1), (c.surf.1 / 2).max(1)); + let p22 = c.r.current_proj_matrix[2][2]; + let p32 = c.r.current_proj_matrix[3][2]; + c.r.record_hiz_chain(c.encoder, c.profiler, hw, hh, p22, p32); + })) + .with_reads(&[PassInput::SceneDepth]) + .with_writes(&[PassOutput::Transient(HIZ_PYRAMID)]) + .with_after(&["translucent"]), + ); + g.push( + PassNode::new("occlusion_capture", Box::new(|c: &mut FrameCtx2| { + if !c.r.ssao_enabled { + return; + } + let (hw, hh) = ((c.surf.0 / 2).max(1), (c.surf.1 / 2).max(1)); + let vp = c.r.vp_matrix(); + let occlusion = &mut c.r.occlusion as *mut OcclusionCuller; + unsafe { + (*occlusion).record(&c.r.device, &c.r.queue, c.encoder, &c.r.hiz_views[0], (hw, hh), vp); + } + })) + .with_reads(&[PassInput::Transient(HIZ_PYRAMID)]) + .with_after(&["hiz_build"]), + ); + g.push( + PassNode::new("gtao", Box::new(|c: &mut FrameCtx2| { + if !c.r.ssao_enabled { + return; + } + let (hw, hh) = ((c.surf.0 / 2).max(1), (c.surf.1 / 2).max(1)); + let p = &c.r.current_proj_matrix; + let (p00, p11, p20, p21) = (p[0][0], p[1][1], p[2][0], p[2][1]); + c.r.record_gtao(c.encoder, c.profiler, hw, hh, p00, p11, p20, p21); + })) + .with_reads(&[PassInput::Transient(HIZ_PYRAMID)]) + .with_after(&["occlusion_capture"]), + ); + g.push( + PassNode::new("ssao_blur", Box::new(|c: &mut FrameCtx2| { + c.r.record_ssao_blur(c.encoder, c.surf.0, c.surf.1); + })) + .with_writes(&[PassOutput::Transient(SSAO_TEX)]) + .with_after(&["gtao"]), + ); + g.push( + PassNode::new("ssr_march", Box::new(|c: &mut FrameCtx2| { + c.r.record_ssr_march(c.encoder, c.profiler); + })) + .with_reads(&[PassInput::SceneColor, PassInput::SceneDepth]) + .with_after(&["ssao_blur"]), + ); + g.push( + PassNode::new("ssr_temporal", Box::new(|c: &mut FrameCtx2| { + c.r.record_ssr_temporal(c.encoder); + })) + .with_writes(&[PassOutput::Transient(SSR_TEX)]) + .with_after(&["ssr_march"]), + ); + g.push( + PassNode::new("ssgi", Box::new(|c: &mut FrameCtx2| { + c.r.record_ssgi_passes(c.encoder, c.profiler, c.surf.0, c.surf.1); + })) + .with_reads(&[PassInput::SceneColor, PassInput::SceneDepth]) + .with_writes(&[PassOutput::Transient(SSGI_TEX)]) + .with_after(&["ssr_temporal"]), + ); + g.push( + PassNode::new("bloom", Box::new(|c: &mut FrameCtx2| { + c.r.record_bloom_chain(c.encoder, c.profiler, c.surf.0, c.surf.1); + })) + .with_reads(&[PassInput::SceneColor]) + .with_writes(&[PassOutput::Transient(BLOOM_CHAIN)]) + .with_after(&["ssgi"]), + ); + g.push( + PassNode::new("compose", Box::new(|c: &mut FrameCtx2| { + c.r.record_scene_compose(c.encoder); + })) + .with_reads(&[ + PassInput::SceneColor, + PassInput::Transient(SSAO_TEX), + PassInput::Transient(SSR_TEX), + PassInput::Transient(SSGI_TEX), + PassInput::Transient(BLOOM_CHAIN), + ]) + .with_writes(&[PassOutput::Transient(COMPOSED)]) + .with_after(&["bloom"]), + ); + g.push( + PassNode::new("postfx_tail", Box::new(|c: &mut FrameCtx2| { + c.r.record_postfx_tail(c.encoder, c.profiler); + })) + .with_reads(&[PassInput::Transient(COMPOSED), PassInput::MotionVectors]) + .with_writes(&[PassOutput::Transient(LDR_FINAL)]) + .with_after(&["compose"]), + ); + g.push( + PassNode::new("auto_exposure", Box::new(|c: &mut FrameCtx2| { + let (src, dst) = c.exposure_idx; + c.r.record_auto_exposure(c.encoder, src, dst); + })) + .with_reads(&[PassInput::Transient(LDR_FINAL)]) + .with_after(&["postfx_tail"]), + ); - // ============================================================ - // Auto-exposure update pass (runs only when auto_exposure is - // on; otherwise the composite reads the old exposure texture - // which is fine since manual_exposure bypasses the read). - // ============================================================ - let exposure_src_idx = self.exposure_current_idx; - let exposure_dst_idx = 1 - self.exposure_current_idx; - // Measurement + adaptation — see record_auto_exposure in - // postfx_chain.rs. Composite reads exposure_views[dst]. - self.record_auto_exposure(&mut encoder, exposure_src_idx, exposure_dst_idx); + let mut ctx = FrameCtx2 { + r: self, + encoder: &mut encoder, + profiler, + scene, + surf: (surf_w, surf_h), + exposure_idx: (exposure_src_idx, exposure_dst_idx), + }; + if let Err(e) = g.execute(&mut ctx) { + // A schedule error means a malformed graph (cycle / + // unknown pin) — a programming error, not a runtime + // condition. Surface loudly; the frame still presents + // whatever was encoded before the failure. + eprintln!("[graph] frame graph failed: {:?}", e); + } + } let composite_src_view = self.composite_source_view(); diff --git a/tools/file-lines-baseline.json b/tools/file-lines-baseline.json index b914210..52aac05 100644 --- a/tools/file-lines-baseline.json +++ b/tools/file-lines-baseline.json @@ -1,3 +1,3 @@ { - "native/shared/src/renderer/mod.rs": 12099 + "native/shared/src/renderer/mod.rs": 11871 } \ No newline at end of file From 4c5d1553cc9e1462a69b32f3c268d66d108eb150 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:32:31 +0200 Subject: [PATCH 6/7] chore: drop debug probe example --- native/shared/examples/probe_compose.rs | 31 ------------------------- 1 file changed, 31 deletions(-) delete mode 100644 native/shared/examples/probe_compose.rs diff --git a/native/shared/examples/probe_compose.rs b/native/shared/examples/probe_compose.rs deleted file mode 100644 index eba063e..0000000 --- a/native/shared/examples/probe_compose.rs +++ /dev/null @@ -1,31 +0,0 @@ -use bloom_shared::engine::EngineState; -use bloom_shared::renderer::Renderer; -fn main() { - let instance = wgpu::Instance::new(wgpu::InstanceDescriptor { - backends: wgpu::Backends::all(), - ..wgpu::InstanceDescriptor::new_without_display_handle() - }); - let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions::default())).unwrap(); - let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { - required_limits: adapter.limits(), ..Default::default() - })).unwrap(); - let renderer = Renderer::new_headless(device, queue, 256, 256); - let mut eng = EngineState::new(renderer); - eng.renderer.set_taa_enabled(true); - - let n: u32 = std::env::args().nth(1).map(|v| v.parse().unwrap()).unwrap_or(1); for i in 0..n { - eng.begin_frame(); - eng.renderer.set_clear_color(13.0, 18.0, 26.0, 255.0); - eng.renderer.begin_mode_3d(4.0, 3.0, 6.0, 0.0, 0.5, 0.0, 0.0, 1.0, 0.0, 45.0, 0.0); - eng.renderer.add_directional_light(-0.5, -1.0, -0.3, 1.0, 0.95, 0.9, 1.2); - eng.renderer.add_point_light(2.0, 2.0, 2.0, 10.0, 0.2, 0.4, 1.0, 2.0); - eng.renderer.draw_plane(0.0, 0.0, 0.0, 10.0, 10.0, 120.0, 120.0, 125.0, 255.0); - eng.renderer.draw_cube(-1.2, 0.5, 0.0, 1.0, 1.0, 1.0, 230.0, 41.0, 55.0, 255.0); - if i + 1 == n { eng.renderer.screenshot_requested = true; } - eng.end_frame(); - } - let (w, h, mut data) = eng.renderer.screenshot_data.take().unwrap(); - for px in data.chunks_exact_mut(4) { px.swap(0, 2); } - image::save_buffer("/tmp/probe_compose.png", &data, w, h, image::ColorType::Rgba8).unwrap(); - println!("sky pixel: {:?}", &data[..4]); -} From 0d98e240ca8ebd49d983420698f5b64cbd0118a8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Fri, 12 Jun 2026 13:54:13 +0200 Subject: [PATCH 7/7] feat(renderer): froxel light clustering for the scene shader (#23) Point lights were O(live lights) per fragment since the 256-light cap raise. A 16x9x24 froxel grid now restores O(cluster lights) on every backend with fragment-stage storage buffers: - froxel.rs: compute pass assigns lights to clusters (sphere vs froxel AABB, log-distributed depth slices); cluster counts + index list in storage buffers; light data stays in the existing UBO. - The clustered point-light loop is spliced into SCENE_SHADER between the BEGIN/END-POINT-LIGHT-LOOP markers at pipeline build; the plain loop remains the WebGL2 fallback and semantic reference. - lighting_layout grows bindings 10-12 when clustered; pipeline_3d shares the layout unaffected (immediate-mode keeps the plain loop). - froxel_assign runs as a render-graph node before hdr_scene. - BLOOM_DISABLE_FROXEL=1 kill-switch forces the reference loop. Parity: new golden many_point_lights_clustered_scene drives the retained scene path (scene_pipeline) under 40 lights; its golden was generated with the reference loop and the clustered path reproduces it bit-exactly (mean/max diff 0.0). The assignment is strictly conservative, so the two paths are mathematically identical. Also extracts lighting.rs: the group-1 bind-group entry list existed in four hand-synced copies (a binding-drift hazard); all lighting bind groups now go through one builder. mod.rs ratchets 11871->11775. --- native/shared/src/renderer/froxel.rs | 421 ++++++++++++++++++ native/shared/src/renderer/lighting.rs | 153 +++++++ native/shared/src/renderer/mod.rs | 224 +++------- native/shared/src/renderer/shaders/core.rs | 5 + .../many_point_lights_clustered_scene.png | Bin 0 -> 4754 bytes native/shared/tests/golden_render.rs | 135 ++++-- tools/file-lines-baseline.json | 4 +- 7 files changed, 748 insertions(+), 194 deletions(-) create mode 100644 native/shared/src/renderer/froxel.rs create mode 100644 native/shared/src/renderer/lighting.rs create mode 100644 native/shared/tests/golden/many_point_lights_clustered_scene.png diff --git a/native/shared/src/renderer/froxel.rs b/native/shared/src/renderer/froxel.rs new file mode 100644 index 0000000..943b726 --- /dev/null +++ b/native/shared/src/renderer/froxel.rs @@ -0,0 +1,421 @@ +//! Froxel light clustering — task #23 of the architecture audit. +//! +//! The 8+256 light-cap raise removed the capability ceiling but left +//! the scene shader paying O(live point lights) per fragment. This +//! module restores O(cluster lights): a compute pass assigns the point +//! lights (read from the same lighting UBO the shaders already use) to +//! a 16×9×24 view-frustum froxel grid each frame, and a clustered +//! variant of the scene shader loops only its froxel's index list. +//! +//! Backend split, by capability rather than cfg: storage buffers in +//! fragment shaders don't exist on WebGL2, so [`FroxelPass::supported`] +//! gates on the device limits. Unsupported backends keep the plain +//! count-driven loop (the semantic reference — the clustered path must +//! match it exactly, which the many_point_lights golden enforces). +//! +//! Memory: counts 3456×4 B ≈ 14 KB; index list 3456×256×4 B ≈ 3.5 MB +//! (256 = worst-case every light in one froxel — exact parity with the +//! reference loop, no truncation). + +use wgpu::util::DeviceExt; + +pub(super) const GRID_X: u32 = 16; +pub(super) const GRID_Y: u32 = 9; +pub(super) const GRID_Z: u32 = 24; +pub(super) const CLUSTER_COUNT: u32 = GRID_X * GRID_Y * GRID_Z; +pub(super) const MAX_LIGHTS_PER_CLUSTER: u32 = 256; + +/// Uniform parameters shared by the assignment compute pass and the +/// clustered fragment loop. Layout mirrored in WGSL below and in the +/// fragment include. +#[repr(C)] +#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +pub(super) struct FroxelParams { + /// View matrix (world → view) for light-position transform. + pub view: [[f32; 4]; 4], + /// x = grid_x, y = grid_y, z = grid_z, w = live point-light count. + pub grid: [u32; 4], + /// x = znear, y = zfar, z = log(zfar/znear), w = unused. + pub depth_range: [f32; 4], + /// x = 1/tile_w_px, y = 1/tile_h_px (fragment tile lookup), + /// z = p22, w = p32 (depth linearization, same convention as Hi-Z). + pub screen: [f32; 4], + /// Inverse projection — froxel corner reconstruction in the + /// assignment pass. + pub inv_proj: [[f32; 4]; 4], +} + +const ASSIGN_SHADER: &str = " +struct FroxelParams { + view: mat4x4, + grid: vec4, + depth_range: vec4, + screen: vec4, + inv_proj: mat4x4, +}; +struct PointLight { position: vec4, color: vec4 }; +struct Lights { + // Mirrors the tail of the Lighting UBO relevant here. The host + // binds a dedicated compact UBO (positions+ranges only would do, + // but reusing PointLight keeps one struct). + count: vec4, + lights: array, +}; + +@group(0) @binding(0) var p: FroxelParams; +@group(0) @binding(1) var l: Lights; +@group(0) @binding(2) var cluster_counts: array; +@group(0) @binding(3) var cluster_indices: array; + +// View-space Z of slice boundary k (logarithmic distribution). +fn slice_z(k: u32) -> f32 { + let t = f32(k) / f32(p.grid.z); + return p.depth_range.x * exp(t * p.depth_range.z); +} + +@compute @workgroup_size(4, 4, 4) +fn cs_main(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= p.grid.x || gid.y >= p.grid.y || gid.z >= p.grid.z) { return; } + let cluster = gid.x + gid.y * p.grid.x + gid.z * p.grid.x * p.grid.y; + + // Froxel AABB in view space: reconstruct the tile's corner rays on + // the near plane of the projection and scale to the slice depths. + // NDC tile extents: + let x0 = (f32(gid.x) / f32(p.grid.x)) * 2.0 - 1.0; + let x1 = (f32(gid.x + 1u) / f32(p.grid.x)) * 2.0 - 1.0; + // NDC y is up; tile row 0 is the TOP of the screen. + let y1 = 1.0 - (f32(gid.y) / f32(p.grid.y)) * 2.0; + let y0 = 1.0 - (f32(gid.y + 1u) / f32(p.grid.y)) * 2.0; + + // Unproject the four corners at an arbitrary depth and normalize to + // rays through the camera (view space, looking down -Z). + var mn = vec3( 1e30, 1e30, 1e30); + var mx = vec3(-1e30, -1e30, -1e30); + let z_near_s = slice_z(gid.z); + let z_far_s = slice_z(gid.z + 1u); + for (var cx = 0u; cx < 2u; cx++) { + for (var cy = 0u; cy < 2u; cy++) { + let nx = select(x0, x1, cx == 1u); + let ny = select(y0, y1, cy == 1u); + let h = p.inv_proj * vec4(nx, ny, 0.5, 1.0); + let dir = h.xyz / h.w; // a point on the ray (view space) + let ray = dir / max(-dir.z, 1e-6); // scale so z == -1 + // corner at both slice depths (view z is negative forward) + let a = ray * z_near_s; + let b = ray * z_far_s; + mn = min(mn, min(vec3(a.xy, -z_near_s), vec3(b.xy, -z_far_s))); + mx = max(mx, max(vec3(a.xy, -z_near_s), vec3(b.xy, -z_far_s))); + } + } + + // Sphere/AABB tests against every live light. + var count = 0u; + let n = u32(l.count.x); + let base = cluster * 256u; + for (var i = 0u; i < n; i++) { + let pos_w = l.lights[i].position; + let pos_v = (p.view * vec4(pos_w.xyz, 1.0)).xyz; + let r = pos_w.w; + let closest = clamp(pos_v, mn, mx); + let d = pos_v - closest; + if (dot(d, d) <= r * r) { + cluster_indices[base + count] = i; + count++; + } + } + cluster_counts[cluster] = count; +} +"; + +/// The fragment-side replacement for the plain point-light loop, plus +/// the bindings it needs. Spliced into SCENE_SHADER between the +/// BEGIN/END-POINT-LIGHT-LOOP markers by [`clustered_scene_shader`]. +const CLUSTERED_BINDINGS: &str = " +struct FroxelParams { + view: mat4x4, + grid: vec4, + depth_range: vec4, + screen: vec4, + inv_proj: mat4x4, +}; +@group(1) @binding(10) var froxel: FroxelParams; +@group(1) @binding(11) var cluster_counts: array; +@group(1) @binding(12) var cluster_indices: array; +"; + +const CLUSTERED_LOOP: &str = " + // Froxel-clustered point lights: identical shading math to the + // reference loop, restricted to this fragment's cluster list. + let view_z = -froxel.screen.w / (in.clip_position.z + froxel.screen.z); + let slice = clamp( + u32(log(max(view_z, froxel.depth_range.x) / froxel.depth_range.x) + / froxel.depth_range.z * f32(froxel.grid.z)), + 0u, froxel.grid.z - 1u); + let tile_x = min(u32(in.clip_position.x * froxel.screen.x), froxel.grid.x - 1u); + let tile_y = min(u32(in.clip_position.y * froxel.screen.y), froxel.grid.y - 1u); + let cluster = tile_x + tile_y * froxel.grid.x + slice * froxel.grid.x * froxel.grid.y; + let cl_count = cluster_counts[cluster]; + let cl_base = cluster * 256u; + for (var ci = 0u; ci < cl_count; ci++) { + let pl = lighting.point_lights[cluster_indices[cl_base + ci]]; + let to_light = pl.position.xyz - in.world_pos; + let dist = length(to_light); + let range = pl.position.w; + if (dist < range && dist > 0.0) { + let l = to_light / dist; + let atten = 1.0 - (dist / range); + let atten2 = atten * atten; + lit += shade_pbr(n, v, l, pl.color.rgb, pl.color.w * atten2, + base_color, metallic, roughness); + } + } +"; + +/// Build the clustered SCENE_SHADER variant from the canonical source. +pub(super) fn clustered_scene_shader(source: &str) -> String { + let begin = source + .find("// BEGIN-POINT-LIGHT-LOOP") + .expect("scene shader missing BEGIN-POINT-LIGHT-LOOP marker"); + let end_marker = "// END-POINT-LIGHT-LOOP"; + let end = source.find(end_marker).expect("scene shader missing END marker") + end_marker.len(); + format!( + "{}{}{}{}", + CLUSTERED_BINDINGS, + &source[..begin], + CLUSTERED_LOOP, + &source[end..] + ) +} + +/// The three entries appended to `lighting_layout` (group 1) when the +/// device supports the clustered path. Pipelines whose shaders don't +/// reference them (SHADER_3D's pipeline_3d) are unaffected — extra +/// layout entries are legal as long as the bind group provides them. +pub(super) fn extra_lighting_layout_entries() -> [wgpu::BindGroupLayoutEntry; 3] { + let storage_ro = wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, + min_binding_size: None, + }; + [ + wgpu::BindGroupLayoutEntry { + binding: 10, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { binding: 11, visibility: wgpu::ShaderStages::FRAGMENT, ty: storage_ro, count: None }, + wgpu::BindGroupLayoutEntry { binding: 12, visibility: wgpu::ShaderStages::FRAGMENT, ty: storage_ro, count: None }, + ] +} + +pub struct FroxelPass { + pub assign_pipeline: wgpu::ComputePipeline, + pub assign_layout: wgpu::BindGroupLayout, + pub params_buffer: wgpu::Buffer, + /// Compact point-light UBO for the compute pass (count + 256 lights). + pub lights_buffer: wgpu::Buffer, + pub counts_buffer: wgpu::Buffer, + pub indices_buffer: wgpu::Buffer, + assign_bg: wgpu::BindGroup, +} + +impl FroxelPass { + /// Storage buffers must be available in BOTH compute and fragment + /// stages (WebGL2 has neither). `BLOOM_DISABLE_FROXEL=1` forces the + /// reference loop — used to (re)generate the clustered-parity + /// golden and to bisect suspected clustering bugs in the field. + pub fn supported(device: &wgpu::Device) -> bool { + if std::env::var_os("BLOOM_DISABLE_FROXEL").is_some_and(|v| v == "1") { + return false; + } + let l = device.limits(); + l.max_storage_buffers_per_shader_stage >= 2 + && l.max_storage_buffer_binding_size as u64 + >= (CLUSTER_COUNT * MAX_LIGHTS_PER_CLUSTER * 4) as u64 + } + + pub fn new(device: &wgpu::Device) -> Self { + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("froxel_assign_shader"), + source: wgpu::ShaderSource::Wgsl(ASSIGN_SHADER.into()), + }); + let assign_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("froxel_assign_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::Uniform, + 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::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + ], + }); + let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("froxel_assign_pl"), + bind_group_layouts: &[Some(&assign_layout)], + ..Default::default() + }); + let assign_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("froxel_assign_pipeline"), + layout: Some(&pl), + module: &shader, + entry_point: Some("cs_main"), + compilation_options: Default::default(), + cache: None, + }); + let params_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("froxel_params"), + contents: &[0u8; std::mem::size_of::()], + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); + // count vec4 + 256 lights × 2 vec4 + let lights_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("froxel_lights"), + size: 16 + 256 * 32, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let counts_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("froxel_counts"), + size: (CLUSTER_COUNT * 4) as u64, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let indices_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("froxel_indices"), + size: (CLUSTER_COUNT * MAX_LIGHTS_PER_CLUSTER * 4) as u64, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + let assign_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("froxel_assign_bg"), + layout: &assign_layout, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: params_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: lights_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: counts_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: indices_buffer.as_entire_binding() }, + ], + }); + Self { + assign_pipeline, + assign_layout, + params_buffer, + lights_buffer, + counts_buffer, + indices_buffer, + assign_bg, + } + } + + /// The bind-group entries matching [`extra_lighting_layout_entries`], + /// appended to every lighting bind group the renderer builds. + pub(super) fn extra_lighting_bind_entries(&self) -> [wgpu::BindGroupEntry<'_>; 3] { + [ + wgpu::BindGroupEntry { binding: 10, resource: self.params_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 11, resource: self.counts_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 12, resource: self.indices_buffer.as_entire_binding() }, + ] + } + + /// Record the per-frame assignment dispatch. The caller uploads + /// params + lights first (see Renderer::record_froxel_assign). + pub fn record(&self, encoder: &mut wgpu::CommandEncoder) { + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("froxel_assign_pass"), + timestamp_writes: None, + }); + pass.set_pipeline(&self.assign_pipeline); + pass.set_bind_group(0, &self.assign_bg, &[]); + pass.dispatch_workgroups(GRID_X / 4, GRID_Y.div_ceil(4), GRID_Z / 4); + } +} + +impl super::Renderer { + /// Upload froxel params + the compact light list and dispatch the + /// assignment pass. Runs every 3D frame on supported devices — + /// even with zero lights, so `cluster_counts` never carries stale + /// data from a previous frame's camera. + pub(super) fn record_froxel_assign(&mut self, encoder: &mut wgpu::CommandEncoder) { + let Some(froxel) = &self.froxel else { return }; + + let proj = self.current_proj_matrix; + let p22 = proj[2][2]; + let p32 = proj[3][2]; + // Same linearization as Hi-Z: view_z(depth) = -p32/(depth + p22), + // positive forward. Evaluate at depth 0 and 1; min/max makes + // this hold for reversed-Z too, and the clamps keep an + // infinite-far projection (division by ~0) finite. + let z_at = |d: f32| -p32 / (d + p22); + let (z0, z1) = (z_at(0.0), z_at(1.0)); + let znear = z0.min(z1).max(1e-3); + let zfar = z0.max(z1).clamp(znear * 1.001, 1e9); + + // clip_position.xy is in render-target pixels — the HDR scene + // pass runs at render_extent (render_scale-aware), not surface + // size. + let (rw, rh) = self.render_extent(); + let n = (self.lighting_uniforms.point_light_count[0] as u32) + .min(MAX_LIGHTS_PER_CLUSTER); + let params = FroxelParams { + view: self.current_view_matrix, + grid: [GRID_X, GRID_Y, GRID_Z, n], + depth_range: [znear, zfar, (zfar / znear).ln(), 0.0], + screen: [ + GRID_X as f32 / rw.max(1) as f32, + GRID_Y as f32 / rh.max(1) as f32, + p22, + p32, + ], + inv_proj: self.current_inv_proj_matrix, + }; + self.queue.write_buffer(&froxel.params_buffer, 0, bytemuck::bytes_of(¶ms)); + let count = [n as f32, 0.0, 0.0, 0.0_f32]; + self.queue.write_buffer(&froxel.lights_buffer, 0, bytemuck::bytes_of(&count)); + self.queue.write_buffer( + &froxel.lights_buffer, + 16, + bytemuck::cast_slice(&self.lighting_uniforms.point_lights), + ); + froxel.record(encoder); + } +} diff --git a/native/shared/src/renderer/lighting.rs b/native/shared/src/renderer/lighting.rs new file mode 100644 index 0000000..453887b --- /dev/null +++ b/native/shared/src/renderer/lighting.rs @@ -0,0 +1,153 @@ +//! Group-1 lighting bind group — layout and construction. +//! +//! The scene + immediate-mode 3D pipelines share one bind-group layout +//! for lighting data: the Lighting UBO, env/IBL textures, the shadow +//! cascade, and (on clustered devices) the froxel buffers at bindings +//! 10-12. The bind group is rebuilt whenever the env source changes +//! (HDR load, panorama, procedural sky); every rebuild goes through +//! [`Renderer::make_lighting_bind_group`] so the entry list exists in +//! exactly one place and cannot drift between call sites. + +use super::{froxel, Renderer}; + +/// Create the group-1 layout. `clustered` appends the froxel bindings — +/// set when [`froxel::FroxelPass::supported`] holds for the device. +/// Pipelines whose shaders don't reference bindings 10-12 (pipeline_3d) +/// share the layout unaffected; extra entries are legal as long as the +/// bind group provides them. +pub(super) fn create_lighting_layout( + device: &wgpu::Device, + clustered: bool, +) -> wgpu::BindGroupLayout { + let tex_float = wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }; + let tex_depth = wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }; + let frag = wgpu::ShaderStages::FRAGMENT; + let mut entries = vec![ + // 0: Lighting UBO + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: frag, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + // 1/2: env (IBL specular) texture + sampler + wgpu::BindGroupLayoutEntry { binding: 1, visibility: frag, ty: tex_float, count: None }, + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: frag, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + // 3/4: BRDF LUT + sampler + wgpu::BindGroupLayoutEntry { binding: 3, visibility: frag, ty: tex_float, count: None }, + wgpu::BindGroupLayoutEntry { + binding: 4, + visibility: frag, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + // 5-7: shadow cascades, 8: comparison sampler + wgpu::BindGroupLayoutEntry { binding: 5, visibility: frag, ty: tex_depth, count: None }, + wgpu::BindGroupLayoutEntry { binding: 6, visibility: frag, ty: tex_depth, count: None }, + wgpu::BindGroupLayoutEntry { binding: 7, visibility: frag, ty: tex_depth, count: None }, + wgpu::BindGroupLayoutEntry { + binding: 8, + visibility: frag, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Comparison), + count: None, + }, + // 9: env diffuse (IBL irradiance) + wgpu::BindGroupLayoutEntry { binding: 9, visibility: frag, ty: tex_float, count: None }, + ]; + if clustered { + entries.extend(froxel::extra_lighting_layout_entries()); + } + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("lighting_layout"), + entries: &entries, + }) +} + +/// Everything a lighting bind group references besides the env views. +/// `Renderer::new` builds one from constructor locals (before `self` +/// exists); [`Renderer::make_lighting_bind_group`] from fields. +pub(super) struct LightingBindSources<'a> { + pub lighting_buffer: &'a wgpu::Buffer, + pub env_sampler: &'a wgpu::Sampler, + pub brdf_lut_view: &'a wgpu::TextureView, + pub brdf_lut_sampler: &'a wgpu::Sampler, + pub shadow_map: &'a crate::shadows::ShadowMap, + pub froxel: Option<&'a froxel::FroxelPass>, +} + +/// The single source of truth for the group-1 entry list — every +/// lighting bind group the renderer ever creates goes through here. +pub(super) fn create_lighting_bind_group( + device: &wgpu::Device, + layout: &wgpu::BindGroupLayout, + label: &str, + src: &LightingBindSources<'_>, + env_view: &wgpu::TextureView, + diffuse_view: &wgpu::TextureView, +) -> wgpu::BindGroup { + let mut entries = vec![ + wgpu::BindGroupEntry { binding: 0, resource: src.lighting_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(env_view) }, + wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(src.env_sampler) }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(src.brdf_lut_view) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(src.brdf_lut_sampler) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&src.shadow_map.depth_views[0]) }, + wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&src.shadow_map.depth_views[1]) }, + wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&src.shadow_map.depth_views[2]) }, + wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::Sampler(&src.shadow_map.sampler) }, + wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(diffuse_view) }, + ]; + if let Some(f) = src.froxel { + entries.extend(f.extra_lighting_bind_entries()); + } + device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some(label), + layout, + entries: &entries, + }) +} + +impl Renderer { + /// Build a group-1 lighting bind group for the given env-specular / + /// env-diffuse views. Everything else (UBO, BRDF LUT, shadow + /// cascade, froxel buffers when clustered) comes from `self`. + pub(super) fn make_lighting_bind_group( + &self, + label: &str, + env_view: &wgpu::TextureView, + diffuse_view: &wgpu::TextureView, + ) -> wgpu::BindGroup { + create_lighting_bind_group( + &self.device, + &self.lighting_layout, + label, + &LightingBindSources { + lighting_buffer: &self.lighting_buffer, + env_sampler: &self.env_sampler, + brdf_lut_view: &self.brdf_lut_view, + brdf_lut_sampler: &self.brdf_lut_sampler, + shadow_map: &self.shadow_map, + froxel: self.froxel.as_ref(), + }, + env_view, + diffuse_view, + ) + } +} diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index d4e306d..7bc3ed1 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -11,6 +11,8 @@ mod ssgi_pass; mod shadow_pass; mod postfx_chain; mod scene_pass; +mod froxel; +mod lighting; pub use occlusion::OcclusionCuller; use shaders::*; @@ -1394,6 +1396,12 @@ pub struct Renderer { // don't have to carry tangent vertex data or normal-map bindings. pub scene_pipeline: wgpu::RenderPipeline, pub scene_material_layout: wgpu::BindGroupLayout, + /// Froxel light clustering (task #23). `Some` when the device has + /// fragment-stage storage buffers (everything but WebGL2); the + /// scene shader is then compiled with the clustered point-light + /// loop and `lighting_layout` gains bindings 10–12. `None` keeps + /// the plain count-driven loop. + pub froxel: Option, /// 1×1 gray env fallback and its sampler — bound in the lighting /// bind group before any HDR is loaded. `load_env_from_hdr` /// rebuilds the lighting bind group to swap in the real env @@ -1674,99 +1682,13 @@ impl Renderer { // higher device limit). pipeline_3d doesn't reference the env // / BRDF bindings — WGSL lets bind group layouts expose more // than a shader consumes. - let lighting_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("lighting_layout"), - entries: &[ - wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: wgpu::ShaderStages::FRAGMENT, - 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::FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Float { filterable: true }, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 2, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 3, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Float { filterable: true }, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 4, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 5, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Depth, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 6, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Depth, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 7, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Depth, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 8, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Comparison), - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 9, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Float { filterable: true }, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, - }, - ], - }); + // Froxel clustering first — its presence decides whether the + // lighting layout grows bindings 10-12 and which point-light + // loop the scene shader is compiled with. + let froxel = froxel::FroxelPass::supported(&device) + .then(|| froxel::FroxelPass::new(&device)); + + let lighting_layout = lighting::create_lighting_layout(&device, froxel.is_some()); let lighting_uniforms = LightingUniforms::defaults(); let lighting_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("lighting_buffer"), @@ -2020,22 +1942,21 @@ impl Renderer { // group since the bind group binds the shadow depth view. let shadow_map = crate::shadows::ShadowMap::new(&device, Vertex3D::desc()); - let lighting_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("lighting_bg"), - layout: &lighting_layout, - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: lighting_buffer.as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&scene_env_default_view) }, - wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&env_sampler) }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&brdf_lut_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(&brdf_lut_sampler) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&shadow_map.depth_views[0]) }, - wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&shadow_map.depth_views[1]) }, - wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&shadow_map.depth_views[2]) }, - wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::Sampler(&shadow_map.sampler) }, - wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&scene_env_default_view) }, - ], - }); + let lighting_bind_group = lighting::create_lighting_bind_group( + &device, + &lighting_layout, + "lighting_bg", + &lighting::LightingBindSources { + lighting_buffer: &lighting_buffer, + env_sampler: &env_sampler, + brdf_lut_view: &brdf_lut_view, + brdf_lut_sampler: &brdf_lut_sampler, + shadow_map: &shadow_map, + froxel: froxel.as_ref(), + }, + &scene_env_default_view, + &scene_env_default_view, + ); // --- Default 1x1 white texture --- let white_data = [255u8, 255, 255, 255]; @@ -2864,9 +2785,17 @@ impl Renderer { // ============================================================ // Scene pipeline (retained scene-graph draws with normal maps) // ============================================================ + // Clustered devices get the froxel point-light loop spliced in + // place of the plain reference loop (same shading math — the + // many_point_lights golden enforces equivalence). + let scene_shader_source: std::borrow::Cow<'static, str> = if froxel.is_some() { + froxel::clustered_scene_shader(SCENE_SHADER).into() + } else { + SCENE_SHADER.into() + }; let scene_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("scene_shader"), - source: wgpu::ShaderSource::Wgsl(SCENE_SHADER.into()), + source: wgpu::ShaderSource::Wgsl(scene_shader_source), }); // Scene material layout: // 0: base_color texture 4: metallic_roughness texture @@ -6274,6 +6203,7 @@ impl Renderer { aerial_perspective_sampler, env_diffuse_texture: None, scene_pipeline, + froxel, scene_material_layout, _scene_env_default_texture: scene_env_default_texture, scene_env_default_view, @@ -8227,22 +8157,7 @@ impl Renderer { // LUT bindings stay put — only env tex/sampler + diffuse view // change. let diffuse_view_bg = diffuse_texture.create_view(&wgpu::TextureViewDescriptor::default()); - let new_lighting_bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("lighting_bg"), - layout: &self.lighting_layout, - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.lighting_buffer.as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&view) }, - wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.env_sampler) }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.brdf_lut_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(&self.brdf_lut_sampler) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[0]) }, - wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[1]) }, - wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[2]) }, - wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::Sampler(&self.shadow_map.sampler) }, - wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&diffuse_view_bg) }, - ], - }); + let new_lighting_bg = self.make_lighting_bind_group("lighting_bg", &view, &diffuse_view_bg); self.sky_texture = Some(texture); self.sky_bind_group = Some(bg); @@ -8555,22 +8470,7 @@ impl Renderer { self._scene_env_default_texture .create_view(&wgpu::TextureViewDescriptor::default()) }); - let new_bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("lighting_bg_panorama"), - layout: &self.lighting_layout, - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.lighting_buffer.as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&env_view) }, - wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.env_sampler) }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.brdf_lut_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(&self.brdf_lut_sampler) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[0]) }, - wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[1]) }, - wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[2]) }, - wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::Sampler(&self.shadow_map.sampler) }, - wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&diffuse_view) }, - ], - }); + let new_bg = self.make_lighting_bind_group("lighting_bg_panorama", &env_view, &diffuse_view); self.lighting_bind_group = new_bg; self.lighting_bg_is_procedural = false; } @@ -8582,22 +8482,11 @@ impl Renderer { /// every sun-move and the bind group's TextureView references /// remain valid. fn swap_lighting_bg_to_procedural(&mut self) { - let new_bg = self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("lighting_bg_procedural"), - layout: &self.lighting_layout, - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.lighting_buffer.as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::TextureView(&self.procedural_sky_equirect_full_view) }, - wgpu::BindGroupEntry { binding: 2, resource: wgpu::BindingResource::Sampler(&self.env_sampler) }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.brdf_lut_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::Sampler(&self.brdf_lut_sampler) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[0]) }, - wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[1]) }, - wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[2]) }, - wgpu::BindGroupEntry { binding: 8, resource: wgpu::BindingResource::Sampler(&self.shadow_map.sampler) }, - wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&self.procedural_env_diffuse_view) }, - ], - }); + let new_bg = self.make_lighting_bind_group( + "lighting_bg_procedural", + &self.procedural_sky_equirect_full_view, + &self.procedural_env_diffuse_view, + ); self.lighting_bind_group = new_bg; self.lighting_bg_is_procedural = true; } @@ -9285,6 +9174,7 @@ impl Renderer { const BLOOM_CHAIN: u32 = 4; const COMPOSED: u32 = 5; const LDR_FINAL: u32 = 6; + const FROXEL_CLUSTERS: u32 = 7; struct FrameCtx2<'a> { r: &'a mut Renderer, @@ -9296,6 +9186,15 @@ impl Renderer { } let mut g: Graph = Graph::new(); + g.push( + PassNode::new("froxel_assign", Box::new(|c: &mut FrameCtx2| { + // No-op when self.froxel is None (the method gates); + // the node stays in the graph so with_after pins + // never dangle. + c.r.record_froxel_assign(c.encoder); + })) + .with_writes(&[PassOutput::Transient(FROXEL_CLUSTERS)]), + ); g.push( PassNode::new("shadow", Box::new(|c: &mut FrameCtx2| { c.r.record_shadow_pass(c.encoder, c.profiler, c.scene); @@ -9306,7 +9205,12 @@ impl Renderer { PassNode::new("hdr_scene", Box::new(|c: &mut FrameCtx2| { c.r.record_hdr_scene_pass(c.encoder, c.profiler, c.scene); })) - .with_reads(&[PassInput::Shadow(0), PassInput::Shadow(1), PassInput::Shadow(2)]) + .with_reads(&[ + PassInput::Shadow(0), + PassInput::Shadow(1), + PassInput::Shadow(2), + PassInput::Transient(FROXEL_CLUSTERS), + ]) .with_writes(&[ PassOutput::HdrColor, PassOutput::MaterialRt, @@ -9314,7 +9218,7 @@ impl Renderer { PassOutput::AlbedoRt, PassOutput::Depth, ]) - .with_after(&["shadow"]), + .with_after(&["shadow", "froxel_assign"]), ); g.push( PassNode::new("translucent", Box::new(|c: &mut FrameCtx2| { diff --git a/native/shared/src/renderer/shaders/core.rs b/native/shared/src/renderer/shaders/core.rs index c7ee590..e3b4aaf 100644 --- a/native/shared/src/renderer/shaders/core.rs +++ b/native/shared/src/renderer/shaders/core.rs @@ -766,6 +766,10 @@ fn fs_main_scene(in: VertexOutputScene) -> SceneOut { base_color, metallic, roughness); } + // BEGIN-POINT-LIGHT-LOOP (replaced by the froxel-clustered variant + // at pipeline build on storage-buffer-capable backends — see + // renderer/froxel.rs; this plain loop is the WebGL fallback and the + // semantic reference the clustered path must match exactly) let pt_count = u32(lighting.point_light_count.x); for (var i = 0u; i < pt_count; i++) { let pl = lighting.point_lights[i]; @@ -780,6 +784,7 @@ fn fs_main_scene(in: VertexOutputScene) -> SceneOut { base_color, metallic, roughness); } } + // END-POINT-LIGHT-LOOP // --- Split-sum IBL (Karis 2013) --- // IBL_diffuse = base_color * (1 - kS_avg) * (1 - metallic) diff --git a/native/shared/tests/golden/many_point_lights_clustered_scene.png b/native/shared/tests/golden/many_point_lights_clustered_scene.png new file mode 100644 index 0000000000000000000000000000000000000000..23e65aa007b3791ac95291dc147e603b91e4a1cf GIT binary patch literal 4754 zcmeHLe^gWF8NLAlCq;JLCQPYpT5Y>ZL#+jDhpAMhj$>I{)5y=DSW$sSp*RHLqRx6+ zS~;gXg;hJ_*5hnRlc4+(1ViUR?9j&JS%4TvI3Qw!7&L}Je%x>G_vOajQ6@Z|lT_j#W8{Vsg?+c%QK<}8>)QB>H5*Vn#9Q37}?pk@cdOJ>nm+bL>Z+J?1j zHkAq{x&lu=l~#M{KKk9gPcC0MU6himDm1jPKRZ7@QTrS9i4`v{zBjYCbywM6kGztv zQ-7b=e@vljH=!j5$;)$(U;djC{7T}|G^ZOwyY}ymfwyZ%>yhN0@$xpK(B;$yOfHi# zwKrv}&(|s#^s-(#{4!Dy>z8KD_>=o<~CobbGWW+(ejXx5bExP*&^M^>GD$t zv{$4<3wEqN-4eqKCQl2g;Xw#Y6(AXK)#a923Tiu!?aL83LTB7W;6=|gQ8TshY zP*ta^f_d7YAKadwUmd^C(t4|s@D|N;^R79=a&48$_H#iEB*Hb4YfO z*eN{AlSPb5xSh5c@o;rszP5j_vdtugAcl6~@$#fx;<9nMvlgv$3$x!hEGoM^Kd<~a zJ8x_fEr$sg?OETjK&ETBjoyowXZ4D;ZtCCk*#ah6^NF8gjA2{|%au27Y&;n7GYG-y9G6xo}2iAAqrV`fA-zLNvDkP)Qz z-fc`cSwHmGT-Wg|h$XH(-;g+k#22)fOPg}@wyV!_Rs>l-ptTYD=vDSMuMDOwZ2_>nLZ{zI~u_Te^QF)5g z4)I4tEIh`W9Qqgdf?xvqlOxHpnj1NWlAVtPBwF!a<=wJ~(JG!r>j|0Xp!lQffvj){ zV!fZMkPgCRvS<0R2v*7s&N?J3B|?Cgrx1^u_aG}^rNHgxwNf4g5E-%kkLHaPz1>6~ z;RzFn;r?aO7<{LhnjPy|fMmO7f$tkR)=+nPDyKuj+*6fhMHg{_5;E=ZU|0t?CiOpj zDr8dcv>yeHB#Y0P=w{XUz~<%C)tc#)xcQ-AjP$@zJ!*#c3}Z$IdJ>uYU#zb9NMj$@ zY{@-ya{Oh0ta}yixsrC(mx|BG=jG+qjzXF2{<$8}32WKVyV>?r3F*(n)J7bHS}(P$ za&ylNnGq9RdUGUNJ2f?5t2CFlw6;zwG|WOcsa-WvugIq0(M{cOX7i+kcEw-Wo@g*+ zI8y*czEqYtxkBiAD=MZACb0OIK)lIK9G3vl9mk0XvC&olEb!M0md}4bs|N;z`b@zj zo7#`6-eF8?^nw1HJRnOeCMIGQ*`|9Mv%@>Cq-cIRvhszaOWjwvi zMM*~6!wT)lB2?NpexPpO4|T&k+HTdO{;%LcX?=f*(`7_~jcK?q^mF@#)8=w=6b5en zV%eoV0Ii#?veP@^WGpAaFRTv2Pp7+eC~_p>!LX3o`s=9|YPQg*3Kdis8#_C@r7o)x z)VNmiBmk2D_{>N`QB&pvEW=%M^rUnPzup>vy`J2=RA{*_>}D9j#n?$XUut9lI??0M zM7Px0U3xcn(M7vV0ws+~IX{>=S`VZ@++r^lnFv+nIxpXNHBBa-%Jg zQ?~GF^)(x@QYS@W&~;Y-9vymi7FZ3EdsFq(ksNpQ@_$mY43}K#RT7Ck9YLGY>sUzj zdZ@R{`&fBlWH1;jsXQgw7L^Ro%AfpQjK-ggsq28Lh|E$tF>@ToyaA(#ggf%|hhL6< zs$lw;oHL={?}=KK{`|hr>nfjxb|XJHfS$P++n}MldOIGb>}kqefs-Z26Z7`b+LUZy z@-Z7m2qBXQ-zml4y|mMvs<*)#QC?Is#@;?9{mA7CWRxBdd`R+Cj6v8yJZiuyCo(vD zlh>*6ET6UsKohLXcXo}c8M|x%@_bnD(?%+8qr#l_@xZ{cEk2bl5+e+!5->FOMZa|GyGU bC+91FEXt3&@e}x86SZO88*5cbIS2m*4BS_t literal 0 HcmV?d00001 diff --git a/native/shared/tests/golden_render.rs b/native/shared/tests/golden_render.rs index acdda08..753079f 100644 --- a/native/shared/tests/golden_render.rs +++ b/native/shared/tests/golden_render.rs @@ -17,7 +17,7 @@ //! regenerated with BLOOM_UPDATE_GOLDEN=1 `cargo test golden`. use bloom_shared::engine::EngineState; -use bloom_shared::renderer::Renderer; +use bloom_shared::renderer::{Renderer, Vertex3D}; const W: u32 = 256; const H: u32 = 256; @@ -219,45 +219,116 @@ fn golden_many_point_lights() { compare_or_update("many_point_lights", w, h, &rgba); } +/// Froxel-clustering parity gate. The golden for this test is generated +/// with `BLOOM_DISABLE_FROXEL=1` (the plain reference loop); the test +/// then runs through the clustered scene shader, so any divergence +/// between the two point-light paths — wrong cluster lookup, lights +/// missed by the sphere/AABB assignment, slice math drift — shows up as +/// a pixel diff. Unlike `golden_many_point_lights` (immediate-mode +/// `pipeline_3d`, which keeps the plain loop), this drives the retained +/// scene graph through `scene_pipeline`, the shader the clustered loop +/// is spliced into. #[test] -fn golden_lod_selection() { - use bloom_shared::renderer::Vertex3D; +fn golden_many_point_lights_clustered_scene() { let Some(mut eng) = try_engine() else { eprintln!("skip: no GPU adapter"); return; }; + // The gate is meaningless if the clustered path silently fell back + // to the reference loop. Storage buffers are available on every + // non-WebGL2 device this test runs on, so demand the froxel path + // unless the kill-switch is set (golden regeneration). + if std::env::var_os("BLOOM_DISABLE_FROXEL").is_none() { + assert!( + eng.renderer.froxel.is_some(), + "froxel clustering inactive on a storage-buffer-capable adapter — \ + parity test would silently test the reference loop against itself" + ); + } - fn cube_verts(half: f32, color: [f32; 4]) -> (Vec, Vec) { - // 6 faces, outward winding (matches scene-node conventions: - // prepare() recomputes bounds from positions). - let h = half; - let faces: [([f32; 3], [[f32; 3]; 4]); 6] = [ - ([0.0, 0.0, -1.0], [[-h,-h,-h],[ h,-h,-h],[ h, h,-h],[-h, h,-h]]), - ([0.0, 0.0, 1.0], [[ h,-h, h],[-h,-h, h],[-h, h, h],[ h, h, h]]), - ([-1.0, 0.0, 0.0], [[-h,-h, h],[-h,-h,-h],[-h, h,-h],[-h, h, h]]), - ([1.0, 0.0, 0.0], [[ h,-h,-h],[ h,-h, h],[ h, h, h],[ h, h,-h]]), - ([0.0, 1.0, 0.0], [[-h, h,-h],[ h, h,-h],[ h, h, h],[-h, h, h]]), - ([0.0, -1.0, 0.0], [[-h,-h, h],[ h,-h, h],[ h,-h,-h],[-h,-h,-h]]), - ]; - let mut verts = Vec::new(); - let mut idx = Vec::new(); - for (normal, vs) in faces { - let base = verts.len() as u32; - for p in vs { - verts.push(Vertex3D { - position: p, - normal, - color, - uv: [0.0, 0.0], - joints: [0.0; 4], - weights: [0.0; 4], - tangent: [0.0; 4], - }); - } - idx.extend_from_slice(&[base, base + 2, base + 1, base, base + 3, base + 2]); + // Floor (squashed cube) + a ring of cubes, lit by 40 colored point + // lights — enough that most froxels see only a few lights, so a + // broken cluster lookup cannot hide. + 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.45, 0.45, 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(14.0, 0.2, 14.0, 0.0, -0.1, 0.0)); + + let (cube_v, cube_i) = cube_verts(0.5, [0.8, 0.8, 0.8, 1.0]); + for i in 0..6u32 { + let t = i as f32 / 6.0 * std::f32::consts::TAU; + let node = eng.scene.create_node(); + eng.scene.update_geometry(node, cube_v.clone(), cube_i.clone()); + eng.scene.set_transform(node, scale_translate(1.0, 1.0, 1.0, t.cos() * 2.2, 0.5, t.sin() * 2.2)); + } + + let (w, h, rgba) = render(&mut eng, 6, |eng| { + let r = &mut eng.renderer; + r.set_clear_color(2.0, 2.0, 4.0, 255.0); + r.begin_mode_3d( + 6.0, 7.0, 6.0, + 0.0, 0.0, 0.0, + 0.0, 1.0, 0.0, + 60.0, 0.0, + ); + for i in 0..40u32 { + let t = i as f32 / 40.0 * std::f32::consts::TAU; + let (sx, sz) = (t.cos() * 4.0, t.sin() * 4.0); + let (lr, lg, lb) = ( + 0.5 + 0.5 * (t).cos(), + 0.5 + 0.5 * (t + 2.094).cos(), + 0.5 + 0.5 * (t + 4.189).cos(), + ); + r.add_point_light(sx, 1.2, sz, 3.5, lr, lg, lb, 1.6); } - (verts, idx) + }); + compare_or_update("many_point_lights_clustered_scene", w, h, &rgba); +} + +/// Unit cube as scene-node geometry — 6 faces, outward winding (matches +/// scene-node conventions: prepare() recomputes bounds from positions). +fn cube_verts(half: f32, color: [f32; 4]) -> (Vec, Vec) { + let h = half; + let faces: [([f32; 3], [[f32; 3]; 4]); 6] = [ + ([0.0, 0.0, -1.0], [[-h,-h,-h],[ h,-h,-h],[ h, h,-h],[-h, h,-h]]), + ([0.0, 0.0, 1.0], [[ h,-h, h],[-h,-h, h],[-h, h, h],[ h, h, h]]), + ([-1.0, 0.0, 0.0], [[-h,-h, h],[-h,-h,-h],[-h, h,-h],[-h, h, h]]), + ([1.0, 0.0, 0.0], [[ h,-h,-h],[ h,-h, h],[ h, h, h],[ h, h,-h]]), + ([0.0, 1.0, 0.0], [[-h, h,-h],[ h, h,-h],[ h, h, h],[-h, h, h]]), + ([0.0, -1.0, 0.0], [[-h,-h, h],[ h,-h, h],[ h,-h,-h],[-h,-h,-h]]), + ]; + let mut verts = Vec::new(); + let mut idx = Vec::new(); + for (normal, vs) in faces { + let base = verts.len() as u32; + for p in vs { + verts.push(Vertex3D { + position: p, + normal, + color, + uv: [0.0, 0.0], + joints: [0.0; 4], + weights: [0.0; 4], + tangent: [0.0; 4], + }); + } + idx.extend_from_slice(&[base, base + 2, base + 1, base, base + 3, base + 2]); } + (verts, idx) +} + +#[test] +fn golden_lod_selection() { + let Some(mut eng) = try_engine() else { + eprintln!("skip: no GPU adapter"); + return; + }; let (red_v, red_i) = cube_verts(0.5, [0.9, 0.1, 0.1, 1.0]); let (green_v, green_i) = cube_verts(0.5, [0.1, 0.9, 0.1, 1.0]); diff --git a/tools/file-lines-baseline.json b/tools/file-lines-baseline.json index 52aac05..e5d34f4 100644 --- a/tools/file-lines-baseline.json +++ b/tools/file-lines-baseline.json @@ -1,3 +1,3 @@ { - "native/shared/src/renderer/mod.rs": 11871 -} \ No newline at end of file + "native/shared/src/renderer/mod.rs": 11775 +}