wtf this is, some parts of this work
This commit is contained in:
parent
b27472fbb3
commit
3a7cee4650
162
src/bin/textured/hit_mesh.wgsl
Normal file
162
src/bin/textured/hit_mesh.wgsl
Normal file
|
|
@ -0,0 +1,162 @@
|
|||
struct VertexOutput {
|
||||
@builtin(position) clip_position: vec4<f32>,
|
||||
@location(0) base: vec3<f32>,
|
||||
@location(1) dir: vec3<f32>,
|
||||
};
|
||||
|
||||
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<f32> {
|
||||
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<f32>,
|
||||
dir: vec3<f32>,
|
||||
}
|
||||
|
||||
struct Face {
|
||||
vertices: array<vec3<f32>, 3>,
|
||||
tex_coords: mat3x2<f32>,
|
||||
}
|
||||
|
||||
struct Input {
|
||||
n_faces: u32,
|
||||
n_rays: u32,
|
||||
}
|
||||
|
||||
// @group(0) @binding(0)
|
||||
// var<uniform> input: Input;
|
||||
|
||||
@group(0) @binding(0)
|
||||
var<storage, read> faces: array<Face>;
|
||||
|
||||
// @group(0) @binding(1)
|
||||
// var<storage, read> rays: array<Ray>;
|
||||
|
||||
// @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<f32>,
|
||||
}
|
||||
|
||||
fn trace(ray: Ray) -> TraceResult {
|
||||
var have_hit = false;
|
||||
var hit_i_face: u32;
|
||||
var hit_rel: vec3<f32>;
|
||||
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<f32>, b: vec3<f32>, ray: Ray) -> f32 {
|
||||
return determinant(mat3x3(b - a, ray.base - a, -ray.dir));
|
||||
}
|
||||
|
||||
fn inverse(m: mat3x3<f32>) -> mat3x3<f32> {
|
||||
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<f32>) -> mat3x3<f32> {
|
||||
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;
|
||||
}
|
||||
|
|
@ -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<Self>) -> 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<Self>) -> 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<dyn Error>> {
|
|||
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::<TracedBindGroup>::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<dyn Error>> {
|
|||
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::<RawFace>()) 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::<HitMeshBindGroup>::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<dyn Error>> {
|
|||
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<T: BindGroup + ?Sized>(wgpu::BindGroupLayout, PhantomData<*const T>);
|
||||
|
||||
impl<T: BindGroup + ?Sized> BindGroupLayout<T> {
|
||||
pub fn create(device: &wgpu::Device) -> Self {
|
||||
Self(device.create_bind_group_layout(&T::LAYOUT), PhantomData::default())
|
||||
}
|
||||
}
|
||||
|
||||
impl<T: BindGroup + ?Sized> Deref for BindGroupLayout<T> {
|
||||
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<Self>) -> wgpu::BindGroup;
|
||||
}
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue
Block a user