From 42045524f2591ef1271a7dd83098a7f120af1f8f Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Mon, 20 Nov 2023 23:52:19 -0800 Subject: [PATCH] Only apply the motion vectors if the scene is fresh --- blade-render/code/blur.wgsl | 13 +++++++++++-- blade-render/code/fill-gbuf.wgsl | 1 + blade-render/code/gbuf.inc.wgsl | 1 + blade-render/code/ray-trace.wgsl | 13 +++++++++++-- blade-render/src/render/mod.rs | 9 +++++++++ examples/scene/main.rs | 1 - 6 files changed, 33 insertions(+), 5 deletions(-) diff --git a/blade-render/code/blur.wgsl b/blade-render/code/blur.wgsl index 0a830299..6c977e2d 100644 --- a/blade-render/code/blur.wgsl +++ b/blade-render/code/blur.wgsl @@ -12,6 +12,7 @@ struct Params { extent: vec2, temporal_weight: f32, iteration: u32, + use_motion_vectors: u32, } var camera: CameraParams; @@ -42,6 +43,15 @@ fn read_prev_surface(pixel: vec2) -> Surface { return surface; } +fn get_prev_pixel(pixel: vec2, pos_world: vec3) -> vec2 { + if (USE_MOTION_VECTORS && params.use_motion_vectors != 0u) { + let motion = textureLoad(t_motion, pixel, 0).xy / MOTION_SCALE; + return vec2(pixel) + 0.5 + motion; + } else { + return get_projected_pixel_float(prev_camera, pos_world); + } +} + @compute @workgroup_size(8, 8) fn temporal_accum(@builtin(global_invocation_id) global_id: vec3) { let pixel = vec2(global_id.xy); @@ -52,8 +62,7 @@ fn temporal_accum(@builtin(global_invocation_id) global_id: vec3) { let surface = read_surface(pixel); let pos_world = camera.position + surface.depth * get_ray_direction(camera, pixel); // considering all samples in 2x2 quad, to help with edges - let motion = textureLoad(t_motion, pixel, 0).xy / MOTION_SCALE; - let center_pixel = vec2(pixel) + 0.5 + motion; + var center_pixel = get_prev_pixel(pixel, pos_world); var prev_pixels = array, 4>( vec2(vec2(center_pixel.x - 0.5, center_pixel.y - 0.5)), vec2(vec2(center_pixel.x + 0.5, center_pixel.y - 0.5)), diff --git a/blade-render/code/fill-gbuf.wgsl b/blade-render/code/fill-gbuf.wgsl index 8ebe3f11..0f143bb7 100644 --- a/blade-render/code/fill-gbuf.wgsl +++ b/blade-render/code/fill-gbuf.wgsl @@ -181,6 +181,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3) { let prev_position = (entry.prev_object_to_world * position_object).xyz; let prev_screen = get_projected_pixel_float(prev_camera, prev_position); //TODO: consider just storing integers here? + //TODO: technically this "0.5" is just a waste compute on both packing and unpacking motion = prev_screen - vec2(global_id.xy) - 0.5; if (debug.view_mode == DebugMode_Motion) { textureStore(out_debug, global_id.xy, vec4(motion * MOTION_SCALE + vec2(0.5), 0.0, 1.0)); diff --git a/blade-render/code/gbuf.inc.wgsl b/blade-render/code/gbuf.inc.wgsl index 3e721716..ecb4642d 100644 --- a/blade-render/code/gbuf.inc.wgsl +++ b/blade-render/code/gbuf.inc.wgsl @@ -1 +1,2 @@ const MOTION_SCALE: f32 = 0.02; +const USE_MOTION_VECTORS: bool = true; \ No newline at end of file diff --git a/blade-render/code/ray-trace.wgsl b/blade-render/code/ray-trace.wgsl index 168e3771..47ebee45 100644 --- a/blade-render/code/ray-trace.wgsl +++ b/blade-render/code/ray-trace.wgsl @@ -31,6 +31,7 @@ struct MainParams { spatial_taps: u32, spatial_tap_history: u32, spatial_radius: i32, + use_motion_vectors: u32, }; var camera: CameraParams; @@ -252,6 +253,15 @@ fn evaluate_reflected_light(surface: Surface, light_index: u32, light_uv: vec2, pos_world: vec3) -> vec2 { + if (USE_MOTION_VECTORS && parameters.use_motion_vectors != 0u) { + let motion = textureLoad(t_motion, pixel, 0).xy / MOTION_SCALE; + return vec2(pixel) + 0.5 + motion; + } else { + return get_projected_pixel_float(prev_camera, pos_world); + } +} + struct TargetScore { color: vec3, score: f32, @@ -364,9 +374,8 @@ fn compute_restir(surface: Surface, pixel: vec2, rng: ptr(motion + vec2(0.5)); + let prev_pixel = vec2(get_prev_pixel(pixel, position)); // First, gather the list of reservoirs to merge with var accepted_reservoir_indices = array(); diff --git a/blade-render/src/render/mod.rs b/blade-render/src/render/mod.rs index e7cacfc4..83efda76 100644 --- a/blade-render/src/render/mod.rs +++ b/blade-render/src/render/mod.rs @@ -321,6 +321,7 @@ pub struct Renderer { screen_size: blade_graphics::Extent, screen_format: blade_graphics::TextureFormat, frame_index: usize, + frame_scene_built: usize, //TODO: refactor `ResourceArray` to not carry the freelist logic // This way we can embed user info into the allocator. texture_resource_lookup: @@ -357,6 +358,7 @@ struct MainParams { spatial_taps: u32, spatial_tap_history: u32, spatial_radius: u32, + use_motion_vectors: u32, } #[derive(blade_macros::ShaderData)] @@ -410,6 +412,8 @@ struct BlurParams { extent: [u32; 2], temporal_weight: f32, iteration: i32, + use_motion_vectors: u32, + pad: u32, } #[derive(blade_macros::ShaderData)] @@ -696,6 +700,7 @@ impl Renderer { screen_size: config.screen_size, screen_format: config.surface_format, frame_index: 0, + frame_scene_built: 0, texture_resource_lookup: HashMap::default(), } } @@ -991,6 +996,7 @@ impl Renderer { temp.buffers.push(instance_buf); temp.buffers.push(scratch_buf); + self.frame_scene_built = self.frame_index + 1; } fn make_debug_params(&self, config: &DebugConfig) -> DebugParams { @@ -1123,6 +1129,7 @@ impl Renderer { spatial_taps: ray_config.spatial_taps, spatial_tap_history: ray_config.spatial_tap_history, spatial_radius: ray_config.spatial_radius, + use_motion_vectors: (self.frame_scene_built == self.frame_index) as u32, }, acc_struct: self.acceleration_structure, sampler_linear: self.samplers.linear, @@ -1158,6 +1165,8 @@ impl Renderer { extent: [self.screen_size.width, self.screen_size.height], temporal_weight: denoiser_config.temporal_weight, iteration: 0, + use_motion_vectors: (self.frame_scene_built == self.frame_index) as u32, + pad: 0, }; let cur = self.frame_index % 2; let prev = cur ^ 1; diff --git a/examples/scene/main.rs b/examples/scene/main.rs index 11e739b3..68401777 100644 --- a/examples/scene/main.rs +++ b/examples/scene/main.rs @@ -548,7 +548,6 @@ impl Example { } .to_blade(); if object.transform != t1 { - object.prev_transform = object.transform; object.transform = t1; self.have_objects_changed = true; }