Skip to content
Open
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
1 change: 1 addition & 0 deletions crates/bevy_solari/src/realtime/mod.rs
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@ pub struct SolariLightingPlugin;
impl Plugin for SolariLightingPlugin {
fn build(&self, app: &mut App) {
load_shader_library!(app, "gbuffer_utils.wgsl");
load_shader_library!(app, "realtime_bindings.wgsl");
load_shader_library!(app, "presample_light_tiles.wgsl");
embedded_asset!(app, "restir_di.wgsl");
embedded_asset!(app, "restir_gi.wgsl");
Expand Down
16 changes: 1 addition & 15 deletions crates/bevy_solari/src/realtime/presample_light_tiles.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -6,12 +6,7 @@
#import bevy_pbr::utils::{octahedral_encode, octahedral_decode}
#import bevy_render::view::View
#import bevy_solari::sampling::{generate_random_light_sample, LightSample, ResolvedLightSample}

@group(1) @binding(1) var<storage, read_write> light_tile_samples: array<LightSample>;
@group(1) @binding(2) var<storage, read_write> light_tile_resolved_samples: array<ResolvedLightSamplePacked>;
@group(1) @binding(12) var<uniform> view: View;
struct PushConstants { frame_index: u32, reset: u32 }
var<push_constant> constants: PushConstants;
#import bevy_solari::realtime_bindings::{light_tile_samples, light_tile_resolved_samples, view, constants, ResolvedLightSamplePacked}

@compute @workgroup_size(1024, 1, 1)
fn presample_light_tiles(@builtin(workgroup_id) workgroup_id: vec3<u32>, @builtin(local_invocation_index) sample_index: u32) {
Expand All @@ -25,15 +20,6 @@ fn presample_light_tiles(@builtin(workgroup_id) workgroup_id: vec3<u32>, @builti
light_tile_resolved_samples[i] = pack_resolved_light_sample(sample.resolved_light_sample);
}

struct ResolvedLightSamplePacked {
world_position_x: f32,
world_position_y: f32,
world_position_z: f32,
world_normal: u32,
radiance: u32,
inverse_pdf: f32,
}

fn pack_resolved_light_sample(sample: ResolvedLightSample) -> ResolvedLightSamplePacked {
return ResolvedLightSamplePacked(
sample.world_position.x,
Expand Down
63 changes: 63 additions & 0 deletions crates/bevy_solari/src/realtime/realtime_bindings.wgsl
Original file line number Diff line number Diff line change
@@ -0,0 +1,63 @@
#define_import_path bevy_solari::realtime_bindings

#import bevy_render::view::View
#import bevy_pbr::prepass_bindings::PreviousViewUniforms
#import bevy_solari::sampling::LightSample

@group(1) @binding(0) var view_output: texture_storage_2d<rgba16float, read_write>;
@group(1) @binding(1) var<storage, read_write> light_tile_samples: array<LightSample>;
@group(1) @binding(2) var<storage, read_write> light_tile_resolved_samples: array<ResolvedLightSamplePacked>;
@group(1) @binding(3) var di_reservoirs_a: texture_storage_2d<rgba32uint, read_write>;
@group(1) @binding(4) var di_reservoirs_b: texture_storage_2d<rgba32uint, read_write>;
@group(1) @binding(5) var<storage, read_write> gi_reservoirs_a: array<Reservoir>;
@group(1) @binding(6) var<storage, read_write> gi_reservoirs_b: array<Reservoir>;
@group(1) @binding(7) var gbuffer: texture_2d<u32>;
@group(1) @binding(8) var depth_buffer: texture_depth_2d;
@group(1) @binding(9) var motion_vectors: texture_2d<f32>;
@group(1) @binding(10) var previous_gbuffer: texture_2d<u32>;
@group(1) @binding(11) var previous_depth_buffer: texture_depth_2d;
@group(1) @binding(12) var<uniform> view: View;
@group(1) @binding(13) var<uniform> previous_view: PreviousViewUniforms;
@group(1) @binding(14) var<storage, read_write> world_cache_checksums: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
#ifdef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER
@group(1) @binding(15) var<storage, read_write> world_cache_life: array<u32, #{WORLD_CACHE_SIZE}>;
#else
@group(1) @binding(15) var<storage, read_write> world_cache_life: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
#endif
@group(1) @binding(16) var<storage, read_write> world_cache_radiance: array<vec4<f32>, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(17) var<storage, read_write> world_cache_geometry_data: array<WorldCacheGeometryData, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(18) var<storage, read_write> world_cache_luminance_deltas: array<f32, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(19) var<storage, read_write> world_cache_active_cells_new_radiance: array<vec3<f32>, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(20) var<storage, read_write> world_cache_a: array<u32, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(21) var<storage, read_write> world_cache_b: array<u32, 1024u>;
@group(1) @binding(22) var<storage, read_write> world_cache_active_cell_indices: array<u32, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(23) var<storage, read_write> world_cache_active_cells_count: u32;
struct PushConstants { frame_index: u32, reset: u32 }
var<push_constant> constants: PushConstants;

// Don't adjust the size of this struct without also adjusting `prepare::RESOLVED_LIGHT_SAMPLE_STRUCT_SIZE`.
struct ResolvedLightSamplePacked {
world_position_x: f32,
world_position_y: f32,
world_position_z: f32,
world_normal: u32,
radiance: u32,
inverse_pdf: f32,
}

// Don't adjust the size of this struct without also adjusting `prepare::GI_RESERVOIR_STRUCT_SIZE`.
struct Reservoir {
sample_point_world_position: vec3<f32>,
weight_sum: f32,
radiance: vec3<f32>,
confidence_weight: f32,
sample_point_world_normal: vec3<f32>,
unbiased_contribution_weight: f32,
}

struct WorldCacheGeometryData {
world_position: vec3<f32>,
padding_a: u32,
world_normal: vec3<f32>,
padding_b: u32
}
Original file line number Diff line number Diff line change
@@ -1,10 +1,7 @@
#import bevy_pbr::pbr_functions::{calculate_diffuse_color, calculate_F0}
#import bevy_render::view::View
#import bevy_solari::gbuffer_utils::gpixel_resolve

@group(1) @binding(7) var gbuffer: texture_2d<u32>;
@group(1) @binding(8) var depth_buffer: texture_depth_2d;
@group(1) @binding(12) var<uniform> view: View;
#import bevy_solari::realtime_bindings::{gbuffer, depth_buffer, view}

@group(2) @binding(0) var diffuse_albedo: texture_storage_2d<rgba8unorm, write>;
@group(2) @binding(1) var specular_albedo: texture_storage_2d<rgba8unorm, write>;
Expand Down
18 changes: 2 additions & 16 deletions crates/bevy_solari/src/realtime/restir_di.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -8,25 +8,11 @@
#import bevy_render::view::View
#import bevy_solari::brdf::{evaluate_brdf, evaluate_diffuse_brdf}
#import bevy_solari::gbuffer_utils::{gpixel_resolve, pixel_dissimilar, permute_pixel}
#import bevy_solari::presample_light_tiles::{ResolvedLightSamplePacked, unpack_resolved_light_sample}
#import bevy_solari::presample_light_tiles::unpack_resolved_light_sample
#import bevy_solari::sampling::{LightSample, calculate_resolved_light_contribution, resolve_and_calculate_light_contribution, resolve_light_sample, trace_light_visibility, balance_heuristic}
#import bevy_solari::scene_bindings::{light_sources, previous_frame_light_id_translations, LIGHT_NOT_PRESENT_THIS_FRAME}
#import bevy_solari::specular_gi::SPECULAR_GI_FOR_DI_ROUGHNESS_THRESHOLD

@group(1) @binding(0) var view_output: texture_storage_2d<rgba16float, read_write>;
@group(1) @binding(1) var<storage, read_write> light_tile_samples: array<LightSample>;
@group(1) @binding(2) var<storage, read_write> light_tile_resolved_samples: array<ResolvedLightSamplePacked>;
@group(1) @binding(3) var di_reservoirs_a: texture_storage_2d<rgba32uint, read_write>;
@group(1) @binding(4) var di_reservoirs_b: texture_storage_2d<rgba32uint, read_write>;
@group(1) @binding(7) var gbuffer: texture_2d<u32>;
@group(1) @binding(8) var depth_buffer: texture_depth_2d;
@group(1) @binding(9) var motion_vectors: texture_2d<f32>;
@group(1) @binding(10) var previous_gbuffer: texture_2d<u32>;
@group(1) @binding(11) var previous_depth_buffer: texture_depth_2d;
@group(1) @binding(12) var<uniform> view: View;
@group(1) @binding(13) var<uniform> previous_view: PreviousViewUniforms;
struct PushConstants { frame_index: u32, reset: u32 }
var<push_constant> constants: PushConstants;
#import bevy_solari::realtime_bindings::{view_output, light_tile_samples, light_tile_resolved_samples, di_reservoirs_a, di_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants, ResolvedLightSamplePacked}

const INITIAL_SAMPLES = 8u;
const SPATIAL_REUSE_RADIUS_PIXELS = 30.0;
Expand Down
24 changes: 1 addition & 23 deletions crates/bevy_solari/src/realtime/restir_gi.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -10,19 +10,7 @@
#import bevy_solari::sampling::{sample_random_light, trace_point_visibility, balance_heuristic}
#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, RAY_T_MIN, RAY_T_MAX}
#import bevy_solari::world_cache::{query_world_cache, WORLD_CACHE_CELL_LIFETIME}

@group(1) @binding(0) var view_output: texture_storage_2d<rgba16float, read_write>;
@group(1) @binding(5) var<storage, read_write> gi_reservoirs_a: array<Reservoir>;
@group(1) @binding(6) var<storage, read_write> gi_reservoirs_b: array<Reservoir>;
@group(1) @binding(7) var gbuffer: texture_2d<u32>;
@group(1) @binding(8) var depth_buffer: texture_depth_2d;
@group(1) @binding(9) var motion_vectors: texture_2d<f32>;
@group(1) @binding(10) var previous_gbuffer: texture_2d<u32>;
@group(1) @binding(11) var previous_depth_buffer: texture_depth_2d;
@group(1) @binding(12) var<uniform> view: View;
@group(1) @binding(13) var<uniform> previous_view: PreviousViewUniforms;
struct PushConstants { frame_index: u32, reset: u32 }
var<push_constant> constants: PushConstants;
#import bevy_solari::realtime_bindings::{view_output, gi_reservoirs_a, gi_reservoirs_b, gbuffer, depth_buffer, motion_vectors, previous_gbuffer, previous_depth_buffer, view, previous_view, constants, Reservoir}

const SPATIAL_REUSE_RADIUS_PIXELS = 30.0;
const CONFIDENCE_WEIGHT_CAP = 8.0;
Expand Down Expand Up @@ -217,16 +205,6 @@ fn isnan(x: f32) -> bool {
return (bitcast<u32>(x) & 0x7fffffffu) > 0x7f800000u;
}

// Don't adjust the size of this struct without also adjusting GI_RESERVOIR_STRUCT_SIZE.
struct Reservoir {
sample_point_world_position: vec3<f32>,
weight_sum: f32,
radiance: vec3<f32>,
confidence_weight: f32,
sample_point_world_normal: vec3<f32>,
unbiased_contribution_weight: f32,
}

fn empty_reservoir() -> Reservoir {
return Reservoir(
vec3(0.0),
Expand Down
9 changes: 1 addition & 8 deletions crates/bevy_solari/src/realtime/specular_gi.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -8,14 +8,7 @@
#import bevy_solari::sampling::{sample_random_light, random_emissive_light_pdf, sample_ggx_vndf, ggx_vndf_pdf, power_heuristic}
#import bevy_solari::scene_bindings::{trace_ray, resolve_ray_hit_full, ResolvedRayHitFull, RAY_T_MIN, RAY_T_MAX}
#import bevy_solari::world_cache::{query_world_cache, get_cell_size, WORLD_CACHE_CELL_LIFETIME}

@group(1) @binding(0) var view_output: texture_storage_2d<rgba16float, read_write>;
@group(1) @binding(5) var<storage, read_write> gi_reservoirs_a: array<Reservoir>;
@group(1) @binding(7) var gbuffer: texture_2d<u32>;
@group(1) @binding(8) var depth_buffer: texture_depth_2d;
@group(1) @binding(12) var<uniform> view: View;
struct PushConstants { frame_index: u32, reset: u32 }
var<push_constant> constants: PushConstants;
#import bevy_solari::realtime_bindings::{view_output, gi_reservoirs_a, gbuffer, depth_buffer, view, constants}

const DIFFUSE_GI_REUSE_ROUGHNESS_THRESHOLD: f32 = 0.4;
const SPECULAR_GI_FOR_DI_ROUGHNESS_THRESHOLD: f32 = 0.0225;
Expand Down
4 changes: 2 additions & 2 deletions crates/bevy_solari/src/realtime/world_cache_compact.wgsl
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
#import bevy_solari::world_cache::{
WORLD_CACHE_EMPTY_CELL,
#import bevy_solari::world_cache::WORLD_CACHE_EMPTY_CELL
#import bevy_solari::realtime_bindings::{
world_cache_life,
world_cache_checksums,
world_cache_radiance,
Expand Down
34 changes: 12 additions & 22 deletions crates/bevy_solari/src/realtime/world_cache_query.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -2,6 +2,18 @@

#import bevy_pbr::utils::rand_vec2f
#import bevy_render::maths::orthonormalize
#import bevy_solari::realtime_bindings::{
world_cache_life,
world_cache_checksums,
world_cache_radiance,
world_cache_geometry_data,
world_cache_luminance_deltas,
world_cache_a,
world_cache_b,
world_cache_active_cell_indices,
world_cache_active_cells_count,
WorldCacheGeometryData,
}

/// How responsive the world cache is to changes in lighting (higher is less responsive, lower is more responsive)
const WORLD_CACHE_MAX_TEMPORAL_SAMPLES: f32 = 32.0;
Expand All @@ -23,28 +35,6 @@ const WORLD_CACHE_POSITION_LOD_SCALE: f32 = 8.0;
/// Marker value for an empty cell
const WORLD_CACHE_EMPTY_CELL: u32 = 0u;

struct WorldCacheGeometryData {
world_position: vec3<f32>,
padding_a: u32,
world_normal: vec3<f32>,
padding_b: u32
}

@group(1) @binding(14) var<storage, read_write> world_cache_checksums: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
#ifdef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER
@group(1) @binding(15) var<storage, read_write> world_cache_life: array<u32, #{WORLD_CACHE_SIZE}>;
#else
@group(1) @binding(15) var<storage, read_write> world_cache_life: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
#endif
@group(1) @binding(16) var<storage, read_write> world_cache_radiance: array<vec4<f32>, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(17) var<storage, read_write> world_cache_geometry_data: array<WorldCacheGeometryData, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(18) var<storage, read_write> world_cache_luminance_deltas: array<f32, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(19) var<storage, read_write> world_cache_active_cells_new_radiance: array<vec3<f32>, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(20) var<storage, read_write> world_cache_a: array<u32, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(21) var<storage, read_write> world_cache_b: array<u32, 1024u>;
@group(1) @binding(22) var<storage, read_write> world_cache_active_cell_indices: array<u32, #{WORLD_CACHE_SIZE}>;
@group(1) @binding(23) var<storage, read_write> world_cache_active_cells_count: u32;

#ifndef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER
fn query_world_cache(world_position_in: vec3<f32>, world_normal: vec3<f32>, view_position: vec3<f32>, cell_lifetime: u32, rng: ptr<function, u32>) -> vec3<f32> {
var world_position = world_position_in;
Expand Down
10 changes: 5 additions & 5 deletions crates/bevy_solari/src/realtime/world_cache_update.wgsl
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,11 @@
WORLD_CACHE_DIRECT_LIGHT_SAMPLE_COUNT,
WORLD_CACHE_MAX_GI_RAY_DISTANCE,
query_world_cache,
}
#import bevy_solari::realtime_bindings::{
light_tile_resolved_samples,
view,
constants,
world_cache_active_cells_count,
world_cache_active_cell_indices,
world_cache_life,
Expand All @@ -18,11 +23,6 @@
world_cache_active_cells_new_radiance,
}

@group(1) @binding(2) var<storage, read_write> light_tile_resolved_samples: array<ResolvedLightSamplePacked>;
@group(1) @binding(12) var<uniform> view: View;
struct PushConstants { frame_index: u32, reset: u32 }
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 {
Expand Down