Support motion vectors

blade-graphics/Cargo.toml

@@ -35,8 +35,7 @@ naga = { workspace = true, features = ["spv-out"] }
 slab = { workspace = true }
 
 [target.'cfg(any(gles, target_arch = "wasm32"))'.dependencies]
-# Version contains `glGetProgramResource`
-glow = "0.12.2"
+glow = "0.13"
 naga = { workspace = true, features = ["glsl-out"] }
 
 [target.'cfg(all(gles, not(target_arch = "wasm32")))'.dependencies]

blade-graphics/src/gles/mod.rs

@@ -451,6 +451,8 @@ struct FormatInfo {
 fn describe_texture_format(format: crate::TextureFormat) -> FormatInfo {
     use crate::TextureFormat as Tf;
     let (internal, external, data_type) = match format {
+        Tf::Rg8Unorm => (glow::RG8, glow::RG, glow::UNSIGNED_BYTE),
+        Tf::Rg8Snorm => (glow::RG8, glow::RG, glow::BYTE),
         Tf::Rgba8Unorm => (glow::RGBA8, glow::RGBA, glow::UNSIGNED_BYTE),
         Tf::Rgba8UnormSrgb => (glow::SRGB8_ALPHA8, glow::RGBA, glow::UNSIGNED_BYTE),
         Tf::Bgra8UnormSrgb => (glow::SRGB8_ALPHA8, glow::BGRA, glow::UNSIGNED_BYTE),

blade-graphics/src/lib.rs

@@ -214,6 +214,8 @@ impl From<Texture> for TexturePiece {
 #[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
 pub enum TextureFormat {
     // color
+    Rg8Unorm,
+    Rg8Snorm,
     Rgba8Unorm,
     Rgba8UnormSrgb,
     Bgra8UnormSrgb,

blade-graphics/src/metal/mod.rs

@@ -278,6 +278,8 @@ fn map_texture_format(format: crate::TextureFormat) -> metal::MTLPixelFormat {
     use crate::TextureFormat as Tf;
     use metal::MTLPixelFormat::*;
     match format {
+        Tf::Rg8Unorm => RG8Unorm,
+        Tf::Rg8Snorm => RG8Snorm,
         Tf::Rgba8Unorm => RGBA8Unorm,
         Tf::Rgba8UnormSrgb => RGBA8Unorm_sRGB,
         Tf::Bgra8UnormSrgb => BGRA8Unorm_sRGB,

blade-graphics/src/util.rs

@@ -56,6 +56,8 @@ impl super::TextureFormat {
             }
         }
         match *self {
+            Self::Rg8Unorm => uncompressed(2),
+            Self::Rg8Snorm => uncompressed(2),
             Self::Rgba8Unorm => uncompressed(4),
             Self::Rgba8UnormSrgb => uncompressed(4),
             Self::Bgra8UnormSrgb => uncompressed(4),

blade-graphics/src/vulkan/mod.rs

@@ -474,6 +474,8 @@ impl crate::traits::CommandDevice for Context {
 fn map_texture_format(format: crate::TextureFormat) -> vk::Format {
     use crate::TextureFormat as Tf;
     match format {
+        Tf::Rg8Unorm => vk::Format::R8G8_UNORM,
+        Tf::Rg8Snorm => vk::Format::R8G8_SNORM,
         Tf::Rgba8Unorm => vk::Format::R8G8B8A8_UNORM,
         Tf::Rgba8UnormSrgb => vk::Format::R8G8B8A8_SRGB,
         Tf::Bgra8UnormSrgb => vk::Format::B8G8R8A8_SRGB,

blade-render/code/blur.wgsl

@@ -1,30 +1,34 @@
 #include "camera.inc.wgsl"
+#include "gbuf.inc.wgsl"
 #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 luminanc"
+// Note: using "ilm" in place of "illumination and the 2nd moment of its luminance"
 
 struct Params {
     extent: vec2<i32>,
     temporal_weight: f32,
     iteration: u32,
+    use_motion_vectors: u32,
 }
 
 var<uniform> camera: CameraParams;
 var<uniform> prev_camera: CameraParams;
 var<uniform> params: Params;
-var t_flat_normal: texture_2d<f32>;
-var t_prev_flat_normal: texture_2d<f32>;
 var t_depth: texture_2d<f32>;
 var t_prev_depth: texture_2d<f32>;
+var t_flat_normal: texture_2d<f32>;
+var t_prev_flat_normal: texture_2d<f32>;
+var t_motion: texture_2d<f32>;
 var input: texture_2d<f32>;
 var prev_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();
@@ -39,18 +43,26 @@ fn read_prev_surface(pixel: vec2<i32>) -> Surface {
     return surface;
 }
 
+fn get_prev_pixel(pixel: vec2<i32>, pos_world: vec3<f32>) -> vec2<f32> {
+    if (USE_MOTION_VECTORS && params.use_motion_vectors != 0u) {
+        let motion = textureLoad(t_motion, pixel, 0).xy / MOTION_SCALE;
+        return vec2<f32>(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<u32>) {
     let pixel = vec2<i32>(global_id.xy);
     if (any(pixel >= params.extent)) {
         return;
     }
 
-    //TODO: use motion vectors
     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 center_pixel = get_projected_pixel_float(prev_camera, pos_world);
+    var center_pixel = get_prev_pixel(pixel, pos_world);
     var prev_pixels = array<vec2<i32>, 4>(
         vec2<i32>(vec2<f32>(center_pixel.x - 0.5, center_pixel.y - 0.5)),
         vec2<i32>(vec2<f32>(center_pixel.x + 0.5, center_pixel.y - 0.5)),
@@ -91,7 +103,7 @@ fn temporal_accum(@builtin(global_invocation_id) global_id: vec3<u32>) {
     let cur_illumination = textureLoad(input, pixel, 0).xyz;
     let cur_luminocity = dot(cur_illumination, LUMA);
     var mixed_ilm = vec4<f32>(cur_illumination, cur_luminocity * cur_luminocity);
-    if (sum_weight > 0.01) {
+    if (sum_weight > MIN_WEIGHT) {
         let prev_ilm = sum_ilm / w4(sum_weight);
         mixed_ilm = mix(mixed_ilm, prev_ilm, sum_weight * (1.0 - params.temporal_weight));
     }
@@ -145,6 +157,6 @@ fn atrous3x3(@builtin(global_invocation_id) global_id: vec3<u32>) {
         }
     }
 
-    let filtered_ilm = sum_ilm / w4(sum_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/fill-gbuf.wgsl

@@ -2,6 +2,7 @@
 #include "camera.inc.wgsl"
 #include "debug.inc.wgsl"
 #include "debug-param.inc.wgsl"
+#include "gbuf.inc.wgsl"
 
 //TODO: use proper WGSL
 const RAY_FLAG_CULL_NO_OPAQUE: u32 = 0x80u;
@@ -33,6 +34,7 @@ struct HitEntry {
     geometry_to_world_rotation: u32,
     pad: u32,
     geometry_to_object: mat4x3<f32>,
+    prev_object_to_world: mat4x3<f32>,
     base_color_texture: u32,
     // packed color factor
     base_color_factor: u32,
@@ -41,13 +43,15 @@ struct HitEntry {
 var<storage, read> hit_entries: array<HitEntry>;
 
 var<uniform> camera: CameraParams;
+var<uniform> prev_camera: CameraParams;
 var<uniform> debug: DebugParams;
 var acc_struct: acceleration_structure;
 
 var out_depth: texture_storage_2d<r32float, write>;
 var out_flat_normal: texture_storage_2d<rgba8snorm, write>;
 var out_basis: texture_storage_2d<rgba8snorm, write>;
 var out_albedo: texture_storage_2d<rgba8unorm, write>;
+var out_motion: texture_storage_2d<rg8snorm, write>;
 var out_debug: texture_storage_2d<rgba8unorm, write>;
 
 fn decode_normal(raw: u32) -> vec3<f32> {
@@ -75,6 +79,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     var basis = vec4<f32>(0.0);
     var flat_normal = vec3<f32>(0.0);
     var albedo = vec3<f32>(1.0);
+    var motion = vec2<f32>(0.0);
     let enable_debug = all(global_id.xy == debug.mouse_pos);
 
     if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) {
@@ -94,15 +99,16 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
             (*vptr)[indices.z],
         );
 
-        let pos_object = entry.geometry_to_object * mat3x4(
+        let positions_object = entry.geometry_to_object * mat3x4(
             vec4<f32>(vertices[0].pos, 1.0), vec4<f32>(vertices[1].pos, 1.0), vec4<f32>(vertices[2].pos, 1.0)
         );
         let positions = intersection.object_to_world * mat3x4(
-            vec4<f32>(pos_object[0], 1.0), vec4<f32>(pos_object[1], 1.0), vec4<f32>(pos_object[2], 1.0)
+            vec4<f32>(positions_object[0], 1.0), vec4<f32>(positions_object[1], 1.0), vec4<f32>(positions_object[2], 1.0)
         );
         flat_normal = normalize(cross(positions[1].xyz - positions[0].xyz, positions[2].xyz - positions[0].xyz));
 
         let barycentrics = vec3<f32>(1.0 - intersection.barycentrics.x - intersection.barycentrics.y, intersection.barycentrics);
+        let position_object = vec4<f32>(positions_object * barycentrics, 1.0);
         let tex_coords = mat3x2(vertices[0].tex_coords, vertices[1].tex_coords, vertices[2].tex_coords) * barycentrics;
         let normal_geo = normalize(mat3x3(decode_normal(vertices[0].normal), decode_normal(vertices[1].normal), decode_normal(vertices[2].normal)) * barycentrics);
         let tangent_geo = normalize(mat3x3(decode_normal(vertices[0].tangent), decode_normal(vertices[1].tangent), decode_normal(vertices[2].tangent)) * barycentrics);
@@ -163,13 +169,23 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
             let base_color_sample = textureSampleLevel(textures[entry.base_color_texture], sampler_linear, tex_coords, lod);
             albedo = (base_color_factor * base_color_sample).xyz;
         }
+
         if (debug.view_mode == DebugMode_HitConsistency) {
             let reprojected = get_projected_pixel(camera, hit_position);
-            let barycentrics_pos_diff = positions * barycentrics - hit_position;
+            let barycentrics_pos_diff = (intersection.object_to_world * position_object).xyz - hit_position;
             let camera_projection_diff = vec2<f32>(global_id.xy) - vec2<f32>(reprojected);
             let consistency = vec4<f32>(length(barycentrics_pos_diff), length(camera_projection_diff), 0.0, 0.0);
             textureStore(out_debug, global_id.xy, consistency);
         }
+
+        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<f32>(global_id.xy) - 0.5;
+        if (debug.view_mode == DebugMode_Motion) {
+            textureStore(out_debug, global_id.xy, vec4<f32>(motion * MOTION_SCALE + vec2<f32>(0.5), 0.0, 1.0));
+        }
     } else {
         if (enable_debug) {
             debug_buf.entry = DebugEntry();
@@ -179,8 +195,10 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
         }
     }
 
+    // TODO: option to avoid writing data for the sky
     textureStore(out_depth, global_id.xy, vec4<f32>(depth, 0.0, 0.0, 0.0));
     textureStore(out_basis, global_id.xy, basis);
     textureStore(out_flat_normal, global_id.xy, vec4<f32>(flat_normal, 0.0));
     textureStore(out_albedo, global_id.xy, vec4<f32>(albedo, 0.0));
+    textureStore(out_motion, global_id.xy, vec4<f32>(motion * MOTION_SCALE, 0.0, 0.0));
 }

blade-render/code/gbuf.inc.wgsl

@@ -0,0 +1,2 @@
+const MOTION_SCALE: f32 = 0.02;
+const USE_MOTION_VECTORS: bool = true;

blade-render/code/ray-trace.wgsl

@@ -5,6 +5,7 @@
 #include "debug-param.inc.wgsl"
 #include "camera.inc.wgsl"
 #include "surface.inc.wgsl"
+#include "gbuf.inc.wgsl"
 
 //TODO: use proper WGSL
 const RAY_FLAG_CULL_NO_OPAQUE: u32 = 0x80u;
@@ -30,6 +31,7 @@ struct MainParams {
     spatial_taps: u32,
     spatial_tap_history: u32,
     spatial_radius: i32,
+    use_motion_vectors: u32,
 };
 
 var<uniform> camera: CameraParams;
@@ -141,6 +143,7 @@ var t_basis: texture_2d<f32>;
 var t_prev_basis: texture_2d<f32>;
 var t_flat_normal: texture_2d<f32>;
 var t_prev_flat_normal: texture_2d<f32>;
+var t_motion: texture_2d<f32>;
 var out_diffuse: texture_storage_2d<rgba16float, write>;
 var out_debug: texture_storage_2d<rgba8unorm, write>;
 
@@ -250,6 +253,15 @@ fn evaluate_reflected_light(surface: Surface, light_index: u32, light_uv: vec2<f
     return radiance * brdf;
 }
 
+fn get_prev_pixel(pixel: vec2<i32>, pos_world: vec3<f32>) -> vec2<f32> {
+    if (USE_MOTION_VECTORS && parameters.use_motion_vectors != 0u) {
+        let motion = textureLoad(t_motion, pixel, 0).xy / MOTION_SCALE;
+        return vec2<f32>(pixel) + 0.5 + motion;
+    } else {
+        return get_projected_pixel_float(prev_camera, pos_world);
+    }
+}
+
 struct TargetScore {
     color: vec3<f32>,
     score: f32,
@@ -362,7 +374,8 @@ fn compute_restir(surface: Surface, pixel: vec2<i32>, rng: ptr<function, RandomS
         }
     }
 
-    let prev_pixel = get_projected_pixel(prev_camera, position);
+    //TODO: find best match in a 2x2 grid
+    let prev_pixel = vec2<i32>(get_prev_pixel(pixel, position));
 
     // First, gather the list of reservoirs to merge with
     var accepted_reservoir_indices = array<i32, MAX_RESERVOIRS>();

blade-render/src/lib.rs

@@ -41,15 +41,17 @@ pub struct Camera {
 }
 
 pub struct Object {
-    pub transform: blade_graphics::Transform,
     pub model: blade_asset::Handle<Model>,
+    pub transform: blade_graphics::Transform,
+    pub prev_transform: blade_graphics::Transform,
 }
 
 impl From<blade_asset::Handle<Model>> for Object {
     fn from(model: blade_asset::Handle<Model>) -> Self {
         Self {
-            transform: blade_graphics::IDENTITY_TRANSFORM,
             model,
+            transform: blade_graphics::IDENTITY_TRANSFORM,
+            prev_transform: blade_graphics::IDENTITY_TRANSFORM,
         }
     }
 }