From 90e02c7b385261156e1288fc633b5b37b0403129 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Wed, 9 Jul 2025 10:18:54 -0400 Subject: [PATCH 01/31] Adjust specular_multiscatter to not take LightingInput --- crates/bevy_pbr/src/render/pbr_lighting.wgsl | 23 +++++++++----------- 1 file changed, 10 insertions(+), 13 deletions(-) diff --git a/crates/bevy_pbr/src/render/pbr_lighting.wgsl b/crates/bevy_pbr/src/render/pbr_lighting.wgsl index 17cae13b920f2..6bc24d8af01d6 100644 --- a/crates/bevy_pbr/src/render/pbr_lighting.wgsl +++ b/crates/bevy_pbr/src/render/pbr_lighting.wgsl @@ -235,16 +235,13 @@ fn fresnel(f0: vec3, LdotH: f32) -> vec3 { // Multiscattering approximation: // fn specular_multiscatter( - input: ptr, D: f32, V: f32, F: vec3, + F0: vec3, + F_ab: vec2, specular_intensity: f32, ) -> vec3 { - // Unpack. - let F0 = (*input).F0_; - let F_ab = (*input).F_ab; - var Fr = (specular_intensity * D * V) * F; Fr *= 1.0 + F0 * (1.0 / F_ab.x - 1.0); return Fr; @@ -329,7 +326,7 @@ fn specular( let F = fresnel(F0, LdotH); // Calculate the specular light. - let Fr = specular_multiscatter(input, D, V, F, specular_intensity); + let Fr = specular_multiscatter(D, V, F, F0, (*input).F_ab, specular_intensity); return Fr; } @@ -397,7 +394,7 @@ fn specular_anisotropy( let Fa = fresnel(F0, LdotH); // Calculate the specular light. - let Fr = specular_multiscatter(input, Da, Va, Fa, specular_intensity); + let Fr = specular_multiscatter(Da, Va, Fa, F0, (*input).F_ab, specular_intensity); return Fr; } @@ -482,7 +479,7 @@ fn cubemap_uv(direction: vec3, cubemap_type: u32) -> vec2 { ), max_axis != abs_direction.x ); - + var face_uv: vec2; var divisor: f32; var corner_uv: vec2 = vec2(0, 0); @@ -500,12 +497,12 @@ fn cubemap_uv(direction: vec3, cubemap_type: u32) -> vec2 { face_uv = (face_uv / divisor) * 0.5 + 0.5; switch cubemap_type { - case CUBEMAP_TYPE_CROSS_VERTICAL: { - face_size = vec2(1.0/3.0, 1.0/4.0); + case CUBEMAP_TYPE_CROSS_VERTICAL: { + face_size = vec2(1.0/3.0, 1.0/4.0); corner_uv = vec2((0x111102u >> (4 * face_index)) & 0xFu, (0x132011u >> (4 * face_index)) & 0xFu); } - case CUBEMAP_TYPE_CROSS_HORIZONTAL: { - face_size = vec2(1.0/4.0, 1.0/3.0); + case CUBEMAP_TYPE_CROSS_HORIZONTAL: { + face_size = vec2(1.0/4.0, 1.0/3.0); corner_uv = vec2((0x131102u >> (4 * face_index)) & 0xFu, (0x112011u >> (4 * face_index)) & 0xFu); } case CUBEMAP_TYPE_SEQUENCE_HORIZONTAL: { @@ -765,7 +762,7 @@ fn directional_light( view_bindings::clustered_decal_sampler, decal_uv - floor(decal_uv), 0.0 - ).r; + ).r; } else { texture_sample = 0f; } From c73405b9e4cb6fce9a3d84cc87ced4d91a7af2bd Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Wed, 9 Jul 2025 10:37:45 -0400 Subject: [PATCH 02/31] Fix SSAO specular occlusion roughness bug --- crates/bevy_pbr/src/render/pbr_fragment.wgsl | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/crates/bevy_pbr/src/render/pbr_fragment.wgsl b/crates/bevy_pbr/src/render/pbr_fragment.wgsl index 779546f8bd67d..3c69c4405f984 100644 --- a/crates/bevy_pbr/src/render/pbr_fragment.wgsl +++ b/crates/bevy_pbr/src/render/pbr_fragment.wgsl @@ -377,7 +377,6 @@ fn pbr_input_from_standard_material( var perceptual_roughness: f32 = pbr_bindings::material.perceptual_roughness; #endif // BINDLESS - let roughness = lighting::perceptualRoughnessToRoughness(perceptual_roughness); #ifdef VERTEX_UVS if ((flags & pbr_types::STANDARD_MATERIAL_FLAGS_METALLIC_ROUGHNESS_TEXTURE_BIT) != 0u) { let metallic_roughness = @@ -627,7 +626,7 @@ fn pbr_input_from_standard_material( var specular_occlusion: f32 = 1.0; #ifdef VERTEX_UVS if ((flags & pbr_types::STANDARD_MATERIAL_FLAGS_OCCLUSION_TEXTURE_BIT) != 0u) { - diffuse_occlusion *= + diffuse_occlusion *= #ifdef MESHLET_MESH_MATERIAL_PASS textureSampleGrad( #else // MESHLET_MESH_MATERIAL_PASS @@ -660,7 +659,8 @@ fn pbr_input_from_standard_material( diffuse_occlusion = min(diffuse_occlusion, ssao_multibounce); // Use SSAO to estimate the specular occlusion. // Lagarde and Rousiers 2014, "Moving Frostbite to Physically Based Rendering" - specular_occlusion = saturate(pow(NdotV + ssao, exp2(-16.0 * roughness - 1.0)) - 1.0 + ssao); + let roughness = lighting::perceptualRoughnessToRoughness(pbr_input.material.perceptual_roughness); + specular_occlusion = saturate(pow(NdotV + ssao, exp2(-16.0 * roughness - 1.0)) - 1.0 + ssao); #endif pbr_input.diffuse_occlusion = diffuse_occlusion; pbr_input.specular_occlusion = specular_occlusion; From f340c77f5b0d36f0e6990b92c4a5d3b7560f58fb Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Wed, 9 Jul 2025 11:02:09 -0400 Subject: [PATCH 03/31] Add specular material properties to solari bindings --- crates/bevy_solari/src/scene/binder.rs | 31 ++++++++++++++----- .../src/scene/raytracing_scene_bindings.wgsl | 26 +++++++++++++--- 2 files changed, 45 insertions(+), 12 deletions(-) diff --git a/crates/bevy_solari/src/scene/binder.rs b/crates/bevy_solari/src/scene/binder.rs index f14b5dbe23b6f..0f71d0e6e6c47 100644 --- a/crates/bevy_solari/src/scene/binder.rs +++ b/crates/bevy_solari/src/scene/binder.rs @@ -115,13 +115,23 @@ pub fn prepare_raytracing_scene_bindings( let Some(emissive_texture_id) = process_texture(&material.emissive_texture) else { continue; }; + let Some(metallic_roughness_texture_id) = + process_texture(&material.metallic_roughness_texture) + else { + continue; + }; materials.get_mut().push(GpuMaterial { - base_color: material.base_color.to_linear(), - emissive: material.emissive, - base_color_texture_id, normal_map_texture_id, + base_color_texture_id, emissive_texture_id, + metallic_roughness_texture_id, + + base_color: LinearRgba::from(material.base_color).to_vec3(), + perceptual_roughness: material.perceptual_roughness, + emissive: material.emissive.to_vec3(), + metallic: material.metallic, + reflectance: LinearRgba::from(material.specular_tint).to_vec3() * material.reflectance, _padding: Default::default(), }); @@ -184,7 +194,7 @@ pub fn prepare_raytracing_scene_bindings( material_ids.get_mut().push(material_id); - if material.emissive != LinearRgba::BLACK { + if material.emissive != Vec3::ZERO { light_sources .get_mut() .push(GpuLightSource::new_emissive_mesh_light( @@ -342,12 +352,17 @@ struct GpuInstanceGeometryIds { #[derive(ShaderType)] struct GpuMaterial { - base_color: LinearRgba, - emissive: LinearRgba, - base_color_texture_id: u32, normal_map_texture_id: u32, + base_color_texture_id: u32, emissive_texture_id: u32, - _padding: u32, + metallic_roughness_texture_id: u32, + + base_color: Vec3, + perceptual_roughness: f32, + emissive: Vec3, + metallic: f32, + reflectance: Vec3, + _padding: f32, } #[derive(ShaderType)] diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index eeed96ad8e818..a59b38e952e05 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -34,12 +34,17 @@ fn unpack_vertex(packed: PackedVertex) -> Vertex { } struct Material { - base_color: vec4, - emissive: vec4, - base_color_texture_id: u32, normal_map_texture_id: u32, + base_color_texture_id: u32, emissive_texture_id: u32, - _padding: u32, + metallic_roughness_texture_id: u32, + + base_color: vec3, + perceptual_roughness: f32, + emissive: vec3, + metallic: f32, + reflectance: vec3, + _padding: f32, } const TEXTURE_MAP_NONE = 0xFFFFFFFFu; @@ -94,6 +99,9 @@ fn sample_texture(id: u32, uv: vec2) -> vec3 { struct ResolvedMaterial { base_color: vec3, emissive: vec3, + reflectance: vec3, + perceptual_roughness: f32, + metallic: f32, } struct ResolvedRayHitFull { @@ -118,6 +126,16 @@ fn resolve_material(material: Material, uv: vec2) -> ResolvedMaterial { m.emissive *= sample_texture(material.emissive_texture_id, uv); } + m.reflectance = material.reflectance; + + m.perceptual_roughness = material.perceptual_roughness; + m.metallic = material.metallic; + if material.metallic_roughness_texture_id != TEXTURE_MAP_NONE { + let metallic_roughness = sample_texture(material.metallic_roughness_texture_id, uv); + m.perceptual_roughness *= metallic_roughness.g; + m.metallic *= metallic_roughness.b; + } + return m; } From dfb3401dd906b5ee6ed405c7b5c88792e8011238 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Wed, 9 Jul 2025 11:20:47 -0400 Subject: [PATCH 04/31] Fix inverted ray in path tracer (dosen't actually affect results) --- crates/bevy_solari/src/pathtracer/pathtracer.wgsl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index c67b53e58e727..8ba5603518e8b 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -57,7 +57,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { ray_t_min = RAY_T_MIN; // Update throughput for next bounce - let cos_theta = dot(-ray_direction, ray_hit.world_normal); + let cos_theta = dot(ray_direction, ray_hit.world_normal); let cosine_hemisphere_pdf = cos_theta / PI; // Weight for the next bounce because we importance sampled the diffuse BRDF for the next ray direction throughput *= (diffuse_brdf * cos_theta) / cosine_hemisphere_pdf; From bf5ed11a845dff9ecf35fe00bf9c1e60e1b3ec31 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Wed, 9 Jul 2025 16:46:15 -0400 Subject: [PATCH 05/31] WIP specular material --- .../src/pathtracer/pathtracer.wgsl | 30 +++++++--- crates/bevy_solari/src/scene/brdf.wgsl | 57 +++++++++++++++++++ crates/bevy_solari/src/scene/mod.rs | 1 + crates/bevy_solari/src/scene/sampling.wgsl | 11 ++++ examples/3d/solari.rs | 8 ++- 5 files changed, 97 insertions(+), 10 deletions(-) create mode 100644 crates/bevy_solari/src/scene/brdf.wgsl diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 8ba5603518e8b..15f6117bacf3e 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -2,7 +2,8 @@ #import bevy_pbr::utils::{rand_f, rand_vec2f} #import bevy_render::maths::PI #import bevy_render::view::View -#import bevy_solari::sampling::{sample_random_light, sample_cosine_hemisphere} +#import bevy_solari::brdf::evaluate_brdf +#import bevy_solari::sampling::{sample_random_light, sample_uniform_hemisphere} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var accumulation_texture: texture_storage_2d; @@ -40,26 +41,37 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { let ray_hit = resolve_ray_hit_full(ray_hit); + // Sample new ray direction for the next bounce + let wo = -ray_direction; + let wi = sample_uniform_hemisphere(ray_hit.world_normal, &rng); + ray_direction = wi; + // Evaluate material BRDF - let diffuse_brdf = ray_hit.material.base_color / PI; + let brdf = evaluate_brdf( + ray_hit.world_normal, + wo, + wi, + ray_hit.material.base_color, + ray_hit.material.metallic, + ray_hit.material.reflectance, + ray_hit.material.perceptual_roughness, + ); // Use emissive only on the first ray (coming from the camera) if ray_t_min == 0.0 { radiance = ray_hit.material.emissive; } // Sample direct lighting - radiance += throughput * diffuse_brdf * sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); - - // Sample new ray direction from the material BRDF for next bounce - ray_direction = sample_cosine_hemisphere(ray_hit.world_normal, &rng); + // TODO: Wrong BRDF here when it comes to specular? + radiance += throughput * brdf * sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); // Update other variables for next bounce ray_origin = ray_hit.world_position; ray_t_min = RAY_T_MIN; // Update throughput for next bounce - let cos_theta = dot(ray_direction, ray_hit.world_normal); - let cosine_hemisphere_pdf = cos_theta / PI; // Weight for the next bounce because we importance sampled the diffuse BRDF for the next ray direction - throughput *= (diffuse_brdf * cos_theta) / cosine_hemisphere_pdf; + let cos_theta = dot(wi, ray_hit.world_normal); + let uniform_hemisphere_pdf = 1.0 / (2.0 * PI); // Weight for the next bounce because we importance sampled the diffuse BRDF for the next ray direction + throughput *= (brdf * cos_theta) / uniform_hemisphere_pdf; // Russian roulette for early termination let p = luminance(throughput); diff --git a/crates/bevy_solari/src/scene/brdf.wgsl b/crates/bevy_solari/src/scene/brdf.wgsl new file mode 100644 index 0000000000000..13ab7b46d4574 --- /dev/null +++ b/crates/bevy_solari/src/scene/brdf.wgsl @@ -0,0 +1,57 @@ +#define_import_path bevy_solari::brdf + +#import bevy_pbr::lighting::{F_AB, perceptualRoughnessToRoughness, D_GGX, V_SmithGGXCorrelated, fresnel, specular_multiscatter} +#import bevy_pbr::pbr_functions::{calculate_diffuse_color, calculate_F0} +#import bevy_render::maths::PI + +fn evaluate_brdf( + world_normal: vec3, + wo: vec3, + wi: vec3, + base_color: vec3, + metallic: f32, + reflectance: vec3, + perceptual_roughness: f32, +) -> vec3 { + let diffuse_brdf = diffuse_brdf(base_color, metallic); + let specular_brdf = specular_brdf( + world_normal, + wo, + wi, + base_color, + metallic, + reflectance, + perceptual_roughness, + ); + return diffuse_brdf + specular_brdf; +} + +fn diffuse_brdf(base_color: vec3, metallic: f32) -> vec3 { + let diffuse_color = calculate_diffuse_color(base_color, metallic, 0.0, 0.0); + return diffuse_color / PI; +} + +fn specular_brdf( + N: vec3, + V: vec3, + L: vec3, + base_color: vec3, + metallic: f32, + reflectance: vec3, + perceptual_roughness: f32, +) -> vec3 { + let H = normalize(L + V); + let NdotL = saturate(dot(N, L)); + let NdotH = saturate(dot(N, H)); + let LdotH = saturate(dot(L, H)); + let NdotV = max(dot(N, V), 0.0001); + + let F0 = calculate_F0(base_color, metallic, reflectance); + let F_ab = F_AB(perceptual_roughness, NdotV); + let roughness = perceptualRoughnessToRoughness(perceptual_roughness); + + let D = D_GGX(roughness, NdotH, H); + let Vs = V_SmithGGXCorrelated(roughness, NdotV, NdotL); + let F = fresnel(F0, LdotH); + return specular_multiscatter(D, Vs, F, F0, F_ab, 1.0); +} diff --git a/crates/bevy_solari/src/scene/mod.rs b/crates/bevy_solari/src/scene/mod.rs index a68e126480a34..ee02f08f9c9fa 100644 --- a/crates/bevy_solari/src/scene/mod.rs +++ b/crates/bevy_solari/src/scene/mod.rs @@ -31,6 +31,7 @@ pub struct RaytracingScenePlugin; impl Plugin for RaytracingScenePlugin { fn build(&self, app: &mut App) { + load_shader_library!(app, "brdf.wgsl"); load_shader_library!(app, "raytracing_scene_bindings.wgsl"); load_shader_library!(app, "sampling.wgsl"); diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index be709f0bc8dd1..c6ad92af49054 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -15,6 +15,17 @@ fn sample_cosine_hemisphere(normal: vec3, rng: ptr) -> vec3< return vec3(x, y, z); } +// https://www.pbr-book.org/3ed-2018/Monte_Carlo_Integration/2D_Sampling_with_Multidimensional_Transformations#UniformlySamplingaHemisphere +fn sample_uniform_hemisphere(normal: vec3, rng: ptr) -> vec3 { + let cos_theta = rand_f(rng); + let phi = PI_2 * rand_f(rng); + let sin_theta = sqrt(max(1.0 - cos_theta * cos_theta, 0.0)); + let x = sin_theta * cos(phi); + let y = sin_theta * sin(phi); + let z = cos_theta; + return build_orthonormal_basis(normal) * vec3(x, y, z); +} + // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec19%3A294 fn sample_disk(disk_radius: f32, rng: ptr) -> vec2 { let ab = 2.0 * rand_vec2f(rng) - 1.0; diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index 895df4d6fd413..bddc15e406994 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -113,8 +113,14 @@ fn add_raytracing_meshes_on_scene_load( } } - // Increase material emissive intensity to make it prettier for the example + // Adjust scene materials to better demo bevy_solari features for (_, material) in materials.iter_mut() { material.emissive *= 200.0; + + if material.base_color.to_linear() == LinearRgba::new(0.5, 0.5, 0.5, 1.0) { + material.metallic = 1.0; + material.perceptual_roughness = 0.0; + material.reflectance = 1.0; + } } } From 0ecbd8f57f683014a5cc0342951fee54f6aa0ce8 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Thu, 17 Jul 2025 11:45:14 -0400 Subject: [PATCH 06/31] Fix merge --- crates/bevy_solari/src/pathtracer/pathtracer.wgsl | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 662fcd992dfd5..15f6117bacf3e 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -61,10 +61,8 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { if ray_t_min == 0.0 { radiance = ray_hit.material.emissive; } // Sample direct lighting - radiance += throughput * diffuse_brdf * sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); - - // Sample new ray direction from the material BRDF for next bounce - ray_direction = sample_cosine_hemisphere(ray_hit.world_normal, &rng); + // TODO: Wrong BRDF here when it comes to specular? + radiance += throughput * brdf * sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); // Update other variables for next bounce ray_origin = ray_hit.world_position; From 18b5d1f3f65eb3b937045f89ece68dcdac600fdf Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 18 Jul 2025 16:43:47 -0400 Subject: [PATCH 07/31] The good parts --- crates/bevy_solari/src/realtime/mod.rs | 1 + crates/bevy_solari/src/realtime/node.rs | 43 ++++++- crates/bevy_solari/src/realtime/prepare.rs | 35 +++++- .../src/realtime/presample_light_tiles.wgsl | 98 ++++++++++++++++ .../bevy_solari/src/realtime/restir_di.wgsl | 105 ++++++++++-------- .../bevy_solari/src/realtime/restir_gi.wgsl | 18 +-- crates/bevy_solari/src/scene/sampling.wgsl | 43 +++---- 7 files changed, 259 insertions(+), 84 deletions(-) create mode 100644 crates/bevy_solari/src/realtime/presample_light_tiles.wgsl diff --git a/crates/bevy_solari/src/realtime/mod.rs b/crates/bevy_solari/src/realtime/mod.rs index a2c17c7269039..472bce67d8ead 100644 --- a/crates/bevy_solari/src/realtime/mod.rs +++ b/crates/bevy_solari/src/realtime/mod.rs @@ -31,6 +31,7 @@ pub struct SolariLightingPlugin; impl Plugin for SolariLightingPlugin { fn build(&self, app: &mut App) { + embedded_asset!(app, "presample_light_tiles.wgsl"); embedded_asset!(app, "restir_di.wgsl"); embedded_asset!(app, "restir_gi.wgsl"); diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index eaa432d8cbace..164c0fc9a1d49 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -1,4 +1,7 @@ -use super::{prepare::SolariLightingResources, SolariLighting}; +use super::{ + prepare::{SolariLightingResources, LIGHT_TILE_BLOCKS}, + SolariLighting, +}; use crate::scene::RaytracingSceneBindings; use bevy_asset::load_embedded_asset; use bevy_core_pipeline::prepass::{ @@ -36,6 +39,7 @@ pub mod graph { pub struct SolariLightingNode { bind_group_layout: BindGroupLayout, + presample_light_tiles_pipeline: CachedComputePipelineId, di_initial_and_temporal_pipeline: CachedComputePipelineId, di_spatial_and_shade_pipeline: CachedComputePipelineId, gi_initial_and_temporal_pipeline: CachedComputePipelineId, @@ -74,6 +78,7 @@ impl ViewNode for SolariLightingNode { let previous_view_uniforms = world.resource::(); let frame_count = world.resource::(); let ( + Some(presample_light_tiles_pipeline), Some(di_initial_and_temporal_pipeline), Some(di_spatial_and_shade_pipeline), Some(gi_initial_and_temporal_pipeline), @@ -86,6 +91,7 @@ impl ViewNode for SolariLightingNode { Some(view_uniforms), Some(previous_view_uniforms), ) = ( + 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.gi_initial_and_temporal_pipeline), @@ -107,6 +113,12 @@ impl ViewNode for SolariLightingNode { &self.bind_group_layout, &BindGroupEntries::sequential(( view_target.get_unsampled_color_attachment().view, + solari_lighting_resources + .light_tile_samples + .as_entire_binding(), + solari_lighting_resources + .light_tile_resolved_samples + .as_entire_binding(), solari_lighting_resources .di_reservoirs_a .as_entire_binding(), @@ -151,11 +163,14 @@ impl ViewNode for SolariLightingNode { ], ); - pass.set_pipeline(di_initial_and_temporal_pipeline); + 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_pipeline(di_initial_and_temporal_pipeline); pass.dispatch_workgroups(viewport.x.div_ceil(8), viewport.y.div_ceil(8), 1); pass.set_pipeline(di_spatial_and_shade_pipeline); @@ -217,6 +232,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), texture_2d(TextureSampleType::Uint), texture_depth_2d(), texture_2d(TextureSampleType::Float { filterable: true }), @@ -228,6 +245,21 @@ impl FromWorld for SolariLightingNode { ), ); + let presample_light_tiles_pipeline = + pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { + label: Some("solari_lighting_presample_light_tiles_pipeline".into()), + layout: vec![ + scene_bindings.bind_group_layout.clone(), + bind_group_layout.clone(), + ], + push_constant_ranges: vec![PushConstantRange { + stages: ShaderStages::COMPUTE, + range: 0..8, + }], + shader: load_embedded_asset!(world, "presample_light_tiles.wgsl"), + ..default() + }); + let di_initial_and_temporal_pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { label: Some("solari_lighting_di_initial_and_temporal_pipeline".into()), @@ -272,9 +304,8 @@ impl FromWorld for SolariLightingNode { range: 0..8, }], shader: load_embedded_asset!(world, "restir_gi.wgsl"), - shader_defs: vec![], entry_point: Some("initial_and_temporal".into()), - zero_initialize_workgroup_memory: false, + ..default() }); let gi_spatial_and_shade_pipeline = @@ -289,13 +320,13 @@ impl FromWorld for SolariLightingNode { range: 0..8, }], shader: load_embedded_asset!(world, "restir_gi.wgsl"), - shader_defs: vec![], entry_point: Some("spatial_and_shade".into()), - zero_initialize_workgroup_memory: false, + ..default() }); Self { bind_group_layout, + presample_light_tiles_pipeline, di_initial_and_temporal_pipeline, di_spatial_and_shade_pipeline, gi_initial_and_temporal_pipeline, diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 46a94a3ca2477..0aee54816ed12 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -17,15 +17,26 @@ use bevy_render::{ renderer::RenderDevice, }; -/// Size of a DI Reservoir shader struct in bytes. -const DI_RESERVOIR_STRUCT_SIZE: u64 = 32; +/// Size of the LightSample shader struct in bytes. +const LIGHT_SAMPLE_STRUCT_SIZE: u64 = 8; -/// Size of a GI Reservoir shader struct in bytes. +/// Size of the ResolvedLightSamplePacked shader struct in bytes. +const RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE: u64 = 24; + +/// Size of the DI Reservoir shader struct in bytes. +const DI_RESERVOIR_STRUCT_SIZE: u64 = 16; + +/// 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; + /// 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: Buffer, pub di_reservoirs_b: Buffer, pub gi_reservoirs_a: Buffer, @@ -52,6 +63,22 @@ 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_a = render_device.create_buffer(&BufferDescriptor { label: Some("solari_lighting_di_reservoirs_a"), size: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE, @@ -105,6 +132,8 @@ pub fn prepare_solari_lighting_resources( let previous_depth_view = previous_depth.create_view(&TextureViewDescriptor::default()); commands.entity(entity).insert(SolariLightingResources { + light_tile_samples, + light_tile_resolved_samples, di_reservoirs_a, di_reservoirs_b, gi_reservoirs_a, diff --git a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl new file mode 100644 index 0000000000000..3d2574333231f --- /dev/null +++ b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl @@ -0,0 +1,98 @@ +#define_import_path bevy_solari::presample_light_tiles + +#import bevy_pbr::rgb9e5::{vec3_to_rgb9e5_, rgb9e5_to_vec3_} +#import bevy_pbr::utils::{rand_u, rand_range_u, octahedral_encode, octahedral_decode} +#import bevy_render::view::View +#import bevy_solari::sampling::{LightSample, ResolvedLightSample, triangle_barycentrics} +#import bevy_solari::scene_bindings::{light_sources, directional_lights, resolve_triangle_data_full, LIGHT_SOURCE_KIND_DIRECTIONAL} + +@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 light_count = arrayLength(&light_sources); + let light_id = rand_range_u(light_count, &rng); + let seed = rand_u(&rng); + + let light_source = light_sources[light_id]; + + var triangle_id = 0u; + if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { + let triangle_count = light_source.kind >> 1u; + triangle_id = seed % triangle_count; + } + + let light_sample = LightSample((light_id << 16u) | triangle_id, seed); + + var resolved_light_sample: ResolvedLightSample; + if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { + // TODO: Add support for DIRECTIONAL_LIGHT_SOFT_SHADOWS + let directional_light = directional_lights[light_source.id]; + + resolved_light_sample = ResolvedLightSample( + vec4(directional_light.direction_to_light, 0.0), + -directional_light.direction_to_light, + directional_light.luminance, + directional_light.inverse_pdf, + ); + } else { + let triangle_count = light_source.kind >> 1u; + let barycentrics = triangle_barycentrics(seed); + let triangle_data = resolve_triangle_data_full(light_source.id, triangle_id, barycentrics); + + resolved_light_sample = ResolvedLightSample( + vec4(triangle_data.world_position, 1.0), + triangle_data.world_normal, + triangle_data.material.emissive.rgb, + f32(triangle_count) * triangle_data.triangle_area, + ); + } + resolved_light_sample.inverse_pdf *= f32(light_count); + + let i = (tile_id * 1024u) + sample_index; + light_tile_samples[i] = light_sample; + light_tile_resolved_samples[i] = pack_resolved_light_sample(resolved_light_sample); +} + +struct ResolvedLightSample { + world_position: vec4, + world_normal: vec3, + radiance: vec3, + inverse_pdf: f32, +} + +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/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index b9a5bfa60cccf..5f04c10cd6c39 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -4,22 +4,25 @@ #import bevy_pbr::pbr_deferred_types::unpack_24bit_normal #import bevy_pbr::prepass_bindings::PreviousViewUniforms #import bevy_pbr::rgb9e5::rgb9e5_to_vec3_ -#import bevy_pbr::utils::{rand_f, octahedral_decode} +#import bevy_pbr::utils::{rand_f, rand_range_u, octahedral_decode} #import bevy_render::maths::PI #import bevy_render::view::View -#import bevy_solari::sampling::{LightSample, generate_random_light_sample, calculate_light_contribution, trace_light_visibility, sample_disk} +#import bevy_solari::presample_light_tiles::{ResolvedLightSamplePacked, unpack_resolved_light_sample} +#import bevy_solari::sampling::{LightSample, calculate_light_contribution, trace_light_visibility, sample_disk} #import bevy_solari::scene_bindings::{previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME} @group(1) @binding(0) var view_output: texture_storage_2d; -@group(1) @binding(1) var di_reservoirs_a: array; -@group(1) @binding(2) var di_reservoirs_b: array; -@group(1) @binding(5) var gbuffer: texture_2d; -@group(1) @binding(6) var depth_buffer: texture_depth_2d; -@group(1) @binding(7) var motion_vectors: texture_2d; -@group(1) @binding(8) var previous_gbuffer: texture_2d; -@group(1) @binding(9) var previous_depth_buffer: texture_depth_2d; -@group(1) @binding(10) var view: View; -@group(1) @binding(11) var previous_view: PreviousViewUniforms; +@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: array; +@group(1) @binding(4) var di_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; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; @@ -30,7 +33,7 @@ const CONFIDENCE_WEIGHT_CAP = 20.0; const NULL_RESERVOIR_SAMPLE = 0xFFFFFFFFu; @compute @workgroup_size(8, 8, 1) -fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { +fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) global_id: vec3) { if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } let pixel_index = global_id.x + global_id.y * u32(view.viewport.z); @@ -47,7 +50,7 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); let diffuse_brdf = base_color / PI; - let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng); + let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng, workgroup_id.xy); 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); @@ -81,37 +84,49 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { di_reservoirs_a[pixel_index] = combined_reservoir; - var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight * combined_reservoir.visibility; + 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, rng: ptr) -> Reservoir{ +fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr, workgroup_id: vec2) -> 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 reservoir_target_function = 0.0; - for (var i = 0u; i < INITIAL_SAMPLES; i++) { - let light_sample = generate_random_light_sample(rng); - - let mis_weight = 1.0 / f32(INITIAL_SAMPLES); - let light_contribution = calculate_light_contribution(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); - - reservoir.weight_sum += resampling_weight; - - if rand_f(rng) < resampling_weight / reservoir.weight_sum { - reservoir.sample = light_sample; + var weight_sum = 0.0; + let mis_weight = 1.0 / f32(INITIAL_SAMPLES); + var start = light_tile_start + rand_range_u(1024u - INITIAL_SAMPLES + 1u, rng); + for (var i = start; i < start + INITIAL_SAMPLES; i++) { + let resolved_light_sample = unpack_resolved_light_sample(light_tile_resolved_samples[i], view.exposure); + + let ray = resolved_light_sample.world_position.xyz - (resolved_light_sample.world_position.w * world_position); + let light_distance = length(ray); + let ray_direction = ray / light_distance; + let cos_theta_origin = saturate(dot(ray_direction, world_normal)); + let cos_theta_light = saturate(dot(-ray_direction, resolved_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 target_function = luminance(radiance * diffuse_brdf); + let resampling_weight = mis_weight * (target_function * resolved_light_sample.inverse_pdf); + + weight_sum += resampling_weight; + + if rand_f(rng) < resampling_weight / weight_sum { + reservoir.sample = light_tile_samples[i]; reservoir_target_function = target_function; } } 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 = reservoir.weight_sum * inverse_target_function; + reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - reservoir.visibility = trace_light_visibility(reservoir.sample, world_position); + reservoir.unbiased_contribution_weight *= trace_light_visibility(reservoir.sample, world_position); } reservoir.confidence_weight = 1.0; @@ -142,10 +157,13 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 var temporal_reservoir = di_reservoirs_a[temporal_pixel_index]; // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) - temporal_reservoir.sample.light_id.x = previous_frame_light_id_translations[temporal_reservoir.sample.light_id.x]; - if temporal_reservoir.sample.light_id.x == LIGHT_NOT_PRESENT_THIS_FRAME { + 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); @@ -167,7 +185,7 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< var spatial_reservoir = di_reservoirs_b[spatial_pixel_index]; if reservoir_valid(spatial_reservoir) { - spatial_reservoir.visibility = trace_light_visibility(spatial_reservoir.sample, world_position); + spatial_reservoir.unbiased_contribution_weight *= trace_light_visibility(spatial_reservoir.sample, world_position); } return spatial_reservoir; @@ -216,24 +234,20 @@ fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { // Don't adjust the size of this struct without also adjusting DI_RESERVOIR_STRUCT_SIZE. struct Reservoir { sample: LightSample, - weight_sum: f32, confidence_weight: f32, unbiased_contribution_weight: f32, - visibility: f32, } fn empty_reservoir() -> Reservoir { return Reservoir( - LightSample(vec2(NULL_RESERVOIR_SAMPLE, 0u), vec2(0.0)), + LightSample(NULL_RESERVOIR_SAMPLE, 0u), 0.0, 0.0, - 0.0, - 0.0 ); } fn reservoir_valid(reservoir: Reservoir) -> bool { - return reservoir.sample.light_id.x != NULL_RESERVOIR_SAMPLE; + return reservoir.sample.light_id != NULL_RESERVOIR_SAMPLE; } struct ReservoirMergeResult { @@ -260,26 +274,23 @@ fn merge_reservoirs( 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.weight_sum = canonical_resampling_weight + other_resampling_weight; combined_reservoir.confidence_weight = canonical_reservoir.confidence_weight + other_reservoir.confidence_weight; - // https://yusuketokuyoshi.com/papers/2024/Efficient_Visibility_Reuse_for_Real-time_ReSTIR_(Supplementary_Document).pdf - combined_reservoir.visibility = max(0.0, (canonical_reservoir.visibility * canonical_resampling_weight - + other_reservoir.visibility * other_resampling_weight) / combined_reservoir.weight_sum); - - if rand_f(rng) < other_resampling_weight / combined_reservoir.weight_sum { + 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 = combined_reservoir.weight_sum * inverse_target_function; + 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 = combined_reservoir.weight_sum * inverse_target_function; + combined_reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; return ReservoirMergeResult(combined_reservoir, canonical_target_function.rgb); } @@ -287,7 +298,7 @@ fn merge_reservoirs( 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 = calculate_light_contribution(reservoir.sample, world_position, world_normal).radiance * reservoir.visibility; + let light_contribution = 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 2b0cff5de751b..3fd83d0587519 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::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var view_output: texture_storage_2d; -@group(1) @binding(3) var gi_reservoirs_a: array; -@group(1) @binding(4) var gi_reservoirs_b: array; -@group(1) @binding(5) var gbuffer: texture_2d; -@group(1) @binding(6) var depth_buffer: texture_depth_2d; -@group(1) @binding(7) var motion_vectors: texture_2d; -@group(1) @binding(8) var previous_gbuffer: texture_2d; -@group(1) @binding(9) var previous_depth_buffer: texture_depth_2d; -@group(1) @binding(10) var view: View; -@group(1) @binding(11) var previous_view: PreviousViewUniforms; +@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; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index e1f67ac1ed102..291e4d48576ce 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -1,6 +1,6 @@ #define_import_path bevy_solari::sampling -#import bevy_pbr::utils::{rand_f, rand_vec2f, rand_range_u} +#import bevy_pbr::utils::{rand_f, rand_vec2f, rand_u, rand_range_u} #import bevy_render::maths::{PI, PI_2} #import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full} @@ -61,8 +61,8 @@ fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rn } struct LightSample { - light_id: vec2, - random: vec2, + light_id: u32, + seed: u32, } struct LightContribution { @@ -73,21 +73,21 @@ struct LightContribution { fn generate_random_light_sample(rng: ptr) -> LightSample { let light_count = arrayLength(&light_sources); let light_id = rand_range_u(light_count, rng); - let random = rand_vec2f(rng); + let seed = rand_u(rng); let light_source = light_sources[light_id]; - var triangle_id = 0u; + var triangle_id = 0u; if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { let triangle_count = light_source.kind >> 1u; - triangle_id = rand_range_u(triangle_count, rng); + triangle_id = seed % triangle_count; } - return LightSample(vec2(light_id, triangle_id), random); + return LightSample((light_id << 16u) | triangle_id, seed); } fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { - let light_id = light_sample.light_id.x; + let light_id = light_sample.light_id >> 16u; let light_source = light_sources[light_id]; var light_contribution: LightContribution; @@ -110,9 +110,11 @@ fn calculate_directional_light_contribution(light_sample: LightSample, direction #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 - let cos_theta = (1.0 - light_sample.random.x) + light_sample.random.x * directional_light.cos_theta_max; + var rng = light_sample.seed; + 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 = light_sample.random.y * PI_2; + let phi = random.y * PI_2; let x = cos(phi) * sin_theta; let y = sin(phi) * sin_theta; var ray_direction = vec3(x, y, cos_theta); @@ -130,8 +132,8 @@ fn calculate_directional_light_contribution(light_sample: LightSample, direction } fn calculate_emissive_mesh_contribution(light_sample: LightSample, instance_id: u32, triangle_count: u32, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { - let barycentrics = triangle_barycentrics(light_sample.random); - let triangle_id = light_sample.light_id.y; + let barycentrics = triangle_barycentrics(light_sample.seed); + let triangle_id = light_sample.light_id & 0xFFFFu; let triangle_data = resolve_triangle_data_full(instance_id, triangle_id, barycentrics); @@ -148,7 +150,7 @@ fn calculate_emissive_mesh_contribution(light_sample: LightSample, instance_id: } fn trace_light_visibility(light_sample: LightSample, ray_origin: vec3) -> f32 { - let light_id = light_sample.light_id.x; + let light_id = light_sample.light_id >> 16u; let light_source = light_sources[light_id]; if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { @@ -164,9 +166,11 @@ fn trace_directional_light_visibility(light_sample: LightSample, directional_lig #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 - let cos_theta = (1.0 - light_sample.random.x) + light_sample.random.x * directional_light.cos_theta_max; + var rng = light_sample.seed; + 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 = light_sample.random.y * PI_2; + let phi = random.y * PI_2; let x = cos(phi) * sin_theta; let y = sin(phi) * sin_theta; var ray_direction = vec3(x, y, cos_theta); @@ -182,8 +186,8 @@ fn trace_directional_light_visibility(light_sample: LightSample, directional_lig } fn trace_emissive_mesh_visibility(light_sample: LightSample, instance_id: u32, ray_origin: vec3) -> f32 { - let barycentrics = triangle_barycentrics(light_sample.random); - let triangle_id = light_sample.light_id.y; + let barycentrics = triangle_barycentrics(light_sample.seed); + let triangle_id = light_sample.light_id & 0xFFFFu; let triangle_data = resolve_triangle_data_full(instance_id, triangle_id, barycentrics); @@ -203,8 +207,9 @@ fn trace_point_visibility(ray_origin: vec3, point: vec3) -> f32 { } // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec22%3A297 -fn triangle_barycentrics(random: vec2) -> vec3 { - var barycentrics = random; +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 c474114f3acaec165af57c96faecf90a49721ff1 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 18 Jul 2025 20:19:14 -0400 Subject: [PATCH 08/31] Remove extraneous h parameter for D_GGX --- crates/bevy_pbr/src/render/pbr_lighting.wgsl | 8 +++----- crates/bevy_solari/src/scene/brdf.wgsl | 2 +- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/crates/bevy_pbr/src/render/pbr_lighting.wgsl b/crates/bevy_pbr/src/render/pbr_lighting.wgsl index 6bc24d8af01d6..09329b900750b 100644 --- a/crates/bevy_pbr/src/render/pbr_lighting.wgsl +++ b/crates/bevy_pbr/src/render/pbr_lighting.wgsl @@ -137,7 +137,7 @@ fn getDistanceAttenuation(distanceSquare: f32, inverseRangeSquared: f32) -> f32 // Simple implementation, has precision problems when using fp16 instead of fp32 // see https://google.github.io/filament/Filament.html#listing_speculardfp16 -fn D_GGX(roughness: f32, NdotH: f32, h: vec3) -> f32 { +fn D_GGX(roughness: f32, NdotH: f32) -> f32 { let oneMinusNdotHSquared = 1.0 - NdotH * NdotH; let a = NdotH * roughness; let k = roughness / (oneMinusNdotHSquared + a * a); @@ -313,13 +313,12 @@ fn specular( let roughness = (*input).layers[LAYER_BASE].roughness; let NdotV = (*input).layers[LAYER_BASE].NdotV; let F0 = (*input).F0_; - let H = (*derived_input).H; let NdotL = (*derived_input).NdotL; let NdotH = (*derived_input).NdotH; let LdotH = (*derived_input).LdotH; // Calculate distribution. - let D = D_GGX(roughness, NdotH, H); + let D = D_GGX(roughness, NdotH); // Calculate visibility. let V = V_SmithGGXCorrelated(roughness, NdotV, NdotL); // Calculate the Fresnel term. @@ -343,12 +342,11 @@ fn specular_clearcoat( ) -> vec2 { // Unpack. let roughness = (*input).layers[LAYER_CLEARCOAT].roughness; - let H = (*derived_input).H; let NdotH = (*derived_input).NdotH; let LdotH = (*derived_input).LdotH; // Calculate distribution. - let Dc = D_GGX(roughness, NdotH, H); + let Dc = D_GGX(roughness, NdotH); // Calculate visibility. let Vc = V_Kelemen(LdotH); // Calculate the Fresnel term. diff --git a/crates/bevy_solari/src/scene/brdf.wgsl b/crates/bevy_solari/src/scene/brdf.wgsl index 13ab7b46d4574..acba5c2792e77 100644 --- a/crates/bevy_solari/src/scene/brdf.wgsl +++ b/crates/bevy_solari/src/scene/brdf.wgsl @@ -50,7 +50,7 @@ fn specular_brdf( let F_ab = F_AB(perceptual_roughness, NdotV); let roughness = perceptualRoughnessToRoughness(perceptual_roughness); - let D = D_GGX(roughness, NdotH, H); + let D = D_GGX(roughness, NdotH); let Vs = V_SmithGGXCorrelated(roughness, NdotV, NdotL); let F = fresnel(F0, LdotH); return specular_multiscatter(D, Vs, F, F0, F_ab, 1.0); From b143c1046dc7ffa3c370a1b940ac6241d882fd7f Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 18 Jul 2025 20:44:02 -0400 Subject: [PATCH 09/31] Add VNDF sampling routines --- crates/bevy_solari/src/scene/sampling.wgsl | 37 ++++++++++++++++++++++ 1 file changed, 37 insertions(+) diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index e1f67ac1ed102..15e2c40383ed0 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -1,5 +1,6 @@ #define_import_path bevy_solari::sampling +#import bevy_pbr::lighting::D_GGX #import bevy_pbr::utils::{rand_f, rand_vec2f, rand_range_u} #import bevy_render::maths::{PI, PI_2} #import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full} @@ -26,6 +27,42 @@ fn sample_uniform_hemisphere(normal: vec3, rng: ptr) -> vec3 return build_orthonormal_basis(normal) * vec3(x, y, z); } +// https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 1) +fn sample_ggx_vndf(i: vec3, roughness: f32, rng: ptr) -> vec3 { + let rand = rand_vec2f(rng); + let i_std = normalize(vec3(i.xy * roughness, i.z)); + let phi = 2.0 * PI * rand.x; + let a = saturate(roughness); + let s = 1.0 + length(vec2(i.xy)); + let a2 = a * a; + let s2 = s * s; + let k = (1.0 - a2) * s2 / (s2 + a2 * i.z * i.z); + let b = select(i_std.z, k * i_std.z, i.z > 0.0); + let z = fma(1.0 - rand.y, 1.0 + b, -b); + let sin_theta = sqrt(saturate(1.0 - z * z)); + let o_std = vec3(sin_theta * cos(phi), sin_theta * sin(phi), z); + let m_std = i_std + o_std; + let m = normalize(vec3(m_std.xy * roughness, m_std.z)); + return 2.0 * dot(i, m) * m - i; +} + +// https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 2) +fn ggx_vndf_pdf(i: vec3, NdotH: f32, roughness: f32) -> f32 { + let ndf = D_GGX(roughness, NdotH); + let ai = roughness * i.xy; + let len2 = dot(ai, ai); + let t = sqrt(len2 + i.z * i.z); + if i.z >= 0.0 { + let a = saturate(roughness); + let s = 1.0 + length(i.xy); + let a2 = a * a; + let s2 = s * s; + let k = (1.0 - a2) * s2 / (s2 + a2 * i.z * i.z); + return ndf / (2.0 * (k * i.z + t)); + } + return ndf * (t - i.z) / (2.0 * len2); +} + // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec19%3A294 fn sample_disk(disk_radius: f32, rng: ptr) -> vec2 { let ab = 2.0 * rand_vec2f(rng) - 1.0; From b4a2f77edae7558d07f1e4a6a4aeedbba3030e12 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 18 Jul 2025 21:34:35 -0400 Subject: [PATCH 10/31] WIP specular pathtracer (removed NEE for now) --- .../src/pathtracer/pathtracer.wgsl | 65 +++++++++++-------- crates/bevy_solari/src/scene/brdf.wgsl | 21 +++--- .../src/scene/raytracing_scene_bindings.wgsl | 4 ++ crates/bevy_solari/src/scene/sampling.wgsl | 5 +- 4 files changed, 54 insertions(+), 41 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 15f6117bacf3e..9c1e93c728c3b 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -3,8 +3,8 @@ #import bevy_render::maths::PI #import bevy_render::view::View #import bevy_solari::brdf::evaluate_brdf -#import bevy_solari::sampling::{sample_random_light, sample_uniform_hemisphere} -#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} +#import bevy_solari::sampling::{sample_cosine_hemisphere, sample_ggx_vndf, ggx_vndf_pdf} +#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; @group(1) @binding(1) var view_output: texture_storage_2d; @@ -40,38 +40,23 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { 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); - - // Sample new ray direction for the next bounce let wo = -ray_direction; - let wi = sample_uniform_hemisphere(ray_hit.world_normal, &rng); - ray_direction = wi; - // Evaluate material BRDF - let brdf = evaluate_brdf( - ray_hit.world_normal, - wo, - wi, - ray_hit.material.base_color, - ray_hit.material.metallic, - ray_hit.material.reflectance, - ray_hit.material.perceptual_roughness, - ); - - // Use emissive only on the first ray (coming from the camera) - if ray_t_min == 0.0 { radiance = ray_hit.material.emissive; } - - // Sample direct lighting - // TODO: Wrong BRDF here when it comes to specular? - radiance += throughput * brdf * sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); - - // Update other variables for next bounce + // Calculate emissive contribution + radiance += throughput * ray_hit.material.emissive; + + // Sample new ray direction from the material BRDF for next bounce + let next_bounce = importance_sample_next_bounce(wo, ray_hit, &rng); + ray_direction = next_bounce.wi; ray_origin = ray_hit.world_position; ray_t_min = RAY_T_MIN; + // Evaluate material BRDF + let brdf = evaluate_brdf(ray_hit.world_normal, wo, next_bounce.wi, ray_hit.material); + // Update throughput for next bounce - let cos_theta = dot(wi, ray_hit.world_normal); - let uniform_hemisphere_pdf = 1.0 / (2.0 * PI); // Weight for the next bounce because we importance sampled the diffuse BRDF for the next ray direction - throughput *= (brdf * cos_theta) / uniform_hemisphere_pdf; + let cos_theta = dot(next_bounce.wi, ray_hit.world_normal); + throughput *= (brdf * cos_theta) / next_bounce.pdf; // Russian roulette for early termination let p = luminance(throughput); @@ -88,3 +73,27 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { textureStore(accumulation_texture, global_id.xy, vec4(new_color, old_color.a + 1.0)); textureStore(view_output, global_id.xy, vec4(new_color, 1.0)); } + +struct NextBounce { + wi: vec3, + pdf: f32, +} + +fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng: ptr) -> NextBounce { + let diffuse_weight = 1.0 - ray_hit.material.metallic; + let specular_weight = ray_hit.material.metallic; + let total_weight = diffuse_weight + specular_weight; + + var wi: vec3; + if rand_f(rng) * total_weight < diffuse_weight { + wi = sample_cosine_hemisphere(ray_hit.world_normal, rng); + } else { + wi = sample_ggx_vndf(wo, ray_hit.material.roughness, rng); + } + + let diffuse_pdf = dot(wi, ray_hit.world_normal) / PI; + let specular_pdf = ggx_vndf_pdf(wo, wi, ray_hit.world_normal, ray_hit.material.roughness); + let pdf = ((diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf)) / total_weight; + + return NextBounce(wi, pdf); +} diff --git a/crates/bevy_solari/src/scene/brdf.wgsl b/crates/bevy_solari/src/scene/brdf.wgsl index acba5c2792e77..bc42203481928 100644 --- a/crates/bevy_solari/src/scene/brdf.wgsl +++ b/crates/bevy_solari/src/scene/brdf.wgsl @@ -1,27 +1,26 @@ #define_import_path bevy_solari::brdf -#import bevy_pbr::lighting::{F_AB, perceptualRoughnessToRoughness, D_GGX, V_SmithGGXCorrelated, fresnel, specular_multiscatter} +#import bevy_pbr::lighting::{F_AB, D_GGX, V_SmithGGXCorrelated, fresnel, specular_multiscatter} #import bevy_pbr::pbr_functions::{calculate_diffuse_color, calculate_F0} #import bevy_render::maths::PI +#import bevy_solari::scene_bindings::ResolvedMaterial fn evaluate_brdf( world_normal: vec3, wo: vec3, wi: vec3, - base_color: vec3, - metallic: f32, - reflectance: vec3, - perceptual_roughness: f32, + material: ResolvedMaterial, ) -> vec3 { - let diffuse_brdf = diffuse_brdf(base_color, metallic); + let diffuse_brdf = diffuse_brdf(material.base_color, material.metallic); let specular_brdf = specular_brdf( world_normal, wo, wi, - base_color, - metallic, - reflectance, - perceptual_roughness, + material.base_color, + material.metallic, + material.reflectance, + material.perceptual_roughness, + material.roughness, ); return diffuse_brdf + specular_brdf; } @@ -39,6 +38,7 @@ fn specular_brdf( metallic: f32, reflectance: vec3, perceptual_roughness: f32, + roughness: f32, ) -> vec3 { let H = normalize(L + V); let NdotL = saturate(dot(N, L)); @@ -48,7 +48,6 @@ fn specular_brdf( let F0 = calculate_F0(base_color, metallic, reflectance); let F_ab = F_AB(perceptual_roughness, NdotV); - let roughness = perceptualRoughnessToRoughness(perceptual_roughness); let D = D_GGX(roughness, NdotH); let Vs = V_SmithGGXCorrelated(roughness, NdotV, NdotL); diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index a59b38e952e05..1b135a8535740 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -1,5 +1,7 @@ #define_import_path bevy_solari::scene_bindings +#import bevy_pbr::lighting::perceptualRoughnessToRoughness + struct InstanceGeometryIds { vertex_buffer_id: u32, vertex_buffer_offset: u32, @@ -101,6 +103,7 @@ struct ResolvedMaterial { emissive: vec3, reflectance: vec3, perceptual_roughness: f32, + roughness: f32, metallic: f32, } @@ -135,6 +138,7 @@ fn resolve_material(material: Material, uv: vec2) -> ResolvedMaterial { m.perceptual_roughness *= metallic_roughness.g; m.metallic *= metallic_roughness.b; } + m.roughness = perceptualRoughnessToRoughness(m.perceptual_roughness); return m; } diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 15e2c40383ed0..e45e7ccc3fda2 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -47,8 +47,9 @@ fn sample_ggx_vndf(i: vec3, roughness: f32, rng: ptr) -> vec } // https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 2) -fn ggx_vndf_pdf(i: vec3, NdotH: f32, roughness: f32) -> f32 { - let ndf = D_GGX(roughness, NdotH); +fn ggx_vndf_pdf(i: vec3, o: vec3, world_normal: vec3, roughness: f32) -> f32 { + let m = normalize(i + o); + let ndf = D_GGX(roughness, saturate(dot(world_normal, m))); let ai = roughness * i.xy; let len2 = dot(ai, ai); let t = sqrt(len2 + i.z * i.z); From 09f36c4f93ff3fce5ee50df7f5499e0954ed5cf7 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Fri, 18 Jul 2025 21:54:15 -0400 Subject: [PATCH 11/31] Diffuse/specular weights always sum to 1 --- crates/bevy_solari/src/pathtracer/pathtracer.wgsl | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 9c1e93c728c3b..6611e31f01e9f 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -82,10 +82,9 @@ struct NextBounce { fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng: ptr) -> NextBounce { let diffuse_weight = 1.0 - ray_hit.material.metallic; let specular_weight = ray_hit.material.metallic; - let total_weight = diffuse_weight + specular_weight; var wi: vec3; - if rand_f(rng) * total_weight < diffuse_weight { + if rand_f(rng) < diffuse_weight { wi = sample_cosine_hemisphere(ray_hit.world_normal, rng); } else { wi = sample_ggx_vndf(wo, ray_hit.material.roughness, rng); @@ -93,7 +92,7 @@ fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng let diffuse_pdf = dot(wi, ray_hit.world_normal) / PI; let specular_pdf = ggx_vndf_pdf(wo, wi, ray_hit.world_normal, ray_hit.material.roughness); - let pdf = ((diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf)) / total_weight; + let pdf = (diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf); return NextBounce(wi, pdf); } From 82519a5c9f93d5e94dbe9e2b04e29750b495babc Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 19 Jul 2025 11:19:04 -0400 Subject: [PATCH 12/31] Add world_tangent to ResolvedRayHitFull --- .../src/scene/raytracing_scene_bindings.wgsl | 19 +++++++++++++------ 1 file changed, 13 insertions(+), 6 deletions(-) diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index 1b135a8535740..4fc8af56ee1bd 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -1,6 +1,7 @@ #define_import_path bevy_solari::scene_bindings #import bevy_pbr::lighting::perceptualRoughnessToRoughness +#import bevy_pbr::pbr_functions::calculate_tbn_mikktspace struct InstanceGeometryIds { vertex_buffer_id: u32, @@ -111,6 +112,7 @@ struct ResolvedRayHitFull { world_position: vec3, world_normal: vec3, geometric_world_normal: vec3, + world_tangent: vec4, uv: vec2, triangle_area: f32, material: ResolvedMaterial, @@ -166,15 +168,20 @@ fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, 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 var world_normal = normalize(mat3x3(transform[0].xyz, transform[1].xyz, transform[2].xyz) * local_normal); let geometric_world_normal = world_normal; if material.normal_map_texture_id != TEXTURE_MAP_NONE { - let local_tangent = mat3x3(vertices[0].tangent.xyz, vertices[1].tangent.xyz, vertices[2].tangent.xyz) * barycentrics; - let world_tangent = normalize(mat3x3(transform[0].xyz, transform[1].xyz, transform[2].xyz) * local_tangent); - let N = world_normal; - let T = world_tangent; - let B = vertices[0].tangent.w * cross(N, T); + let TBN = calculate_tbn_mikktspace(world_normal, world_tangent); + let T = TBN[0]; + let B = TBN[1]; + let N = TBN[2]; let Nt = sample_texture(material.normal_map_texture_id, uv); world_normal = normalize(Nt.x * T + Nt.y * B + Nt.z * N); } @@ -185,5 +192,5 @@ fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: let resolved_material = resolve_material(material, uv); - return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, uv, triangle_area, resolved_material); + return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, world_tangent, uv, triangle_area, resolved_material); } From 02841d7d46f52c09d04d86526580c9c36e911826 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 19 Jul 2025 11:36:57 -0400 Subject: [PATCH 13/31] Fix specular --- crates/bevy_solari/src/pathtracer/pathtracer.wgsl | 15 +++++++++++++-- crates/bevy_solari/src/scene/sampling.wgsl | 9 ++++++--- 2 files changed, 19 insertions(+), 5 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 6611e31f01e9f..0300a3271ff7e 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -1,4 +1,5 @@ #import bevy_core_pipeline::tonemapping::tonemapping_luminance as luminance +#import bevy_pbr::pbr_functions::calculate_tbn_mikktspace #import bevy_pbr::utils::{rand_f, rand_vec2f} #import bevy_render::maths::PI #import bevy_render::view::View @@ -83,15 +84,25 @@ fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng let diffuse_weight = 1.0 - ray_hit.material.metallic; let specular_weight = ray_hit.material.metallic; + let TBN = calculate_tbn_mikktspace(ray_hit.world_normal, ray_hit.world_tangent); + let T = TBN[0]; + let B = TBN[1]; + let N = TBN[2]; + + let wo_tangent = vec3(dot(wo, T), dot(wo, B), dot(wo, N)); + var wi: vec3; + var wi_tangent: vec3; if rand_f(rng) < diffuse_weight { wi = sample_cosine_hemisphere(ray_hit.world_normal, rng); + wi_tangent = vec3(dot(wi, T), dot(wi, B), dot(wi, N)); } else { - wi = sample_ggx_vndf(wo, ray_hit.material.roughness, rng); + wi_tangent = sample_ggx_vndf(wo_tangent, ray_hit.material.roughness, rng); + wi = wi_tangent.x * T + wi_tangent.y * B + wi_tangent.z * N; } let diffuse_pdf = dot(wi, ray_hit.world_normal) / PI; - let specular_pdf = ggx_vndf_pdf(wo, wi, ray_hit.world_normal, ray_hit.material.roughness); + let specular_pdf = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); let pdf = (diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf); return NextBounce(wi, pdf); diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index e45e7ccc3fda2..b4c24e33babed 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -28,7 +28,8 @@ fn sample_uniform_hemisphere(normal: vec3, rng: ptr) -> vec3 } // https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 1) -fn sample_ggx_vndf(i: vec3, roughness: f32, rng: ptr) -> vec3 { +fn sample_ggx_vndf(wi_tangent: vec3, roughness: f32, rng: ptr) -> vec3 { + let i = wi_tangent; let rand = rand_vec2f(rng); let i_std = normalize(vec3(i.xy * roughness, i.z)); let phi = 2.0 * PI * rand.x; @@ -47,9 +48,11 @@ fn sample_ggx_vndf(i: vec3, roughness: f32, rng: ptr) -> vec } // https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 2) -fn ggx_vndf_pdf(i: vec3, o: vec3, world_normal: vec3, roughness: f32) -> f32 { +fn ggx_vndf_pdf(wi_tangent: vec3, wo_tangent: vec3, roughness: f32) -> f32 { + let i = wi_tangent; + let o = wo_tangent; let m = normalize(i + o); - let ndf = D_GGX(roughness, saturate(dot(world_normal, m))); + let ndf = D_GGX(roughness, saturate(m.z)); let ai = roughness * i.xy; let len2 = dot(ai, ai); let t = sqrt(len2 + i.z * i.z); From 3ebb7d4d72f9ed03aa3d072abcf1763860184190 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 19 Jul 2025 12:07:44 -0400 Subject: [PATCH 14/31] Allow smoother specular materials --- crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index 4fc8af56ee1bd..6582100711836 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -140,7 +140,7 @@ fn resolve_material(material: Material, uv: vec2) -> ResolvedMaterial { m.perceptual_roughness *= metallic_roughness.g; m.metallic *= metallic_roughness.b; } - m.roughness = perceptualRoughnessToRoughness(m.perceptual_roughness); + m.roughness = clamp(m.perceptual_roughness * m.perceptual_roughness, 0.001, 1.0); return m; } From 3df77e2d6ba84ee375746d95cac1de80a1c5e16c Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 19 Jul 2025 12:16:03 -0400 Subject: [PATCH 15/31] Remove uneeded saturates --- crates/bevy_solari/src/scene/sampling.wgsl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index b4c24e33babed..2c34ed1df0f21 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -33,7 +33,7 @@ fn sample_ggx_vndf(wi_tangent: vec3, roughness: f32, rng: ptr, wo_tangent: vec3, roughness: f32) -> let len2 = dot(ai, ai); let t = sqrt(len2 + i.z * i.z); if i.z >= 0.0 { - let a = saturate(roughness); + let a = roughness; let s = 1.0 + length(i.xy); let a2 = a * a; let s2 = s * s; From 6f9b1b5ac4bbb78c8039d02c7d62d75e99ddad06 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sat, 19 Jul 2025 12:30:39 -0400 Subject: [PATCH 16/31] Add back NEE --- .../src/pathtracer/pathtracer.wgsl | 14 ++++++----- crates/bevy_solari/src/scene/sampling.wgsl | 25 +++++++++++-------- 2 files changed, 23 insertions(+), 16 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 0300a3271ff7e..881ff7d1f069c 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_cosine_hemisphere, sample_ggx_vndf, ggx_vndf_pdf} +#import bevy_solari::sampling::{sample_random_light, sample_cosine_hemisphere, sample_ggx_vndf, ggx_vndf_pdf} #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; @@ -43,8 +43,12 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { let ray_hit = resolve_ray_hit_full(ray_hit); let wo = -ray_direction; - // Calculate emissive contribution - radiance += throughput * ray_hit.material.emissive; + // Use emissive only on the first ray (coming from the camera) + if ray_t_min == 0.0 { radiance = ray_hit.material.emissive; } + + // Sample direct lighting + let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, wo, ray_hit.material, &rng); + radiance += throughput * direct_lighting.radiance * direct_lighting.inverse_pdf; // Sample new ray direction from the material BRDF for next bounce let next_bounce = importance_sample_next_bounce(wo, ray_hit, &rng); @@ -52,10 +56,8 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { ray_origin = ray_hit.world_position; ray_t_min = RAY_T_MIN; - // Evaluate material BRDF - let brdf = evaluate_brdf(ray_hit.world_normal, wo, next_bounce.wi, ray_hit.material); - // Update throughput for next bounce + let brdf = evaluate_brdf(ray_hit.world_normal, wo, next_bounce.wi, ray_hit.material); let cos_theta = dot(next_bounce.wi, ray_hit.world_normal); throughput *= (brdf * cos_theta) / next_bounce.pdf; diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 2c34ed1df0f21..46a77784ea2ad 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -3,7 +3,8 @@ #import bevy_pbr::lighting::D_GGX #import bevy_pbr::utils::{rand_f, rand_vec2f, rand_range_u} #import bevy_render::maths::{PI, PI_2} -#import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full} +#import bevy_solari::brdf::evaluate_brdf +#import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full, ResolvedMaterial} // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec28%3A303 fn sample_cosine_hemisphere(normal: vec3, rng: ptr) -> vec3 { @@ -94,9 +95,9 @@ struct SampleRandomLightResult { inverse_pdf: f32, } -fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rng: ptr) -> SampleRandomLightResult { +fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, wo: vec3, material: ResolvedMaterial, rng: ptr) -> SampleRandomLightResult { let light_sample = generate_random_light_sample(rng); - let light_contribution = calculate_light_contribution(light_sample, ray_origin, origin_world_normal); + let light_contribution = calculate_light_contribution(light_sample, ray_origin, origin_world_normal, wo, material); let visibility = trace_light_visibility(light_sample, ray_origin); return SampleRandomLightResult(light_contribution.radiance * visibility, light_contribution.inverse_pdf); } @@ -127,16 +128,16 @@ fn generate_random_light_sample(rng: ptr) -> LightSample { return LightSample(vec2(light_id, triangle_id), random); } -fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { +fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3, wo: vec3, material: ResolvedMaterial) -> LightContribution { let light_id = light_sample.light_id.x; let light_source = light_sources[light_id]; var light_contribution: LightContribution; if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { - light_contribution = calculate_directional_light_contribution(light_sample, light_source.id, origin_world_normal); + light_contribution = calculate_directional_light_contribution(light_sample, light_source.id, origin_world_normal, wo, material); } else { let triangle_count = light_source.kind >> 1u; - light_contribution = calculate_emissive_mesh_contribution(light_sample, light_source.id, triangle_count, ray_origin, origin_world_normal); + light_contribution = calculate_emissive_mesh_contribution(light_sample, light_source.id, triangle_count, ray_origin, origin_world_normal, wo, material); } let light_count = arrayLength(&light_sources); @@ -145,7 +146,7 @@ fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3 return light_contribution; } -fn calculate_directional_light_contribution(light_sample: LightSample, directional_light_id: u32, origin_world_normal: vec3) -> LightContribution { +fn calculate_directional_light_contribution(light_sample: LightSample, directional_light_id: u32, origin_world_normal: vec3, wo: vec3, material: ResolvedMaterial) -> LightContribution { let directional_light = directional_lights[directional_light_id]; #ifdef DIRECTIONAL_LIGHT_SOFT_SHADOWS @@ -167,10 +168,12 @@ fn calculate_directional_light_contribution(light_sample: LightSample, direction let cos_theta_origin = saturate(dot(ray_direction, origin_world_normal)); let radiance = directional_light.luminance * cos_theta_origin; - return LightContribution(radiance, directional_light.inverse_pdf); + let brdf = evaluate_brdf(origin_world_normal, wo, ray_direction, material); + + return LightContribution(radiance * brdf, directional_light.inverse_pdf); } -fn calculate_emissive_mesh_contribution(light_sample: LightSample, instance_id: u32, triangle_count: u32, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { +fn calculate_emissive_mesh_contribution(light_sample: LightSample, instance_id: u32, triangle_count: u32, ray_origin: vec3, origin_world_normal: vec3, wo: vec3, material: ResolvedMaterial) -> LightContribution { let barycentrics = triangle_barycentrics(light_sample.random); let triangle_id = light_sample.light_id.y; @@ -185,7 +188,9 @@ fn calculate_emissive_mesh_contribution(light_sample: LightSample, instance_id: let radiance = triangle_data.material.emissive.rgb * cos_theta_origin * (cos_theta_light / light_distance_squared); let inverse_pdf = f32(triangle_count) * triangle_data.triangle_area; - return LightContribution(radiance, inverse_pdf); + let brdf = evaluate_brdf(origin_world_normal, wo, ray_direction, material); + + return LightContribution(radiance * brdf, inverse_pdf); } fn trace_light_visibility(light_sample: LightSample, ray_origin: vec3) -> f32 { From 92bc08b17cad0c6537ada9972df5081fa4e5b6a0 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 14:31:59 -0400 Subject: [PATCH 17/31] Refactoring --- crates/bevy_solari/src/realtime/node.rs | 131 ++++-------- .../src/realtime/presample_light_tiles.wgsl | 55 +---- .../bevy_solari/src/realtime/restir_di.wgsl | 39 ++-- .../bevy_solari/src/realtime/restir_gi.wgsl | 2 +- crates/bevy_solari/src/scene/sampling.wgsl | 192 +++++++++--------- release-content/release-notes/bevy_solari.md | 2 +- 6 files changed, 161 insertions(+), 260 deletions(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index 164c0fc9a1d49..b873d7bee567f 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -3,7 +3,7 @@ use super::{ SolariLighting, }; use crate::scene::RaytracingSceneBindings; -use bevy_asset::load_embedded_asset; +use bevy_asset::{load_embedded_asset, Handle}; use bevy_core_pipeline::prepass::{ PreviousViewData, PreviousViewUniformOffset, PreviousViewUniforms, ViewPrepassTextures, }; @@ -22,7 +22,7 @@ use bevy_render::{ storage_buffer_sized, texture_2d, texture_depth_2d, texture_storage_2d, uniform_buffer, }, BindGroupEntries, BindGroupLayout, BindGroupLayoutEntries, CachedComputePipelineId, - ComputePassDescriptor, ComputePipelineDescriptor, PipelineCache, PushConstantRange, + ComputePassDescriptor, ComputePipelineDescriptor, PipelineCache, PushConstantRange, Shader, ShaderStages, StorageTextureAccess, TextureSampleType, }, renderer::{RenderContext, RenderDevice}, @@ -245,92 +245,51 @@ impl FromWorld for SolariLightingNode { ), ); - let presample_light_tiles_pipeline = - pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_presample_light_tiles_pipeline".into()), - layout: vec![ - scene_bindings.bind_group_layout.clone(), - bind_group_layout.clone(), - ], - push_constant_ranges: vec![PushConstantRange { - stages: ShaderStages::COMPUTE, - range: 0..8, - }], - shader: load_embedded_asset!(world, "presample_light_tiles.wgsl"), - ..default() - }); - - let di_initial_and_temporal_pipeline = - pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_di_initial_and_temporal_pipeline".into()), - layout: vec![ - scene_bindings.bind_group_layout.clone(), - bind_group_layout.clone(), - ], - push_constant_ranges: vec![PushConstantRange { - stages: ShaderStages::COMPUTE, - range: 0..8, - }], - shader: load_embedded_asset!(world, "restir_di.wgsl"), - entry_point: Some("initial_and_temporal".into()), - ..default() - }); - - let di_spatial_and_shade_pipeline = - pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_di_spatial_and_shade_pipeline".into()), - layout: vec![ - scene_bindings.bind_group_layout.clone(), - bind_group_layout.clone(), - ], - push_constant_ranges: vec![PushConstantRange { - stages: ShaderStages::COMPUTE, - range: 0..8, - }], - shader: load_embedded_asset!(world, "restir_di.wgsl"), - entry_point: Some("spatial_and_shade".into()), - ..default() - }); - - let gi_initial_and_temporal_pipeline = - pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_gi_initial_and_temporal_pipeline".into()), - layout: vec![ - scene_bindings.bind_group_layout.clone(), - bind_group_layout.clone(), - ], - push_constant_ranges: vec![PushConstantRange { - stages: ShaderStages::COMPUTE, - range: 0..8, - }], - shader: load_embedded_asset!(world, "restir_gi.wgsl"), - entry_point: Some("initial_and_temporal".into()), - ..default() - }); - - let gi_spatial_and_shade_pipeline = - pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { - label: Some("solari_lighting_gi_spatial_and_shade_pipeline".into()), - layout: vec![ - scene_bindings.bind_group_layout.clone(), - bind_group_layout.clone(), - ], - push_constant_ranges: vec![PushConstantRange { - stages: ShaderStages::COMPUTE, - range: 0..8, - }], - shader: load_embedded_asset!(world, "restir_gi.wgsl"), - entry_point: Some("spatial_and_shade".into()), - ..default() - }); + let create_pipeline = + |label: &'static str, entry_point: &'static str, shader: Handle| { + pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { + label: Some(label.into()), + layout: vec![ + scene_bindings.bind_group_layout.clone(), + bind_group_layout.clone(), + ], + push_constant_ranges: vec![PushConstantRange { + stages: ShaderStages::COMPUTE, + range: 0..8, + }], + shader, + entry_point: Some(entry_point.into()), + ..default() + }) + }; Self { - bind_group_layout, - presample_light_tiles_pipeline, - di_initial_and_temporal_pipeline, - di_spatial_and_shade_pipeline, - gi_initial_and_temporal_pipeline, - gi_spatial_and_shade_pipeline, + bind_group_layout: bind_group_layout.clone(), + presample_light_tiles_pipeline: create_pipeline( + "solari_lighting_presample_light_tiles_pipeline", + "presample_light_tiles", + load_embedded_asset!(world, "presample_light_tiles.wgsl"), + ), + di_initial_and_temporal_pipeline: create_pipeline( + "solari_lighting_di_initial_and_temporal_pipeline", + "initial_and_temporal", + load_embedded_asset!(world, "restir_di.wgsl"), + ), + di_spatial_and_shade_pipeline: create_pipeline( + "solari_lighting_di_spatial_and_shade_pipeline", + "spatial_and_shade", + load_embedded_asset!(world, "restir_di.wgsl"), + ), + gi_initial_and_temporal_pipeline: create_pipeline( + "solari_lighting_gi_initial_and_temporal_pipeline", + "initial_and_temporal", + load_embedded_asset!(world, "restir_gi.wgsl"), + ), + gi_spatial_and_shade_pipeline: create_pipeline( + "solari_lighting_gi_spatial_and_shade_pipeline", + "spatial_and_shade", + load_embedded_asset!(world, "restir_gi.wgsl"), + ), } } } diff --git a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl index 3d2574333231f..755b0813df8fd 100644 --- a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl +++ b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl @@ -1,10 +1,9 @@ #define_import_path bevy_solari::presample_light_tiles #import bevy_pbr::rgb9e5::{vec3_to_rgb9e5_, rgb9e5_to_vec3_} -#import bevy_pbr::utils::{rand_u, rand_range_u, octahedral_encode, octahedral_decode} +#import bevy_pbr::utils::{octahedral_encode, octahedral_decode} #import bevy_render::view::View -#import bevy_solari::sampling::{LightSample, ResolvedLightSample, triangle_barycentrics} -#import bevy_solari::scene_bindings::{light_sources, directional_lights, resolve_triangle_data_full, LIGHT_SOURCE_KIND_DIRECTIONAL} +#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; @@ -17,55 +16,11 @@ fn presample_light_tiles(@builtin(workgroup_id) workgroup_id: vec3, @builti let tile_id = workgroup_id.x; var rng = (tile_id * 5782582u) + sample_index + constants.frame_index; - let light_count = arrayLength(&light_sources); - let light_id = rand_range_u(light_count, &rng); - let seed = rand_u(&rng); - - let light_source = light_sources[light_id]; - - var triangle_id = 0u; - if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { - let triangle_count = light_source.kind >> 1u; - triangle_id = seed % triangle_count; - } - - let light_sample = LightSample((light_id << 16u) | triangle_id, seed); - - var resolved_light_sample: ResolvedLightSample; - if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { - // TODO: Add support for DIRECTIONAL_LIGHT_SOFT_SHADOWS - let directional_light = directional_lights[light_source.id]; - - resolved_light_sample = ResolvedLightSample( - vec4(directional_light.direction_to_light, 0.0), - -directional_light.direction_to_light, - directional_light.luminance, - directional_light.inverse_pdf, - ); - } else { - let triangle_count = light_source.kind >> 1u; - let barycentrics = triangle_barycentrics(seed); - let triangle_data = resolve_triangle_data_full(light_source.id, triangle_id, barycentrics); - - resolved_light_sample = ResolvedLightSample( - vec4(triangle_data.world_position, 1.0), - triangle_data.world_normal, - triangle_data.material.emissive.rgb, - f32(triangle_count) * triangle_data.triangle_area, - ); - } - resolved_light_sample.inverse_pdf *= f32(light_count); + let sample = generate_random_light_sample(&rng); let i = (tile_id * 1024u) + sample_index; - light_tile_samples[i] = light_sample; - light_tile_resolved_samples[i] = pack_resolved_light_sample(resolved_light_sample); -} - -struct ResolvedLightSample { - world_position: vec4, - world_normal: vec3, - radiance: vec3, - inverse_pdf: f32, + light_tile_samples[i] = sample.light_sample; + light_tile_resolved_samples[i] = pack_resolved_light_sample(sample.resolved_light_sample); } struct ResolvedLightSamplePacked { diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 5f04c10cd6c39..c3aedb96c214e 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -8,8 +8,8 @@ #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_light_contribution, trace_light_visibility, sample_disk} -#import bevy_solari::scene_bindings::{previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME} +#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility, sample_disk} +#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; @@ -50,7 +50,7 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); let diffuse_brdf = base_color / PI; - let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng, workgroup_id.xy); + 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); @@ -91,34 +91,29 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { textureStore(view_output, global_id.xy, vec4(pixel_color, 1.0)); } -fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr, workgroup_id: vec2) -> Reservoir{ +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 reservoir_target_function = 0.0; + var light_sample_world_position = vec4(0.0); var weight_sum = 0.0; let mis_weight = 1.0 / f32(INITIAL_SAMPLES); - var start = light_tile_start + rand_range_u(1024u - INITIAL_SAMPLES + 1u, rng); - for (var i = start; i < start + INITIAL_SAMPLES; i++) { - let resolved_light_sample = unpack_resolved_light_sample(light_tile_resolved_samples[i], view.exposure); + 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 ray = resolved_light_sample.world_position.xyz - (resolved_light_sample.world_position.w * world_position); - let light_distance = length(ray); - let ray_direction = ray / light_distance; - let cos_theta_origin = saturate(dot(ray_direction, world_normal)); - let cos_theta_light = saturate(dot(-ray_direction, resolved_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 target_function = luminance(radiance * diffuse_brdf); - let resampling_weight = mis_weight * (target_function * resolved_light_sample.inverse_pdf); + 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.sample = light_tile_samples[i]; + reservoir.sample = light_tile_samples[tile_sample]; reservoir_target_function = target_function; + light_sample_world_position = resolved_light_sample.world_position; } } @@ -126,7 +121,7 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 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(reservoir.sample, world_position); + reservoir.unbiased_contribution_weight *= trace_light_visibility(world_position, light_sample_world_position); } reservoir.confidence_weight = 1.0; @@ -185,7 +180,8 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< var spatial_reservoir = di_reservoirs_b[spatial_pixel_index]; if reservoir_valid(spatial_reservoir) { - spatial_reservoir.unbiased_contribution_weight *= trace_light_visibility(spatial_reservoir.sample, world_position); + 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; @@ -296,9 +292,10 @@ fn merge_reservoirs( } } +// 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 = calculate_light_contribution(reservoir.sample, world_position, world_normal).radiance; + 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 3fd83d0587519..2be5e8554dba9 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -7,7 +7,7 @@ #import bevy_pbr::utils::{rand_f, octahedral_decode} #import bevy_render::maths::{PI, PI_2} #import bevy_render::view::View -#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light, sample_disk, trace_point_visibility} +#import bevy_solari::sampling::{sample_uniform_hemisphere, sample_random_light, trace_point_visibility, sample_disk} #import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX} @group(1) @binding(0) var view_output: texture_storage_2d; diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index f500310cf698e..f402c1ea419dc 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -2,7 +2,7 @@ #import bevy_pbr::utils::{rand_f, rand_vec2f, rand_u, rand_range_u} #import bevy_render::maths::{PI, PI_2, orthonormalize} -#import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full} +#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} // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec28%3A303 fn sample_cosine_hemisphere(normal: vec3, rng: ptr) -> vec3 { @@ -48,29 +48,42 @@ fn sample_disk(disk_radius: f32, rng: ptr) -> vec2 { return vec2(x, y); } -struct SampleRandomLightResult { - radiance: vec3, - inverse_pdf: f32, -} - -fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rng: ptr) -> SampleRandomLightResult { - let light_sample = generate_random_light_sample(rng); - let light_contribution = calculate_light_contribution(light_sample, ray_origin, origin_world_normal); - let visibility = trace_light_visibility(light_sample, ray_origin); - return SampleRandomLightResult(light_contribution.radiance * visibility, light_contribution.inverse_pdf); -} - struct LightSample { light_id: u32, seed: u32, } +struct ResolvedLightSample { + world_position: vec4, + world_normal: vec3, + radiance: vec3, + inverse_pdf: f32, +} + struct LightContribution { radiance: vec3, inverse_pdf: f32, + wi: vec3, +} + +struct LightContributionNoPdf { + radiance: vec3, + wi: vec3, } -fn generate_random_light_sample(rng: ptr) -> LightSample { +struct GenerateRandomLightSampleResult { + light_sample: LightSample, + resolved_light_sample: ResolvedLightSample, +} + +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 generate_random_light_sample(rng: ptr) -> GenerateRandomLightSampleResult { let light_count = arrayLength(&light_sources); let light_id = rand_range_u(light_count, rng); let seed = rand_u(rng); @@ -83,115 +96,92 @@ fn generate_random_light_sample(rng: ptr) -> LightSample { triangle_id = seed % triangle_count; } - return LightSample((light_id << 16u) | triangle_id, seed); -} - -fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { - let light_id = light_sample.light_id >> 16u; - let light_source = light_sources[light_id]; - - var light_contribution: LightContribution; - if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { - light_contribution = calculate_directional_light_contribution(light_sample, light_source.id, origin_world_normal); - } else { - let triangle_count = light_source.kind >> 1u; - light_contribution = calculate_emissive_mesh_contribution(light_sample, light_source.id, triangle_count, ray_origin, origin_world_normal); - } + let light_sample = LightSample((light_id << 16u) | triangle_id, seed); - let light_count = arrayLength(&light_sources); - light_contribution.inverse_pdf *= f32(light_count); + var resolved_light_sample = resolve_light_sample(light_sample, light_source); + resolved_light_sample.inverse_pdf *= f32(light_count); - return light_contribution; + return GenerateRandomLightSampleResult(light_sample, resolved_light_sample); } -fn calculate_directional_light_contribution(light_sample: LightSample, directional_light_id: u32, origin_world_normal: vec3) -> LightContribution { - let directional_light = directional_lights[directional_light_id]; +fn resolve_light_sample(light_sample: LightSample, light_source: LightSource) -> ResolvedLightSample { + 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 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; - let x = cos(phi) * sin_theta; - let y = sin(phi) * sin_theta; - var ray_direction = vec3(x, y, cos_theta); - - // Rotate the ray so that the cone it was sampled from is aligned with the light direction - ray_direction = orthonormalize(directional_light.direction_to_light) * ray_direction; + // 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 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; + let x = cos(phi) * sin_theta; + let y = sin(phi) * sin_theta; + var direction_to_light = vec3(x, y, cos_theta); + + // Rotate the ray so that the cone it was sampled from is aligned with the light direction + direction_to_light = orthonormalize(directional_light.direction_to_light) * direction_to_light; #else - let ray_direction = directional_light.direction_to_light; + let direction_to_light = directional_light.direction_to_light; #endif - let cos_theta_origin = saturate(dot(ray_direction, origin_world_normal)); - let radiance = directional_light.luminance * cos_theta_origin; - - return LightContribution(radiance, directional_light.inverse_pdf); + return ResolvedLightSample( + 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); + + return ResolvedLightSample( + vec4(triangle_data.world_position, 1.0), + triangle_data.world_normal, + triangle_data.material.emissive.rgb, + f32(triangle_count) * triangle_data.triangle_area, + ); + } } -fn calculate_emissive_mesh_contribution(light_sample: LightSample, instance_id: u32, triangle_count: u32, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { - let barycentrics = triangle_barycentrics(light_sample.seed); - let triangle_id = light_sample.light_id & 0xFFFFu; - - let triangle_data = resolve_triangle_data_full(instance_id, triangle_id, barycentrics); +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); + let light_distance = length(ray); + let wi = ray / light_distance; - let light_distance = distance(ray_origin, triangle_data.world_position); - let ray_direction = (triangle_data.world_position - ray_origin) / light_distance; - let cos_theta_origin = saturate(dot(ray_direction, origin_world_normal)); - let cos_theta_light = saturate(dot(-ray_direction, triangle_data.world_normal)); + let cos_theta_origin = saturate(dot(wi, origin_world_normal)); + let cos_theta_light = saturate(dot(-wi, resolved_light_sample.world_normal)); let light_distance_squared = light_distance * light_distance; - let radiance = triangle_data.material.emissive.rgb * cos_theta_origin * (cos_theta_light / light_distance_squared); - let inverse_pdf = f32(triangle_count) * triangle_data.triangle_area; + let radiance = resolved_light_sample.radiance * cos_theta_origin * (cos_theta_light / light_distance_squared); - return LightContribution(radiance, inverse_pdf); + return LightContribution(radiance, resolved_light_sample.inverse_pdf, wi); } -fn trace_light_visibility(light_sample: LightSample, ray_origin: vec3) -> f32 { - let light_id = light_sample.light_id >> 16u; - let light_source = light_sources[light_id]; - - if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { - return trace_directional_light_visibility(light_sample, light_source.id, ray_origin); - } else { - return trace_emissive_mesh_visibility(light_sample, light_source.id, ray_origin); - } +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); } -fn trace_directional_light_visibility(light_sample: LightSample, directional_light_id: u32, ray_origin: vec3) -> f32 { - let directional_light = directional_lights[directional_light_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 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; - let x = cos(phi) * sin_theta; - let y = sin(phi) * sin_theta; - var ray_direction = vec3(x, y, cos_theta); - - // Rotate the ray so that the cone it was sampled from is aligned with the light direction - ray_direction = orthonormalize(directional_light.direction_to_light) * ray_direction; -#else - let ray_direction = directional_light.direction_to_light; -#endif - - 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); -} +fn trace_light_visibility(ray_origin: vec3, light_sample_world_position: vec4) -> f32 { + var ray_direction = light_sample_world_position.xyz; + var ray_t_max = RAY_T_MAX; -fn trace_emissive_mesh_visibility(light_sample: LightSample, instance_id: u32, ray_origin: vec3) -> f32 { - let barycentrics = triangle_barycentrics(light_sample.seed); - let triangle_id = light_sample.light_id & 0xFFFFu; + if light_sample_world_position.w == 1.0 { + let ray = ray_direction - ray_origin; + let dist = length(ray); + ray_direction = ray / dist; + ray_t_max = dist - RAY_T_MIN - RAY_T_MIN; + } - let triangle_data = resolve_triangle_data_full(instance_id, triangle_id, barycentrics); + if ray_t_max < RAY_T_MIN { return 0.0; } - return trace_point_visibility(ray_origin, triangle_data.world_position); + 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); } fn trace_point_visibility(ray_origin: vec3, point: vec3) -> f32 { diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index 66f258eeb1fdc..e2bab8cfd7499 100644 --- a/release-content/release-notes/bevy_solari.md +++ b/release-content/release-notes/bevy_solari.md @@ -1,7 +1,7 @@ --- title: Initial raytraced lighting progress (bevy_solari) authors: ["@JMS55"] -pull_requests: [19058, 19620, 19790, 20020, 20113] +pull_requests: [19058, 19620, 19790, 20020, 20113, TODO] --- (TODO: Embed solari example screenshot here) From 454a861ed4ab667877ecd663cf1d6adf060be8e0 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 14:34:18 -0400 Subject: [PATCH 18/31] Paper links --- crates/bevy_solari/src/realtime/presample_light_tiles.wgsl | 2 ++ crates/bevy_solari/src/realtime/restir_di.wgsl | 1 + 2 files changed, 3 insertions(+) diff --git a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl index 755b0813df8fd..75d5c0ba80710 100644 --- a/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl +++ b/crates/bevy_solari/src/realtime/presample_light_tiles.wgsl @@ -1,3 +1,5 @@ +// 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_} diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index c3aedb96c214e..236bf777c28dd 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -1,4 +1,5 @@ // 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 From 548f73f9b6ce638e1d46f632626dfbd7f478b719 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 14:38:17 -0400 Subject: [PATCH 19/31] Update release notes with PR number --- release-content/release-notes/bevy_solari.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index e2bab8cfd7499..e7c25e3c8e6dc 100644 --- a/release-content/release-notes/bevy_solari.md +++ b/release-content/release-notes/bevy_solari.md @@ -1,7 +1,7 @@ --- title: Initial raytraced lighting progress (bevy_solari) authors: ["@JMS55"] -pull_requests: [19058, 19620, 19790, 20020, 20113, TODO] +pull_requests: [19058, 19620, 19790, 20020, 20113, 20213] --- (TODO: Embed solari example screenshot here) From aa4c54506b6c54bd15f0a2fddb34f4c8a794a2d0 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 14:45:13 -0400 Subject: [PATCH 20/31] CI --- crates/bevy_solari/src/realtime/prepare.rs | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 0aee54816ed12..ea5ce3cf8fbf4 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -17,16 +17,16 @@ use bevy_render::{ renderer::RenderDevice, }; -/// Size of the LightSample shader struct in bytes. +/// Size of the `LightSample` shader struct in bytes. const LIGHT_SAMPLE_STRUCT_SIZE: u64 = 8; -/// Size of the ResolvedLightSamplePacked shader struct in bytes. +/// Size of the `ResolvedLightSamplePacked` shader struct in bytes. const RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE: u64 = 24; -/// Size of the DI Reservoir shader struct in bytes. +/// Size of the DI `Reservoir` shader struct in bytes. const DI_RESERVOIR_STRUCT_SIZE: u64 = 16; -/// Size of the GI Reservoir shader struct in bytes. +/// Size of the GI `Reservoir` shader struct in bytes. const GI_RESERVOIR_STRUCT_SIZE: u64 = 48; pub const LIGHT_TILE_BLOCKS: u64 = 128; From eb8429828eaebfdde8ea9158a94efc8eacff5dcd Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 16:03:05 -0400 Subject: [PATCH 21/31] Small fix --- crates/bevy_solari/src/scene/sampling.wgsl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index f402c1ea419dc..677d2649a8cac 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -86,16 +86,16 @@ fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rn fn generate_random_light_sample(rng: ptr) -> GenerateRandomLightSampleResult { let light_count = arrayLength(&light_sources); let light_id = rand_range_u(light_count, rng); - let seed = rand_u(rng); let light_source = light_sources[light_id]; var triangle_id = 0u; if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { let triangle_count = light_source.kind >> 1u; - triangle_id = seed % triangle_count; + 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); From 5c5f0994c10aacd0fa69ff87c76cfa3927275e9e Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 16:43:06 -0400 Subject: [PATCH 22/31] Add scene limit checks --- crates/bevy_solari/src/scene/binder.rs | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/crates/bevy_solari/src/scene/binder.rs b/crates/bevy_solari/src/scene/binder.rs index f14b5dbe23b6f..45e879a600e2a 100644 --- a/crates/bevy_solari/src/scene/binder.rs +++ b/crates/bevy_solari/src/scene/binder.rs @@ -231,6 +231,10 @@ pub fn prepare_raytracing_scene_bindings( .push(current_frame_index); } + if light_sources.get().len() > u16::MAX as usize { + panic!("Too many light sources in the scene, maximum is 2^16."); + } + materials.write_buffer(&render_device, &render_queue); transforms.write_buffer(&render_device, &render_queue); geometry_ids.write_buffer(&render_device, &render_queue); @@ -358,6 +362,10 @@ struct GpuLightSource { impl GpuLightSource { fn new_emissive_mesh_light(instance_id: u32, triangle_count: u32) -> GpuLightSource { + if triangle_count > u16::MAX as u32 { + panic!("Too triangles in an emissive mesh, maximum is 2^16."); + } + Self { kind: triangle_count << 1, id: instance_id, From aefcddf2087b941e13b69a2a410cd1d35524dc9d Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 16:44:05 -0400 Subject: [PATCH 23/31] Fix error message --- crates/bevy_solari/src/scene/binder.rs | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/crates/bevy_solari/src/scene/binder.rs b/crates/bevy_solari/src/scene/binder.rs index 45e879a600e2a..4e75e7db6b3ce 100644 --- a/crates/bevy_solari/src/scene/binder.rs +++ b/crates/bevy_solari/src/scene/binder.rs @@ -232,7 +232,7 @@ pub fn prepare_raytracing_scene_bindings( } if light_sources.get().len() > u16::MAX as usize { - panic!("Too many light sources in the scene, maximum is 2^16."); + panic!("Too many light sources in the scene, maximum is 65536."); } materials.write_buffer(&render_device, &render_queue); @@ -363,7 +363,7 @@ struct GpuLightSource { impl GpuLightSource { fn new_emissive_mesh_light(instance_id: u32, triangle_count: u32) -> GpuLightSource { if triangle_count > u16::MAX as u32 { - panic!("Too triangles in an emissive mesh, maximum is 2^16."); + panic!("Too triangles in an emissive mesh, maximum is 65535."); } Self { From 7430d907c0f5dd081132b706704bc54861286d62 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Sun, 20 Jul 2025 21:25:06 -0400 Subject: [PATCH 24/31] Adjust example --- examples/3d/solari.rs | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/examples/3d/solari.rs b/examples/3d/solari.rs index 9991ff7badb1b..88422d093dcca 100644 --- a/examples/3d/solari.rs +++ b/examples/3d/solari.rs @@ -87,7 +87,7 @@ fn add_raytracing_meshes_on_scene_load( mut commands: Commands, args: Res, ) { - // Ensure meshes are bevy_solari compatible + // Ensure meshes are Solari compatible for (_, mesh) in meshes.iter_mut() { mesh.remove_attribute(Mesh::ATTRIBUTE_UV_1.id); mesh.remove_attribute(Mesh::ATTRIBUTE_COLOR.id); @@ -113,14 +113,13 @@ fn add_raytracing_meshes_on_scene_load( } } - // Adjust scene materials to better demo bevy_solari features + // Adjust scene materials to better demo Solari features for (_, material) in materials.iter_mut() { material.emissive *= 200.0; if material.base_color.to_linear() == LinearRgba::new(0.5, 0.5, 0.5, 1.0) { material.metallic = 1.0; - material.perceptual_roughness = 0.0; - material.reflectance = 1.0; + material.perceptual_roughness = 0.15; } } } From 2f8a1a46d32c8411e3009bb0ec0b15d18e1aff85 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 22 Jul 2025 00:12:17 -0400 Subject: [PATCH 25/31] Adjust release notes --- release-content/release-notes/bevy_solari.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index e7c25e3c8e6dc..f5326b32879a5 100644 --- a/release-content/release-notes/bevy_solari.md +++ b/release-content/release-notes/bevy_solari.md @@ -1,7 +1,7 @@ --- title: Initial raytraced lighting progress (bevy_solari) authors: ["@JMS55"] -pull_requests: [19058, 19620, 19790, 20020, 20113, 20213] +pull_requests: [19058, 19620, 19790, 20020, 20113, 20213, 20242] --- (TODO: Embed solari example screenshot here) From 65b9cafcbbf39ecaaaf01725af4990ae3cbd7a1d Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 22 Jul 2025 22:20:29 -0400 Subject: [PATCH 26/31] Misc --- crates/bevy_solari/src/realtime/restir_gi.wgsl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 2be5e8554dba9..f22c634a91e3a 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -155,7 +155,7 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.viewport.z); var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; - var jacobian = jacobian( + let jacobian = jacobian( world_position, spatial_world_position, spatial_reservoir.sample_point_world_position, From 47a4d045b1b5981e26b8f9aaed73de424d52bd72 Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Tue, 22 Jul 2025 23:24:50 -0400 Subject: [PATCH 27/31] Misc --- crates/bevy_solari/src/realtime/restir_gi.wgsl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index f22c634a91e3a..2be5e8554dba9 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -155,7 +155,7 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.viewport.z); var spatial_reservoir = gi_reservoirs_b[spatial_pixel_index]; - let jacobian = jacobian( + var jacobian = jacobian( world_position, spatial_world_position, spatial_reservoir.sample_point_world_position, From e79d7ae0e0fd10fa30f6499153d917a1412b012c Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Thu, 24 Jul 2025 22:03:02 +0100 Subject: [PATCH 28/31] implement mis into pt --- .../src/pathtracer/pathtracer.wgsl | 51 +++++++++++++++---- crates/bevy_solari/src/scene/binder.rs | 2 + .../src/scene/raytracing_scene_bindings.wgsl | 4 +- crates/bevy_solari/src/scene/sampling.wgsl | 16 +++++- 4 files changed, 62 insertions(+), 11 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index dc80bb4799047..e8c5cfbd3bf77 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, sample_ggx_vndf, ggx_vndf_pdf} +#import bevy_solari::sampling::{sample_random_light, 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; @@ -37,30 +37,42 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { // Path trace var radiance = vec3(0.0); var throughput = vec3(1.0); + var p_bounce = 0.0; + var is_perfectly_specular = true; + var previous_normal = vec3(0.0); loop { 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); let wo = -ray_direction; - // Use emissive only on the first ray (coming from the camera) - if ray_t_min == 0.0 { radiance = ray_hit.material.emissive; } + var mis_weight = 1.0; + if !is_perfectly_specular { + let p_light = random_light_pdf(ray_hit); + mis_weight = power_heuristic(p_bounce, p_light); + } + radiance += mis_weight * throughput * ray_hit.material.emissive; // Sample direct lighting let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); + let p_would_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); + mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, p_would_bounce); let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); - radiance += throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; + radiance += mis_weight * throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; // Sample new ray direction from the material BRDF for next bounce let next_bounce = importance_sample_next_bounce(wo, ray_hit, &rng); ray_direction = next_bounce.wi; ray_origin = ray_hit.world_position; ray_t_min = RAY_T_MIN; + p_bounce = next_bounce.pdf; + is_perfectly_specular = ray_hit.material.roughness < 0.0001 && ray_hit.material.metallic > 0.9999; + previous_normal = ray_hit.world_normal; // Update throughput for next bounce let brdf = evaluate_brdf(ray_hit.world_normal, wo, next_bounce.wi, ray_hit.material); let cos_theta = dot(next_bounce.wi, ray_hit.world_normal); - throughput *= (brdf * cos_theta) / next_bounce.pdf; + throughput *= next_bounce.mis_weight * (brdf * cos_theta) / next_bounce.pdf; // Russian roulette for early termination let p = luminance(throughput); @@ -80,12 +92,13 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { struct NextBounce { wi: vec3, + mis_weight: f32, pdf: f32, } fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng: ptr) -> NextBounce { - let diffuse_weight = 1.0 - ray_hit.material.metallic; - let specular_weight = ray_hit.material.metallic; + let diffuse_weight = mix(mix(0.4f, 0.9f, ray_hit.material.roughness), 0.f, ray_hit.material.metallic); + let specular_weight = 1.0 - diffuse_weight; let TBN = calculate_tbn_mikktspace(ray_hit.world_normal, ray_hit.world_tangent); let T = TBN[0]; @@ -96,7 +109,8 @@ fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng var wi: vec3; var wi_tangent: vec3; - if rand_f(rng) < diffuse_weight { + let diffuse_selected = rand_f(rng) < diffuse_weight; + if diffuse_selected { wi = sample_cosine_hemisphere(ray_hit.world_normal, rng); wi_tangent = vec3(dot(wi, T), dot(wi, B), dot(wi, N)); } else { @@ -107,6 +121,25 @@ fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng let diffuse_pdf = dot(wi, ray_hit.world_normal) / PI; let specular_pdf = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); let pdf = (diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf); + let mis_weight = select(balance_heuristic(specular_pdf, diffuse_pdf), balance_heuristic(diffuse_pdf, specular_pdf), diffuse_selected); - return NextBounce(wi, pdf); + return NextBounce(wi, mis_weight, pdf); +} + +fn brdf_pdf(wo: vec3, wi: vec3, ray_hit: ResolvedRayHitFull) -> f32 { + let diffuse_weight = mix(mix(0.4f, 0.9f, ray_hit.material.roughness), 0.f, ray_hit.material.metallic); + let specular_weight = 1.0 - diffuse_weight; + + let TBN = calculate_tbn_mikktspace(ray_hit.world_normal, ray_hit.world_tangent); + let T = TBN[0]; + let B = TBN[1]; + let N = TBN[2]; + + let wo_tangent = vec3(dot(wo, T), dot(wo, B), dot(wo, N)); + let wi_tangent = vec3(dot(wi, T), dot(wi, B), dot(wi, N)); + + let diffuse_pdf = wi_tangent.z / PI; + let specular_pdf = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); + let pdf = (diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf); + return pdf; } diff --git a/crates/bevy_solari/src/scene/binder.rs b/crates/bevy_solari/src/scene/binder.rs index 45ec5f13ee6e7..1aed9310e2325 100644 --- a/crates/bevy_solari/src/scene/binder.rs +++ b/crates/bevy_solari/src/scene/binder.rs @@ -190,6 +190,7 @@ pub fn prepare_raytracing_scene_bindings( vertex_buffer_offset: vertex_slice.range.start, index_buffer_id, index_buffer_offset: index_slice.range.start, + triangle_count: (index_slice.range.len() / 3) as u32, }); material_ids.get_mut().push(material_id); @@ -352,6 +353,7 @@ struct GpuInstanceGeometryIds { vertex_buffer_offset: u32, index_buffer_id: u32, index_buffer_offset: u32, + triangle_count: u32, } #[derive(ShaderType)] diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index 6582100711836..7359ad9063e2d 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -8,6 +8,7 @@ struct InstanceGeometryIds { vertex_buffer_offset: u32, index_buffer_id: u32, index_buffer_offset: u32, + triangle_count: u32, } struct VertexBuffer { vertices: array } @@ -115,6 +116,7 @@ struct ResolvedRayHitFull { world_tangent: vec4, uv: vec2, triangle_area: f32, + triangle_count: u32, material: ResolvedMaterial, } @@ -192,5 +194,5 @@ fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: let resolved_material = resolve_material(material, uv); - return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, world_tangent, uv, triangle_area, resolved_material); + return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, world_tangent, uv, triangle_area, instance_geometry_ids.triangle_count, resolved_material); } diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index fe97bbc41d172..8385f9b3999e6 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -3,7 +3,15 @@ #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} +#import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LightSource, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full, ResolvedRayHitFull} + +fn power_heuristic(f: f32, g: f32) -> f32 { + return f * f / (f * f + g * g); +} + +fn balance_heuristic(f: f32, g: f32) -> f32 { + return f / (f + g); +} // https://gpuopen.com/download/Bounded_VNDF_Sampling_for_Smith-GGX_Reflections.pdf (Listing 1) fn sample_ggx_vndf(wi_tangent: vec3, roughness: f32, rng: ptr) -> vec3 { @@ -80,6 +88,12 @@ fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rn return light_contribution; } +fn 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)); +} + fn generate_random_light_sample(rng: ptr) -> GenerateRandomLightSampleResult { let light_count = arrayLength(&light_sources); let light_id = rand_range_u(light_count, rng); From 9aea6433b835bebbaed3514265e6fd39a2d24a20 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Thu, 24 Jul 2025 23:56:58 +0100 Subject: [PATCH 29/31] add to release notes --- release-content/release-notes/bevy_solari.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index f5326b32879a5..c5a3881f00043 100644 --- a/release-content/release-notes/bevy_solari.md +++ b/release-content/release-notes/bevy_solari.md @@ -1,6 +1,6 @@ --- title: Initial raytraced lighting progress (bevy_solari) -authors: ["@JMS55"] +authors: ["@JMS55", "@SparkyPotato"] pull_requests: [19058, 19620, 19790, 20020, 20113, 20213, 20242] --- From 70f8ad88024939c5de9ee83c53ffd21a9d4b574c Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Fri, 25 Jul 2025 00:25:17 +0100 Subject: [PATCH 30/31] address feedback and special case specular --- .../src/pathtracer/pathtracer.wgsl | 19 ++++++++++++------- 1 file changed, 12 insertions(+), 7 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index e8c5cfbd3bf77..0c2cb135a0511 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -38,7 +38,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { var radiance = vec3(0.0); var throughput = vec3(1.0); var p_bounce = 0.0; - var is_perfectly_specular = true; + var bounce_was_perfect_reflection = true; var previous_normal = vec3(0.0); loop { let ray_hit = trace_ray(ray_origin, ray_direction, ray_t_min, RAY_T_MAX, RAY_FLAG_NONE); @@ -47,7 +47,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { let wo = -ray_direction; var mis_weight = 1.0; - if !is_perfectly_specular { + if !bounce_was_perfect_reflection { let p_light = random_light_pdf(ray_hit); mis_weight = power_heuristic(p_bounce, p_light); } @@ -55,8 +55,8 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { // Sample direct lighting let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); - let p_would_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); - mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, p_would_bounce); + let pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); + mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); radiance += mis_weight * throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; @@ -66,7 +66,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { ray_origin = ray_hit.world_position; ray_t_min = RAY_T_MIN; p_bounce = next_bounce.pdf; - is_perfectly_specular = ray_hit.material.roughness < 0.0001 && ray_hit.material.metallic > 0.9999; + bounce_was_perfect_reflection = next_bounce.perfectly_specular_bounce; previous_normal = ray_hit.world_normal; // Update throughput for next bounce @@ -94,10 +94,15 @@ struct NextBounce { wi: vec3, mis_weight: f32, pdf: f32, + perfectly_specular_bounce: bool, } fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng: ptr) -> NextBounce { - let diffuse_weight = mix(mix(0.4f, 0.9f, ray_hit.material.roughness), 0.f, ray_hit.material.metallic); + let is_perfectly_specular = ray_hit.material.roughness < 0.0001 && ray_hit.material.metallic > 0.9999; + if is_perfectly_specular { + return NextBounce(reflect(-wo, ray_hit.world_normal), 1.0, 1.0, true); + } + let diffuse_weight = mix(mix(0.4f, 0.9f, ray_hit.material.perceptual_roughness), 0.f, ray_hit.material.metallic); let specular_weight = 1.0 - diffuse_weight; let TBN = calculate_tbn_mikktspace(ray_hit.world_normal, ray_hit.world_tangent); @@ -123,7 +128,7 @@ fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng let pdf = (diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf); let mis_weight = select(balance_heuristic(specular_pdf, diffuse_pdf), balance_heuristic(diffuse_pdf, specular_pdf), diffuse_selected); - return NextBounce(wi, mis_weight, pdf); + return NextBounce(wi, mis_weight, pdf, false); } fn brdf_pdf(wo: vec3, wi: vec3, ray_hit: ResolvedRayHitFull) -> f32 { From 5761e40a050384067d9397a505435fcbd0999d02 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Sat, 26 Jul 2025 01:07:36 +0100 Subject: [PATCH 31/31] fix lower exposure --- .../src/pathtracer/pathtracer.wgsl | 23 ++++++++++--------- 1 file changed, 12 insertions(+), 11 deletions(-) diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index 0c2cb135a0511..da216da959201 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -53,12 +53,15 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { } radiance += mis_weight * throughput * ray_hit.material.emissive; - // Sample direct lighting - let direct_lighting = sample_random_light(ray_hit.world_position, ray_hit.world_normal, &rng); - let pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); - mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); - 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; + // 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 pdf_of_bounce = brdf_pdf(wo, direct_lighting.wi, ray_hit); + mis_weight = power_heuristic(1.0 / direct_lighting.inverse_pdf, pdf_of_bounce); + let direct_lighting_brdf = evaluate_brdf(ray_hit.world_normal, wo, direct_lighting.wi, ray_hit.material); + radiance += mis_weight * throughput * direct_lighting.radiance * direct_lighting.inverse_pdf * direct_lighting_brdf; + } // Sample new ray direction from the material BRDF for next bounce let next_bounce = importance_sample_next_bounce(wo, ray_hit, &rng); @@ -72,7 +75,7 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { // Update throughput for next bounce let brdf = evaluate_brdf(ray_hit.world_normal, wo, next_bounce.wi, ray_hit.material); let cos_theta = dot(next_bounce.wi, ray_hit.world_normal); - throughput *= next_bounce.mis_weight * (brdf * cos_theta) / next_bounce.pdf; + throughput *= (brdf * cos_theta) / next_bounce.pdf; // Russian roulette for early termination let p = luminance(throughput); @@ -92,7 +95,6 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { struct NextBounce { wi: vec3, - mis_weight: f32, pdf: f32, perfectly_specular_bounce: bool, } @@ -100,7 +102,7 @@ struct NextBounce { fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng: ptr) -> NextBounce { let is_perfectly_specular = ray_hit.material.roughness < 0.0001 && ray_hit.material.metallic > 0.9999; if is_perfectly_specular { - return NextBounce(reflect(-wo, ray_hit.world_normal), 1.0, 1.0, true); + return NextBounce(reflect(-wo, ray_hit.world_normal), 1.0, true); } let diffuse_weight = mix(mix(0.4f, 0.9f, ray_hit.material.perceptual_roughness), 0.f, ray_hit.material.metallic); let specular_weight = 1.0 - diffuse_weight; @@ -126,9 +128,8 @@ fn importance_sample_next_bounce(wo: vec3, ray_hit: ResolvedRayHitFull, rng let diffuse_pdf = dot(wi, ray_hit.world_normal) / PI; let specular_pdf = ggx_vndf_pdf(wo_tangent, wi_tangent, ray_hit.material.roughness); let pdf = (diffuse_weight * diffuse_pdf) + (specular_weight * specular_pdf); - let mis_weight = select(balance_heuristic(specular_pdf, diffuse_pdf), balance_heuristic(diffuse_pdf, specular_pdf), diffuse_selected); - return NextBounce(wi, mis_weight, pdf, false); + return NextBounce(wi, pdf, false); } fn brdf_pdf(wo: vec3, wi: vec3, ray_hit: ResolvedRayHitFull) -> f32 {