From 86d5d020f699d4b24670ca691979f632cea1c783 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Fri, 3 Oct 2025 00:15:51 +0100 Subject: [PATCH 1/4] first steps --- crates/bevy_anti_alias/Cargo.toml | 2 +- .../src/pathtracer/pathtracer.wgsl | 12 +- crates/bevy_solari/src/realtime/mod.rs | 3 +- crates/bevy_solari/src/realtime/node.rs | 68 +--- crates/bevy_solari/src/realtime/prepare.rs | 78 ++--- .../src/realtime/presample_light_tiles.wgsl | 55 --- .../realtime/resolve_dlss_rr_textures.wgsl | 4 +- .../bevy_solari/src/realtime/restir_di.wgsl | 329 ------------------ .../bevy_solari/src/realtime/restir_gi.wgsl | 22 +- crates/bevy_solari/src/realtime/shade_di.wgsl | 59 ++++ .../src/realtime/world_cache_compact.wgsl | 19 +- .../src/realtime/world_cache_query.wgsl | 238 +++++++++++-- .../src/realtime/world_cache_update.wgsl | 66 +--- .../src/scene/raytracing_scene_bindings.wgsl | 41 +++ crates/bevy_solari/src/scene/sampling.wgsl | 111 +++--- 15 files changed, 468 insertions(+), 639 deletions(-) delete mode 100644 crates/bevy_solari/src/realtime/presample_light_tiles.wgsl delete mode 100644 crates/bevy_solari/src/realtime/restir_di.wgsl create mode 100644 crates/bevy_solari/src/realtime/shade_di.wgsl diff --git a/crates/bevy_anti_alias/Cargo.toml b/crates/bevy_anti_alias/Cargo.toml index 8fbcf88fdc23e..60668ddb38b1f 100644 --- a/crates/bevy_anti_alias/Cargo.toml +++ b/crates/bevy_anti_alias/Cargo.toml @@ -13,7 +13,7 @@ trace = [] webgl = [] webgpu = [] smaa_luts = ["bevy_image/ktx2", "bevy_image/zstd"] -dlss = ["dep:dlss_wgpu", "dep:uuid", "bevy_render/raw_vulkan_init"] +dlss = ["dep:dlss_wgpu", "dep:uuid", "bevy_render/raw_vulkan_init", "dlss_wgpu/debug_overlay"] force_disable_dlss = ["dlss_wgpu?/mock"] [dependencies] diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index da216da959201..7ed9095d4b2c0 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -4,7 +4,7 @@ #import bevy_render::maths::PI #import bevy_render::view::View #import bevy_solari::brdf::evaluate_brdf -#import bevy_solari::sampling::{sample_random_light, random_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, balance_heuristic, power_heuristic} +#import bevy_solari::sampling::{random_light_contribution, hit_random_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, balance_heuristic, power_heuristic} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, ResolvedRayHitFull, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var accumulation_texture: texture_storage_2d; @@ -40,7 +40,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { var p_bounce = 0.0; var bounce_was_perfect_reflection = true; var previous_normal = vec3(0.0); - loop { + for (var bounces = 0; bounces < 10; bounces++) { let ray_hit = trace_ray(ray_origin, ray_direction, ray_t_min, RAY_T_MAX, RAY_FLAG_NONE); if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { let ray_hit = resolve_ray_hit_full(ray_hit); @@ -48,17 +48,19 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { var mis_weight = 1.0; if !bounce_was_perfect_reflection { - let p_light = random_light_pdf(ray_hit); - mis_weight = power_heuristic(p_bounce, p_light); + let p_light = hit_random_light_pdf(ray_hit); + // mis_weight = power_heuristic(p_bounce, p_light); + mis_weight = 0.0; } radiance += mis_weight * throughput * ray_hit.material.emissive; // Sample direct lighting, but only if the surface is not mirror-like let is_perfectly_specular = ray_hit.material.roughness < 0.0001 && ray_hit.material.metallic > 0.9999; if !is_perfectly_specular { - let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); + let direct_lighting = random_light_contribution(&rng, ray_hit.world_position, ray_hit.world_normal); let pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); + mis_weight = 1.0; let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); radiance += mis_weight * throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; } diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index 04e3528e63004..cf475785c17ba 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -32,8 +32,7 @@ pub struct SolariLightingPlugin; impl Plugin for SolariLightingPlugin { fn build(&self, app: &mut App) { - load_shader_library!(app, "presample_light_tiles.wgsl"); - embedded_asset!(app, "restir_di.wgsl"); + embedded_asset!(app, "shade_di.wgsl"); embedded_asset!(app, "restir_gi.wgsl"); load_shader_library!(app, "world_cache_query.wgsl"); embedded_asset!(app, "world_cache_compact.wgsl"); diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 9d3e8e5d0d9c9..b568bc8e9a70d 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -1,5 +1,5 @@ use super::{ - prepare::{SolariLightingResources, LIGHT_TILE_BLOCKS, WORLD_CACHE_SIZE}, + prepare::{SolariLightingResources, WORLD_CACHE_SIZE}, SolariLighting, }; use crate::scene::RaytracingSceneBindings; @@ -15,6 +15,8 @@ use bevy_ecs::{ world::{FromWorld, World}, }; use bevy_image::ToExtents; +#[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] +use bevy_render::render_resource::TextureFormat; use bevy_render::{ diagnostic::RecordDiagnostics, render_graph::{NodeRunError, RenderGraphContext, ViewNode}, @@ -24,7 +26,7 @@ use bevy_render::{ }, BindGroupEntries, BindGroupLayout, BindGroupLayoutEntries, CachedComputePipelineId, ComputePassDescriptor, ComputePipelineDescriptor, LoadOp, PipelineCache, PushConstantRange, - RenderPassDescriptor, ShaderStages, StorageTextureAccess, TextureFormat, TextureSampleType, + RenderPassDescriptor, ShaderStages, StorageTextureAccess, TextureSampleType, }, renderer::{RenderContext, RenderDevice}, view::{ViewTarget, ViewUniform, ViewUniformOffset, ViewUniforms}, @@ -50,9 +52,7 @@ pub struct SolariLightingNode { compact_world_cache_write_active_cells_pipeline: CachedComputePipelineId, sample_for_world_cache_pipeline: CachedComputePipelineId, blend_new_world_cache_samples_pipeline: CachedComputePipelineId, - presample_light_tiles_pipeline: CachedComputePipelineId, - di_initial_and_temporal_pipeline: CachedComputePipelineId, - di_spatial_and_shade_pipeline: CachedComputePipelineId, + di_shade_pipeline: CachedComputePipelineId, gi_initial_and_temporal_pipeline: CachedComputePipelineId, gi_spatial_and_shade_pipeline: CachedComputePipelineId, #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] @@ -115,9 +115,7 @@ impl ViewNode for SolariLightingNode { Some(compact_world_cache_write_active_cells_pipeline), Some(sample_for_world_cache_pipeline), Some(blend_new_world_cache_samples_pipeline), - Some(presample_light_tiles_pipeline), - Some(di_initial_and_temporal_pipeline), - Some(di_spatial_and_shade_pipeline), + Some(di_shade_pipeline), Some(gi_initial_and_temporal_pipeline), Some(gi_spatial_and_shade_pipeline), Some(scene_bindings), @@ -134,9 +132,7 @@ impl ViewNode for SolariLightingNode { .get_compute_pipeline(self.compact_world_cache_write_active_cells_pipeline), pipeline_cache.get_compute_pipeline(self.sample_for_world_cache_pipeline), pipeline_cache.get_compute_pipeline(self.blend_new_world_cache_samples_pipeline), - pipeline_cache.get_compute_pipeline(self.presample_light_tiles_pipeline), - pipeline_cache.get_compute_pipeline(self.di_initial_and_temporal_pipeline), - pipeline_cache.get_compute_pipeline(self.di_spatial_and_shade_pipeline), + pipeline_cache.get_compute_pipeline(self.di_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, @@ -164,10 +160,6 @@ impl ViewNode for SolariLightingNode { &self.bind_group_layout, &BindGroupEntries::sequential(( view_target.view, - s.light_tile_samples.as_entire_binding(), - s.light_tile_resolved_samples.as_entire_binding(), - &s.di_reservoirs_a.1, - &s.di_reservoirs_b.1, s.gi_reservoirs_a.as_entire_binding(), s.gi_reservoirs_b.as_entire_binding(), gbuffer, @@ -181,6 +173,8 @@ impl ViewNode for SolariLightingNode { s.world_cache_life.as_entire_binding(), s.world_cache_radiance.as_entire_binding(), s.world_cache_geometry_data.as_entire_binding(), + s.world_cache_light_data.as_entire_binding(), + s.world_cache_light_data_new_lights.as_entire_binding(), s.world_cache_active_cells_new_radiance.as_entire_binding(), s.world_cache_a.as_entire_binding(), s.world_cache_b.as_entire_binding(), @@ -251,13 +245,6 @@ impl ViewNode for SolariLightingNode { pass.dispatch_workgroups(dx, dy, 1); } - pass.set_pipeline(presample_light_tiles_pipeline); - pass.set_push_constants( - 0, - bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), - ); - pass.dispatch_workgroups(LIGHT_TILE_BLOCKS as u32, 1, 1); - pass.set_bind_group(2, &bind_group_world_cache_active_cells_dispatch, &[]); pass.set_pipeline(decay_world_cache_pipeline); @@ -290,14 +277,7 @@ impl ViewNode for SolariLightingNode { 0, ); - 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(dx, dy, 1); - - pass.set_pipeline(di_spatial_and_shade_pipeline); + pass.set_pipeline(di_shade_pipeline); pass.set_push_constants( 0, bytemuck::cast_slice(&[frame_index, solari_lighting.reset as u32]), @@ -366,10 +346,6 @@ impl FromWorld for SolariLightingNode { ), storage_buffer_sized(false, None), storage_buffer_sized(false, None), - texture_storage_2d(TextureFormat::Rgba32Uint, StorageTextureAccess::ReadWrite), - texture_storage_2d(TextureFormat::Rgba32Uint, StorageTextureAccess::ReadWrite), - storage_buffer_sized(false, None), - storage_buffer_sized(false, None), texture_2d(TextureSampleType::Uint), texture_depth_2d(), texture_2d(TextureSampleType::Float { filterable: true }), @@ -386,6 +362,8 @@ impl FromWorld for SolariLightingNode { storage_buffer_sized(false, None), storage_buffer_sized(false, None), storage_buffer_sized(false, None), + storage_buffer_sized(false, None), + storage_buffer_sized(false, None), ), ), ); @@ -495,24 +473,10 @@ impl FromWorld for SolariLightingNode { None, vec![], ), - presample_light_tiles_pipeline: create_pipeline( - "solari_lighting_presample_light_tiles_pipeline", - "presample_light_tiles", - load_embedded_asset!(world, "presample_light_tiles.wgsl"), - None, - vec![], - ), - di_initial_and_temporal_pipeline: create_pipeline( - "solari_lighting_di_initial_and_temporal_pipeline", - "initial_and_temporal", - load_embedded_asset!(world, "restir_di.wgsl"), - None, - vec![], - ), - di_spatial_and_shade_pipeline: create_pipeline( - "solari_lighting_di_spatial_and_shade_pipeline", - "spatial_and_shade", - load_embedded_asset!(world, "restir_di.wgsl"), + di_shade_pipeline: create_pipeline( + "solari_lighting_di_shade_pipeline", + "shade", + load_embedded_asset!(world, "shade_di.wgsl"), None, vec![], ), diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index e4d13d3c72425..b56e18a97d583 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -15,28 +15,21 @@ use bevy_ecs::{ }; use bevy_image::ToExtents; use bevy_math::UVec2; -#[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] -use bevy_render::texture::CachedTexture; use bevy_render::{ camera::ExtractedCamera, render_resource::{ Buffer, BufferDescriptor, BufferUsages, Texture, TextureDescriptor, TextureDimension, - TextureFormat, TextureUsages, TextureView, TextureViewDescriptor, + TextureUsages, TextureView, TextureViewDescriptor, }, renderer::RenderDevice, }; - -/// Size of the `LightSample` shader struct in bytes. -const LIGHT_SAMPLE_STRUCT_SIZE: u64 = 8; - -/// Size of the `ResolvedLightSamplePacked` shader struct in bytes. -const RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE: u64 = 24; +#[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))] +use bevy_render::{render_resource::TextureFormat, texture::CachedTexture}; /// Size of the GI `Reservoir` shader struct in bytes. const GI_RESERVOIR_STRUCT_SIZE: u64 = 48; - -pub const LIGHT_TILE_BLOCKS: u64 = 128; -pub const LIGHT_TILE_SAMPLES_PER_BLOCK: u64 = 1024; +/// Number of lights stored per world cache cell. +const WORLD_CACHE_CELL_LIGHT_COUNT: u64 = 8; /// Amount of entries in the world cache (must be a power of 2, and >= 2^10) pub const WORLD_CACHE_SIZE: u64 = 2u64.pow(20); @@ -44,10 +37,6 @@ pub const WORLD_CACHE_SIZE: u64 = 2u64.pow(20); /// Internal rendering resources used for Solari lighting. #[derive(Component)] pub struct SolariLightingResources { - pub light_tile_samples: Buffer, - pub light_tile_resolved_samples: Buffer, - pub di_reservoirs_a: (Texture, TextureView), - pub di_reservoirs_b: (Texture, TextureView), pub gi_reservoirs_a: Buffer, pub gi_reservoirs_b: Buffer, pub previous_gbuffer: (Texture, TextureView), @@ -56,6 +45,8 @@ pub struct SolariLightingResources { pub world_cache_life: Buffer, pub world_cache_radiance: Buffer, pub world_cache_geometry_data: Buffer, + pub world_cache_light_data: Buffer, + pub world_cache_light_data_new_lights: Buffer, pub world_cache_active_cells_new_radiance: Buffer, pub world_cache_a: Buffer, pub world_cache_b: Buffer, @@ -106,39 +97,6 @@ pub fn prepare_solari_lighting_resources( continue; } - let light_tile_samples = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_light_tile_samples"), - size: LIGHT_TILE_BLOCKS * LIGHT_TILE_SAMPLES_PER_BLOCK * LIGHT_SAMPLE_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); - - let light_tile_resolved_samples = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_light_tile_resolved_samples"), - size: LIGHT_TILE_BLOCKS - * LIGHT_TILE_SAMPLES_PER_BLOCK - * RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); - - let di_reservoirs = |name| { - let tex = render_device.create_texture(&TextureDescriptor { - label: Some(name), - size: view_size.to_extents(), - mip_level_count: 1, - sample_count: 1, - dimension: TextureDimension::D2, - format: TextureFormat::Rgba32Uint, - usage: TextureUsages::STORAGE_BINDING, - view_formats: &[], - }); - let view = tex.create_view(&TextureViewDescriptor::default()); - (tex, view) - }; - let di_reservoirs_a = di_reservoirs("solari_lighting_di_reservoirs_a"); - let di_reservoirs_b = di_reservoirs("solari_lighting_di_reservoirs_b"); - let gi_reservoirs = |name| { render_device.create_buffer(&BufferDescriptor { label: Some(name), @@ -202,6 +160,22 @@ pub fn prepare_solari_lighting_resources( mapped_at_creation: false, }); + let world_cache_light_data = render_device.create_buffer(&BufferDescriptor { + label: Some("solari_lighting_world_cache_light_data"), + size: WORLD_CACHE_SIZE + * size_of::<[u64; 1 + WORLD_CACHE_CELL_LIGHT_COUNT as usize]>() as u64, + usage: BufferUsages::STORAGE, + mapped_at_creation: false, + }); + + let world_cache_light_data_new_lights = render_device.create_buffer(&BufferDescriptor { + label: Some("solari_lighting_world_cache_light_data_new_lights"), + size: WORLD_CACHE_SIZE + * size_of::<[u64; 1 + WORLD_CACHE_CELL_LIGHT_COUNT as usize]>() as u64, + usage: BufferUsages::STORAGE, + mapped_at_creation: false, + }); + let world_cache_active_cells_new_radiance = render_device.create_buffer(&BufferDescriptor { label: Some("solari_lighting_world_cache_active_cells_new_irradiance"), @@ -245,10 +219,6 @@ pub fn prepare_solari_lighting_resources( }); commands.entity(entity).insert(SolariLightingResources { - light_tile_samples, - light_tile_resolved_samples, - di_reservoirs_a, - di_reservoirs_b, gi_reservoirs_a, gi_reservoirs_b, previous_gbuffer: (previous_gbuffer, previous_gbuffer_view), @@ -257,6 +227,8 @@ pub fn prepare_solari_lighting_resources( world_cache_life, world_cache_radiance, world_cache_geometry_data, + world_cache_light_data, + world_cache_light_data_new_lights, world_cache_active_cells_new_radiance, world_cache_a, world_cache_b, diff --git a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl deleted file mode 100644 index 75d5c0ba80710..0000000000000 --- a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl +++ /dev/null @@ -1,55 +0,0 @@ -// https://cwyman.org/papers/hpg21_rearchitectingReSTIR.pdf - -#define_import_path bevy_solari::presample_light_tiles - -#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} - -@group(1) @binding(1) var light_tile_samples: array; -@group(1) @binding(2) var light_tile_resolved_samples: array; -@group(1) @binding(12) var view: View; -struct PushConstants { frame_index: u32, reset: u32 } -var constants: PushConstants; - -@compute @workgroup_size(1024, 1, 1) -fn presample_light_tiles(@builtin(workgroup_id) workgroup_id: vec3, @builtin(local_invocation_index) sample_index: u32) { - let tile_id = workgroup_id.x; - var rng = (tile_id * 5782582u) + sample_index + constants.frame_index; - - let sample = generate_random_light_sample(&rng); - - let i = (tile_id * 1024u) + sample_index; - light_tile_samples[i] = sample.light_sample; - light_tile_resolved_samples[i] = pack_resolved_light_sample(sample.resolved_light_sample); -} - -struct ResolvedLightSamplePacked { - world_position_x: f32, - world_position_y: f32, - world_position_z: f32, - world_normal: u32, - radiance: u32, - inverse_pdf: f32, -} - -fn pack_resolved_light_sample(sample: ResolvedLightSample) -> ResolvedLightSamplePacked { - return ResolvedLightSamplePacked( - sample.world_position.x, - sample.world_position.y, - sample.world_position.z, - pack2x16unorm(octahedral_encode(sample.world_normal)), - vec3_to_rgb9e5_(sample.radiance * view.exposure), - sample.inverse_pdf * select(1.0, -1.0, sample.world_position.w == 0.0), - ); -} - -fn unpack_resolved_light_sample(packed: ResolvedLightSamplePacked, exposure: f32) -> ResolvedLightSample { - return ResolvedLightSample( - vec4(packed.world_position_x, packed.world_position_y, packed.world_position_z, select(1.0, 0.0, packed.inverse_pdf < 0.0)), - octahedral_decode(unpack2x16unorm(packed.world_normal)), - rgb9e5_to_vec3_(packed.radiance) / exposure, - abs(packed.inverse_pdf), - ); -} 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 007498968bfb0..39deb7c506e83 100644 --- a/crates/bevy_solari/src/realtime/resolve_dlss_rr_textures.wgsl +++ b/crates/bevy_solari/src/realtime/resolve_dlss_rr_textures.wgsl @@ -2,8 +2,8 @@ #import bevy_pbr::utils::octahedral_decode #import bevy_render::view::View -@group(1) @binding(7) var gbuffer: texture_2d; -@group(1) @binding(12) var view: View; +@group(1) @binding(3) var gbuffer: texture_2d; +@group(1) @binding(8) var view: View; @group(2) @binding(0) var diffuse_albedo: texture_storage_2d; @group(2) @binding(1) var specular_albedo: texture_storage_2d; diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl deleted file mode 100644 index 4512214a01541..0000000000000 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ /dev/null @@ -1,329 +0,0 @@ -// https://intro-to-restir.cwyman.org/presentations/2023ReSTIR_Course_Notes.pdf -// https://d1qx31qr3h6wln.cloudfront.net/publications/ReSTIR%20GI.pdf - -#import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance -#import bevy_pbr::pbr_deferred_types::unpack_24bit_normal -#import bevy_pbr::prepass_bindings::PreviousViewUniforms -#import bevy_pbr::rgb9e5::rgb9e5_to_vec3_ -#import bevy_pbr::utils::{rand_f, rand_range_u, octahedral_decode, sample_disk} -#import bevy_render::maths::PI -#import bevy_render::view::View -#import bevy_solari::presample_light_tiles::{ResolvedLightSamplePacked, unpack_resolved_light_sample} -#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility} -#import bevy_solari::scene_bindings::{light_sources, previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME} - -@group(1) @binding(0) var view_output: texture_storage_2d; -@group(1) @binding(1) var light_tile_samples: array; -@group(1) @binding(2) var light_tile_resolved_samples: array; -@group(1) @binding(3) var di_reservoirs_a: texture_storage_2d; -@group(1) @binding(4) var di_reservoirs_b: texture_storage_2d; -@group(1) @binding(7) var gbuffer: texture_2d; -@group(1) @binding(8) var depth_buffer: texture_depth_2d; -@group(1) @binding(9) var motion_vectors: texture_2d; -@group(1) @binding(10) var previous_gbuffer: texture_2d; -@group(1) @binding(11) var previous_depth_buffer: texture_depth_2d; -@group(1) @binding(12) var view: View; -@group(1) @binding(13) var previous_view: PreviousViewUniforms; -struct PushConstants { frame_index: u32, reset: u32 } -var constants: PushConstants; - -const INITIAL_SAMPLES = 32u; -const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; -const CONFIDENCE_WEIGHT_CAP = 20.0; - -const NULL_RESERVOIR_SAMPLE = 0xFFFFFFFFu; - -@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; } - - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); - var rng = pixel_index + constants.frame_index; - - let depth = textureLoad(depth_buffer, global_id.xy, 0); - if depth == 0.0 { - store_reservoir_b(global_id.xy, empty_reservoir()); - return; - } - let gpixel = textureLoad(gbuffer, global_id.xy, 0); - let world_position = reconstruct_world_position(global_id.xy, depth); - let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); - let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); - let diffuse_brdf = base_color / PI; - - let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, workgroup_id.xy, &rng); - let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal); - let merge_result = merge_reservoirs(initial_reservoir, temporal_reservoir, world_position, world_normal, diffuse_brdf, &rng); - - store_reservoir_b(global_id.xy, merge_result.merged_reservoir); -} - -@compute @workgroup_size(8, 8, 1) -fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { - if any(global_id.xy >= vec2u(view.main_pass_viewport.zw)) { return; } - - let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); - var rng = pixel_index + constants.frame_index; - - let depth = textureLoad(depth_buffer, global_id.xy, 0); - if depth == 0.0 { - store_reservoir_a(global_id.xy, empty_reservoir()); - return; - } - let gpixel = textureLoad(gbuffer, global_id.xy, 0); - let world_position = reconstruct_world_position(global_id.xy, depth); - let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); - let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); - let diffuse_brdf = base_color / PI; - let emissive = rgb9e5_to_vec3_(gpixel.g); - - let input_reservoir = load_reservoir_b(global_id.xy); - 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; - - store_reservoir_a(global_id.xy, combined_reservoir); - - var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight; - pixel_color *= view.exposure; - pixel_color *= diffuse_brdf; - pixel_color += emissive; - textureStore(view_output, global_id.xy, vec4(pixel_color, 1.0)); -} - -fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, workgroup_id: vec2, rng: ptr) -> Reservoir { - var workgroup_rng = (workgroup_id.x * 5782582u) + workgroup_id.y; - let light_tile_start = rand_range_u(128u, &workgroup_rng) * 1024u; - - var reservoir = empty_reservoir(); - var weight_sum = 0.0; - let mis_weight = 1.0 / f32(INITIAL_SAMPLES); - - var reservoir_target_function = 0.0; - var light_sample_world_position = vec4(0.0); - var selected_tile_sample = 0u; - for (var i = 0u; i < INITIAL_SAMPLES; i++) { - let tile_sample = light_tile_start + rand_range_u(1024u, rng); - let resolved_light_sample = unpack_resolved_light_sample(light_tile_resolved_samples[tile_sample], view.exposure); - let light_contribution = calculate_resolved_light_contribution(resolved_light_sample, world_position, world_normal); - - let target_function = luminance(light_contribution.radiance * diffuse_brdf); - let resampling_weight = mis_weight * (target_function * light_contribution.inverse_pdf); - - weight_sum += resampling_weight; - - if rand_f(rng) < resampling_weight / weight_sum { - reservoir_target_function = target_function; - light_sample_world_position = resolved_light_sample.world_position; - selected_tile_sample = tile_sample; - } - } - - if reservoir_target_function != 0.0 { - reservoir.sample = light_tile_samples[selected_tile_sample]; - } - - if reservoir_valid(reservoir) { - let inverse_target_function = select(0.0, 1.0 / reservoir_target_function, reservoir_target_function > 0.0); - reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - - reservoir.unbiased_contribution_weight *= trace_light_visibility(world_position, light_sample_world_position); - } - - reservoir.confidence_weight = 1.0; - return reservoir; -} - -fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3) -> Reservoir { - let motion_vector = textureLoad(motion_vectors, pixel_id, 0).xy; - let temporal_pixel_id_float = round(vec2(pixel_id) - (motion_vector * view.main_pass_viewport.zw)); - let temporal_pixel_id = vec2(temporal_pixel_id_float); - - // Check if the current pixel was off screen during the previous frame (current pixel is newly visible), - // or if all temporal history should assumed to be invalid - if any(temporal_pixel_id_float < vec2(0.0)) || any(temporal_pixel_id_float >= view.main_pass_viewport.zw) || bool(constants.reset) { - return empty_reservoir(); - } - - // 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); - let temporal_world_normal = octahedral_decode(unpack_24bit_normal(temporal_gpixel.a)); - if pixel_dissimilar(depth, world_position, temporal_world_position, world_normal, temporal_world_normal) { - return empty_reservoir(); - } - - var temporal_reservoir = load_reservoir_a(temporal_pixel_id); - - // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) - let previous_light_id = temporal_reservoir.sample.light_id >> 16u; - let triangle_id = temporal_reservoir.sample.light_id & 0xFFFFu; - let light_id = previous_frame_light_id_translations[previous_light_id]; - if light_id == LIGHT_NOT_PRESENT_THIS_FRAME { - return empty_reservoir(); - } - temporal_reservoir.sample.light_id = (light_id << 16u) | triangle_id; - - temporal_reservoir.confidence_weight = min(temporal_reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); - - return temporal_reservoir; -} - -fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3, world_normal: vec3, rng: ptr) -> Reservoir { - let spatial_pixel_id = get_neighbor_pixel_id(pixel_id, rng); - - let spatial_depth = textureLoad(depth_buffer, spatial_pixel_id, 0); - let spatial_gpixel = textureLoad(gbuffer, spatial_pixel_id, 0); - let spatial_world_position = reconstruct_world_position(spatial_pixel_id, spatial_depth); - let spatial_world_normal = octahedral_decode(unpack_24bit_normal(spatial_gpixel.a)); - if pixel_dissimilar(depth, world_position, spatial_world_position, world_normal, spatial_world_normal) { - return empty_reservoir(); - } - - var spatial_reservoir = load_reservoir_b(spatial_pixel_id); - - if reservoir_valid(spatial_reservoir) { - let resolved_light_sample = resolve_light_sample(spatial_reservoir.sample, light_sources[spatial_reservoir.sample.light_id >> 16u]); - spatial_reservoir.unbiased_contribution_weight *= trace_light_visibility(world_position, resolved_light_sample.world_position); - } - - return spatial_reservoir; -} - -fn get_neighbor_pixel_id(center_pixel_id: vec2, rng: ptr) -> vec2 { - var spatial_id = vec2(center_pixel_id) + sample_disk(SPATIAL_REUSE_RADIUS_PIXELS, rng); - spatial_id = clamp(spatial_id, vec2(0.0), view.main_pass_viewport.zw - 1.0); - return vec2(spatial_id); -} - -fn reconstruct_world_position(pixel_id: vec2, depth: f32) -> vec3 { - let uv = (vec2(pixel_id) + 0.5) / view.main_pass_viewport.zw; - let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); - let world_pos = view.world_from_clip * vec4(xy_ndc, depth, 1.0); - return world_pos.xyz / world_pos.w; -} - -fn reconstruct_previous_world_position(pixel_id: vec2, depth: f32) -> vec3 { - let uv = (vec2(pixel_id) + 0.5) / view.main_pass_viewport.zw; - let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); - let world_pos = previous_view.world_from_clip * vec4(xy_ndc, depth, 1.0); - return world_pos.xyz / world_pos.w; -} - -// Reject if tangent plane difference difference more than 0.3% or angle between normals more than 25 degrees -fn pixel_dissimilar(depth: f32, world_position: vec3, other_world_position: vec3, normal: vec3, other_normal: vec3) -> bool { - // https://developer.download.nvidia.com/video/gputechconf/gtc/2020/presentations/s22699-fast-denoising-with-self-stabilizing-recurrent-blurs.pdf#page=45 - let tangent_plane_distance = abs(dot(normal, other_world_position - world_position)); - let view_z = -depth_ndc_to_view_z(depth); - - return tangent_plane_distance / view_z > 0.003 || dot(normal, other_normal) < 0.906; -} - -fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { -#ifdef VIEW_PROJECTION_PERSPECTIVE - return -view.clip_from_view[3][2]() / ndc_depth; -#else ifdef VIEW_PROJECTION_ORTHOGRAPHIC - return -(view.clip_from_view[3][2] - ndc_depth) / view.clip_from_view[2][2]; -#else - let view_pos = view.view_from_clip * vec4(0.0, 0.0, ndc_depth, 1.0); - return view_pos.z / view_pos.w; -#endif -} - -struct Reservoir { - sample: LightSample, - confidence_weight: f32, - unbiased_contribution_weight: f32, -} - -fn empty_reservoir() -> Reservoir { - return Reservoir( - LightSample(NULL_RESERVOIR_SAMPLE, 0u), - 0.0, - 0.0, - ); -} - -fn reservoir_valid(reservoir: Reservoir) -> bool { - return reservoir.sample.light_id != NULL_RESERVOIR_SAMPLE; -} - -fn pack_reservoir(reservoir: Reservoir) -> vec4 { - let weights = bitcast>(vec2(reservoir.confidence_weight, reservoir.unbiased_contribution_weight)); - return vec4(reservoir.sample.light_id, reservoir.sample.seed, weights); -} - -fn store_reservoir_a(pixel: vec2, reservoir: Reservoir) { - textureStore(di_reservoirs_a, pixel, pack_reservoir(reservoir)); -} - -fn store_reservoir_b(pixel: vec2, reservoir: Reservoir) { - textureStore(di_reservoirs_b, pixel, pack_reservoir(reservoir)); -} - -fn unpack_reservoir(packed: vec4) -> Reservoir { - let weights = bitcast>(packed.zw); - return Reservoir(LightSample(packed.x, packed.y), weights.x, weights.y); -} - -fn load_reservoir_a(pixel: vec2) -> Reservoir { - return unpack_reservoir(textureLoad(di_reservoirs_a, pixel)); -} - -fn load_reservoir_b(pixel: vec2) -> Reservoir { - return unpack_reservoir(textureLoad(di_reservoirs_b, pixel)); -} - -struct ReservoirMergeResult { - merged_reservoir: Reservoir, - selected_sample_radiance: vec3, -} - -fn merge_reservoirs( - canonical_reservoir: Reservoir, - other_reservoir: Reservoir, - world_position: vec3, - world_normal: vec3, - diffuse_brdf: vec3, - rng: ptr, -) -> ReservoirMergeResult { - let mis_weight_denominator = 1.0 / (canonical_reservoir.confidence_weight + other_reservoir.confidence_weight); - - let canonical_mis_weight = canonical_reservoir.confidence_weight * mis_weight_denominator; - let canonical_target_function = reservoir_target_function(canonical_reservoir, world_position, world_normal, diffuse_brdf); - let canonical_resampling_weight = canonical_mis_weight * (canonical_target_function.a * canonical_reservoir.unbiased_contribution_weight); - - let other_mis_weight = other_reservoir.confidence_weight * mis_weight_denominator; - let other_target_function = reservoir_target_function(other_reservoir, world_position, world_normal, diffuse_brdf); - let other_resampling_weight = other_mis_weight * (other_target_function.a * other_reservoir.unbiased_contribution_weight); - - let weight_sum = canonical_resampling_weight + other_resampling_weight; - - var combined_reservoir = empty_reservoir(); - combined_reservoir.confidence_weight = canonical_reservoir.confidence_weight + other_reservoir.confidence_weight; - - if rand_f(rng) < other_resampling_weight / weight_sum { - combined_reservoir.sample = other_reservoir.sample; - - let inverse_target_function = select(0.0, 1.0 / other_target_function.a, other_target_function.a > 0.0); - combined_reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - - return ReservoirMergeResult(combined_reservoir, other_target_function.rgb); - } else { - combined_reservoir.sample = canonical_reservoir.sample; - - let inverse_target_function = select(0.0, 1.0 / canonical_target_function.a, canonical_target_function.a > 0.0); - combined_reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - - return ReservoirMergeResult(combined_reservoir, canonical_target_function.rgb); - } -} - -// TODO: Have input take ResolvedLightSample instead of reservoir.light_sample -fn reservoir_target_function(reservoir: Reservoir, world_position: vec3, world_normal: vec3, diffuse_brdf: vec3) -> vec4 { - if !reservoir_valid(reservoir) { return vec4(0.0); } - let light_contribution = resolve_and_calculate_light_contribution(reservoir.sample, world_position, world_normal).radiance; - let target_function = luminance(light_contribution * diffuse_brdf); - return vec4(light_contribution, target_function); -} diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 66ececed93a16..e3ea84b8ad5d2 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -11,15 +11,15 @@ #import bevy_solari::world_cache::query_world_cache @group(1) @binding(0) var view_output: texture_storage_2d; -@group(1) @binding(5) var gi_reservoirs_a: array; -@group(1) @binding(6) var gi_reservoirs_b: array; -@group(1) @binding(7) var gbuffer: texture_2d; -@group(1) @binding(8) var depth_buffer: texture_depth_2d; -@group(1) @binding(9) var motion_vectors: texture_2d; -@group(1) @binding(10) var previous_gbuffer: texture_2d; -@group(1) @binding(11) var previous_depth_buffer: texture_depth_2d; -@group(1) @binding(12) var view: View; -@group(1) @binding(13) var previous_view: PreviousViewUniforms; +@group(1) @binding(1) var gi_reservoirs_a: array; +@group(1) @binding(2) var gi_reservoirs_b: array; +@group(1) @binding(3) var gbuffer: texture_2d; +@group(1) @binding(4) var depth_buffer: texture_depth_2d; +@group(1) @binding(5) var motion_vectors: texture_2d; +@group(1) @binding(6) var previous_gbuffer: texture_2d; +@group(1) @binding(7) var previous_depth_buffer: texture_depth_2d; +@group(1) @binding(8) var view: View; +@group(1) @binding(9) var previous_view: PreviousViewUniforms; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; @@ -80,7 +80,7 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { var pixel_color = textureLoad(view_output, global_id.xy); pixel_color += vec4(merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight * view.exposure, 0.0); - textureStore(view_output, global_id.xy, pixel_color); + // textureStore(view_output, global_id.xy, pixel_color); #ifdef VISUALIZE_WORLD_CACHE textureStore(view_output, global_id.xy, vec4(query_world_cache(world_position, world_normal, view.world_position) * view.exposure, 1.0)); @@ -112,7 +112,7 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 reservoir.radiance = direct_lighting.radiance; reservoir.unbiased_contribution_weight = direct_lighting.inverse_pdf * uniform_hemisphere_inverse_pdf(); #else - reservoir.radiance = query_world_cache(sample_point.world_position, sample_point.geometric_world_normal, view.world_position); + reservoir.radiance = query_world_cache(sample_point.world_position, sample_point.geometric_world_normal, view.world_position).radiance; reservoir.unbiased_contribution_weight = uniform_hemisphere_inverse_pdf(); #endif diff --git a/crates/bevy_solari/src/realtime/shade_di.wgsl b/crates/bevy_solari/src/realtime/shade_di.wgsl new file mode 100644 index 0000000000000..8b73ccaf98127 --- /dev/null +++ b/crates/bevy_solari/src/realtime/shade_di.wgsl @@ -0,0 +1,59 @@ +#import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance +#import bevy_pbr::pbr_deferred_types::{unpack_unorm4x8_, unpack_24bit_normal} +#import bevy_pbr::prepass_bindings::PreviousViewUniforms +#import bevy_pbr::rgb9e5::rgb9e5_to_vec3_ +#import bevy_pbr::utils::{rand_f, rand_range_u, octahedral_decode, sample_disk} +#import bevy_render::maths::PI +#import bevy_render::view::View +#import bevy_solari::scene_bindings::ResolvedMaterial +#import bevy_solari::world_cache::{query_world_cache, evaluate_lighting_from_cache, write_world_cache_light} + +@group(1) @binding(0) var view_output: texture_storage_2d; +@group(1) @binding(3) var gbuffer: texture_2d; +@group(1) @binding(4) var depth_buffer: texture_depth_2d; +@group(1) @binding(8) var view: View; +@group(1) @binding(9) var previous_view: PreviousViewUniforms; +struct PushConstants { frame_index: u32, reset: u32 } +var constants: PushConstants; + +@compute @workgroup_size(8, 8, 1) +fn shade(@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; } + + let pixel_index = global_id.x + global_id.y * u32(view.main_pass_viewport.z); + var rng = pixel_index + constants.frame_index; + + let depth = textureLoad(depth_buffer, global_id.xy, 0); + if depth == 0.0 { + return; + } + + var material: ResolvedMaterial; + let gpixel = textureLoad(gbuffer, global_id.xy, 0); + let world_position = reconstruct_world_position(global_id.xy, depth); + let world_normal = octahedral_decode(unpack_24bit_normal(gpixel.a)); + let wo = normalize(view.world_position - world_position); + let base_rough = unpack4x8unorm(gpixel.r); + let props = unpack_unorm4x8_(gpixel.b); + + material.base_color = pow(base_rough.rgb, vec3(2.2)); + material.emissive = rgb9e5_to_vec3_(gpixel.g); + material.reflectance = vec3(props.r); + material.perceptual_roughness = base_rough.a; + material.roughness = clamp(base_rough.a * base_rough.a, 0.001, 1.0); + material.metallic = props.g; + + let cell = query_world_cache(world_position, world_normal, view.world_position); + let direct_lighting = evaluate_lighting_from_cache(&rng, cell, world_position, world_normal, wo, material, view.exposure); + write_world_cache_light(direct_lighting, world_position, world_normal, view.world_position); + + let pixel_color = direct_lighting.radiance * direct_lighting.inverse_pdf * view.exposure + material.emissive; + textureStore(view_output, global_id.xy, vec4(pixel_color, 1.0)); +} + +fn reconstruct_world_position(pixel_id: vec2, depth: f32) -> vec3 { + let uv = (vec2(pixel_id) + 0.5) / view.main_pass_viewport.zw; + let xy_ndc = (uv - vec2(0.5)) * vec2(2.0, -2.0); + let world_pos = view.world_from_clip * vec4(xy_ndc, depth, 1.0); + return world_pos.xyz / world_pos.w; +} diff --git a/crates/bevy_solari/src/realtime/world_cache_compact.wgsl b/crates/bevy_solari/src/realtime/world_cache_compact.wgsl index 71585223a5e44..36aeef6d321aa 100644 --- a/crates/bevy_solari/src/realtime/world_cache_compact.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_compact.wgsl @@ -1,4 +1,4 @@ -#import bevy_solari::world_cache::{WORLD_CACHE_EMPTY_CELL, world_cache_life, world_cache_checksums, world_cache_radiance, world_cache_a, world_cache_b, world_cache_active_cell_indices, world_cache_active_cells_count} +#import bevy_solari::world_cache::{WORLD_CACHE_EMPTY_CELL, world_cache_life, world_cache_checksums, world_cache_radiance, world_cache_a, world_cache_b, world_cache_active_cell_indices, world_cache_active_cells_count, world_cache_light_data, world_cache_light_data_new_lights, WORLD_CACHE_CELL_LIGHT_COUNT} @group(2) @binding(0) var world_cache_active_cells_dispatch: vec3; @@ -15,6 +15,23 @@ fn decay_world_cache(@builtin(global_invocation_id) global_id: vec3) { if life == 0u { world_cache_checksums[global_id.x] = WORLD_CACHE_EMPTY_CELL; world_cache_radiance[global_id.x] = vec4(0.0); + } else { + var count = min(WORLD_CACHE_CELL_LIGHT_COUNT, atomicLoad(&world_cache_light_data_new_lights[global_id.x].visible_light_count)); + atomicStore(&world_cache_light_data_new_lights[global_id.x].visible_light_count, 0u); + for (var i = 0u; i < count; i++) { + let data = atomicLoad(&world_cache_light_data_new_lights[global_id.x].visible_lights[i]); + atomicStore(&world_cache_light_data_new_lights[global_id.x].visible_lights[i], 0); + if data == 0 { + count = i; + break; + } + + let light = u32(data & 0xffffffff); + let weight = bitcast(u32(data >> 32u)); + world_cache_light_data[global_id.x].visible_lights[i].light = light; + world_cache_light_data[global_id.x].visible_lights[i].weight = weight; + } + world_cache_light_data[global_id.x].visible_light_count = count; } } } diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index 46f0fe920f2ce..eee9f330349f8 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -1,11 +1,21 @@ #define_import_path bevy_solari::world_cache +#import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance +#import bevy_pbr::utils::rand_f +#import bevy_solari::brdf::evaluate_brdf +#import bevy_solari::sampling::{light_contribution_no_trace, select_random_light, select_random_light_inverse_pdf, trace_light_visibility} +#import bevy_solari::scene_bindings::ResolvedMaterial + /// How responsive the world cache is to changes in lighting (higher is less responsive, lower is more responsive) const WORLD_CACHE_MAX_TEMPORAL_SAMPLES: f32 = 20.0; /// Maximum amount of frames a cell can live for without being queried const WORLD_CACHE_CELL_LIFETIME: u32 = 30u; /// Maximum amount of attempts to find a cache entry after a hash collision const WORLD_CACHE_MAX_SEARCH_STEPS: u32 = 3u; +/// Maximum lights stored in each cache cell +const WORLD_CACHE_CELL_LIGHT_COUNT: u32 = 8u; +/// Lights searched that aren't in the cell +const WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT: u32 = 2u; /// The size of a cache cell at the lowest LOD in meters const WORLD_CACHE_POSITION_BASE_CELL_SIZE: f32 = 0.25; @@ -22,25 +32,51 @@ struct WorldCacheGeometryData { padding_b: u32 } -@group(1) @binding(14) var world_cache_checksums: array, #{WORLD_CACHE_SIZE}>; +struct WorldCacheSingleLightData { + light: u32, + weight: f32, +} + +// The size of these structs should match `WORLD_CACHE_CELL_LIGHT_COUNT` in `realtime/prepare.rs`! +struct WorldCacheLightDataRead { + visible_light_count: u32, + padding: u32, + visible_lights: array, +} + +struct WorldCacheLightDataWrite { + visible_light_count: atomic, + padding: u32, + visible_lights: array, WORLD_CACHE_CELL_LIGHT_COUNT>, +} + +struct WorldCacheData { + radiance: vec3, + visible_light_count: u32, + visible_lights: array, +} + +@group(1) @binding(10) var world_cache_checksums: array, #{WORLD_CACHE_SIZE}>; #ifdef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER -@group(1) @binding(15) var world_cache_life: array; +@group(1) @binding(11) var world_cache_life: array; #else -@group(1) @binding(15) var world_cache_life: array, #{WORLD_CACHE_SIZE}>; +@group(1) @binding(11) var world_cache_life: array, #{WORLD_CACHE_SIZE}>; #endif -@group(1) @binding(16) var world_cache_radiance: array, #{WORLD_CACHE_SIZE}>; -@group(1) @binding(17) var world_cache_geometry_data: array; -@group(1) @binding(18) var world_cache_active_cells_new_radiance: array, #{WORLD_CACHE_SIZE}>; -@group(1) @binding(19) var world_cache_a: array; -@group(1) @binding(20) var world_cache_b: array; -@group(1) @binding(21) var world_cache_active_cell_indices: array; -@group(1) @binding(22) var world_cache_active_cells_count: u32; +@group(1) @binding(12) var world_cache_radiance: array, #{WORLD_CACHE_SIZE}>; +@group(1) @binding(13) var world_cache_geometry_data: array; +@group(1) @binding(14) var world_cache_light_data: array; +@group(1) @binding(15) var world_cache_light_data_new_lights: array; +@group(1) @binding(16) var world_cache_active_cells_new_radiance: array, #{WORLD_CACHE_SIZE}>; +@group(1) @binding(17) var world_cache_a: array; +@group(1) @binding(18) var world_cache_b: array; +@group(1) @binding(19) var world_cache_active_cell_indices: array; +@group(1) @binding(20) var world_cache_active_cells_count: u32; #ifndef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER -fn query_world_cache(world_position: vec3, world_normal: vec3, view_position: vec3) -> vec3 { +fn query_world_cache(world_position: vec3, world_normal: vec3, view_position: vec3) -> WorldCacheData { let cell_size = get_cell_size(world_position, view_position); - let world_position_quantized = bitcast>(quantize_position(world_position, cell_size)); - let world_normal_quantized = bitcast>(quantize_normal(world_normal)); + let world_position_quantized = quantize_position(world_position, cell_size); + let world_normal_quantized = quantize_normal(world_normal); var key = compute_key(world_position_quantized, world_normal_quantized); let checksum = compute_checksum(world_position_quantized, world_normal_quantized); @@ -49,35 +85,197 @@ fn query_world_cache(world_position: vec3, world_normal: vec3, view_po if existing_checksum == checksum { // Cache entry already exists - get radiance and reset cell lifetime atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); - return world_cache_radiance[key].rgb; + let radiance = world_cache_radiance[key].rgb; + let data = world_cache_light_data[key]; + return WorldCacheData(radiance, data.visible_light_count, data.visible_lights); } else if existing_checksum == WORLD_CACHE_EMPTY_CELL { // Cell is empty - reset cell lifetime so that it starts getting updated next frame atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); world_cache_geometry_data[key].world_position = world_position; world_cache_geometry_data[key].world_normal = world_normal; - return vec3(0.0); + world_cache_light_data[key].visible_light_count = 0u; + return WorldCacheData(vec3(0.0), 0u, array()); } else { // Collision - jump to another entry key = wrap_key(pcg_hash(key)); } } - return vec3(0.0); + return WorldCacheData(vec3(0.0), 0u, array()); } #endif +fn write_world_cache_light(cell: EvaluatedLighting, world_position: vec3, world_normal: vec3, view_position: vec3) { + let cell_size = get_cell_size(world_position, view_position); + let world_position_quantized = quantize_position(world_position, cell_size); + let world_normal_quantized = quantize_normal(world_normal); + var key = compute_key(world_position_quantized, world_normal_quantized); + let checksum = compute_checksum(world_position_quantized, world_normal_quantized); + + for (var i = 0u; i < WORLD_CACHE_MAX_SEARCH_STEPS; i++) { + // Don't need a CAS because we know the cell is alive + let existing_checksum = atomicLoad(&world_cache_checksums[key]); + if existing_checksum == checksum { + let index = atomicAdd(&world_cache_light_data_new_lights[key].visible_light_count, 1u) & (WORLD_CACHE_CELL_LIGHT_COUNT - 1u); + let packed = (u64(bitcast(cell.data.weight)) << 32u) | u64(cell.data.light); + atomicMax(&world_cache_light_data_new_lights[key].visible_lights[index], packed); + } else { + // Collision - jump to another entry + key = wrap_key(pcg_hash(key)); + } + } +} + +struct EvaluatedLighting { + radiance: vec3, + inverse_pdf: f32, + data: WorldCacheSingleLightData, +} + +fn evaluate_lighting_from_cache( + rng: ptr, + cell: WorldCacheData, + world_position: vec3, + world_normal: vec3, + wo: vec3, + material: ResolvedMaterial, + exposure: f32, +) -> EvaluatedLighting { + let cell_selected_light = select_light_from_cache_cell(rng, cell, world_position, world_normal, wo, material); + let cell_selected_weight = cell_selected_light.weight + log2(exposure); + let cell_confidence = smoothstep(0.1, 0.3, cell_selected_weight); + + // Sample more random lights if our cell has bad lights + let random_sample_count = u32(round(mix(f32(WORLD_CACHE_CELL_LIGHT_COUNT), f32(WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT), cell_confidence))); + let random_selected_light = select_light_randomly(rng, world_position, world_normal, wo, material, random_sample_count); + + let p_cell_selection = select(p_wrs(cell_selected_light), 0.0, cell_selected_light.weight_sum < 0.0001); + let p_random_selection = select(p_wrs(random_selected_light), 0.0, random_selected_light.weight_sum < 0.0001); + let p_random_selection_clamped = min(mix(1.0, 0.25 * p_cell_selection, cell_confidence), p_random_selection); + + let weight_sum = p_cell_selection + p_random_selection_clamped; + if weight_sum < 0.0001 { + return EvaluatedLighting(vec3(0.0), 0.0, WorldCacheSingleLightData(0, 0.0)); + } + + let p_should_choose_cell = p_cell_selection / weight_sum; + let p_should_choose_random = 1.0 - p_should_choose_cell; + var sel: u32; + var sel_weight: f32; + var pdf: f32; + if rand_f(rng) < p_should_choose_cell { + sel = cell_selected_light.light; + sel_weight = cell_selected_light.weight; + pdf = p_should_choose_cell * p_cell_selection; + } else { + sel = random_selected_light.light; + sel_weight = random_selected_light.weight; + pdf = p_should_choose_random * p_random_selection; + } + + // TODO: reuse the eval that we did for light selection somehow + let direct_lighting = light_contribution_no_trace(rng, sel, world_position, world_normal); + let brdf = evaluate_brdf(world_normal, wo, direct_lighting.wi, material); + let visibility = trace_light_visibility(world_position, direct_lighting.world_position); + let radiance = direct_lighting.radiance * brdf * visibility; + let inverse_pdf = direct_lighting.inverse_pdf / pdf; + return EvaluatedLighting(radiance, inverse_pdf, WorldCacheSingleLightData(sel, sel_weight * visibility)); +} + +struct SelectedLight { + light: u32, + weight: f32, + weight_sum: f32, + base_pdf: f32, +} + +fn p_wrs(selection: SelectedLight) -> f32 { + return (selection.weight / selection.weight_sum) * selection.base_pdf; +} + +fn select_light_from_cache_cell( + rng: ptr, + cell: WorldCacheData, + world_position: vec3, + world_normal: vec3, + wo: vec3, + material: ResolvedMaterial +) -> SelectedLight { + var p = rand_f(rng); + + var selected = 0u; + var selected_weight = 0.0; + var weight_sum = 0.0; + // WRS to select the light based on unshadowed contribution + for (var i = 0u; i < cell.visible_light_count; i++) { + let light_id = cell.visible_lights[i].light; + let direct_lighting = light_contribution_no_trace(rng, light_id, world_position, world_normal); + let brdf = evaluate_brdf(world_normal, wo, direct_lighting.wi, material); + // Weight by inverse_pdf to bias towards larger triangles + let radiance = direct_lighting.radiance * direct_lighting.inverse_pdf * brdf; + + let weight = log2(luminance(radiance) + 1.0); + weight_sum += weight; + + let prob = weight / weight_sum; + if p < prob { + selected = light_id; + selected_weight = weight; + p /= prob; + } else { + p = (p - prob) / (1.0 - prob); + } + } + return SelectedLight(selected, selected_weight, weight_sum, 1.0); +} + +fn select_light_randomly( + rng: ptr, + world_position: vec3, + world_normal: vec3, + wo: vec3, + material: ResolvedMaterial, + samples: u32, +) -> SelectedLight { + var p = rand_f(rng); + + var selected = 0u; + var selected_weight = 0.0; + var weight_sum = 0.0; + for (var i = 0u; i < samples; i++) { + let light_id = select_random_light(rng); + let direct_lighting = light_contribution_no_trace(rng, light_id, world_position, world_normal); + let brdf = evaluate_brdf(world_normal, wo, direct_lighting.wi, material); + let radiance = direct_lighting.radiance * direct_lighting.inverse_pdf * brdf; + + let weight = log2(luminance(radiance) + 1.0); + weight_sum += weight; + + let prob = weight / weight_sum; + if p < prob { + selected = light_id; + selected_weight = weight; + p /= prob; + } else { + p = (p - prob) / (1.0 - prob); + } + } + let base_pdf = f32(samples) / select_random_light_inverse_pdf(selected); + return SelectedLight(selected, selected_weight, weight_sum, base_pdf); +} + fn get_cell_size(world_position: vec3, view_position: vec3) -> f32 { let camera_distance = distance(view_position, world_position) / WORLD_CACHE_POSITION_LOD_SCALE; let lod = exp2(floor(log2(1.0 + camera_distance))); return WORLD_CACHE_POSITION_BASE_CELL_SIZE * lod; } -fn quantize_position(world_position: vec3, quantization_factor: f32) -> vec3 { - return floor(world_position / quantization_factor + 0.0001); +fn quantize_position(world_position: vec3, quantization_factor: f32) -> vec3 { + return vec3(vec3(floor(world_position / quantization_factor + 0.0001))); } -fn quantize_normal(world_normal: vec3) -> vec3 { - return floor(world_normal + 0.0001); +fn quantize_normal(world_normal: vec3) -> vec3 { + return vec3(vec3(floor(world_normal + 0.0001))); } // TODO: Clustering diff --git a/crates/bevy_solari/src/realtime/world_cache_update.wgsl b/crates/bevy_solari/src/realtime/world_cache_update.wgsl index 43b11f279f51b..178db2cd283d4 100644 --- a/crates/bevy_solari/src/realtime/world_cache_update.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_update.wgsl @@ -1,26 +1,24 @@ #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::sampling::{calculate_resolved_light_contribution, trace_light_visibility} -#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} +#import bevy_solari::scene_bindings::{ResolvedMaterial, trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} #import bevy_solari::world_cache::{ WORLD_CACHE_MAX_TEMPORAL_SAMPLES, - query_world_cache, world_cache_active_cells_count, world_cache_active_cell_indices, world_cache_geometry_data, world_cache_radiance, world_cache_active_cells_new_radiance, + query_world_cache, + evaluate_lighting_from_cache, + write_world_cache_light, } -@group(1) @binding(2) var light_tile_resolved_samples: array; -@group(1) @binding(12) var view: View; +@group(1) @binding(8) var view: View; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; -const DIRECT_LIGHT_SAMPLE_COUNT: u32 = 32u; - @compute @workgroup_size(1024, 1, 1) fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) active_cell_id: vec3) { if active_cell_id.x < world_cache_active_cells_count { @@ -29,15 +27,26 @@ fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3, @builtin(glob var rng = cell_index + constants.frame_index; // TODO: Initialize newly active cells with data from an adjacent LOD - - var new_radiance = sample_random_light_ris(geometry_data.world_position, geometry_data.world_normal, workgroup_id.xy, &rng); + + var material: ResolvedMaterial; + material.base_color = vec3(1.0); + material.emissive = vec3(0.0); + material.reflectance = vec3(0.0); + material.perceptual_roughness = 1.0; + material.roughness = 1.0; + material.metallic = 0.0; + + let cell = query_world_cache(geometry_data.world_position, geometry_data.world_normal, view.world_position); + let direct_lighting = evaluate_lighting_from_cache(&rng, cell, geometry_data.world_position, geometry_data.world_normal, geometry_data.world_normal, material, view.exposure); + write_world_cache_light(direct_lighting, geometry_data.world_position, geometry_data.world_normal, view.world_position); + var new_radiance = direct_lighting.radiance * direct_lighting.inverse_pdf; #ifndef NO_MULTIBOUNCE let ray_direction = sample_cosine_hemisphere(geometry_data.world_normal, &rng); let ray_hit = trace_ray(geometry_data.world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { let ray_hit = resolve_ray_hit_full(ray_hit); - new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position); + new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position).radiance; } #endif @@ -59,40 +68,3 @@ fn blend_new_samples(@builtin(global_invocation_id) active_cell_id: vec3) { world_cache_radiance[cell_index] = vec4(blended_radiance, sample_count); } } - -fn sample_random_light_ris(world_position: vec3, world_normal: vec3, workgroup_id: vec2, rng: ptr) -> vec3 { - var workgroup_rng = (workgroup_id.x * 5782582u) + workgroup_id.y; - let light_tile_start = rand_range_u(128u, &workgroup_rng) * 1024u; - - var weight_sum = 0.0; - var selected_sample_radiance = vec3(0.0); - var selected_sample_target_function = 0.0; - var selected_sample_world_position = vec4(0.0); - let mis_weight = 1.0 / f32(DIRECT_LIGHT_SAMPLE_COUNT); - for (var i = 0u; i < DIRECT_LIGHT_SAMPLE_COUNT; i++) { - let tile_sample = light_tile_start + rand_range_u(1024u, rng); - let resolved_light_sample = unpack_resolved_light_sample(light_tile_resolved_samples[tile_sample], view.exposure); - let light_contribution = calculate_resolved_light_contribution(resolved_light_sample, world_position, world_normal); - - let target_function = luminance(light_contribution.radiance); - let resampling_weight = mis_weight * (target_function * light_contribution.inverse_pdf); - - weight_sum += resampling_weight; - - if rand_f(rng) < resampling_weight / weight_sum { - selected_sample_radiance = light_contribution.radiance; - selected_sample_target_function = target_function; - selected_sample_world_position = resolved_light_sample.world_position; - } - } - - var unbiased_contribution_weight = 0.0; - if all(selected_sample_radiance != vec3(0.0)) { - let inverse_target_function = select(0.0, 1.0 / selected_sample_target_function, selected_sample_target_function > 0.0); - unbiased_contribution_weight = weight_sum * inverse_target_function; - - unbiased_contribution_weight *= trace_light_visibility(world_position, selected_sample_world_position); - } - - return selected_sample_radiance * unbiased_contribution_weight; -} diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index 9e346bbf2d783..62f3ae4607b77 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -120,6 +120,13 @@ struct ResolvedRayHitFull { material: ResolvedMaterial, } +struct ResolvedRayHitEmissive { + world_position: vec3, + world_normal: vec3, + triangle_area: f32, + emissive: vec3, +} + fn resolve_material(material: Material, uv: vec2) -> ResolvedMaterial { var m: ResolvedMaterial; @@ -213,3 +220,37 @@ fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, world_tangent, uv, triangle_area, instance_geometry_ids.triangle_count, resolved_material); } + +fn resolve_triangle_data_emissive(instance_id: u32, triangle_id: u32, barycentrics: vec3) -> ResolvedRayHitEmissive { + let material_id = material_ids[instance_id]; + let material = materials[material_id]; + + let instance_geometry_ids = geometry_ids[instance_id]; + let vertices = load_vertices(instance_geometry_ids, triangle_id); + let transform = transforms[instance_id]; + let world_vertices = transform_positions(transform, vertices); + + let world_position = mat3x3(world_vertices[0], world_vertices[1], world_vertices[2]) * barycentrics; + + let uv = mat3x2(vertices[0].uv, vertices[1].uv, vertices[2].uv) * barycentrics; + + let local_tangent = mat3x3(vertices[0].tangent.xyz, vertices[1].tangent.xyz, vertices[2].tangent.xyz) * barycentrics; + let world_tangent = vec4( + normalize(mat3x3(transform[0].xyz, transform[1].xyz, transform[2].xyz) * local_tangent), + vertices[0].tangent.w, + ); + + let local_normal = mat3x3(vertices[0].normal, vertices[1].normal, vertices[2].normal) * barycentrics; // TODO: Use barycentric lerp, ray_hit.object_to_world, cross product geo normal + let world_normal = normalize(mat3x3(transform[0].xyz, transform[1].xyz, transform[2].xyz) * local_normal); + + let triangle_edge0 = world_vertices[0] - world_vertices[1]; + let triangle_edge1 = world_vertices[0] - world_vertices[2]; + let triangle_area = length(cross(triangle_edge0, triangle_edge1)) / 2.0; + + var emissive = material.emissive.rgb; + if material.emissive_texture_id != TEXTURE_MAP_NONE { + emissive *= sample_texture(material.emissive_texture_id, uv); + } + + return ResolvedRayHitEmissive(world_position, world_normal, triangle_area, emissive); +} diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 8385f9b3999e6..cd343b6bd7cee 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -3,7 +3,7 @@ #import bevy_pbr::lighting::D_GGX #import bevy_pbr::utils::{rand_f, 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} +#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, resolve_triangle_data_emissive, ResolvedRayHitFull} fn power_heuristic(f: f32, g: f32) -> f32 { return f * f / (f * f + g * g); @@ -55,10 +55,6 @@ fn ggx_vndf_pdf(wi_tangent: vec3, wo_tangent: vec3, roughness: f32) -> struct LightSample { light_id: u32, - seed: u32, -} - -struct ResolvedLightSample { world_position: vec4, world_normal: vec3, radiance: vec3, @@ -66,39 +62,34 @@ struct ResolvedLightSample { } struct LightContribution { + light_id: u32, + world_position: vec4, radiance: vec3, inverse_pdf: f32, wi: vec3, } -struct LightContributionNoPdf { - radiance: vec3, - wi: vec3, -} - -struct GenerateRandomLightSampleResult { - light_sample: LightSample, - resolved_light_sample: ResolvedLightSample, +fn random_light_contribution(rng: ptr, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { + let light_id = select_random_light(rng); + var light_contribution = light_contribution_no_trace(rng, light_id, ray_origin, origin_world_normal); + light_contribution.radiance *= trace_light_visibility(ray_origin, light_contribution.world_position); + light_contribution.inverse_pdf *= select_random_light_inverse_pdf(light_id); + return light_contribution; } -fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rng: ptr) -> LightContribution { - let sample = generate_random_light_sample(rng); - var light_contribution = calculate_resolved_light_contribution(sample.resolved_light_sample, ray_origin, origin_world_normal); - light_contribution.radiance *= trace_light_visibility(ray_origin, sample.resolved_light_sample.world_position); - return light_contribution; +fn light_contribution_no_trace(rng: ptr, light_id: u32, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { + return calculate_light_contribution(sample_light(rng, light_id), ray_origin, origin_world_normal); } -fn random_light_pdf(hit: ResolvedRayHitFull) -> f32 { +fn hit_random_light_pdf(hit: ResolvedRayHitFull) -> f32 { let light_count = arrayLength(&light_sources); - let p_light = 1.0 / f32(light_count); - return p_light / (hit.triangle_area * f32(hit.triangle_count)); + return 1.0 / (hit.triangle_area * f32(hit.triangle_count) * f32(light_count)); } -fn generate_random_light_sample(rng: ptr) -> GenerateRandomLightSampleResult { +fn select_random_light(rng: ptr) -> u32 { let light_count = arrayLength(&light_sources); - let light_id = rand_range_u(light_count, rng); - - let light_source = light_sources[light_id]; + let light_index = rand_range_u(light_count, rng); + let light_source = light_sources[light_index]; var triangle_id = 0u; if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { @@ -106,24 +97,28 @@ fn generate_random_light_sample(rng: ptr) -> GenerateRandomLightS triangle_id = rand_range_u(triangle_count, rng); } - let seed = rand_u(rng); - let light_sample = LightSample((light_id << 16u) | triangle_id, seed); - - var resolved_light_sample = resolve_light_sample(light_sample, light_source); - resolved_light_sample.inverse_pdf *= f32(light_count); + return (light_index << 16u) | triangle_id; +} - return GenerateRandomLightSampleResult(light_sample, resolved_light_sample); +fn select_random_light_inverse_pdf(light_id: u32) -> f32 { + let light_count = arrayLength(&light_sources); + let light_source = light_sources[light_id >> 16u]; + var triangle_count = 1u; + if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { + triangle_count = light_source.kind >> 1u; + } + return f32(light_count) * f32(triangle_count); } -fn resolve_light_sample(light_sample: LightSample, light_source: LightSource) -> ResolvedLightSample { +fn sample_light(rng: ptr, light_id: u32) -> LightSample { + let light_source = light_sources[light_id >> 16u]; if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { let directional_light = directional_lights[light_source.id]; #ifdef DIRECTIONAL_LIGHT_SOFT_SHADOWS // Sample a random direction within a cone whose base is the sun approximated as a disk // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec30%3A305 - var rng = light_sample.seed; - let random = rand_vec2f(&rng); + let random = rand_vec2f(rng); let cos_theta = (1.0 - random.x) + random.x * directional_light.cos_theta_max; let sin_theta = sqrt(1.0 - cos_theta * cos_theta); let phi = random.y * PI_2; @@ -137,45 +132,47 @@ fn resolve_light_sample(light_sample: LightSample, light_source: LightSource) -> let direction_to_light = directional_light.direction_to_light; #endif - return ResolvedLightSample( + return LightSample( + light_id, vec4(direction_to_light, 0.0), -direction_to_light, directional_light.luminance, directional_light.inverse_pdf, ); } else { - let triangle_count = light_source.kind >> 1u; - let triangle_id = light_sample.light_id & 0xFFFFu; - let barycentrics = triangle_barycentrics(light_sample.seed); - let triangle_data = resolve_triangle_data_full(light_source.id, triangle_id, barycentrics); + let triangle_id = light_id & 0xFFFFu; + let barycentrics = sample_triangle(rng); + let triangle_data = resolve_triangle_data_emissive(light_source.id, triangle_id, barycentrics); - return ResolvedLightSample( + return LightSample( + light_id, vec4(triangle_data.world_position, 1.0), triangle_data.world_normal, - triangle_data.material.emissive.rgb, - f32(triangle_count) * triangle_data.triangle_area, + triangle_data.emissive, + triangle_data.triangle_area, ); } } -fn calculate_resolved_light_contribution(resolved_light_sample: ResolvedLightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { - let ray = resolved_light_sample.world_position.xyz - (resolved_light_sample.world_position.w * ray_origin); +// https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec22%3A297 +fn sample_triangle(rng: ptr) -> vec3 { + var barycentrics = rand_vec2f(rng); + if barycentrics.x + barycentrics.y > 1.0 { barycentrics = 1.0 - barycentrics; } + return vec3(1.0 - barycentrics.x - barycentrics.y, barycentrics); +} + +fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { + let ray = light_sample.world_position.xyz - (light_sample.world_position.w * ray_origin); let light_distance = length(ray); let wi = ray / light_distance; let cos_theta_origin = saturate(dot(wi, origin_world_normal)); - let cos_theta_light = saturate(dot(-wi, resolved_light_sample.world_normal)); + let cos_theta_light = saturate(dot(-wi, light_sample.world_normal)); let light_distance_squared = light_distance * light_distance; - let radiance = resolved_light_sample.radiance * cos_theta_origin * (cos_theta_light / light_distance_squared); + let radiance = light_sample.radiance * cos_theta_origin * (cos_theta_light / light_distance_squared); - return LightContribution(radiance, resolved_light_sample.inverse_pdf, wi); -} - -fn resolve_and_calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContributionNoPdf { - let resolved_light_sample = resolve_light_sample(light_sample, light_sources[light_sample.light_id >> 16u]); - let light_contribution = calculate_resolved_light_contribution(resolved_light_sample, ray_origin, origin_world_normal); - return LightContributionNoPdf(light_contribution.radiance, light_contribution.wi); + return LightContribution(light_sample.light_id, light_sample.world_position, radiance, light_sample.inverse_pdf, wi); } fn trace_light_visibility(ray_origin: vec3, light_sample_world_position: vec4) -> f32 { @@ -206,11 +203,3 @@ fn trace_point_visibility(ray_origin: vec3, point: vec3) -> f32 { let ray_hit = trace_ray(ray_origin, ray_direction, RAY_T_MIN, ray_t_max, RAY_FLAG_TERMINATE_ON_FIRST_HIT); return f32(ray_hit.kind == RAY_QUERY_INTERSECTION_NONE); } - -// https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec22%3A297 -fn triangle_barycentrics(seed: u32) -> vec3 { - var rng = seed; - var barycentrics = rand_vec2f(&rng); - if barycentrics.x + barycentrics.y > 1.0 { barycentrics = 1.0 - barycentrics; } - return vec3(1.0 - barycentrics.x - barycentrics.y, barycentrics); -} From 0b099aa5509f5cb24e69bcf47127f349089f9767 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Fri, 3 Oct 2025 17:46:08 +0100 Subject: [PATCH 2/4] improve stability and fix artifacting --- .../src/pathtracer/pathtracer.wgsl | 4 +- .../bevy_solari/src/realtime/restir_gi.wgsl | 6 +- crates/bevy_solari/src/realtime/shade_di.wgsl | 6 +- .../src/realtime/world_cache_compact.wgsl | 43 ++++++-- .../src/realtime/world_cache_query.wgsl | 101 +++++++++++++----- .../src/realtime/world_cache_update.wgsl | 9 +- 6 files changed, 124 insertions(+), 45 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 7ed9095d4b2c0..84bc2e324e584 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -49,8 +49,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { var mis_weight = 1.0; if !bounce_was_perfect_reflection { let p_light = hit_random_light_pdf(ray_hit); - // mis_weight = power_heuristic(p_bounce, p_light); - mis_weight = 0.0; + mis_weight = power_heuristic(p_bounce, p_light); } radiance += mis_weight * throughput * ray_hit.material.emissive; @@ -60,7 +59,6 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { let direct_lighting = random_light_contribution(&rng, ray_hit.world_position, ray_hit.world_normal); let pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); - mis_weight = 1.0; let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); radiance += mis_weight * throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; } diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index e3ea84b8ad5d2..42db57e4a4d3d 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -8,7 +8,7 @@ #import bevy_render::view::View #import bevy_solari::sampling::{sample_random_light, trace_point_visibility} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} -#import bevy_solari::world_cache::query_world_cache +#import bevy_solari::world_cache::query_world_cache_radiance @group(1) @binding(0) var view_output: texture_storage_2d; @group(1) @binding(1) var gi_reservoirs_a: array; @@ -83,7 +83,7 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { // textureStore(view_output, global_id.xy, pixel_color); #ifdef VISUALIZE_WORLD_CACHE - textureStore(view_output, global_id.xy, vec4(query_world_cache(world_position, world_normal, view.world_position) * view.exposure, 1.0)); + textureStore(view_output, global_id.xy, vec4(query_world_cache_radiance(world_position, world_normal, view.world_position) * view.exposure, 1.0)); #endif } @@ -112,7 +112,7 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 reservoir.radiance = direct_lighting.radiance; reservoir.unbiased_contribution_weight = direct_lighting.inverse_pdf * uniform_hemisphere_inverse_pdf(); #else - reservoir.radiance = query_world_cache(sample_point.world_position, sample_point.geometric_world_normal, view.world_position).radiance; + reservoir.radiance = query_world_cache_radiance(sample_point.world_position, sample_point.geometric_world_normal, view.world_position); reservoir.unbiased_contribution_weight = uniform_hemisphere_inverse_pdf(); #endif diff --git a/crates/bevy_solari/src/realtime/shade_di.wgsl b/crates/bevy_solari/src/realtime/shade_di.wgsl index 8b73ccaf98127..f5e184832f2b8 100644 --- a/crates/bevy_solari/src/realtime/shade_di.wgsl +++ b/crates/bevy_solari/src/realtime/shade_di.wgsl @@ -6,7 +6,7 @@ #import bevy_render::maths::PI #import bevy_render::view::View #import bevy_solari::scene_bindings::ResolvedMaterial -#import bevy_solari::world_cache::{query_world_cache, evaluate_lighting_from_cache, write_world_cache_light} +#import bevy_solari::world_cache::{query_world_cache_lights, evaluate_lighting_from_cache, write_world_cache_light} @group(1) @binding(0) var view_output: texture_storage_2d; @group(1) @binding(3) var gbuffer: texture_2d; @@ -43,9 +43,9 @@ fn shade(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocat material.roughness = clamp(base_rough.a * base_rough.a, 0.001, 1.0); material.metallic = props.g; - let cell = query_world_cache(world_position, world_normal, view.world_position); + let cell = query_world_cache_lights(&rng, world_position, world_normal, view.world_position); let direct_lighting = evaluate_lighting_from_cache(&rng, cell, world_position, world_normal, wo, material, view.exposure); - write_world_cache_light(direct_lighting, world_position, world_normal, view.world_position); + write_world_cache_light(direct_lighting, world_position, world_normal, view.world_position, view.exposure); let pixel_color = direct_lighting.radiance * direct_lighting.inverse_pdf * view.exposure + material.emissive; textureStore(view_output, global_id.xy, vec4(pixel_color, 1.0)); diff --git a/crates/bevy_solari/src/realtime/world_cache_compact.wgsl b/crates/bevy_solari/src/realtime/world_cache_compact.wgsl index 36aeef6d321aa..b83d51082e85e 100644 --- a/crates/bevy_solari/src/realtime/world_cache_compact.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_compact.wgsl @@ -1,4 +1,4 @@ -#import bevy_solari::world_cache::{WORLD_CACHE_EMPTY_CELL, world_cache_life, world_cache_checksums, world_cache_radiance, world_cache_a, world_cache_b, world_cache_active_cell_indices, world_cache_active_cells_count, world_cache_light_data, world_cache_light_data_new_lights, WORLD_CACHE_CELL_LIGHT_COUNT} +#import bevy_solari::world_cache::{WORLD_CACHE_EMPTY_CELL, world_cache_life, world_cache_checksums, world_cache_radiance, world_cache_a, world_cache_b, world_cache_active_cell_indices, world_cache_active_cells_count, world_cache_light_data, world_cache_light_data_new_lights, WorldCacheSingleLightData, WORLD_CACHE_CELL_LIGHT_COUNT} @group(2) @binding(0) var world_cache_active_cells_dispatch: vec3; @@ -16,24 +16,53 @@ fn decay_world_cache(@builtin(global_invocation_id) global_id: vec3) { world_cache_checksums[global_id.x] = WORLD_CACHE_EMPTY_CELL; world_cache_radiance[global_id.x] = vec4(0.0); } else { - var count = min(WORLD_CACHE_CELL_LIGHT_COUNT, atomicLoad(&world_cache_light_data_new_lights[global_id.x].visible_light_count)); + let old_count = world_cache_light_data[global_id.x].visible_light_count; + let old_lights = world_cache_light_data[global_id.x].visible_lights; + let new_count = min(WORLD_CACHE_CELL_LIGHT_COUNT, atomicLoad(&world_cache_light_data_new_lights[global_id.x].visible_light_count)); atomicStore(&world_cache_light_data_new_lights[global_id.x].visible_light_count, 0u); - for (var i = 0u; i < count; i++) { + var out_i = 0u; + var out_lights: array; + + for (var i = 0u; i < new_count; i++) { let data = atomicLoad(&world_cache_light_data_new_lights[global_id.x].visible_lights[i]); atomicStore(&world_cache_light_data_new_lights[global_id.x].visible_lights[i], 0); if data == 0 { - count = i; break; } let light = u32(data & 0xffffffff); let weight = bitcast(u32(data >> 32u)); - world_cache_light_data[global_id.x].visible_lights[i].light = light; - world_cache_light_data[global_id.x].visible_lights[i].weight = weight; + var exist_index = 0u; + if is_light_in_array(out_lights, out_i, light, &exist_index) { + out_lights[exist_index].weight = max(out_lights[exist_index].weight, weight); + } else { + out_lights[out_i] = WorldCacheSingleLightData(light, weight); + out_i++; + } + } + for (var i = 0u; i < old_count && out_i < WORLD_CACHE_CELL_LIGHT_COUNT; i++) { + var exist_index = 0u; + if is_light_in_array(out_lights, out_i, old_lights[i].light, &exist_index) { + out_lights[exist_index].weight = max(out_lights[exist_index].weight, old_lights[i].weight); + } else { + out_lights[out_i] = old_lights[i]; + out_i++; + } } - world_cache_light_data[global_id.x].visible_light_count = count; + world_cache_light_data[global_id.x].visible_light_count = out_i; + world_cache_light_data[global_id.x].visible_lights = out_lights; + } + } +} + +fn is_light_in_array(arr: array, len: u32, light: u32, out_index: ptr) -> bool { + for (var i = 0u; i < len; i++) { + if arr[i].light == light { + *out_index = i; + return true; } } + return false; } @compute @workgroup_size(1024, 1, 1) diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index eee9f330349f8..d10f302cb5093 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -50,12 +50,6 @@ struct WorldCacheLightDataWrite { visible_lights: array, WORLD_CACHE_CELL_LIGHT_COUNT>, } -struct WorldCacheData { - radiance: vec3, - visible_light_count: u32, - visible_lights: array, -} - @group(1) @binding(10) var world_cache_checksums: array, #{WORLD_CACHE_SIZE}>; #ifdef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER @group(1) @binding(11) var world_cache_life: array; @@ -73,7 +67,7 @@ struct WorldCacheData { @group(1) @binding(20) var world_cache_active_cells_count: u32; #ifndef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER -fn query_world_cache(world_position: vec3, world_normal: vec3, view_position: vec3) -> WorldCacheData { +fn query_world_cache_radiance(world_position: vec3, world_normal: vec3, view_position: vec3) -> vec3 { let cell_size = get_cell_size(world_position, view_position); let world_position_quantized = quantize_position(world_position, cell_size); let world_normal_quantized = quantize_normal(world_normal); @@ -85,46 +79,89 @@ fn query_world_cache(world_position: vec3, world_normal: vec3, view_po if existing_checksum == checksum { // Cache entry already exists - get radiance and reset cell lifetime atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); - let radiance = world_cache_radiance[key].rgb; - let data = world_cache_light_data[key]; - return WorldCacheData(radiance, data.visible_light_count, data.visible_lights); + return world_cache_radiance[key].rgb; } else if existing_checksum == WORLD_CACHE_EMPTY_CELL { // Cell is empty - reset cell lifetime so that it starts getting updated next frame atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); world_cache_geometry_data[key].world_position = world_position; world_cache_geometry_data[key].world_normal = world_normal; world_cache_light_data[key].visible_light_count = 0u; - return WorldCacheData(vec3(0.0), 0u, array()); + return vec3(0.0); } else { // Collision - jump to another entry key = wrap_key(pcg_hash(key)); } } - return WorldCacheData(vec3(0.0), 0u, array()); + return vec3(0.0); } -#endif -fn write_world_cache_light(cell: EvaluatedLighting, world_position: vec3, world_normal: vec3, view_position: vec3) { +fn query_world_cache_lights(rng: ptr, world_position: vec3, world_normal: vec3, view_position: vec3) -> WorldCacheLightDataRead { let cell_size = get_cell_size(world_position, view_position); - let world_position_quantized = quantize_position(world_position, cell_size); + var world_position_quantized = quantize_position(world_position, cell_size); + let center_offset = quantize_position_fract(world_position, cell_size) - vec3(0.5); + let direction = vec3(sign(center_offset)); + let lerp = vec3(rand_f(rng), rand_f(rng), rand_f(rng)); + let p_lerp_away = abs(center_offset); + world_position_quantized += vec3(select(vec3(0), direction, lerp > p_lerp_away)); + let world_normal_quantized = quantize_normal(world_normal); var key = compute_key(world_position_quantized, world_normal_quantized); let checksum = compute_checksum(world_position_quantized, world_normal_quantized); for (var i = 0u; i < WORLD_CACHE_MAX_SEARCH_STEPS; i++) { - // Don't need a CAS because we know the cell is alive - let existing_checksum = atomicLoad(&world_cache_checksums[key]); + let existing_checksum = atomicCompareExchangeWeak(&world_cache_checksums[key], WORLD_CACHE_EMPTY_CELL, checksum).old_value; if existing_checksum == checksum { + // Cache entry already exists - get radiance and reset cell lifetime + atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); + return world_cache_light_data[key]; + } else if existing_checksum == WORLD_CACHE_EMPTY_CELL { + // Cell is empty - reset cell lifetime so that it starts getting updated next frame + atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); + world_cache_geometry_data[key].world_position = world_position; + world_cache_geometry_data[key].world_normal = world_normal; + world_cache_light_data[key].visible_light_count = 0u; + return WorldCacheLightDataRead(0u, 0u, array()); + } else { + // Collision - jump to another entry + key = wrap_key(pcg_hash(key)); + } + } + + return WorldCacheLightDataRead(0u, 0u, array()); +} + +fn write_world_cache_light(cell: EvaluatedLighting, world_position: vec3, world_normal: vec3, view_position: vec3, exposure: f32) { + let cell_selected_weight = cell.data.weight + log2(exposure); + if cell_selected_weight < 0.0001 { return; } + + let cell_size = get_cell_size(world_position, view_position); + let world_position_quantized = quantize_position(world_position, cell_size); + let world_normal_quantized = quantize_normal(world_normal); + var key = compute_key(world_position_quantized, world_normal_quantized); + let checksum = compute_checksum(world_position_quantized, world_normal_quantized); + + for (var i = 0u; i < WORLD_CACHE_MAX_SEARCH_STEPS; i++) { + let existing_checksum = atomicCompareExchangeWeak(&world_cache_checksums[key], WORLD_CACHE_EMPTY_CELL, checksum).old_value; + if existing_checksum == checksum || existing_checksum == WORLD_CACHE_EMPTY_CELL { let index = atomicAdd(&world_cache_light_data_new_lights[key].visible_light_count, 1u) & (WORLD_CACHE_CELL_LIGHT_COUNT - 1u); let packed = (u64(bitcast(cell.data.weight)) << 32u) | u64(cell.data.light); - atomicMax(&world_cache_light_data_new_lights[key].visible_lights[index], packed); - } else { + atomicStore(&world_cache_light_data_new_lights[key].visible_lights[index], packed); + + if existing_checksum == WORLD_CACHE_EMPTY_CELL { + // Cell is empty - reset cell lifetime so that it starts getting updated next frame + atomicStore(&world_cache_life[key], WORLD_CACHE_CELL_LIFETIME); + world_cache_geometry_data[key].world_position = world_position; + world_cache_geometry_data[key].world_normal = world_normal; + world_cache_light_data[key].visible_light_count = 0u; + } + } else { // Collision - jump to another entry key = wrap_key(pcg_hash(key)); } } } +#endif struct EvaluatedLighting { radiance: vec3, @@ -134,7 +171,7 @@ struct EvaluatedLighting { fn evaluate_lighting_from_cache( rng: ptr, - cell: WorldCacheData, + cell: WorldCacheLightDataRead, world_position: vec3, world_normal: vec3, wo: vec3, @@ -143,7 +180,7 @@ fn evaluate_lighting_from_cache( ) -> EvaluatedLighting { let cell_selected_light = select_light_from_cache_cell(rng, cell, world_position, world_normal, wo, material); let cell_selected_weight = cell_selected_light.weight + log2(exposure); - let cell_confidence = smoothstep(0.1, 0.3, cell_selected_weight); + let cell_confidence = smoothstep(0.0001, 0.3, cell_selected_weight); // Sample more random lights if our cell has bad lights let random_sample_count = u32(round(mix(f32(WORLD_CACHE_CELL_LIGHT_COUNT), f32(WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT), cell_confidence))); @@ -195,7 +232,7 @@ fn p_wrs(selection: SelectedLight) -> f32 { fn select_light_from_cache_cell( rng: ptr, - cell: WorldCacheData, + cell: WorldCacheLightDataRead, world_position: vec3, world_normal: vec3, wo: vec3, @@ -271,11 +308,27 @@ fn get_cell_size(world_position: vec3, view_position: vec3) -> f32 { } fn quantize_position(world_position: vec3, quantization_factor: f32) -> vec3 { - return vec3(vec3(floor(world_position / quantization_factor + 0.0001))); + return vec3(vec3(floor(world_position / quantization_factor))); +} + +fn quantize_position_fract(world_position: vec3, quantization_factor: f32) -> vec3 { + return fract(world_position / quantization_factor); } fn quantize_normal(world_normal: vec3) -> vec3 { - return vec3(vec3(floor(world_normal + 0.0001))); + let x = vec3(1.0, 0.0, 0.0); + let y = vec3(0.0, 1.0, 0.0); + let z = vec3(0.0, 0.0, 1.0); + let dot_x = dot(world_normal, x); + let dot_y = dot(world_normal, y); + let dot_z = dot(world_normal, z); + let max_dot = max(max(abs(dot_x), abs(dot_y)), abs(dot_z)); + + var sel = select(vec3(0.0), x, max_dot == dot_x); + sel = select(sel, y, max_dot == dot_y); + sel = select(sel, z, max_dot == dot_z); + + return vec3(vec3(sel * sign(max_dot))); } // TODO: Clustering diff --git a/crates/bevy_solari/src/realtime/world_cache_update.wgsl b/crates/bevy_solari/src/realtime/world_cache_update.wgsl index 178db2cd283d4..cc4aa64b94776 100644 --- a/crates/bevy_solari/src/realtime/world_cache_update.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_update.wgsl @@ -10,9 +10,9 @@ world_cache_geometry_data, world_cache_radiance, world_cache_active_cells_new_radiance, - query_world_cache, + query_world_cache_radiance, + query_world_cache_lights, evaluate_lighting_from_cache, - write_world_cache_light, } @group(1) @binding(8) var view: View; @@ -36,9 +36,8 @@ fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3, @builtin(glob material.roughness = 1.0; material.metallic = 0.0; - let cell = query_world_cache(geometry_data.world_position, geometry_data.world_normal, view.world_position); + let cell = query_world_cache_lights(&rng, geometry_data.world_position, geometry_data.world_normal, view.world_position); let direct_lighting = evaluate_lighting_from_cache(&rng, cell, geometry_data.world_position, geometry_data.world_normal, geometry_data.world_normal, material, view.exposure); - write_world_cache_light(direct_lighting, geometry_data.world_position, geometry_data.world_normal, view.world_position); var new_radiance = direct_lighting.radiance * direct_lighting.inverse_pdf; #ifndef NO_MULTIBOUNCE @@ -46,7 +45,7 @@ fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3, @builtin(glob let ray_hit = trace_ray(geometry_data.world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { let ray_hit = resolve_ray_hit_full(ray_hit); - new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position).radiance; + new_radiance += ray_hit.material.base_color * query_world_cache_radiance(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position); } #endif From c94c7a126621e36c4b1f489460855cbb5a2b0e07 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Fri, 3 Oct 2025 18:28:26 +0100 Subject: [PATCH 3/4] tune numbers a bit --- .../src/realtime/world_cache_query.wgsl | 69 +++++++++++-------- 1 file changed, 41 insertions(+), 28 deletions(-) diff --git a/crates/bevy_solari/src/realtime/world_cache_query.wgsl b/crates/bevy_solari/src/realtime/world_cache_query.wgsl index d10f302cb5093..79364a04ee37a 100644 --- a/crates/bevy_solari/src/realtime/world_cache_query.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_query.wgsl @@ -2,7 +2,8 @@ #import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance #import bevy_pbr::utils::rand_f -#import bevy_solari::brdf::evaluate_brdf +#import bevy_render::maths::PI +// #import bevy_solari::brdf::evaluate_brdf #import bevy_solari::sampling::{light_contribution_no_trace, select_random_light, select_random_light_inverse_pdf, trace_light_visibility} #import bevy_solari::scene_bindings::ResolvedMaterial @@ -13,12 +14,17 @@ const WORLD_CACHE_CELL_LIFETIME: u32 = 30u; /// Maximum amount of attempts to find a cache entry after a hash collision const WORLD_CACHE_MAX_SEARCH_STEPS: u32 = 3u; /// Maximum lights stored in each cache cell +/// This should match `WORLD_CACHE_CELL_LIGHT_COUNT` in `realtime/prepare.rs`! const WORLD_CACHE_CELL_LIGHT_COUNT: u32 = 8u; /// Lights searched that aren't in the cell -const WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT: u32 = 2u; +const WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT_MIN: u32 = 4u; +const WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT_MAX: u32 = 10u; +const WORLD_CACHE_EXPLORATORY_SAMPLE_RATIO: f32 = 0.20; +const WORLD_CACHE_CELL_CONFIDENCE_LUM_MIN: f32 = 0.0001; +const WORLD_CACHE_CELL_CONFIDENCE_LUM_MAX: f32 = 0.1; /// The size of a cache cell at the lowest LOD in meters -const WORLD_CACHE_POSITION_BASE_CELL_SIZE: f32 = 0.25; +const WORLD_CACHE_POSITION_BASE_CELL_SIZE: f32 = 0.1; /// How fast the world cache transitions between LODs as a function of distance to the camera const WORLD_CACHE_POSITION_LOD_SCALE: f32 = 30.0; @@ -37,7 +43,6 @@ struct WorldCacheSingleLightData { weight: f32, } -// The size of these structs should match `WORLD_CACHE_CELL_LIGHT_COUNT` in `realtime/prepare.rs`! struct WorldCacheLightDataRead { visible_light_count: u32, padding: u32, @@ -103,7 +108,7 @@ fn query_world_cache_lights(rng: ptr, world_position: vec3, let direction = vec3(sign(center_offset)); let lerp = vec3(rand_f(rng), rand_f(rng), rand_f(rng)); let p_lerp_away = abs(center_offset); - world_position_quantized += vec3(select(vec3(0), direction, lerp > p_lerp_away)); + world_position_quantized += select(vec3(0), direction, lerp > p_lerp_away); let world_normal_quantized = quantize_normal(world_normal); var key = compute_key(world_position_quantized, world_normal_quantized); @@ -133,7 +138,7 @@ fn query_world_cache_lights(rng: ptr, world_position: vec3, fn write_world_cache_light(cell: EvaluatedLighting, world_position: vec3, world_normal: vec3, view_position: vec3, exposure: f32) { let cell_selected_weight = cell.data.weight + log2(exposure); - if cell_selected_weight < 0.0001 { return; } + if cell_selected_weight < WORLD_CACHE_CELL_CONFIDENCE_LUM_MIN { return; } let cell_size = get_cell_size(world_position, view_position); let world_position_quantized = quantize_position(world_position, cell_size); @@ -180,15 +185,15 @@ fn evaluate_lighting_from_cache( ) -> EvaluatedLighting { let cell_selected_light = select_light_from_cache_cell(rng, cell, world_position, world_normal, wo, material); let cell_selected_weight = cell_selected_light.weight + log2(exposure); - let cell_confidence = smoothstep(0.0001, 0.3, cell_selected_weight); + let cell_confidence = smoothstep(WORLD_CACHE_CELL_CONFIDENCE_LUM_MIN, WORLD_CACHE_CELL_CONFIDENCE_LUM_MAX, cell_selected_weight); // Sample more random lights if our cell has bad lights - let random_sample_count = u32(round(mix(f32(WORLD_CACHE_CELL_LIGHT_COUNT), f32(WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT), cell_confidence))); + let random_sample_count = u32(round(mix(f32(WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT_MAX), f32(WORLD_CACHE_NEW_LIGHTS_SEARCH_COUNT_MIN), cell_confidence))); let random_selected_light = select_light_randomly(rng, world_position, world_normal, wo, material, random_sample_count); let p_cell_selection = select(p_wrs(cell_selected_light), 0.0, cell_selected_light.weight_sum < 0.0001); let p_random_selection = select(p_wrs(random_selected_light), 0.0, random_selected_light.weight_sum < 0.0001); - let p_random_selection_clamped = min(mix(1.0, 0.25 * p_cell_selection, cell_confidence), p_random_selection); + let p_random_selection_clamped = min(mix(1.0, WORLD_CACHE_EXPLORATORY_SAMPLE_RATIO * p_cell_selection, cell_confidence), p_random_selection); let weight_sum = p_cell_selection + p_random_selection_clamped; if weight_sum < 0.0001 { @@ -301,21 +306,25 @@ fn select_light_randomly( return SelectedLight(selected, selected_weight, weight_sum, base_pdf); } +fn evaluate_brdf(normal: vec3, wo: vec3, wi: vec3, material: ResolvedMaterial) -> vec3 { + return material.base_color / PI; +} + fn get_cell_size(world_position: vec3, view_position: vec3) -> f32 { let camera_distance = distance(view_position, world_position) / WORLD_CACHE_POSITION_LOD_SCALE; let lod = exp2(floor(log2(1.0 + camera_distance))); return WORLD_CACHE_POSITION_BASE_CELL_SIZE * lod; } -fn quantize_position(world_position: vec3, quantization_factor: f32) -> vec3 { - return vec3(vec3(floor(world_position / quantization_factor))); +fn quantize_position(world_position: vec3, quantization_factor: f32) -> vec3 { + return vec3(floor(world_position / quantization_factor)); } fn quantize_position_fract(world_position: vec3, quantization_factor: f32) -> vec3 { return fract(world_position / quantization_factor); } -fn quantize_normal(world_normal: vec3) -> vec3 { +fn quantize_normal(world_normal: vec3) -> vec3 { let x = vec3(1.0, 0.0, 0.0); let y = vec3(0.0, 1.0, 0.0); let z = vec3(0.0, 0.0, 1.0); @@ -328,28 +337,32 @@ fn quantize_normal(world_normal: vec3) -> vec3 { sel = select(sel, y, max_dot == dot_y); sel = select(sel, z, max_dot == dot_z); - return vec3(vec3(sel * sign(max_dot))); + return vec3(sel * sign(max_dot)); } // TODO: Clustering -fn compute_key(world_position: vec3, world_normal: vec3) -> u32 { - var key = pcg_hash(world_position.x); - key = pcg_hash(key + world_position.y); - key = pcg_hash(key + world_position.z); - key = pcg_hash(key + world_normal.x); - key = pcg_hash(key + world_normal.y); - key = pcg_hash(key + world_normal.z); +fn compute_key(world_position: vec3, world_normal: vec3) -> u32 { + let pos = vec3(world_position); + let norm = vec3(world_normal); + var key = pcg_hash(pos.x); + key = pcg_hash(key + pos.y); + key = pcg_hash(key + pos.z); + key = pcg_hash(key + norm.x); + key = pcg_hash(key + norm.y); + key = pcg_hash(key + norm.z); return wrap_key(key); } -fn compute_checksum(world_position: vec3, world_normal: vec3) -> u32 { - var key = iqint_hash(world_position.x); - key = iqint_hash(key + world_position.y); - key = iqint_hash(key + world_position.z); - key = iqint_hash(key + world_normal.x); - key = iqint_hash(key + world_normal.y); - key = iqint_hash(key + world_normal.z); - return key; +fn compute_checksum(world_position: vec3, world_normal: vec3) -> u32 { + let pos = vec3(world_position); + let norm = vec3(world_normal); + var key = iqint_hash(pos.x); + key = iqint_hash(key + pos.y); + key = iqint_hash(key + pos.z); + key = iqint_hash(key + norm.x); + key = iqint_hash(key + norm.y); + key = iqint_hash(key + norm.z); + return u32(key); } fn pcg_hash(input: u32) -> u32 { From 064c05b0acbb42747c954281b2a9cc6acb888a87 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Sat, 4 Oct 2025 20:56:01 +0100 Subject: [PATCH 4/4] add required wgpu features --- crates/bevy_solari/src/lib.rs | 2 ++ 1 file changed, 2 insertions(+) diff --git a/crates/bevy_solari/src/lib.rs b/crates/bevy_solari/src/lib.rs index 6aaea2180fe0b..5e144fad08c6d 100644 --- a/crates/bevy_solari/src/lib.rs +++ b/crates/bevy_solari/src/lib.rs @@ -55,5 +55,7 @@ impl SolariPlugins { | WgpuFeatures::TEXTURE_BINDING_ARRAY | WgpuFeatures::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING | WgpuFeatures::PARTIALLY_BOUND_BINDING_ARRAY + | WgpuFeatures::SHADER_INT64 + | WgpuFeatures::SHADER_INT64_ATOMIC_MIN_MAX } }