From 8ab52ba34158052d62b78c4d4247145e70cb5876 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Thu, 30 Mar 2023 23:52:28 -0700 Subject: [PATCH] render: activate the fill stage, break the basis --- blade-graphics/src/gles/mod.rs | 1 + blade-graphics/src/lib.rs | 1 + blade-graphics/src/metal/mod.rs | 1 + blade-graphics/src/util.rs | 1 + blade-graphics/src/vulkan/mod.rs | 1 + blade-render/shader.wgsl | 177 ++++++++++++++++--------------- blade-render/src/renderer.rs | 66 +++++++----- 7 files changed, 139 insertions(+), 109 deletions(-) diff --git a/blade-graphics/src/gles/mod.rs b/blade-graphics/src/gles/mod.rs index 9296d1dc..7d5b9ea1 100644 --- a/blade-graphics/src/gles/mod.rs +++ b/blade-graphics/src/gles/mod.rs @@ -451,6 +451,7 @@ fn describe_texture_format(format: crate::TextureFormat) -> FormatInfo { 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), + Tf::Rgba8Snorm => (glow::RGBA8, glow::RGBA, glow::BYTE), Tf::Rgba16Float => (glow::RGBA16F, glow::RGBA, glow::FLOAT), Tf::R32Float => (glow::R32F, glow::RED, glow::FLOAT), Tf::Depth32Float => (glow::DEPTH_COMPONENT32F, glow::DEPTH_COMPONENT, glow::FLOAT), diff --git a/blade-graphics/src/lib.rs b/blade-graphics/src/lib.rs index 246c157a..58c3f00c 100644 --- a/blade-graphics/src/lib.rs +++ b/blade-graphics/src/lib.rs @@ -195,6 +195,7 @@ pub enum TextureFormat { Rgba8Unorm, Rgba8UnormSrgb, Bgra8UnormSrgb, + Rgba8Snorm, Rgba16Float, R32Float, Depth32Float, diff --git a/blade-graphics/src/metal/mod.rs b/blade-graphics/src/metal/mod.rs index c09100a7..9568b435 100644 --- a/blade-graphics/src/metal/mod.rs +++ b/blade-graphics/src/metal/mod.rs @@ -263,6 +263,7 @@ fn map_texture_format(format: crate::TextureFormat) -> metal::MTLPixelFormat { Tf::Rgba8Unorm => RGBA8Unorm, Tf::Rgba8UnormSrgb => RGBA8Unorm_sRGB, Tf::Bgra8UnormSrgb => BGRA8Unorm_sRGB, + Tf::Rgba8Snorm => RGBA8Snorm, Tf::Rgba16Float => RGBA16Float, Tf::R32Float => R32Float, Tf::Depth32Float => Depth32Float, diff --git a/blade-graphics/src/util.rs b/blade-graphics/src/util.rs index 8ccbcbb8..395425ac 100644 --- a/blade-graphics/src/util.rs +++ b/blade-graphics/src/util.rs @@ -53,6 +53,7 @@ impl super::TextureFormat { Self::Rgba8Unorm => uncompressed(4), Self::Rgba8UnormSrgb => uncompressed(4), Self::Bgra8UnormSrgb => uncompressed(4), + Self::Rgba8Snorm => uncompressed(4), Self::Rgba16Float => uncompressed(8), Self::R32Float => uncompressed(4), Self::Depth32Float => uncompressed(4), diff --git a/blade-graphics/src/vulkan/mod.rs b/blade-graphics/src/vulkan/mod.rs index dd5d8394..4898fd9b 100644 --- a/blade-graphics/src/vulkan/mod.rs +++ b/blade-graphics/src/vulkan/mod.rs @@ -429,6 +429,7 @@ fn map_texture_format(format: crate::TextureFormat) -> vk::Format { Tf::Rgba8Unorm => vk::Format::R8G8B8A8_UNORM, Tf::Rgba8UnormSrgb => vk::Format::R8G8B8A8_SRGB, Tf::Bgra8UnormSrgb => vk::Format::B8G8R8A8_SRGB, + Tf::Rgba8Snorm => vk::Format::R8G8B8A8_SNORM, Tf::Rgba16Float => vk::Format::R16G16B16A16_SFLOAT, Tf::R32Float => vk::Format::R32_SFLOAT, Tf::Depth32Float => vk::Format::D32_SFLOAT, diff --git a/blade-render/shader.wgsl b/blade-render/shader.wgsl index cbf75d5b..7fd90a30 100644 --- a/blade-render/shader.wgsl +++ b/blade-render/shader.wgsl @@ -6,11 +6,11 @@ struct CameraParams { depth: f32, orientation: vec4, fov: vec2, + mouse_pos: vec2, }; struct MainParams { frame_index: u32, debug_mode: u32, - mouse_pos: vec2, num_environment_samples: u32, }; @@ -117,34 +117,93 @@ fn debug_fs(in: DebugVarying) -> @location(0) vec4 { } var out_depth: texture_storage_2d; +var out_basis: texture_storage_2d; +var out_albedo: texture_storage_2d; + +fn get_ray_direction(global_id: vec2, target_size: vec2) -> vec3 { + let half_size = vec2(target_size >> vec2(1u)); + let ndc = (vec2(global_id) - half_size) / half_size; + let local_dir = vec3(ndc * tan(camera.fov), 1.0); + return normalize(qrot(camera.orientation, local_dir)); +} @compute @workgroup_size(8, 8) fn fill_gbuf(@builtin(global_invocation_id) global_id: vec3) { let target_size = textureDimensions(out_depth); - let half_size = vec2(target_size >> vec2(1u)); - let ndc = (vec2(global_id.xy) - half_size) / half_size; if (any(global_id.xy > target_size)) { return; } - let global_index = global_id.y * target_size.x + global_id.x; - let local_dir = vec3(ndc * tan(camera.fov), 1.0); - let world_dir = normalize(qrot(camera.orientation, local_dir)); - var rq: ray_query; - var ray_pos = camera.position; - var ray_dir = world_dir; - rayQueryInitialize(&rq, acc_struct, RayDesc(0x10u, 0xFFu, 0.0, camera.depth, ray_pos, ray_dir)); + let ray_dir = get_ray_direction(global_id.xy, target_size); + rayQueryInitialize(&rq, acc_struct, RayDesc(0x10u, 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 albedo = vec3(0.0); if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) { + let enable_debug = all(global_id.xy == camera.mouse_pos); + 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 barycentrics = vec3(1.0 - intersection.barycentrics.x - intersection.barycentrics.y, intersection.barycentrics); + let tex_coords = mat3x2(vertices[0].tex_coords, vertices[1].tex_coords, vertices[2].tex_coords) * barycentrics; + let normal_rough = mat3x2(unpack2x16snorm(vertices[0].normal), unpack2x16snorm(vertices[1].normal), unpack2x16snorm(vertices[2].normal)) * barycentrics; + let normal_object = vec3(normal_rough, sqrt(max(0.0, 1.0 - dot(normal_rough, normal_rough)))); + let object_to_world_rot = normalize(unpack4x8snorm(entry.rotation)); + let normal_world = qrot(object_to_world_rot, normal_object); + + if (enable_debug) { + let debug_len = intersection.t * 0.2; + let positions = intersection.object_to_world * mat3x4( + vec4(vertices[0].pos, 1.0), vec4(vertices[1].pos, 1.0), vec4(vertices[2].pos, 1.0) + ); + debug_line(positions[0].xyz, positions[1].xyz, 0x00FF00u); + debug_line(positions[1].xyz, positions[2].xyz, 0x00FF00u); + debug_line(positions[2].xyz, positions[0].xyz, 0x00FF00u); + let poly_normal = normalize(cross(positions[1].xyz - positions[0].xyz, positions[2].xyz - positions[0].xyz)); + let poly_center = (positions[0].xyz + positions[1].xyz + positions[2].xyz) / 3.0; + debug_line(poly_center, poly_center + debug_len * poly_normal, 0x0000FFu); + let pos_world = camera.position + intersection.t * ray_dir; + debug_line(pos_world, pos_world + debug_len * normal_world, 0xFF0000u); + } + + let pre_tangent = select( + vec3(0.0, 0.0, 1.0), + vec3(0.0, 1.0, 0.0), + abs(dot(normal_world, vec3(0.0, 1.0, 0.0))) < abs(dot(normal_world, vec3(0.0, 0.0, 1.0)))); + let bitangent = normalize(cross(normal_world, pre_tangent)); + let tangent = normalize(cross(bitangent, normal_world)); + basis = vec4(normal_world, 0.0); //TODO: construct quaternion at point + + let base_color_factor = unpack4x8unorm(entry.base_color_factor); + let lod = 0.0; //TODO: this is actually complicated + let base_color_sample = textureSampleLevel(textures[entry.base_color_texture], sampler_linear, tex_coords, lod); + albedo = (base_color_factor * base_color_sample).xyz; } textureStore(out_depth, global_id.xy, vec4(depth, 0.0, 0.0, 0.0)); + textureStore(out_basis, global_id.xy, basis); + textureStore(out_albedo, global_id.xy, vec4(albedo, 0.0)); } +var in_depth: texture_2d; +var in_basis: texture_2d; +var in_albedo: texture_2d; var output: texture_storage_2d; struct RandomState { @@ -234,73 +293,41 @@ fn evaluate_environment(dir: vec3) -> vec3 { return vec3(max(2.0 * dir.y, 0.0)); } -fn compute_hit_color(ri: RayIntersection, ray_dir: vec3, hit_world: vec3, rng: ptr, enable_debug: bool) -> vec3 { - let entry = hit_entries[ri.instance_custom_index + ri.geometry_index]; - if (parameters.debug_mode == DEBUG_MODE_DEPTH) { - return vec3(ri.t / camera.depth); - } +struct Surface { + basis: vec4, + albedo: vec3, +} - var indices = ri.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]); +fn compute_hit_color(ray_dir: vec3, depth: f32, surface: Surface, rng: ptr) -> vec3 { + if (parameters.debug_mode == DEBUG_MODE_DEPTH) { + return vec3(depth / camera.depth); } - - let vptr = &vertex_buffers[entry.vertex_buf].data; - let vertices = array( - (*vptr)[indices.x], - (*vptr)[indices.y], - (*vptr)[indices.z], - ); - - let barycentrics = vec3(1.0 - ri.barycentrics.x - ri.barycentrics.y, ri.barycentrics); - //let pos_object = mat3x3(vertices[0].pos, vertices[1].pos, vertices[2].pos) * barycentrics; - //let pos_world = ri.object_to_world * vec4(pos_object, 1.0); - let tex_coords = mat3x2(vertices[0].tex_coords, vertices[1].tex_coords, vertices[2].tex_coords) * barycentrics; - let normal_rough = mat3x2(unpack2x16snorm(vertices[0].normal), unpack2x16snorm(vertices[1].normal), unpack2x16snorm(vertices[2].normal)) * barycentrics; - let normal_object = vec3(normal_rough, sqrt(max(0.0, 1.0 - dot(normal_rough, normal_rough)))); - let object_to_world_rot = normalize(unpack4x8snorm(entry.rotation)); - let normal_world = qrot(object_to_world_rot, normal_object); - - let debug_len = ri.t * 0.2; - if (enable_debug) { - let positions = ri.object_to_world * mat3x4(vec4(vertices[0].pos, 1.0), vec4(vertices[1].pos, 1.0), vec4(vertices[2].pos, 1.0)); - debug_line(positions[0].xyz, positions[1].xyz, 0x00FF00u); - debug_line(positions[1].xyz, positions[2].xyz, 0x00FF00u); - debug_line(positions[2].xyz, positions[0].xyz, 0x00FF00u); - let poly_normal = normalize(cross(positions[1].xyz - positions[0].xyz, positions[2].xyz - positions[0].xyz)); - let poly_center = (positions[0].xyz + positions[1].xyz + positions[2].xyz) / 3.0; - debug_line(poly_center, poly_center + debug_len * poly_normal, 0x0000FFu); - debug_line(hit_world, hit_world + debug_len * normal_world, 0xFF0000u); + if (depth == 0.0) { + return evaluate_environment(ray_dir); } - // Note: this line allows to check the correctness of data passed in - //return abs(pos_world - hit_world) * 1000000.0; + let position = camera.position + depth * ray_dir; + let normal = qrot(surface.basis, vec3(0.0, 0.0, 1.0)); if (parameters.debug_mode == DEBUG_MODE_NORMAL) { - return normal_world; + return normal; } - let pre_tangent = select( - vec3(0.0, 0.0, 1.0), - vec3(0.0, 1.0, 0.0), - abs(dot(normal_world, vec3(0.0, 1.0, 0.0))) < abs(dot(normal_world, vec3(0.0, 0.0, 1.0)))); - let bitangent = normalize(cross(normal_world, pre_tangent)); - let tangent = normalize(cross(bitangent, normal_world)); - let steradians_in_hemisphere = 2.0 * PI; let lambert_brdf = 1.0 / PI; + let start_t = 0.5; // some offset required to avoid self-shadowing + var color = vec3(0.0); var rq: ray_query; for (var i = 0u; i < parameters.num_environment_samples; i += 1u) { let light_dir_tbn = sample_uniform_hemisphere(rng); - let light_dir = light_dir_tbn.x * tangent + light_dir_tbn.y * bitangent + light_dir_tbn.z * normal_world; + let light_dir = qrot(surface.basis, light_dir_tbn); let lambert_term = light_dir_tbn.z; let estimate_color = evaluate_environment(light_dir) * lambert_term; if (dot(estimate_color, estimate_color) < 0.01) { continue; } - rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, 0.5, camera.depth, hit_world, light_dir)); + rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, start_t, camera.depth, position, light_dir)); rayQueryProceed(&rq); let intersection = rayQueryGetCommittedIntersection(&rq); if (intersection.kind == RAY_QUERY_INTERSECTION_NONE) { @@ -308,45 +335,25 @@ fn compute_hit_color(ri: RayIntersection, ray_dir: vec3, hit_world: vec3) { let target_size = textureDimensions(output); - let half_size = vec2(target_size >> vec2(1u)); - let ndc = (vec2(global_id.xy) - half_size) / half_size; if (any(global_id.xy > target_size)) { return; } let global_index = global_id.y * target_size.x + global_id.x; var rng = random_init(global_index, parameters.frame_index); - let local_dir = vec3(ndc * tan(camera.fov), 1.0); - let world_dir = normalize(qrot(camera.orientation, local_dir)); - var rq: ray_query; - var ray_pos = camera.position; - var ray_dir = world_dir; - rayQueryInitialize(&rq, acc_struct, RayDesc(0x10u, 0xFFu, 0.0, camera.depth, ray_pos, ray_dir)); - var iterations = 0u; - while (rayQueryProceed(&rq)) {iterations += 1u;} - let intersection = rayQueryGetCommittedIntersection(&rq); - - var color = vec3(0.0); - if (intersection.kind == RAY_QUERY_INTERSECTION_NONE) { - color = evaluate_environment(ray_dir); - } else { - let expected = ray_pos + intersection.t * ray_dir; - let enable_debug = all(global_id.xy == parameters.mouse_pos); - color = compute_hit_color(intersection, ray_dir, expected, &rng, enable_debug); - } + var surface: Surface; + let ray_dir = get_ray_direction(global_id.xy, target_size); + let depth = textureLoad(in_depth, global_id.xy, 0).x; + surface.basis = normalize(textureLoad(in_basis, global_id.xy, 0)); + surface.albedo = textureLoad(in_albedo, global_id.xy, 0).xyz; + let color = compute_hit_color(ray_dir, depth, surface, &rng); textureStore(output, global_id.xy, vec4(color, 1.0)); } diff --git a/blade-render/src/renderer.rs b/blade-render/src/renderer.rs index 3b0f0e4d..8027d053 100644 --- a/blade-render/src/renderer.rs +++ b/blade-render/src/renderer.rs @@ -43,6 +43,10 @@ struct Targets { main_view: blade::TextureView, depth: blade::Texture, depth_view: blade::TextureView, + basis: blade::Texture, + basis_view: blade::TextureView, + albedo: blade::Texture, + albedo_view: blade::TextureView, } impl Targets { @@ -78,11 +82,21 @@ impl Targets { let (depth, depth_view) = Self::create_target("depth", blade::TextureFormat::R32Float, size, gpu); encoder.init_texture(depth); + let (basis, basis_view) = + Self::create_target("basis", blade::TextureFormat::Rgba8Snorm, size, gpu); + encoder.init_texture(basis); + let (albedo, albedo_view) = + Self::create_target("basis", blade::TextureFormat::Rgba8Unorm, size, gpu); + encoder.init_texture(albedo); Self { main, main_view, depth, depth_view, + basis, + basis_view, + albedo, + albedo_view, } } @@ -91,6 +105,10 @@ impl Targets { gpu.destroy_texture(self.main); gpu.destroy_texture_view(self.depth_view); gpu.destroy_texture(self.depth); + gpu.destroy_texture_view(self.basis_view); + gpu.destroy_texture(self.basis); + gpu.destroy_texture_view(self.albedo_view); + gpu.destroy_texture(self.albedo); } } @@ -122,7 +140,7 @@ struct CameraParams { depth: f32, orientation: [f32; 4], fov: [f32; 2], - pad: [u32; 2], + mouse_pos: [i32; 2], } #[repr(C)] @@ -130,9 +148,7 @@ struct CameraParams { struct MainParams { frame_index: u32, debug_mode: u32, - mouse_pos: [i32; 2], num_environment_samples: u32, - pad: u32, } #[derive(blade_macros::ShaderData)] @@ -146,18 +162,18 @@ struct FillData<'a> { sampler_linear: blade::Sampler, debug_buf: blade::BufferPiece, out_depth: blade::TextureView, + out_basis: blade::TextureView, + out_albedo: blade::TextureView, } #[derive(blade_macros::ShaderData)] -struct MainData<'a> { +struct MainData { camera: CameraParams, parameters: MainParams, acc_struct: blade::AccelerationStructure, - hit_entries: blade::BufferPiece, - index_buffers: &'a blade::BufferArray, - vertex_buffers: &'a blade::BufferArray, - textures: &'a blade::TextureArray, - sampler_linear: blade::Sampler, + in_depth: blade::TextureView, + in_basis: blade::TextureView, + in_albedo: blade::TextureView, debug_buf: blade::BufferPiece, output: blade::TextureView, } @@ -537,14 +553,21 @@ impl Renderer { } } - fn make_camera_params(&self, camera: &super::Camera) -> CameraParams { + fn make_camera_params( + &self, + camera: &super::Camera, + mouse_pos: Option<[i32; 2]>, + ) -> CameraParams { let fov_x = camera.fov_y * self.screen_size.width as f32 / self.screen_size.height as f32; CameraParams { position: camera.pos.into(), depth: camera.depth, orientation: camera.rot.into(), fov: [fov_x, camera.fov_y], - pad: [0; 2], + mouse_pos: match mouse_pos { + Some(p) => [p[0], self.screen_size.height as i32 - p[1]], + None => [-1; 2], + }, } } @@ -570,7 +593,7 @@ impl Renderer { pc.bind( 0, &FillData { - camera: self.make_camera_params(camera), + camera: self.make_camera_params(camera, mouse_pos), acc_struct: self.acceleration_structure, hit_entries: self.hit_buffer.into(), index_buffers: &self.index_buffers, @@ -579,6 +602,8 @@ impl Renderer { sampler_linear: self.samplers.linear, debug_buf: self.debug.buffer.into(), out_depth: self.targets.depth_view, + out_basis: self.targets.basis_view, + out_albedo: self.targets.albedo_view, }, ); pc.dispatch(group_count); @@ -596,23 +621,16 @@ impl Renderer { pc.bind( 0, &MainData { - camera: self.make_camera_params(camera), + camera: self.make_camera_params(camera, mouse_pos), parameters: MainParams { frame_index: self.frame_index, debug_mode: debug_mode as u32, - mouse_pos: match mouse_pos { - Some(p) => [p[0], self.screen_size.height as i32 - p[1]], - None => [-1; 2], - }, num_environment_samples: ray_config.num_environment_samples, - pad: 0, }, 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, + in_depth: self.targets.depth_view, + in_basis: self.targets.basis_view, + in_albedo: self.targets.albedo_view, debug_buf: self.debug.buffer.into(), output: self.targets.main_view, }, @@ -635,7 +653,7 @@ impl Renderer { pc.bind( 0, &DebugData { - camera: self.make_camera_params(camera), + camera: self.make_camera_params(camera, None), debug_buf: self.debug.buffer.into(), }, );