diff --git a/src/bin/textured/hit_mesh.wgsl b/src/bin/textured/hit_mesh.wgsl new file mode 100644 index 0000000..ba5e5c8 --- /dev/null +++ b/src/bin/textured/hit_mesh.wgsl @@ -0,0 +1,162 @@ +struct VertexOutput { + @builtin(position) clip_position: vec4, + @location(0) base: vec3, + @location(1) dir: vec3, +}; + +const dist: f32 = 40.0; +const off: f32 = 10.0; + +@vertex +fn vs_persp(@builtin(vertex_index) vn: u32) -> VertexOutput { + var out = VertexOutput(vec4(0.0), vec3(0.0), vec3(0.0)); + switch (vn) { + case 0u: { out = VertexOutput(vec4(-1.0, 1.0, 0.0, 1.0), vec3(0.0, 0.0, -dist), vec3(-off, -off, dist)); } + case 1u: { out = VertexOutput(vec4(1.0, 1.0, 0.0, 1.0), vec3(0.0, 0.0, -dist), vec3(off, -off, dist)); } + case 2u: { out = VertexOutput(vec4(-1.0, -1.0, 0.0, 1.0), vec3(0.0, 0.0, -dist), vec3(-off, off, dist)); } + case 3u: { out = VertexOutput(vec4(1.0, -1.0, 0.0, 1.0), vec3(0.0, 0.0, -dist), vec3(off, off, dist)); } + default: { } + } + let PI = 3.1416; + let m_view = ypr_to_mat(vec3(135.0 * PI / 180.0, -30.0 * PI / 180.0, 0.0)); + let m_camera = transpose(m_view); + out.base = m_camera * out.base; + out.dir = m_camera * out.dir; + return out; +} + +@vertex +fn vs_ortho(@builtin(vertex_index) vn: u32) -> VertexOutput { + switch (vn) { + case 0u: { return VertexOutput(vec4(-1.0, 1.0, 0.0, 1.0), vec3(-off, -off, dist), vec3(0.0, 0.0, 1.0)); } + case 1u: { return VertexOutput(vec4(1.0, 1.0, 0.0, 1.0), vec3(off, -off, dist), vec3(0.0, 0.0, 1.0)); } + case 2u: { return VertexOutput(vec4(-1.0, -1.0, 0.0, 1.0), vec3(-off, off, dist), vec3(0.0, 0.0, 1.0)); } + case 3u: { return VertexOutput(vec4(1.0, -1.0, 0.0, 1.0), vec3(off, off, dist), vec3(0.0, 0.0, 1.0)); } + default: { return VertexOutput(vec4(0.0), vec3(0.0), vec3(0.0)); } + } +} + +@fragment +fn fs_main(in: VertexOutput) -> @location(0) vec4 { + let ret = trace(Ray(in.base, in.dir)); + if (ret.have_hit) { + return vec4(ret.tc, 0.0, 1.0); + } else { + return vec4(0.0, 0.0, 0.5, 1.0); + } +// return textureSample(t_traced, s_traced, in.rel_coords); +} + +struct Ray { + base: vec3, + dir: vec3, +} + +struct Face { + vertices: array, 3>, + tex_coords: mat3x2, +} + +struct Input { + n_faces: u32, + n_rays: u32, +} + +// @group(0) @binding(0) +// var input: Input; + +@group(0) @binding(0) +var faces: array; + +// @group(0) @binding(1) +// var rays: array; + +// @compute @workgroup_size(64) +// fn main( +// @builtin(global_invocation_id) +// global_id : vec3u, +// ) { +// if (global_id.x >= input.n_rays) { +// return; +// } +// let ray = rays[global_id.x]; +// // output[global_id.x] = f32(global_id.x) * 1000. + f32(local_id.x); +// } + +struct TraceResult { + have_hit: bool, + face_id: u32, + tc: vec2, +} + +fn trace(ray: Ray) -> TraceResult { + var have_hit = false; + var hit_i_face: u32; + var hit_rel: vec3; + for (var i_face = 0u; i_face < arrayLength(&faces); i_face++) { + let f = faces[i_face]; + let fs = vec3( + edge_dist(f.vertices[0], f.vertices[1], ray), + edge_dist(f.vertices[1], f.vertices[2], ray), + edge_dist(f.vertices[2], f.vertices[0], ray), + ); + if (any(fs < vec3(0.0))) { + continue; + } + let m1 = mat3x3(f.vertices[1] - f.vertices[0], f.vertices[2] - f.vertices[0], -ray.dir); + let m = inverse(m1); + let rel = m * (ray.base - f.vertices[0]); + if (!have_hit || rel.z < hit_rel.z) { + have_hit = true; + hit_i_face = i_face; + hit_rel = rel; + } + } + if (!have_hit) { + return TraceResult(false, 0, vec2(0.0)); + } + let f = faces[hit_i_face]; + let w = vec3(1. - hit_rel.x - hit_rel.y, hit_rel.x, hit_rel.y); + let tc = f.tex_coords * w; + return TraceResult( + true, + hit_i_face, + tc, + ); +} + +fn edge_dist(a: vec3, b: vec3, ray: Ray) -> f32 { + return determinant(mat3x3(b - a, ray.base - a, -ray.dir)); +} + +fn inverse(m: mat3x3) -> mat3x3 { + let tmp0 = cross(m.y, m.z); + let tmp1 = cross(m.z, m.x); + let tmp2 = cross(m.x, m.y); + let det = dot(m.z, tmp2); +// glam_assert!(det != 0.0); + let inv_det = 1.0 / det; + return transpose(inv_det * mat3x3(tmp0, tmp1, tmp2)); +} + +fn ypr_to_mat(ypr: vec3) -> mat3x3 { + let yaw = ypr.x; + let pitch = ypr.y; + let roll = ypr.z; + let m_roll = mat3x3( + vec3(cos(roll), sin(roll), 0.0), + vec3(-sin(roll), cos(roll), 0.0), + vec3(0.0, 0.0, 1.0), + ); + let m_yaw = mat3x3( + vec3(cos(yaw), 0.0, sin(yaw)), + vec3(0.0, 1.0, 0.0), + vec3(-sin(yaw), 0.0, cos(yaw)), + ); + let m_pitch = mat3x3( + vec3(1.0, 0.0, 0.0), + vec3(0.0, cos(pitch), -sin(pitch)), + vec3(0.0, sin(pitch), cos(pitch)), + ); + return m_roll * m_pitch * m_yaw; +} diff --git a/src/bin/textured/main.rs b/src/bin/textured/main.rs index a24454a..195f7ca 100644 --- a/src/bin/textured/main.rs +++ b/src/bin/textured/main.rs @@ -1,5 +1,7 @@ +use bind_group::{BindGroup, BindGroupLayout}; +use bytemuck::{Pod, Zeroable}; use glam::*; -use glam::{vec3, Vec3}; +use glam::{vec3, Vec2, Vec3}; use image::RgbImage; use refraction::mesh_loader::load_mesh; use refraction::mesh_tracer::{trace_to_mesh, Mesh}; @@ -8,6 +10,7 @@ use std::error::Error; use std::f32::consts::PI; use std::fs::File; use std::io::BufReader; +use std::marker::PhantomData; use std::process::exit; use std::sync::atomic::{AtomicUsize, Ordering::Relaxed}; use winit::keyboard::{KeyCode, PhysicalKey}; @@ -119,6 +122,83 @@ fn test_projs() { check(ortho, 9., 1., 8.); } +struct TracedBindGroup<'a> { + t_traced: &'a wgpu::TextureView, + s_traced: &'a wgpu::Sampler, +} + +impl<'a> BindGroup for TracedBindGroup<'a> { + const LAYOUT: wgpu::BindGroupLayoutDescriptor<'static> = wgpu::BindGroupLayoutDescriptor { + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Texture { + multisampled: false, + view_dimension: wgpu::TextureViewDimension::D2, + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + ], + label: Some("TracedBindGroup"), + }; + + fn create(self, device: &wgpu::Device, layout: &BindGroupLayout) -> wgpu::BindGroup { + device.create_bind_group(&wgpu::BindGroupDescriptor { + layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(self.t_traced), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::Sampler(self.s_traced), + }, + ], + label: Some("TracedBindGroup"), + }) + } +} + +struct HitMeshBindGroup<'a> { + buf_faces: wgpu::BufferBinding<'a>, +} + +impl<'a> BindGroup for HitMeshBindGroup<'a> { + const LAYOUT: wgpu::BindGroupLayoutDescriptor<'static> = wgpu::BindGroupLayoutDescriptor { + entries: &[wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }], + label: Some("HitMeshBindGroup"), + }; + + fn create(self, device: &wgpu::Device, layout: &BindGroupLayout) -> wgpu::BindGroup { + device.create_bind_group(&wgpu::BindGroupDescriptor { + layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::Buffer(self.buf_faces), + }], + label: Some("HitMeshBindGroup"), + }) + } +} + // add_event_handler wants 'static + Send. Let it be so. static PROJ_INDEX: AtomicUsize = AtomicUsize::new(0); static PROJS: [fn(dist: f32, off: Vec2) -> (Vec3, Vec3); 2] = [persp, ortho]; @@ -169,41 +249,12 @@ fn main() -> Result<(), Box> { mipmap_filter: wgpu::FilterMode::Nearest, ..Default::default() }); - let traced_bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - entries: &[ - wgpu::BindGroupLayoutEntry { - binding: 0, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Texture { - multisampled: false, - view_dimension: wgpu::TextureViewDimension::D2, - sample_type: wgpu::TextureSampleType::Float { filterable: true }, - }, - count: None, - }, - wgpu::BindGroupLayoutEntry { - binding: 1, - visibility: wgpu::ShaderStages::FRAGMENT, - ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), - count: None, - }, - ], - label: Some("traced_bind_group_layout"), - }); - let traced_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - layout: &traced_bind_group_layout, - entries: &[ - wgpu::BindGroupEntry { - binding: 0, - resource: wgpu::BindingResource::TextureView(&traced_view), - }, - wgpu::BindGroupEntry { - binding: 1, - resource: wgpu::BindingResource::Sampler(&traced_sampler), - }, - ], - label: Some("traced_bind_group"), - }); + let traced_bind_group_layout = BindGroupLayout::::create(&device); + let traced_bind_group = TracedBindGroup { + t_traced: &traced_view, + s_traced: &traced_sampler, + } + .create(&device, &traced_bind_group_layout); let present_shader = device.create_shader_module(wgpu::include_wgsl!("present.wgsl")); let render_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("Render Pipeline Layout"), @@ -239,6 +290,73 @@ fn main() -> Result<(), Box> { cache: None, }); + #[repr(C, align(16))] + #[derive(Clone, Copy, Pod, Zeroable)] + struct RawFace { + vertices: [[f32; 4]; 3], + tex_coords: [[f32; 2]; 3], + pad: [u32; 2], + } + + let faces_buf = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: (mesh.len() * size_of::()) as u64, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let faces_content: Vec<_> = mesh + .iter() + .map(|face| RawFace { + vertices: face.vertices.map(|v| [v.x, v.y, v.z, 1.0]), + tex_coords: face.tex_coords.map(|v| v.into()), + pad: [0; 2], + }) + .collect(); + queue.write_buffer(&faces_buf, 0, bytemuck::cast_slice(&faces_content)); + queue.submit([]); + + let hit_mesh_bind_group_layout = BindGroupLayout::::create(&device); + let hit_mesh_bind_group = HitMeshBindGroup { + buf_faces: faces_buf.as_entire_buffer_binding(), + } + .create(&device, &hit_mesh_bind_group_layout); + + let hit_mesh_shader = device.create_shader_module(wgpu::include_wgsl!("hit_mesh.wgsl")); + + let render_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("Render Pipeline Layout"), + bind_group_layouts: &[&hit_mesh_bind_group_layout], + push_constant_ranges: &[], + }); + let render_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some("Render Pipeline"), + layout: Some(&render_pipeline_layout), + vertex: wgpu::VertexState { + module: &hit_mesh_shader, + entry_point: "vs_persp", + buffers: &[], + compilation_options: wgpu::PipelineCompilationOptions::default(), + }, + fragment: Some(wgpu::FragmentState { + module: &hit_mesh_shader, + entry_point: "fs_main", + targets: &[Some(wgpu::ColorTargetState { + format: wgpu::TextureFormat::Bgra8UnormSrgb, + blend: Some(wgpu::BlendState::REPLACE), + write_mask: wgpu::ColorWrites::ALL, + })], + compilation_options: wgpu::PipelineCompilationOptions::default(), + }), + primitive: wgpu::PrimitiveState { + topology: wgpu::PrimitiveTopology::TriangleStrip, + ..Default::default() + }, + depth_stencil: None, + multisample: wgpu::MultisampleState::default(), + multiview: None, + cache: None, + }); + let mut surface_configured = false; let mut phi = 0; event_loop.run(move |event, control_flow| match event { @@ -318,7 +436,7 @@ fn main() -> Result<(), Box> { timestamp_writes: None, }); render_pass.set_pipeline(&render_pipeline); - render_pass.set_bind_group(0, &traced_bind_group, &[]); + render_pass.set_bind_group(0, &hit_mesh_bind_group, &[]); render_pass.draw(0..4, 0..1); drop(render_pass); queue.submit(std::iter::once(encoder.finish())); @@ -359,3 +477,28 @@ async fn init_gpu(wnd: &Window) -> Result<(wgpu::Device, wgpu::Queue, wgpu::Surf .unwrap(); Ok((device, queue, surface)) } + +mod bind_group { + use std::{marker::PhantomData, ops::Deref}; + + pub struct BindGroupLayout(wgpu::BindGroupLayout, PhantomData<*const T>); + + impl BindGroupLayout { + pub fn create(device: &wgpu::Device) -> Self { + Self(device.create_bind_group_layout(&T::LAYOUT), PhantomData::default()) + } + } + + impl Deref for BindGroupLayout { + type Target = wgpu::BindGroupLayout; + + fn deref(&self) -> &Self::Target { + &self.0 + } + } + + pub trait BindGroup { + const LAYOUT: wgpu::BindGroupLayoutDescriptor<'_>; + fn create(self, device: &wgpu::Device, layout: &BindGroupLayout) -> wgpu::BindGroup; + } +}