Solari initial GI

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.
 ///

crates/bevy_solari/src/pathtracer/pathtracer.wgsl

@@ -47,7 +47,8 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3<u32>) {
             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);

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::<SolariLighting>()
             .insert_resource(DefaultOpaqueRendererMethod::deferred());

crates/bevy_solari/src/realtime/node.rs

@@ -36,8 +36,10 @@ 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,
+    gi_spatial_and_shade_pipeline: CachedComputePipelineId,
 }
 
 impl ViewNode for SolariLightingNode {
@@ -72,8 +74,10 @@ impl ViewNode for SolariLightingNode {
         let previous_view_uniforms = world.resource::<PreviousViewUniforms>();
         let frame_count = world.resource::<FrameCount>();
         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(gi_spatial_and_shade_pipeline),
             Some(scene_bindings),
             Some(viewport),
             Some(gbuffer),
@@ -82,8 +86,10 @@ 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),
+            pipeline_cache.get_compute_pipeline(self.gi_spatial_and_shade_pipeline),
             &scene_bindings.bind_group,
             camera.physical_viewport_size,
             view_prepass_textures.deferred_view(),
@@ -101,8 +107,18 @@ 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(),
+                solari_lighting_resources
+                    .gi_reservoirs_a
+                    .as_entire_binding(),
+                solari_lighting_resources
+                    .gi_reservoirs_b
+                    .as_entire_binding(),
                 gbuffer,
                 depth_buffer,
                 motion_vectors,
@@ -135,14 +151,20 @@ 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.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);
@@ -189,10 +211,12 @@ 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),
+                    storage_buffer_sized(false, None),
+                    storage_buffer_sized(false, None),
                     texture_2d(TextureSampleType::Uint),
                     texture_depth_2d(),
                     texture_2d(TextureSampleType::Float { filterable: true }),
@@ -204,9 +228,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 +244,9 @@ impl FromWorld for SolariLightingNode {
                 ..default()
             });
 
-        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(),
@@ -236,10 +260,46 @@ impl FromWorld for SolariLightingNode {
                 ..default()
             });
 
+        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: Some("initial_and_temporal".into()),
+                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: Some("spatial_and_shade".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,
+            gi_spatial_and_shade_pipeline,
         }
     }
 }

crates/bevy_solari/src/realtime/prepare.rs

@@ -17,14 +17,19 @@ use bevy_render::{
     renderer::RenderDevice,
 };
 
-/// Size of a Reservoir shader struct in bytes.
-const RESERVOIR_STRUCT_SIZE: u64 = 32;
+/// 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 reservoirs_a: Buffer,
-    pub reservoirs_b: Buffer,
+    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 * RESERVOIR_STRUCT_SIZE;
+        let di_reservoirs_a = render_device.create_buffer(&BufferDescriptor {
+            label: Some("solari_lighting_di_reservoirs_a"),
+            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: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE,
+            usage: BufferUsages::STORAGE,
+            mapped_at_creation: false,
+        });
 
-        let reservoirs_a = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_reservoirs_a"),
-            size,
+        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 reservoirs_b = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_reservoirs_b"),
-            size,
+        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,
         });
@@ -88,8 +105,10 @@ 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,
+            gi_reservoirs_a,
+            gi_reservoirs_b,
             previous_gbuffer: (previous_gbuffer, previous_gbuffer_view),
             previous_depth: (previous_depth, previous_depth_view),
             view_size,

crates/bevy_solari/src/realtime/restir_di.wgsl

@@ -10,16 +10,16 @@
 #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<rgba16float, write>;
-@group(1) @binding(1) var<storage, read_write> reservoirs_a: array<Reservoir>;
-@group(1) @binding(2) var<storage, read_write> reservoirs_b: array<Reservoir>;
-@group(1) @binding(3) var gbuffer: texture_2d<u32>;
-@group(1) @binding(4) var depth_buffer: texture_depth_2d;
-@group(1) @binding(5) var motion_vectors: texture_2d<f32>;
-@group(1) @binding(6) var previous_gbuffer: texture_2d<u32>;
-@group(1) @binding(7) var previous_depth_buffer: texture_depth_2d;
-@group(1) @binding(8) var<uniform> view: View;
-@group(1) @binding(9) var<uniform> previous_view: PreviousViewUniforms;
+@group(1) @binding(0) var view_output: texture_storage_2d<rgba16float, read_write>;
+@group(1) @binding(1) var<storage, read_write> di_reservoirs_a: array<Reservoir>;
+@group(1) @binding(2) var<storage, read_write> di_reservoirs_b: array<Reservoir>;
+@group(1) @binding(5) var gbuffer: texture_2d<u32>;
+@group(1) @binding(6) var depth_buffer: texture_depth_2d;
+@group(1) @binding(7) var motion_vectors: texture_2d<f32>;
+@group(1) @binding(8) var previous_gbuffer: texture_2d<u32>;
+@group(1) @binding(9) var previous_depth_buffer: texture_depth_2d;
+@group(1) @binding(10) var<uniform> view: View;
+@group(1) @binding(11) var<uniform> previous_view: PreviousViewUniforms;
 struct PushConstants { frame_index: u32, reset: u32 }
 var<push_constant> constants: PushConstants;
 
@@ -38,7 +38,7 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3<u32>) {
 
     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);
@@ -48,10 +48,10 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3<u32>) {
     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);
 
-    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<u32>) {
 
     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<u32>) {
     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;
@@ -119,14 +119,18 @@ fn generate_initial_reservoir(world_position: vec3<f32>, world_normal: vec3<f32>
     return reservoir;
 }
 
-fn load_temporal_reservoir(pixel_id: vec2<u32>, depth: f32, world_position: vec3<f32>, world_normal: vec3<f32>) -> Reservoir {
+fn load_temporal_reservoir(pixel_id: vec2<u32>, depth: f32, world_position: vec3<f32>, world_normal: vec3<f32>, rng: ptr<function, u32>) -> Reservoir {
     let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy;
     let temporal_pixel_id_float = round(vec2<f32>(pixel_id) - (motion_vector * view.viewport.zw));
     let temporal_pixel_id = vec2<u32>(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
     if any(temporal_pixel_id_float < vec2(0.0)) || any(temporal_pixel_id_float >= view.viewport.zw) || bool(constants.reset) {
         return empty_reservoir();
     }
 
+    // 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);
@@ -136,8 +140,9 @@ fn load_temporal_reservoir(pixel_id: vec2<u32>, 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];
 
+    // 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();
@@ -160,7 +165,7 @@ fn load_spatial_reservoir(pixel_id: vec2<u32>, 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 +214,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,