From 8929fa2d63bc639daddbf8d6228768ce5d4108c6 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Tue, 15 Jul 2025 23:49:47 +0100 Subject: [PATCH 1/6] optimize initial and temporal DI --- crates/bevy_solari/src/realtime/node.rs | 18 +-- crates/bevy_solari/src/realtime/prepare.rs | 53 ++++--- .../bevy_solari/src/realtime/restir_di.wgsl | 135 +++++++++++++----- .../bevy_solari/src/realtime/restir_gi.wgsl | 18 +-- .../src/scene/raytracing_scene_bindings.wgsl | 67 +++++++-- crates/bevy_solari/src/scene/sampling.wgsl | 111 +++++++++++--- 6 files changed, 303 insertions(+), 99 deletions(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index eaa432d8cbace..c76840b61bff7 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -20,7 +20,7 @@ use bevy_render::{ }, BindGroupEntries, BindGroupLayout, BindGroupLayoutEntries, CachedComputePipelineId, ComputePassDescriptor, ComputePipelineDescriptor, PipelineCache, PushConstantRange, - ShaderStages, StorageTextureAccess, TextureSampleType, + ShaderStages, StorageTextureAccess, TextureFormat, TextureSampleType, }, renderer::{RenderContext, RenderDevice}, view::{ViewTarget, ViewUniform, ViewUniformOffset, ViewUniforms}, @@ -107,12 +107,10 @@ impl ViewNode for SolariLightingNode { &self.bind_group_layout, &BindGroupEntries::sequential(( view_target.get_unsampled_color_attachment().view, - solari_lighting_resources - .di_reservoirs_a - .as_entire_binding(), - solari_lighting_resources - .di_reservoirs_b - .as_entire_binding(), + &solari_lighting_resources.di_reservoirs_a.1, + &solari_lighting_resources.di_reservoirs_a.3, + &solari_lighting_resources.di_reservoirs_b.1, + &solari_lighting_resources.di_reservoirs_b.3, solari_lighting_resources .gi_reservoirs_a .as_entire_binding(), @@ -213,8 +211,10 @@ impl FromWorld for SolariLightingNode { ViewTarget::TEXTURE_FORMAT_HDR, StorageTextureAccess::ReadWrite, ), - storage_buffer_sized(false, None), - storage_buffer_sized(false, None), + texture_storage_2d(TextureFormat::Rgba32Float, StorageTextureAccess::ReadWrite), + texture_storage_2d(TextureFormat::Rgba32Float, StorageTextureAccess::ReadWrite), + texture_storage_2d(TextureFormat::Rgba32Float, StorageTextureAccess::ReadWrite), + texture_storage_2d(TextureFormat::Rgba32Float, StorageTextureAccess::ReadWrite), storage_buffer_sized(false, None), storage_buffer_sized(false, None), texture_2d(TextureSampleType::Uint), diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 46a94a3ca2477..0c694d7a5329c 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -12,22 +12,19 @@ use bevy_render::{ camera::ExtractedCamera, render_resource::{ Buffer, BufferDescriptor, BufferUsages, Texture, TextureDescriptor, TextureDimension, - TextureUsages, TextureView, TextureViewDescriptor, + TextureFormat, TextureUsages, TextureView, TextureViewDescriptor, }, renderer::RenderDevice, }; -/// Size of a DI Reservoir shader struct in bytes. -const DI_RESERVOIR_STRUCT_SIZE: u64 = 32; - /// Size of a GI Reservoir shader struct in bytes. const GI_RESERVOIR_STRUCT_SIZE: u64 = 48; /// Internal rendering resources used for Solari lighting. #[derive(Component)] pub struct SolariLightingResources { - pub di_reservoirs_a: Buffer, - pub di_reservoirs_b: Buffer, + pub di_reservoirs_a: (Texture, TextureView, Texture, TextureView), + pub di_reservoirs_b: (Texture, TextureView, Texture, TextureView), pub gi_reservoirs_a: Buffer, pub gi_reservoirs_b: Buffer, pub previous_gbuffer: (Texture, TextureView), @@ -52,19 +49,25 @@ pub fn prepare_solari_lighting_resources( continue; } - let di_reservoirs_a = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_di_reservoirs_a"), - size: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); + let di_reservoir = |label| { + let tex = render_device.create_texture(&TextureDescriptor { + label: Some(label), + size: view_size.to_extents(), + mip_level_count: 1, + sample_count: 1, + dimension: TextureDimension::D2, + format: TextureFormat::Rgba32Float, + usage: TextureUsages::STORAGE_BINDING, + view_formats: &[], + }); + let view = tex.create_view(&TextureViewDescriptor::default()); + (tex, view) + }; - let di_reservoirs_b = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_di_reservoirs_b"), - size: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); + let di_reservoirs_a_1 = di_reservoir("solari_lighting_di_reservoirs_a_1"); + let di_reservoirs_a_2 = di_reservoir("solari_lighting_di_reservoirs_a_2"); + let di_reservoirs_b_1 = di_reservoir("solari_lighting_di_reservoirs_b_1"); + let di_reservoirs_b_2 = di_reservoir("solari_lighting_di_reservoirs_b_2"); let gi_reservoirs_a = render_device.create_buffer(&BufferDescriptor { label: Some("solari_lighting_gi_reservoirs_a"), @@ -105,8 +108,18 @@ pub fn prepare_solari_lighting_resources( let previous_depth_view = previous_depth.create_view(&TextureViewDescriptor::default()); commands.entity(entity).insert(SolariLightingResources { - di_reservoirs_a, - di_reservoirs_b, + di_reservoirs_a: ( + di_reservoirs_a_1.0, + di_reservoirs_a_1.1, + di_reservoirs_a_2.0, + di_reservoirs_a_2.1, + ), + di_reservoirs_b: ( + di_reservoirs_b_1.0, + di_reservoirs_b_1.1, + di_reservoirs_b_2.0, + di_reservoirs_b_2.1, + ), gi_reservoirs_a, gi_reservoirs_b, previous_gbuffer: (previous_gbuffer, previous_gbuffer_view), diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index b9a5bfa60cccf..19543666e59c2 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -4,41 +4,60 @@ #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::sampling::{ + LightSample, IndependentLightSample, generate_point_independent_sample, calculate_light_contribution_from_independent, calculate_light_contribution, + trace_visibility, 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 di_reservoirs_a_1: texture_storage_2d; +@group(1) @binding(2) var di_reservoirs_a_2: texture_storage_2d; +@group(1) @binding(3) var di_reservoirs_b_1: texture_storage_2d; +@group(1) @binding(4) var di_reservoirs_b_2: texture_storage_2d; +@group(1) @binding(7) var gbuffer: texture_2d; +@group(1) @binding(8) var depth_buffer: texture_depth_2d; +@group(1) @binding(9) var motion_vectors: texture_2d; +@group(1) @binding(10) var previous_gbuffer: texture_2d; +@group(1) @binding(11) var previous_depth_buffer: texture_depth_2d; +@group(1) @binding(12) var view: View; +@group(1) @binding(13) var previous_view: PreviousViewUniforms; struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; -const INITIAL_SAMPLES = 32u; const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 20.0; const NULL_RESERVOIR_SAMPLE = 0xFFFFFFFFu; +// https://www.reedbeta.com/blog/quick-and-easy-gpu-random-numbers-in-d3d11/ +fn compute_seed(global_id: vec3) -> u32 { + let pixel_index = global_id.x + global_id.y * u32(view.viewport.z); + var seed = pixel_index + constants.frame_index; + seed = (seed ^ 61) ^ (seed >> 16); + seed *= 9; + seed = seed ^ (seed >> 4); + seed *= 0x27d4eb2d; + seed = seed ^ (seed >> 15); + return seed; +} + @compute @workgroup_size(8, 8, 1) -fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { +fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3, + @builtin(local_invocation_index) local_id: u32) { if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } - let pixel_index = global_id.x + global_id.y * u32(view.viewport.z); - var rng = pixel_index + constants.frame_index; + var rng = compute_seed(global_id); let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - di_reservoirs_b[pixel_index] = empty_reservoir(); + let packed = pack_reservoir(empty_reservoir()); + textureStore(di_reservoirs_b_1, global_id.xy, packed.a); + textureStore(di_reservoirs_b_2, global_id.xy, packed.b); return; } let gpixel = textureLoad(gbuffer, global_id.xy, 0); @@ -47,11 +66,13 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { let base_color = pow(unpack4x8unorm(gpixel.r).rgb, vec3(2.2)); let diffuse_brdf = base_color / PI; - let 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, local_id); 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); - di_reservoirs_b[pixel_index] = merge_result.merged_reservoir; + let packed = pack_reservoir(merge_result.merged_reservoir); + textureStore(di_reservoirs_b_1, global_id.xy, packed.a); + textureStore(di_reservoirs_b_2, global_id.xy, packed.b); } @compute @workgroup_size(8, 8, 1) @@ -63,7 +84,9 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - di_reservoirs_a[pixel_index] = empty_reservoir(); + let packed = pack_reservoir(empty_reservoir()); + textureStore(di_reservoirs_a_1, global_id.xy, packed.a); + textureStore(di_reservoirs_a_2, global_id.xy, packed.b); textureStore(view_output, global_id.xy, vec4(vec3(0.0), 1.0)); return; } @@ -74,12 +97,16 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = base_color / PI; let emissive = rgb9e5_to_vec3_(gpixel.g); - let input_reservoir = di_reservoirs_b[pixel_index]; + let packed_a = textureLoad(di_reservoirs_b_1, global_id.xy); + let packed_b = textureLoad(di_reservoirs_b_2, global_id.xy); + let input_reservoir = unpack_reservoir(PackedReservoir(packed_a, packed_b)); let spatial_reservoir = load_spatial_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let merge_result = merge_reservoirs(input_reservoir, spatial_reservoir, world_position, world_normal, diffuse_brdf, &rng); let combined_reservoir = merge_result.merged_reservoir; - di_reservoirs_a[pixel_index] = combined_reservoir; + let packed = pack_reservoir(combined_reservoir); + textureStore(di_reservoirs_a_1, global_id.xy, packed.a); + textureStore(di_reservoirs_a_2, global_id.xy, packed.b); var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight * combined_reservoir.visibility; pixel_color *= view.exposure; @@ -88,30 +115,35 @@ 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) -> Reservoir{ +var position_independent_initial_samples: array; + +fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr, local_id: u32) -> Reservoir { + position_independent_initial_samples[local_id] = generate_point_independent_sample(rng); + workgroupBarrier(); + 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); + var ray_direction = vec4(0.0); + for (var i = 0u; i < 64u; i++) { + let light_contribution = calculate_light_contribution_from_independent(position_independent_initial_samples[i], world_position, world_normal); - let mis_weight = 1.0 / f32(INITIAL_SAMPLES); - let light_contribution = calculate_light_contribution(light_sample, world_position, world_normal); + let mis_weight = 1.0 / 64.0; 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; + reservoir.sample = light_contribution.sample; reservoir_target_function = target_function; + ray_direction = light_contribution.ray_direction; } } 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.visibility = trace_light_visibility(reservoir.sample, world_position); + reservoir.visibility = trace_visibility(world_position, ray_direction); } reservoir.confidence_weight = 1.0; @@ -138,8 +170,9 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return empty_reservoir(); } - let temporal_pixel_index = temporal_pixel_id.x + temporal_pixel_id.y * u32(view.viewport.z); - var temporal_reservoir = di_reservoirs_a[temporal_pixel_index]; + let packed_a = textureLoad(di_reservoirs_a_1, temporal_pixel_id); + let packed_b = textureLoad(di_reservoirs_a_2, temporal_pixel_id); + var temporal_reservoir = unpack_reservoir(PackedReservoir(packed_a, packed_b)); // 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]; @@ -163,8 +196,9 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< return empty_reservoir(); } - let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.viewport.z); - var spatial_reservoir = di_reservoirs_b[spatial_pixel_index]; + let packed_a = textureLoad(di_reservoirs_b_1, spatial_pixel_id); + let packed_b = textureLoad(di_reservoirs_b_2, spatial_pixel_id); + var spatial_reservoir = unpack_reservoir(PackedReservoir(packed_a, packed_b)); if reservoir_valid(spatial_reservoir) { spatial_reservoir.visibility = trace_light_visibility(spatial_reservoir.sample, world_position); @@ -213,7 +247,6 @@ fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { #endif } -// Don't adjust the size of this struct without also adjusting DI_RESERVOIR_STRUCT_SIZE. struct Reservoir { sample: LightSample, weight_sum: f32, @@ -236,6 +269,40 @@ fn reservoir_valid(reservoir: Reservoir) -> bool { return reservoir.sample.light_id.x != NULL_RESERVOIR_SAMPLE; } +struct PackedReservoir { + a: vec4, + b: vec4, +} + +fn pack_reservoir(reservoir: Reservoir) -> PackedReservoir { + let sample_light_id = bitcast>(reservoir.sample.light_id); + let packed_a = vec4(sample_light_id, reservoir.sample.random); + let packed_b = vec4( + reservoir.weight_sum, + reservoir.confidence_weight, + reservoir.unbiased_contribution_weight, + reservoir.visibility + ); + return PackedReservoir(packed_a, packed_b); +} + +fn unpack_reservoir(packed: PackedReservoir) -> Reservoir { + let sample_light_id = bitcast>(packed.a.xy); + let sample_random = packed.a.zw; + let weight_sum = packed.b.x; + let confidence_weight = packed.b.y; + let unbiased_contribution_weight = packed.b.z; + let visibility = packed.b.w; + + return Reservoir( + LightSample(sample_light_id, sample_random), + weight_sum, + confidence_weight, + unbiased_contribution_weight, + visibility + ); +} + struct ReservoirMergeResult { merged_reservoir: Reservoir, selected_sample_radiance: vec3, 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/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index eeed96ad8e818..b5c1632386f20 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -126,21 +126,38 @@ fn resolve_ray_hit_full(ray_hit: RayIntersection) -> ResolvedRayHitFull { return resolve_triangle_data_full(ray_hit.instance_index, ray_hit.primitive_index, barycentrics); } -fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: vec3) -> ResolvedRayHitFull { +fn load_vertices(instance_id: u32, triangle_id: u32) -> array { let instance_geometry_ids = geometry_ids[instance_id]; - let material_id = material_ids[instance_id]; - let index_buffer = &index_buffers[instance_geometry_ids.index_buffer_id].indices; let vertex_buffer = &vertex_buffers[instance_geometry_ids.vertex_buffer_id].vertices; - let material = materials[material_id]; let indices_i = (triangle_id * 3u) + vec3(0u, 1u, 2u) + instance_geometry_ids.index_buffer_offset; let indices = vec3((*index_buffer)[indices_i.x], (*index_buffer)[indices_i.y], (*index_buffer)[indices_i.z]) + instance_geometry_ids.vertex_buffer_offset; - let vertices = array(unpack_vertex((*vertex_buffer)[indices.x]), unpack_vertex((*vertex_buffer)[indices.y]), unpack_vertex((*vertex_buffer)[indices.z])); + return array( + unpack_vertex((*vertex_buffer)[indices.x]), + unpack_vertex((*vertex_buffer)[indices.y]), + unpack_vertex((*vertex_buffer)[indices.z]) + ); +} + +fn transform_positions(transform: mat4x4, vertices: array) -> array, 3> { + return array, 3>( + (transform * vec4(vertices[0].position, 1.0)).xyz, + (transform * vec4(vertices[1].position, 1.0)).xyz, + (transform * vec4(vertices[2].position, 1.0)).xyz + ); +} + +fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: vec3) -> ResolvedRayHitFull { + let material_id = material_ids[instance_id]; + let material = materials[material_id]; + + let vertices = load_vertices(instance_id, triangle_id); let transform = transforms[instance_id]; - let local_position = mat3x3(vertices[0].position, vertices[1].position, vertices[2].position) * barycentrics; - let world_position = (transform * vec4(local_position, 1.0)).xyz; + let world_vertices = transform_positions(transform, vertices); + + let world_position = mat3x3(world_vertices[0], world_vertices[1], world_vertices[2]) * barycentrics; let uv = mat3x2(vertices[0].uv, vertices[1].uv, vertices[2].uv) * barycentrics; @@ -157,11 +174,43 @@ fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: world_normal = normalize(Nt.x * T + Nt.y * B + Nt.z * N); } - let triangle_edge0 = vertices[0].position - vertices[1].position; - let triangle_edge1 = vertices[0].position - vertices[2].position; + let triangle_edge0 = world_vertices[0] - world_vertices[1]; + let triangle_edge1 = world_vertices[0] - world_vertices[2]; let triangle_area = length(cross(triangle_edge0, triangle_edge1)) / 2.0; let resolved_material = resolve_material(material, uv); return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, uv, triangle_area, resolved_material); } + +struct ResolvedDISample { + world_position: vec3, + world_normal: vec3, + emissive: vec3, + triangle_area: f32, +} + +fn resolve_triangle_for_di(instance_id: u32, triangle_id: u32, barycentrics: vec3) -> ResolvedDISample { + let material_id = material_ids[instance_id]; + var emissive = materials[material_id].emissive.rgb; + let emissive_texture_id = materials[material_id].emissive_texture_id; + + let vertices = load_vertices(instance_id, triangle_id); + let transform = transforms[instance_id]; + let world_vertices = transform_positions(transform, vertices); + + let triangle_edge0 = world_vertices[0] - world_vertices[1]; + let triangle_edge1 = world_vertices[0] - world_vertices[2]; + let triangle_area = length(cross(triangle_edge0, triangle_edge1)) / 2.0; + + let world_position = mat3x3(world_vertices[0], world_vertices[1], world_vertices[2]) * barycentrics; + let uv = mat3x2(vertices[0].uv, vertices[1].uv, vertices[2].uv) * barycentrics; + let local_normal = mat3x3(vertices[0].normal, vertices[1].normal, vertices[2].normal) * barycentrics; // TODO: Use barycentric lerp, ray_hit.object_to_world, cross product geo normal + let world_normal = normalize(mat3x3(transform[0].xyz, transform[1].xyz, transform[2].xyz) * local_normal); + + if emissive_texture_id != TEXTURE_MAP_NONE { + emissive *= sample_texture(emissive_texture_id, uv); + } + + return ResolvedDISample(world_position, world_normal, emissive, triangle_area); +} diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index e1f67ac1ed102..43ea00193c05d 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -2,7 +2,11 @@ #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::scene_bindings::{ + trace_ray, RAY_T_MIN, RAY_T_MAX, + light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, + resolve_triangle_data_full, resolve_triangle_for_di +} // https://www.realtimerendering.com/raytracinggems/unofficial_RayTracingGems_v1.9.pdf#0004286901.INDD%3ASec28%3A303 fn sample_cosine_hemisphere(normal: vec3, rng: ptr) -> vec3 { @@ -54,9 +58,9 @@ struct SampleRandomLightResult { } 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); + let pos_independent_sample = generate_point_independent_sample(rng); + let light_contribution = calculate_light_contribution_from_independent(pos_independent_sample, ray_origin, origin_world_normal); + let visibility = trace_visibility(ray_origin, light_contribution.ray_direction); return SampleRandomLightResult(light_contribution.radiance * visibility, light_contribution.inverse_pdf); } @@ -65,25 +69,90 @@ struct LightSample { random: vec2, } +struct IndependentLightSample { + sample: LightSample, + radiance: vec3, + inverse_pdf: f32, + pos_or_direction: vec4, + normal: vec3, +} + struct LightContribution { + sample: LightSample, radiance: vec3, inverse_pdf: f32, + ray_direction: vec4, } -fn generate_random_light_sample(rng: ptr) -> LightSample { +fn generate_point_independent_sample(rng: ptr) -> IndependentLightSample { let light_count = arrayLength(&light_sources); let light_id = rand_range_u(light_count, rng); - let random = rand_vec2f(rng); - let light_source = light_sources[light_id]; - var triangle_id = 0u; + var light_sample = LightSample(vec2(light_id, 0), rand_vec2f(rng)); - if light_source.kind != LIGHT_SOURCE_KIND_DIRECTIONAL { + var sample: IndependentLightSample; + if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { + sample = calculate_directional_light_independent_contribution(light_sample, light_source.id); + } else { let triangle_count = light_source.kind >> 1u; - triangle_id = rand_range_u(triangle_count, rng); + light_sample.light_id.y = rand_range_u(triangle_count, rng); + sample = calculate_emissive_mesh_independent_contribution(light_sample, light_source.id, triangle_count); } - return LightSample(vec2(light_id, triangle_id), random); + sample.inverse_pdf *= f32(light_count); + return sample; +} + +fn calculate_directional_light_independent_contribution(light_sample: LightSample, directional_light_id: u32) -> IndependentLightSample { + 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 + let cos_theta = (1.0 - light_sample.random.x) + light_sample.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 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 = build_orthonormal_basis(directional_light.direction_to_light) * ray_direction; +#else + let ray_direction = directional_light.direction_to_light; +#endif + return IndependentLightSample(light_sample, directional_light.luminance, directional_light.inverse_pdf, vec4(ray_direction, 0.0), vec3(0.0)); +} + +fn calculate_emissive_mesh_independent_contribution(light_sample: LightSample, instance_id: u32, triangle_count: u32) -> IndependentLightSample { + let barycentrics = triangle_barycentrics(light_sample.random); + let triangle_id = light_sample.light_id.y; + let triangle_data = resolve_triangle_for_di(instance_id, triangle_id, barycentrics); + let inverse_pdf = f32(triangle_count) * triangle_data.triangle_area; + return IndependentLightSample(light_sample, triangle_data.emissive, inverse_pdf, vec4(triangle_data.world_position, 1.0), triangle_data.world_normal); +} + +fn calculate_light_contribution_from_independent(sample: IndependentLightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { + var radiance = vec3(0.0); + var ray_direction = vec4(0.0); + if sample.pos_or_direction.w == 0.0 { + // Directional light + ray_direction = vec4(sample.pos_or_direction.xyz, RAY_T_MAX); + let cos_theta_origin = saturate(dot(ray_direction.xyz, origin_world_normal)); + radiance = sample.radiance * cos_theta_origin; + } else { + // Emissive mesh + let end_point = sample.pos_or_direction.xyz; + ray_direction = vec4(end_point - ray_origin, 0.0); + let light_distance_squared = dot(ray_direction.xyz, ray_direction.xyz); + ray_direction.w = sqrt(light_distance_squared); + ray_direction = vec4(ray_direction.xyz / ray_direction.w, ray_direction.w); + let cos_theta_origin = saturate(dot(ray_direction.xyz, origin_world_normal)); + let cos_theta_light = saturate(dot(-ray_direction.xyz, sample.normal)); + + radiance = sample.radiance * cos_theta_origin * (cos_theta_light / light_distance_squared); + } + return LightContribution(sample.sample, radiance, sample.inverse_pdf, ray_direction); } fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { @@ -126,25 +195,26 @@ 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); + return LightContribution(light_sample, radiance, directional_light.inverse_pdf, vec4(ray_direction, RAY_T_MAX)); } 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 triangle_data = resolve_triangle_data_full(instance_id, triangle_id, barycentrics); + let triangle_data = resolve_triangle_for_di(instance_id, triangle_id, barycentrics); - let light_distance = distance(ray_origin, triangle_data.world_position); - let ray_direction = (triangle_data.world_position - ray_origin) / light_distance; + var ray_direction = triangle_data.world_position - ray_origin; + let light_distance_squared = dot(ray_direction, ray_direction); + let light_distance = sqrt(light_distance_squared); + ray_direction /= 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 light_distance_squared = light_distance * light_distance; - let radiance = triangle_data.material.emissive.rgb * cos_theta_origin * (cos_theta_light / light_distance_squared); + let radiance = triangle_data.emissive * cos_theta_origin * (cos_theta_light / light_distance_squared); let inverse_pdf = f32(triangle_count) * triangle_data.triangle_area; - return LightContribution(radiance, inverse_pdf); + return LightContribution(light_sample, radiance, inverse_pdf, vec4(ray_direction, light_distance)); } fn trace_light_visibility(light_sample: LightSample, ray_origin: vec3) -> f32 { @@ -190,6 +260,11 @@ fn trace_emissive_mesh_visibility(light_sample: LightSample, instance_id: u32, r return trace_point_visibility(ray_origin, triangle_data.world_position); } +fn trace_visibility(ray_origin: vec3, ray_direction: vec4) -> f32 { + let ray_hit = trace_ray(ray_origin, ray_direction.xyz, RAY_T_MIN, ray_direction.w, 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 { let ray = point - ray_origin; let dist = length(ray); From e12b27c2be4c81c9eff00c0698c2d9443b1b11ed Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Wed, 16 Jul 2025 20:10:11 +0100 Subject: [PATCH 2/6] remove workgroup optimization --- .../bevy_solari/src/realtime/restir_di.wgsl | 21 +++--- crates/bevy_solari/src/scene/sampling.wgsl | 68 ++----------------- 2 files changed, 15 insertions(+), 74 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 19543666e59c2..c101157d75013 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -8,7 +8,7 @@ #import bevy_render::maths::PI #import bevy_render::view::View #import bevy_solari::sampling::{ - LightSample, IndependentLightSample, generate_point_independent_sample, calculate_light_contribution_from_independent, calculate_light_contribution, + LightSample, IndependentLightSample, sample_random_light_contribution, calculate_light_contribution, trace_visibility, trace_light_visibility, sample_disk } @@ -29,6 +29,7 @@ struct PushConstants { frame_index: u32, reset: u32 } var constants: PushConstants; +const INITIAL_SAMPLES = 32u; const SPATIAL_REUSE_RADIUS_PIXELS = 30.0; const CONFIDENCE_WEIGHT_CAP = 20.0; @@ -47,8 +48,7 @@ fn compute_seed(global_id: vec3) -> u32 { } @compute @workgroup_size(8, 8, 1) -fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3, - @builtin(local_invocation_index) local_id: u32) { +fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { if any(global_id.xy >= vec2u(view.viewport.zw)) { return; } var rng = compute_seed(global_id); @@ -66,7 +66,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, local_id); + let initial_reservoir = generate_initial_reservoir(world_position, world_normal, diffuse_brdf, &rng); let temporal_reservoir = load_temporal_reservoir(global_id.xy, depth, world_position, world_normal); let merge_result = merge_reservoirs(initial_reservoir, temporal_reservoir, world_position, world_normal, diffuse_brdf, &rng); @@ -115,19 +115,14 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { textureStore(view_output, global_id.xy, vec4(pixel_color, 1.0)); } -var position_independent_initial_samples: array; - -fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr, local_id: u32) -> Reservoir { - position_independent_initial_samples[local_id] = generate_point_independent_sample(rng); - workgroupBarrier(); - +fn generate_initial_reservoir(world_position: vec3, world_normal: vec3, diffuse_brdf: vec3, rng: ptr) -> Reservoir { var reservoir = empty_reservoir(); var reservoir_target_function = 0.0; var ray_direction = vec4(0.0); - for (var i = 0u; i < 64u; i++) { - let light_contribution = calculate_light_contribution_from_independent(position_independent_initial_samples[i], world_position, world_normal); + for (var i = 0u; i < INITIAL_SAMPLES; i++) { + let light_contribution = sample_random_light_contribution(rng, world_position, world_normal); - let mis_weight = 1.0 / 64.0; + let mis_weight = 1.0 / f32(INITIAL_SAMPLES); let target_function = luminance(light_contribution.radiance * diffuse_brdf); let resampling_weight = mis_weight * (target_function * light_contribution.inverse_pdf); diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 43ea00193c05d..1f4ef7e9c941a 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -58,8 +58,7 @@ struct SampleRandomLightResult { } fn sample_random_light(ray_origin: vec3, origin_world_normal: vec3, rng: ptr) -> SampleRandomLightResult { - let pos_independent_sample = generate_point_independent_sample(rng); - let light_contribution = calculate_light_contribution_from_independent(pos_independent_sample, ray_origin, origin_world_normal); + let light_contribution = sample_random_light_contribution(rng, ray_origin, origin_world_normal); let visibility = trace_visibility(ray_origin, light_contribution.ray_direction); return SampleRandomLightResult(light_contribution.radiance * visibility, light_contribution.inverse_pdf); } @@ -84,75 +83,22 @@ struct LightContribution { ray_direction: vec4, } -fn generate_point_independent_sample(rng: ptr) -> IndependentLightSample { +fn sample_random_light_contribution(rng: ptr, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { let light_count = arrayLength(&light_sources); let light_id = rand_range_u(light_count, rng); let light_source = light_sources[light_id]; var light_sample = LightSample(vec2(light_id, 0), rand_vec2f(rng)); - var sample: IndependentLightSample; + var light_contribution: LightContribution; if light_source.kind == LIGHT_SOURCE_KIND_DIRECTIONAL { - sample = calculate_directional_light_independent_contribution(light_sample, light_source.id); + light_contribution = calculate_directional_light_contribution(light_sample, light_source.id, origin_world_normal); } else { let triangle_count = light_source.kind >> 1u; - light_sample.light_id.y = rand_range_u(triangle_count, rng); - sample = calculate_emissive_mesh_independent_contribution(light_sample, light_source.id, triangle_count); + light_contribution = calculate_emissive_mesh_contribution(light_sample, light_source.id, triangle_count, ray_origin, origin_world_normal); } + light_contribution.inverse_pdf *= f32(light_count); - sample.inverse_pdf *= f32(light_count); - return sample; -} - -fn calculate_directional_light_independent_contribution(light_sample: LightSample, directional_light_id: u32) -> IndependentLightSample { - 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 - let cos_theta = (1.0 - light_sample.random.x) + light_sample.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 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 = build_orthonormal_basis(directional_light.direction_to_light) * ray_direction; -#else - let ray_direction = directional_light.direction_to_light; -#endif - return IndependentLightSample(light_sample, directional_light.luminance, directional_light.inverse_pdf, vec4(ray_direction, 0.0), vec3(0.0)); -} - -fn calculate_emissive_mesh_independent_contribution(light_sample: LightSample, instance_id: u32, triangle_count: u32) -> IndependentLightSample { - let barycentrics = triangle_barycentrics(light_sample.random); - let triangle_id = light_sample.light_id.y; - let triangle_data = resolve_triangle_for_di(instance_id, triangle_id, barycentrics); - let inverse_pdf = f32(triangle_count) * triangle_data.triangle_area; - return IndependentLightSample(light_sample, triangle_data.emissive, inverse_pdf, vec4(triangle_data.world_position, 1.0), triangle_data.world_normal); -} - -fn calculate_light_contribution_from_independent(sample: IndependentLightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { - var radiance = vec3(0.0); - var ray_direction = vec4(0.0); - if sample.pos_or_direction.w == 0.0 { - // Directional light - ray_direction = vec4(sample.pos_or_direction.xyz, RAY_T_MAX); - let cos_theta_origin = saturate(dot(ray_direction.xyz, origin_world_normal)); - radiance = sample.radiance * cos_theta_origin; - } else { - // Emissive mesh - let end_point = sample.pos_or_direction.xyz; - ray_direction = vec4(end_point - ray_origin, 0.0); - let light_distance_squared = dot(ray_direction.xyz, ray_direction.xyz); - ray_direction.w = sqrt(light_distance_squared); - ray_direction = vec4(ray_direction.xyz / ray_direction.w, ray_direction.w); - let cos_theta_origin = saturate(dot(ray_direction.xyz, origin_world_normal)); - let cos_theta_light = saturate(dot(-ray_direction.xyz, sample.normal)); - - radiance = sample.radiance * cos_theta_origin * (cos_theta_light / light_distance_squared); - } - return LightContribution(sample.sample, radiance, sample.inverse_pdf, ray_direction); + return light_contribution; } fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { From 4cdcc8d9f11751aa4985fa046891d2e8536ccb13 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Wed, 16 Jul 2025 20:31:38 +0100 Subject: [PATCH 3/6] update release notes --- release-content/release-notes/bevy_solari.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index 66f258eeb1fdc..cf41c2a0bf9e9 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] +authors: ["@JMS55", "@SparkyPotato"] +pull_requests: [19058, 19620, 19790, 20020, 20113, 20156] --- (TODO: Embed solari example screenshot here) From c7cf3f4b224f7d5e650cd93fc433e32901357ed3 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Tue, 29 Jul 2025 22:23:46 +0100 Subject: [PATCH 4/6] cleanup --- crates/bevy_solari/src/realtime/prepare.rs | 2 +- .../bevy_solari/src/realtime/restir_di.wgsl | 4 +- .../src/scene/raytracing_scene_bindings.wgsl | 32 -------------- crates/bevy_solari/src/scene/sampling.wgsl | 44 ------------------- 4 files changed, 3 insertions(+), 79 deletions(-) diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index 6188f80235bec..ea5ce3cf8fbf4 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -12,7 +12,7 @@ use bevy_render::{ camera::ExtractedCamera, render_resource::{ Buffer, BufferDescriptor, BufferUsages, Texture, TextureDescriptor, TextureDimension, - TextureFormat, TextureUsages, TextureView, TextureViewDescriptor, + TextureUsages, TextureView, TextureViewDescriptor, }, renderer::RenderDevice, }; diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 175b9ded96c09..bea79ef42ccb5 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -5,11 +5,11 @@ #import bevy_pbr::pbr_deferred_types::unpack_24bit_normal #import bevy_pbr::prepass_bindings::PreviousViewUniforms #import bevy_pbr::rgb9e5::rgb9e5_to_vec3_ -#import bevy_pbr::utils::{rand_f, rand_range_u, octahedral_decode} +#import bevy_pbr::utils::{rand_f, rand_range_u, octahedral_decode, sample_disk} #import bevy_render::maths::PI #import bevy_render::view::View #import bevy_solari::presample_light_tiles::{ResolvedLightSamplePacked, unpack_resolved_light_sample} -#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility, sample_disk} +#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility} #import bevy_solari::scene_bindings::{light_sources, previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME} @group(1) @binding(0) var view_output: texture_storage_2d; diff --git a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl index b5c1632386f20..c97475b200574 100644 --- a/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl +++ b/crates/bevy_solari/src/scene/raytracing_scene_bindings.wgsl @@ -182,35 +182,3 @@ fn resolve_triangle_data_full(instance_id: u32, triangle_id: u32, barycentrics: return ResolvedRayHitFull(world_position, world_normal, geometric_world_normal, uv, triangle_area, resolved_material); } - -struct ResolvedDISample { - world_position: vec3, - world_normal: vec3, - emissive: vec3, - triangle_area: f32, -} - -fn resolve_triangle_for_di(instance_id: u32, triangle_id: u32, barycentrics: vec3) -> ResolvedDISample { - let material_id = material_ids[instance_id]; - var emissive = materials[material_id].emissive.rgb; - let emissive_texture_id = materials[material_id].emissive_texture_id; - - let vertices = load_vertices(instance_id, triangle_id); - let transform = transforms[instance_id]; - let world_vertices = transform_positions(transform, vertices); - - let triangle_edge0 = world_vertices[0] - world_vertices[1]; - let triangle_edge1 = world_vertices[0] - world_vertices[2]; - let triangle_area = length(cross(triangle_edge0, triangle_edge1)) / 2.0; - - let world_position = mat3x3(world_vertices[0], world_vertices[1], world_vertices[2]) * barycentrics; - let uv = mat3x2(vertices[0].uv, vertices[1].uv, vertices[2].uv) * barycentrics; - let local_normal = mat3x3(vertices[0].normal, vertices[1].normal, vertices[2].normal) * barycentrics; // TODO: Use barycentric lerp, ray_hit.object_to_world, cross product geo normal - let world_normal = normalize(mat3x3(transform[0].xyz, transform[1].xyz, transform[2].xyz) * local_normal); - - if emissive_texture_id != TEXTURE_MAP_NONE { - emissive *= sample_texture(emissive_texture_id, uv); - } - - return ResolvedDISample(world_position, world_normal, emissive, triangle_area); -} diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 677d2649a8cac..f11ad77a46789 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -4,50 +4,6 @@ #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, 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 { - let cos_theta = 1.0 - 2.0 * 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 = normal.x + sin_theta * cos(phi); - let y = normal.y + sin_theta * sin(phi); - let z = normal.z + cos_theta; - 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 orthonormalize(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; - let a = ab.x; - var b = ab.y; - if (b == 0.0) { b = 1.0; } - - var phi: f32; - var r: f32; - if (a * a > b * b) { - r = disk_radius * a; - phi = (PI / 4.0) * (b / a); - } else { - r = disk_radius * b; - phi = (PI / 2.0) - (PI / 4.0) * (a / b); - } - - let x = r * cos(phi); - let y = r * sin(phi); - return vec2(x, y); -} - struct LightSample { light_id: u32, seed: u32, From bdfe6ce43b5ff21c1b23a9354d72b27b305ae0d6 Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Tue, 29 Jul 2025 22:52:00 +0100 Subject: [PATCH 5/6] readd reservoir texture --- crates/bevy_solari/src/realtime/node.rs | 14 ++--- crates/bevy_solari/src/realtime/prepare.rs | 63 +++++++++---------- .../bevy_solari/src/realtime/restir_di.wgsl | 34 ++++++---- crates/bevy_solari/src/scene/sampling.wgsl | 2 +- 4 files changed, 58 insertions(+), 55 deletions(-) diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index b873d7bee567f..66c64835b5fa3 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -23,7 +23,7 @@ use bevy_render::{ }, BindGroupEntries, BindGroupLayout, BindGroupLayoutEntries, CachedComputePipelineId, ComputePassDescriptor, ComputePipelineDescriptor, PipelineCache, PushConstantRange, Shader, - ShaderStages, StorageTextureAccess, TextureSampleType, + ShaderStages, StorageTextureAccess, TextureFormat, TextureSampleType, }, renderer::{RenderContext, RenderDevice}, view::{ViewTarget, ViewUniform, ViewUniformOffset, ViewUniforms}, @@ -119,12 +119,8 @@ impl ViewNode for SolariLightingNode { solari_lighting_resources .light_tile_resolved_samples .as_entire_binding(), - solari_lighting_resources - .di_reservoirs_a - .as_entire_binding(), - solari_lighting_resources - .di_reservoirs_b - .as_entire_binding(), + &solari_lighting_resources.di_reservoirs_a.1, + &solari_lighting_resources.di_reservoirs_b.1, solari_lighting_resources .gi_reservoirs_a .as_entire_binding(), @@ -230,8 +226,8 @@ impl FromWorld for SolariLightingNode { ), storage_buffer_sized(false, None), storage_buffer_sized(false, None), - storage_buffer_sized(false, None), - storage_buffer_sized(false, None), + texture_storage_2d(TextureFormat::Rgba32Uint, StorageTextureAccess::ReadWrite), + texture_storage_2d(TextureFormat::Rgba32Uint, StorageTextureAccess::ReadWrite), storage_buffer_sized(false, None), storage_buffer_sized(false, None), texture_2d(TextureSampleType::Uint), diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index ea5ce3cf8fbf4..f31f4daeb5180 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -12,7 +12,7 @@ use bevy_render::{ camera::ExtractedCamera, render_resource::{ Buffer, BufferDescriptor, BufferUsages, Texture, TextureDescriptor, TextureDimension, - TextureUsages, TextureView, TextureViewDescriptor, + TextureFormat, TextureUsages, TextureView, TextureViewDescriptor, }, renderer::RenderDevice, }; @@ -23,9 +23,6 @@ const LIGHT_SAMPLE_STRUCT_SIZE: u64 = 8; /// 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; @@ -37,8 +34,8 @@ pub const LIGHT_TILE_SAMPLES_PER_BLOCK: u64 = 1024; 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 di_reservoirs_a: (Texture, TextureView), + pub di_reservoirs_b: (Texture, TextureView), pub gi_reservoirs_a: Buffer, pub gi_reservoirs_b: Buffer, pub previous_gbuffer: (Texture, TextureView), @@ -79,33 +76,33 @@ pub fn prepare_solari_lighting_resources( 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, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); - - let di_reservoirs_b = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_di_reservoirs_b"), - size: (view_size.x * view_size.y) as u64 * DI_RESERVOIR_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); - - let gi_reservoirs_a = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_gi_reservoirs_a"), - size: (view_size.x * view_size.y) as u64 * GI_RESERVOIR_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); - - let gi_reservoirs_b = render_device.create_buffer(&BufferDescriptor { - label: Some("solari_lighting_gi_reservoirs_b"), - size: (view_size.x * view_size.y) as u64 * GI_RESERVOIR_STRUCT_SIZE, - usage: BufferUsages::STORAGE, - mapped_at_creation: false, - }); + let di_reservoirs = |name| { + let tex = render_device.create_texture(&TextureDescriptor { + label: Some(name), + size: view_size.to_extents(), + mip_level_count: 1, + sample_count: 1, + dimension: TextureDimension::D2, + format: TextureFormat::Rgba32Uint, + usage: TextureUsages::STORAGE_BINDING, + view_formats: &[], + }); + let view = tex.create_view(&TextureViewDescriptor::default()); + (tex, view) + }; + let di_reservoirs_a = di_reservoirs("solari_lighting_di_reservoirs_a"); + let di_reservoirs_b = di_reservoirs("solari_lighting_di_reservoirs_b"); + + let gi_reservoirs = |name| { + render_device.create_buffer(&BufferDescriptor { + label: Some(name), + size: (view_size.x * view_size.y) as u64 * GI_RESERVOIR_STRUCT_SIZE, + usage: BufferUsages::STORAGE, + mapped_at_creation: false, + }) + }; + let gi_reservoirs_a = gi_reservoirs("solari_lighting_gi_reservoirs_a"); + let gi_reservoirs_b = gi_reservoirs("solari_lighting_gi_reservoirs_b"); let previous_gbuffer = render_device.create_texture(&TextureDescriptor { label: Some("solari_lighting_previous_gbuffer"), diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index bea79ef42ccb5..12d2f9bf5e5bd 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -15,8 +15,8 @@ @group(1) @binding(0) var view_output: texture_storage_2d; @group(1) @binding(1) var light_tile_samples: array; @group(1) @binding(2) var light_tile_resolved_samples: array; -@group(1) @binding(3) var di_reservoirs_a: array; -@group(1) @binding(4) var di_reservoirs_b: array; +@group(1) @binding(3) var di_reservoirs_a: texture_storage_2d; +@group(1) @binding(4) var di_reservoirs_b: texture_storage_2d; @group(1) @binding(7) var gbuffer: texture_2d; @group(1) @binding(8) var depth_buffer: texture_depth_2d; @group(1) @binding(9) var motion_vectors: texture_2d; @@ -42,7 +42,7 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - di_reservoirs_b[pixel_index] = empty_reservoir(); + textureStore(di_reservoirs_b, global_id.xy, pack_reservoir(empty_reservoir())); return; } let gpixel = textureLoad(gbuffer, global_id.xy, 0); @@ -55,7 +55,7 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin 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); - di_reservoirs_b[pixel_index] = merge_result.merged_reservoir; + textureStore(di_reservoirs_b, global_id.xy, pack_reservoir(merge_result.merged_reservoir)); } @compute @workgroup_size(8, 8, 1) @@ -67,7 +67,7 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - di_reservoirs_a[pixel_index] = empty_reservoir(); + textureStore(di_reservoirs_a, global_id.xy, pack_reservoir(empty_reservoir())); textureStore(view_output, global_id.xy, vec4(vec3(0.0), 1.0)); return; } @@ -78,12 +78,13 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = base_color / PI; let emissive = rgb9e5_to_vec3_(gpixel.g); - let input_reservoir = di_reservoirs_b[pixel_index]; + let packed_input_reservoir = textureLoad(di_reservoirs_a, global_id.xy); + let input_reservoir = unpack_reservoir(packed_input_reservoir); let spatial_reservoir = load_spatial_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let merge_result = merge_reservoirs(input_reservoir, spatial_reservoir, world_position, world_normal, diffuse_brdf, &rng); let combined_reservoir = merge_result.merged_reservoir; - di_reservoirs_a[pixel_index] = combined_reservoir; + textureStore(di_reservoirs_a, global_id.xy, pack_reservoir(combined_reservoir)); var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight; pixel_color *= view.exposure; @@ -155,8 +156,8 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return empty_reservoir(); } - let temporal_pixel_index = temporal_pixel_id.x + temporal_pixel_id.y * u32(view.viewport.z); - var temporal_reservoir = di_reservoirs_a[temporal_pixel_index]; + let packed_reservoir = textureLoad(di_reservoirs_a, temporal_pixel_id); + var temporal_reservoir = unpack_reservoir(packed_reservoir); // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) let previous_light_id = temporal_reservoir.sample.light_id >> 16u; @@ -183,8 +184,8 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< return empty_reservoir(); } - let spatial_pixel_index = spatial_pixel_id.x + spatial_pixel_id.y * u32(view.viewport.z); - var spatial_reservoir = di_reservoirs_b[spatial_pixel_index]; + let packed_reservoir = textureLoad(di_reservoirs_b, spatial_pixel_id); + var spatial_reservoir = unpack_reservoir(packed_reservoir); if reservoir_valid(spatial_reservoir) { let resolved_light_sample = resolve_light_sample(spatial_reservoir.sample, light_sources[spatial_reservoir.sample.light_id >> 16u]); @@ -234,7 +235,6 @@ fn depth_ndc_to_view_z(ndc_depth: f32) -> f32 { #endif } -// Don't adjust the size of this struct without also adjusting DI_RESERVOIR_STRUCT_SIZE. struct Reservoir { sample: LightSample, confidence_weight: f32, @@ -253,6 +253,16 @@ fn reservoir_valid(reservoir: Reservoir) -> bool { return reservoir.sample.light_id != NULL_RESERVOIR_SAMPLE; } +fn pack_reservoir(reservoir: Reservoir) -> vec4 { + let weights = bitcast>(vec2(reservoir.confidence_weight, reservoir.unbiased_contribution_weight)); + return vec4(reservoir.sample.light_id, reservoir.sample.seed, weights); +} + +fn unpack_reservoir(packed: vec4) -> Reservoir { + let weights = bitcast>(packed.zw); + return Reservoir(LightSample(packed.x, packed.y), weights.x, weights.y); +} + struct ReservoirMergeResult { merged_reservoir: Reservoir, selected_sample_radiance: vec3, diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index f11ad77a46789..1df1268d9ca87 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -1,7 +1,7 @@ #define_import_path bevy_solari::sampling #import bevy_pbr::utils::{rand_f, rand_vec2f, rand_u, rand_range_u} -#import bevy_render::maths::{PI, PI_2, orthonormalize} +#import bevy_render::maths::{PI_2} #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} struct LightSample { From da6e6b5689fce2303fec81ebf2a14fc2351e2eed Mon Sep 17 00:00:00 2001 From: SparkyPotato Date: Thu, 31 Jul 2025 19:35:59 +0100 Subject: [PATCH 6/6] fix nans --- .../bevy_solari/src/realtime/restir_di.wgsl | 35 ++++++++++++------- crates/bevy_solari/src/scene/sampling.wgsl | 2 +- 2 files changed, 24 insertions(+), 13 deletions(-) diff --git a/crates/bevy_solari/src/realtime/restir_di.wgsl b/crates/bevy_solari/src/realtime/restir_di.wgsl index 12d2f9bf5e5bd..6845ff11bd242 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -42,7 +42,7 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - textureStore(di_reservoirs_b, global_id.xy, pack_reservoir(empty_reservoir())); + store_reservoir_b(global_id.xy, empty_reservoir()); return; } let gpixel = textureLoad(gbuffer, global_id.xy, 0); @@ -55,7 +55,7 @@ fn initial_and_temporal(@builtin(workgroup_id) workgroup_id: vec3, @builtin 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); - textureStore(di_reservoirs_b, global_id.xy, pack_reservoir(merge_result.merged_reservoir)); + store_reservoir_b(global_id.xy, merge_result.merged_reservoir); } @compute @workgroup_size(8, 8, 1) @@ -67,7 +67,7 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let depth = textureLoad(depth_buffer, global_id.xy, 0); if depth == 0.0 { - textureStore(di_reservoirs_a, global_id.xy, pack_reservoir(empty_reservoir())); + store_reservoir_a(global_id.xy, empty_reservoir()); textureStore(view_output, global_id.xy, vec4(vec3(0.0), 1.0)); return; } @@ -78,13 +78,12 @@ fn spatial_and_shade(@builtin(global_invocation_id) global_id: vec3) { let diffuse_brdf = base_color / PI; let emissive = rgb9e5_to_vec3_(gpixel.g); - let packed_input_reservoir = textureLoad(di_reservoirs_a, global_id.xy); - let input_reservoir = unpack_reservoir(packed_input_reservoir); + let input_reservoir = load_reservoir_b(global_id.xy); let spatial_reservoir = load_spatial_reservoir(global_id.xy, depth, world_position, world_normal, &rng); let merge_result = merge_reservoirs(input_reservoir, spatial_reservoir, world_position, world_normal, diffuse_brdf, &rng); let combined_reservoir = merge_result.merged_reservoir; - textureStore(di_reservoirs_a, global_id.xy, pack_reservoir(combined_reservoir)); + store_reservoir_a(global_id.xy, combined_reservoir); var pixel_color = merge_result.selected_sample_radiance * combined_reservoir.unbiased_contribution_weight; pixel_color *= view.exposure; @@ -128,7 +127,6 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 if reservoir_valid(reservoir) { let inverse_target_function = select(0.0, 1.0 / reservoir_target_function, reservoir_target_function > 0.0); reservoir.unbiased_contribution_weight = weight_sum * inverse_target_function; - reservoir.unbiased_contribution_weight *= trace_light_visibility(world_position, light_sample_world_position); } @@ -156,8 +154,7 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return empty_reservoir(); } - let packed_reservoir = textureLoad(di_reservoirs_a, temporal_pixel_id); - var temporal_reservoir = unpack_reservoir(packed_reservoir); + var temporal_reservoir = load_reservoir_a(temporal_pixel_id); // Check if the light selected in the previous frame no longer exists in the current frame (e.g. entity despawned) let previous_light_id = temporal_reservoir.sample.light_id >> 16u; @@ -167,7 +164,6 @@ fn load_temporal_reservoir(pixel_id: vec2, depth: f32, world_position: vec3 return empty_reservoir(); } temporal_reservoir.sample.light_id = (light_id << 16u) | triangle_id; - temporal_reservoir.confidence_weight = min(temporal_reservoir.confidence_weight, CONFIDENCE_WEIGHT_CAP); return temporal_reservoir; @@ -184,8 +180,7 @@ fn load_spatial_reservoir(pixel_id: vec2, depth: f32, world_position: vec3< return empty_reservoir(); } - let packed_reservoir = textureLoad(di_reservoirs_b, spatial_pixel_id); - var spatial_reservoir = unpack_reservoir(packed_reservoir); + var spatial_reservoir = load_reservoir_b(spatial_pixel_id); if reservoir_valid(spatial_reservoir) { let resolved_light_sample = resolve_light_sample(spatial_reservoir.sample, light_sources[spatial_reservoir.sample.light_id >> 16u]); @@ -258,11 +253,27 @@ fn pack_reservoir(reservoir: Reservoir) -> vec4 { return vec4(reservoir.sample.light_id, reservoir.sample.seed, weights); } +fn store_reservoir_a(pixel: vec2, reservoir: Reservoir) { + textureStore(di_reservoirs_a, pixel, pack_reservoir(reservoir)); +} + +fn store_reservoir_b(pixel: vec2, reservoir: Reservoir) { + textureStore(di_reservoirs_b, pixel, pack_reservoir(reservoir)); +} + fn unpack_reservoir(packed: vec4) -> Reservoir { let weights = bitcast>(packed.zw); return Reservoir(LightSample(packed.x, packed.y), weights.x, weights.y); } +fn load_reservoir_a(pixel: vec2) -> Reservoir { + return unpack_reservoir(textureLoad(di_reservoirs_a, pixel)); +} + +fn load_reservoir_b(pixel: vec2) -> Reservoir { + return unpack_reservoir(textureLoad(di_reservoirs_b, pixel)); +} + struct ReservoirMergeResult { merged_reservoir: Reservoir, selected_sample_radiance: vec3, diff --git a/crates/bevy_solari/src/scene/sampling.wgsl b/crates/bevy_solari/src/scene/sampling.wgsl index 1df1268d9ca87..298cf8ad679a0 100644 --- a/crates/bevy_solari/src/scene/sampling.wgsl +++ b/crates/bevy_solari/src/scene/sampling.wgsl @@ -1,7 +1,7 @@ #define_import_path bevy_solari::sampling #import bevy_pbr::utils::{rand_f, rand_vec2f, rand_u, rand_range_u} -#import bevy_render::maths::{PI_2} +#import bevy_render::maths::PI_2 #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} struct LightSample {