project ideal point

This reverts commit 0ddd41c1b6.

Cargo.lock

@@ -3892,6 +3892,7 @@ dependencies = [
  "glam",
  "id",
  "strafesnet_common",
+ "strafesnet_physics",
  "strafesnet_session",
  "strafesnet_settings",
  "wgpu",

engine/graphics/Cargo.toml

@@ -9,6 +9,7 @@ ddsfile = "0.5.1"
 glam = "0.30.0"
 id = { version = "0.1.0", registry = "strafesnet" }
 strafesnet_common = { path = "../../lib/common", registry = "strafesnet" }
+strafesnet_physics = { path = "../physics", registry = "strafesnet" }
 strafesnet_session = { path = "../session", registry = "strafesnet" }
 strafesnet_settings = { path = "../settings", registry = "strafesnet" }
 wgpu = "28.0.0"

engine/graphics/src/graphics.rs

@@ -5,10 +5,20 @@ use strafesnet_settings::settings;
 use strafesnet_session::session;
 use strafesnet_common::model::{self, ColorId, NormalId, PolygonIter, PositionId, RenderConfigId, TextureCoordinateId, VertexId};
 use wgpu::{util::DeviceExt,AstcBlock,AstcChannel};
-use crate::model::{self as model_graphics,IndexedGraphicsMeshOwnedRenderConfig,IndexedGraphicsMeshOwnedRenderConfigId,GraphicsMeshOwnedRenderConfig,GraphicsModelColor4,GraphicsModelOwned,GraphicsVertex};
+use crate::model::{self as model_graphics,IndexedGraphicsMeshOwnedRenderConfig,IndexedGraphicsMeshOwnedRenderConfigId,GraphicsMeshOwnedRenderConfig,GraphicsModelColor4,GraphicsModelOwned,GraphicsVertex,DebugGraphicsVertex};
 
 pub fn required_limits()->wgpu::Limits{
-	wgpu::Limits::default()
+	let mut limits=wgpu::Limits::default();
+	limits.max_task_invocations_per_dimension=1;
+	limits.max_task_invocations_per_workgroup=1;
+	limits.max_mesh_invocations_per_dimension=1;
+	limits.max_mesh_invocations_per_workgroup=1;
+	limits.max_task_mesh_workgroup_total_count=1;
+	limits.max_task_mesh_workgroups_per_dimension=1;
+	limits.max_task_payload_size=4;
+	limits.max_mesh_output_vertices=2*(3+2+4+8);
+	limits.max_mesh_output_primitives=2*(1+2+4+8)+2;
+	limits
 }
 
 struct Indices{
@@ -36,12 +46,32 @@ struct GraphicsModel{
 	instance_count:u32,
 }
 
+struct DebugGraphicsSubmesh{
+	verts:Vec<strafesnet_physics::model::MeshVertId>,
+	edges:Vec<[strafesnet_physics::model::MeshVertId;2]>,
+	faces:Vec<Indices>,
+}
+struct DebugGraphicsMesh{
+	vertices:Vec<DebugGraphicsVertex>,
+	submeshes:Vec<DebugGraphicsSubmesh>,
+	vertex_buf:wgpu::Buffer,
+}
+struct DebugGraphicsModel{
+	debug_mesh_id:u32,
+	bind_group:wgpu::BindGroup,
+	// 32 bytes used to tell the mesh shader where to draw
+	// Vert: [vec4,_]
+	// Edge: [vec4,vec4]
+	debug_buf:wgpu::Buffer,
+}
+
 struct GraphicsSamplers{
 	repeat:wgpu::Sampler,
 }
 
 struct GraphicsBindGroupLayouts{
 	model:wgpu::BindGroupLayout,
+	debug_model:wgpu::BindGroupLayout,
 }
 
 struct GraphicsBindGroups{
@@ -52,6 +82,9 @@ struct GraphicsBindGroups{
 struct GraphicsPipelines{
 	skybox:wgpu::RenderPipeline,
 	model:wgpu::RenderPipeline,
+	debug_face:wgpu::RenderPipeline,
+	debug_edge:wgpu::RenderPipeline,
+	debug_vert:wgpu::RenderPipeline,
 }
 
 struct GraphicsCamera{
@@ -132,6 +165,8 @@ pub struct GraphicsState{
 	camera_buf:wgpu::Buffer,
 	temp_squid_texture_view:wgpu::TextureView,
 	models:Vec<GraphicsModel>,
+	debug_meshes:Vec<DebugGraphicsMesh>,
+	debug_models:Vec<DebugGraphicsModel>,
 	depth_view:wgpu::TextureView,
 	staging_belt:wgpu::util::StagingBelt,
 }
@@ -166,6 +201,98 @@ impl GraphicsState{
 		self.camera.fov=user_settings.calculate_fov(1.0,&self.camera.screen_size).as_vec2();
 	}
 	pub fn generate_models(&mut self,device:&wgpu::Device,queue:&wgpu::Queue,map:&map::CompleteMap){
+		//generate debug meshes, each debug model refers to one
+		self.debug_meshes=map.meshes.iter().map(|mesh|{
+			let vertices:Vec<DebugGraphicsVertex>=mesh.unique_pos.iter().copied().map(|pos|{
+				DebugGraphicsVertex{
+					pos:pos.to_array().map(Into::into),
+				}
+			}).collect();
+			let vertex_buf=device.create_buffer_init(&wgpu::util::BufferInitDescriptor{
+				label:Some("Vertex"),
+				contents:bytemuck::cast_slice(&vertices),
+				usage:wgpu::BufferUsages::VERTEX,
+			});
+
+			macro_rules! indices{
+				($indices:expr)=>{
+					if (u32::MAX as usize)<vertices.len(){
+						panic!("Model has too many vertices!");
+					}else if (u16::MAX as usize)<vertices.len(){
+						Indices::new(device,&$indices.into_iter().map(|vertex_idx|vertex_idx.get() as u32).collect(),wgpu::IndexFormat::Uint32)
+					}else{
+						Indices::new(device,&$indices.into_iter().map(|vertex_idx|vertex_idx.get() as u16).collect(),wgpu::IndexFormat::Uint16)
+					}
+				};
+			}
+
+			let submeshes=if let Ok(physics_mesh)=strafesnet_physics::model::PhysicsMesh::try_from(mesh){
+				physics_mesh.submesh_views().into_iter().map(|submesh_view|DebugGraphicsSubmesh{
+					verts:submesh_view.verts().to_owned(),
+					edges:submesh_view.edge_vert_ids_iter().collect(),
+					faces:submesh_view.face_vert_ids_iter().map(|face_verts|{
+						// triangulate
+						let mut indices=Vec::new();
+						let mut poly_vertices=face_verts.into_iter();
+						if let (Some(a),Some(mut b))=(poly_vertices.next(),poly_vertices.next()){
+							for c in poly_vertices{
+								indices.extend([a,b,c]);
+								b=c;
+							}
+						}
+						indices!(indices)
+					}).collect(),
+				}).collect()
+			}else{
+				//idc
+				Vec::new()
+			};
+
+			DebugGraphicsMesh{
+				vertices,
+				submeshes,
+				vertex_buf,
+			}
+		}).collect();
+
+		//generate debug models, only one will be rendered at a time
+		self.debug_models=map.models.iter().enumerate().map(|(model_id,model)|{
+			let model_uniforms=get_instances_buffer_data(&[GraphicsModelOwned{
+				transform:model.transform.into(),
+				normal_transform:glam::Mat3::from_cols_array_2d(&model.transform.matrix3.to_array().map(|row|row.map(Into::into))).inverse().transpose(),
+				color:GraphicsModelColor4::new(glam::vec4(1.0,0.0,0.0,0.2)),
+			}]);
+			let model_buf=device.create_buffer_init(&wgpu::util::BufferInitDescriptor{
+				label:Some(format!("Debug Model{} Buf",model_id).as_str()),
+				contents:bytemuck::cast_slice(&model_uniforms),
+				usage:wgpu::BufferUsages::UNIFORM|wgpu::BufferUsages::COPY_DST,
+			});
+			let debug_buf=device.create_buffer_init(&wgpu::util::BufferInitDescriptor{
+				label:Some(format!("Debug Model{} EV Buf",model_id).as_str()),
+				contents:bytemuck::cast_slice(&[0u8;32]),
+				usage:wgpu::BufferUsages::UNIFORM|wgpu::BufferUsages::COPY_DST,
+			});
+			let bind_group=device.create_bind_group(&wgpu::BindGroupDescriptor{
+				layout:&self.bind_group_layouts.debug_model,
+				entries:&[
+					wgpu::BindGroupEntry{
+						binding:0,
+						resource:model_buf.as_entire_binding(),
+					},
+					wgpu::BindGroupEntry{
+						binding:1,
+						resource:debug_buf.as_entire_binding(),
+					},
+				],
+				label:Some(format!("Debug Model{} Bind Group",model_id).as_str()),
+			});
+			DebugGraphicsModel{
+				debug_mesh_id:model.mesh.get(),
+				bind_group,
+				debug_buf,
+			}
+		}).collect();
+
 		//generate texture view per texture
 		let texture_views:HashMap<model::TextureId,wgpu::TextureView>=map.textures.iter().enumerate().filter_map(|(texture_id,texture_data)|{
 			let texture_id=model::TextureId::new(texture_id as u32);
@@ -546,7 +673,7 @@ impl GraphicsState{
 			entries:&[
 				wgpu::BindGroupLayoutEntry{
 					binding:0,
-					visibility:wgpu::ShaderStages::VERTEX,
+					visibility:wgpu::ShaderStages::VERTEX|wgpu::ShaderStages::MESH,
 					ty:wgpu::BindingType::Buffer{
 						ty:wgpu::BufferBindingType::Uniform,
 						has_dynamic_offset:false,
@@ -608,6 +735,31 @@ impl GraphicsState{
 				},
 			],
 		});
+		let debug_model_bind_group_layout=device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor{
+			label:Some("Debug Model Bind Group Layout"),
+			entries:&[
+				wgpu::BindGroupLayoutEntry{
+					binding:0,
+					visibility:wgpu::ShaderStages::VERTEX_FRAGMENT|wgpu::ShaderStages::MESH,
+					ty:wgpu::BindingType::Buffer{
+						ty:wgpu::BufferBindingType::Uniform,
+						has_dynamic_offset:false,
+						min_binding_size:None,
+					},
+					count:None,
+				},
+				wgpu::BindGroupLayoutEntry{
+					binding:1,
+					visibility:wgpu::ShaderStages::MESH,
+					ty:wgpu::BindingType::Buffer{
+						ty:wgpu::BufferBindingType::Uniform,
+						has_dynamic_offset:false,
+						min_binding_size:None,
+					},
+					count:None,
+				},
+			],
+		});
 
 		let clamp_sampler=device.create_sampler(&wgpu::SamplerDescriptor{
 			label:Some("Clamp Sampler"),
@@ -756,6 +908,14 @@ impl GraphicsState{
 			],
 			immediate_size:0,
 		});
+		let debug_model_pipeline_layout=device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor{
+				label:None,
+				bind_group_layouts:&[
+					&camera_bind_group_layout,
+					&debug_model_bind_group_layout,
+				],
+				immediate_size:0,
+			});
 		let sky_pipeline_layout=device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor{
 			label:None,
 			bind_group_layouts:&[
@@ -831,6 +991,92 @@ impl GraphicsState{
 			multiview_mask:None,
 			cache:None,
 		});
+		let debug_model_pipeline_face=device.create_render_pipeline(&wgpu::RenderPipelineDescriptor{
+			label:Some("Debug Face Pipeline"),
+			layout:Some(&debug_model_pipeline_layout),
+			vertex:wgpu::VertexState{
+				module:&shader,
+				entry_point:Some("vs_debug_face"),
+				buffers:&[wgpu::VertexBufferLayout{
+					array_stride:size_of::<DebugGraphicsVertex>() as wgpu::BufferAddress,
+					step_mode:wgpu::VertexStepMode::Vertex,
+					attributes:&wgpu::vertex_attr_array![0=>Float32x3],
+				}],
+				compilation_options:wgpu::PipelineCompilationOptions::default(),
+			},
+			fragment:Some(wgpu::FragmentState{
+				module:&shader,
+				entry_point:Some("fs_debug"),
+				targets:&[Some(wgpu::ColorTargetState{
+					format:config.view_formats[0],
+					blend:Some(wgpu::BlendState::ALPHA_BLENDING),
+					write_mask:wgpu::ColorWrites::default(),
+				})],
+				compilation_options:wgpu::PipelineCompilationOptions::default(),
+			}),
+			primitive:wgpu::PrimitiveState{
+				topology:wgpu::PrimitiveTopology::TriangleList,
+				front_face:wgpu::FrontFace::Cw,
+				cull_mode:Some(wgpu::Face::Front),
+				..Default::default()
+			},
+			depth_stencil:Some(wgpu::DepthStencilState{
+				format:Self::DEPTH_FORMAT,
+				depth_write_enabled:true,
+				depth_compare:wgpu::CompareFunction::Always,
+				stencil:wgpu::StencilState::default(),
+				bias:wgpu::DepthBiasState::default(),
+			}),
+			multisample:wgpu::MultisampleState::default(),
+			multiview_mask:None,
+			cache:None,
+		});
+		let mut debug_model_pipeline=wgpu::MeshPipelineDescriptor{
+			label:None,//filled in below
+			layout:Some(&debug_model_pipeline_layout),
+			task:Some(wgpu::TaskState{
+				module:&shader,
+				entry_point:Some("ts_main"),
+				compilation_options:wgpu::PipelineCompilationOptions::default(),
+			}),
+			mesh:wgpu::MeshState{
+				module:&shader,
+				entry_point:None,//filled in below
+				compilation_options:wgpu::PipelineCompilationOptions::default(),
+			},
+			fragment:Some(wgpu::FragmentState{
+				module:&shader,
+				entry_point:Some("fs_debug"),
+				targets:&[Some(wgpu::ColorTargetState{
+					format:config.view_formats[0],
+					blend:Some(wgpu::BlendState::ALPHA_BLENDING),
+					write_mask:wgpu::ColorWrites::default(),
+				})],
+				compilation_options:wgpu::PipelineCompilationOptions::default(),
+			}),
+			primitive:wgpu::PrimitiveState{
+				topology:wgpu::PrimitiveTopology::TriangleList,
+				front_face:wgpu::FrontFace::Cw,
+				cull_mode:None,
+				..Default::default()
+			},
+			depth_stencil:Some(wgpu::DepthStencilState{
+				format:Self::DEPTH_FORMAT,
+				depth_write_enabled:true,
+				depth_compare:wgpu::CompareFunction::Always,
+				stencil:wgpu::StencilState::default(),
+				bias:wgpu::DepthBiasState::default(),
+			}),
+			multisample:wgpu::MultisampleState::default(),
+			multiview:None,
+			cache:None,
+		};
+		debug_model_pipeline.label=Some("Debug Vert Pipeline");
+		debug_model_pipeline.mesh.entry_point=Some("ms_debug_vert");
+		let debug_model_pipeline_vert=device.create_mesh_pipeline(&debug_model_pipeline);
+		debug_model_pipeline.label=Some("Debug Edge Pipeline");
+		debug_model_pipeline.mesh.entry_point=Some("ms_debug_edge");
+		let debug_model_pipeline_edge=device.create_mesh_pipeline(&debug_model_pipeline);
 
 		let camera=GraphicsCamera::default();
 		let camera_uniforms=camera.to_uniform_data(glam::Vec3::ZERO,glam::Vec2::ZERO);
@@ -870,7 +1116,10 @@ impl GraphicsState{
 		Self{
 			pipelines:GraphicsPipelines{
 				skybox:sky_pipeline,
-				model:model_pipeline
+				model:model_pipeline,
+				debug_face:debug_model_pipeline_face,
+				debug_edge:debug_model_pipeline_edge,
+				debug_vert:debug_model_pipeline_vert,
 			},
 			bind_groups:GraphicsBindGroups{
 				camera:camera_bind_group,
@@ -879,9 +1128,14 @@ impl GraphicsState{
 			camera,
 			camera_buf,
 			models:Vec::new(),
+			debug_meshes:Vec::new(),
+			debug_models:Vec::new(),
 			depth_view,
 			staging_belt:wgpu::util::StagingBelt::new(device.clone(),0x100),
-			bind_group_layouts:GraphicsBindGroupLayouts{model:model_bind_group_layout},
+			bind_group_layouts:GraphicsBindGroupLayouts{
+				model:model_bind_group_layout,
+				debug_model:debug_model_bind_group_layout,
+			},
 			samplers:GraphicsSamplers{repeat:repeat_sampler},
 			temp_squid_texture_view:squid_texture_view,
 		}
@@ -935,6 +1189,49 @@ impl GraphicsState{
 				.copy_from_slice(bytemuck::cast_slice(&model_uniforms));
 		}
 		*/
+		// upload the edge or vertex for teh mesh shader to highlight
+		if let Some(hit)=&frame_state.hit{
+			if let Some(closest_fev)=&hit.closest_fev{
+				let model_id:model::ModelId=hit.convex_mesh_id.model_id.into();
+				if let Some(model)=self.debug_models.get(model_id.get() as usize){
+					let mesh=&self.debug_meshes[model.debug_mesh_id as usize];
+					match closest_fev{
+						strafesnet_physics::model::FEV::Face(_face)=>{
+							// face is rendered normally
+						},
+						strafesnet_physics::model::FEV::Edge(edge)=>{
+							let [v0_id,v1_id]=mesh.submeshes[hit.convex_mesh_id.submesh_id.get() as usize].edges[edge.get() as usize];
+							let v0_pos=mesh.vertices[v0_id.get() as usize].pos;
+							let v1_pos=mesh.vertices[v1_id.get() as usize].pos;
+							let debug_data=[glam::Vec3A::from_array(v0_pos).extend(1.0).to_array(),glam::Vec3A::from_array(v1_pos).extend(1.0).to_array()];
+							let debug_slice=bytemuck::cast_slice(&debug_data);
+							self.staging_belt
+								.write_buffer(
+									&mut encoder,
+									&model.debug_buf,
+									0,
+									wgpu::BufferSize::new(debug_slice.len() as wgpu::BufferAddress).unwrap(),
+								)
+								.copy_from_slice(debug_slice);
+						},
+						strafesnet_physics::model::FEV::Vert(vert)=>{
+							let vert_id=mesh.submeshes[hit.convex_mesh_id.submesh_id.get() as usize].verts[vert.get() as usize].get();
+							let pos=mesh.vertices[vert_id as usize].pos;
+							let debug_data=[glam::Vec3A::from_array(pos).extend(1.0).to_array()];
+							let debug_slice=bytemuck::cast_slice(&debug_data);
+							self.staging_belt
+								.write_buffer(
+									&mut encoder,
+									&model.debug_buf,
+									0,
+									wgpu::BufferSize::new(debug_slice.len() as wgpu::BufferAddress).unwrap(),
+								)
+								.copy_from_slice(debug_slice);
+						},
+					}
+				}
+			}
+		}
 		self.staging_belt.finish();
 
 		{
@@ -970,6 +1267,7 @@ impl GraphicsState{
 			rpass.set_bind_group(0,&self.bind_groups.camera,&[]);
 			rpass.set_bind_group(1,&self.bind_groups.skybox_texture,&[]);
 
+			// Draw all models.
 			rpass.set_pipeline(&self.pipelines.model);
 			for model in &self.models{
 				rpass.set_bind_group(2,&model.bind_group,&[]);
@@ -981,6 +1279,37 @@ impl GraphicsState{
 
 			rpass.set_pipeline(&self.pipelines.skybox);
 			rpass.draw(0..3,0..1);
+
+			// render a single debug_model in red
+			if let Some(hit)=&frame_state.hit{
+				if let Some(closest_fev)=&hit.closest_fev{
+					let model_id:model::ModelId=hit.convex_mesh_id.model_id.into();
+					if let Some(model)=self.debug_models.get(model_id.get() as usize){
+						let mesh=&self.debug_meshes[model.debug_mesh_id as usize];
+						rpass.set_bind_group(1,&model.bind_group,&[]);
+						rpass.set_vertex_buffer(0,mesh.vertex_buf.slice(..));
+						match closest_fev{
+							strafesnet_physics::model::FEV::Face(face)=>{
+								rpass.set_pipeline(&self.pipelines.debug_face);
+								let indices=&mesh.submeshes[hit.convex_mesh_id.submesh_id.get() as usize].faces[face.get() as usize];
+								rpass.set_index_buffer(indices.buf.slice(..),indices.format);
+								//TODO: loop over triangle strips
+								rpass.draw_indexed(0..indices.count,0,0..1);
+							},
+							strafesnet_physics::model::FEV::Edge(_edge)=>{
+								rpass.set_pipeline(&self.pipelines.debug_edge);
+								// the data has already been primed by the staging belt
+								rpass.draw_mesh_tasks(1, 1, 1);
+							},
+							strafesnet_physics::model::FEV::Vert(_vert)=>{
+								rpass.set_pipeline(&self.pipelines.debug_vert);
+								// the data has already been primed by the staging belt
+								rpass.draw_mesh_tasks(1, 1, 1);
+							},
+						}
+					}
+				}
+			}
 		}
 
 		queue.submit(std::iter::once(encoder.finish()));

engine/graphics/src/model.rs

@@ -8,6 +8,11 @@ pub struct GraphicsVertex{
 	pub normal:[f32;3],
 	pub color:[f32;4],
 }
+#[derive(Clone,Copy,Pod,Zeroable)]
+#[repr(C)]
+pub struct DebugGraphicsVertex{
+	pub pos:[f32;3],
+}
 #[derive(Clone,Copy,id::Id)]
 pub struct IndexedGraphicsMeshOwnedRenderConfigId(u32);
 pub struct IndexedGraphicsMeshOwnedRenderConfig{

engine/physics/src/lib.rs

@@ -1,6 +1,6 @@
 mod body;
 mod face_crawler;
-mod model;
+pub mod model;
 mod push_solve;
 mod minimum_difference;
 

engine/physics/src/model.rs

@@ -60,7 +60,7 @@ impl DirectedEdge for SubmeshDirectedEdgeId{
 }
 
 //Vertex <-> Edge <-> Face -> Collide
-#[derive(Debug)]
+#[derive(Debug,Clone)]
 pub enum FEV<M:MeshQuery>{
 	Face(M::Face),
 	Edge(<M::Edge as DirectedEdge>::UndirectedEdge),
@@ -445,6 +445,24 @@ pub struct PhysicsMeshView<'a>{
 	data:&'a PhysicsMeshData,
 	topology:&'a PhysicsMeshTopology,
 }
+impl PhysicsMeshView<'_>{
+	pub fn verts(&self)->&[MeshVertId]{
+		&self.topology.verts
+	}
+	pub fn edge_vert_ids_iter(&self)->impl Iterator<Item=[MeshVertId;2]>+'_{
+		self.topology.edge_topology.iter().map(|edge|{
+			edge.verts.map(|vert_id|self.topology.verts[vert_id.get() as usize])
+		})
+	}
+	pub fn face_vert_ids_iter(&self)->impl Iterator<Item=impl Iterator<Item=MeshVertId>>+'_{
+		self.topology.face_topology.iter().map(|face|{
+			face.edges.iter().map(|edge_id|{
+				let vert_id=self.topology.edge_topology[edge_id.as_undirected().get() as usize].verts[edge_id.parity() as usize];
+				self.topology.verts[vert_id.get() as usize]
+			})
+		})
+	}
+}
 impl MeshQuery for PhysicsMeshView<'_>{
 	type Face=SubmeshFaceId;
 	type Edge=SubmeshDirectedEdgeId;
@@ -527,6 +545,9 @@ impl TransformedMesh<'_>{
 	pub fn verts<'a>(&'a self)->impl Iterator<Item=Vector3<Fixed<2,64>>>+'a{
 		self.view.data.verts.iter().map(|&Vert(pos)|self.transform.vertex.transform_point3(pos))
 	}
+	pub fn faces(&self)->impl Iterator<Item=SubmeshFaceId>{
+		(0..self.view.topology.faces.len() as u32).map(SubmeshFaceId::new)
+	}
 }
 impl MeshQuery for TransformedMesh<'_>{
 	type Face=SubmeshFaceId;

engine/physics/src/physics.rs

@@ -1,4 +1,5 @@
 use std::collections::{HashMap,HashSet};
+use crate::model::DirectedEdge;
 use crate::model::{self as model_physics,PhysicsMesh,PhysicsMeshTransform,TransformedMesh,MeshQuery,PhysicsMeshId,PhysicsSubmeshId};
 use strafesnet_common::bvh;
 use strafesnet_common::map;
@@ -286,7 +287,8 @@ impl PhysicsCamera{
 		.clamp(Self::ANGLE_PITCH_LOWER_LIMIT,Self::ANGLE_PITCH_UPPER_LIMIT);
 		mat3::from_rotation_yx(ax,ay)
 	}
-	fn rotation(&self)->Planar64Mat3{
+	#[inline]
+	pub fn rotation(&self)->Planar64Mat3{
 		self.get_rotation(self.clamped_mouse_pos)
 	}
 	#[expect(dead_code)]
@@ -668,7 +670,7 @@ impl From<CollisionAttributesId> for IntersectAttributesId{
 	}
 }
 #[derive(Debug,Clone,Copy,Hash,id::Id,Eq,PartialEq)]
-struct ContactModelId(u32);
+pub struct ContactModelId(u32);
 impl From<ContactModelId> for ModelId{
 	fn from(value:ContactModelId)->ModelId{
 		ModelId::new(value.get())
@@ -680,7 +682,7 @@ impl From<ModelId> for ContactModelId{
 	}
 }
 #[derive(Debug,Clone,Copy,Hash,id::Id,Eq,PartialEq)]
-struct IntersectModelId(u32);
+pub struct IntersectModelId(u32);
 impl From<IntersectModelId> for ModelId{
 	fn from(value:IntersectModelId)->ModelId{
 		ModelId::new(value.get())
@@ -692,7 +694,7 @@ impl From<ModelId> for IntersectModelId{
 	}
 }
 #[derive(Debug,Clone,Copy,Hash,Eq,PartialEq)]
-enum PhysicsModelId{
+pub enum PhysicsModelId{
 	Contact(ContactModelId),
 	Intersect(IntersectModelId),
 }
@@ -706,9 +708,9 @@ impl From<PhysicsModelId> for ModelId{
 }
 //unique physics meshes indexed by this
 #[derive(Debug,Clone,Copy,Eq,Hash,PartialEq)]
-struct ConvexMeshId<Id>{
+pub struct ConvexMeshId<Id>{
-	model_id:Id,
+	pub model_id:Id,
-	submesh_id:PhysicsSubmeshId,
+	pub submesh_id:PhysicsSubmeshId,
 }
 impl<Id> ConvexMeshId<Id>{
 	fn map<NewId>(self,model_id:NewId)->ConvexMeshId<NewId>{
@@ -997,6 +999,70 @@ impl PhysicsData{
 			hitbox_mesh:StyleModifiers::default().calculate_mesh(),
 		}
 	}
+	pub fn trace_ray(&self,ray:strafesnet_common::ray::Ray)->Option<ConvexMeshId<PhysicsModelId>>{
+		let (_time,convex_mesh_id)=self.bvh.sample_ray(&ray,Time::ZERO,Time::MAX/4,|&model,ray|{
+			let mesh=self.models.mesh(model);
+			// brute force trace every face
+			let faces=mesh.faces().filter_map(|face_id|{
+				let (n,d)=mesh.face_nd(face_id);
+				// trace ray onto face
+				// n.(o+d*t)==n.p
+				// n.o + n.d * t == n.p
+				// t == (n.p - n.o)/n.d
+				let nd=n.dot(ray.direction);
+				if nd.is_zero(){
+					return None;
+				}
+				let t=(d-n.dot(ray.origin))/nd;
+				Some((face_id,t,n))
+			});
+			let mut min=None;
+			for (face_id,t,n) in faces{
+				if let Some(min)=min&&min<t{
+					continue;
+				}
+				// check if point of intersection is behind face edges
+				// *2 because average of 2 vertices
+				let p=ray.extrapolate(t)*2;
+				let is_inside_face=mesh.face_edges(face_id).as_ref().iter().all(|&directed_edge_id|{
+					let edge_n=mesh.directed_edge_n(directed_edge_id);
+					let cross_n=edge_n.cross(n);
+					let &[vert0,vert1]=mesh.edge_verts(directed_edge_id.as_undirected()).as_ref();
+					cross_n.dot(p)<cross_n.dot(mesh.vert(vert0)+mesh.vert(vert1))
+				});
+				if is_inside_face{
+					min=Some(t);
+				}
+			}
+			min.map(Into::into)
+		})?;
+		Some(*convex_mesh_id)
+	}
+	pub fn closest_fev_not_inside(&self,convex_mesh_id:ConvexMeshId<PhysicsModelId>,point:Planar64Vec3)->Option<model_physics::FEV<TransformedMesh<'static>>>{
+		let model_mesh=self.models.mesh(convex_mesh_id);
+		let minkowski=model_physics::MinkowskiMesh::minkowski_sum(model_mesh,self.hitbox_mesh.transformed_mesh());
+		let fev=crate::minimum_difference::closest_fev_not_inside(&minkowski,point)?;
+		Some(match fev{
+			model_physics::FEV::Face(face)=>{
+				match face{
+					model_physics::MinkowskiFace::VertFace(submesh_vert_id,_)=>model_physics::FEV::Vert(submesh_vert_id),
+					model_physics::MinkowskiFace::EdgeEdge(submesh_edge_id,..)=>model_physics::FEV::Edge(submesh_edge_id),
+					model_physics::MinkowskiFace::FaceVert(submesh_face_id,_)=>model_physics::FEV::Face(submesh_face_id),
+				}
+			},
+			model_physics::FEV::Edge(edge)=>{
+				match edge{
+					model_physics::MinkowskiEdge::VertEdge(submesh_vert_id,_)=>model_physics::FEV::Vert(submesh_vert_id),
+					model_physics::MinkowskiEdge::EdgeVert(submesh_edge_id,_)=>model_physics::FEV::Edge(submesh_edge_id),
+				}
+			},
+			model_physics::FEV::Vert(vert)=>{
+				match vert{
+					model_physics::MinkowskiVert::VertVert(submesh_vert_id,_)=>model_physics::FEV::Vert(submesh_vert_id),
+				}
+			},
+		})
+	}
 	pub fn new(map:&map::CompleteMap)->Self{
 		let modes=map.modes.clone().denormalize();
 		let mut used_contact_attributes=Vec::new();

engine/session/src/session.rs

@@ -61,6 +61,7 @@ pub struct FrameState{
 	pub body:physics::Body,
 	pub camera:physics::PhysicsCamera,
 	pub time:PhysicsTime,
+	pub hit:Option<Hit>,
 }
 
 pub struct Simulation{
@@ -77,11 +78,12 @@ impl Simulation{
 			physics,
 		}
 	}
-	pub fn get_frame_state(&self,time:SessionTime)->FrameState{
+	pub fn get_frame_state(&self,time:SessionTime,debug_model:Option<Hit>)->FrameState{
 		FrameState{
 			body:self.physics.camera_body(),
 			camera:self.physics.camera(),
 			time:self.timer.time(time),
+			hit: debug_model,
 		}
 	}
 }
@@ -149,6 +151,12 @@ enum ViewState{
 	Replay(BotId),
 }
 
+#[derive(Clone)]
+pub struct Hit{
+	pub convex_mesh_id:physics::ConvexMeshId<physics::PhysicsModelId>,
+	pub closest_fev:Option<strafesnet_physics::model::FEV<strafesnet_physics::model::TransformedMesh<'static>>>
+}
+
 pub struct Session{
 	directories:Directories,
 	user_settings:UserSettings,
@@ -161,6 +169,7 @@ pub struct Session{
 	recording:Recording,
 	//players:HashMap<PlayerId,Simulation>,
 	replays:HashMap<BotId,Replay>,
+	last_ray_hit:Option<Hit>,
 }
 impl Session{
 	pub fn new(
@@ -177,6 +186,7 @@ impl Session{
 			view_state:ViewState::Play,
 			recording:Default::default(),
 			replays:HashMap::new(),
+			last_ray_hit:None,
 		}
 	}
 	fn clear_recording(&mut self){
@@ -188,12 +198,30 @@ impl Session{
 	}
 	pub fn get_frame_state(&self,time:SessionTime)->Option<FrameState>{
 		match &self.view_state{
-			ViewState::Play=>Some(self.simulation.get_frame_state(time)),
+			ViewState::Play=>Some(self.simulation.get_frame_state(time,self.last_ray_hit.clone())),
 			ViewState::Replay(bot_id)=>self.replays.get(bot_id).map(|replay|
-				replay.simulation.get_frame_state(time)
+				replay.simulation.get_frame_state(time,None)
 			),
 		}
 	}
+	pub fn debug_raycast_print_model_id_if_changed(&mut self,time:SessionTime){
+		if let Some(frame_state)=self.get_frame_state(time){
+			let ray=strafesnet_common::ray::Ray{
+				origin:frame_state.body.extrapolated_position(self.simulation.timer.time(time)),
+				direction:-frame_state.camera.rotation().z_axis,
+			};
+			match self.geometry_shared.trace_ray(ray){
+				Some(convex_mesh_id)=>{
+					let closest_fev=self.geometry_shared.closest_fev_not_inside(convex_mesh_id,self.simulation.physics.body().position);
+					self.last_ray_hit=Some(Hit{
+						convex_mesh_id,
+						closest_fev,
+					});
+				},
+				None=>self.last_ray_hit=None,
+			}
+		}
+	}
 	pub fn user_settings(&self)->&UserSettings{
 		&self.user_settings
 	}

strafe-client/src/physics_worker.rs

@@ -77,5 +77,8 @@ pub fn new<'a>(
 				run_session_instruction!(ins.time,SessionInstruction::LoadReplay(bot));
 			}
 		}
+
+		//whatever just do it
+		session.debug_raycast_print_model_id_if_changed(ins.time);
 	})
 }

strafe-client/src/setup.rs

@@ -3,7 +3,7 @@ fn optional_features()->wgpu::Features{
 	|wgpu::Features::TEXTURE_COMPRESSION_ETC2
 }
 fn required_features()->wgpu::Features{
-	wgpu::Features::TEXTURE_COMPRESSION_BC
+	wgpu::Features::TEXTURE_COMPRESSION_BC|wgpu::Features::EXPERIMENTAL_MESH_SHADER
 }
 fn required_downlevel_capabilities()->wgpu::DownlevelCapabilities{
 	wgpu::DownlevelCapabilities{
@@ -125,7 +125,7 @@ impl<'a> SetupContextPartial3<'a>{
 					required_limits:needed_limits,
 					memory_hints:wgpu::MemoryHints::Performance,
 					trace:wgpu::Trace::Off,
-					experimental_features:wgpu::ExperimentalFeatures::disabled(),
+					experimental_features:unsafe{wgpu::ExperimentalFeatures::enabled()},
 				},
 			))
 			.expect("Unable to find a suitable GPU adapter!");

strafe-client/src/shader.wgsl

@@ -1,3 +1,5 @@
+enable wgpu_mesh_shader;
+
 struct Camera {
 	// from camera to screen
 	proj: mat4x4<f32>,
@@ -86,6 +88,166 @@ fn vs_entity_texture(
 	return result;
 }
 
+@group(1)
+@binding(0)
+var<uniform> model_instance: ModelInstance;
+
+@group(1)
+@binding(1)
+var<uniform> ve_verts: array<vec4<f32>, 2>;
+
+struct DebugEntityOutput {
+	@builtin(position) position: vec4<f32>,
+};
+
+@vertex
+fn vs_debug_face(
+	@location(0) pos: vec3<f32>,
+) -> DebugEntityOutput {
+	var position: vec4<f32> = model_instance.transform * vec4<f32>(pos, 1.0);
+	var result: DebugEntityOutput;
+	result.position = camera.proj * camera.view * position;
+	return result;
+}
+
+struct TaskPayload {
+    four_bytes: u32,
+}
+
+var<task_payload> taskPayload: TaskPayload;
+
+@task
+@payload(taskPayload)
+@workgroup_size(1)
+fn ts_main() -> @builtin(mesh_task_size) vec3<u32> {
+	taskPayload.four_bytes = 0;
+	return vec3(1, 1, 1);
+}
+
+struct VertexOutput {
+	@builtin(position) position: vec4<f32>,
+	@location(0) color: vec4<f32>,
+}
+struct PrimitiveOutput {
+	@builtin(triangle_indices) indices: vec3<u32>,
+}
+
+struct CircleOutput {
+	@builtin(vertices) vertices: array<VertexOutput, 24>,
+	@builtin(primitives) primitives: array<PrimitiveOutput, 22>,
+	@builtin(vertex_count) vertex_count: u32,
+	@builtin(primitive_count) primitive_count: u32,
+}
+
+var<workgroup> mesh_output: CircleOutput;
+
+const tau: f32 = 3.141592653589793 * 2.0;
+
+fn modulo(value:u32,modulus:u32)->u32{
+	return value-value/modulus*modulus;
+}
+
+@mesh(mesh_output)
+@payload(taskPayload)
+@workgroup_size(1)
+fn ms_debug_vert(){
+	// circle with 24 vertices.
+	const LAYERS:u32 = 3;
+	const N:u32 = 3*(1<<LAYERS);
+	mesh_output.vertex_count = N;
+	mesh_output.primitive_count = N-2;
+
+	var vertex_world_position: vec4<f32> = model_instance.transform * ve_verts[0];
+	var vertex_screen_position: vec4<f32> = camera.proj * camera.view * vertex_world_position;
+
+	for (var i:u32 = 0; i<N/4; i++){
+		// draw a 1 unit redius circle
+		var theta: f32 = f32(i) * tau / f32(N);
+		var cos_sin: vec2<f32> = vec2(cos(theta), sin(theta));
+		var offset: vec2<f32> = 0.5 * cos_sin;
+		mesh_output.vertices[i].position = vertex_screen_position + vec4<f32>(offset, 0.0, 0.0);
+		mesh_output.vertices[i+N/4].position = vertex_screen_position + vec4<f32>(-offset.y, offset.x, 0.0, 0.0);
+		mesh_output.vertices[i+N/4*2].position = vertex_screen_position + vec4<f32>(-offset, 0.0, 0.0);
+		mesh_output.vertices[i+N/4*3].position = vertex_screen_position + vec4<f32>(offset.y, -offset.x, 0.0, 0.0);
+	}
+
+	// max area triangle indices
+
+	// the big triangle
+	mesh_output.primitives[0].indices = vec3<u32>(0, N/3, N/3*2);
+
+	// 3 layers of infill triangles to approximate circle better than 1 triangle.
+	// we start on the outer layer because it's easier to construct this way
+	var count:u32=N;
+	var base:u32=1;
+	for (var layer:u32 = 0; layer<LAYERS; layer++){
+		count=count>>1;
+		var step:u32=N/count;
+		for (var i:u32 = 0; i<count; i++){
+			mesh_output.primitives[base+i].indices = vec3<u32>(i*step, i*step+(step>>1), modulo(i*step+step,N));
+		}
+		base+=count;
+	}
+}
+
+@mesh(mesh_output)
+@payload(taskPayload)
+@workgroup_size(1)
+fn ms_debug_edge(){
+	// draw two circles for now.
+	const LAYERS:u32 = 3;
+	const N:u32 = 2*(1<<LAYERS);
+	mesh_output.vertex_count = 2*(3+2+4+8);
+	mesh_output.primitive_count = 2*(1+2+4+8)+2;
+
+	var v0_world_position: vec4<f32> = model_instance.transform * ve_verts[0];
+	var v1_world_position: vec4<f32> = model_instance.transform * ve_verts[1];
+
+	var v0_screen_position: vec4<f32> = camera.proj * camera.view * v0_world_position;
+	var v1_screen_position: vec4<f32> = camera.proj * camera.view * v1_world_position;
+
+	var edge_dir_world: vec4<f32> = normalize(v0_world_position - v1_world_position);
+	var edge_dir_screen: vec4<f32> = camera.proj * camera.view * edge_dir_world;
+
+	for (var i:u32 = 0; i<=N/2; i++){
+		// two half circles that make a whole
+		var theta: f32 = f32(i) * tau / f32(N);
+		var cos_sin: vec2<f32> = vec2(cos(theta), sin(theta));
+
+		// construct basis vectors
+		var y_axis: vec2<f32> = edge_dir_screen.xy;
+		var x_axis: vec2<f32> = y_axis.yx;
+		x_axis.x = -x_axis.x;
+
+		var offset: vec4<f32> = vec4<f32>(0.5 * (x_axis * cos_sin.x + y_axis * cos_sin.y), 0.0, 0.0);;
+
+		mesh_output.vertices[i].position = v0_screen_position + offset;
+		mesh_output.vertices[N/2+1+i].position = v1_screen_position - offset;
+	}
+
+	// max area triangle indices
+	// number of primitives per circle half
+	const P:u32 = N/2;
+	// the big triangles between the circles
+	mesh_output.primitives[0].indices = vec3<u32>(0, N/2+1, P);
+	mesh_output.primitives[P].indices = vec3<u32>(N/2+1, 0, P + N/2+1);
+
+	// 3 layers of infill triangles to approximate circle better than 1 triangle.
+	// we start on the outer layer because it's easier to construct this way
+	var count:u32=P;
+	var base:u32=1;
+	for (var layer:u32 = 0; layer<LAYERS; layer++){
+		count=count>>1;
+		var step:u32=P/count;
+		for (var i:u32 = 0; i<count; i++){
+			var indices = vec3<u32>(i*step, i*step+(step>>1), i*step+step);
+			mesh_output.primitives[base+i].indices = indices;
+			mesh_output.primitives[P+base+i].indices = indices + N/2+1;
+		}
+		base+=count;
+	}
+}
+
 //group 2 is the skybox texture
 @group(1)
 @binding(0)
@@ -110,3 +272,8 @@ fn fs_entity_texture(vertex: EntityOutputTexture) -> @location(0) vec4<f32> {
 	let reflected_color = textureSample(cube_texture, cube_sampler, reflected).rgb;
 	return mix(vec4<f32>(vec3<f32>(0.05) + 0.2 * reflected_color,1.0),mix(vertex.model_color,vec4<f32>(fragment_color.rgb,1.0),fragment_color.a),0.5+0.5*abs(d));
 }
+
+@fragment
+fn fs_debug(vertex: DebugEntityOutput) -> @location(0) vec4<f32> {
+	return model_instance.color;
+}