Group-local ReSTIR

Cargo.toml

@@ -22,6 +22,7 @@ gltf = { version = "1.1", default-features = false }
 log = "0.4"
 mint = "0.5"
 naga = { version = "22", features = ["wgsl-in"] }
+nanorand = { version = "0.7", default-features = false }
 profiling = "1"
 slab = "0.4"
 strum = { version = "0.25", features = ["derive"] }
@@ -79,7 +80,7 @@ glam = { workspace = true }
 log = { workspace = true }
 mint = { workspace = true, features = ["serde"] }
 naga = { workspace = true }
-nanorand = { version = "0.7", default-features = false, features = ["wyrand"] }
+nanorand = { workspace = true, features = ["wyrand"] }
 profiling = { workspace = true }
 ron = "0.8"
 serde = { version = "1", features = ["serde_derive"] }
@@ -95,7 +96,6 @@ egui-winit = "0.28"
 console_error_panic_hook = "0.1.7"
 console_log = "1"
 web-sys = { workspace = true, features = ["Window"] }
-getrandom = { version = "0.2", features = ["js"] }
 
 [target.'cfg(any(target_os = "windows", target_os = "linux"))'.dev-dependencies]
 renderdoc = "0.12"

blade-graphics/src/util.rs

@@ -94,14 +94,19 @@ impl super::TextureFormat {
     }
 }
 
+impl super::Extent {
+    pub fn group_by(&self, size: [u32; 3]) -> [u32; 3] {
+        [
+            (self.width + size[0] - 1) / size[0],
+            (self.height + size[1] - 1) / size[1],
+            (self.depth + size[2] - 1) / size[2],
+        ]
+    }
+}
+
 impl super::ComputePipeline {
     /// Return the dispatch group counts sufficient to cover the given extent.
     pub fn get_dispatch_for(&self, extent: super::Extent) -> [u32; 3] {
-        let wg_size = self.get_workgroup_size();
-        [
-            (extent.width + wg_size[0] - 1) / wg_size[0],
-            (extent.height + wg_size[1] - 1) / wg_size[1],
-            (extent.depth + wg_size[2] - 1) / wg_size[2],
-        ]
+        extent.group_by(self.get_workgroup_size())
     }
 }

blade-helpers/src/camera.rs

@@ -1,5 +1,7 @@
 use super::ExposeHud;
 
+const MAX_FLY_SPEED: f32 = 1000000.0;
+
 pub struct ControlledCamera {
     pub inner: blade_render::Camera,
     pub fly_speed: f32,
@@ -86,6 +88,14 @@ impl ControlledCamera {
 
         true
     }
+
+    pub fn on_wheel(&mut self, delta: winit::event::MouseScrollDelta) {
+        let shift = match delta {
+            winit::event::MouseScrollDelta::LineDelta(_, lines) => lines,
+            winit::event::MouseScrollDelta::PixelDelta(position) => position.y as f32,
+        };
+        self.fly_speed = (self.fly_speed * shift.exp()).clamp(1.0, MAX_FLY_SPEED);
+    }
 }
 
 impl ExposeHud for ControlledCamera {
@@ -105,7 +115,7 @@ impl ExposeHud for ControlledCamera {
         });
         ui.add(egui::Slider::new(&mut self.inner.fov_y, 0.5f32..=2.0f32).text("FOV"));
         ui.add(
-            egui::Slider::new(&mut self.fly_speed, 1f32..=100000f32)
+            egui::Slider::new(&mut self.fly_speed, 1f32..=MAX_FLY_SPEED)
                 .text("Fly speed")
                 .logarithmic(true),
         );

blade-helpers/src/hud.rs

@@ -15,27 +15,34 @@ impl ExposeHud for blade_render::RayConfig {
         );
         ui.checkbox(&mut self.temporal_tap, "Temporal tap");
         ui.add(
-            egui::widgets::Slider::new(&mut self.temporal_history, 0..=50).text("Temporal history"),
+            egui::widgets::Slider::new(&mut self.temporal_confidence, 0.0..=50.0)
+                .text("Temporal confidence"),
         );
         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)
-                .text("Spatial tap history"),
+            egui::widgets::Slider::new(&mut self.spatial_confidence, 0.0..=50.0)
+                .text("Spatial confidence"),
         );
+        ui.add(egui::widgets::Slider::new(&mut self.group_mixer, 1..=10).text("Group mixer"));
         ui.add(
-            egui::widgets::Slider::new(&mut self.spatial_radius, 1..=50)
-                .text("Spatial radius (px)"),
+            egui::widgets::Slider::new(&mut self.spatial_min_distance, 1..=10)
+                .text("Spatial minimum distance (px)"),
         );
         ui.add(
             egui::widgets::Slider::new(&mut self.t_start, 0.001..=0.5)
                 .text("T min")
                 .logarithmic(true),
         );
+        ui.checkbox(&mut self.pairwise_mis, "Pairwise MIS");
+        ui.add(
+            egui::widgets::Slider::new(&mut self.defensive_mis, 0.0..=1.0).text("Defensive MIS"),
+        );
     }
 }
 
 impl ExposeHud for blade_render::DenoiserConfig {
     fn populate_hud(&mut self, ui: &mut egui::Ui) {
+        ui.checkbox(&mut self.enabled, "Enable denoiser");
         ui.add(egui::Slider::new(&mut self.temporal_weight, 0.0..=1.0f32).text("Temporal weight"));
         ui.add(egui::Slider::new(&mut self.num_passes, 0..=5u32).text("A-trous passes"));
     }
@@ -57,6 +64,15 @@ impl ExposeHud for blade_render::PostProcConfig {
     }
 }
 
+impl ExposeHud for blade_render::FrameConfig {
+    fn populate_hud(&mut self, ui: &mut egui::Ui) {
+        ui.horizontal(|ui| {
+            self.reset_reservoirs |= ui.button("Reset Accumulation").clicked();
+            ui.toggle_value(&mut self.frozen, "Freeze");
+        });
+    }
+}
+
 impl ExposeHud for blade_render::DebugConfig {
     fn populate_hud(&mut self, ui: &mut egui::Ui) {
         use strum::IntoEnumIterator as _;

blade-render/code/a-trous.wgsl

@@ -0,0 +1,79 @@
+#include "quaternion.inc.wgsl"
+#include "surface.inc.wgsl"
+
+// Spatio-temporal variance-guided filtering
+// https://research.nvidia.com/sites/default/files/pubs/2017-07_Spatiotemporal-Variance-Guided-Filtering%3A//svgf_preprint.pdf
+
+// Note: using "ilm" in place of "illumination and the 2nd moment of its luminance"
+
+struct Params {
+    extent: vec2<i32>,
+    iteration: u32,
+}
+
+var<uniform> params: Params;
+var t_depth: texture_2d<f32>;
+var t_flat_normal: texture_2d<f32>;
+var input: texture_2d<f32>;
+var output: texture_storage_2d<rgba16float, write>;
+
+const LUMA: vec3<f32> = vec3<f32>(0.2126, 0.7152, 0.0722);
+const MIN_WEIGHT: f32 = 0.01;
+
+fn read_surface(pixel: vec2<i32>) -> Surface {
+    var surface = Surface();
+    surface.flat_normal = normalize(textureLoad(t_flat_normal, pixel, 0).xyz);
+    surface.depth = textureLoad(t_depth, pixel, 0).x;
+    return surface;
+}
+
+const GAUSSIAN_WEIGHTS = vec2<f32>(0.44198, 0.27901);
+const SIGMA_L: f32 = 4.0;
+const EPSILON: f32 = 0.001;
+
+fn compare_luminance(a_lum: f32, b_lum: f32, variance: f32) -> f32 {
+    return exp(-abs(a_lum - b_lum) / (SIGMA_L * variance + EPSILON));
+}
+
+fn w4(w: f32) -> vec4<f32> {
+    return vec4<f32>(vec3<f32>(w), w * w);
+}
+
+@compute @workgroup_size(8, 8)
+fn atrous3x3(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let center = vec2<i32>(global_id.xy);
+    if (any(center >= params.extent)) {
+        return;
+    }
+
+    let center_ilm = textureLoad(input, center, 0);
+    let center_luma = dot(center_ilm.xyz, LUMA);
+    let variance = sqrt(center_ilm.w);
+    let center_suf = read_surface(center);
+    var sum_weight = GAUSSIAN_WEIGHTS[0] * GAUSSIAN_WEIGHTS[0];
+    var sum_ilm = w4(sum_weight) * center_ilm;
+
+    for (var yy=-1; yy<=1; yy+=1) {
+        for (var xx=-1; xx<=1; xx+=1) {
+            let p = center + vec2<i32>(xx, yy) * (1i << params.iteration);
+            if (all(p == center) || any(p < vec2<i32>(0)) || any(p >= params.extent)) {
+                continue;
+            }
+
+            //TODO: store in group-shared memory
+            let surface = read_surface(p);
+            var weight = GAUSSIAN_WEIGHTS[abs(xx)] * GAUSSIAN_WEIGHTS[abs(yy)];
+            //TODO: make it stricter on higher iterations
+            weight *= compare_flat_normals(surface.flat_normal, center_suf.flat_normal);
+            //Note: should we use a projected depth instead of the surface one?
+            weight *= compare_depths(surface.depth, center_suf.depth);
+            let other_ilm = textureLoad(input, p, 0);
+            weight *= compare_luminance(center_luma, dot(other_ilm.xyz, LUMA), variance);
+            sum_ilm += w4(weight) * other_ilm;
+            sum_weight += weight;
+        }
+    }
+
+    let filtered_ilm = select(center_ilm, sum_ilm / w4(sum_weight), sum_weight > MIN_WEIGHT);
+    textureStore(output, global_id.xy, filtered_ilm);
+}

blade-render/code/accum.inc.wgsl

@@ -0,0 +1,21 @@
+const LUMA: vec3<f32> = vec3<f32>(0.2126, 0.7152, 0.0722);
+const MOTION_FACTOR: f32 = 0.1;
+
+var inout_diffuse: texture_storage_2d<rgba16float, read_write>;
+
+fn accumulate_temporal(
+    pixel: vec2<i32>, cur_illumination: vec3<f32>,
+    temporal_weight: f32, prev_pixel: vec2<i32>,
+    motion_sqr: f32,
+) {
+    var illumination = cur_illumination;
+    if (prev_pixel.x >= 0 && temporal_weight < 1.0) {
+        let factor = mix(temporal_weight, 1.0, min(pow(motion_sqr, 0.25) * MOTION_FACTOR, 1.0));
+        let prev_illumination = textureLoad(inout_diffuse, prev_pixel).xyz;
+        illumination = mix(prev_illumination, illumination, factor);
+    }
+
+    let luminocity = dot(illumination, LUMA);
+    let ilm = vec4<f32>(illumination, luminocity * luminocity);
+    textureStore(inout_diffuse, pixel, ilm);
+}