From 542f32b8f488eb6186855f1a616c19a69762f02a Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Thu, 1 Jan 2026 22:59:46 -0500 Subject: [PATCH 1/3] Solari: Light leak reduction and stochastic WC update --- crates/bevy_solari/src/realtime/node.rs | 42 +++++++-- crates/bevy_solari/src/realtime/prepare.rs | 2 +- .../bevy_solari/src/realtime/restir_gi.wgsl | 2 +- .../bevy_solari/src/realtime/specular_gi.wgsl | 4 +- .../src/realtime/world_cache_query.wgsl | 18 ++-- .../src/realtime/world_cache_update.wgsl | 87 +++++++++++-------- examples/3d/solari.rs | 15 +++- 7 files changed, 118 insertions(+), 52 deletions(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index ac8fc69456e63..9f8ba959cca9c 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -48,7 +48,8 @@ pub struct SolariLightingNode { compact_world_cache_single_block_pipeline: CachedComputePipelineId, compact_world_cache_blocks_pipeline: CachedComputePipelineId, compact_world_cache_write_active_cells_pipeline: CachedComputePipelineId, - sample_for_world_cache_pipeline: CachedComputePipelineId, + sample_di_for_world_cache_pipeline: CachedComputePipelineId, + sample_gi_for_world_cache_pipeline: CachedComputePipelineId, blend_new_world_cache_samples_pipeline: CachedComputePipelineId, presample_light_tiles_pipeline: CachedComputePipelineId, di_initial_and_temporal_pipeline: CachedComputePipelineId, @@ -114,7 +115,8 @@ impl ViewNode for SolariLightingNode { Some(compact_world_cache_single_block_pipeline), Some(compact_world_cache_blocks_pipeline), Some(compact_world_cache_write_active_cells_pipeline), - Some(sample_for_world_cache_pipeline), + Some(sample_di_for_world_cache_pipeline), + Some(sample_gi_for_world_cache_pipeline), Some(blend_new_world_cache_samples_pipeline), Some(presample_light_tiles_pipeline), Some(di_initial_and_temporal_pipeline), @@ -136,7 +138,8 @@ impl ViewNode for SolariLightingNode { pipeline_cache.get_compute_pipeline(self.compact_world_cache_blocks_pipeline), pipeline_cache .get_compute_pipeline(self.compact_world_cache_write_active_cells_pipeline), - pipeline_cache.get_compute_pipeline(self.sample_for_world_cache_pipeline), + pipeline_cache.get_compute_pipeline(self.sample_di_for_world_cache_pipeline), + pipeline_cache.get_compute_pipeline(self.sample_gi_for_world_cache_pipeline), pipeline_cache.get_compute_pipeline(self.blend_new_world_cache_samples_pipeline), pipeline_cache.get_compute_pipeline(self.presample_light_tiles_pipeline), pipeline_cache.get_compute_pipeline(self.di_initial_and_temporal_pipeline), @@ -288,7 +291,17 @@ impl ViewNode for SolariLightingNode { pass.set_bind_group(2, None, &[]); - pass.set_pipeline(sample_for_world_cache_pipeline); + pass.set_pipeline(sample_di_for_world_cache_pipeline); + pass.set_push_constants( + 0, + bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), + ); + pass.dispatch_workgroups_indirect( + &solari_lighting_resources.world_cache_active_cells_dispatch, + 0, + ); + + pass.set_pipeline(sample_gi_for_world_cache_pipeline); pass.set_push_constants( 0, bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), @@ -351,6 +364,14 @@ impl ViewNode for SolariLightingNode { pass.dispatch_workgroups(dx, dy, 1); d.end(&mut pass); + drop(pass); + + diagnostics.record_u32( + render_context.command_encoder(), + &s.world_cache_active_cells_count.slice(..), + "solari_lighting/world_cache_active_cells_count", + ); + Ok(()) } } @@ -486,9 +507,16 @@ impl FromWorld for SolariLightingNode { Some(&bind_group_layout_world_cache_active_cells_dispatch), vec!["WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER".into()], ), - sample_for_world_cache_pipeline: create_pipeline( - "solari_lighting_sample_for_world_cache_pipeline", - "sample_radiance", + sample_di_for_world_cache_pipeline: create_pipeline( + "solari_lighting_sample_di_for_world_cache_pipeline", + "sample_di", + load_embedded_asset!(world, "world_cache_update.wgsl"), + None, + vec![], + ), + sample_gi_for_world_cache_pipeline: create_pipeline( + "solari_lighting_sample_gi_for_world_cache_pipeline", + "sample_gi", load_embedded_asset!(world, "world_cache_update.wgsl"), None, vec!["WORLD_CACHE_QUERY_ATOMIC_MAX_LIFETIME".into()], diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index b166c93b1754a..5bca6e7185f36 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -214,7 +214,7 @@ pub fn prepare_solari_lighting_resources( let world_cache_active_cells_count = render_device.create_buffer(&BufferDescriptor { label: Some("solari_lighting_world_cache_active_cells_count"), size: size_of::() as u64, - usage: BufferUsages::STORAGE, + usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC, mapped_at_creation: false, }); diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 8fc3089da0b33..de6e14fdf8dc2 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -113,7 +113,7 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 reservoir.radiance = direct_lighting.radiance; reservoir.unbiased_contribution_weight = direct_lighting.inverse_pdf * uniform_hemisphere_inverse_pdf(); #else - reservoir.radiance = query_world_cache(sample_point.world_position, sample_point.geometric_world_normal, view.world_position, WORLD_CACHE_CELL_LIFETIME, rng); + reservoir.radiance = query_world_cache(sample_point.world_position, sample_point.geometric_world_normal, view.world_position, ray.t, WORLD_CACHE_CELL_LIFETIME, rng); reservoir.unbiased_contribution_weight = uniform_hemisphere_inverse_pdf(); #endif diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index 8c3d0e2496cbe..364ccc2f6c6c4 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -73,7 +73,7 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { textureStore(view_output, global_id.xy, pixel_color); #ifdef VISUALIZE_WORLD_CACHE - textureStore(view_output, global_id.xy, vec4(query_world_cache(surface.world_position, surface.world_normal, view.world_position, WORLD_CACHE_CELL_LIFETIME, &rng) * view.exposure, 1.0)); + textureStore(view_output, global_id.xy, vec4(query_world_cache(surface.world_position, surface.world_normal, view.world_position, RAY_T_MAX, WORLD_CACHE_CELL_LIFETIME, &rng) * view.exposure, 1.0)); #endif } @@ -114,7 +114,7 @@ fn trace_glossy_path(initial_ray_origin: vec3, initial_wi: vec3, initi if path_spread * path_spread > a0 * get_cell_size(ray_hit.world_position, view.world_position) { // Path spread is wide enough, terminate path in the world cache let diffuse_brdf = ray_hit.material.base_color / PI; - radiance += throughput * diffuse_brdf * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, WORLD_CACHE_CELL_LIFETIME, rng); + radiance += throughput * diffuse_brdf * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, ray.t, WORLD_CACHE_CELL_LIFETIME, rng); break; } else if !surface_perfectly_specular { // Sample direct lighting (NEE) diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index c4c03b9aba340..53486d26d3f62 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -9,16 +9,18 @@ const WORLD_CACHE_MAX_TEMPORAL_SAMPLES: f32 = 32.0; const WORLD_CACHE_DIRECT_LIGHT_SAMPLE_COUNT: u32 = 32u; /// Maximum amount of distance to trace GI rays between two cache cells const WORLD_CACHE_MAX_GI_RAY_DISTANCE: f32 = 50.0; +/// Target upper limit on the amount of cache cells to update each frame +const WORLD_CACHE_TARGET_CELL_UPDATES: u32 = 40000u; /// Maximum amount of frames a cell can live for without being queried -const WORLD_CACHE_CELL_LIFETIME: u32 = 30u; +const WORLD_CACHE_CELL_LIFETIME: u32 = 10u; /// Maximum amount of attempts to find a cache entry after a hash collision const WORLD_CACHE_MAX_SEARCH_STEPS: u32 = 3u; /// Size of a cache cell at the lowest LOD in meters -const WORLD_CACHE_POSITION_BASE_CELL_SIZE: f32 = 0.25; +const WORLD_CACHE_POSITION_BASE_CELL_SIZE: f32 = 0.15; /// How fast the world cache transitions between LODs as a function of distance to the camera -const WORLD_CACHE_POSITION_LOD_SCALE: f32 = 8.0; +const WORLD_CACHE_POSITION_LOD_SCALE: f32 = 15.0; /// Marker value for an empty cell const WORLD_CACHE_EMPTY_CELL: u32 = 0u; @@ -46,12 +48,18 @@ struct WorldCacheGeometryData { @group(1) @binding(23) var world_cache_active_cells_count: u32; #ifndef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER -fn query_world_cache(world_position_in: vec3, world_normal: vec3, view_position: vec3, cell_lifetime: u32, rng: ptr) -> vec3 { +fn query_world_cache(world_position_in: vec3, world_normal: vec3, view_position: vec3, ray_t: f32, cell_lifetime: u32, rng: ptr) -> vec3 { var world_position = world_position_in; var cell_size = get_cell_size(world_position, view_position); - // https://tomclabault.github.io/blog/2025/regir, jitter_world_position_tangent_plane + if ray_t < cell_size { + // Prevent light leaks + cell_size = WORLD_CACHE_POSITION_BASE_CELL_SIZE; + } + #ifdef JITTER_WORLD_CACHE + // Jitter query point, which essentially blurs the cache a bit so it's not so grid-like + // https://tomclabault.github.io/blog/2025/regir, jitter_world_position_tangent_plane let TBN = orthonormalize(world_normal); let offset = (rand_vec2f(rng) * 2.0 - 1.0) * cell_size * 0.5; world_position += offset.x * TBN[0] + offset.y * TBN[1]; diff --git a/crates/bevy_solari/src/realtime/world_cache_update.wgsl b/crates/bevy_solari/src/realtime/world_cache_update.wgsl index 4f4e98f84d857..894e01270bcd0 100644 --- a/crates/bevy_solari/src/realtime/world_cache_update.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_update.wgsl @@ -8,6 +8,7 @@ WORLD_CACHE_MAX_TEMPORAL_SAMPLES, WORLD_CACHE_DIRECT_LIGHT_SAMPLE_COUNT, WORLD_CACHE_MAX_GI_RAY_DISTANCE, + WORLD_CACHE_TARGET_CELL_UPDATES, query_world_cache, world_cache_active_cells_count, world_cache_active_cell_indices, @@ -24,51 +25,67 @@ struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; @compute @workgroup_size(64, 1, 1) -fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) active_cell_id: vec3) { - if active_cell_id.x < world_cache_active_cells_count { - let cell_index = world_cache_active_cell_indices[active_cell_id.x]; - let geometry_data = world_cache_geometry_data[cell_index]; - var rng = cell_index + constants.frame_index; - - // TODO: Initialize newly active cells with data from an adjacent LOD - - var new_radiance = sample_random_light_ris(geometry_data.world_position, geometry_data.world_normal, workgroup_id.xy, &rng); - -#ifndef NO_MULTIBOUNCE - let ray_direction = sample_cosine_hemisphere(geometry_data.world_normal, &rng); - let ray_hit = trace_ray(geometry_data.world_position, ray_direction, RAY_T_MIN, WORLD_CACHE_MAX_GI_RAY_DISTANCE, RAY_FLAG_NONE); - if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { - let ray_hit = resolve_ray_hit_full(ray_hit); - let cell_life = atomicLoad(&world_cache_life[cell_index]); - new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, cell_life, &rng); - } -#endif +fn sample_di(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) active_cell_id: vec3) { + if active_cell_id.x >= world_cache_active_cells_count { return; } + + let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + let geometry_data = world_cache_geometry_data[cell_index]; + var rng = cell_index + constants.frame_index; + + if rand_f(&rng) >= f32(WORLD_CACHE_TARGET_CELL_UPDATES) / f32(world_cache_active_cells_count) { return; } + + let new_radiance = sample_random_light_ris(geometry_data.world_position, geometry_data.world_normal, workgroup_id.xy, &rng); + + world_cache_active_cells_new_radiance[active_cell_id.x] = new_radiance; +} + +@compute @workgroup_size(64, 1, 1) +fn sample_gi(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) active_cell_id: vec3) { + if active_cell_id.x >= world_cache_active_cells_count { return; } - world_cache_active_cells_new_radiance[active_cell_id.x] = new_radiance; + let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + let geometry_data = world_cache_geometry_data[cell_index]; + var rng = cell_index + constants.frame_index; + + var new_radiance = world_cache_active_cells_new_radiance[active_cell_id.x]; + + if rand_f(&rng) >= f32(WORLD_CACHE_TARGET_CELL_UPDATES) / f32(world_cache_active_cells_count) { return; } + + let ray_direction = sample_cosine_hemisphere(geometry_data.world_normal, &rng); + let ray = trace_ray(geometry_data.world_position, ray_direction, RAY_T_MIN, WORLD_CACHE_MAX_GI_RAY_DISTANCE, RAY_FLAG_NONE); + if ray.kind != RAY_QUERY_INTERSECTION_NONE { + let ray_hit = resolve_ray_hit_full(ray); + let cell_life = atomicLoad(&world_cache_life[cell_index]); + new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, ray.t, cell_life, &rng); } + + world_cache_active_cells_new_radiance[active_cell_id.x] = new_radiance; } @compute @workgroup_size(64, 1, 1) fn blend_new_samples(@builtin(global_invocation_id) active_cell_id: vec3) { - if active_cell_id.x < world_cache_active_cells_count { - let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + if active_cell_id.x >= world_cache_active_cells_count { return; } - let old_radiance = world_cache_radiance[cell_index]; - let new_radiance = world_cache_active_cells_new_radiance[active_cell_id.x]; - let luminance_delta = world_cache_luminance_deltas[cell_index]; + let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + var rng = cell_index + constants.frame_index; - // https://bsky.app/profile/gboisse.bsky.social/post/3m5blga3ftk2a - let sample_count = min(old_radiance.a + 1.0, WORLD_CACHE_MAX_TEMPORAL_SAMPLES); - let alpha = abs(luminance_delta) / max(luminance(old_radiance.rgb), 0.001); - let max_sample_count = mix(WORLD_CACHE_MAX_TEMPORAL_SAMPLES, 1.0, pow(saturate(alpha), 1.0 / 8.0)); - let blend_amount = 1.0 / min(sample_count, max_sample_count); + if rand_f(&rng) >= f32(WORLD_CACHE_TARGET_CELL_UPDATES) / f32(world_cache_active_cells_count) { return; } - let blended_radiance = mix(old_radiance.rgb, new_radiance, blend_amount); - let blended_luminance_delta = mix(luminance_delta, luminance(blended_radiance) - luminance(old_radiance.rgb), 1.0 / 8.0); + let old_radiance = world_cache_radiance[cell_index]; + let new_radiance = world_cache_active_cells_new_radiance[active_cell_id.x]; + let luminance_delta = world_cache_luminance_deltas[cell_index]; - world_cache_radiance[cell_index] = vec4(blended_radiance, sample_count); - world_cache_luminance_deltas[cell_index] = blended_luminance_delta; - } + // https://bsky.app/profile/gboisse.bsky.social/post/3m5blga3ftk2a + let sample_count = min(old_radiance.a + 1.0, WORLD_CACHE_MAX_TEMPORAL_SAMPLES); + let alpha = abs(luminance_delta) / max(luminance(old_radiance.rgb), 0.001); + let max_sample_count = mix(WORLD_CACHE_MAX_TEMPORAL_SAMPLES, 1.0, pow(saturate(alpha), 1.0 / 8.0)); + let blend_amount = 1.0 / min(sample_count, max_sample_count); + + let blended_radiance = mix(old_radiance.rgb, new_radiance, blend_amount); + let blended_luminance_delta = mix(luminance_delta, luminance(blended_radiance) - luminance(old_radiance.rgb), 1.0 / 8.0); + + world_cache_radiance[cell_index] = vec4(blended_radiance, sample_count); + world_cache_luminance_deltas[cell_index] = blended_luminance_delta; } fn sample_random_light_ris(world_position: vec3, world_normal: vec3, workgroup_id: vec2, rng: ptr) -> vec3 { diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index 4f74e33eceb0e..f7d84c4e73a7c 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -602,5 +602,18 @@ fn update_performance_text( "render/solari_lighting/specular_indirect_lighting/elapsed_gpu", ); text.push_str(&format!("{:17} TODO\n", "DLSS-RR")); - text.push_str(&format!("\n{:17} {total:.2} ms", "Total")); + text.push_str(&format!("{:17} {total:.2} ms\n", "Total")); + + if let Some(world_cache_active_cells_count) = diagnostics + .get(&DiagnosticPath::new( + "render/solari_lighting/world_cache_active_cells_count", + )) + .and_then(Diagnostic::average) + { + text.push_str(&format!( + "\nWorld cache cells {} ({:.0}%)", + world_cache_active_cells_count as u32, + (world_cache_active_cells_count * 100.0) / (2u64.pow(20) as f64) + )); + } } From 0b772512ddb5b66655d2d513479ceebf0bb72cec Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 2 Jan 2026 21:57:40 -0500 Subject: [PATCH 2/3] Misc refactors --- crates/bevy_solari/src/realtime/world_cache_update.wgsl | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/crates/bevy_solari/src/realtime/world_cache_update.wgsl b/crates/bevy_solari/src/realtime/world_cache_update.wgsl index 894e01270bcd0..c56bd9f55e992 100644 --- a/crates/bevy_solari/src/realtime/world_cache_update.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_update.wgsl @@ -47,8 +47,6 @@ fn sample_gi(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_inv let geometry_data = world_cache_geometry_data[cell_index]; var rng = cell_index + constants.frame_index; - var new_radiance = world_cache_active_cells_new_radiance[active_cell_id.x]; - if rand_f(&rng) >= f32(WORLD_CACHE_TARGET_CELL_UPDATES) / f32(world_cache_active_cells_count) { return; } let ray_direction = sample_cosine_hemisphere(geometry_data.world_normal, &rng); @@ -56,10 +54,9 @@ fn sample_gi(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_inv if ray.kind != RAY_QUERY_INTERSECTION_NONE { let ray_hit = resolve_ray_hit_full(ray); let cell_life = atomicLoad(&world_cache_life[cell_index]); - new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, ray.t, cell_life, &rng); + let radiance = query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, ray.t, cell_life, &rng); + world_cache_active_cells_new_radiance[active_cell_id.x] += ray_hit.material.base_color * radiance; } - - world_cache_active_cells_new_radiance[active_cell_id.x] = new_radiance; } @compute @workgroup_size(64, 1, 1) From 51f950a08cf436cffbfbe294701970c8c0b6c5e1 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 2 Jan 2026 21:58:39 -0500 Subject: [PATCH 3/3] Only limit cell size on first bounce --- crates/bevy_solari/src/realtime/node.rs | 2 +- crates/bevy_solari/src/realtime/world_cache_query.wgsl | 2 ++ 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 9f8ba959cca9c..65f6942883646 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -554,7 +554,7 @@ impl FromWorld for SolariLightingNode { "initial_and_temporal", load_embedded_asset!(world, "restir_gi.wgsl"), None, - vec![], + vec!["WORLD_CACHE_FIRST_BOUNCE_LIGHT_LEAK_PREVENTION".into()], ), gi_spatial_and_shade_pipeline: create_pipeline( "solari_lighting_gi_spatial_and_shade_pipeline", diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index 53486d26d3f62..378d2b25a638c 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -52,10 +52,12 @@ fn query_world_cache(world_position_in: vec3, world_normal: vec3, view var world_position = world_position_in; var cell_size = get_cell_size(world_position, view_position); +#ifdef WORLD_CACHE_FIRST_BOUNCE_LIGHT_LEAK_PREVENTION if ray_t < cell_size { // Prevent light leaks cell_size = WORLD_CACHE_POSITION_BASE_CELL_SIZE; } +#endif #ifdef JITTER_WORLD_CACHE // Jitter query point, which essentially blurs the cache a bit so it's not so grid-like