diff --git a/blade-helpers/src/hud.rs b/blade-helpers/src/hud.rs index 880b745f..6273083b 100644 --- a/blade-helpers/src/hud.rs +++ b/blade-helpers/src/hud.rs @@ -17,6 +17,7 @@ impl ExposeHud for blade_render::RayConfig { ui.add( egui::widgets::Slider::new(&mut self.temporal_history, 0..=50).text("Temporal history"), ); + ui.checkbox(&mut self.spatial_pass, "Spatial pass"); ui.add(egui::widgets::Slider::new(&mut self.spatial_taps, 0..=10).text("Spatial taps")); ui.add( egui::widgets::Slider::new(&mut self.spatial_tap_history, 0..=50) diff --git a/blade-helpers/src/lib.rs b/blade-helpers/src/lib.rs index db7f2ec8..bf42be0e 100644 --- a/blade-helpers/src/lib.rs +++ b/blade-helpers/src/lib.rs @@ -13,6 +13,7 @@ pub fn default_ray_config() -> blade_render::RayConfig { environment_importance_sampling: false, temporal_tap: true, temporal_history: 10, + spatial_pass: true, spatial_taps: 1, spatial_tap_history: 10, spatial_radius: 20, diff --git a/blade-render/code/fill-gbuf.wgsl b/blade-render/code/fill-gbuf.wgsl index 180cd628..be7b7831 100644 --- a/blade-render/code/fill-gbuf.wgsl +++ b/blade-render/code/fill-gbuf.wgsl @@ -177,6 +177,9 @@ fn main(@builtin(global_invocation_id) global_id: vec3) { } if (WRITE_DEBUG_IMAGE) { + if (debug.view_mode == DebugMode_Depth) { + textureStore(out_debug, global_id.xy, vec4(1.0 / depth)); + } if (debug.view_mode == DebugMode_DiffuseAlbedoTexture) { textureStore(out_debug, global_id.xy, vec4(albedo, 0.0)); } diff --git a/blade-render/code/ray-trace.wgsl b/blade-render/code/ray-trace.wgsl index 168b93cd..f88e6d2e 100644 --- a/blade-render/code/ray-trace.wgsl +++ b/blade-render/code/ray-trace.wgsl @@ -68,13 +68,14 @@ struct LiveReservoir { selected_uv: vec2, selected_light_index: u32, selected_target_score: f32, - selected_radiance: vec3, + /// Material output without visibility. + selected_color: vec3, weight_sum: f32, history: f32, } fn compute_target_score(radiance: vec3) -> f32 { - return dot(radiance, vec3(0.3, 0.4, 0.3)); + return dot(radiance, vec3(0.212, 0.7152, 0.0722)); } fn get_reservoir_index(pixel: vec2, camera: CameraParams) -> i32 { @@ -96,10 +97,10 @@ fn bump_reservoir(r: ptr, history: f32) { } fn make_reservoir(ls: LightSample, light_index: u32, brdf: vec3) -> LiveReservoir { var r: LiveReservoir; - r.selected_radiance = ls.radiance * brdf; + r.selected_color = ls.radiance * brdf; r.selected_uv = ls.uv; r.selected_light_index = light_index; - r.selected_target_score = compute_target_score(r.selected_radiance); + r.selected_target_score = compute_target_score(r.selected_color); r.weight_sum = r.selected_target_score / ls.pdf; r.history = 1.0; return r; @@ -111,7 +112,7 @@ fn merge_reservoir(r: ptr, other: LiveReservoir, random (*r).selected_light_index = other.selected_light_index; (*r).selected_uv = other.selected_uv; (*r).selected_target_score = other.selected_target_score; - (*r).selected_radiance = other.selected_radiance; + (*r).selected_color = other.selected_color; return true; } else { return false; @@ -129,7 +130,7 @@ fn unpack_reservoir(f: StoredReservoir, max_history: u32, radiance: vec3) - r.selected_light_index = f.light_index; r.selected_uv = f.light_uv; r.selected_target_score = f.target_score; - r.selected_radiance = radiance; + r.selected_color = radiance; let history = min(f.confidence, f32(max_history)); r.weight_sum = f.contribution_weight * f.target_score * history; r.history = history; @@ -158,6 +159,7 @@ var t_flat_normal: texture_2d; var t_prev_flat_normal: texture_2d; var t_motion: texture_2d; var out_diffuse: texture_storage_2d; +var in_diffuse: texture_2d; var out_debug: texture_storage_2d; fn sample_circle(random: f32) -> vec2 { @@ -181,11 +183,6 @@ fn map_equirect_uv_to_dir(uv: vec2) -> vec3 { return vec3(cos(yaw) * sin(pitch), sin(yaw), cos(yaw) * cos(pitch)); } -fn evaluate_environment(dir: vec3) -> vec3 { - let uv = map_equirect_dir_to_uv(dir); - return textureSampleLevel(env_map, sampler_linear, uv, 0.0).xyz; -} - fn sample_light_from_sphere(rng: ptr) -> LightSample { let a = random_gen(rng); let h = 1.0 - 2.0 * random_gen(rng); // make sure to allow h==1 @@ -227,13 +224,32 @@ fn read_prev_surface(pixel: vec2) -> Surface { return surface; } -fn evaluate_brdf(surface: Surface, dir: vec3) -> f32 { +fn evaluate_material(surface: Surface, dir: vec3) -> f32 { let lambert_brdf = 1.0 / PI; let lambert_term = qrot(qinv(surface.basis), dir).z; //Note: albedo not modulated return lambert_brdf * max(0.0, lambert_term); } +fn evaluate_environment(direction: vec3) -> vec3 { + let uv = map_equirect_dir_to_uv(direction); + return textureSampleLevel(env_map, sampler_nearest, uv, 0.0).xyz; +} + +fn evaluate_reflected_light(surface: Surface, light_index: u32, light_uv: vec2) -> vec3 { + if (light_index != 0u) { + return vec3(0.0); + } + let direction = map_equirect_uv_to_dir(light_uv); + let brdf = evaluate_material(surface, direction); + if (brdf <= 0.0) { + return vec3(0.0); + } + // Note: returns radiance not modulated by albedo + let radiance = textureSampleLevel(env_map, sampler_nearest, light_uv, 0.0).xyz; + return radiance * brdf; +} + fn check_ray_occluded(acs: acceleration_structure, position: vec3, direction: vec3, debug_len: f32) -> bool { var rq: ray_query; let flags = RAY_FLAG_TERMINATE_ON_FIRST_HIT | RAY_FLAG_CULL_NO_OPAQUE; @@ -251,20 +267,6 @@ fn check_ray_occluded(acs: acceleration_structure, position: vec3, directio return occluded; } -fn evaluate_reflected_light(surface: Surface, light_index: u32, light_uv: vec2) -> vec3 { - if (light_index != 0u) { - return vec3(0.0); - } - let direction = map_equirect_uv_to_dir(light_uv); - let brdf = evaluate_brdf(surface, direction); - if (brdf <= 0.0) { - return vec3(0.0); - } - // Note: returns radiance not modulated by albedo - let radiance = textureSampleLevel(env_map, sampler_nearest, light_uv, 0.0).xyz; - return radiance * brdf; -} - fn get_prev_pixel(pixel: vec2, pos_world: vec3) -> vec2 { if (USE_MOTION_VECTORS && parameters.use_motion_vectors != 0u) { let motion = textureLoad(t_motion, pixel, 0).xy / MOTION_SCALE; @@ -293,7 +295,7 @@ fn estimate_target_score_with_occlusion( if (dot(direction, surface.flat_normal) <= 0.0) { return TargetScore(); } - let brdf = evaluate_brdf(surface, direction); + let brdf = evaluate_material(surface, direction); if (brdf <= 0.0) { return TargetScore(); } @@ -313,7 +315,7 @@ fn evaluate_sample(ls: LightSample, surface: Surface, start_pos: vec3, debu return 0.0; } - let brdf = evaluate_brdf(surface, dir); + let brdf = evaluate_material(surface, dir); if (brdf <= 0.0) { return 0.0; } @@ -330,31 +332,9 @@ fn evaluate_sample(ls: LightSample, surface: Surface, start_pos: vec3, debu return brdf; } -fn ratio(a: f32, b: f32) -> f32 { - return select(0.0, a / (a+b), a+b > 0.0); -} - -struct RestirOutput { - radiance: vec3, -} - -fn compute_restir(surface: Surface, pixel: vec2, rng: ptr, enable_debug: bool) -> RestirOutput { - let ray_dir = get_ray_direction(camera, pixel); - let pixel_index = get_reservoir_index(pixel, camera); - if (surface.depth == 0.0) { - reservoirs[pixel_index] = StoredReservoir(); - let env = evaluate_environment(ray_dir); - return RestirOutput(env); - } - - if (WRITE_DEBUG_IMAGE && debug.view_mode == DebugMode_Depth) { - textureStore(out_debug, pixel, vec4(1.0 / surface.depth)); - } +fn produce_canonical(surface: Surface, position: vec3, rng: ptr, enable_debug: bool) -> LiveReservoir { + var reservoir = LiveReservoir(); let debug_len = select(0.0, surface.depth * 0.2, enable_debug); - let position = camera.position + surface.depth * ray_dir; - let normal = qrot(surface.basis, vec3(0.0, 0.0, 1.0)); - - var canonical = LiveReservoir(); for (var i = 0u; i < parameters.num_environment_samples; i += 1u) { var ls: LightSample; if (parameters.environment_importance_sampling != 0u) { @@ -366,12 +346,27 @@ fn compute_restir(surface: Surface, pixel: vec2, rng: ptr 0.0) { let other = make_reservoir(ls, 0u, vec3(brdf)); - merge_reservoir(&canonical, other, random_gen(rng)); + merge_reservoir(&reservoir, other, random_gen(rng)); } else { - bump_reservoir(&canonical, 1.0); + bump_reservoir(&reservoir, 1.0); } } + return reservoir; +} + +struct Neighborhood { + reservoir_indices: array, + count: u32, +} + +fn gather_neighborhood_temporal( + surface: Surface, position: vec3, pixel: vec2, rng: ptr +) -> Neighborhood { + if (surface.depth == 0.0 || parameters.temporal_tap == 0u) { + return Neighborhood(); + } + let center_coord = get_prev_pixel(pixel, position); let center_pixel = vec2(center_coord); // Trick to start with closer pixels: we derive the "further" @@ -379,30 +374,12 @@ fn compute_restir(surface: Surface, pixel: vec2, rng: ptr(center_coord - 0.5) + vec2(center_coord + 0.5) - center_pixel; // First, gather the list of reservoirs to merge with - var accepted_reservoir_indices = array(); - var accepted_count = 0u; - var temporal_index = ~0u; - let num_temporal_candidates = parameters.temporal_tap * FACTOR_TEMPORAL_CANDIDATES; - let num_candidates = num_temporal_candidates + parameters.spatial_taps * FACTOR_SPATIAL_CANDIDATES; - let max_samples = min(MAX_RESERVOIRS, 1u + parameters.spatial_taps); - - for (var tap = 0u; tap < num_candidates && accepted_count < max_samples; tap += 1u) { - var other_pixel = center_pixel; - if (tap < num_temporal_candidates) { - if (temporal_index < tap) { - continue; - } - let mask = vec2(tap) & vec2(1u, 2u); - other_pixel = select(center_pixel, further_pixel, mask != vec2(0u)); - } else { - let r0 = max(center_pixel - vec2(parameters.spatial_radius), vec2(0)); - let r1 = min(center_pixel + vec2(parameters.spatial_radius + 1), vec2(prev_camera.target_size)); - other_pixel = vec2(mix(vec2(r0), vec2(r1), vec2(random_gen(rng), random_gen(rng)))); - let diff = other_pixel - center_pixel; - if (dot(diff, diff) < MIN_SPATIAL_REUSE_DISTANCE) { - continue; - } - } + var nh = Neighborhood(); + let num_candidates = parameters.temporal_tap * FACTOR_TEMPORAL_CANDIDATES; + + for (var tap = 0u; tap < num_candidates && nh.count == 0u; tap += 1u) { + let mask = vec2(tap) & vec2(1u, 2u); + let other_pixel = select(center_pixel, further_pixel, mask != vec2(0u)); let other_index = get_reservoir_index(other_pixel, prev_camera); if (other_index < 0) { @@ -419,33 +396,98 @@ fn compute_restir(surface: Surface, pixel: vec2, rng: ptr(f32(nh.count)); + textureStore(out_debug, pixel, color); + } + + return nh; +} + +fn gather_neighborhood_spatial( + surface: Surface, pixel: vec2, rng: ptr +) -> Neighborhood { + if (surface.depth == 0.0 || parameters.spatial_taps == 0u) { + return Neighborhood(); + } + + // First, gather the list of reservoirs to merge with + var nh = Neighborhood(); + let num_candidates = parameters.spatial_taps * FACTOR_SPATIAL_CANDIDATES; + let max_samples = min(MAX_RESERVOIRS, parameters.spatial_taps); + + for (var tap = 0u; tap < num_candidates && nh.count < max_samples; tap += 1u) { + let r0 = max(pixel - vec2(parameters.spatial_radius), vec2(0)); + let r1 = min(pixel + vec2(parameters.spatial_radius + 1), vec2(prev_camera.target_size)); + let other_pixel = vec2(mix(vec2(r0), vec2(r1), vec2(random_gen(rng), random_gen(rng)))); + let diff = other_pixel - pixel; + if (dot(diff, diff) < MIN_SPATIAL_REUSE_DISTANCE) { + continue; + } + + let other_index = get_reservoir_index(other_pixel, prev_camera); + if (other_index < 0) { + continue; + } + if (prev_reservoirs[other_index].confidence == 0.0) { + continue; + } + + let other_surface = read_prev_surface(other_pixel); + let compatibility = compare_surfaces(surface, other_surface); + if (compatibility < 0.1) { + // if the surfaces are too different, there is no trust in this sample + continue; } - accepted_reservoir_indices[accepted_count] = other_index; - accepted_count += 1u; + + nh.reservoir_indices[nh.count] = other_index; + nh.count += 1u; } - if (WRITE_DEBUG_IMAGE && debug.view_mode == DebugMode_SampleReuse) { + if (WRITE_DEBUG_IMAGE && debug.view_mode == DebugMode_SpatialReuse) { var color = vec4(0.0); - for (var i = 0u; i < min(3u, accepted_count); i += 1u) { + for (var i = 0u; i < min(3u, nh.count); i += 1u) { color[i] = 1.0; } textureStore(out_debug, pixel, color); } - // Next, evaluate the MIS of each of the samples versus the canonical one. + return nh; +} + +struct RestirOutput { + reservoir: StoredReservoir, + radiance: vec3, +} + +fn ratio(a: f32, b: f32) -> f32 { + return select(0.0, a / (a+b), a+b > 0.0); +} + +fn compute_restir( + surface: Surface, ray_dir: vec3, + canonical: LiveReservoir, nh: Neighborhood, max_history: u32, + rng: ptr, enable_debug: bool, +) -> RestirOutput { + let enable_restir_debug = (debug.draw_flags & DebugDrawFlags_RESTIR) != 0u && enable_debug; + let position = camera.position + surface.depth * ray_dir; + let debug_len = select(0.0, surface.depth * 0.2, enable_restir_debug); + var accepted_reservoir_indices = nh.reservoir_indices; + // evaluate the MIS of each of the samples versus the canonical one. var reservoir = LiveReservoir(); var color_and_weight = vec4(0.0); - let mis_scale = 1.0 / (f32(accepted_count) + parameters.defensive_mis); - var mis_canonical = select(mis_scale * parameters.defensive_mis, 1.0, accepted_count == 0u || parameters.use_pairwise_mis == 0u); - let inv_count = 1.0 / f32(accepted_count); + let mis_scale = 1.0 / (f32(nh.count) + parameters.defensive_mis); + var mis_canonical = select(mis_scale * parameters.defensive_mis, 1.0, nh.count == 0u || parameters.use_pairwise_mis == 0u); + let inv_count = 1.0 / f32(nh.count); - for (var rid = 0u; rid < accepted_count; rid += 1u) { + for (var rid = 0u; rid < nh.count; rid += 1u) { let neighbor_index = accepted_reservoir_indices[rid]; let neighbor = prev_reservoirs[neighbor_index]; - let max_history = select(parameters.spatial_tap_history, parameters.temporal_history, rid == temporal_index); var other: LiveReservoir; if (parameters.use_pairwise_mis != 0u) { let neighbor_pixel = get_pixel_from_reservoir_index(neighbor_index, prev_camera); @@ -470,15 +512,15 @@ fn compute_restir(surface: Surface, pixel: vec2, rng: ptr(color, 1.0); } if (other.weight_sum <= 0.0) { @@ -489,23 +531,23 @@ fn compute_restir(surface: Surface, pixel: vec2, rng: ptr(cw * canonical.selected_radiance, 1.0); + let cw = canonical_mod.weight_sum / max(canonical_mod.selected_target_score, 0.1); + color_and_weight += canonical_mod.weight_sum * vec4(cw * canonical_mod.selected_color, 1.0); } - merge_reservoir(&reservoir, canonical, random_gen(rng)); + merge_reservoir(&reservoir, canonical_mod, random_gen(rng)); let effective_history = select(reservoir.history, 1.0, parameters.use_pairwise_mis != 0); - let stored = pack_reservoir_detail(reservoir, effective_history); - reservoirs[pixel_index] = stored; var ro = RestirOutput(); + ro.reservoir = pack_reservoir_detail(reservoir, effective_history); if (DECOUPLED_SHADING) { ro.radiance = color_and_weight.xyz / max(color_and_weight.w, 0.001); } else { - ro.radiance = stored.contribution_weight * reservoir.selected_radiance; + ro.radiance = ro.reservoir.contribution_weight * reservoir.selected_color; } return ro; } @@ -520,15 +562,61 @@ fn main(@builtin(global_invocation_id) global_id: vec3) { var rng = random_init(global_index, parameters.frame_index); let surface = read_surface(vec2(global_id.xy)); + let pixel = vec2(global_id.xy); + let pixel_index = get_reservoir_index(pixel, camera); + let ray_dir = get_ray_direction(camera, pixel); let enable_debug = DEBUG_MODE && all(global_id.xy == debug.mouse_pos); - let enable_restir_debug = (debug.draw_flags & DebugDrawFlags_RESTIR) != 0u && enable_debug; - let ro = compute_restir(surface, vec2(global_id.xy), &rng, enable_restir_debug); + if (surface.depth == 0.0) { + reservoirs[pixel_index] = StoredReservoir(); + let radiance = evaluate_environment(ray_dir); + textureStore(out_diffuse, pixel, vec4(radiance, 1.0)); + return; + } + + let position = camera.position + surface.depth * ray_dir; + let neighborhood = gather_neighborhood_temporal(surface, position, pixel, &rng); + let canonical = produce_canonical(surface, position, &rng, enable_debug); + let ro = compute_restir(surface, ray_dir, canonical, neighborhood, parameters.temporal_history, &rng, enable_debug); + + reservoirs[pixel_index] = ro.reservoir; + textureStore(out_diffuse, pixel, vec4(ro.radiance, 1.0)); +} + +@compute @workgroup_size(8, 4) +fn main_spatial(@builtin(global_invocation_id) global_id: vec3) { + if (any(global_id.xy >= camera.target_size)) { + return; + } + + let surface = read_surface(vec2(global_id.xy)); + let pixel = vec2(global_id.xy); + let pixel_index = get_reservoir_index(pixel, camera); + let in_radiance = textureLoad(in_diffuse, pixel, 0).xyz; + + if (surface.depth == 0.0) { + reservoirs[pixel_index] = StoredReservoir(); + textureStore(out_diffuse, pixel, vec4(in_radiance, 1.0)); + return; + } + + let enable_debug = DEBUG_MODE && all(global_id.xy == debug.mouse_pos); + let global_index = global_id.y * camera.target_size.x + global_id.x; + var rng = random_init(global_index, parameters.frame_index * 2u); + + let neighborhood = gather_neighborhood_spatial(surface, pixel, &rng); + let old_reservoir = prev_reservoirs[pixel_index]; + let old_selected_color = in_radiance / max(old_reservoir.contribution_weight, 0.01); + let canonical = unpack_reservoir(old_reservoir, ~0u, old_selected_color); + let ray_dir = get_ray_direction(camera, pixel); + let ro = compute_restir(surface, ray_dir, canonical, neighborhood, parameters.spatial_tap_history, &rng, enable_debug); + reservoirs[pixel_index] = ro.reservoir; let color = ro.radiance; + if (enable_debug) { debug_buf.variance.color_sum += color; debug_buf.variance.color2_sum += color * color; debug_buf.variance.count += 1u; } - textureStore(out_diffuse, global_id.xy, vec4(color, 1.0)); + textureStore(out_diffuse, pixel, vec4(color, 1.0)); } diff --git a/blade-render/src/render/mod.rs b/blade-render/src/render/mod.rs index 8c9677bf..67952b75 100644 --- a/blade-render/src/render/mod.rs +++ b/blade-render/src/render/mod.rs @@ -56,7 +56,8 @@ pub enum DebugMode { ShadingNormal = 7, Motion = 8, HitConsistency = 9, - SampleReuse = 10, + TemporalReuse = 10, + SpatialReuse = 11, Variance = 15, } @@ -97,6 +98,7 @@ pub struct RayConfig { pub environment_importance_sampling: bool, pub temporal_tap: bool, pub temporal_history: u32, + pub spatial_pass: bool, pub spatial_taps: u32, pub spatial_tap_history: u32, pub spatial_radius: u32, @@ -319,9 +321,9 @@ struct Blur { pub struct Renderer { shaders: Shaders, targets: RestirTargets, - post_proc_input_index: usize, fill_pipeline: blade_graphics::ComputePipeline, main_pipeline: blade_graphics::ComputePipeline, + spatial_pipeline: blade_graphics::ComputePipeline, post_proc_pipeline: blade_graphics::RenderPipeline, blur: Blur, acceleration_structure: blade_graphics::AccelerationStructure, @@ -339,6 +341,7 @@ pub struct Renderer { surface_info: blade_graphics::SurfaceInfo, frame_index: usize, frame_scene_built: usize, + frame_resize: usize, is_frozen: bool, //TODO: refactor `ResourceArray` to not carry the freelist logic // This way we can embed user info into the allocator. @@ -426,6 +429,7 @@ struct MainData { reservoirs: blade_graphics::BufferPiece, prev_reservoirs: blade_graphics::BufferPiece, out_diffuse: blade_graphics::TextureView, + in_diffuse: blade_graphics::TextureView, out_debug: blade_graphics::TextureView, } @@ -527,6 +531,7 @@ impl Shaders { struct ShaderPipelines { fill: blade_graphics::ComputePipeline, main: blade_graphics::ComputePipeline, + spatial: blade_graphics::ComputePipeline, temporal_accum: blade_graphics::ComputePipeline, a_trous: blade_graphics::ComputePipeline, post_proc: blade_graphics::RenderPipeline, @@ -548,7 +553,7 @@ impl ShaderPipelines { compute: shader.at("main"), }) } - fn create_ray_trace( + fn create_ray_trace_main( shader: &blade_graphics::Shader, gpu: &blade_graphics::Context, ) -> blade_graphics::ComputePipeline { @@ -565,6 +570,21 @@ impl ShaderPipelines { }) } + fn create_ray_trace_spatial( + shader: &blade_graphics::Shader, + gpu: &blade_graphics::Context, + ) -> blade_graphics::ComputePipeline { + //Note: technically the spatial reuse shader could use a reduced + // set of resources, but in terms of code reuse it's much easier + // to just bind the new data as the old one. + let layout = ::layout(); + gpu.create_compute_pipeline(blade_graphics::ComputePipelineDesc { + name: "ray-trace-spatial", + data_layouts: &[&layout], + compute: shader.at("main_spatial"), + }) + } + fn create_temporal_accum( shader: &blade_graphics::Shader, gpu: &blade_graphics::Context, @@ -620,7 +640,8 @@ impl ShaderPipelines { let sh_a_trous = shader_man[shaders.a_trous].raw.as_ref().unwrap(); Ok(Self { fill: Self::create_gbuf_fill(shader_man[shaders.fill_gbuf].raw.as_ref().unwrap(), gpu), - main: Self::create_ray_trace(sh_main, gpu), + main: Self::create_ray_trace_main(sh_main, gpu), + spatial: Self::create_ray_trace_spatial(sh_main, gpu), temporal_accum: Self::create_temporal_accum(sh_a_trous, gpu), a_trous: Self::create_a_trous(sh_a_trous, gpu), post_proc: Self::create_post_proc( @@ -709,9 +730,9 @@ impl Renderer { Self { shaders, targets, - post_proc_input_index: 0, fill_pipeline: sp.fill, main_pipeline: sp.main, + spatial_pipeline: sp.spatial, post_proc_pipeline: sp.post_proc, blur: Blur { temporal_accum_pipeline: sp.temporal_accum, @@ -732,6 +753,7 @@ impl Renderer { surface_info: config.surface_info, frame_index: 0, frame_scene_built: 0, + frame_resize: 0, is_frozen: false, texture_resource_lookup: HashMap::default(), } @@ -760,6 +782,7 @@ impl Renderer { gpu.destroy_compute_pipeline(&mut self.blur.a_trous_pipeline); gpu.destroy_compute_pipeline(&mut self.fill_pipeline); gpu.destroy_compute_pipeline(&mut self.main_pipeline); + gpu.destroy_compute_pipeline(&mut self.spatial_pipeline); gpu.destroy_render_pipeline(&mut self.post_proc_pipeline); } @@ -801,7 +824,8 @@ impl Renderer { shader.get_struct_size("StoredReservoir"), self.reservoir_size ); - self.main_pipeline = ShaderPipelines::create_ray_trace(shader, gpu); + self.main_pipeline = ShaderPipelines::create_ray_trace_main(shader, gpu); + self.spatial_pipeline = ShaderPipelines::create_ray_trace_spatial(shader, gpu); } } if self.shaders.a_trous != old.a_trous { @@ -855,6 +879,7 @@ impl Renderer { self.surface_size = size; self.targets.destroy(gpu); self.targets = RestirTargets::new(size, self.reservoir_size, encoder, gpu); + self.frame_resize = self.frame_index + 1; } #[profiling::function] @@ -1095,7 +1120,7 @@ impl Renderer { } self.debug.update_entry(&mut transfer); - if config.reset_reservoirs { + if config.reset_reservoirs || self.frame_resize == self.frame_index { if !config.debug_draw { self.debug.reset_lines(&mut transfer); } @@ -1114,7 +1139,6 @@ impl Renderer { } self.is_frozen = config.frozen; self.targets.camera_params[self.frame_index % 2] = self.make_camera_params(camera); - self.post_proc_input_index = self.frame_index % 2; } /// Ray trace the scene. @@ -1122,14 +1146,13 @@ impl Renderer { /// The result is stored internally in an HDR render target. #[profiling::function] pub fn ray_trace( - &self, + &mut self, command_encoder: &mut blade_graphics::CommandEncoder, debug_config: DebugConfig, ray_config: RayConfig, ) { let debug = self.make_debug_params(&debug_config); let (cur, prev) = self.work_indices(); - assert_eq!(cur, self.post_proc_input_index); if let mut pass = command_encoder.compute("fill-gbuf") { let mut pc = pass.with(&self.fill_pipeline); @@ -1158,7 +1181,22 @@ impl Renderer { pc.dispatch(groups); } - if let mut pass = command_encoder.compute("ray-trace") { + let parameters = MainParams { + frame_index: self.frame_index as u32, + num_environment_samples: ray_config.num_environment_samples, + environment_importance_sampling: ray_config.environment_importance_sampling as u32, + temporal_tap: ray_config.temporal_tap as u32, + temporal_history: ray_config.temporal_history, + spatial_taps: ray_config.spatial_taps, + spatial_tap_history: ray_config.spatial_tap_history, + spatial_radius: ray_config.spatial_radius, + t_start: ray_config.t_start, + use_pairwise_mis: ray_config.pairwise_mis as u32, + defensive_mis: ray_config.defensive_mis, + use_motion_vectors: (self.frame_scene_built >= self.frame_index) as u32, + }; + + if let mut pass = command_encoder.compute("ray-trace-main") { let mut pc = pass.with(&self.main_pipeline); let groups = self.main_pipeline.get_dispatch_for(self.surface_size); pc.bind( @@ -1167,21 +1205,7 @@ impl Renderer { camera: self.targets.camera_params[cur], prev_camera: self.targets.camera_params[prev], debug, - parameters: MainParams { - frame_index: self.frame_index as u32, - num_environment_samples: ray_config.num_environment_samples, - environment_importance_sampling: ray_config.environment_importance_sampling - as u32, - temporal_tap: ray_config.temporal_tap as u32, - temporal_history: ray_config.temporal_history, - spatial_taps: ray_config.spatial_taps, - spatial_tap_history: ray_config.spatial_tap_history, - spatial_radius: ray_config.spatial_radius, - t_start: ray_config.t_start, - use_pairwise_mis: ray_config.pairwise_mis as u32, - defensive_mis: ray_config.defensive_mis, - use_motion_vectors: (self.frame_scene_built >= self.frame_index) as u32, - }, + parameters, acc_struct: self.acceleration_structure, prev_acc_struct: if self.frame_scene_built < self.frame_index || self.prev_acceleration_structure @@ -1203,26 +1227,71 @@ impl Renderer { t_prev_flat_normal: self.targets.flat_normal.views[prev], t_motion: self.targets.motion.views[0], debug_buf: self.debug.buffer_resource(), - reservoirs: self.targets.reservoir_buf[cur].into(), - prev_reservoirs: self.targets.reservoir_buf[prev].into(), - out_diffuse: self.targets.light_diffuse.views[cur], + reservoirs: self.targets.reservoir_buf[1].into(), + prev_reservoirs: self.targets.reservoir_buf[0].into(), + out_diffuse: self.targets.light_diffuse.views[1], + in_diffuse: self.targets.light_diffuse.views[0], + out_debug: self.targets.debug.views[0], + }, + ); + pc.dispatch(groups); + } + //Note: output reservoirs and lighting is in [1] + + if ray_config.spatial_pass { + let mut pass = command_encoder.compute("ray-trace-spatial"); + let mut pc = pass.with(&self.spatial_pipeline); + let groups = self.spatial_pipeline.get_dispatch_for(self.surface_size); + pc.bind( + 0, + &MainData { + camera: self.targets.camera_params[cur], + prev_camera: self.targets.camera_params[cur], + debug, + parameters, + acc_struct: self.acceleration_structure, + prev_acc_struct: self.acceleration_structure, + sampler_linear: self.samplers.linear, + sampler_nearest: self.samplers.nearest, + env_map: self.env_map.main_view, + env_weights: self.env_map.weight_view, + t_depth: self.targets.depth.views[cur], + t_prev_depth: self.targets.depth.views[cur], + t_basis: self.targets.basis.views[cur], + t_prev_basis: self.targets.basis.views[cur], + t_flat_normal: self.targets.flat_normal.views[cur], + t_prev_flat_normal: self.targets.flat_normal.views[cur], + t_motion: self.targets.motion.views[0], + debug_buf: self.debug.buffer_resource(), + reservoirs: self.targets.reservoir_buf[0].into(), + prev_reservoirs: self.targets.reservoir_buf[1].into(), + out_diffuse: self.targets.light_diffuse.views[0], + in_diffuse: self.targets.light_diffuse.views[1], out_debug: self.targets.debug.views[0], }, ); pc.dispatch(groups); + } else { + self.targets.reservoir_buf.swap(0, 1); + self.targets.light_diffuse.views.swap(0, 1); } + //Note: output reservoirs and lighting is in [0] } /// Perform noise reduction using SVGF. #[profiling::function] pub fn denoise( - &mut self, //TODO: borrow immutably + &mut self, command_encoder: &mut blade_graphics::CommandEncoder, denoiser_config: DenoiserConfig, ) { let mut params = BlurParams { extent: [self.surface_size.width, self.surface_size.height], - temporal_weight: denoiser_config.temporal_weight, + temporal_weight: if self.frame_resize < self.frame_index { + denoiser_config.temporal_weight + } else { + 1.0 + }, iteration: 0, use_motion_vectors: (self.frame_scene_built >= self.frame_index) as u32, pad: 0, @@ -1242,21 +1311,24 @@ impl Renderer { camera: self.targets.camera_params[cur], prev_camera: self.targets.camera_params[prev], params, - input: self.targets.light_diffuse.views[prev], + input: self.targets.light_diffuse.views[2], t_depth: self.targets.depth.views[cur], t_prev_depth: self.targets.depth.views[prev], t_flat_normal: self.targets.flat_normal.views[cur], t_prev_flat_normal: self.targets.flat_normal.views[prev], t_motion: self.targets.motion.views[0], - output: self.targets.light_diffuse.views[cur], + output: self.targets.light_diffuse.views[0], }, ); pc.dispatch(groups); } - assert_eq!(cur, self.post_proc_input_index); - let mut ping_pong = [2, if self.is_frozen { cur } else { prev }]; + // Make sure the accumulated result is in [2] + self.targets.light_diffuse.views.swap(0, 2); + + let mut input_index = 2; for _ in 0..denoiser_config.num_passes { + let output_index = if input_index == 0 { 1 } else { 0 }; let mut pass = command_encoder.compute("a-trous"); let mut pc = pass.with(&self.blur.a_trous_pipeline); let groups = self @@ -1267,17 +1339,20 @@ impl Renderer { 0, &ATrousData { params, - input: self.targets.light_diffuse.views[self.post_proc_input_index], + input: self.targets.light_diffuse.views[input_index], t_depth: self.targets.depth.views[cur], t_flat_normal: self.targets.flat_normal.views[cur], - output: self.targets.light_diffuse.views[ping_pong[0]], + output: self.targets.light_diffuse.views[output_index], }, ); pc.dispatch(groups); - self.post_proc_input_index = ping_pong[0]; - ping_pong.swap(0, 1); params.iteration += 1; + input_index = output_index; + } + if input_index != 0 { + self.targets.light_diffuse.views.swap(0, input_index); } + //Note: output lighting is in [0] } /// Blit the rendering result into a specified render pass. @@ -1297,7 +1372,7 @@ impl Renderer { 0, &PostProcData { t_albedo: self.targets.albedo.views[0], - light_diffuse: self.targets.light_diffuse.views[self.post_proc_input_index], + light_diffuse: self.targets.light_diffuse.views[0], t_debug: self.targets.debug.views[0], tone_map_params: ToneMapParams { enabled: 1,