From 4bdae6a860a1389b024934f21deab06ea2bbb83d Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 1 Jul 2025 22:38:04 -0400 Subject: [PATCH 01/12] 1-bounce GI (no ReSTIR) --- crates/bevy_solari/src/realtime/mod.rs | 1 + crates/bevy_solari/src/realtime/node.rs | 62 +++++++++++++----- crates/bevy_solari/src/realtime/prepare.rs | 20 +++--- .../bevy_solari/src/realtime/restir_di.wgsl | 22 +++---- .../bevy_solari/src/realtime/restir_gi.wgsl | 63 +++++++++++++++++++ crates/bevy_solari/src/scene/sampling.wgsl | 11 ++++ 6 files changed, 141 insertions(+), 38 deletions(-) create mode 100644 crates/bevy_solari/src/realtime/restir_gi.wgsl diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index b6a8f27a31dce..96bedd0b9714c 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -28,6 +28,7 @@ pub struct SolariLightingPlugin; impl Plugin for SolariLightingPlugin { fn build(&self, app: &mut App) { embedded_asset!(app, "restir_di.wgsl"); + embedded_asset!(app, "restir_gi.wgsl"); app.register_type::() .insert_resource(DefaultOpaqueRendererMethod::deferred()); diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 8fb97f84b21db..2561a897e65aa 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -35,8 +35,9 @@ pub mod graph { pub struct SolariLightingNode { bind_group_layout: BindGroupLayout, - initial_and_temporal_pipeline: CachedComputePipelineId, - spatial_and_shade_pipeline: CachedComputePipelineId, + di_initial_and_temporal_pipeline: CachedComputePipelineId, + di_spatial_and_shade_pipeline: CachedComputePipelineId, + gi_initial_and_temporal_pipeline: CachedComputePipelineId, } impl ViewNode for SolariLightingNode { @@ -71,8 +72,9 @@ impl ViewNode for SolariLightingNode { let previous_view_uniforms = world.resource::(); let frame_count = world.resource::(); let ( - Some(initial_and_temporal_pipeline), - Some(spatial_and_shade_pipeline), + Some(di_initial_and_temporal_pipeline), + Some(di_spatial_and_shade_pipeline), + Some(gi_initial_and_temporal_pipeline), Some(scene_bindings), Some(viewport), Some(gbuffer), @@ -81,8 +83,9 @@ impl ViewNode for SolariLightingNode { Some(view_uniforms), Some(previous_view_uniforms), ) = ( - pipeline_cache.get_compute_pipeline(self.initial_and_temporal_pipeline), - pipeline_cache.get_compute_pipeline(self.spatial_and_shade_pipeline), + pipeline_cache.get_compute_pipeline(self.di_initial_and_temporal_pipeline), + pipeline_cache.get_compute_pipeline(self.di_spatial_and_shade_pipeline), + pipeline_cache.get_compute_pipeline(self.gi_initial_and_temporal_pipeline), &scene_bindings.bind_group, camera.physical_viewport_size, view_prepass_textures.deferred_view(), @@ -100,8 +103,12 @@ impl ViewNode for SolariLightingNode { &self.bind_group_layout, &BindGroupEntries::sequential(( view_target.get_unsampled_color_attachment().view, - solari_lighting_resources.reservoirs_a.as_entire_binding(), - solari_lighting_resources.reservoirs_b.as_entire_binding(), + solari_lighting_resources + .di_reservoirs_a + .as_entire_binding(), + solari_lighting_resources + .di_reservoirs_b + .as_entire_binding(), gbuffer, depth_buffer, motion_vectors, @@ -134,14 +141,17 @@ impl ViewNode for SolariLightingNode { ], ); - pass.set_pipeline(initial_and_temporal_pipeline); + pass.set_pipeline(di_initial_and_temporal_pipeline); pass.set_push_constants( 0, bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), ); pass.dispatch_workgroups(viewport.x.div_ceil(8), viewport.y.div_ceil(8), 1); - pass.set_pipeline(spatial_and_shade_pipeline); + pass.set_pipeline(di_spatial_and_shade_pipeline); + pass.dispatch_workgroups(viewport.x.div_ceil(8), viewport.y.div_ceil(8), 1); + + pass.set_pipeline(gi_initial_and_temporal_pipeline); pass.dispatch_workgroups(viewport.x.div_ceil(8), viewport.y.div_ceil(8), 1); pass_span.end(&mut pass); @@ -188,7 +198,7 @@ impl FromWorld for SolariLightingNode { ( texture_storage_2d( ViewTarget::TEXTURE_FORMAT_HDR, - StorageTextureAccess::WriteOnly, + StorageTextureAccess::ReadWrite, ), storage_buffer_sized(false, None), storage_buffer_sized(false, None), @@ -203,9 +213,9 @@ impl FromWorld for SolariLightingNode { ), ); - let initial_and_temporal_pipeline = + let di_initial_and_temporal_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_initial_and_temporal_pipeline".into()), + label: Some("solari_lighting_di_initial_and_temporal_pipeline".into()), layout: vec![ scene_bindings.bind_group_layout.clone(), bind_group_layout.clone(), @@ -220,9 +230,9 @@ impl FromWorld for SolariLightingNode { zero_initialize_workgroup_memory: false, }); - let spatial_and_shade_pipeline = + let di_spatial_and_shade_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_spatial_and_shade_pipeline".into()), + label: Some("solari_lighting_di_spatial_and_shade_pipeline".into()), layout: vec![ scene_bindings.bind_group_layout.clone(), bind_group_layout.clone(), @@ -237,10 +247,28 @@ impl FromWorld for SolariLightingNode { zero_initialize_workgroup_memory: false, }); + let gi_initial_and_temporal_pipeline = + pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { + label: Some("solari_lighting_gi_initial_and_temporal_pipeline".into()), + layout: vec![ + scene_bindings.bind_group_layout.clone(), + bind_group_layout.clone(), + ], + push_constant_ranges: vec![PushConstantRange { + stages: ShaderStages::COMPUTE, + range: 0..8, + }], + shader: load_embedded_asset!(world, "restir_gi.wgsl"), + shader_defs: vec![], + entry_point: "initial_and_temporal".into(), + zero_initialize_workgroup_memory: false, + }); + Self { bind_group_layout, - initial_and_temporal_pipeline, - spatial_and_shade_pipeline, + di_initial_and_temporal_pipeline, + di_spatial_and_shade_pipeline, + gi_initial_and_temporal_pipeline, } } } diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 992a75c451134..3c454727bd8cc 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -18,13 +18,13 @@ use bevy_render::{ }; /// Size of a Reservoir shader struct in bytes. -const RESERVOIR_STRUCT_SIZE: u64 = 32; +const DI_RESERVOIR_STRUCT_SIZE: u64 = 32; /// Internal rendering resources used for Solari lighting. #[derive(Component)] pub struct SolariLightingResources { - pub reservoirs_a: Buffer, - pub reservoirs_b: Buffer, + pub di_reservoirs_a: Buffer, + pub di_reservoirs_b: Buffer, pub previous_gbuffer: (Texture, TextureView), pub previous_depth: (Texture, TextureView), pub view_size: UVec2, @@ -47,17 +47,17 @@ pub fn prepare_solari_lighting_resources( continue; } - let size = (view_size.x * view_size.y) as u64 * RESERVOIR_STRUCT_SIZE; + let size = (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE; - let reservoirs_a = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_reservoirs_a"), + let di_reservoirs_a = render_device.create_buffer(&BufferDescriptor { + label: Some("solari_lighting_di_reservoirs_a"), size, usage: BufferUsages::STORAGE, mapped_at_creation: false, }); - let reservoirs_b = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_reservoirs_b"), + let di_reservoirs_b = render_device.create_buffer(&BufferDescriptor { + label: Some("solari_lighting_di_reservoirs_b"), size, usage: BufferUsages::STORAGE, mapped_at_creation: false, @@ -88,8 +88,8 @@ pub fn prepare_solari_lighting_resources( let previous_depth_view = previous_depth.create_view(&TextureViewDescriptor::default()); commands.entity(entity).insert(SolariLightingResources { - reservoirs_a, - reservoirs_b, + di_reservoirs_a, + di_reservoirs_b, previous_gbuffer: (previous_gbuffer, previous_gbuffer_view), previous_depth: (previous_depth, previous_depth_view), view_size, diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 70de4564cc7d1..c9efc6fb205d5 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -10,9 +10,9 @@ #import bevy_solari::sampling::{LightSample, generate_random_light_sample, calculate_light_contribution, trace_light_visibility, sample_disk} #import bevy_solari::scene_bindings::{previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME} -@group(1) @binding(0) var view_output: texture_storage_2d; -@group(1) @binding(1) var reservoirs_a: array; -@group(1) @binding(2) var reservoirs_b: array; +@group(1) @binding(0) var view_output: texture_storage_2d; +@group(1) @binding(1) var di_reservoirs_a: array; +@group(1) @binding(2) var di_reservoirs_b: array; @group(1) @binding(3) var gbuffer: texture_2d; @group(1) @binding(4) var depth_buffer: texture_depth_2d; @group(1) @binding(5) var motion_vectors: texture_2d; @@ -38,7 +38,7 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - reservoirs_b[pixel_index] = empty_reservoir(); + di_reservoirs_b[pixel_index] = empty_reservoir(); return; } let gpixel = textureLoad(gbuffer, global_id.xy, 0); @@ -51,7 +51,7 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal); let combined_reservoir = merge_reservoirs(initial_reservoir, temporal_reservoir, world_position, world_normal, diffuse_brdf, &rng); - reservoirs_b[pixel_index] = combined_reservoir.merged_reservoir; + di_reservoirs_b[pixel_index] = combined_reservoir.merged_reservoir; } @compute @workgroup_size(8, 8, 1) @@ -63,7 +63,7 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - reservoirs_a[pixel_index] = empty_reservoir(); + di_reservoirs_a[pixel_index] = empty_reservoir(); textureStore(view_output, global_id.xy, vec4(vec3(0.0), 1.0)); return; } @@ -74,12 +74,12 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = base_color / PI; let emissive = rgb9e5_to_vec3_(gpixel.g); - let input_reservoir = reservoirs_b[pixel_index]; + let input_reservoir = di_reservoirs_b[pixel_index]; let spatial_reservoir = load_spatial_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let merge_result = merge_reservoirs(input_reservoir, spatial_reservoir, world_position, world_normal, diffuse_brdf, &rng); let combined_reservoir = merge_result.merged_reservoir; - reservoirs_a[pixel_index] = combined_reservoir; + di_reservoirs_a[pixel_index] = combined_reservoir; var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight * combined_reservoir.visibility; pixel_color *= view.exposure; @@ -136,7 +136,7 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 } let temporal_pixel_index = temporal_pixel_id.x + temporal_pixel_id.y * u32(view.viewport.z); - var temporal_reservoir = reservoirs_a[temporal_pixel_index]; + var temporal_reservoir = di_reservoirs_a[temporal_pixel_index]; temporal_reservoir.sample.light_id.x = previous_frame_light_id_translations[temporal_reservoir.sample.light_id.x]; if temporal_reservoir.sample.light_id.x == LIGHT_NOT_PRESENT_THIS_FRAME { @@ -160,7 +160,7 @@ 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.viewport.z); - var spatial_reservoir = reservoirs_b[spatial_pixel_index]; + var spatial_reservoir = di_reservoirs_b[spatial_pixel_index]; if reservoir_valid(spatial_reservoir) { spatial_reservoir.visibility = trace_light_visibility(spatial_reservoir.sample, world_position); @@ -209,7 +209,7 @@ fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { #endif } -// Don't adjust the size of this struct without also adjusting RESERVOIR_STRUCT_SIZE. +// Don't adjust the size of this struct without also adjusting DI_RESERVOIR_STRUCT_SIZE. struct Reservoir { sample: LightSample, weight_sum: f32, diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl new file mode 100644 index 0000000000000..7ef919d845449 --- /dev/null +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -0,0 +1,63 @@ +// https://intro-to-restir.cwyman.org/presentations/2023ReSTIR_Course_Notes.pdf + +#import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance +#import bevy_pbr::pbr_deferred_types::unpack_24bit_normal +#import bevy_pbr::prepass_bindings::PreviousViewUniforms +#import bevy_pbr::rgb9e5::rgb9e5_to_vec3_ +#import bevy_pbr::utils::{rand_f, octahedral_decode} +#import bevy_render::maths::{PI, PI_2} +#import bevy_render::view::View +#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light} +#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} + +@group(1) @binding(0) var view_output: texture_storage_2d; +// @group(1) @binding(1) var di_reservoirs_a: array; +// @group(1) @binding(2) var di_reservoirs_b: array; +@group(1) @binding(3) var gbuffer: texture_2d; +@group(1) @binding(4) var depth_buffer: texture_depth_2d; +@group(1) @binding(5) var motion_vectors: texture_2d; +@group(1) @binding(6) var previous_gbuffer: texture_2d; +@group(1) @binding(7) var previous_depth_buffer: texture_depth_2d; +@group(1) @binding(8) var view: View; +@group(1) @binding(9) var previous_view: PreviousViewUniforms; +struct PushConstants { frame_index: u32, reset: u32 } +var constants: PushConstants; + +@compute @workgroup_size(8, 8, 1) +fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { + if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } + + let pixel_index = global_id.x + global_id.y * u32(view.viewport.z); + var rng = pixel_index + constants.frame_index; + + let depth = textureLoad(depth_buffer, global_id.xy, 0); + if depth == 0.0 { return; } + let gpixel = textureLoad(gbuffer, global_id.xy, 0); + let world_position = reconstruct_world_position(global_id.xy, depth); + let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); + let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); + let diffuse_brdf = base_color / PI; + + let ray_direction = sample_uniform_hemisphere(world_normal, &rng); + let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); + if ray_hit.kind == RAY_QUERY_INTERSECTION_NONE { return; } + let sample_point = resolve_ray_hit_full(ray_hit); + if all(sample_point.material.emissive != vec3(0.0)) { return; } + let sample_point_diffuse_brdf = sample_point.material.base_color / PI; + let radiance = sample_random_light(sample_point.world_position, sample_point.world_normal, &rng); + + let cos_theta = dot(ray_direction, world_normal); + let inverse_uniform_hemisphere_pdf = PI_2; + let contribution = (radiance * sample_point_diffuse_brdf * diffuse_brdf * cos_theta * inverse_uniform_hemisphere_pdf); + + var pixel_color = textureLoad(view_output, global_id.xy); + pixel_color += vec4(contribution * view.exposure, 0.0); + textureStore(view_output, global_id.xy, pixel_color); +} + +fn reconstruct_world_position(pixel_id: vec2, depth: f32) -> vec3 { + let uv = (vec2(pixel_id) + 0.5) / view.viewport.zw; + let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); + let world_pos = view.world_from_clip * vec4(xy_ndc, depth, 1.0); + return world_pos.xyz / world_pos.w; +} diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index be709f0bc8dd1..c6ad92af49054 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -15,6 +15,17 @@ fn sample_cosine_hemisphere(normal: vec3, rng: ptr) -> vec3< return vec3(x, y, z); } +// https://www.pbr-book.org/3ed-2018/Monte_Carlo_Integration/2D_Sampling_with_Multidimensional_Transformations#UniformlySamplingaHemisphere +fn sample_uniform_hemisphere(normal: vec3, rng: ptr) -> vec3 { + let cos_theta = rand_f(rng); + let phi = PI_2 * rand_f(rng); + let sin_theta = sqrt(max(1.0 - cos_theta * cos_theta, 0.0)); + let x = sin_theta * cos(phi); + let y = sin_theta * sin(phi); + let z = cos_theta; + return build_orthonormal_basis(normal) * vec3(x, y, z); +} + // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec19%3A294 fn sample_disk(disk_radius: f32, rng: ptr) -> vec2 { let ab = 2.0 * rand_vec2f(rng) - 1.0; From 749681f552918e46378345ec142d5cd1731b575b Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 5 Jul 2025 15:12:24 -0400 Subject: [PATCH 02/12] ReSTIR GI (temporal reuse only, no spatial reuse yet) --- .../src/pathtracer/pathtracer.wgsl | 3 +- crates/bevy_solari/src/realtime/node.rs | 32 ++++ crates/bevy_solari/src/realtime/prepare.rs | 29 ++- .../bevy_solari/src/realtime/restir_di.wgsl | 14 +- .../bevy_solari/src/realtime/restir_gi.wgsl | 177 ++++++++++++++++-- crates/bevy_solari/src/scene/sampling.wgsl | 9 +- 6 files changed, 234 insertions(+), 30 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index c67b53e58e727..be92b67bfae9f 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -47,7 +47,8 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { if ray_t_min == 0.0 { radiance = ray_hit.material.emissive; } // Sample direct lighting - radiance += throughput * diffuse_brdf * sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); + let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); + radiance += throughput * diffuse_brdf * direct_lighting.radiance * direct_lighting.inverse_pdf; // Sample new ray direction from the material BRDF for next bounce ray_direction = sample_cosine_hemisphere(ray_hit.world_normal, &rng); diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 2561a897e65aa..056993d3c41bc 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -38,6 +38,7 @@ pub struct SolariLightingNode { di_initial_and_temporal_pipeline: CachedComputePipelineId, di_spatial_and_shade_pipeline: CachedComputePipelineId, gi_initial_and_temporal_pipeline: CachedComputePipelineId, + gi_spatial_and_shade_pipeline: CachedComputePipelineId, } impl ViewNode for SolariLightingNode { @@ -75,6 +76,7 @@ impl ViewNode for SolariLightingNode { Some(di_initial_and_temporal_pipeline), Some(di_spatial_and_shade_pipeline), Some(gi_initial_and_temporal_pipeline), + Some(gi_spatial_and_shade_pipeline), Some(scene_bindings), Some(viewport), Some(gbuffer), @@ -86,6 +88,7 @@ impl ViewNode for SolariLightingNode { pipeline_cache.get_compute_pipeline(self.di_initial_and_temporal_pipeline), pipeline_cache.get_compute_pipeline(self.di_spatial_and_shade_pipeline), pipeline_cache.get_compute_pipeline(self.gi_initial_and_temporal_pipeline), + pipeline_cache.get_compute_pipeline(self.gi_spatial_and_shade_pipeline), &scene_bindings.bind_group, camera.physical_viewport_size, view_prepass_textures.deferred_view(), @@ -109,6 +112,12 @@ impl ViewNode for SolariLightingNode { solari_lighting_resources .di_reservoirs_b .as_entire_binding(), + solari_lighting_resources + .gi_reservoirs_a + .as_entire_binding(), + solari_lighting_resources + .gi_reservoirs_b + .as_entire_binding(), gbuffer, depth_buffer, motion_vectors, @@ -154,6 +163,9 @@ impl ViewNode for SolariLightingNode { pass.set_pipeline(gi_initial_and_temporal_pipeline); pass.dispatch_workgroups(viewport.x.div_ceil(8), viewport.y.div_ceil(8), 1); + pass.set_pipeline(gi_spatial_and_shade_pipeline); + pass.dispatch_workgroups(viewport.x.div_ceil(8), viewport.y.div_ceil(8), 1); + pass_span.end(&mut pass); drop(pass); @@ -202,6 +214,8 @@ impl FromWorld for SolariLightingNode { ), storage_buffer_sized(false, None), storage_buffer_sized(false, None), + storage_buffer_sized(false, None), + storage_buffer_sized(false, None), texture_2d(TextureSampleType::Uint), texture_depth_2d(), texture_2d(TextureSampleType::Float { filterable: true }), @@ -264,11 +278,29 @@ impl FromWorld for SolariLightingNode { zero_initialize_workgroup_memory: false, }); + let gi_spatial_and_shade_pipeline = + pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { + label: Some("solari_lighting_gi_spatial_and_shade_pipeline".into()), + layout: vec![ + scene_bindings.bind_group_layout.clone(), + bind_group_layout.clone(), + ], + push_constant_ranges: vec![PushConstantRange { + stages: ShaderStages::COMPUTE, + range: 0..8, + }], + shader: load_embedded_asset!(world, "restir_gi.wgsl"), + shader_defs: vec![], + entry_point: "spatial_and_shade".into(), + zero_initialize_workgroup_memory: false, + }); + Self { bind_group_layout, di_initial_and_temporal_pipeline, di_spatial_and_shade_pipeline, gi_initial_and_temporal_pipeline, + gi_spatial_and_shade_pipeline, } } } diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 3c454727bd8cc..46a94a3ca2477 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -17,14 +17,19 @@ use bevy_render::{ renderer::RenderDevice, }; -/// Size of a Reservoir shader struct in bytes. +/// Size of a DI Reservoir shader struct in bytes. const DI_RESERVOIR_STRUCT_SIZE: u64 = 32; +/// Size of a GI Reservoir shader struct in bytes. +const GI_RESERVOIR_STRUCT_SIZE: u64 = 48; + /// Internal rendering resources used for Solari lighting. #[derive(Component)] pub struct SolariLightingResources { pub di_reservoirs_a: Buffer, pub di_reservoirs_b: Buffer, + pub gi_reservoirs_a: Buffer, + pub gi_reservoirs_b: Buffer, pub previous_gbuffer: (Texture, TextureView), pub previous_depth: (Texture, TextureView), pub view_size: UVec2, @@ -47,18 +52,30 @@ pub fn prepare_solari_lighting_resources( continue; } - let size = (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE; - let di_reservoirs_a = render_device.create_buffer(&BufferDescriptor { label: Some("solari_lighting_di_reservoirs_a"), - size, + size: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE, usage: BufferUsages::STORAGE, mapped_at_creation: false, }); let di_reservoirs_b = render_device.create_buffer(&BufferDescriptor { label: Some("solari_lighting_di_reservoirs_b"), - size, + size: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE, + usage: BufferUsages::STORAGE, + mapped_at_creation: false, + }); + + let gi_reservoirs_a = render_device.create_buffer(&BufferDescriptor { + label: Some("solari_lighting_gi_reservoirs_a"), + size: (view_size.x * view_size.y) as u64 * GI_RESERVOIR_STRUCT_SIZE, + usage: BufferUsages::STORAGE, + mapped_at_creation: false, + }); + + let gi_reservoirs_b = render_device.create_buffer(&BufferDescriptor { + label: Some("solari_lighting_gi_reservoirs_b"), + size: (view_size.x * view_size.y) as u64 * GI_RESERVOIR_STRUCT_SIZE, usage: BufferUsages::STORAGE, mapped_at_creation: false, }); @@ -90,6 +107,8 @@ pub fn prepare_solari_lighting_resources( commands.entity(entity).insert(SolariLightingResources { di_reservoirs_a, di_reservoirs_b, + gi_reservoirs_a, + gi_reservoirs_b, previous_gbuffer: (previous_gbuffer, previous_gbuffer_view), previous_depth: (previous_depth, previous_depth_view), view_size, diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index c9efc6fb205d5..6dae89d21b999 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -13,13 +13,13 @@ @group(1) @binding(0) var view_output: texture_storage_2d; @group(1) @binding(1) var di_reservoirs_a: array; @group(1) @binding(2) var di_reservoirs_b: array; -@group(1) @binding(3) var gbuffer: texture_2d; -@group(1) @binding(4) var depth_buffer: texture_depth_2d; -@group(1) @binding(5) var motion_vectors: texture_2d; -@group(1) @binding(6) var previous_gbuffer: texture_2d; -@group(1) @binding(7) var previous_depth_buffer: texture_depth_2d; -@group(1) @binding(8) var view: View; -@group(1) @binding(9) var previous_view: PreviousViewUniforms; +@group(1) @binding(5) var gbuffer: texture_2d; +@group(1) @binding(6) var depth_buffer: texture_depth_2d; +@group(1) @binding(7) var motion_vectors: texture_2d; +@group(1) @binding(8) var previous_gbuffer: texture_2d; +@group(1) @binding(9) var previous_depth_buffer: texture_depth_2d; +@group(1) @binding(10) var view: View; +@group(1) @binding(11) var previous_view: PreviousViewUniforms; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 7ef919d845449..8e090ca29cfb6 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -11,18 +11,20 @@ #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var view_output: texture_storage_2d; -// @group(1) @binding(1) var di_reservoirs_a: array; -// @group(1) @binding(2) var di_reservoirs_b: array; -@group(1) @binding(3) var gbuffer: texture_2d; -@group(1) @binding(4) var depth_buffer: texture_depth_2d; -@group(1) @binding(5) var motion_vectors: texture_2d; -@group(1) @binding(6) var previous_gbuffer: texture_2d; -@group(1) @binding(7) var previous_depth_buffer: texture_depth_2d; -@group(1) @binding(8) var view: View; -@group(1) @binding(9) var previous_view: PreviousViewUniforms; +@group(1) @binding(3) var gi_reservoirs_a: array; +@group(1) @binding(4) var gi_reservoirs_b: array; +@group(1) @binding(5) var gbuffer: texture_2d; +@group(1) @binding(6) var depth_buffer: texture_depth_2d; +@group(1) @binding(7) var motion_vectors: texture_2d; +@group(1) @binding(8) var previous_gbuffer: texture_2d; +@group(1) @binding(9) var previous_depth_buffer: texture_depth_2d; +@group(1) @binding(10) var view: View; +@group(1) @binding(11) var previous_view: PreviousViewUniforms; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; +const CONFIDENCE_WEIGHT_CAP = 30.0; + @compute @workgroup_size(8, 8, 1) fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } @@ -31,33 +33,178 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { var rng = pixel_index + constants.frame_index; let depth = textureLoad(depth_buffer, global_id.xy, 0); - if depth == 0.0 { return; } + if depth == 0.0 { + gi_reservoirs_b[pixel_index] = empty_reservoir(); + return; + } let gpixel = textureLoad(gbuffer, global_id.xy, 0); let world_position = reconstruct_world_position(global_id.xy, depth); let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); let diffuse_brdf = base_color / PI; + let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal); + let ray_direction = sample_uniform_hemisphere(world_normal, &rng); let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); - if ray_hit.kind == RAY_QUERY_INTERSECTION_NONE { return; } + if ray_hit.kind == RAY_QUERY_INTERSECTION_NONE { + gi_reservoirs_b[pixel_index] = temporal_reservoir; + return; + } let sample_point = resolve_ray_hit_full(ray_hit); - if all(sample_point.material.emissive != vec3(0.0)) { return; } + if all(sample_point.material.emissive != vec3(0.0)) { + gi_reservoirs_b[pixel_index] = temporal_reservoir; + return; + } let sample_point_diffuse_brdf = sample_point.material.base_color / PI; - let radiance = sample_random_light(sample_point.world_position, sample_point.world_normal, &rng); + let direct_lighting = sample_random_light(sample_point.world_position, sample_point.world_normal, &rng); + let sample_point_radiance = direct_lighting.radiance * sample_point_diffuse_brdf; let cos_theta = dot(ray_direction, world_normal); let inverse_uniform_hemisphere_pdf = PI_2; - let contribution = (radiance * sample_point_diffuse_brdf * diffuse_brdf * cos_theta * inverse_uniform_hemisphere_pdf); + + var combined_reservoir = empty_reservoir(); + combined_reservoir.confidence_weight = 1.0 + temporal_reservoir.confidence_weight; + + let mis_weight_denominator = 1.0 / combined_reservoir.confidence_weight; + + let new_mis_weight = mis_weight_denominator; + let new_target_function = luminance(sample_point_radiance * diffuse_brdf * cos_theta); + let new_inverse_pdf = direct_lighting.inverse_pdf * inverse_uniform_hemisphere_pdf; + let new_resampling_weight = new_mis_weight * (new_target_function * new_inverse_pdf); + + let temporal_mis_weight = temporal_reservoir.confidence_weight * mis_weight_denominator; + let temporal_cos_theta = dot(normalize(temporal_reservoir.sample_point_world_position - world_position), world_normal); + let temporal_target_function = luminance(temporal_reservoir.radiance * diffuse_brdf * temporal_cos_theta); + let temporal_resampling_weight = temporal_mis_weight * (temporal_target_function * temporal_reservoir.unbiased_contribution_weight); + + combined_reservoir.weight_sum = new_resampling_weight + temporal_resampling_weight; + + if rand_f(&rng) < temporal_resampling_weight / combined_reservoir.weight_sum { + combined_reservoir.sample_point_world_position = temporal_reservoir.sample_point_world_position; + combined_reservoir.radiance = temporal_reservoir.radiance; + + let inverse_target_function = select(0.0, 1.0 / temporal_target_function, temporal_target_function > 0.0); + combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; + } else { + combined_reservoir.sample_point_world_position = sample_point.world_position; + combined_reservoir.radiance = sample_point_radiance; + + let inverse_target_function = select(0.0, 1.0 / new_target_function, new_target_function > 0.0); + combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; + } + + gi_reservoirs_b[pixel_index] = combined_reservoir; +} + + +@compute @workgroup_size(8, 8, 1) +fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { + if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } + + let pixel_index = global_id.x + global_id.y * u32(view.viewport.z); + var rng = pixel_index + constants.frame_index; + + let depth = textureLoad(depth_buffer, global_id.xy, 0); + if depth == 0.0 { + gi_reservoirs_a[pixel_index] = empty_reservoir(); + return; + } + let gpixel = textureLoad(gbuffer, global_id.xy, 0); + let world_position = reconstruct_world_position(global_id.xy, depth); + let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); + let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); + let diffuse_brdf = base_color / PI; + + let input_reservoir = gi_reservoirs_b[pixel_index]; + let cos_theta = dot(normalize(input_reservoir.sample_point_world_position - world_position), world_normal); + let radiance = input_reservoir.radiance * diffuse_brdf * cos_theta; + + gi_reservoirs_a[pixel_index] = input_reservoir; var pixel_color = textureLoad(view_output, global_id.xy); - pixel_color += vec4(contribution * view.exposure, 0.0); + pixel_color += vec4(radiance * input_reservoir.unbiased_contribution_weight * view.exposure, 0.0); textureStore(view_output, global_id.xy, pixel_color); } +fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> Reservoir { + let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; + let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.viewport.zw)); + let temporal_pixel_id = vec2(temporal_pixel_id_float); + if any(temporal_pixel_id_float < vec2(0.0)) || any(temporal_pixel_id_float >= view.viewport.zw) || bool(constants.reset) { + return empty_reservoir(); + } + + let temporal_depth = textureLoad(previous_depth_buffer, temporal_pixel_id, 0); + let temporal_gpixel = textureLoad(previous_gbuffer, temporal_pixel_id, 0); + let temporal_world_position = reconstruct_previous_world_position(temporal_pixel_id, temporal_depth); + let temporal_world_normal = octahedral_decode(unpack_24bit_normal(temporal_gpixel.a)); + if pixel_dissimilar(depth, world_position, temporal_world_position, world_normal, temporal_world_normal) { + return empty_reservoir(); + } + + let temporal_pixel_index = temporal_pixel_id.x + temporal_pixel_id.y * u32(view.viewport.z); + var temporal_reservoir = gi_reservoirs_a[temporal_pixel_index]; + + temporal_reservoir.confidence_weight = min(temporal_reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); + + return temporal_reservoir; +} + fn reconstruct_world_position(pixel_id: vec2, depth: f32) -> vec3 { let uv = (vec2(pixel_id) + 0.5) / view.viewport.zw; let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); let world_pos = view.world_from_clip * vec4(xy_ndc, depth, 1.0); return world_pos.xyz / world_pos.w; } + +fn reconstruct_previous_world_position(pixel_id: vec2, depth: f32) -> vec3 { + let uv = (vec2(pixel_id) + 0.5) / view.viewport.zw; + let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); + let world_pos = previous_view.world_from_clip * vec4(xy_ndc, depth, 1.0); + return world_pos.xyz / world_pos.w; +} + +// Reject if tangent plane difference difference more than 0.3% or angle between normals more than 25 degrees +fn pixel_dissimilar(depth: f32, world_position: vec3, other_world_position: vec3, normal: vec3, other_normal: vec3) -> bool { + // https://developer.download.nvidia.com/video/gputechconf/gtc/2020/presentations/s22699-fast-denoising-with-self-stabilizing-recurrent-blurs.pdf#page=45 + let tangent_plane_distance = abs(dot(normal, other_world_position - world_position)); + let view_z = -depth_ndc_to_view_z(depth); + + return tangent_plane_distance / view_z > 0.003 || dot(normal, other_normal) < 0.906; +} + +fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { +#ifdef VIEW_PROJECTION_PERSPECTIVE + return -view.clip_from_view[3][2]() / ndc_depth; +#else ifdef VIEW_PROJECTION_ORTHOGRAPHIC + return -(view.clip_from_view[3][2] - ndc_depth) / view.clip_from_view[2][2]; +#else + let view_pos = view.view_from_clip * vec4(0.0, 0.0, ndc_depth, 1.0); + return view_pos.z / view_pos.w; +#endif +} + +struct Reservoir { + sample_point_world_position: vec3, + weight_sum: f32, + radiance: vec3, + confidence_weight: f32, + unbiased_contribution_weight: f32, + padding1: f32, + padding2: f32, + padding3: f32, +} + +fn empty_reservoir() -> Reservoir { + return Reservoir( + vec3(0.0), + 0.0, + vec3(0.0), + 0.0, + 0.0, + 0.0, + 0.0, + 0.0, + ); +} diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index c6ad92af49054..a707fbb1b3054 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -48,11 +48,16 @@ fn sample_disk(disk_radius: f32, rng: ptr) -> vec2 { return vec2(x, y); } -fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rng: ptr) -> vec3 { +struct SampleRandomLightResult { + radiance: vec3, + inverse_pdf: f32, +} + +fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rng: ptr) -> SampleRandomLightResult { let light_sample = generate_random_light_sample(rng); let light_contribution = calculate_light_contribution(light_sample, ray_origin, origin_world_normal); let visibility = trace_light_visibility(light_sample, ray_origin); - return light_contribution.radiance * visibility * light_contribution.inverse_pdf; + return SampleRandomLightResult(light_contribution.radiance * visibility, light_contribution.inverse_pdf); } struct LightSample { From 5769fed9de7445feb5e420754f2f276b1a57eedb Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 5 Jul 2025 15:28:22 -0400 Subject: [PATCH 03/12] Remove outdated comment --- crates/bevy_solari/src/lib.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/lib.rs b/crates/bevy_solari/src/lib.rs index d5a22e014b8c5..65765d31a4b65 100644 --- a/crates/bevy_solari/src/lib.rs +++ b/crates/bevy_solari/src/lib.rs @@ -26,7 +26,7 @@ use bevy_render::settings::WgpuFeatures; /// An experimental plugin for raytraced lighting. /// /// This plugin provides: -/// * [`SolariLightingPlugin`] - Raytraced direct and indirect lighting (indirect lighting not yet implemented). +/// * [`SolariLightingPlugin`] - Raytraced direct and indirect lighting. /// * [`RaytracingScenePlugin`] - BLAS building, resource and lighting binding. /// * [`pathtracer::PathtracingPlugin`] - A non-realtime pathtracer for validation purposes. /// From c12670b0be1fa8f1fc6693d33728bf6ab3677e60 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 5 Jul 2025 18:50:37 -0400 Subject: [PATCH 04/12] Experimental bilinear reprojection --- .../bevy_solari/src/realtime/restir_di.wgsl | 42 +++++++++++++++++-- .../bevy_solari/src/realtime/restir_gi.wgsl | 41 ++++++++++++++++-- 2 files changed, 75 insertions(+), 8 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 6dae89d21b999..7bc3d32fcfb8c 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -48,7 +48,7 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = base_color / PI; let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng); - let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal); + let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let combined_reservoir = merge_reservoirs(initial_reservoir, temporal_reservoir, world_position, world_normal, diffuse_brdf, &rng); di_reservoirs_b[pixel_index] = combined_reservoir.merged_reservoir; @@ -119,14 +119,47 @@ 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, rng: ptr) -> Reservoir { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; - let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.viewport.zw)); - let temporal_pixel_id = vec2(temporal_pixel_id_float); + let temporal_pixel_id_float = vec2(pixel_id) - (motion_vector * view.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.viewport.zw) || bool(constants.reset) { return empty_reservoir(); } + // https://en.wikipedia.org/wiki/Bilinear_interpolation#On_the_unit_square + let tl = vec2(temporal_pixel_id_float); + let tr = tl + vec2(1u, 0u); + let bl = tl + vec2(0u, 1u); + let br = tl + vec2(1u, 1u); + let f = fract(temporal_pixel_id_float); + let tl_w = (1.0 - f.x) * (1.0 - f.y); + let tr_w = f.x * (1.0 - f.y); + let bl_w = (1.0 - f.x) * f.y; + + // Choose a random pixel from the 2x2 quad, weighted by the bilinear weights + // This gives better results than always using the nearest pixel + var temporal_pixel_id = tl; + var weight_sum = tl_w; + let r = rand_f(rng); + if (r > weight_sum) { + temporal_pixel_id = tr; + weight_sum += tr_w; + } + if (r > weight_sum) { + temporal_pixel_id = bl; + weight_sum += bl_w; + } + if (r > weight_sum) { + temporal_pixel_id = br; + } + + // Clamp to view size, since 2x2 quad may go off screen + temporal_pixel_id = min(temporal_pixel_id, vec2(view.viewport.zw - 1.0)); + + // 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_gpixel = textureLoad(previous_gbuffer, temporal_pixel_id, 0); let temporal_world_position = reconstruct_previous_world_position(temporal_pixel_id, temporal_depth); @@ -138,6 +171,7 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 let temporal_pixel_index = temporal_pixel_id.x + temporal_pixel_id.y * u32(view.viewport.z); var temporal_reservoir = di_reservoirs_a[temporal_pixel_index]; + // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) temporal_reservoir.sample.light_id.x = previous_frame_light_id_translations[temporal_reservoir.sample.light_id.x]; if temporal_reservoir.sample.light_id.x == LIGHT_NOT_PRESENT_THIS_FRAME { return empty_reservoir(); diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 8e090ca29cfb6..aae0db19420cb 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -43,7 +43,7 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); let diffuse_brdf = base_color / PI; - let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal); + let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let ray_direction = sample_uniform_hemisphere(world_normal, &rng); let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); @@ -127,14 +127,47 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { textureStore(view_output, global_id.xy, pixel_color); } -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, rng: ptr) -> Reservoir { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; - let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.viewport.zw)); - let temporal_pixel_id = vec2(temporal_pixel_id_float); + let temporal_pixel_id_float = vec2(pixel_id) - (motion_vector * view.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.viewport.zw) || bool(constants.reset) { return empty_reservoir(); } + // https://en.wikipedia.org/wiki/Bilinear_interpolation#On_the_unit_square + let tl = vec2(temporal_pixel_id_float); + let tr = tl + vec2(1u, 0u); + let bl = tl + vec2(0u, 1u); + let br = tl + vec2(1u, 1u); + let f = fract(temporal_pixel_id_float); + let tl_w = (1.0 - f.x) * (1.0 - f.y); + let tr_w = f.x * (1.0 - f.y); + let bl_w = (1.0 - f.x) * f.y; + + // Choose a random pixel from the 2x2 quad, weighted by the bilinear weights + // This gives better results than always using the nearest pixel + var temporal_pixel_id = tl; + var weight_sum = tl_w; + let r = rand_f(rng); + if (r > weight_sum) { + temporal_pixel_id = tr; + weight_sum += tr_w; + } + if (r > weight_sum) { + temporal_pixel_id = bl; + weight_sum += bl_w; + } + if (r > weight_sum) { + temporal_pixel_id = br; + } + + // Clamp to view size, since 2x2 quad may go off screen + temporal_pixel_id = min(temporal_pixel_id, vec2(view.viewport.zw - 1.0)); + + // 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_gpixel = textureLoad(previous_gbuffer, temporal_pixel_id, 0); let temporal_world_position = reconstruct_previous_world_position(temporal_pixel_id, temporal_depth); From cb2a1fdddb11cf69f90c0da55262324efd08eeee Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 6 Jul 2025 11:20:31 -0400 Subject: [PATCH 05/12] Cleanup GI --- .../bevy_solari/src/realtime/restir_gi.wgsl | 123 ++++++++++-------- 1 file changed, 70 insertions(+), 53 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index aae0db19420cb..189121cb9a229 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -43,61 +43,13 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); let diffuse_brdf = base_color / PI; - let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal, &rng); - - let ray_direction = sample_uniform_hemisphere(world_normal, &rng); - let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); - if ray_hit.kind == RAY_QUERY_INTERSECTION_NONE { - gi_reservoirs_b[pixel_index] = temporal_reservoir; - return; - } - let sample_point = resolve_ray_hit_full(ray_hit); - if all(sample_point.material.emissive != vec3(0.0)) { - gi_reservoirs_b[pixel_index] = temporal_reservoir; - return; - } - let sample_point_diffuse_brdf = sample_point.material.base_color / PI; - let direct_lighting = sample_random_light(sample_point.world_position, sample_point.world_normal, &rng); - let sample_point_radiance = direct_lighting.radiance * sample_point_diffuse_brdf; - - let cos_theta = dot(ray_direction, world_normal); - let inverse_uniform_hemisphere_pdf = PI_2; - - var combined_reservoir = empty_reservoir(); - combined_reservoir.confidence_weight = 1.0 + temporal_reservoir.confidence_weight; - - let mis_weight_denominator = 1.0 / combined_reservoir.confidence_weight; - - let new_mis_weight = mis_weight_denominator; - let new_target_function = luminance(sample_point_radiance * diffuse_brdf * cos_theta); - let new_inverse_pdf = direct_lighting.inverse_pdf * inverse_uniform_hemisphere_pdf; - let new_resampling_weight = new_mis_weight * (new_target_function * new_inverse_pdf); - - let temporal_mis_weight = temporal_reservoir.confidence_weight * mis_weight_denominator; - let temporal_cos_theta = dot(normalize(temporal_reservoir.sample_point_world_position - world_position), world_normal); - let temporal_target_function = luminance(temporal_reservoir.radiance * diffuse_brdf * temporal_cos_theta); - let temporal_resampling_weight = temporal_mis_weight * (temporal_target_function * temporal_reservoir.unbiased_contribution_weight); - - combined_reservoir.weight_sum = new_resampling_weight + temporal_resampling_weight; - - if rand_f(&rng) < temporal_resampling_weight / combined_reservoir.weight_sum { - combined_reservoir.sample_point_world_position = temporal_reservoir.sample_point_world_position; - combined_reservoir.radiance = temporal_reservoir.radiance; - - let inverse_target_function = select(0.0, 1.0 / temporal_target_function, temporal_target_function > 0.0); - combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; - } else { - combined_reservoir.sample_point_world_position = sample_point.world_position; - combined_reservoir.radiance = sample_point_radiance; - - let inverse_target_function = select(0.0, 1.0 / new_target_function, new_target_function > 0.0); - combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; - } + let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng); + let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal, diffuse_brdf, &rng); + let combined_reservoir = merge_reservoirs(initial_reservoir, temporal_reservoir, &rng); gi_reservoirs_b[pixel_index] = combined_reservoir; } - @compute @workgroup_size(8, 8, 1) fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } @@ -127,7 +79,39 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { textureStore(view_output, global_id.xy, pixel_color); } -fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { +fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr) -> Reservoir{ + var reservoir = empty_reservoir(); + + let ray_direction = sample_uniform_hemisphere(world_normal, rng); + let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); + + if ray_hit.kind == RAY_QUERY_INTERSECTION_NONE { + return reservoir; + } + + let sample_point = resolve_ray_hit_full(ray_hit); + + if all(sample_point.material.emissive != vec3(0.0)) { + return reservoir; + } + + reservoir.sample_point_world_position = sample_point.world_position; + reservoir.confidence_weight = 1.0; + + let sample_point_diffuse_brdf = sample_point.material.base_color / PI; + let direct_lighting = sample_random_light(sample_point.world_position, sample_point.world_normal, rng); + reservoir.radiance = direct_lighting.radiance * sample_point_diffuse_brdf; + + let inverse_uniform_hemisphere_pdf = PI_2; + reservoir.unbiased_contribution_weight = direct_lighting.inverse_pdf * inverse_uniform_hemisphere_pdf; + + let cos_theta = dot(ray_direction, world_normal); + reservoir.target_function = luminance(reservoir.radiance * diffuse_brdf * cos_theta); + + return reservoir; +} + +fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr) -> Reservoir { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; let temporal_pixel_id_float = vec2(pixel_id) - (motion_vector * view.viewport.zw); @@ -181,6 +165,9 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 temporal_reservoir.confidence_weight = min(temporal_reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); + let temporal_cos_theta = dot(normalize(temporal_reservoir.sample_point_world_position - world_position), world_normal); + temporal_reservoir.target_function = luminance(temporal_reservoir.radiance * diffuse_brdf * temporal_cos_theta); + return temporal_reservoir; } @@ -224,7 +211,7 @@ struct Reservoir { radiance: vec3, confidence_weight: f32, unbiased_contribution_weight: f32, - padding1: f32, + target_function: f32, padding2: f32, padding3: f32, } @@ -241,3 +228,33 @@ fn empty_reservoir() -> Reservoir { 0.0, ); } + +fn merge_reservoirs(canonical_reservoir: Reservoir,other_reservoir: Reservoir,rng: ptr) -> Reservoir { + // TODO: Balance heuristic MIS weights + 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_reservoir.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_reservoir.target_function * other_reservoir.unbiased_contribution_weight); + + var combined_reservoir = empty_reservoir(); + combined_reservoir.weight_sum = canonical_resampling_weight + other_resampling_weight; + combined_reservoir.confidence_weight = canonical_reservoir.confidence_weight + other_reservoir.confidence_weight; + + if rand_f(rng) < other_resampling_weight / combined_reservoir.weight_sum { + combined_reservoir.sample_point_world_position = other_reservoir.sample_point_world_position; + combined_reservoir.radiance = other_reservoir.radiance; + combined_reservoir.target_function = other_reservoir.target_function; + } else { + combined_reservoir.sample_point_world_position = canonical_reservoir.sample_point_world_position; + combined_reservoir.radiance = canonical_reservoir.radiance; + combined_reservoir.target_function = canonical_reservoir.target_function; + } + + let inverse_target_function = select(0.0, 1.0 / combined_reservoir.target_function, combined_reservoir.target_function > 0.0); + combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; + + return combined_reservoir; +} From 438cfef3c01a6d0d6ff90b93216fd0aebc10ccc3 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 6 Jul 2025 12:01:26 -0400 Subject: [PATCH 06/12] Initial spatial resampling (not working) --- .../bevy_solari/src/realtime/restir_gi.wgsl | 79 +++++++++++++------ 1 file changed, 53 insertions(+), 26 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 189121cb9a229..5df11090dbad5 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_pbr::utils::{rand_f, octahedral_decode} #import bevy_render::maths::{PI, PI_2} #import bevy_render::view::View -#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light} +#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light, sample_disk} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var view_output: texture_storage_2d; @@ -23,6 +23,7 @@ struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; +const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 30.0; @compute @workgroup_size(8, 8, 1) @@ -40,11 +41,9 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let gpixel = textureLoad(gbuffer, global_id.xy, 0); let world_position = reconstruct_world_position(global_id.xy, depth); let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); - let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); - let diffuse_brdf = base_color / PI; - let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng); - let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal, diffuse_brdf, &rng); + let initial_reservoir = generate_initial_reservoir(world_position, world_normal, &rng); + let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let combined_reservoir = merge_reservoirs(initial_reservoir, temporal_reservoir, &rng); gi_reservoirs_b[pixel_index] = combined_reservoir; @@ -69,17 +68,20 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = base_color / PI; let input_reservoir = gi_reservoirs_b[pixel_index]; - let cos_theta = dot(normalize(input_reservoir.sample_point_world_position - world_position), world_normal); - let radiance = input_reservoir.radiance * diffuse_brdf * cos_theta; + let spatial_reservoir = load_spatial_reservoir(global_id.xy, depth, world_position, world_normal, &rng); + let combined_reservoir = merge_reservoirs(input_reservoir, spatial_reservoir, &rng); + + gi_reservoirs_a[pixel_index] = combined_reservoir; - gi_reservoirs_a[pixel_index] = input_reservoir; + let cos_theta = dot(normalize(combined_reservoir.sample_point_world_position - world_position), world_normal); + let radiance = combined_reservoir.radiance * diffuse_brdf * cos_theta; var pixel_color = textureLoad(view_output, global_id.xy); - pixel_color += vec4(radiance * input_reservoir.unbiased_contribution_weight * view.exposure, 0.0); + pixel_color += vec4(radiance * combined_reservoir.unbiased_contribution_weight * view.exposure, 0.0); textureStore(view_output, global_id.xy, pixel_color); } -fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr) -> Reservoir{ +fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir{ var reservoir = empty_reservoir(); let ray_direction = sample_uniform_hemisphere(world_normal, rng); @@ -105,13 +107,10 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 let inverse_uniform_hemisphere_pdf = PI_2; reservoir.unbiased_contribution_weight = direct_lighting.inverse_pdf * inverse_uniform_hemisphere_pdf; - let cos_theta = dot(ray_direction, world_normal); - reservoir.target_function = luminance(reservoir.radiance * diffuse_brdf * cos_theta); - return reservoir; } -fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr) -> Reservoir { +fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; let temporal_pixel_id_float = vec2(pixel_id) - (motion_vector * view.viewport.zw); @@ -165,12 +164,37 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 temporal_reservoir.confidence_weight = min(temporal_reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); - let temporal_cos_theta = dot(normalize(temporal_reservoir.sample_point_world_position - world_position), world_normal); - temporal_reservoir.target_function = luminance(temporal_reservoir.radiance * diffuse_brdf * temporal_cos_theta); - return temporal_reservoir; } +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); + + let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); + let spatial_gpixel = textureLoad(gbuffer, spatial_pixel_id, 0); + let spatial_world_position = reconstruct_world_position(spatial_pixel_id, spatial_depth); + let spatial_world_normal = octahedral_decode(unpack_24bit_normal(spatial_gpixel.a)); + if pixel_dissimilar(depth, world_position, spatial_world_position, world_normal, spatial_world_normal) { + return empty_reservoir(); + } + + let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.viewport.z); + var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; + + let ray_direction = normalize(spatial_reservoir.sample_point_world_position - world_position); + let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); + spatial_reservoir.unbiased_contribution_weight *= f32(ray_hit.kind == RAY_QUERY_INTERSECTION_NONE); + + return spatial_reservoir; +} + + +fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> vec2 { + var spatial_id = vec2(center_pixel_id) + vec2(sample_disk(SPATIAL_REUSE_RADIUS_PIXELS, rng)); + spatial_id = clamp(spatial_id, vec2(0i), vec2(view.viewport.zw) - 1i); + return vec2(spatial_id); +} + fn reconstruct_world_position(pixel_id: vec2, depth: f32) -> vec3 { let uv = (vec2(pixel_id) + 0.5) / view.viewport.zw; let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); @@ -211,7 +235,7 @@ struct Reservoir { radiance: vec3, confidence_weight: f32, unbiased_contribution_weight: f32, - target_function: f32, + padding1: f32, padding2: f32, padding3: f32, } @@ -229,15 +253,17 @@ fn empty_reservoir() -> Reservoir { ); } -fn merge_reservoirs(canonical_reservoir: Reservoir,other_reservoir: Reservoir,rng: ptr) -> Reservoir { +fn merge_reservoirs(canonical_reservoir: Reservoir, other_reservoir: Reservoir, rng: ptr) -> Reservoir { // TODO: Balance heuristic MIS weights 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_reservoir.target_function * canonical_reservoir.unbiased_contribution_weight); + let canonical_target_function = luminance(canonical_reservoir.radiance); + let canonical_resampling_weight = canonical_mis_weight * (canonical_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_reservoir.target_function * other_reservoir.unbiased_contribution_weight); + let other_target_function = luminance(other_reservoir.radiance); + let other_resampling_weight = other_mis_weight * (other_target_function * other_reservoir.unbiased_contribution_weight); var combined_reservoir = empty_reservoir(); combined_reservoir.weight_sum = canonical_resampling_weight + other_resampling_weight; @@ -246,15 +272,16 @@ fn merge_reservoirs(canonical_reservoir: Reservoir,other_reservoir: Reservoir,rn if rand_f(rng) < other_resampling_weight / combined_reservoir.weight_sum { combined_reservoir.sample_point_world_position = other_reservoir.sample_point_world_position; combined_reservoir.radiance = other_reservoir.radiance; - combined_reservoir.target_function = other_reservoir.target_function; + + let inverse_target_function = select(0.0, 1.0 / other_target_function, other_target_function > 0.0); + combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; } else { combined_reservoir.sample_point_world_position = canonical_reservoir.sample_point_world_position; combined_reservoir.radiance = canonical_reservoir.radiance; - combined_reservoir.target_function = canonical_reservoir.target_function; - } - let inverse_target_function = select(0.0, 1.0 / combined_reservoir.target_function, combined_reservoir.target_function > 0.0); - combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; + let inverse_target_function = select(0.0, 1.0 / canonical_target_function, canonical_target_function > 0.0); + combined_reservoir.unbiased_contribution_weight = combined_reservoir.weight_sum * inverse_target_function; + } return combined_reservoir; } From b858feaebafcb7917dae48c6de07ffa69ba11b03 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 6 Jul 2025 12:19:07 -0400 Subject: [PATCH 07/12] Fix GI spatial resampling --- crates/bevy_solari/src/realtime/restir_gi.wgsl | 15 ++++++++------- crates/bevy_solari/src/scene/sampling.wgsl | 10 +++++++--- 2 files changed, 15 insertions(+), 10 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 5df11090dbad5..3e009e1bd78c1 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_pbr::utils::{rand_f, octahedral_decode} #import bevy_render::maths::{PI, PI_2} #import bevy_render::view::View -#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light, sample_disk} +#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light, sample_disk, trace_point_visibility} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var view_output: texture_storage_2d; @@ -181,9 +181,7 @@ 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.viewport.z); var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; - let ray_direction = normalize(spatial_reservoir.sample_point_world_position - world_position); - let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); - spatial_reservoir.unbiased_contribution_weight *= f32(ray_hit.kind == RAY_QUERY_INTERSECTION_NONE); + spatial_reservoir.unbiased_contribution_weight *= trace_point_visibility(world_position, spatial_reservoir.sample_point_world_position); return spatial_reservoir; } @@ -254,8 +252,13 @@ fn empty_reservoir() -> Reservoir { } fn merge_reservoirs(canonical_reservoir: Reservoir, other_reservoir: Reservoir, rng: ptr) -> Reservoir { + var combined_reservoir = empty_reservoir(); + combined_reservoir.confidence_weight = canonical_reservoir.confidence_weight + other_reservoir.confidence_weight; + + if combined_reservoir.confidence_weight == 0.0 { return combined_reservoir; } + // TODO: Balance heuristic MIS weights - let mis_weight_denominator = 1.0 / (canonical_reservoir.confidence_weight + other_reservoir.confidence_weight); + let mis_weight_denominator = 1.0 / combined_reservoir.confidence_weight; let canonical_mis_weight = canonical_reservoir.confidence_weight * mis_weight_denominator; let canonical_target_function = luminance(canonical_reservoir.radiance); @@ -265,9 +268,7 @@ fn merge_reservoirs(canonical_reservoir: Reservoir, other_reservoir: Reservoir, let other_target_function = luminance(other_reservoir.radiance); let other_resampling_weight = other_mis_weight * (other_target_function * other_reservoir.unbiased_contribution_weight); - var combined_reservoir = empty_reservoir(); combined_reservoir.weight_sum = canonical_resampling_weight + other_resampling_weight; - combined_reservoir.confidence_weight = canonical_reservoir.confidence_weight + other_reservoir.confidence_weight; if rand_f(rng) < other_resampling_weight / combined_reservoir.weight_sum { combined_reservoir.sample_point_world_position = other_reservoir.sample_point_world_position; diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index a707fbb1b3054..740e44aa6a4d4 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -187,10 +187,14 @@ fn trace_emissive_mesh_visibility(light_sample: LightSample, instance_id: u32, r let triangle_data = resolve_triangle_data_full(instance_id, triangle_id, barycentrics); - let light_distance = distance(ray_origin, triangle_data.world_position); - let ray_direction = (triangle_data.world_position - ray_origin) / light_distance; + return trace_point_visibility(ray_origin, triangle_data.world_position); +} + +fn trace_point_visibility(ray_origin: vec3, point: vec3) -> f32 { + let dist = distance(ray_origin, point); + let ray_direction = (point - ray_origin) / dist; - let ray_t_max = light_distance - RAY_T_MIN - RAY_T_MIN; + let ray_t_max = dist - RAY_T_MIN - RAY_T_MIN; if ray_t_max < RAY_T_MIN { return 0.0; } let ray_hit = trace_ray(ray_origin, ray_direction, RAY_T_MIN, ray_t_max, RAY_FLAG_TERMINATE_ON_FIRST_HIT); From fe55cd7588c8906cc6cccc868bf863ffe3e7e540 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 7 Jul 2025 16:44:56 -0400 Subject: [PATCH 08/12] Fix rebase --- crates/bevy_solari/src/realtime/node.rs | 4 ++-- crates/bevy_solari/src/realtime/restir_gi.wgsl | 1 + 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 3c3ce4d23bb14..eaa432d8cbace 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -273,7 +273,7 @@ impl FromWorld for SolariLightingNode { }], shader: load_embedded_asset!(world, "restir_gi.wgsl"), shader_defs: vec![], - entry_point: "initial_and_temporal".into(), + entry_point: Some("initial_and_temporal".into()), zero_initialize_workgroup_memory: false, }); @@ -290,7 +290,7 @@ impl FromWorld for SolariLightingNode { }], shader: load_embedded_asset!(world, "restir_gi.wgsl"), shader_defs: vec![], - entry_point: "spatial_and_shade".into(), + entry_point: Some("spatial_and_shade".into()), zero_initialize_workgroup_memory: false, }); diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 3e009e1bd78c1..039f2e398100b 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -227,6 +227,7 @@ fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { #endif } +// Don't adjust the size of this struct without also adjusting GI_RESERVOIR_STRUCT_SIZE. struct Reservoir { sample_point_world_position: vec3, weight_sum: f32, From 31c2477247fafa32a91b2fbbfc0c9f6dfa98d26e Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 7 Jul 2025 16:52:56 -0400 Subject: [PATCH 09/12] Revert bilinear temporal reprojection for now, didn't make a big difference --- .../bevy_solari/src/realtime/restir_di.wgsl | 33 ++----------------- .../bevy_solari/src/realtime/restir_gi.wgsl | 33 ++----------------- 2 files changed, 4 insertions(+), 62 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 7bc3d32fcfb8c..3efb655b9e424 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -121,7 +121,8 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; - let temporal_pixel_id_float = vec2(pixel_id) - (motion_vector * view.viewport.zw); + let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.viewport.zw)); + let temporal_pixel_id = vec2(temporal_pixel_id_float); // 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 @@ -129,36 +130,6 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return empty_reservoir(); } - // https://en.wikipedia.org/wiki/Bilinear_interpolation#On_the_unit_square - let tl = vec2(temporal_pixel_id_float); - let tr = tl + vec2(1u, 0u); - let bl = tl + vec2(0u, 1u); - let br = tl + vec2(1u, 1u); - let f = fract(temporal_pixel_id_float); - let tl_w = (1.0 - f.x) * (1.0 - f.y); - let tr_w = f.x * (1.0 - f.y); - let bl_w = (1.0 - f.x) * f.y; - - // Choose a random pixel from the 2x2 quad, weighted by the bilinear weights - // This gives better results than always using the nearest pixel - var temporal_pixel_id = tl; - var weight_sum = tl_w; - let r = rand_f(rng); - if (r > weight_sum) { - temporal_pixel_id = tr; - weight_sum += tr_w; - } - if (r > weight_sum) { - temporal_pixel_id = bl; - weight_sum += bl_w; - } - if (r > weight_sum) { - temporal_pixel_id = br; - } - - // Clamp to view size, since 2x2 quad may go off screen - temporal_pixel_id = min(temporal_pixel_id, vec2(view.viewport.zw - 1.0)); - // 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_gpixel = textureLoad(previous_gbuffer, temporal_pixel_id, 0); diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 039f2e398100b..ea43ee7a1a8d1 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -112,7 +112,8 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; - let temporal_pixel_id_float = vec2(pixel_id) - (motion_vector * view.viewport.zw); + let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.viewport.zw)); + let temporal_pixel_id = vec2(temporal_pixel_id_float); // 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 @@ -120,36 +121,6 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return empty_reservoir(); } - // https://en.wikipedia.org/wiki/Bilinear_interpolation#On_the_unit_square - let tl = vec2(temporal_pixel_id_float); - let tr = tl + vec2(1u, 0u); - let bl = tl + vec2(0u, 1u); - let br = tl + vec2(1u, 1u); - let f = fract(temporal_pixel_id_float); - let tl_w = (1.0 - f.x) * (1.0 - f.y); - let tr_w = f.x * (1.0 - f.y); - let bl_w = (1.0 - f.x) * f.y; - - // Choose a random pixel from the 2x2 quad, weighted by the bilinear weights - // This gives better results than always using the nearest pixel - var temporal_pixel_id = tl; - var weight_sum = tl_w; - let r = rand_f(rng); - if (r > weight_sum) { - temporal_pixel_id = tr; - weight_sum += tr_w; - } - if (r > weight_sum) { - temporal_pixel_id = bl; - weight_sum += bl_w; - } - if (r > weight_sum) { - temporal_pixel_id = br; - } - - // Clamp to view size, since 2x2 quad may go off screen - temporal_pixel_id = min(temporal_pixel_id, vec2(view.viewport.zw - 1.0)); - // 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_gpixel = textureLoad(previous_gbuffer, temporal_pixel_id, 0); From f176daea21088a4b67d97dbbec36552e21ab30c9 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 7 Jul 2025 16:55:31 -0400 Subject: [PATCH 10/12] Update release notes --- release-content/release-notes/bevy_solari.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index ac1b1abe1cbb9..86098290bbdde 100644 --- a/release-content/release-notes/bevy_solari.md +++ b/release-content/release-notes/bevy_solari.md @@ -19,13 +19,13 @@ In Bevy, direct lighting comes from analytical light components (`DirectionalLig The problem with these methods is that they all have large downsides: * Emissive meshes do not cast light onto other objects, either direct or indirect. -* Shadow maps are very expensive to render and consume a lot of memory, so you're limited to using only a few shadow casting lights. Good quality can be difficult to obtain in large scenes. +* Shadow maps are very expensive to render and consume a lot of memory, so you're limited to using only a few shadow casting lights. Good shadow quality can be difficult to obtain in large scenes. * Baked lighting does not update in realtime as objects and lights move around, is low resolution/quality, and requires time to bake, slowing down game production. * Screen-space methods have low quality and do not capture off-screen geometry and light. Bevy Solari is intended as a completely alternate, high-end lighting solution for Bevy that uses GPU-accelerated raytracing to fix all of the above problems. Emissive meshes will properly cast light and shadows, you will be able to have hundreds of shadow casting lights, quality will be much better, it will require no baking time, and it will support _fully_ dynamic scenes! -While Bevy 0.17 adds the bevy_solari crate, it's intended as a long-term project. It is not yet usable by game developers. However, feel free to run the solari example (`cargo run --release --example solari --features bevy_solari` (realtime direct lighting, no denoising) or `cargo run --release --example solari --features bevy_solari -- --pathtracer` (non-realtime pathtracing)) to check out the progress we've made, and look forward to more work on Bevy Solari in future releases! +While Bevy 0.17 adds the bevy_solari crate, it's intended as a long-term project. It is not yet usable by game developers. However, feel free to run the solari example (`cargo run --release --example solari --features bevy_solari` (realtime direct and 1-bounce indirect lighting, no denoising) or `cargo run --release --example solari --features bevy_solari -- --pathtracer` (non-realtime pathtracing)) to check out the progress we've made, and look forward to more work on Bevy Solari in future releases! (TODO: Embed bevy_solari logo here, or somewhere else that looks good) From 3d93fc835fde54fbde9787450bcd8a6b64d35aee Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 7 Jul 2025 17:09:22 -0400 Subject: [PATCH 11/12] Add new PR to release notes --- release-content/release-notes/bevy_solari.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index 86098290bbdde..e6727deb47b70 100644 --- a/release-content/release-notes/bevy_solari.md +++ b/release-content/release-notes/bevy_solari.md @@ -1,7 +1,7 @@ --- title: Initial raytraced lighting progress (bevy_solari) authors: ["@JMS55"] -pull_requests: [19058, 19620, 19790] +pull_requests: [19058, 19620, 19790, 20020] --- (TODO: Embed solari example screenshot here) From bc29cb418e77f09dd8fa180dc8ad5569d5da4498 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 8 Jul 2025 22:18:21 -0400 Subject: [PATCH 12/12] Formatting --- crates/bevy_solari/src/realtime/restir_gi.wgsl | 1 - 1 file changed, 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index ea43ee7a1a8d1..e91b267b16cd6 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -157,7 +157,6 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< return spatial_reservoir; } - fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> vec2 { var spatial_id = vec2(center_pixel_id) + vec2(sample_disk(SPATIAL_REUSE_RADIUS_PIXELS, rng)); spatial_id = clamp(spatial_id, vec2(0i), vec2(view.viewport.zw) - 1i);