From 4aae6af65f547060e047cbcf5a38af228e6f1f68 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 23 Nov 2025 16:18:02 -0500 Subject: [PATCH 1/8] Fix solari GI shadow regression --- crates/bevy_solari/src/realtime/restir_gi.wgsl | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 7d3a60ac9e6d3..d41b3b70b76a3 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -67,9 +67,7 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let spatial = load_spatial_reservoir(global_id.xy, depth, surface.world_position, surface.world_normal, &rng); let merge_result = merge_reservoirs(input_reservoir, surface.world_position, surface.world_normal, surface.material.base_color / PI, spatial.reservoir, spatial.world_position, spatial.world_normal, spatial.diffuse_brdf, &rng); - var combined_reservoir = merge_result.merged_reservoir; - - combined_reservoir.radiance *= trace_point_visibility(surface.world_position, combined_reservoir.sample_point_world_position); + let combined_reservoir = merge_result.merged_reservoir; gi_reservoirs_a[pixel_index] = combined_reservoir; @@ -164,7 +162,9 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< } let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.main_pass_viewport.z); - let spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; + var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; + + spatial_reservoir.radiance *= trace_point_visibility(world_position, spatial_reservoir.sample_point_world_position); return NeighborInfo(spatial_reservoir, spatial_surface.world_position, spatial_surface.world_normal, spatial_diffuse_brdf); } From 7e165b550831afbaa4b49ea77dc968144878f6f8 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Wed, 26 Nov 2025 23:45:49 -0500 Subject: [PATCH 2/8] Solari: Misc improvements --- .../bevy_solari/src/realtime/restir_di.wgsl | 16 +++++++----- .../bevy_solari/src/realtime/restir_gi.wgsl | 26 +++++++++++-------- .../src/realtime/world_cache_query.wgsl | 3 ++- .../src/scene/raytracing_scene_bindings.wgsl | 2 +- 4 files changed, 28 insertions(+), 19 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index b31f70f49963e..e7492f3c47ff3 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -179,15 +179,19 @@ fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world } fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { - let spatial_pixel_id = get_neighbor_pixel_id(pixel_id, rng); + for (var i = 0u; i < 5u; i++) { + let spatial_pixel_id = get_neighbor_pixel_id(pixel_id, rng); - let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); - let spatial_surface = gpixel_resolve(textureLoad(gbuffer, spatial_pixel_id, 0), spatial_depth, spatial_pixel_id, view.main_pass_viewport.zw, view.world_from_clip); - if pixel_dissimilar(depth, world_position, spatial_surface.world_position, world_normal, spatial_surface.world_normal, view) { - return empty_reservoir(); + let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); + let spatial_surface = gpixel_resolve(textureLoad(gbuffer, spatial_pixel_id, 0), spatial_depth, spatial_pixel_id, view.main_pass_viewport.zw, view.world_from_clip); + if pixel_dissimilar(depth, world_position, spatial_surface.world_position, world_normal, spatial_surface.world_normal, view) { + continue; + } + + return load_reservoir_b(spatial_pixel_id); } - return load_reservoir_b(spatial_pixel_id); + return empty_reservoir(); } fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> vec2 { diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index b0bbfa90ba421..a6355517af5c2 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -152,21 +152,25 @@ fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world } fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> NeighborInfo { - let spatial_pixel_id = get_neighbor_pixel_id(pixel_id, rng); + for (var i = 0u; i < 5u; i++) { + let spatial_pixel_id = get_neighbor_pixel_id(pixel_id, rng); - let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); - let spatial_surface = gpixel_resolve(textureLoad(gbuffer, spatial_pixel_id, 0), spatial_depth, spatial_pixel_id, view.main_pass_viewport.zw, view.world_from_clip); - let spatial_diffuse_brdf = spatial_surface.material.base_color / PI; - if pixel_dissimilar(depth, world_position, spatial_surface.world_position, world_normal, spatial_surface.world_normal, view) { - return NeighborInfo(empty_reservoir(), spatial_surface.world_position, spatial_surface.world_normal, spatial_diffuse_brdf); - } + let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); + let spatial_surface = gpixel_resolve(textureLoad(gbuffer, spatial_pixel_id, 0), spatial_depth, spatial_pixel_id, view.main_pass_viewport.zw, view.world_from_clip); + let spatial_diffuse_brdf = spatial_surface.material.base_color / PI; + if pixel_dissimilar(depth, world_position, spatial_surface.world_position, world_normal, spatial_surface.world_normal, view) { + continue; + } + + let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.main_pass_viewport.z); + var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; - let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.main_pass_viewport.z); - var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; + spatial_reservoir.radiance *= trace_point_visibility(world_position, spatial_reservoir.sample_point_world_position); - spatial_reservoir.radiance *= trace_point_visibility(world_position, spatial_reservoir.sample_point_world_position); + return NeighborInfo(spatial_reservoir, spatial_surface.world_position, spatial_surface.world_normal, spatial_diffuse_brdf); + } - return NeighborInfo(spatial_reservoir, spatial_surface.world_position, spatial_surface.world_normal, spatial_diffuse_brdf); + return NeighborInfo(empty_reservoir(), world_position, world_normal, vec3(0.0)); } fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> vec2 { diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index 64a7d67fdd4ee..d6644919b255e 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -47,12 +47,13 @@ struct WorldCacheGeometryData { #ifndef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER fn query_world_cache(world_position: vec3, world_normal: vec3, view_position: vec3, cell_lifetime: u32, rng: ptr) -> vec3 { - let cell_size = get_cell_size(world_position, view_position); + var cell_size = get_cell_size(world_position, view_position); // 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; let jittered_position = world_position + offset.x * TBN[0] + offset.y * TBN[1]; + cell_size = get_cell_size(jittered_position, view_position); let world_position_quantized = bitcast>(quantize_position(jittered_position, cell_size)); let world_normal_quantized = bitcast>(quantize_normal(world_normal)); diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index 8b1e875fc5423..4d5ec4584db33 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -83,7 +83,7 @@ const LIGHT_NOT_PRESENT_THIS_FRAME = 0xFFFFFFFFu; @group(0) @binding(10) var directional_lights: array; @group(0) @binding(11) var previous_frame_light_id_translations: array; -const RAY_T_MIN = 0.01f; +const RAY_T_MIN = 0.001f; const RAY_T_MAX = 100000.0f; const RAY_NO_CULL = 0xFFu; From b2a892b072b0fe7c9b68ea3acb2edb0441c39329 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Thu, 27 Nov 2025 00:11:47 -0500 Subject: [PATCH 3/8] Checksum check --- crates/bevy_solari/src/realtime/world_cache_query.wgsl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index d6644919b255e..cafefffac58a8 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -122,7 +122,7 @@ fn compute_checksum(world_position: vec3, world_normal: vec3) -> u32 { key = iqint_hash(key + world_normal.x); key = iqint_hash(key + world_normal.y); key = iqint_hash(key + world_normal.z); - return key; + return max(key, 1u); // 0u is reserved for WORLD_CACHE_EMPTY_CELL } fn pcg_hash(input: u32) -> u32 { From e7dfcef7c37638911f83827144d3137d465be4ef Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 30 Nov 2025 15:18:48 -0500 Subject: [PATCH 4/8] Solari: MIS for ReSTIR DI resampling --- .../bevy_solari/src/realtime/restir_di.wgsl | 109 +++++++++++------- .../bevy_solari/src/realtime/restir_gi.wgsl | 10 +- crates/bevy_solari/src/scene/sampling.wgsl | 8 +- 3 files changed, 75 insertions(+), 52 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index e7492f3c47ff3..2d31c57f81f4f 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -9,7 +9,7 @@ #import bevy_solari::brdf::evaluate_brdf #import bevy_solari::gbuffer_utils::{gpixel_resolve, pixel_dissimilar, permute_pixel} #import bevy_solari::presample_light_tiles::{ResolvedLightSamplePacked, unpack_resolved_light_sample} -#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility} +#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility, balance_heuristic} #import bevy_solari::scene_bindings::{light_sources, previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME} @group(1) @binding(0) var view_output: texture_storage_2d; @@ -49,8 +49,9 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin let diffuse_brdf = surface.material.base_color / PI; let initial_reservoir = generate_initial_reservoir(surface.world_position, surface.world_normal, diffuse_brdf, workgroup_id.xy, &rng); - let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, surface.world_position, surface.world_normal); - let merge_result = merge_reservoirs(initial_reservoir, temporal_reservoir, surface.world_position, surface.world_normal, diffuse_brdf, &rng); + let temporal = load_temporal_reservoir(global_id.xy, depth, surface.world_position, surface.world_normal); + let merge_result = merge_reservoirs(initial_reservoir, surface.world_position, surface.world_normal, diffuse_brdf, + temporal.reservoir, temporal.world_position, temporal.world_normal, temporal.diffuse_brdf, &rng); store_reservoir_b(global_id.xy, merge_result.merged_reservoir); } @@ -71,8 +72,9 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = surface.material.base_color / PI; let input_reservoir = load_reservoir_b(global_id.xy); - let spatial_reservoir = load_spatial_reservoir(global_id.xy, depth, surface.world_position, surface.world_normal, &rng); - let merge_result = merge_reservoirs(input_reservoir, spatial_reservoir, surface.world_position, surface.world_normal, diffuse_brdf, &rng); + let spatial = load_spatial_reservoir(global_id.xy, depth, surface.world_position, surface.world_normal, &rng); + let merge_result = merge_reservoirs(input_reservoir, surface.world_position, surface.world_normal, diffuse_brdf, + spatial.reservoir, spatial.world_position, spatial.world_normal, spatial.diffuse_brdf, &rng); var combined_reservoir = merge_result.merged_reservoir; if reservoir_valid(combined_reservoir) { @@ -135,63 +137,67 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 return reservoir; } -fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> Reservoir { +fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.main_pass_viewport.zw)); // Check if the current pixel was off screen during the previous frame (current pixel is newly visible), // or if all temporal history should assumed to be invalid if any(temporal_pixel_id_float < vec2(0.0)) || any(temporal_pixel_id_float >= view.main_pass_viewport.zw) || bool(constants.reset) { - return empty_reservoir(); + return NeighborInfo(empty_reservoir(), vec3(0.0), vec3(0.0), vec3(0.0)); } let permuted_temporal_pixel_id = permute_pixel(vec2(temporal_pixel_id_float), constants.frame_index, view.viewport.zw); - var temporal_reservoir = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); + var temporal = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); // If permuted reprojection failed (tends to happen on object edges), try point reprojection - if !reservoir_valid(temporal_reservoir) { - temporal_reservoir = load_temporal_reservoir_inner(vec2(temporal_pixel_id_float), depth, world_position, world_normal); + if !reservoir_valid(temporal.reservoir) { + temporal = load_temporal_reservoir_inner(vec2(temporal_pixel_id_float), depth, world_position, world_normal); } // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) - let previous_light_id = temporal_reservoir.sample.light_id >> 16u; - let triangle_id = temporal_reservoir.sample.light_id & 0xFFFFu; + let previous_light_id = temporal.reservoir.sample.light_id >> 16u; + let triangle_id = temporal.reservoir.sample.light_id & 0xFFFFu; let light_id = previous_frame_light_id_translations[previous_light_id]; if light_id == LIGHT_NOT_PRESENT_THIS_FRAME { - return empty_reservoir(); + return NeighborInfo(empty_reservoir(), vec3(0.0), vec3(0.0), vec3(0.0)); } - temporal_reservoir.sample.light_id = (light_id << 16u) | triangle_id; + temporal.reservoir.sample.light_id = (light_id << 16u) | triangle_id; - temporal_reservoir.confidence_weight = min(temporal_reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); + temporal.reservoir.confidence_weight = min(temporal.reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); - return temporal_reservoir; + return temporal; } -fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> Reservoir { +fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { // Check if the pixel features have changed heavily between the current and previous frame let temporal_depth = textureLoad(previous_depth_buffer, temporal_pixel_id, 0); let temporal_surface = gpixel_resolve(textureLoad(previous_gbuffer, temporal_pixel_id, 0), temporal_depth, temporal_pixel_id, view.main_pass_viewport.zw, previous_view.world_from_clip); + let temporal_diffuse_brdf = temporal_surface.material.base_color / PI; if pixel_dissimilar(depth, world_position, temporal_surface.world_position, world_normal, temporal_surface.world_normal, view) { - return empty_reservoir(); + return NeighborInfo(empty_reservoir(), vec3(0.0), vec3(0.0), vec3(0.0)); } - return load_reservoir_a(temporal_pixel_id); + let temporal_reservoir = load_reservoir_a(temporal_pixel_id); + return NeighborInfo(temporal_reservoir, temporal_surface.world_position, temporal_surface.world_normal, temporal_diffuse_brdf); } -fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { +fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> NeighborInfo { for (var i = 0u; i < 5u; i++) { let spatial_pixel_id = get_neighbor_pixel_id(pixel_id, rng); let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); let spatial_surface = gpixel_resolve(textureLoad(gbuffer, spatial_pixel_id, 0), spatial_depth, spatial_pixel_id, view.main_pass_viewport.zw, view.world_from_clip); + let spatial_diffuse_brdf = spatial_surface.material.base_color / PI; if pixel_dissimilar(depth, world_position, spatial_surface.world_position, world_normal, spatial_surface.world_normal, view) { continue; } - return load_reservoir_b(spatial_pixel_id); + let spatial_reservoir = load_reservoir_b(spatial_pixel_id); + return NeighborInfo(spatial_reservoir, spatial_surface.world_position, spatial_surface.world_normal, spatial_diffuse_brdf); } - return empty_reservoir(); + return NeighborInfo(empty_reservoir(), world_position, world_normal, vec3(0.0)); } fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> vec2 { @@ -200,6 +206,13 @@ fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> return vec2(spatial_id); } +struct NeighborInfo { + reservoir: Reservoir, + world_position: vec3, + world_normal: vec3, + diffuse_brdf: vec3, +} + struct Reservoir { sample: LightSample, confidence_weight: f32, @@ -252,42 +265,56 @@ struct ReservoirMergeResult { fn merge_reservoirs( canonical_reservoir: Reservoir, + canonical_world_position: vec3, + canonical_world_normal: vec3, + canonical_diffuse_brdf: vec3, other_reservoir: Reservoir, - world_position: vec3, - world_normal: vec3, - diffuse_brdf: vec3, + other_world_position: vec3, + other_world_normal: vec3, + other_diffuse_brdf: vec3, rng: ptr, ) -> ReservoirMergeResult { - let canonical_contribution = reservoir_contribution(canonical_reservoir, world_position, world_normal, diffuse_brdf); - let other_contribution = reservoir_contribution(other_reservoir, world_position, world_normal, diffuse_brdf); - - let mis_weight_denominator = 1.0 / (canonical_reservoir.confidence_weight + other_reservoir.confidence_weight); - - let canonical_mis_weight = canonical_reservoir.confidence_weight * mis_weight_denominator; - let canonical_resampling_weight = canonical_mis_weight * (canonical_contribution.target_function * canonical_reservoir.unbiased_contribution_weight); - - let other_mis_weight = other_reservoir.confidence_weight * mis_weight_denominator; - let other_resampling_weight = other_mis_weight * (other_contribution.target_function * other_reservoir.unbiased_contribution_weight); + // Contributions for resampling + let canonical_contribution_canonical_sample = reservoir_contribution(canonical_reservoir, canonical_world_position, canonical_world_normal, canonical_diffuse_brdf); + let canonical_contribution_other_sample = reservoir_contribution(other_reservoir, canonical_world_position, canonical_world_normal, canonical_diffuse_brdf); + + // Extra contributions for MIS + let other_contribution_canonical_sample = reservoir_contribution(canonical_reservoir, other_world_position, other_world_normal, other_diffuse_brdf); + let other_contribution_other_sample = reservoir_contribution(other_reservoir, other_world_position, other_world_normal, other_diffuse_brdf); + + // Resampling weight for canonical sample + let canonical_sample_mis_weight = balance_heuristic( + canonical_reservoir.confidence_weight * canonical_contribution_canonical_sample.target_function, + other_reservoir.confidence_weight * other_contribution_canonical_sample.target_function, + ); + let canonical_sample_resampling_weight = canonical_sample_mis_weight * canonical_contribution_canonical_sample.target_function * canonical_reservoir.unbiased_contribution_weight; - let weight_sum = canonical_resampling_weight + other_resampling_weight; + // Resampling weight for other sample + let other_sample_mis_weight = balance_heuristic( + other_reservoir.confidence_weight * other_contribution_other_sample.target_function, + canonical_reservoir.confidence_weight * canonical_contribution_other_sample.target_function, + ); + let other_sample_resampling_weight = other_sample_mis_weight * canonical_contribution_other_sample.target_function * other_reservoir.unbiased_contribution_weight; + // Perform resampling var combined_reservoir = empty_reservoir(); combined_reservoir.confidence_weight = canonical_reservoir.confidence_weight + other_reservoir.confidence_weight; + let weight_sum = canonical_sample_resampling_weight + other_sample_resampling_weight; - if rand_f(rng) < other_resampling_weight / weight_sum { + if rand_f(rng) < other_sample_resampling_weight / weight_sum { combined_reservoir.sample = other_reservoir.sample; - let inverse_target_function = select(0.0, 1.0 / other_contribution.target_function, other_contribution.target_function > 0.0); + let inverse_target_function = select(0.0, 1.0 / canonical_contribution_other_sample.target_function, canonical_contribution_other_sample.target_function > 0.0); combined_reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - return ReservoirMergeResult(combined_reservoir, other_contribution.radiance, other_contribution.wi); + return ReservoirMergeResult(combined_reservoir, canonical_contribution_other_sample.radiance, canonical_contribution_other_sample.wi); } else { combined_reservoir.sample = canonical_reservoir.sample; - let inverse_target_function = select(0.0, 1.0 / canonical_contribution.target_function, canonical_contribution.target_function > 0.0); + let inverse_target_function = select(0.0, 1.0 / canonical_contribution_canonical_sample.target_function, canonical_contribution_canonical_sample.target_function > 0.0); combined_reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - return ReservoirMergeResult(combined_reservoir, canonical_contribution.radiance, canonical_contribution.wi); + return ReservoirMergeResult(combined_reservoir, canonical_contribution_canonical_sample.radiance, canonical_contribution_canonical_sample.wi); } } diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index a6355517af5c2..233c20897857b 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -7,7 +7,7 @@ #import bevy_render::view::View #import bevy_solari::brdf::evaluate_diffuse_brdf #import bevy_solari::gbuffer_utils::{gpixel_resolve, pixel_dissimilar, permute_pixel} -#import bevy_solari::sampling::{sample_random_light, trace_point_visibility} +#import bevy_solari::sampling::{sample_random_light, trace_point_visibility, balance_heuristic} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} #import bevy_solari::world_cache::{query_world_cache, WORLD_CACHE_CELL_LIFETIME} @@ -321,11 +321,3 @@ fn merge_reservoirs( return ReservoirMergeResult(combined_reservoir, canonical_sample_radiance); } } - -fn balance_heuristic(x: f32, y: f32) -> f32 { - let sum = x + y; - if sum == 0.0 { - return 0.0; - } - return x / sum; -} diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 8385f9b3999e6..3030db514eeaf 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -6,11 +6,15 @@ #import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LightSource, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full, ResolvedRayHitFull} fn power_heuristic(f: f32, g: f32) -> f32 { - return f * f / (f * f + g * g); + return balance_heuristic(f * f, g * g); } fn balance_heuristic(f: f32, g: f32) -> f32 { - return f / (f + g); + let sum = f + g; + if sum == 0.0 { + return 0.0; + } + return f / sum; } // https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 1) From 742c79a24607ad59a97b33df00b696a4fcddc0e4 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 2 Dec 2025 21:32:38 -0500 Subject: [PATCH 5/8] Prevent NaNs --- .../src/pathtracer/pathtracer.wgsl | 13 ++-- .../bevy_solari/src/realtime/specular_gi.wgsl | 69 +++++++++++++++---- crates/bevy_solari/src/scene/sampling.wgsl | 14 ++-- 3 files changed, 74 insertions(+), 22 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 77fc834b2f5a5..c2402deef7e40 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -4,7 +4,7 @@ #import bevy_render::maths::PI #import bevy_render::view::View #import bevy_solari::brdf::evaluate_brdf -#import bevy_solari::sampling::{sample_random_light, random_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, balance_heuristic, power_heuristic} +#import bevy_solari::sampling::{sample_random_light, random_emissive_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, power_heuristic} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, ResolvedRayHitFull, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var accumulation_texture: texture_storage_2d; @@ -48,7 +48,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { var mis_weight = 1.0; if !bounce_was_perfect_reflection { - let p_light = random_light_pdf(ray_hit); + let p_light = random_emissive_light_pdf(ray_hit); mis_weight = power_heuristic(p_bounce, p_light); } radiance += mis_weight * throughput * ray_hit.material.emissive; @@ -57,8 +57,13 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { let is_perfectly_specular = ray_hit.material.roughness <= 0.001 && ray_hit.material.metallic > 0.9999; if !is_perfectly_specular { let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); - let pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); - mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); + + mis_weight = 1.0; + if direct_lighting.brdf_rays_can_hit { + let pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); + mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); + } + let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); radiance += mis_weight * throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; } diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index 335c5588271de..176be1c619258 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -3,8 +3,8 @@ #import bevy_render::view::View #import bevy_solari::brdf::{evaluate_brdf, evaluate_specular_brdf} #import bevy_solari::gbuffer_utils::gpixel_resolve -#import bevy_solari::sampling::{sample_ggx_vndf, ggx_vndf_pdf} -#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} +#import bevy_solari::sampling::{sample_random_light, random_emissive_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, power_heuristic} +#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, ResolvedRayHitFull, RAY_T_MIN, RAY_T_MAX} #import bevy_solari::world_cache::{query_world_cache, WORLD_CACHE_CELL_LIFETIME} @group(1) @binding(0) var view_output: texture_storage_2d; @@ -15,6 +15,9 @@ struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; +const DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD: f32 = 0.4; +const WORLD_CACHE_TERMINATION_ROUGHNESS_THRESHOLD: f32 = 0.4; + @compute @workgroup_size(8, 8, 1) fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { if any(global_id.xy >= vec2u(view.main_pass_viewport.zw)) { return; } @@ -32,7 +35,7 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { var radiance: vec3; var wi: vec3; - if surface.material.roughness > 0.1 { + if surface.material.roughness > DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD { // Surface is very rough, reuse the ReSTIR GI reservoir let gi_reservoir = gi_reservoirs_a[pixel_index]; wi = normalize(gi_reservoir.sample_point_world_position - surface.world_position); @@ -68,6 +71,8 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { fn trace_glossy_path(initial_ray_origin: vec3, initial_wi: vec3, rng: ptr) -> vec3 { var ray_origin = initial_ray_origin; var wi = initial_wi; + var surface_perfectly_specular = false; + var p_bounce = 0.0; // Trace up to three bounces, getting the net throughput from them var radiance = vec3(0.0); @@ -78,34 +83,72 @@ fn trace_glossy_path(initial_ray_origin: vec3, initial_wi: vec3, rng: if ray.kind == RAY_QUERY_INTERSECTION_NONE { break; } let ray_hit = resolve_ray_hit_full(ray); - // Add world cache contribution - 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); - - // Surface is very rough, terminate path in the world cache - if ray_hit.material.roughness > 0.1 && i != 0u { break; } - - // Sample new ray direction from the GGX BRDF for next bounce let TBN = calculate_tbn_mikktspace(ray_hit.world_normal, ray_hit.world_tangent); let T = TBN[0]; let B = TBN[1]; let N = TBN[2]; + let wo = -wi; let wo_tangent = vec3(dot(wo, T), dot(wo, B), dot(wo, N)); + + // Add emissive contribution (but not on the first bounce, since ReSTIR DI handles that) + if i != 0u { + radiance += throughput * emissive_mis_weight(p_bounce, ray_hit, surface_perfectly_specular) * ray_hit.material.emissive; + } + + // Should not perform NEE for mirror-like surfaces + surface_perfectly_specular = ray_hit.material.roughness <= 0.001 && ray_hit.material.metallic > 0.9999; + + if ray_hit.material.roughness > WORLD_CACHE_TERMINATION_ROUGHNESS_THRESHOLD && i != 0u { + // Surface is very rough, 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); + break; + } else if !surface_perfectly_specular { + // Sample direct lighting (NEE) + let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, rng); + let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); + let mis_weight = nee_mis_weight(direct_lighting.inverse_pdf, direct_lighting.brdf_rays_can_hit, wo_tangent, direct_lighting.wi, ray_hit, TBN); + radiance += throughput * mis_weight * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; + } + + // Sample new ray direction from the GGX BRDF for next bounce let wi_tangent = sample_ggx_vndf(wo_tangent, ray_hit.material.roughness, rng); wi = wi_tangent.x * T + wi_tangent.y * B + wi_tangent.z * N; ray_origin = ray_hit.world_position; // Update throughput for next bounce - let pdf = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); + p_bounce = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); let brdf = evaluate_brdf(N, wo, wi, ray_hit.material); let cos_theta = saturate(dot(wi, N)); - throughput *= (brdf * cos_theta) / pdf; + throughput *= (brdf * cos_theta) / p_bounce; } return radiance; } +fn emissive_mis_weight(p_bounce: f32, ray_hit: ResolvedRayHitFull, previous_surface_perfectly_specular: bool) -> f32 { + if previous_surface_perfectly_specular { return 1.0; } + + let p_light = random_emissive_light_pdf(ray_hit); + return power_heuristic(p_bounce, p_light); +} + +fn nee_mis_weight(inverse_p_light: f32, brdf_rays_can_hit: bool, wo_tangent: vec3, wi: vec3, ray_hit: ResolvedRayHitFull, TBN: mat3x3) -> f32 { + if !brdf_rays_can_hit { + return 1.0; + } + + let T = TBN[0]; + let B = TBN[1]; + let N = TBN[2]; + let wi_tangent = vec3(dot(wi, T), dot(wi, B), dot(wi, N)); + + let p_light = 1.0 / inverse_p_light; + let p_bounce = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); + return power_heuristic(p_light, p_bounce); +} + // Don't adjust the size of this struct without also adjusting GI_RESERVOIR_STRUCT_SIZE. struct Reservoir { sample_point_world_position: vec3, diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 3030db514eeaf..d0f8d8bdb3ce3 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -14,11 +14,15 @@ fn balance_heuristic(f: f32, g: f32) -> f32 { if sum == 0.0 { return 0.0; } - return f / sum; + return max(0.0, f / sum); } // https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 1) fn sample_ggx_vndf(wi_tangent: vec3, roughness: f32, rng: ptr) -> vec3 { + if roughness <= 0.01 { + return vec3(-wi_tangent.xy, wi_tangent.z); + } + let i = wi_tangent; let rand = rand_vec2f(rng); let i_std = normalize(vec3(i.xy * roughness, i.z)); @@ -73,6 +77,7 @@ struct LightContribution { radiance: vec3, inverse_pdf: f32, wi: vec3, + brdf_rays_can_hit: bool, } struct LightContributionNoPdf { @@ -92,10 +97,9 @@ fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rn return light_contribution; } -fn random_light_pdf(hit: ResolvedRayHitFull) -> f32 { +fn random_emissive_light_pdf(hit: ResolvedRayHitFull) -> f32 { let light_count = arrayLength(&light_sources); - let p_light = 1.0 / f32(light_count); - return p_light / (hit.triangle_area * f32(hit.triangle_count)); + return 1.0 / (f32(light_count) * f32(hit.triangle_count) * hit.triangle_area); } fn generate_random_light_sample(rng: ptr) -> GenerateRandomLightSampleResult { @@ -173,7 +177,7 @@ fn calculate_resolved_light_contribution(resolved_light_sample: ResolvedLightSam let radiance = resolved_light_sample.radiance * cos_theta_origin * (cos_theta_light / light_distance_squared); - return LightContribution(radiance, resolved_light_sample.inverse_pdf, wi); + return LightContribution(radiance, resolved_light_sample.inverse_pdf, wi, resolved_light_sample.world_position.w == 1.0); } fn resolve_and_calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContributionNoPdf { From 79c7b9aa66156bdf756d30d42cadcfe991654ac8 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 2 Dec 2025 21:38:28 -0500 Subject: [PATCH 6/8] Move visibility test to after reservoir storage to avoid bias --- crates/bevy_solari/src/realtime/restir_di.wgsl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 2d31c57f81f4f..878768dfc9e13 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -77,13 +77,13 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { spatial.reservoir, spatial.world_position, spatial.world_normal, spatial.diffuse_brdf, &rng); var combined_reservoir = merge_result.merged_reservoir; + store_reservoir_a(global_id.xy, combined_reservoir); + if reservoir_valid(combined_reservoir) { let resolved_light_sample = resolve_light_sample(combined_reservoir.sample, light_sources[combined_reservoir.sample.light_id >> 16u]); combined_reservoir.unbiased_contribution_weight *= trace_light_visibility(surface.world_position, resolved_light_sample.world_position); } - store_reservoir_a(global_id.xy, combined_reservoir); - let wo = normalize(view.world_position - surface.world_position); let brdf = evaluate_brdf(surface.world_normal, wo, merge_result.wi, surface.material); From 430bfc81b042c3b741f5598cc5754435fb036181 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 5 Dec 2025 09:41:31 -0500 Subject: [PATCH 7/8] Add comments --- crates/bevy_solari/src/realtime/restir_di.wgsl | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 878768dfc9e13..0601a6f153a6f 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -77,13 +77,21 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { spatial.reservoir, spatial.world_position, spatial.world_normal, spatial.diffuse_brdf, &rng); var combined_reservoir = merge_result.merged_reservoir; + // More accuracy, less stability +#ifndef BIASED_RESAMPLING store_reservoir_a(global_id.xy, combined_reservoir); +#endif if reservoir_valid(combined_reservoir) { let resolved_light_sample = resolve_light_sample(combined_reservoir.sample, light_sources[combined_reservoir.sample.light_id >> 16u]); combined_reservoir.unbiased_contribution_weight *= trace_light_visibility(surface.world_position, resolved_light_sample.world_position); } + // More stability, less accuracy (shadows extend further out than they should) +#ifdef BIASED_RESAMPLING + store_reservoir_a(global_id.xy, combined_reservoir); +#endif + let wo = normalize(view.world_position - surface.world_position); let brdf = evaluate_brdf(surface.world_normal, wo, merge_result.wi, surface.material); From 576865f9054bb155dec714c12a81a986e6612ce6 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 5 Dec 2025 10:03:09 -0500 Subject: [PATCH 8/8] Make it easier to disable world cache jitter if needed --- .../bevy_solari/src/realtime/world_cache_query.wgsl | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index cafefffac58a8..6414583dafe0e 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -46,16 +46,19 @@ 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: 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, 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 +#ifndef NO_JITTER_WORLD_CACHE let TBN = orthonormalize(world_normal); let offset = (rand_vec2f(rng) * 2.0 - 1.0) * cell_size * 0.5; - let jittered_position = world_position + offset.x * TBN[0] + offset.y * TBN[1]; - cell_size = get_cell_size(jittered_position, view_position); + world_position += offset.x * TBN[0] + offset.y * TBN[1]; + cell_size = get_cell_size(world_position, view_position); +#endif - let world_position_quantized = bitcast>(quantize_position(jittered_position, cell_size)); + let world_position_quantized = bitcast>(quantize_position(world_position, cell_size)); let world_normal_quantized = bitcast>(quantize_normal(world_normal)); var key = compute_key(world_position_quantized, world_normal_quantized); let checksum = compute_checksum(world_position_quantized, world_normal_quantized); @@ -77,7 +80,7 @@ fn query_world_cache(world_position: vec3, world_normal: vec3, view_po return world_cache_radiance[key].rgb; } else if existing_checksum == WORLD_CACHE_EMPTY_CELL { // Cell is empty - initialize it - world_cache_geometry_data[key].world_position = jittered_position; + world_cache_geometry_data[key].world_position = world_position; world_cache_geometry_data[key].world_normal = world_normal; return vec3(0.0); } else {