diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index eaa432d8cb..c76840b61b 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 46a94a3ca2..0c694d7a53 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 b9a5bfa60c..c101157d75 100644 --- a/crates/bevy_solari/src/realtime/restir_di.wgsl +++ b/crates/bevy_solari/src/realtime/restir_di.wgsl @@ -4,22 +4,28 @@ #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, sample_random_light_contribution, 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; @@ -29,16 +35,29 @@ 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) { 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); @@ -51,7 +70,9 @@ fn initial_and_temporal(@builtin(global_invocation_id) global_id: vec3) { 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,30 @@ 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{ +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 < INITIAL_SAMPLES; i++) { - let light_sample = generate_random_light_sample(rng); + let light_contribution = sample_random_light_contribution(rng, 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 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 +165,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 +191,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 +242,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 +264,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 2b0cff5de7..3fd83d0587 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 eeed96ad8e..b5c1632386 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 e1f67ac1ed..1f4ef7e9c9 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,8 @@ 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 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); } @@ -65,25 +68,37 @@ struct LightSample { random: vec2, } -struct LightContribution { +struct IndependentLightSample { + sample: LightSample, radiance: vec3, inverse_pdf: f32, + pos_or_direction: vec4, + normal: vec3, } -fn generate_random_light_sample(rng: ptr) -> LightSample { +struct LightContribution { + sample: LightSample, + radiance: vec3, + inverse_pdf: f32, + ray_direction: vec4, +} + +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 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 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; - triangle_id = rand_range_u(triangle_count, rng); + 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); - return LightSample(vec2(light_id, triangle_id), random); + return light_contribution; } fn calculate_light_contribution(light_sample: LightSample, ray_origin: vec3, origin_world_normal: vec3) -> LightContribution { @@ -126,25 +141,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 +206,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); diff --git a/release-content/release-notes/bevy_solari.md b/release-content/release-notes/bevy_solari.md index 66f258eeb1..cf41c2a0bf 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)