diff --git a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl index c2402deef7e40..e3fdcbd9a751c 100644 --- a/crates/bevy_solari/src/pathtracer/pathtracer.wgsl +++ b/crates/bevy_solari/src/pathtracer/pathtracer.wgsl @@ -41,9 +41,9 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3) { var bounce_was_perfect_reflection = true; var previous_normal = vec3(0.0); loop { - let ray_hit = trace_ray(ray_origin, ray_direction, ray_t_min, RAY_T_MAX, RAY_FLAG_NONE); - if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { - let ray_hit = resolve_ray_hit_full(ray_hit); + let ray = trace_ray(ray_origin, ray_direction, ray_t_min, RAY_T_MAX, RAY_FLAG_NONE); + if ray.kind != RAY_QUERY_INTERSECTION_NONE { + let ray_hit = resolve_ray_hit_full(ray); let wo = -ray_direction; var mis_weight = 1.0; diff --git a/crates/bevy_solari/src/realtime/node.rs b/crates/bevy_solari/src/realtime/node.rs index ac8fc69456e63..2900fa2e06c1f 100644 --- a/crates/bevy_solari/src/realtime/node.rs +++ b/crates/bevy_solari/src/realtime/node.rs @@ -173,8 +173,8 @@ impl ViewNode for SolariLightingNode { view_target.view, s.light_tile_samples.as_entire_binding(), s.light_tile_resolved_samples.as_entire_binding(), - &s.di_reservoirs_a.1, - &s.di_reservoirs_b.1, + &s.di_reservoirs_a, + &s.di_reservoirs_b, s.gi_reservoirs_a.as_entire_binding(), s.gi_reservoirs_b.as_entire_binding(), gbuffer, diff --git a/crates/bevy_solari/src/realtime/prepare.rs b/crates/bevy_solari/src/realtime/prepare.rs index b166c93b1754a..56d662064cee9 100644 --- a/crates/bevy_solari/src/realtime/prepare.rs +++ b/crates/bevy_solari/src/realtime/prepare.rs @@ -19,8 +19,8 @@ use bevy_render::texture::CachedTexture; use bevy_render::{ camera::ExtractedCamera, render_resource::{ - Buffer, BufferDescriptor, BufferUsages, Texture, TextureDescriptor, TextureDimension, - TextureFormat, TextureUsages, TextureView, TextureViewDescriptor, + Buffer, BufferDescriptor, BufferUsages, TextureDescriptor, TextureDimension, TextureFormat, + TextureUsages, TextureView, TextureViewDescriptor, }, renderer::RenderDevice, }; @@ -45,8 +45,8 @@ pub const WORLD_CACHE_SIZE: u64 = 2u64.pow(20); pub struct SolariLightingResources { pub light_tile_samples: Buffer, pub light_tile_resolved_samples: Buffer, - pub di_reservoirs_a: (Texture, TextureView), - pub di_reservoirs_b: (Texture, TextureView), + pub di_reservoirs_a: TextureView, + pub di_reservoirs_b: TextureView, pub gi_reservoirs_a: Buffer, pub gi_reservoirs_b: Buffer, pub world_cache_checksums: Buffer, @@ -121,18 +121,18 @@ pub fn prepare_solari_lighting_resources( }); 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) + 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: &[], + }) + .create_view(&TextureViewDescriptor::default()) }; let di_reservoirs_a = di_reservoirs("solari_lighting_di_reservoirs_a"); let di_reservoirs_b = di_reservoirs("solari_lighting_di_reservoirs_b"); diff --git a/crates/bevy_solari/src/realtime/restir_gi.wgsl b/crates/bevy_solari/src/realtime/restir_gi.wgsl index 8fc3089da0b33..971e67784ebb4 100644 --- a/crates/bevy_solari/src/realtime/restir_gi.wgsl +++ b/crates/bevy_solari/src/realtime/restir_gi.wgsl @@ -92,13 +92,13 @@ fn generate_initial_reservoir(world_position: vec3, world_normal: vec3 var reservoir = empty_reservoir(); let ray_direction = sample_uniform_hemisphere(world_normal, rng); - let ray_hit = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); + let ray = trace_ray(world_position, ray_direction, RAY_T_MIN, RAY_T_MAX, RAY_FLAG_NONE); - if ray_hit.kind == RAY_QUERY_INTERSECTION_NONE { + if ray.kind == RAY_QUERY_INTERSECTION_NONE { return reservoir; } - let sample_point = resolve_ray_hit_full(ray_hit); + let sample_point = resolve_ray_hit_full(ray); if all(sample_point.material.emissive != vec3(0.0)) { return reservoir; diff --git a/crates/bevy_solari/src/realtime/world_cache_compact.wgsl b/crates/bevy_solari/src/realtime/world_cache_compact.wgsl index a83ca817967b2..fd8164a63cb4e 100644 --- a/crates/bevy_solari/src/realtime/world_cache_compact.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_compact.wgsl @@ -77,7 +77,8 @@ fn compact_world_cache_write_active_cells( } if thread_index == 1023u && workgroup_id.x == 1023u { - world_cache_active_cells_count = compacted_index + u32(cell_active); - world_cache_active_cells_dispatch = vec3((world_cache_active_cells_count + 63u) / 64u, 1u, 1u); + let active_cell_count = compacted_index + u32(cell_active); + world_cache_active_cells_count = active_cell_count; + world_cache_active_cells_dispatch = vec3((active_cell_count + 63u) / 64u, 1u, 1u); } } diff --git a/crates/bevy_solari/src/realtime/world_cache_update.wgsl b/crates/bevy_solari/src/realtime/world_cache_update.wgsl index 4f4e98f84d857..e4b4bbfe36d82 100644 --- a/crates/bevy_solari/src/realtime/world_cache_update.wgsl +++ b/crates/bevy_solari/src/realtime/world_cache_update.wgsl @@ -25,50 +25,49 @@ var constants: PushConstants; @compute @workgroup_size(64, 1, 1) fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3, @builtin(global_invocation_id) active_cell_id: vec3) { - if active_cell_id.x < world_cache_active_cells_count { - let cell_index = world_cache_active_cell_indices[active_cell_id.x]; - let geometry_data = world_cache_geometry_data[cell_index]; - var rng = cell_index + constants.frame_index; + if active_cell_id.x >= world_cache_active_cells_count { return; } - // TODO: Initialize newly active cells with data from an adjacent LOD + let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + let geometry_data = world_cache_geometry_data[cell_index]; + var rng = cell_index + constants.frame_index; - var new_radiance = sample_random_light_ris(geometry_data.world_position, geometry_data.world_normal, workgroup_id.xy, &rng); + var new_radiance = sample_random_light_ris(geometry_data.world_position, geometry_data.world_normal, workgroup_id.xy, &rng); #ifndef NO_MULTIBOUNCE - let ray_direction = sample_cosine_hemisphere(geometry_data.world_normal, &rng); - let ray_hit = trace_ray(geometry_data.world_position, ray_direction, RAY_T_MIN, WORLD_CACHE_MAX_GI_RAY_DISTANCE, RAY_FLAG_NONE); - if ray_hit.kind != RAY_QUERY_INTERSECTION_NONE { - let ray_hit = resolve_ray_hit_full(ray_hit); - let cell_life = atomicLoad(&world_cache_life[cell_index]); - new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, cell_life, &rng); - } + let ray_direction = sample_cosine_hemisphere(geometry_data.world_normal, &rng); + let ray = trace_ray(geometry_data.world_position, ray_direction, RAY_T_MIN, WORLD_CACHE_MAX_GI_RAY_DISTANCE, RAY_FLAG_NONE); + if ray.kind != RAY_QUERY_INTERSECTION_NONE { + let ray_hit = resolve_ray_hit_full(ray); + let cell_life = atomicLoad(&world_cache_life[cell_index]); + new_radiance += ray_hit.material.base_color * query_world_cache(ray_hit.world_position, ray_hit.geometric_world_normal, view.world_position, cell_life, &rng); + } #endif - world_cache_active_cells_new_radiance[active_cell_id.x] = new_radiance; - } + world_cache_active_cells_new_radiance[active_cell_id.x] = new_radiance; } @compute @workgroup_size(64, 1, 1) fn blend_new_samples(@builtin(global_invocation_id) active_cell_id: vec3) { - if active_cell_id.x < world_cache_active_cells_count { - let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + if active_cell_id.x >= world_cache_active_cells_count { return; } - let old_radiance = world_cache_radiance[cell_index]; - let new_radiance = world_cache_active_cells_new_radiance[active_cell_id.x]; - let luminance_delta = world_cache_luminance_deltas[cell_index]; + let cell_index = world_cache_active_cell_indices[active_cell_id.x]; + var rng = cell_index + constants.frame_index; - // https://bsky.app/profile/gboisse.bsky.social/post/3m5blga3ftk2a - let sample_count = min(old_radiance.a + 1.0, WORLD_CACHE_MAX_TEMPORAL_SAMPLES); - let alpha = abs(luminance_delta) / max(luminance(old_radiance.rgb), 0.001); - let max_sample_count = mix(WORLD_CACHE_MAX_TEMPORAL_SAMPLES, 1.0, pow(saturate(alpha), 1.0 / 8.0)); - let blend_amount = 1.0 / min(sample_count, max_sample_count); + let old_radiance = world_cache_radiance[cell_index]; + let new_radiance = world_cache_active_cells_new_radiance[active_cell_id.x]; + let luminance_delta = world_cache_luminance_deltas[cell_index]; - let blended_radiance = mix(old_radiance.rgb, new_radiance, blend_amount); - let blended_luminance_delta = mix(luminance_delta, luminance(blended_radiance) - luminance(old_radiance.rgb), 1.0 / 8.0); + // https://bsky.app/profile/gboisse.bsky.social/post/3m5blga3ftk2a + let sample_count = min(old_radiance.a + 1.0, WORLD_CACHE_MAX_TEMPORAL_SAMPLES); + let alpha = abs(luminance_delta) / max(luminance(old_radiance.rgb), 0.001); + let max_sample_count = mix(WORLD_CACHE_MAX_TEMPORAL_SAMPLES, 1.0, pow(saturate(alpha), 1.0 / 8.0)); + let blend_amount = 1.0 / min(sample_count, max_sample_count); - world_cache_radiance[cell_index] = vec4(blended_radiance, sample_count); - world_cache_luminance_deltas[cell_index] = blended_luminance_delta; - } + let blended_radiance = mix(old_radiance.rgb, new_radiance, blend_amount); + let blended_luminance_delta = mix(luminance_delta, luminance(blended_radiance - old_radiance.rgb), 1.0 / 8.0); + + world_cache_radiance[cell_index] = vec4(blended_radiance, sample_count); + world_cache_luminance_deltas[cell_index] = blended_luminance_delta; } fn sample_random_light_ris(world_position: vec3, world_normal: vec3, workgroup_id: vec2, rng: ptr) -> vec3 {