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

@@ -8,7 +8,17 @@ use wgpu::{util::DeviceExt,AstcBlock,AstcChannel};
 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,13 +46,23 @@ 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{
-	indices:Indices,
+	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{
@@ -62,7 +82,9 @@ struct GraphicsBindGroups{
 struct GraphicsPipelines{
 	skybox:wgpu::RenderPipeline,
 	model:wgpu::RenderPipeline,
-	debug:wgpu::RenderPipeline,
+	debug_face:wgpu::RenderPipeline,
+	debug_edge:wgpu::RenderPipeline,
+	debug_vert:wgpu::RenderPipeline,
 }
 
 struct GraphicsCamera{
@@ -181,10 +203,9 @@ impl GraphicsState{
 	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_vertices.iter().map(|vertex|{
+			let vertices:Vec<DebugGraphicsVertex>=mesh.unique_pos.iter().copied().map(|pos|{
 				DebugGraphicsVertex{
-					pos:mesh.unique_pos[vertex.pos.get() as usize].to_array().map(Into::into),
-					normal:mesh.unique_normal[vertex.normal.get() as usize].to_array().map(Into::into),
+					pos:pos.to_array().map(Into::into),
 				}
 			}).collect();
 			let vertex_buf=device.create_buffer_init(&wgpu::util::BufferInitDescriptor{
@@ -193,30 +214,43 @@ impl GraphicsState{
 				usage:wgpu::BufferUsages::VERTEX,
 			});
 
-			let mut indices=Vec::new();
-			for physics_group in &mesh.physics_groups{
-				for polygon_group_id in &physics_group.groups{
-					for poly in mesh.polygon_groups[polygon_group_id.get() as usize].polys(){
+			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 poly_vertices=poly.into_iter().copied();
+						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{
-				indices: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)
-				},
+				vertices,
+				submeshes,
 				vertex_buf,
 			}
 		}).collect();
@@ -233,6 +267,11 @@ impl GraphicsState{
 				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:&[
@@ -240,12 +279,17 @@ impl GraphicsState{
 						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();
 
@@ -629,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,
@@ -696,7 +740,17 @@ impl GraphicsState{
 			entries:&[
 				wgpu::BindGroupLayoutEntry{
 					binding:0,
-					visibility:wgpu::ShaderStages::VERTEX_FRAGMENT,
+					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,
@@ -860,7 +914,7 @@ impl GraphicsState{
 					&camera_bind_group_layout,
 					&debug_model_bind_group_layout,
 				],
-				push_constant_ranges:&[],
+				immediate_size:0,
 			});
 		let sky_pipeline_layout=device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor{
 			label:None,
@@ -937,16 +991,16 @@ impl GraphicsState{
 			multiview_mask:None,
 			cache:None,
 		});
-		let debug_model_pipeline=device.create_render_pipeline(&wgpu::RenderPipelineDescriptor{
-			label:Some("Debug Model Pipeline"),
+		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"),
+				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,1=>Float32x3],
+					attributes:&wgpu::vertex_attr_array![0=>Float32x3],
 				}],
 				compilation_options:wgpu::PipelineCompilationOptions::default(),
 			},
@@ -961,6 +1015,7 @@ impl GraphicsState{
 				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()
@@ -973,9 +1028,55 @@ impl GraphicsState{
 				bias:wgpu::DepthBiasState::default(),
 			}),
 			multisample:wgpu::MultisampleState::default(),
-			multiview:None,
+			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);
@@ -1016,7 +1117,9 @@ impl GraphicsState{
 			pipelines:GraphicsPipelines{
 				skybox:sky_pipeline,
 				model:model_pipeline,
-				debug:debug_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,
@@ -1086,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();
 
 		{
@@ -1135,15 +1281,33 @@ impl GraphicsState{
 			rpass.draw(0..3,0..1);
 
 			// render a single debug_model in red
-			if let Some(model_id)=frame_state.debug_model{
-				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_pipeline(&self.pipelines.debug);
-					rpass.set_bind_group(1,&model.bind_group,&[]);
-					rpass.set_vertex_buffer(0,mesh.vertex_buf.slice(..));
-					rpass.set_index_buffer(mesh.indices.buf.slice(..),mesh.indices.format);
-					//TODO: loop over triangle strips
-					rpass.draw_indexed(0..mesh.indices.count,0,0..1);
+			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);
+							},
+						}
+					}
 				}
 			}
 		}

engine/graphics/src/model.rs

@@ -12,7 +12,6 @@ pub struct GraphicsVertex{
 #[repr(C)]
 pub struct DebugGraphicsVertex{
 	pub pos:[f32;3],
-	pub normal:[f32;3],
 }
 #[derive(Clone,Copy,id::Id)]
 pub struct IndexedGraphicsMeshOwnedRenderConfigId(u32);

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;

engine/physics/src/physics.rs

@@ -670,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())
@@ -682,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())
@@ -694,7 +694,7 @@ impl From<ModelId> for IntersectModelId{
 	}
 }
 #[derive(Debug,Clone,Copy,Hash,Eq,PartialEq)]
-enum PhysicsModelId{
+pub enum PhysicsModelId{
 	Contact(ContactModelId),
 	Intersect(IntersectModelId),
 }
@@ -708,9 +708,9 @@ impl From<PhysicsModelId> for ModelId{
 }
 //unique physics meshes indexed by this
 #[derive(Debug,Clone,Copy,Eq,Hash,PartialEq)]
-struct ConvexMeshId<Id>{
-	model_id:Id,
-	submesh_id:PhysicsSubmeshId,
+pub struct ConvexMeshId<Id>{
+	pub model_id:Id,
+	pub submesh_id:PhysicsSubmeshId,
 }
 impl<Id> ConvexMeshId<Id>{
 	fn map<NewId>(self,model_id:NewId)->ConvexMeshId<NewId>{
@@ -999,7 +999,7 @@ impl PhysicsData{
 			hitbox_mesh:StyleModifiers::default().calculate_mesh(),
 		}
 	}
-	pub fn trace_ray(&self,ray:strafesnet_common::ray::Ray)->Option<ModelId>{
+	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
@@ -1036,7 +1036,32 @@ impl PhysicsData{
 			}
 			min.map(Into::into)
 		})?;
-		Some(convex_mesh_id.model_id.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();

engine/session/src/session.rs

@@ -2,7 +2,6 @@ use std::collections::HashMap;
 
 use strafesnet_common::gameplay_modes::{ModeId,StageId};
 use strafesnet_common::instruction::{InstructionConsumer,InstructionEmitter,InstructionFeedback,TimedInstruction};
-use strafesnet_common::model::ModelId;
 // session represents the non-hardware state of the client.
 // Ideally it is a deterministic state which is atomically updated by instructions, same as the simulation state.
 use strafesnet_common::physics::{
@@ -62,7 +61,7 @@ pub struct FrameState{
 	pub body:physics::Body,
 	pub camera:physics::PhysicsCamera,
 	pub time:PhysicsTime,
-	pub debug_model:Option<ModelId>,
+	pub hit:Option<Hit>,
 }
 
 pub struct Simulation{
@@ -79,12 +78,12 @@ impl Simulation{
 			physics,
 		}
 	}
-	pub fn get_frame_state(&self,time:SessionTime,debug_model:Option<ModelId>)->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),
-			debug_model,
+			hit: debug_model,
 		}
 	}
 }
@@ -152,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,
@@ -164,7 +169,7 @@ pub struct Session{
 	recording:Recording,
 	//players:HashMap<PlayerId,Simulation>,
 	replays:HashMap<BotId,Replay>,
-	last_ray_hit:Option<ModelId>,
+	last_ray_hit:Option<Hit>,
 }
 impl Session{
 	pub fn new(
@@ -193,7 +198,7 @@ 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,self.last_ray_hit)),
+			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,None)
 			),
@@ -205,10 +210,15 @@ impl Session{
 				origin:frame_state.body.extrapolated_position(self.simulation.timer.time(time)),
 				direction:-frame_state.camera.rotation().z_axis,
 			};
-			let model_id=self.geometry_shared.trace_ray(ray);
-			if model_id!=self.last_ray_hit{
-				println!("hit={model_id:?}");
-				self.last_ray_hit=model_id;
+			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,
 			}
 		}
 	}

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>,
@@ -90,25 +92,162 @@ fn vs_entity_texture(
 @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>,
-	@location(1) normal: vec3<f32>,
-	@location(2) view: vec3<f32>,
 };
 
 @vertex
-fn vs_debug(
+fn vs_debug_face(
 	@location(0) pos: vec3<f32>,
-	@location(1) normal: vec3<f32>,
 ) -> DebugEntityOutput {
 	var position: vec4<f32> = model_instance.transform * vec4<f32>(pos, 1.0);
 	var result: DebugEntityOutput;
-	result.normal = model_instance.normal_transform * normal;
-	result.view = position.xyz - camera.view_inv[3].xyz;//col(3)
 	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)