From 27b03513c973f9ab03ddddd23cd4ea5b60698c07 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 18 Apr 2026 11:28:23 -0400 Subject: [PATCH 01/11] Solari: Minor bugfix --- crates/bevy_solari/src/realtime/restir_gi.wgsl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index ed7c197c97fc8..ba78f02e55366 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -99,7 +99,7 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 let sample_point = resolve_ray_hit_full(ray); - if all(sample_point.material.emissive != vec3(0.0)) { + if any(sample_point.material.emissive != vec3(0.0)) { return reservoir; } From 0f93c0c0e6211b466635d2faa50dc765a6976207 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 22:46:33 -0400 Subject: [PATCH 02/11] Half res diffuse GI prototype --- crates/bevy_solari/src/realtime/mod.rs | 3 + crates/bevy_solari/src/realtime/node.rs | 58 ++++------ crates/bevy_solari/src/realtime/prepare.rs | 64 ++++++---- .../src/realtime/presample_light_tiles.wgsl | 3 +- .../src/realtime/realtime_bindings.wgsl | 2 +- .../realtime/resolve_dlss_rr_textures.wgsl | 1 - .../bevy_solari/src/realtime/restir_di.wgsl | 6 +- .../bevy_solari/src/realtime/restir_gi.wgsl | 109 ++++++++++++++---- .../bevy_solari/src/realtime/specular_gi.wgsl | 38 +++++- .../src/realtime/world_cache_query.wgsl | 6 - .../src/realtime/world_cache_update.wgsl | 3 +- .../src/scene/raytracing_scene_bindings.wgsl | 1 - crates/bevy_solari/src/scene/sampling.wgsl | 2 +- 13 files changed, 192 insertions(+), 104 deletions(-) diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index bf8350eb16138..85aa21f423239 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -93,6 +93,8 @@ impl Plugin for SolariLightingPlugin { DepthPrepassDoubleBuffer )] pub struct SolariLighting { + /// Set to true to greatly improve performance at the cost of quality. + pub quarter_resolution_indirect_lighting: bool, /// Set to true to delete the saved temporal history (past frames). /// /// Useful for preventing ghosting when the history is no longer @@ -106,6 +108,7 @@ pub struct SolariLighting { impl Default for SolariLighting { fn default() -> Self { Self { + quarter_resolution_indirect_lighting: true, reset: true, // No temporal history on the first frame } } diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 3b36a38f9588c..93a185b3cb924 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -234,6 +234,12 @@ pub fn solari_lighting( // Choice of number here is arbitrary let frame_index = frame_count.0.wrapping_mul(5782582); + let immediates = [ + frame_index, + solari_lighting.reset as u32, + solari_lighting.quarter_resolution_indirect_lighting as u32, + ]; + let immediates = bytemuck::cast_slice(&immediates); let diagnostics = ctx.diagnostic_recorder(); let diagnostics = diagnostics.as_deref(); @@ -259,6 +265,12 @@ pub fn solari_lighting( let dx = solari_lighting_resources.view_size.x.div_ceil(8); let dy = solari_lighting_resources.view_size.y.div_ceil(8); + let mut gi_dx = solari_lighting_resources.view_size.x.div_ceil(8); + let mut gi_dy = solari_lighting_resources.view_size.y.div_ceil(8); + if solari_lighting.quarter_resolution_indirect_lighting { + gi_dx = gi_dx.div_ceil(2); + gi_dy = gi_dy.div_ceil(2); + } pass.set_bind_group(0, scene_bind_group, &[]); pass.set_bind_group( @@ -279,10 +291,7 @@ pub fn solari_lighting( let d = diagnostics.time_span(&mut pass, "solari_lighting/presample_light_tiles"); pass.set_pipeline(presample_light_tiles_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); + pass.set_immediates(0, immediates); pass.dispatch_workgroups(LIGHT_TILE_BLOCKS as u32, 1, 1); d.end(&mut pass); @@ -305,20 +314,14 @@ pub fn solari_lighting( pass.set_bind_group(2, None, &[]); pass.set_pipeline(sample_di_for_world_cache_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); + pass.set_immediates(0, immediates); pass.dispatch_workgroups_indirect( &solari_lighting_resources.world_cache_active_cells_dispatch, 0, ); pass.set_pipeline(sample_gi_for_world_cache_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); + pass.set_immediates(0, immediates); pass.dispatch_workgroups_indirect( &solari_lighting_resources.world_cache_active_cells_dispatch, 0, @@ -335,17 +338,11 @@ pub fn solari_lighting( let d = diagnostics.time_span(&mut pass, "solari_lighting/direct_lighting"); pass.set_pipeline(di_initial_and_temporal_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); + pass.set_immediates(0, immediates); pass.dispatch_workgroups(dx, dy, 1); pass.set_pipeline(di_spatial_and_shade_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); + pass.set_immediates(0, immediates); pass.dispatch_workgroups(dx, dy, 1); d.end(&mut pass); @@ -353,18 +350,12 @@ pub fn solari_lighting( let d = diagnostics.time_span(&mut pass, "solari_lighting/diffuse_indirect_lighting"); pass.set_pipeline(gi_initial_and_temporal_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); - pass.dispatch_workgroups(dx, dy, 1); + pass.set_immediates(0, immediates); + pass.dispatch_workgroups(gi_dx, gi_dy, 1); pass.set_pipeline(gi_spatial_and_shade_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); - pass.dispatch_workgroups(dx, dy, 1); + pass.set_immediates(0, immediates); + pass.dispatch_workgroups(gi_dx, gi_dy, 1); d.end(&mut pass); @@ -374,10 +365,7 @@ pub fn solari_lighting( pass.set_bind_group(2, bind_group_resolve_dlss_rr_textures, &[]); } pass.set_pipeline(specular_gi_pipeline); - pass.set_immediates( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); + pass.set_immediates(0, immediates); pass.dispatch_workgroups(dx, dy, 1); d.end(&mut pass); @@ -471,7 +459,7 @@ pub fn init_solari_lighting_pipelines( pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { label: Some(label.into()), layout, - immediate_size: 8, + immediate_size: 12, shader, shader_defs, entry_point: Some(entry_point.into()), diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 865c8c8fba807..b0025c7d30e64 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -9,7 +9,6 @@ use bevy_ecs::query::Has; use bevy_ecs::{ component::Component, entity::Entity, - query::With, system::{Commands, Query, Res}, }; use bevy_image::ToExtents; @@ -61,37 +60,41 @@ pub struct SolariLightingResources { pub world_cache_active_cells_count: Buffer, pub world_cache_active_cells_dispatch: Buffer, pub view_size: UVec2, + pub quarter_resolution_indirect_lighting: bool, } pub fn prepare_solari_lighting_resources( - #[cfg(any(not(feature = "dlss"), feature = "force_disable_dlss"))] query: Query< - ( - Entity, - &ExtractedCamera, - Option<&SolariLightingResources>, - Option<&MainPassResolutionOverride>, - ), - With, - >, - #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] query: Query< - ( - Entity, - &ExtractedCamera, - Option<&SolariLightingResources>, - Option<&MainPassResolutionOverride>, - Has>, - ), - With, - >, + #[cfg(any(not(feature = "dlss"), feature = "force_disable_dlss"))] query: Query<( + Entity, + &SolariLighting, + &ExtractedCamera, + Option<&SolariLightingResources>, + Option<&MainPassResolutionOverride>, + )>, + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] query: Query<( + Entity, + &SolariLighting, + &ExtractedCamera, + Option<&SolariLightingResources>, + Option<&MainPassResolutionOverride>, + Has>, + )>, render_device: Res, mut commands: Commands, ) { for query_item in &query { #[cfg(any(not(feature = "dlss"), feature = "force_disable_dlss"))] - let (entity, camera, solari_lighting_resources, resolution_override) = query_item; - #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] - let (entity, camera, solari_lighting_resources, resolution_override, has_dlss_rr) = + let (entity, solari_lighting, camera, solari_lighting_resources, resolution_override) = query_item; + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] + let ( + entity, + solari_lighting, + camera, + solari_lighting_resources, + resolution_override, + has_dlss_rr, + ) = query_item; let Some(mut view_size) = camera.physical_viewport_size else { continue; @@ -100,7 +103,12 @@ pub fn prepare_solari_lighting_resources( view_size = *resolution_override; } - if solari_lighting_resources.map(|r| r.view_size) == Some(view_size) { + if solari_lighting_resources.map(|r| (r.view_size, r.quarter_resolution_indirect_lighting)) + == Some(( + view_size, + solari_lighting.quarter_resolution_indirect_lighting, + )) + { continue; } @@ -138,6 +146,12 @@ pub fn prepare_solari_lighting_resources( let di_reservoirs_b = di_reservoirs("solari_lighting_di_reservoirs_b"); let gi_reservoirs = |name| { + let mut view_size = view_size; + if solari_lighting.quarter_resolution_indirect_lighting { + view_size.x = view_size.x.div_ceil(2); + view_size.y = view_size.y.div_ceil(2); + } + render_device.create_buffer(&BufferDescriptor { label: Some(name), size: (view_size.x * view_size.y) as u64 * GI_RESERVOIR_STRUCT_SIZE, @@ -244,6 +258,8 @@ pub fn prepare_solari_lighting_resources( world_cache_active_cells_count, world_cache_active_cells_dispatch, view_size, + quarter_resolution_indirect_lighting: solari_lighting + .quarter_resolution_indirect_lighting, }); #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] diff --git a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl index 1b67a4c71f35c..5954458370c9e 100644 --- a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl +++ b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl @@ -6,8 +6,7 @@ enable wgpu_ray_query; #import bevy_pbr::rgb9e5::{vec3_to_rgb9e5_, rgb9e5_to_vec3_} #import bevy_pbr::utils::{octahedral_encode, octahedral_decode} -#import bevy_render::view::View -#import bevy_solari::sampling::{generate_random_light_sample, LightSample, ResolvedLightSample} +#import bevy_solari::sampling::{generate_random_light_sample, ResolvedLightSample} #import bevy_solari::realtime_bindings::{light_tile_samples, light_tile_resolved_samples, view, constants, ResolvedLightSamplePacked} @compute @workgroup_size(1024, 1, 1) diff --git a/crates/bevy_solari/src/realtime/realtime_bindings.wgsl b/crates/bevy_solari/src/realtime/realtime_bindings.wgsl index 5dbbc65a3a470..624dcfb62b59d 100644 --- a/crates/bevy_solari/src/realtime/realtime_bindings.wgsl +++ b/crates/bevy_solari/src/realtime/realtime_bindings.wgsl @@ -43,7 +43,7 @@ enable wgpu_ray_query; @group(2) @binding(3) var specular_motion_vectors: texture_storage_2d; #endif -struct PushConstants { frame_index: u32, reset: u32 } +struct PushConstants { frame_index: u32, reset: u32, quarter_resolution_indirect_lighting: u32 } var constants: PushConstants; // Don't adjust the size of this struct without also adjusting `prepare::RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE`. diff --git a/crates/bevy_solari/src/realtime/resolve_dlss_rr_textures.wgsl b/crates/bevy_solari/src/realtime/resolve_dlss_rr_textures.wgsl index c5799c4078a09..041538b940820 100644 --- a/crates/bevy_solari/src/realtime/resolve_dlss_rr_textures.wgsl +++ b/crates/bevy_solari/src/realtime/resolve_dlss_rr_textures.wgsl @@ -2,7 +2,6 @@ enable wgpu_ray_query; #define_import_path bevy_solari::resolve_dlss_rr_textures #import bevy_pbr::pbr_functions::{calculate_diffuse_color, calculate_F0} -#import bevy_render::view::View #import bevy_solari::gbuffer_utils::gpixel_resolve #import bevy_solari::realtime_bindings::{gbuffer, depth_buffer, view, diffuse_albedo, specular_albedo, normal_roughness, specular_motion_vectors} diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index c9abc93faec50..4335deefed4df 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -4,17 +4,15 @@ enable wgpu_ray_query; #import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance -#import bevy_pbr::prepass_bindings::PreviousViewUniforms #import bevy_pbr::utils::{rand_f, rand_range_u, sample_disk} #import bevy_render::maths::PI -#import bevy_render::view::View #import bevy_solari::brdf::{evaluate_diffuse_brdf, evaluate_specular_brdf} #import bevy_solari::gbuffer_utils::{gpixel_resolve, pixel_dissimilar, permute_pixel} #import bevy_solari::presample_light_tiles::unpack_resolved_light_sample -#import bevy_solari::sampling::{LightSample, ResolvedLightSample, NULL_LIGHT_ID, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility, balance_heuristic} +#import bevy_solari::sampling::{LightSample, ResolvedLightSample, NULL_LIGHT_ID, calculate_resolved_light_contribution, resolve_light_sample, trace_light_visibility, balance_heuristic} #import bevy_solari::scene_bindings::{light_sources, previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME, RAY_T_MIN} #import bevy_solari::specular_gi::SPECULAR_GI_FOR_DI_ROUGHNESS_THRESHOLD -#import bevy_solari::realtime_bindings::{view_output, light_tile_samples, light_tile_resolved_samples, di_reservoirs_a, di_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants, ResolvedLightSamplePacked} +#import bevy_solari::realtime_bindings::{view_output, light_tile_samples, light_tile_resolved_samples, di_reservoirs_a, di_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants} const INITIAL_SAMPLES = 8u; const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index ba78f02e55366..10b2e69687cce 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -2,10 +2,8 @@ enable wgpu_ray_query; #import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance -#import bevy_pbr::prepass_bindings::PreviousViewUniforms #import bevy_pbr::utils::{rand_f, sample_uniform_hemisphere, uniform_hemisphere_inverse_pdf, sample_disk} #import bevy_render::maths::PI -#import bevy_render::view::View #import bevy_solari::brdf::evaluate_diffuse_brdf #import bevy_solari::gbuffer_utils::{gpixel_resolve, pixel_dissimilar, permute_pixel} #import bevy_solari::sampling::{sample_random_light, trace_point_visibility, balance_heuristic, isnan} @@ -18,11 +16,12 @@ const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 8.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.main_pass_viewport.zw)) { return; } +fn initial_and_temporal(@builtin(global_invocation_id) thread_id: vec3) { + if any(thread_id.xy >= gi_resolution()) { return; } - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); - var rng = pixel_index + constants.frame_index; + let global_id = vec3(gi_thread_to_full_resolution_pixel(thread_id.xy), 0u); + let pixel_index = gi_reservoir_index(global_id.xy); + var rng = (global_id.x + global_id.y * u32(view.main_pass_viewport.z)) + constants.frame_index; let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { @@ -44,11 +43,12 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { } @compute @workgroup_size(8, 8, 1) -fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { - if any(global_id.xy >= vec2u(view.main_pass_viewport.zw)) { return; } +fn spatial_and_shade(@builtin(global_invocation_id) thread_id: vec3) { + if any(thread_id.xy >= gi_resolution()) { return; } - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); - var rng = pixel_index + constants.frame_index; + let global_id = vec3(gi_thread_to_full_resolution_pixel(thread_id.xy), 0u); + let pixel_index = gi_reservoir_index(global_id.xy); + var rng = (global_id.x + global_id.y * u32(view.main_pass_viewport.z)) + constants.frame_index; let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { @@ -82,6 +82,10 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let wo = normalize(view.world_position - surface.world_position); let brdf = evaluate_diffuse_brdf(wo, merge_result.wi, surface.world_normal, surface.material); + if bool(constants.quarter_resolution_indirect_lighting) { + combined_reservoir.unbiased_contribution_weight *= 4.0; + } + var pixel_color = textureLoad(view_output, global_id.xy); pixel_color += vec4(merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight * view.exposure * brdf, 0.0); textureStore(view_output, global_id.xy, pixel_color); @@ -132,12 +136,18 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return NeighborInfo(empty_reservoir(), vec3(0.0), vec3(0.0), vec3(0.0)); } - let permuted_temporal_pixel_id = permute_pixel(vec2(temporal_pixel_id_float), constants.frame_index, view.main_pass_viewport.zw); - var temporal = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); + let point_temporal_pixel_id = vec2(temporal_pixel_id_float); + var temporal: NeighborInfo; + if bool(constants.quarter_resolution_indirect_lighting) { + temporal = load_temporal_reservoir_inner(point_temporal_pixel_id, depth, world_position, world_normal); + } else { + let permuted_temporal_pixel_id = permute_pixel(point_temporal_pixel_id, constants.frame_index, view.main_pass_viewport.zw); + temporal = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); - // If permuted reprojection failed (tends to happen on object edges), try point reprojection - if all(temporal.reservoir.radiance == vec3(0.0)) { - temporal = load_temporal_reservoir_inner(vec2(temporal_pixel_id_float), depth, world_position, world_normal); + // If permuted reprojection failed (tends to happen on object edges), try point reprojection + if all(temporal.reservoir.radiance == vec3(0.0)) { + temporal = load_temporal_reservoir_inner(point_temporal_pixel_id, depth, world_position, world_normal); + } } temporal.reservoir.confidence_weight = min(temporal.reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); @@ -145,8 +155,9 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return temporal; } -fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { - // Check if the pixel features have changed heavily between the current and previous frame +fn load_temporal_reservoir_inner(temporal_pixel_id_in: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { + let temporal_pixel_id = gi_snap_to_quad_pixel_previous_frame(temporal_pixel_id_in); + let temporal_depth = textureLoad(previous_depth_buffer, temporal_pixel_id, 0); let temporal_surface = gpixel_resolve(textureLoad(previous_gbuffer, temporal_pixel_id, 0), temporal_depth, temporal_pixel_id, view.main_pass_viewport.zw, previous_view.world_from_clip); let temporal_diffuse_brdf = temporal_surface.material.base_color / PI; @@ -154,8 +165,7 @@ fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world return NeighborInfo(empty_reservoir(), vec3(0.0), vec3(0.0), vec3(0.0)); } - let temporal_pixel_index = temporal_pixel_id.x + temporal_pixel_id.y * u32(view.main_pass_viewport.z); - let temporal_reservoir = gi_reservoirs_a[temporal_pixel_index]; + let temporal_reservoir = gi_reservoirs_a[gi_reservoir_index(temporal_pixel_id)]; return NeighborInfo(temporal_reservoir, temporal_surface.world_position, temporal_surface.world_normal, temporal_diffuse_brdf); } @@ -173,8 +183,7 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< continue; } - let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.main_pass_viewport.z); - let spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; + let spatial_reservoir = gi_reservoirs_b[gi_reservoir_index(spatial_pixel_id)]; return NeighborInfo(spatial_reservoir, spatial_surface.world_position, spatial_surface.world_normal, spatial_diffuse_brdf); } @@ -182,9 +191,15 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< } fn get_neighbor_pixel_id(center_pixel_id: vec2, search_radius: f32, rng: ptr) -> vec2 { - var spatial_id = vec2(center_pixel_id) + sample_disk(search_radius, rng); - spatial_id = clamp(spatial_id, vec2(0.0), view.main_pass_viewport.zw - 1.0); - return vec2(spatial_id); + if bool(constants.quarter_resolution_indirect_lighting) { + var spatial_id = vec2(center_pixel_id / vec2(2u)) + sample_disk(search_radius, rng) * 0.5; + spatial_id = clamp(spatial_id, vec2(0.0), vec2(quarter_resolution_dimensions()) - 1.0); + return quarter_to_full_resolution_pixel(vec2(spatial_id), constants.frame_index); + } else { + var spatial_id = vec2(center_pixel_id) + sample_disk(search_radius, rng); + spatial_id = clamp(spatial_id, vec2(0.0), view.main_pass_viewport.zw - 1.0); + return vec2(spatial_id); + } } struct NeighborInfo { @@ -318,3 +333,49 @@ fn merge_reservoirs( return ReservoirMergeResult(combined_reservoir, canonical_sample_radiance, canonical_sample_wi); } } + +fn gi_resolution() -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_resolution_dimensions(); + } else { + return vec2u(view.main_pass_viewport.zw); + } +} + +fn gi_thread_to_full_resolution_pixel(thread_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_to_full_resolution_pixel(thread_xy, constants.frame_index); + } else { + return thread_xy; + } +} + +fn gi_reservoir_index(full_xy: vec2) -> u32 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_resolution_index(full_xy / vec2(2u)); + } else { + return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); + } +} + +fn gi_snap_to_quad_pixel_previous_frame(full_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_to_full_resolution_pixel(full_xy / vec2(2u), constants.frame_index - 5782582u); + } else { + return full_xy; + } +} + +fn quarter_resolution_dimensions() -> vec2 { + return (vec2u(view.main_pass_viewport.zw) + vec2(1u)) / vec2(2u); +} + +fn quarter_resolution_index(quarter_xy: vec2) -> u32 { + return quarter_xy.x + quarter_xy.y * quarter_resolution_dimensions().x; +} + +fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { + var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; + let qi = u32(rand_f(&rng) * 4.0); + return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - vec2(1u)); +} diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index 3fbea809c4ebe..6126eff79d087 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -3,8 +3,8 @@ enable wgpu_ray_query; #define_import_path bevy_solari::specular_gi #import bevy_pbr::pbr_functions::{calculate_tbn_mikktspace, calculate_diffuse_color, calculate_F0} +#import bevy_pbr::utils::rand_f #import bevy_render::maths::{orthonormalize, PI} -#import bevy_render::view::View #import bevy_solari::brdf::{evaluate_brdf, evaluate_specular_brdf} #import bevy_solari::gbuffer_utils::{gpixel_resolve, ResolvedGPixel} #import bevy_solari::sampling::{sample_random_light, random_emissive_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, ggx_vndf_sample_invalid, power_heuristic} @@ -40,9 +40,11 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { var wi: vec3; if surface.material.roughness > DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD { // Surface is very rough, reuse the ReSTIR GI reservoir - let gi_reservoir = gi_reservoirs_a[pixel_index]; + if any(global_id.xy != gi_snap_to_quad_pixel(global_id.xy)) { return; } + let gi_reservoir = gi_reservoirs_a[gi_reservoir_index(global_id.xy)]; + let ucw = select(1.0, 4.0, bool(constants.quarter_resolution_indirect_lighting)); wi = normalize(gi_reservoir.sample_point_world_position - surface.world_position); - radiance = gi_reservoir.radiance * gi_reservoir.unbiased_contribution_weight; + radiance = gi_reservoir.radiance * gi_reservoir.unbiased_contribution_weight * ucw; } else { // Surface is glossy or mirror-like, trace a new path let TBN = orthonormalize(surface.world_normal); @@ -241,3 +243,33 @@ fn calculate_motion_vector(world_position: vec3, previous_world_position: v return (clip_position - previous_clip_position) * vec2(0.5, -0.5); } #endif + +fn gi_snap_to_quad_pixel(full_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_to_full_resolution_pixel(full_xy / vec2(2u), constants.frame_index); + } else { + return full_xy; + } +} + +fn gi_reservoir_index(full_xy: vec2) -> u32 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_resolution_index(full_xy / vec2(2u)); + } else { + return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); + } +} + +fn quarter_resolution_dimensions() -> vec2 { + return (vec2u(view.main_pass_viewport.zw) + vec2(1u)) / vec2(2u); +} + +fn quarter_resolution_index(quarter_xy: vec2) -> u32 { + return quarter_xy.x + quarter_xy.y * quarter_resolution_dimensions().x; +} + +fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { + var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; + let qi = u32(rand_f(&rng) * 4.0); + return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - vec2(1u)); +} diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index 9245d0064197f..943f56315f375 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -9,12 +9,6 @@ enable wgpu_ray_query; world_cache_checksums, world_cache_radiance, world_cache_geometry_data, - world_cache_luminance_deltas, - world_cache_a, - world_cache_b, - world_cache_active_cell_indices, - world_cache_active_cells_count, - WorldCacheGeometryData, } /// How responsive the world cache is to changes in lighting (higher is less responsive, lower is more responsive) diff --git a/crates/bevy_solari/src/realtime/world_cache_update.wgsl b/crates/bevy_solari/src/realtime/world_cache_update.wgsl index 46c1d4497381f..2d92ad16d4a94 100644 --- a/crates/bevy_solari/src/realtime/world_cache_update.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_update.wgsl @@ -2,8 +2,7 @@ enable wgpu_ray_query; #import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance #import bevy_pbr::utils::{rand_f, rand_range_u, sample_cosine_hemisphere} -#import bevy_render::view::View -#import bevy_solari::presample_light_tiles::{ResolvedLightSamplePacked, unpack_resolved_light_sample} +#import bevy_solari::presample_light_tiles::unpack_resolved_light_sample #import bevy_solari::sampling::{calculate_resolved_light_contribution, trace_light_visibility} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN} #import bevy_solari::world_cache::{ diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index 0e7f7a977b73a..dfc00bc4a2271 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -2,7 +2,6 @@ enable wgpu_ray_query; #define_import_path bevy_solari::scene_bindings -#import bevy_pbr::lighting::perceptualRoughnessToRoughness #import bevy_pbr::pbr_functions::calculate_tbn_mikktspace struct InstanceGeometryIds { diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index e37a28b7ba240..b39727e023899 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -3,7 +3,7 @@ enable wgpu_ray_query; #define_import_path bevy_solari::sampling #import bevy_pbr::lighting::D_GGX -#import bevy_pbr::utils::{rand_f, rand_vec2f, rand_u, rand_range_u} +#import bevy_pbr::utils::{rand_vec2f, rand_u, rand_range_u} #import bevy_render::maths::{PI_2, orthonormalize} #import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LightSource, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full, ResolvedRayHitFull, MIRROR_ROUGHNESS_THRESHOLD} From b3bc686de6c16d11a22f03c2166104e8c5edafb9 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 23:04:43 -0400 Subject: [PATCH 03/11] Turns out quarter-res still needs permutation --- .../bevy_solari/src/realtime/restir_gi.wgsl | 26 ++++++++----------- .../bevy_solari/src/realtime/specular_gi.wgsl | 8 +++--- 2 files changed, 15 insertions(+), 19 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 10b2e69687cce..d44ae7019e5dd 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -137,17 +137,13 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 } let point_temporal_pixel_id = vec2(temporal_pixel_id_float); - var temporal: NeighborInfo; - if bool(constants.quarter_resolution_indirect_lighting) { - temporal = load_temporal_reservoir_inner(point_temporal_pixel_id, depth, world_position, world_normal); - } else { - let permuted_temporal_pixel_id = permute_pixel(point_temporal_pixel_id, constants.frame_index, view.main_pass_viewport.zw); - temporal = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); + let scale = select(1u, 2u, bool(constants.quarter_resolution_indirect_lighting)); + let permuted_temporal_pixel_id = permute_pixel(point_temporal_pixel_id / scale, constants.frame_index, vec2(gi_resolution())) * scale; + var temporal = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); - // If permuted reprojection failed (tends to happen on object edges), try point reprojection - if all(temporal.reservoir.radiance == vec3(0.0)) { - temporal = load_temporal_reservoir_inner(point_temporal_pixel_id, depth, world_position, world_normal); - } + // If permuted reprojection failed (tends to happen on object edges), try point reprojection + if all(temporal.reservoir.radiance == vec3(0.0)) { + temporal = load_temporal_reservoir_inner(point_temporal_pixel_id, depth, world_position, world_normal); } temporal.reservoir.confidence_weight = min(temporal.reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); @@ -192,7 +188,7 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< fn get_neighbor_pixel_id(center_pixel_id: vec2, search_radius: f32, rng: ptr) -> vec2 { if bool(constants.quarter_resolution_indirect_lighting) { - var spatial_id = vec2(center_pixel_id / vec2(2u)) + sample_disk(search_radius, rng) * 0.5; + var spatial_id = vec2(center_pixel_id / 2u) + sample_disk(search_radius, rng) * 0.5; spatial_id = clamp(spatial_id, vec2(0.0), vec2(quarter_resolution_dimensions()) - 1.0); return quarter_to_full_resolution_pixel(vec2(spatial_id), constants.frame_index); } else { @@ -352,7 +348,7 @@ fn gi_thread_to_full_resolution_pixel(thread_xy: vec2) -> vec2 { fn gi_reservoir_index(full_xy: vec2) -> u32 { if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_resolution_index(full_xy / vec2(2u)); + return quarter_resolution_index(full_xy / 2u); } else { return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); } @@ -360,14 +356,14 @@ fn gi_reservoir_index(full_xy: vec2) -> u32 { fn gi_snap_to_quad_pixel_previous_frame(full_xy: vec2) -> vec2 { if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_to_full_resolution_pixel(full_xy / vec2(2u), constants.frame_index - 5782582u); + return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index - 5782582u); } else { return full_xy; } } fn quarter_resolution_dimensions() -> vec2 { - return (vec2u(view.main_pass_viewport.zw) + vec2(1u)) / vec2(2u); + return (vec2u(view.main_pass_viewport.zw) + 1u) / 2u; } fn quarter_resolution_index(quarter_xy: vec2) -> u32 { @@ -377,5 +373,5 @@ fn quarter_resolution_index(quarter_xy: vec2) -> u32 { fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; let qi = u32(rand_f(&rng) * 4.0); - return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - vec2(1u)); + return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - 1u); } diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index 6126eff79d087..c5af6affa1f8e 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -246,7 +246,7 @@ fn calculate_motion_vector(world_position: vec3, previous_world_position: v fn gi_snap_to_quad_pixel(full_xy: vec2) -> vec2 { if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_to_full_resolution_pixel(full_xy / vec2(2u), constants.frame_index); + return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index); } else { return full_xy; } @@ -254,14 +254,14 @@ fn gi_snap_to_quad_pixel(full_xy: vec2) -> vec2 { fn gi_reservoir_index(full_xy: vec2) -> u32 { if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_resolution_index(full_xy / vec2(2u)); + return quarter_resolution_index(full_xy / 2u); } else { return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); } } fn quarter_resolution_dimensions() -> vec2 { - return (vec2u(view.main_pass_viewport.zw) + vec2(1u)) / vec2(2u); + return (vec2u(view.main_pass_viewport.zw) + 1u) / 2u; } fn quarter_resolution_index(quarter_xy: vec2) -> u32 { @@ -271,5 +271,5 @@ fn quarter_resolution_index(quarter_xy: vec2) -> u32 { fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; let qi = u32(rand_f(&rng) * 4.0); - return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - vec2(1u)); + return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - 1u); } From b3d9214f8e3991415c87a7ef0252316f43bc4ff8 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 23:16:59 -0400 Subject: [PATCH 04/11] Quarter-res specular --- crates/bevy_solari/src/realtime/gi_utils.wgsl | 58 +++++++++++++++++++ crates/bevy_solari/src/realtime/mod.rs | 1 + crates/bevy_solari/src/realtime/node.rs | 2 +- .../bevy_solari/src/realtime/restir_gi.wgsl | 46 +-------------- .../bevy_solari/src/realtime/specular_gi.wgsl | 56 +++++------------- 5 files changed, 74 insertions(+), 89 deletions(-) create mode 100644 crates/bevy_solari/src/realtime/gi_utils.wgsl diff --git a/crates/bevy_solari/src/realtime/gi_utils.wgsl b/crates/bevy_solari/src/realtime/gi_utils.wgsl new file mode 100644 index 0000000000000..f2575106e22ff --- /dev/null +++ b/crates/bevy_solari/src/realtime/gi_utils.wgsl @@ -0,0 +1,58 @@ +#define_import_path bevy_solari::gi_utils + +#import bevy_pbr::utils::rand_f +#import bevy_solari::realtime_bindings::{view, constants} + +fn gi_resolution() -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_resolution_dimensions(); + } else { + return vec2u(view.main_pass_viewport.zw); + } +} + +fn gi_thread_to_full_resolution_pixel(thread_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_to_full_resolution_pixel(thread_xy, constants.frame_index); + } else { + return thread_xy; + } +} + +fn gi_snap_to_quad_pixel(full_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index); + } else { + return full_xy; + } +} + +fn gi_snap_to_quad_pixel_previous_frame(full_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index - 5782582u); + } else { + return full_xy; + } +} + +fn gi_reservoir_index(full_xy: vec2) -> u32 { + if bool(constants.quarter_resolution_indirect_lighting) { + return quarter_resolution_index(full_xy / 2u); + } else { + return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); + } +} + +fn quarter_resolution_dimensions() -> vec2 { + return (vec2u(view.main_pass_viewport.zw) + 1u) / 2u; +} + +fn quarter_resolution_index(quarter_xy: vec2) -> u32 { + return quarter_xy.x + quarter_xy.y * quarter_resolution_dimensions().x; +} + +fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { + var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; + let qi = u32(rand_f(&rng) * 4.0); + return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - 1u); +} diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index 85aa21f423239..8d01bfdfd7fa7 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -36,6 +36,7 @@ impl Plugin for SolariLightingPlugin { fn build(&self, app: &mut App) { load_shader_library!(app, "gbuffer_utils.wgsl"); load_shader_library!(app, "realtime_bindings.wgsl"); + load_shader_library!(app, "gi_utils.wgsl"); load_shader_library!(app, "presample_light_tiles.wgsl"); embedded_asset!(app, "restir_di.wgsl"); embedded_asset!(app, "restir_gi.wgsl"); diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 93a185b3cb924..836d7d41fee8a 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -366,7 +366,7 @@ pub fn solari_lighting( } pass.set_pipeline(specular_gi_pipeline); pass.set_immediates(0, immediates); - pass.dispatch_workgroups(dx, dy, 1); + pass.dispatch_workgroups(gi_dx, gi_dy, 1); d.end(&mut pass); drop(pass); diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index d44ae7019e5dd..c8d56539526d8 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -11,6 +11,7 @@ enable wgpu_ray_query; #import bevy_solari::world_cache::{query_world_cache, WORLD_CACHE_CELL_LIFETIME} #import bevy_solari::realtime_bindings::{view_output, gi_reservoirs_a, gi_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants, Reservoir} #import bevy_solari::specular_gi::DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD +#import bevy_solari::gi_utils::{gi_resolution, gi_thread_to_full_resolution_pixel, gi_reservoir_index, gi_snap_to_quad_pixel_previous_frame, quarter_resolution_dimensions, quarter_to_full_resolution_pixel} const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 8.0; @@ -330,48 +331,3 @@ fn merge_reservoirs( } } -fn gi_resolution() -> vec2 { - if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_resolution_dimensions(); - } else { - return vec2u(view.main_pass_viewport.zw); - } -} - -fn gi_thread_to_full_resolution_pixel(thread_xy: vec2) -> vec2 { - if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_to_full_resolution_pixel(thread_xy, constants.frame_index); - } else { - return thread_xy; - } -} - -fn gi_reservoir_index(full_xy: vec2) -> u32 { - if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_resolution_index(full_xy / 2u); - } else { - return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); - } -} - -fn gi_snap_to_quad_pixel_previous_frame(full_xy: vec2) -> vec2 { - if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index - 5782582u); - } else { - return full_xy; - } -} - -fn quarter_resolution_dimensions() -> vec2 { - return (vec2u(view.main_pass_viewport.zw) + 1u) / 2u; -} - -fn quarter_resolution_index(quarter_xy: vec2) -> u32 { - return quarter_xy.x + quarter_xy.y * quarter_resolution_dimensions().x; -} - -fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { - var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; - let qi = u32(rand_f(&rng) * 4.0); - return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - 1u); -} diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index c5af6affa1f8e..e24d90d8870f9 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -3,8 +3,8 @@ enable wgpu_ray_query; #define_import_path bevy_solari::specular_gi #import bevy_pbr::pbr_functions::{calculate_tbn_mikktspace, calculate_diffuse_color, calculate_F0} -#import bevy_pbr::utils::rand_f #import bevy_render::maths::{orthonormalize, PI} +#import bevy_solari::gi_utils::{gi_resolution, gi_thread_to_full_resolution_pixel, gi_reservoir_index} #import bevy_solari::brdf::{evaluate_brdf, evaluate_specular_brdf} #import bevy_solari::gbuffer_utils::{gpixel_resolve, ResolvedGPixel} #import bevy_solari::sampling::{sample_random_light, random_emissive_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, ggx_vndf_sample_invalid, power_heuristic} @@ -21,16 +21,17 @@ const SPECULAR_GI_FOR_DI_ROUGHNESS_THRESHOLD: f32 = 0.0225; @compute @workgroup_size(8, 8, 1) fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { - if any(global_id.xy >= vec2u(view.main_pass_viewport.zw)) { return; } + if any(global_id.xy >= gi_resolution()) { return; } - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); + let pixel_id = gi_thread_to_full_resolution_pixel(global_id.xy); + let pixel_index = pixel_id.x + pixel_id.y * u32(view.main_pass_viewport.z); var rng = pixel_index + constants.frame_index; - let depth = textureLoad(depth_buffer, global_id.xy, 0); + let depth = textureLoad(depth_buffer, pixel_id, 0); if depth == 0.0 { return; } - let surface = gpixel_resolve(textureLoad(gbuffer, global_id.xy, 0), depth, global_id.xy, view.main_pass_viewport.zw, view.world_from_clip); + let surface = gpixel_resolve(textureLoad(gbuffer, pixel_id, 0), depth, pixel_id, view.main_pass_viewport.zw, view.world_from_clip); let wo_unnormalized = view.world_position - surface.world_position; let wo_length = length(wo_unnormalized); @@ -40,11 +41,9 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { var wi: vec3; if surface.material.roughness > DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD { // Surface is very rough, reuse the ReSTIR GI reservoir - if any(global_id.xy != gi_snap_to_quad_pixel(global_id.xy)) { return; } - let gi_reservoir = gi_reservoirs_a[gi_reservoir_index(global_id.xy)]; - let ucw = select(1.0, 4.0, bool(constants.quarter_resolution_indirect_lighting)); + let gi_reservoir = gi_reservoirs_a[gi_reservoir_index(pixel_id)]; wi = normalize(gi_reservoir.sample_point_world_position - surface.world_position); - radiance = gi_reservoir.radiance * gi_reservoir.unbiased_contribution_weight * ucw; + radiance = gi_reservoir.radiance * gi_reservoir.unbiased_contribution_weight; } else { // Surface is glossy or mirror-like, trace a new path let TBN = orthonormalize(surface.world_normal); @@ -60,7 +59,7 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { wi = wi_tangent.x * T + wi_tangent.y * B + wi_tangent.z * N; let pdf = ggx_vndf_pdf(wo_tangent, wi_tangent, surface.material.roughness); - radiance = trace_glossy_path(global_id.xy, surface, wo_length, wi, pdf, &rng); + radiance = trace_glossy_path(pixel_id, surface, wo_length, wi, pdf, &rng); if surface.material.roughness > MIRROR_ROUGHNESS_THRESHOLD { radiance /= pdf; } @@ -69,13 +68,14 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { let brdf = evaluate_specular_brdf(wo, wi, surface.world_normal, surface.material); radiance *= brdf * view.exposure; + radiance *= select(1.0, 4.0, bool(constants.quarter_resolution_indirect_lighting)); - var pixel_color = textureLoad(view_output, global_id.xy); + var pixel_color = textureLoad(view_output, pixel_id); pixel_color += vec4(radiance, 0.0); - textureStore(view_output, global_id.xy, pixel_color); + textureStore(view_output, pixel_id, pixel_color); #ifdef VISUALIZE_WORLD_CACHE - textureStore(view_output, global_id.xy, vec4(query_world_cache(surface.world_position, surface.world_normal, view.world_position, RAY_T_MAX, WORLD_CACHE_CELL_LIFETIME, &rng) * view.exposure, 1.0)); + textureStore(view_output, pixel_id, vec4(query_world_cache(surface.world_position, surface.world_normal, view.world_position, RAY_T_MAX, WORLD_CACHE_CELL_LIFETIME, &rng) * view.exposure, 1.0)); #endif } @@ -243,33 +243,3 @@ fn calculate_motion_vector(world_position: vec3, previous_world_position: v return (clip_position - previous_clip_position) * vec2(0.5, -0.5); } #endif - -fn gi_snap_to_quad_pixel(full_xy: vec2) -> vec2 { - if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index); - } else { - return full_xy; - } -} - -fn gi_reservoir_index(full_xy: vec2) -> u32 { - if bool(constants.quarter_resolution_indirect_lighting) { - return quarter_resolution_index(full_xy / 2u); - } else { - return full_xy.x + full_xy.y * u32(view.main_pass_viewport.z); - } -} - -fn quarter_resolution_dimensions() -> vec2 { - return (vec2u(view.main_pass_viewport.zw) + 1u) / 2u; -} - -fn quarter_resolution_index(quarter_xy: vec2) -> u32 { - return quarter_xy.x + quarter_xy.y * quarter_resolution_dimensions().x; -} - -fn quarter_to_full_resolution_pixel(quarter_xy: vec2, frame: u32) -> vec2 { - var rng = quarter_resolution_index(quarter_xy) * 0x9E3779B9u + frame; - let qi = u32(rand_f(&rng) * 4.0); - return min(quarter_xy * 2u + vec2(qi / 2u, qi % 2u), vec2u(view.main_pass_viewport.zw) - 1u); -} From 650a97408d6b77749dd7a11aab2f4734ea905b80 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 23:22:26 -0400 Subject: [PATCH 05/11] Misc --- crates/bevy_solari/src/realtime/specular_gi.wgsl | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index e24d90d8870f9..37063dadc0036 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -68,7 +68,10 @@ fn specular_gi(@builtin(global_invocation_id) global_id: vec3) { let brdf = evaluate_specular_brdf(wo, wi, surface.world_normal, surface.material); radiance *= brdf * view.exposure; - radiance *= select(1.0, 4.0, bool(constants.quarter_resolution_indirect_lighting)); + + if bool(constants.quarter_resolution_indirect_lighting) { + radiance *= 4.0; + } var pixel_color = textureLoad(view_output, pixel_id); pixel_color += vec4(radiance, 0.0); From f67df031aec4d708194dd0ecbe309cc6e6e8ca92 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 23:23:39 -0400 Subject: [PATCH 06/11] Misc --- 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 c8d56539526d8..87488e8ef018d 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -330,4 +330,3 @@ fn merge_reservoirs( return ReservoirMergeResult(combined_reservoir, canonical_sample_radiance, canonical_sample_wi); } } - From 545f828875abae81bdad61ac94e9d22f9bcf26b4 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 23:41:33 -0400 Subject: [PATCH 07/11] Disable PSR when using quarter-res mode to prevent artifacts --- crates/bevy_solari/src/realtime/node.rs | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 836d7d41fee8a..a6707d892f478 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -114,7 +114,9 @@ pub fn solari_lighting( #[cfg(not(all(feature = "dlss", not(feature = "force_disable_dlss"))))] let specular_gi_pipeline = pipelines.specular_gi_pipeline; #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] - let specular_gi_pipeline = if view_dlss_rr_textures.is_some() { + let specular_gi_pipeline = if view_dlss_rr_textures.is_some() + && !solari_lighting.quarter_resolution_indirect_lighting + { pipelines.specular_gi_with_psr_pipeline } else { pipelines.specular_gi_pipeline @@ -361,7 +363,9 @@ pub fn solari_lighting( let d = diagnostics.time_span(&mut pass, "solari_lighting/specular_indirect_lighting"); #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] - if let Some(bind_group_resolve_dlss_rr_textures) = &bind_group_resolve_dlss_rr_textures { + if let Some(bind_group_resolve_dlss_rr_textures) = &bind_group_resolve_dlss_rr_textures + && !solari_lighting.quarter_resolution_indirect_lighting + { pass.set_bind_group(2, bind_group_resolve_dlss_rr_textures, &[]); } pass.set_pipeline(specular_gi_pipeline); From 900bb6a720d60676f97f06f5cc3052a4c7f2fcd4 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 19 Apr 2026 23:55:41 -0400 Subject: [PATCH 08/11] Disable by default and polish example --- crates/bevy_solari/src/realtime/mod.rs | 2 +- examples/3d/solari.rs | 66 ++++++++++++++++++++++++-- 2 files changed, 63 insertions(+), 5 deletions(-) diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index 8d01bfdfd7fa7..f54c3aec8e494 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -109,7 +109,7 @@ pub struct SolariLighting { impl Default for SolariLighting { fn default() -> Self { Self { - quarter_resolution_indirect_lighting: true, + quarter_resolution_indirect_lighting: false, reset: true, // No temporal history on the first frame } } diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index c976807bd7c4b..74f8e89706c9c 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -22,8 +22,11 @@ use rand::{RngExt, SeedableRng}; use std::f32::consts::PI; #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] -use bevy::anti_alias::dlss::{ - Dlss, DlssProjectId, DlssRayReconstructionFeature, DlssRayReconstructionSupported, +use bevy::{ + anti_alias::dlss::{ + Dlss, DlssProjectId, DlssRayReconstructionFeature, DlssRayReconstructionSupported, + }, + render::camera::{MipBias, TemporalJitter}, }; /// `bevy_solari` demo. @@ -64,6 +67,10 @@ fn main() { if args.pathtracer == Some(true) { app.add_plugins(PathtracingPlugin); } else { + app.add_systems(Update, toggle_quarter_res_indirect); + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] + app.add_systems(Update, toggle_dlss_rr); + if args.many_lights != Some(true) { app.add_systems(Update, (pause_scene, toggle_lights, patrol_path)) .add_systems(PostUpdate, update_control_text); @@ -443,6 +450,41 @@ fn add_raytracing_meshes_on_scene_load( } } +fn toggle_quarter_res_indirect( + key_input: Res>, + mut solari_lighting: Single<&mut SolariLighting>, +) { + if key_input.just_pressed(KeyCode::Digit3) { + solari_lighting.quarter_resolution_indirect_lighting = + !solari_lighting.quarter_resolution_indirect_lighting; + } +} + +#[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] +fn toggle_dlss_rr( + key_input: Res>, + camera: Single<(Entity, Has>), With>, + dlss_rr_supported: Option>, + mut commands: Commands, +) { + if key_input.just_pressed(KeyCode::Digit4) && dlss_rr_supported.is_some() { + let (entity, dlss) = *camera; + if dlss { + commands + .entity(entity) + .remove::<(Dlss, TemporalJitter, MipBias)>(); + } else { + commands + .entity(entity) + .insert(Dlss:: { + perf_quality_mode: Default::default(), + reset: Default::default(), + _phantom_data: Default::default(), + }); + } + } +} + fn pause_scene(mut time: ResMut>, key_input: Res>) { if key_input.just_pressed(KeyCode::Space) { time.toggle(); @@ -530,10 +572,15 @@ fn update_control_text( robot_light_material: Option>, materials: Res>, directional_light: Query>, + solari_lighting: Single<&SolariLighting>, time: Res>, #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] dlss_rr_supported: Option< Res, >, + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] dlss_camera: Query< + Has>, + With, + >, ) { text.0.clear(); @@ -558,10 +605,21 @@ fn update_control_text( } } + if solari_lighting.quarter_resolution_indirect_lighting { + text.0 + .push_str("\n(3): Disable quarter-res indirect lighting"); + } else { + text.0 + .push_str("\n(3): Enable quarter-res indirect lighting"); + } + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] if dlss_rr_supported.is_some() { - text.0 - .push_str("\nDenoising: DLSS Ray Reconstruction enabled"); + if matches!(dlss_camera.single(), Ok(true)) { + text.0.push_str("\n(4): Disable DLSS Ray Reconstruction"); + } else { + text.0.push_str("\n(4): Enable DLSS Ray Reconstruction"); + } } else { text.0 .push_str("\nDenoising: DLSS Ray Reconstruction not supported"); From 42242803bd9d25588865ced2d66b6d76eae3d84c Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 20 Apr 2026 00:00:12 -0400 Subject: [PATCH 09/11] Add controls to many lights example --- examples/3d/solari.rs | 51 +++++++++++++++++++++++++++---------------- 1 file changed, 32 insertions(+), 19 deletions(-) diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index 74f8e89706c9c..7e2b8a345e473 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -72,10 +72,9 @@ fn main() { app.add_systems(Update, toggle_dlss_rr); if args.many_lights != Some(true) { - app.add_systems(Update, (pause_scene, toggle_lights, patrol_path)) - .add_systems(PostUpdate, update_control_text); + app.add_systems(Update, (pause_scene, toggle_lights, patrol_path)); } - app.add_systems(PostUpdate, update_performance_text); + app.add_systems(PostUpdate, (update_control_text, update_performance_text)); } app.run(); @@ -358,6 +357,17 @@ fn setup_many_lights( }); } + commands.spawn(( + ControlText, + Text::default(), + Node { + position_type: PositionType::Absolute, + bottom: px(12.0), + left: px(12.0), + ..default() + }, + )); + commands.spawn(( Node { position_type: PositionType::Absolute, @@ -574,6 +584,7 @@ fn update_control_text( directional_light: Query>, solari_lighting: Single<&SolariLighting>, time: Res>, + args: Res, #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] dlss_rr_supported: Option< Res, >, @@ -584,24 +595,26 @@ fn update_control_text( ) { text.0.clear(); - if time.is_paused() { - text.0.push_str("(Space): Resume"); - } else { - text.0.push_str("(Space): Pause"); - } - - if directional_light.single().is_ok() { - text.0.push_str("\n(1): Disable directional light"); - } else { - text.0.push_str("\n(1): Enable directional light"); - } + if args.many_lights != Some(true) { + if time.is_paused() { + text.0.push_str("(Space): Resume"); + } else { + text.0.push_str("(Space): Pause"); + } - match robot_light_material.and_then(|m| materials.get(&m.0)) { - Some(robot_light_material) if robot_light_material.emissive != LinearRgba::BLACK => { - text.0.push_str("\n(2): Disable robot emissive light"); + if directional_light.single().is_ok() { + text.0.push_str("\n(1): Disable directional light"); + } else { + text.0.push_str("\n(1): Enable directional light"); } - _ => { - text.0.push_str("\n(2): Enable robot emissive light"); + + match robot_light_material.and_then(|m| materials.get(&m.0)) { + Some(robot_light_material) if robot_light_material.emissive != LinearRgba::BLACK => { + text.0.push_str("\n(2): Disable robot emissive light"); + } + _ => { + text.0.push_str("\n(2): Enable robot emissive light"); + } } } From a435d6f49852f9cc2e0bedee49aa1e527a397de1 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 20 Apr 2026 10:55:23 -0400 Subject: [PATCH 10/11] Quarter-res direct lighting --- crates/bevy_solari/src/realtime/mod.rs | 7 ++- crates/bevy_solari/src/realtime/node.rs | 28 +++++++--- crates/bevy_solari/src/realtime/prepare.rs | 24 ++++++--- .../src/realtime/realtime_bindings.wgsl | 2 +- .../{gi_utils.wgsl => resolution_utils.wgsl} | 34 +++++++++++- .../bevy_solari/src/realtime/restir_di.wgsl | 52 ++++++++++++------- .../bevy_solari/src/realtime/restir_gi.wgsl | 3 +- .../bevy_solari/src/realtime/specular_gi.wgsl | 2 +- examples/3d/solari.rs | 25 ++++++--- 9 files changed, 134 insertions(+), 43 deletions(-) rename crates/bevy_solari/src/realtime/{gi_utils.wgsl => resolution_utils.wgsl} (65%) diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index f54c3aec8e494..bff63c508cdf7 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -36,7 +36,7 @@ impl Plugin for SolariLightingPlugin { fn build(&self, app: &mut App) { load_shader_library!(app, "gbuffer_utils.wgsl"); load_shader_library!(app, "realtime_bindings.wgsl"); - load_shader_library!(app, "gi_utils.wgsl"); + load_shader_library!(app, "resolution_utils.wgsl"); load_shader_library!(app, "presample_light_tiles.wgsl"); embedded_asset!(app, "restir_di.wgsl"); embedded_asset!(app, "restir_gi.wgsl"); @@ -94,6 +94,10 @@ impl Plugin for SolariLightingPlugin { DepthPrepassDoubleBuffer )] pub struct SolariLighting { + /// Set to true to greatly improve performance at the cost of quality. + /// + /// May force [`bevy_camera::ClearColor`] to black if not already. + pub quarter_resolution_direct_lighting: bool, /// Set to true to greatly improve performance at the cost of quality. pub quarter_resolution_indirect_lighting: bool, /// Set to true to delete the saved temporal history (past frames). @@ -109,6 +113,7 @@ pub struct SolariLighting { impl Default for SolariLighting { fn default() -> Self { Self { + quarter_resolution_direct_lighting: false, quarter_resolution_indirect_lighting: false, reset: true, // No temporal history on the first frame } diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index a6707d892f478..9e6cb80796148 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -179,7 +179,7 @@ pub fn solari_lighting( return; }; - let view_target_attachment = view_target.get_unsampled_color_attachment(); + let mut view_target_attachment = view_target.get_unsampled_color_attachment(); let s = solari_lighting_resources; let bind_group = render_device.create_bind_group( @@ -239,6 +239,7 @@ pub fn solari_lighting( let immediates = [ frame_index, solari_lighting.reset as u32, + solari_lighting.quarter_resolution_direct_lighting as u32, solari_lighting.quarter_resolution_indirect_lighting as u32, ]; let immediates = bytemuck::cast_slice(&immediates); @@ -249,7 +250,14 @@ pub fn solari_lighting( let command_encoder = ctx.command_encoder(); // Clear the view target if we're the first node to write to it - if matches!(view_target_attachment.ops.load, LoadOp::Clear(_)) { + if let LoadOp::Clear(ref mut clear_color) = view_target_attachment.ops.load { + if solari_lighting.quarter_resolution_direct_lighting { + clear_color.r = 0.0; + clear_color.g = 0.0; + clear_color.b = 0.0; + clear_color.a = 1.0; + } + command_encoder.begin_render_pass(&RenderPassDescriptor { label: Some("solari_lighting_clear"), color_attachments: &[Some(view_target_attachment)], @@ -267,8 +275,14 @@ pub fn solari_lighting( let dx = solari_lighting_resources.view_size.x.div_ceil(8); let dy = solari_lighting_resources.view_size.y.div_ceil(8); - let mut gi_dx = solari_lighting_resources.view_size.x.div_ceil(8); - let mut gi_dy = solari_lighting_resources.view_size.y.div_ceil(8); + let mut di_dx = dx; + let mut di_dy = dy; + if solari_lighting.quarter_resolution_direct_lighting { + di_dx = di_dx.div_ceil(2); + di_dy = di_dy.div_ceil(2); + } + let mut gi_dx = dx; + let mut gi_dy = dy; if solari_lighting.quarter_resolution_indirect_lighting { gi_dx = gi_dx.div_ceil(2); gi_dy = gi_dy.div_ceil(2); @@ -341,11 +355,11 @@ pub fn solari_lighting( pass.set_pipeline(di_initial_and_temporal_pipeline); pass.set_immediates(0, immediates); - pass.dispatch_workgroups(dx, dy, 1); + pass.dispatch_workgroups(di_dx, di_dy, 1); pass.set_pipeline(di_spatial_and_shade_pipeline); pass.set_immediates(0, immediates); - pass.dispatch_workgroups(dx, dy, 1); + pass.dispatch_workgroups(di_dx, di_dy, 1); d.end(&mut pass); @@ -463,7 +477,7 @@ pub fn init_solari_lighting_pipelines( pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { label: Some(label.into()), layout, - immediate_size: 12, + immediate_size: 16, shader, shader_defs, entry_point: Some(entry_point.into()), diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index b0025c7d30e64..38aa03614b5a0 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -60,6 +60,7 @@ pub struct SolariLightingResources { pub world_cache_active_cells_count: Buffer, pub world_cache_active_cells_dispatch: Buffer, pub view_size: UVec2, + pub quarter_resolution_direct_lighting: bool, pub quarter_resolution_indirect_lighting: bool, } @@ -103,12 +104,17 @@ pub fn prepare_solari_lighting_resources( view_size = *resolution_override; } - if solari_lighting_resources.map(|r| (r.view_size, r.quarter_resolution_indirect_lighting)) - == Some(( - view_size, - solari_lighting.quarter_resolution_indirect_lighting, - )) - { + if solari_lighting_resources.map(|r| { + ( + r.view_size, + r.quarter_resolution_direct_lighting, + r.quarter_resolution_indirect_lighting, + ) + }) == Some(( + view_size, + solari_lighting.quarter_resolution_direct_lighting, + solari_lighting.quarter_resolution_indirect_lighting, + )) { continue; } @@ -129,6 +135,11 @@ pub fn prepare_solari_lighting_resources( }); let di_reservoirs = |name| { + let mut view_size = view_size; + if solari_lighting.quarter_resolution_direct_lighting { + view_size.x = view_size.x.div_ceil(2); + view_size.y = view_size.y.div_ceil(2); + } render_device .create_texture(&TextureDescriptor { label: Some(name), @@ -258,6 +269,7 @@ pub fn prepare_solari_lighting_resources( world_cache_active_cells_count, world_cache_active_cells_dispatch, view_size, + quarter_resolution_direct_lighting: solari_lighting.quarter_resolution_direct_lighting, quarter_resolution_indirect_lighting: solari_lighting .quarter_resolution_indirect_lighting, }); diff --git a/crates/bevy_solari/src/realtime/realtime_bindings.wgsl b/crates/bevy_solari/src/realtime/realtime_bindings.wgsl index 624dcfb62b59d..82fd517109fe7 100644 --- a/crates/bevy_solari/src/realtime/realtime_bindings.wgsl +++ b/crates/bevy_solari/src/realtime/realtime_bindings.wgsl @@ -43,7 +43,7 @@ enable wgpu_ray_query; @group(2) @binding(3) var specular_motion_vectors: texture_storage_2d; #endif -struct PushConstants { frame_index: u32, reset: u32, quarter_resolution_indirect_lighting: u32 } +struct PushConstants { frame_index: u32, reset: u32, quarter_resolution_direct_lighting: u32, quarter_resolution_indirect_lighting: u32 } var constants: PushConstants; // Don't adjust the size of this struct without also adjusting `prepare::RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE`. diff --git a/crates/bevy_solari/src/realtime/gi_utils.wgsl b/crates/bevy_solari/src/realtime/resolution_utils.wgsl similarity index 65% rename from crates/bevy_solari/src/realtime/gi_utils.wgsl rename to crates/bevy_solari/src/realtime/resolution_utils.wgsl index f2575106e22ff..e0882508bcef5 100644 --- a/crates/bevy_solari/src/realtime/gi_utils.wgsl +++ b/crates/bevy_solari/src/realtime/resolution_utils.wgsl @@ -1,8 +1,40 @@ -#define_import_path bevy_solari::gi_utils +#define_import_path bevy_solari::resolution_utils #import bevy_pbr::utils::rand_f #import bevy_solari::realtime_bindings::{view, constants} +fn di_resolution() -> vec2 { + if bool(constants.quarter_resolution_direct_lighting) { + return quarter_resolution_dimensions(); + } else { + return vec2u(view.main_pass_viewport.zw); + } +} + +fn di_thread_to_full_resolution_pixel(thread_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_direct_lighting) { + return quarter_to_full_resolution_pixel(thread_xy, constants.frame_index); + } else { + return thread_xy; + } +} + +fn di_reservoir_pixel(full_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_direct_lighting) { + return full_xy / 2u; + } else { + return full_xy; + } +} + +fn di_snap_to_quad_pixel_previous_frame(full_xy: vec2) -> vec2 { + if bool(constants.quarter_resolution_direct_lighting) { + return quarter_to_full_resolution_pixel(full_xy / 2u, constants.frame_index - 5782582u); + } else { + return full_xy; + } +} + fn gi_resolution() -> vec2 { if bool(constants.quarter_resolution_indirect_lighting) { return quarter_resolution_dimensions(); diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 4335deefed4df..c46dc2d26ab3c 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -13,17 +13,18 @@ enable wgpu_ray_query; #import bevy_solari::scene_bindings::{light_sources, previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME, RAY_T_MIN} #import bevy_solari::specular_gi::SPECULAR_GI_FOR_DI_ROUGHNESS_THRESHOLD #import bevy_solari::realtime_bindings::{view_output, light_tile_samples, light_tile_resolved_samples, di_reservoirs_a, di_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants} +#import bevy_solari::resolution_utils::{di_resolution, di_thread_to_full_resolution_pixel, di_reservoir_pixel, di_snap_to_quad_pixel_previous_frame, quarter_resolution_dimensions, quarter_to_full_resolution_pixel} const INITIAL_SAMPLES = 8u; const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 20.0; @compute @workgroup_size(8, 8, 1) -fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) global_id: vec3) { - if any(global_id.xy >= vec2u(view.main_pass_viewport.zw)) { return; } +fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) thread_id: vec3) { + if any(thread_id.xy >= di_resolution()) { return; } - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); - var rng = pixel_index + constants.frame_index; + let global_id = vec3(di_thread_to_full_resolution_pixel(thread_id.xy), 0u); + var rng = (global_id.x + global_id.y * u32(view.main_pass_viewport.z)) + constants.frame_index; let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { @@ -42,11 +43,11 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin } @compute @workgroup_size(8, 8, 1) -fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { - if any(global_id.xy >= vec2u(view.main_pass_viewport.zw)) { return; } +fn spatial_and_shade(@builtin(global_invocation_id) thread_id: vec3) { + if any(thread_id.xy >= di_resolution()) { return; } - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); - var rng = pixel_index + constants.frame_index; + let global_id = vec3(di_thread_to_full_resolution_pixel(thread_id.xy), 0u); + var rng = (global_id.x + global_id.y * u32(view.main_pass_viewport.z)) + constants.frame_index; let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { @@ -83,6 +84,10 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { brdf += evaluate_specular_brdf(wo, merge_result.wi, surface.world_normal, surface.material); } + if bool(constants.quarter_resolution_direct_lighting) { + combined_reservoir.unbiased_contribution_weight *= 4.0; + } + var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight; pixel_color *= brdf; pixel_color += surface.material.emissive; @@ -143,12 +148,14 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return NeighborInfo(empty_reservoir(), vec3(0.0), vec3(0.0), vec3(0.0)); } - let permuted_temporal_pixel_id = permute_pixel(vec2(temporal_pixel_id_float), constants.frame_index, view.main_pass_viewport.zw); + let point_temporal_pixel_id = vec2(temporal_pixel_id_float); + let scale = select(1u, 2u, bool(constants.quarter_resolution_direct_lighting)); + let permuted_temporal_pixel_id = permute_pixel(point_temporal_pixel_id / scale, constants.frame_index, vec2(di_resolution())) * scale; var temporal = load_temporal_reservoir_inner(permuted_temporal_pixel_id, depth, world_position, world_normal); // If permuted reprojection failed (tends to happen on object edges), try point reprojection if !reservoir_valid(temporal.reservoir) { - temporal = load_temporal_reservoir_inner(vec2(temporal_pixel_id_float), depth, world_position, world_normal); + temporal = load_temporal_reservoir_inner(point_temporal_pixel_id, depth, world_position, world_normal); } // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) @@ -165,7 +172,9 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return temporal; } -fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { +fn load_temporal_reservoir_inner(temporal_pixel_id_in: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { + let temporal_pixel_id = di_snap_to_quad_pixel_previous_frame(temporal_pixel_id_in); + // Check if the pixel features have changed heavily between the current and previous frame let temporal_depth = textureLoad(previous_depth_buffer, temporal_pixel_id, 0); let temporal_surface = gpixel_resolve(textureLoad(previous_gbuffer, temporal_pixel_id, 0), temporal_depth, temporal_pixel_id, view.main_pass_viewport.zw, previous_view.world_from_clip); @@ -175,6 +184,7 @@ fn load_temporal_reservoir_inner(temporal_pixel_id: vec2, depth: f32, world } let temporal_reservoir = load_reservoir_a(temporal_pixel_id); + return NeighborInfo(temporal_reservoir, temporal_surface.world_position, temporal_surface.world_normal, temporal_diffuse_brdf); } @@ -199,9 +209,15 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< } fn get_neighbor_pixel_id(center_pixel_id: vec2, search_radius: f32, rng: ptr) -> vec2 { - var spatial_id = vec2(center_pixel_id) + sample_disk(search_radius, rng); - spatial_id = clamp(spatial_id, vec2(0.0), view.main_pass_viewport.zw - 1.0); - return vec2(spatial_id); + if bool(constants.quarter_resolution_direct_lighting) { + var spatial_id = vec2(center_pixel_id / 2u) + sample_disk(search_radius, rng) * 0.5; + spatial_id = clamp(spatial_id, vec2(0.0), vec2(quarter_resolution_dimensions()) - 1.0); + return quarter_to_full_resolution_pixel(vec2(spatial_id), constants.frame_index); + } else { + var spatial_id = vec2(center_pixel_id) + sample_disk(search_radius, rng); + spatial_id = clamp(spatial_id, vec2(0.0), view.main_pass_viewport.zw - 1.0); + return vec2(spatial_id); + } } struct NeighborInfo { @@ -235,11 +251,11 @@ fn pack_reservoir(reservoir: Reservoir) -> vec4 { } fn store_reservoir_a(pixel: vec2, reservoir: Reservoir) { - textureStore(di_reservoirs_a, pixel, pack_reservoir(reservoir)); + textureStore(di_reservoirs_a, di_reservoir_pixel(pixel), pack_reservoir(reservoir)); } fn store_reservoir_b(pixel: vec2, reservoir: Reservoir) { - textureStore(di_reservoirs_b, pixel, pack_reservoir(reservoir)); + textureStore(di_reservoirs_b, di_reservoir_pixel(pixel), pack_reservoir(reservoir)); } fn unpack_reservoir(packed: vec4) -> Reservoir { @@ -248,11 +264,11 @@ fn unpack_reservoir(packed: vec4) -> Reservoir { } fn load_reservoir_a(pixel: vec2) -> Reservoir { - return unpack_reservoir(textureLoad(di_reservoirs_a, pixel)); + return unpack_reservoir(textureLoad(di_reservoirs_a, di_reservoir_pixel(pixel))); } fn load_reservoir_b(pixel: vec2) -> Reservoir { - return unpack_reservoir(textureLoad(di_reservoirs_b, pixel)); + return unpack_reservoir(textureLoad(di_reservoirs_b, di_reservoir_pixel(pixel))); } struct ReservoirMergeResult { diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 87488e8ef018d..168bf0d1ada10 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -11,7 +11,7 @@ enable wgpu_ray_query; #import bevy_solari::world_cache::{query_world_cache, WORLD_CACHE_CELL_LIFETIME} #import bevy_solari::realtime_bindings::{view_output, gi_reservoirs_a, gi_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants, Reservoir} #import bevy_solari::specular_gi::DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD -#import bevy_solari::gi_utils::{gi_resolution, gi_thread_to_full_resolution_pixel, gi_reservoir_index, gi_snap_to_quad_pixel_previous_frame, quarter_resolution_dimensions, quarter_to_full_resolution_pixel} +#import bevy_solari::resolution_utils::{gi_resolution, gi_thread_to_full_resolution_pixel, gi_reservoir_index, gi_snap_to_quad_pixel_previous_frame, quarter_resolution_dimensions, quarter_to_full_resolution_pixel} const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 8.0; @@ -155,6 +155,7 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 fn load_temporal_reservoir_inner(temporal_pixel_id_in: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> NeighborInfo { let temporal_pixel_id = gi_snap_to_quad_pixel_previous_frame(temporal_pixel_id_in); + // Check if the pixel features have changed heavily between the current and previous frame let temporal_depth = textureLoad(previous_depth_buffer, temporal_pixel_id, 0); let temporal_surface = gpixel_resolve(textureLoad(previous_gbuffer, temporal_pixel_id, 0), temporal_depth, temporal_pixel_id, view.main_pass_viewport.zw, previous_view.world_from_clip); let temporal_diffuse_brdf = temporal_surface.material.base_color / PI; diff --git a/crates/bevy_solari/src/realtime/specular_gi.wgsl b/crates/bevy_solari/src/realtime/specular_gi.wgsl index 37063dadc0036..6d9bb8d833d35 100644 --- a/crates/bevy_solari/src/realtime/specular_gi.wgsl +++ b/crates/bevy_solari/src/realtime/specular_gi.wgsl @@ -4,7 +4,7 @@ enable wgpu_ray_query; #import bevy_pbr::pbr_functions::{calculate_tbn_mikktspace, calculate_diffuse_color, calculate_F0} #import bevy_render::maths::{orthonormalize, PI} -#import bevy_solari::gi_utils::{gi_resolution, gi_thread_to_full_resolution_pixel, gi_reservoir_index} +#import bevy_solari::resolution_utils::{gi_resolution, gi_thread_to_full_resolution_pixel, gi_reservoir_index} #import bevy_solari::brdf::{evaluate_brdf, evaluate_specular_brdf} #import bevy_solari::gbuffer_utils::{gpixel_resolve, ResolvedGPixel} #import bevy_solari::sampling::{sample_random_light, random_emissive_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, ggx_vndf_sample_invalid, power_heuristic} diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index 7e2b8a345e473..2dff366cfae33 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -67,7 +67,7 @@ fn main() { if args.pathtracer == Some(true) { app.add_plugins(PathtracingPlugin); } else { - app.add_systems(Update, toggle_quarter_res_indirect); + app.add_systems(Update, toggle_quarter_res); #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] app.add_systems(Update, toggle_dlss_rr); @@ -460,11 +460,15 @@ fn add_raytracing_meshes_on_scene_load( } } -fn toggle_quarter_res_indirect( +fn toggle_quarter_res( key_input: Res>, mut solari_lighting: Single<&mut SolariLighting>, ) { if key_input.just_pressed(KeyCode::Digit3) { + solari_lighting.quarter_resolution_direct_lighting = + !solari_lighting.quarter_resolution_direct_lighting; + } + if key_input.just_pressed(KeyCode::Digit4) { solari_lighting.quarter_resolution_indirect_lighting = !solari_lighting.quarter_resolution_indirect_lighting; } @@ -477,7 +481,7 @@ fn toggle_dlss_rr( dlss_rr_supported: Option>, mut commands: Commands, ) { - if key_input.just_pressed(KeyCode::Digit4) && dlss_rr_supported.is_some() { + if key_input.just_pressed(KeyCode::Digit5) && dlss_rr_supported.is_some() { let (entity, dlss) = *camera; if dlss { commands @@ -618,20 +622,27 @@ fn update_control_text( } } + if solari_lighting.quarter_resolution_direct_lighting { + text.0 + .push_str("\n(3): Disable quarter-res direct lighting"); + } else { + text.0.push_str("\n(3): Enable quarter-res direct lighting"); + } + if solari_lighting.quarter_resolution_indirect_lighting { text.0 - .push_str("\n(3): Disable quarter-res indirect lighting"); + .push_str("\n(4): Disable quarter-res indirect lighting"); } else { text.0 - .push_str("\n(3): Enable quarter-res indirect lighting"); + .push_str("\n(4): Enable quarter-res indirect lighting"); } #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] if dlss_rr_supported.is_some() { if matches!(dlss_camera.single(), Ok(true)) { - text.0.push_str("\n(4): Disable DLSS Ray Reconstruction"); + text.0.push_str("\n(5): Disable DLSS Ray Reconstruction"); } else { - text.0.push_str("\n(4): Enable DLSS Ray Reconstruction"); + text.0.push_str("\n(5): Enable DLSS Ray Reconstruction"); } } else { text.0 From 1065751bfb6d56ec4cde38ec5add1a678b86ba08 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 20 Apr 2026 11:29:23 -0400 Subject: [PATCH 11/11] Gate DLSS-RR diagnostics to prevent reading cached performance data when DLSS is disabled --- examples/3d/solari.rs | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index 2dff366cfae33..43a10f79da08a 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -660,6 +660,10 @@ struct PerformanceText; fn update_performance_text( mut text: Single<&mut Text, With>, diagnostics: Res, + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] dlss_camera: Query< + Has>, + With, + >, ) { text.0.clear(); @@ -692,7 +696,10 @@ fn update_performance_text( "Specular indirect", "render/solari_lighting/specular_indirect_lighting/elapsed_gpu", ); - (add_diagnostic)("DLSS-RR", "render/dlss_ray_reconstruction/elapsed_gpu"); + #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] + if matches!(dlss_camera.single(), Ok(true)) { + (add_diagnostic)("DLSS-RR", "render/dlss_ray_reconstruction/elapsed_gpu"); + } text.push_str(&format!("{:17} {total:.2} ms\n", "Total")); if let Some(world_cache_active_cells_count) = diagnostics