diff --git a/blade-render/code/blur.wgsl b/blade-render/code/blur.wgsl index 3207ef60..3aec665b 100644 --- a/blade-render/code/blur.wgsl +++ b/blade-render/code/blur.wgsl @@ -1,5 +1,5 @@ #include "camera.inc.wgsl" -#include "gbuf.inc.wgsl" +#include "motion.inc.wgsl" #include "quaternion.inc.wgsl" #include "surface.inc.wgsl" diff --git a/blade-render/code/fill-gbuf.wgsl b/blade-render/code/fill-gbuf.wgsl deleted file mode 100644 index 346edf51..00000000 --- a/blade-render/code/fill-gbuf.wgsl +++ /dev/null @@ -1,204 +0,0 @@ -#include "quaternion.inc.wgsl" -#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; - -// Has to match the host! -struct Vertex { - pos: vec3, - bitangent_sign: f32, - tex_coords: vec2, - normal: u32, - tangent: u32, -} -struct VertexBuffer { - data: array, -} -struct IndexBuffer { - data: array, -} -var vertex_buffers: binding_array; -var index_buffers: binding_array; -var textures: binding_array>; -var sampler_linear: sampler; -var sampler_nearest: sampler; - -struct HitEntry { - index_buf: u32, - vertex_buf: u32, - winding: f32, - // packed quaternion - geometry_to_world_rotation: u32, - geometry_to_object: mat4x3, - prev_object_to_world: mat4x3, - base_color_texture: u32, - // packed color factor - base_color_factor: u32, - normal_texture: u32, -} -var hit_entries: array; - -var camera: CameraParams; -var prev_camera: CameraParams; -var debug: DebugParams; -var acc_struct: acceleration_structure; - -var out_depth: texture_storage_2d; -var out_flat_normal: texture_storage_2d; -var out_basis: texture_storage_2d; -var out_albedo: texture_storage_2d; -var out_motion: texture_storage_2d; -var out_debug: texture_storage_2d; - -fn decode_normal(raw: u32) -> vec3 { - return unpack4x8snorm(raw).xyz; -} - -fn debug_raw_normal(pos: vec3, normal_raw: u32, rotation: vec4, debug_len: f32, color: u32) { - let nw = normalize(qrot(rotation, decode_normal(normal_raw))); - debug_line(pos, pos + debug_len * nw, color); -} - -@compute @workgroup_size(8, 4) -fn main(@builtin(global_invocation_id) global_id: vec3) { - if (any(global_id.xy >= camera.target_size)) { - return; - } - - var rq: ray_query; - let ray_dir = get_ray_direction(camera, vec2(global_id.xy)); - rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_CULL_NO_OPAQUE, 0xFFu, 0.0, camera.depth, camera.position, ray_dir)); - rayQueryProceed(&rq); - let intersection = rayQueryGetCommittedIntersection(&rq); - - var depth = 0.0; - var basis = vec4(0.0); - var flat_normal = vec3(0.0); - var albedo = vec3(1.0); - var motion = vec2(0.0); - let enable_debug = all(global_id.xy == debug.mouse_pos); - - if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) { - let entry = hit_entries[intersection.instance_custom_index + intersection.geometry_index]; - depth = intersection.t; - - var indices = intersection.primitive_index * 3u + vec3(0u, 1u, 2u); - if (entry.index_buf != ~0u) { - let iptr = &index_buffers[entry.index_buf].data; - indices = vec3((*iptr)[indices.x], (*iptr)[indices.y], (*iptr)[indices.z]); - } - - let vptr = &vertex_buffers[entry.vertex_buf].data; - let vertices = array( - (*vptr)[indices.x], - (*vptr)[indices.y], - (*vptr)[indices.z], - ); - - let positions_object = entry.geometry_to_object * mat3x4( - vec4(vertices[0].pos, 1.0), vec4(vertices[1].pos, 1.0), vec4(vertices[2].pos, 1.0) - ); - let positions = intersection.object_to_world * mat3x4( - vec4(positions_object[0], 1.0), vec4(positions_object[1], 1.0), vec4(positions_object[2], 1.0) - ); - flat_normal = entry.winding * normalize(cross(positions[1].xyz - positions[0].xyz, positions[2].xyz - positions[0].xyz)); - - let barycentrics = vec3(1.0 - intersection.barycentrics.x - intersection.barycentrics.y, intersection.barycentrics); - let position_object = vec4(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); - let bitangent_geo = normalize(cross(normal_geo, tangent_geo)) * vertices[0].bitangent_sign; - - let lod = 0.0; //TODO: this is actually complicated - - let geo_to_world_rot = normalize(unpack4x8snorm(entry.geometry_to_world_rotation)); - let tangent_space_geo = mat3x3(tangent_geo, bitangent_geo, normal_geo); - var normal_local: vec3; - if ((debug.texture_flags & DebugTextureFlags_NORMAL) != 0u) { - normal_local = vec3(0.0, 0.0, 1.0); // ignore normal map - } else { - let n_xy = textureSampleLevel(textures[entry.normal_texture], sampler_linear, tex_coords, lod).xy; - normal_local = vec3(n_xy, sqrt(max(0.0, 1.0 - dot(n_xy.xy, n_xy.xy)))); - } - var normal = qrot(geo_to_world_rot, tangent_space_geo * normal_local); - basis = shortest_arc_quat(vec3(0.0, 0.0, 1.0), normalize(normal)); - - let hit_position = camera.position + intersection.t * ray_dir; - if (enable_debug) { - debug_buf.entry.custom_index = intersection.instance_custom_index; - debug_buf.entry.depth = intersection.t; - debug_buf.entry.tex_coords = tex_coords; - debug_buf.entry.base_color_texture = entry.base_color_texture; - debug_buf.entry.normal_texture = entry.normal_texture; - debug_buf.entry.position = hit_position; - debug_buf.entry.flat_normal = flat_normal; - } - if (enable_debug && (debug.draw_flags & DebugDrawFlags_SPACE) != 0u) { - let normal_len = 0.15 * intersection.t; - let side = 0.05 * intersection.t; - debug_line(hit_position, hit_position + normal_len * qrot(geo_to_world_rot, normal_geo), 0xFFFFFFu); - debug_line(hit_position - side * tangent_geo, hit_position + side * tangent_geo, 0x808080u); - debug_line(hit_position - side * bitangent_geo, hit_position + side * bitangent_geo, 0x808080u); - } - if (enable_debug && (debug.draw_flags & DebugDrawFlags_GEOMETRY) != 0u) { - let debug_len = intersection.t * 0.2; - debug_line(positions[0].xyz, positions[1].xyz, 0x00FFFFu); - debug_line(positions[1].xyz, positions[2].xyz, 0x00FFFFu); - debug_line(positions[2].xyz, positions[0].xyz, 0x00FFFFu); - let poly_center = (positions[0].xyz + positions[1].xyz + positions[2].xyz) / 3.0; - debug_line(poly_center, poly_center + 0.2 * debug_len * flat_normal, 0xFF00FFu); - // note: dynamic indexing into positions isn't allowed by WGSL yet - debug_raw_normal(positions[0].xyz, vertices[0].normal, geo_to_world_rot, 0.5*debug_len, 0xFFFF00u); - debug_raw_normal(positions[1].xyz, vertices[1].normal, geo_to_world_rot, 0.5*debug_len, 0xFFFF00u); - debug_raw_normal(positions[2].xyz, vertices[2].normal, geo_to_world_rot, 0.5*debug_len, 0xFFFF00u); - // draw tangent space - debug_line(hit_position, hit_position + debug_len * qrot(basis, vec3(1.0, 0.0, 0.0)), 0x0000FFu); - debug_line(hit_position, hit_position + debug_len * qrot(basis, vec3(0.0, 1.0, 0.0)), 0x00FF00u); - debug_line(hit_position, hit_position + debug_len * qrot(basis, vec3(0.0, 0.0, 1.0)), 0xFF0000u); - } - - let base_color_factor = unpack4x8unorm(entry.base_color_factor); - if ((debug.texture_flags & DebugTextureFlags_ALBEDO) != 0u) { - albedo = base_color_factor.xyz; - } else { - 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 = (intersection.object_to_world * position_object).xyz - hit_position; - let camera_projection_diff = vec2(global_id.xy) - vec2(reprojected); - let consistency = vec4(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(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)); - } - } else { - if (enable_debug) { - debug_buf.entry = DebugEntry(); - } - if (debug.view_mode != DebugMode_Final) { - textureStore(out_debug, global_id.xy, vec4(0.0)); - } - } - - // TODO: option to avoid writing data for the sky - textureStore(out_depth, global_id.xy, vec4(depth, 0.0, 0.0, 0.0)); - textureStore(out_basis, global_id.xy, basis); - textureStore(out_flat_normal, global_id.xy, vec4(flat_normal, 0.0)); - textureStore(out_albedo, global_id.xy, vec4(albedo, 0.0)); - textureStore(out_motion, global_id.xy, vec4(motion * MOTION_SCALE, 0.0, 0.0)); -} diff --git a/blade-render/code/gbuf.inc.wgsl b/blade-render/code/gbuf.inc.wgsl deleted file mode 100644 index ecb4642d..00000000 --- a/blade-render/code/gbuf.inc.wgsl +++ /dev/null @@ -1,2 +0,0 @@ -const MOTION_SCALE: f32 = 0.02; -const USE_MOTION_VECTORS: bool = true; \ No newline at end of file diff --git a/blade-render/code/geometry.inc.wgsl b/blade-render/code/geometry.inc.wgsl new file mode 100644 index 00000000..99b771b1 --- /dev/null +++ b/blade-render/code/geometry.inc.wgsl @@ -0,0 +1,182 @@ +//TODO: https://github.com/gfx-rs/wgpu/pull/5429 +const RAY_FLAG_CULL_NO_OPAQUE: u32 = 0x80u; + +// Has to match the host! +struct Vertex { + pos: vec3, + bitangent_sign: f32, + tex_coords: vec2, + normal: u32, + tangent: u32, +} +struct VertexBuffer { + data: array, +} +struct IndexBuffer { + data: array, +} +var vertex_buffers: binding_array; +var index_buffers: binding_array; +var textures: binding_array>; + +struct HitEntry { + index_buf: u32, + vertex_buf: u32, + winding: f32, + // packed quaternion + geometry_to_world_rotation: u32, + geometry_to_object: mat4x3, + prev_object_to_world: mat4x3, + base_color_texture: u32, + // packed color factor + base_color_factor: u32, + normal_texture: u32, +} +var hit_entries: array; + +fn decode_normal(raw: u32) -> vec3 { + return unpack4x8snorm(raw).xyz; +} + +fn debug_raw_normal(pos: vec3, normal_raw: u32, rotation: vec4, debug_len: f32, color: u32) { + let nw = normalize(qrot(rotation, decode_normal(normal_raw))); + debug_line(pos, pos + debug_len * nw, color); +} + +struct RichSurface { + inner: Surface, + position: vec3, + albedo: vec3, + motion: vec2, +} + +fn fetch_geometry(pixel_coord: vec2, enable_debug: bool, is_primary: bool) -> RichSurface { + var rq: ray_query; + let ray_dir = get_ray_direction(camera, pixel_coord); + rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_CULL_NO_OPAQUE, 0xFFu, 0.0, camera.depth, camera.position, ray_dir)); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + + var rs = RichSurface(); + rs.albedo = vec3(1.0); + + if (intersection.kind == RAY_QUERY_INTERSECTION_NONE) { + if (enable_debug) { + debug_buf.entry = DebugEntry(); + } + return rs; + } + + let entry = hit_entries[intersection.instance_custom_index + intersection.geometry_index]; + + var indices = intersection.primitive_index * 3u + vec3(0u, 1u, 2u); + if (entry.index_buf != ~0u) { + let iptr = &index_buffers[entry.index_buf].data; + indices = vec3((*iptr)[indices.x], (*iptr)[indices.y], (*iptr)[indices.z]); + } + + let vptr = &vertex_buffers[entry.vertex_buf].data; + let vertices = array( + (*vptr)[indices.x], + (*vptr)[indices.y], + (*vptr)[indices.z], + ); + + let positions_object = entry.geometry_to_object * mat3x4( + vec4(vertices[0].pos, 1.0), vec4(vertices[1].pos, 1.0), vec4(vertices[2].pos, 1.0) + ); + let positions = intersection.object_to_world * mat3x4( + vec4(positions_object[0], 1.0), vec4(positions_object[1], 1.0), vec4(positions_object[2], 1.0) + ); + let flat_normal = entry.winding * normalize(cross(positions[1].xyz - positions[0].xyz, positions[2].xyz - positions[0].xyz)); + + let barycentrics = vec3(1.0 - intersection.barycentrics.x - intersection.barycentrics.y, intersection.barycentrics); + let position_object = vec4(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); + let bitangent_geo = normalize(cross(normal_geo, tangent_geo)) * vertices[0].bitangent_sign; + + let lod = 0.0; //TODO: this is actually complicated + + let geo_to_world_rot = normalize(unpack4x8snorm(entry.geometry_to_world_rotation)); + let tangent_space_geo = mat3x3(tangent_geo, bitangent_geo, normal_geo); + var normal_local: vec3; + if ((debug.texture_flags & DebugTextureFlags_NORMAL) != 0u) { + normal_local = vec3(0.0, 0.0, 1.0); // ignore normal map + } else { + let n_xy = textureSampleLevel(textures[entry.normal_texture], sampler_linear, tex_coords, lod).xy; + normal_local = vec3(n_xy, sqrt(max(0.0, 1.0 - dot(n_xy.xy, n_xy.xy)))); + } + let normal = qrot(geo_to_world_rot, tangent_space_geo * normal_local); + let basis = shortest_arc_quat(vec3(0.0, 0.0, 1.0), normalize(normal)); + + let hit_position = camera.position + intersection.t * ray_dir; + if (enable_debug && is_primary) { + debug_buf.entry.custom_index = intersection.instance_custom_index; + debug_buf.entry.depth = intersection.t; + debug_buf.entry.tex_coords = tex_coords; + debug_buf.entry.base_color_texture = entry.base_color_texture; + debug_buf.entry.normal_texture = entry.normal_texture; + debug_buf.entry.position = hit_position; + debug_buf.entry.flat_normal = flat_normal; + } + if (enable_debug && (debug.draw_flags & DebugDrawFlags_SPACE) != 0u) { + let normal_len = 0.15 * intersection.t; + let side = 0.05 * intersection.t; + debug_line(hit_position, hit_position + normal_len * qrot(geo_to_world_rot, normal_geo), 0xFFFFFFu); + debug_line(hit_position - side * tangent_geo, hit_position + side * tangent_geo, 0x808080u); + debug_line(hit_position - side * bitangent_geo, hit_position + side * bitangent_geo, 0x808080u); + } + if (enable_debug && (debug.draw_flags & DebugDrawFlags_GEOMETRY) != 0u) { + let debug_len = intersection.t * 0.2; + debug_line(positions[0].xyz, positions[1].xyz, 0x00FFFFu); + debug_line(positions[1].xyz, positions[2].xyz, 0x00FFFFu); + debug_line(positions[2].xyz, positions[0].xyz, 0x00FFFFu); + let poly_center = (positions[0].xyz + positions[1].xyz + positions[2].xyz) / 3.0; + debug_line(poly_center, poly_center + 0.2 * debug_len * flat_normal, 0xFF00FFu); + // note: dynamic indexing into positions isn't allowed by WGSL yet + debug_raw_normal(positions[0].xyz, vertices[0].normal, geo_to_world_rot, 0.5*debug_len, 0xFFFF00u); + debug_raw_normal(positions[1].xyz, vertices[1].normal, geo_to_world_rot, 0.5*debug_len, 0xFFFF00u); + debug_raw_normal(positions[2].xyz, vertices[2].normal, geo_to_world_rot, 0.5*debug_len, 0xFFFF00u); + // draw tangent space + debug_line(hit_position, hit_position + debug_len * qrot(basis, vec3(1.0, 0.0, 0.0)), 0x0000FFu); + debug_line(hit_position, hit_position + debug_len * qrot(basis, vec3(0.0, 1.0, 0.0)), 0x00FF00u); + debug_line(hit_position, hit_position + debug_len * qrot(basis, vec3(0.0, 0.0, 1.0)), 0xFF0000u); + } + + rs.albedo = unpack4x8unorm(entry.base_color_factor).xyz; + if ((debug.texture_flags & DebugTextureFlags_ALBEDO) == 0u) { + let base_color_sample = textureSampleLevel(textures[entry.base_color_texture], sampler_linear, tex_coords, lod); + rs.albedo *= base_color_sample.xyz; + } + + if (is_primary) { + if (debug.view_mode == DebugMode_Depth) { + textureStore(out_debug, pixel_coord, vec4(intersection.t / camera.depth)); + } + if (debug.view_mode == DebugMode_Normal) { + textureStore(out_debug, pixel_coord, vec4(normal, 0.0)); + } + if (debug.view_mode == DebugMode_HitConsistency) { + let reprojected = get_projected_pixel(camera, hit_position); + let barycentrics_pos_diff = (intersection.object_to_world * position_object).xyz - hit_position; + let camera_projection_diff = vec2(pixel_coord - reprojected); + let consistency = vec4(length(barycentrics_pos_diff), length(camera_projection_diff), 0.0, 0.0); + textureStore(out_debug, pixel_coord, 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 + rs.motion = prev_screen - vec2(pixel_coord) - 0.5; + rs.position = hit_position; + + // Write down the Surface + rs.inner.basis = basis; + rs.inner.flat_normal = flat_normal; + rs.inner.depth = intersection.t; + return rs; +} diff --git a/blade-render/code/motion.inc.wgsl b/blade-render/code/motion.inc.wgsl new file mode 100644 index 00000000..a9a9f48a --- /dev/null +++ b/blade-render/code/motion.inc.wgsl @@ -0,0 +1,2 @@ +const MOTION_SCALE: f32 = 0.02; +const USE_MOTION_VECTORS: bool = true; diff --git a/blade-render/code/ray-trace.wgsl b/blade-render/code/ray-trace.wgsl index 776ec338..2abf2147 100644 --- a/blade-render/code/ray-trace.wgsl +++ b/blade-render/code/ray-trace.wgsl @@ -6,10 +6,8 @@ #include "debug-param.inc.wgsl" #include "camera.inc.wgsl" #include "surface.inc.wgsl" -#include "gbuf.inc.wgsl" - -//TODO: https://github.com/gfx-rs/wgpu/pull/5429 -const RAY_FLAG_CULL_NO_OPAQUE: u32 = 0x80u; +#include "geometry.inc.wgsl" +#include "motion.inc.wgsl" const PI: f32 = 3.1415926; const MAX_RESAMPLE: u32 = 4u; @@ -23,7 +21,6 @@ const DECOUPLED_SHADING: bool = false; //TODO: crashes on AMD 6850U if `GROUP_SIZE_TOTAL` > 32 const GROUP_SIZE: vec2 = vec2(8, 4); const GROUP_SIZE_TOTAL: u32 = GROUP_SIZE.x * GROUP_SIZE.y; -const GROUP_VISUALIZE: bool = false; struct MainParams { frame_index: u32, @@ -106,6 +103,7 @@ fn make_reservoir(ls: LightSample, light_index: u32, brdf: vec3) -> LiveRes r.history = 1.0; return r; } + fn merge_reservoir(r: ptr, other: LiveReservoir, random: f32) -> bool { (*r).weight_sum += other.weight_sum; (*r).history += other.history; @@ -144,14 +142,9 @@ fn pack_reservoir(r: LiveReservoir) -> StoredReservoir { return pack_reservoir_detail(r, r.history); } -var t_depth: texture_2d; var t_prev_depth: texture_2d; -var t_basis: texture_2d; -var t_prev_basis: texture_2d -; -var t_flat_normal: texture_2d; +var t_prev_basis: texture_2d; var t_prev_flat_normal: texture_2d; -var t_motion: texture_2d; var out_diffuse: texture_storage_2d; var out_debug: texture_storage_2d; @@ -206,14 +199,6 @@ fn sample_light_from_environment(rng: ptr) -> LightSample return ls; } -fn read_surface(pixel: vec2) -> Surface { - var surface: Surface; - surface.basis = normalize(textureLoad(t_basis, pixel, 0)); - surface.flat_normal = normalize(textureLoad(t_flat_normal, pixel, 0).xyz); - surface.depth = textureLoad(t_depth, pixel, 0).x; - return surface; -} - fn read_prev_surface(pixel: vec2) -> Surface { var surface: Surface; surface.basis = normalize(textureLoad(t_prev_basis, pixel, 0)); @@ -268,9 +253,8 @@ fn evaluate_reflected_light(surface: Surface, light_index: u32, light_uv: vec2, pos_world: vec3) -> vec2 { +fn get_prev_pixel(pixel: vec2, pos_world: vec3, motion: vec2) -> 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); @@ -473,7 +457,7 @@ fn finalize_resampling( } fn resample_temporal( - surface: Surface, cur_pixel: vec2, position: vec3, + surface: Surface, motion: vec2, cur_pixel: vec2, position: vec3, rng: ptr, debug_len: f32, ) -> ResampleOutput { if (debug.view_mode == DebugMode_TemporalMatch || debug.view_mode == DebugMode_TemporalMisCanonical || debug.view_mode == DebugMode_TemporalMisError) { @@ -485,7 +469,7 @@ fn resample_temporal( let canonical = produce_canonical(surface, position, rng, debug_len); //TODO: find best match in a 2x2 grid - let prev_pixel = vec2(get_prev_pixel(cur_pixel, position)); + let prev_pixel = vec2(get_prev_pixel(cur_pixel, position, motion)); let prev_reservoir_index = get_reservoir_index(prev_pixel, prev_camera); if (parameters.temporal_tap == 0u || prev_reservoir_index < 0) { @@ -585,64 +569,71 @@ fn resample_spatial( } fn compute_restir( + rs: RichSurface, pixel: vec2, local_index: u32, group_id: vec3, rng: ptr, enable_debug: bool, ) -> vec3 { - let surface = read_surface(pixel); - if (debug.view_mode == DebugMode_Depth) { - textureStore(out_debug, pixel, vec4(surface.depth / camera.depth)); - } - if (debug.view_mode == DebugMode_Normal) { - let normal = qrot(surface.basis, vec3(0.0, 0.0, 1.0)); - textureStore(out_debug, pixel, vec4(normal, 0.0)); - } - let debug_len = select(0.0, surface.depth * 0.2, enable_debug); - let ray_dir = get_ray_direction(camera, pixel); - let pixel_index = get_reservoir_index(pixel, camera); - let position = camera.position + surface.depth * ray_dir; + let debug_len = select(0.0, rs.inner.depth * 0.2, enable_debug); - let temporal = resample_temporal(surface, pixel, position, rng, debug_len); - pixel_cache[local_index] = PixelCache(surface, temporal.reservoir, position); + let temporal = resample_temporal(rs.inner, rs.motion, pixel, rs.position, rng, debug_len); + pixel_cache[local_index] = PixelCache(rs.inner, temporal.reservoir, rs.position); // sync with the workgroup to ensure all reservoirs are available. workgroupBarrier(); let temporal_live = revive_canonical(temporal); - let spatial = resample_spatial(surface, pixel, position, group_id, temporal_live, rng, debug_len); + let spatial = resample_spatial(rs.inner, pixel, rs.position, group_id, temporal_live, rng, debug_len); + + let pixel_index = get_reservoir_index(pixel, camera); reservoirs[pixel_index] = spatial.reservoir; return spatial.color; } +var out_depth: texture_storage_2d; +var out_basis: texture_storage_2d; +var out_flat_normal: texture_storage_2d; +var out_albedo: texture_storage_2d; +var out_motion: texture_storage_2d; + @compute @workgroup_size(GROUP_SIZE.x, GROUP_SIZE.y) fn main( @builtin(workgroup_id) group_id: vec3, @builtin(local_invocation_index) local_index: u32, ) { - pixel_cache[local_index].reservoir.confidence = 0.0; + pixel_cache[local_index] = PixelCache(); let pixel_coord = thread_index_to_coord(local_index, group_id); if (any(vec2(pixel_coord) >= camera.target_size)) { return; } - if (GROUP_VISUALIZE) - { + + if (debug.view_mode == DebugMode_Grouping) { var rng = random_init(group_id.y * 1000u + group_id.x, 0u); let h = random_gen(&rng) * 360.0; let color = hsv_to_rgb(h, 0.5, 1.0); - textureStore(out_diffuse, pixel_coord, vec4(color, 1.0)); + textureStore(out_debug, pixel_coord, vec4(color, 1.0)); return; } + let enable_debug = all(pixel_coord == vec2(debug.mouse_pos)); + let rs = fetch_geometry(pixel_coord, true, enable_debug); + + // TODO: option to avoid writing data for the sky + textureStore(out_depth, pixel_coord, vec4(rs.inner.depth, 0.0, 0.0, 0.0)); + textureStore(out_basis, pixel_coord, rs.inner.basis); + textureStore(out_flat_normal, pixel_coord, vec4(rs.inner.flat_normal, 0.0)); + textureStore(out_albedo, pixel_coord, vec4(rs.albedo, 0.0)); + textureStore(out_motion, pixel_coord, vec4(rs.motion * MOTION_SCALE, 0.0, 0.0)); + let global_index = u32(pixel_coord.y) * camera.target_size.x + u32(pixel_coord.x); var rng = random_init(global_index, parameters.frame_index); - let enable_debug = all(pixel_coord == vec2(debug.mouse_pos)); let enable_restir_debug = (debug.draw_flags & DebugDrawFlags_RESTIR) != 0u && enable_debug; - let color = compute_restir(pixel_coord, local_index, group_id, &rng, enable_restir_debug); + let color = compute_restir(rs, pixel_coord, local_index, group_id, &rng, enable_restir_debug); + textureStore(out_diffuse, pixel_coord, vec4(color, 1.0)); if (enable_debug) { debug_buf.variance.color_sum += color; debug_buf.variance.color2_sum += color * color; debug_buf.variance.count += 1u; } - textureStore(out_diffuse, pixel_coord, vec4(color, 1.0)); } diff --git a/blade-render/src/render/mod.rs b/blade-render/src/render/mod.rs index 74e3dd9c..e0e724c5 100644 --- a/blade-render/src/render/mod.rs +++ b/blade-render/src/render/mod.rs @@ -51,12 +51,13 @@ pub enum DebugMode { Normal = 2, Motion = 3, HitConsistency = 4, - TemporalMatch = 5, - TemporalMisCanonical = 6, - TemporalMisError = 7, - SpatialMatch = 8, - SpatialMisCanonical = 9, - SpatialMisError = 10, + Grouping = 5, + TemporalMatch = 10, + TemporalMisCanonical = 11, + TemporalMisError = 12, + SpatialMatch = 13, + SpatialMisCanonical = 14, + SpatialMisError = 15, Variance = 100, } @@ -313,7 +314,6 @@ pub struct Renderer { shaders: Shaders, targets: RestirTargets, post_proc_input_index: usize, - fill_pipeline: blade_graphics::ComputePipeline, main_pipeline: blade_graphics::ComputePipeline, post_proc_pipeline: blade_graphics::RenderPipeline, blur: Blur, @@ -375,46 +375,31 @@ struct MainParams { } #[derive(blade_macros::ShaderData)] -struct FillData<'a> { +struct MainData<'a> { camera: CameraParams, prev_camera: CameraParams, debug: DebugParams, + parameters: MainParams, acc_struct: blade_graphics::AccelerationStructure, + prev_acc_struct: blade_graphics::AccelerationStructure, hit_entries: blade_graphics::BufferPiece, index_buffers: &'a blade_graphics::BufferArray, vertex_buffers: &'a blade_graphics::BufferArray, textures: &'a blade_graphics::TextureArray, sampler_linear: blade_graphics::Sampler, - debug_buf: blade_graphics::BufferPiece, - out_depth: blade_graphics::TextureView, - out_basis: blade_graphics::TextureView, - out_flat_normal: blade_graphics::TextureView, - out_albedo: blade_graphics::TextureView, - out_motion: blade_graphics::TextureView, - out_debug: blade_graphics::TextureView, -} - -#[derive(blade_macros::ShaderData)] -struct MainData { - camera: CameraParams, - prev_camera: CameraParams, - debug: DebugParams, - parameters: MainParams, - acc_struct: blade_graphics::AccelerationStructure, - prev_acc_struct: blade_graphics::AccelerationStructure, - sampler_linear: blade_graphics::Sampler, sampler_nearest: blade_graphics::Sampler, env_map: blade_graphics::TextureView, env_weights: blade_graphics::TextureView, - t_depth: blade_graphics::TextureView, t_prev_depth: blade_graphics::TextureView, - t_basis: blade_graphics::TextureView, t_prev_basis: blade_graphics::TextureView, - t_flat_normal: blade_graphics::TextureView, t_prev_flat_normal: blade_graphics::TextureView, - t_motion: blade_graphics::TextureView, debug_buf: blade_graphics::BufferPiece, reservoirs: blade_graphics::BufferPiece, + out_depth: blade_graphics::TextureView, + out_basis: blade_graphics::TextureView, + out_flat_normal: blade_graphics::TextureView, + out_albedo: blade_graphics::TextureView, + out_motion: blade_graphics::TextureView, out_diffuse: blade_graphics::TextureView, out_debug: blade_graphics::TextureView, } @@ -492,7 +477,6 @@ struct HitEntry { #[derive(Clone, PartialEq)] pub struct Shaders { env_prepare: blade_asset::Handle, - fill_gbuf: blade_asset::Handle, ray_trace: blade_asset::Handle, blur: blade_asset::Handle, post_proc: blade_asset::Handle, @@ -505,7 +489,6 @@ impl Shaders { let mut ctx = asset_hub.open_context(path, "shader finish"); let shaders = Self { env_prepare: ctx.load_shader("env-prepare.wgsl"), - fill_gbuf: ctx.load_shader("fill-gbuf.wgsl"), ray_trace: ctx.load_shader("ray-trace.wgsl"), blur: ctx.load_shader("blur.wgsl"), post_proc: ctx.load_shader("post-proc.wgsl"), @@ -517,7 +500,6 @@ impl Shaders { } struct ShaderPipelines { - fill: blade_graphics::ComputePipeline, main: blade_graphics::ComputePipeline, temporal_accum: blade_graphics::ComputePipeline, atrous: blade_graphics::ComputePipeline, @@ -527,19 +509,6 @@ struct ShaderPipelines { } impl ShaderPipelines { - fn create_gbuf_fill( - shader: &blade_graphics::Shader, - gpu: &blade_graphics::Context, - ) -> blade_graphics::ComputePipeline { - shader.check_struct_size::(); - shader.check_struct_size::(); - let layout = ::layout(); - gpu.create_compute_pipeline(blade_graphics::ComputePipelineDesc { - name: "fill-gbuf", - data_layouts: &[&layout], - compute: shader.at("main"), - }) - } fn create_ray_trace( shader: &blade_graphics::Shader, gpu: &blade_graphics::Context, @@ -617,7 +586,6 @@ impl ShaderPipelines { let sh_main = shader_man[shaders.ray_trace].raw.as_ref().unwrap(); let sh_blur = shader_man[shaders.blur].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), temporal_accum: Self::create_temporal_accum(sh_blur, gpu), atrous: Self::create_atrous(sh_blur, gpu), @@ -708,7 +676,6 @@ impl Renderer { shaders, targets, post_proc_input_index: 0, - fill_pipeline: sp.fill, main_pipeline: sp.main, post_proc_pipeline: sp.post_proc, blur: Blur { @@ -755,7 +722,6 @@ impl Renderer { // pipelines gpu.destroy_compute_pipeline(&mut self.blur.temporal_accum_pipeline); gpu.destroy_compute_pipeline(&mut self.blur.atrous_pipeline); - gpu.destroy_compute_pipeline(&mut self.fill_pipeline); gpu.destroy_compute_pipeline(&mut self.main_pipeline); gpu.destroy_render_pipeline(&mut self.post_proc_pipeline); } @@ -770,7 +736,6 @@ impl Renderer { let mut tasks = Vec::new(); let old = self.shaders.clone(); - tasks.extend(asset_hub.shaders.hot_reload(&mut self.shaders.fill_gbuf)); tasks.extend(asset_hub.shaders.hot_reload(&mut self.shaders.ray_trace)); tasks.extend(asset_hub.shaders.hot_reload(&mut self.shaders.blur)); tasks.extend(asset_hub.shaders.hot_reload(&mut self.shaders.post_proc)); @@ -787,11 +752,6 @@ impl Renderer { let _ = task.join(); } - if self.shaders.fill_gbuf != old.fill_gbuf { - if let Ok(ref shader) = asset_hub.shaders[self.shaders.fill_gbuf].raw { - self.fill_pipeline = ShaderPipelines::create_gbuf_fill(shader, gpu); - } - } if self.shaders.ray_trace != old.ray_trace { if let Ok(ref shader) = asset_hub.shaders[self.shaders.ray_trace].raw { assert_eq!( @@ -1127,33 +1087,6 @@ impl Renderer { let debug = self.make_debug_params(&debug_config); let (cur, prev) = self.work_indices(); - if let mut pass = command_encoder.compute() { - let mut pc = pass.with(&self.fill_pipeline); - let groups = self.fill_pipeline.get_dispatch_for(self.surface_size); - pc.bind( - 0, - &FillData { - camera: self.targets.camera_params[cur], - prev_camera: self.targets.camera_params[prev], - debug, - acc_struct: self.acceleration_structure, - hit_entries: self.hit_buffer.into(), - index_buffers: &self.index_buffers, - vertex_buffers: &self.vertex_buffers, - textures: &self.textures, - sampler_linear: self.samplers.linear, - debug_buf: self.debug.buffer_resource(), - out_depth: self.targets.depth.views[cur], - out_basis: self.targets.basis.views[cur], - out_flat_normal: self.targets.flat_normal.views[cur], - out_albedo: self.targets.albedo.views[0], - out_motion: self.targets.motion.views[0], - out_debug: self.targets.debug.views[0], - }, - ); - pc.dispatch(groups); - } - if let mut pass = command_encoder.compute() { let grid_scale = { let limit = ray_config.group_mixer; @@ -1204,19 +1137,24 @@ impl Renderer { } else { self.prev_acceleration_structure }, + hit_entries: self.hit_buffer.into(), + index_buffers: &self.index_buffers, + vertex_buffers: &self.vertex_buffers, + textures: &self.textures, 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[prev], - t_basis: self.targets.basis.views[cur], t_prev_basis: self.targets.basis.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], debug_buf: self.debug.buffer_resource(), reservoirs: self.targets.reservoir_buf.into(), + out_depth: self.targets.depth.views[cur], + out_basis: self.targets.basis.views[cur], + out_flat_normal: self.targets.flat_normal.views[cur], + out_albedo: self.targets.albedo.views[0], + out_motion: self.targets.motion.views[0], out_diffuse: self.targets.light_diffuse.views[cur], out_debug: self.targets.debug.views[0], },