Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions crates/bevy_solari/src/pathtracer/pathtracer.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -41,9 +41,9 @@ fn pathtrace(@builtin(global_invocation_id) global_id: vec3<u32>) {
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;
Expand Down
4 changes: 2 additions & 2 deletions crates/bevy_solari/src/realtime/node.rs
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
32 changes: 16 additions & 16 deletions crates/bevy_solari/src/realtime/prepare.rs
Original file line number Diff line number Diff line change
Expand Up @@ -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,
};
Expand All @@ -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,
Expand Down Expand Up @@ -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");
Expand Down
6 changes: 3 additions & 3 deletions crates/bevy_solari/src/realtime/restir_gi.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -92,13 +92,13 @@ fn generate_initial_reservoir(world_position: vec3<f32>, world_normal: vec3<f32>
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;
Expand Down
5 changes: 3 additions & 2 deletions crates/bevy_solari/src/realtime/world_cache_compact.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
}
59 changes: 29 additions & 30 deletions crates/bevy_solari/src/realtime/world_cache_update.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -25,50 +25,49 @@ var<push_constant> constants: PushConstants;

@compute @workgroup_size(64, 1, 1)
fn sample_radiance(@builtin(workgroup_id) workgroup_id: vec3<u32>, @builtin(global_invocation_id) active_cell_id: vec3<u32>) {
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
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should probably keep the TODO

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nah I'm not a big fan of TODOs in code, and I'm not sure what shape it'll take if we ever add it. E.g. it might make more sense to have a dedicated step each frame to propagate finer LODs into coarser LODs, and skip the coarse LOD update entirely or something.

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<u32>) {
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<f32>, world_normal: vec3<f32>, workgroup_id: vec2<u32>, rng: ptr<function, u32>) -> vec3<f32> {
Expand Down