project ideal point

This reverts commit 0ddd41c1b6.

engine/graphics/Cargo.toml

@@ -9,9 +9,10 @@ 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 = "27.0.0"
+wgpu = "28.0.0"
 
 [lints]
 workspace = true

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"),
@@ -616,7 +768,7 @@ impl GraphicsState{
 			address_mode_w:wgpu::AddressMode::ClampToEdge,
 			mag_filter:wgpu::FilterMode::Linear,
 			min_filter:wgpu::FilterMode::Linear,
-			mipmap_filter:wgpu::FilterMode::Linear,
+			mipmap_filter:wgpu::MipmapFilterMode::Linear,
 			..Default::default()
 		});
 		let repeat_sampler=device.create_sampler(&wgpu::SamplerDescriptor{
@@ -626,7 +778,7 @@ impl GraphicsState{
 			address_mode_w:wgpu::AddressMode::Repeat,
 			mag_filter:wgpu::FilterMode::Linear,
 			min_filter:wgpu::FilterMode::Linear,
-			mipmap_filter:wgpu::FilterMode::Linear,
+			mipmap_filter:wgpu::MipmapFilterMode::Linear,
 			anisotropy_clamp:16,
 			..Default::default()
 		});
@@ -754,15 +906,23 @@ impl GraphicsState{
 				&skybox_texture_bind_group_layout,
 				&model_bind_group_layout,
 			],
-			push_constant_ranges:&[],
+			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:&[
 				&camera_bind_group_layout,
 				&skybox_texture_bind_group_layout,
 			],
-			push_constant_ranges:&[],
+			immediate_size:0,
 		});
 
 		// Create the render pipelines
@@ -793,7 +953,7 @@ impl GraphicsState{
 				bias:wgpu::DepthBiasState::default(),
 			}),
 			multisample:wgpu::MultisampleState::default(),
-			multiview:None,
+			multiview_mask:None,
 			cache:None,
 		});
 		let model_pipeline=device.create_render_pipeline(&wgpu::RenderPipelineDescriptor{
@@ -828,9 +988,95 @@ impl GraphicsState{
 				bias:wgpu::DepthBiasState::default(),
 			}),
 			multisample:wgpu::MultisampleState::default(),
-			multiview:None,
+			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(0x100),
-			bind_group_layouts:GraphicsBindGroupLayouts{model:model_bind_group_layout},
+			staging_belt:wgpu::util::StagingBelt::new(device.clone(),0x100),
+			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,
 		}
@@ -918,7 +1172,6 @@ impl GraphicsState{
 				&self.camera_buf,
 				0,
 				wgpu::BufferSize::new((camera_uniforms.len() * 4) as wgpu::BufferAddress).unwrap(),
-				device,
 			)
 			.copy_from_slice(bytemuck::cast_slice(&camera_uniforms));
 		//This code only needs to run when the uniforms change
@@ -936,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();
 
 		{
@@ -965,11 +1261,13 @@ impl GraphicsState{
 				}),
 				timestamp_writes:Default::default(),
 				occlusion_query_set:Default::default(),
+				multiview_mask:None,
 			});
 
 			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/body.rs

@@ -33,7 +33,7 @@ impl<T:Copy> std::ops::Neg for &Body<T>{
 impl<T> Body<T>
 	where Time<T>:Copy,
 {
-	pub const ZERO:Self=Self::new(vec3::ZERO,vec3::ZERO,vec3::ZERO,Time::ZERO);
+	pub const ZERO:Self=Self::new(vec3::zero(),vec3::zero(),vec3::zero(),Time::ZERO);
 	pub const fn new(position:Planar64Vec3,velocity:Planar64Vec3,acceleration:Planar64Vec3,time:Time<T>)->Self{
 		Self{
 			position,
@@ -107,8 +107,8 @@ impl<T> Body<T>
 		self.time+=dt.into();
 	}
 	pub fn infinity_dir(&self)->Option<Planar64Vec3>{
-		if self.velocity==vec3::ZERO{
-			if self.acceleration==vec3::ZERO{
+		if self.velocity==vec3::zero(){
+			if self.acceleration==vec3::zero(){
 				None
 			}else{
 				Some(self.acceleration)

engine/physics/src/face_crawler.rs

@@ -21,12 +21,6 @@ impl<M:MeshQuery> CrawlResult<M>{
 			CrawlResult::Hit(face,time)=>Some((face,time)),
 		}
 	}
-	pub fn miss(self)->Option<FEV<M>>{
-		match self{
-			CrawlResult::Miss(fev)=>Some(fev),
-			CrawlResult::Hit(_,_)=>None,
-		}
-	}
 }
 
 // TODO: move predict_collision_face_out algorithm in here or something

engine/physics/src/lib.rs

@@ -1,7 +1,8 @@
 mod body;
-mod push_solve;
 mod face_crawler;
-mod model;
+pub mod model;
+mod push_solve;
+mod minimum_difference;
 
 pub mod physics;
 

engine/physics/src/minimum_difference.rs

@@ -0,0 +1,913 @@
+use strafesnet_common::integer::vec3;
+use strafesnet_common::integer::vec3::Vector3;
+use strafesnet_common::integer::{Fixed,Planar64,Planar64Vec3};
+
+use crate::model::{DirectedEdge,FEV,MeshQuery};
+// TODO: remove mesh invert
+use crate::model::{MinkowskiMesh,MinkowskiVert};
+
+// This algorithm is based on Lua code
+// written by Trey Reynolds in 2021
+
+type Simplex<const N:usize,Vert>=[Vert;N];
+#[derive(Clone,Copy)]
+enum Simplex1_3<Vert>{
+	Simplex1(Simplex<1,Vert>),
+	Simplex2(Simplex<2,Vert>),
+	Simplex3(Simplex<3,Vert>),
+}
+impl<Vert> Simplex1_3<Vert>{
+	fn push_front(self,v:Vert)->Simplex2_4<Vert>{
+		match self{
+			Simplex1_3::Simplex1([v0])=>Simplex2_4::Simplex2([v,v0]),
+			Simplex1_3::Simplex2([v0,v1])=>Simplex2_4::Simplex3([v,v0,v1]),
+			Simplex1_3::Simplex3([v0,v1,v2])=>Simplex2_4::Simplex4([v,v0,v1,v2]),
+		}
+	}
+}
+#[derive(Clone,Copy)]
+enum Simplex2_4<Vert>{
+	Simplex2(Simplex<2,Vert>),
+	Simplex3(Simplex<3,Vert>),
+	Simplex4(Simplex<4,Vert>),
+}
+
+/*
+local function absDet(r, u, v, w)
+	if w then
+		return math.abs((u - r):Cross(v - r):Dot(w - r))
+	elseif v then
+		return (u - r):Cross(v - r).magnitude
+	elseif u then
+		return (u - r).magnitude
+	else
+		return 1
+	end
+end
+*/
+impl<Vert> Simplex2_4<Vert>{
+	fn det_is_zero<M:MeshQuery<Vert=Vert>>(self,mesh:&M)->bool{
+		match self{
+			Self::Simplex4([p0,p1,p2,p3])=>{
+				let p0=mesh.vert(p0);
+				let p1=mesh.vert(p1);
+				let p2=mesh.vert(p2);
+				let p3=mesh.vert(p3);
+				(p1-p0).cross(p2-p0).dot(p3-p0)==Fixed::ZERO
+			},
+			Self::Simplex3([p0,p1,p2])=>{
+				let p0=mesh.vert(p0);
+				let p1=mesh.vert(p1);
+				let p2=mesh.vert(p2);
+				(p1-p0).cross(p2-p0)==vec3::zero()
+			},
+			Self::Simplex2([p0,p1])=>{
+				let p0=mesh.vert(p0);
+				let p1=mesh.vert(p1);
+				p1-p0==vec3::zero()
+			}
+		}
+	}
+}
+
+/*
+local function choosePerpendicularDirection(d)
+	local x, y, z = d.x, d.y, d.z
+	local best = math.min(x*x, y*y, z*z)
+	if x*x == best then
+		return Vector3.new(y*y + z*z, -x*y, -x*z)
+	elseif y*y == best then
+		return Vector3.new(-x*y, x*x + z*z, -y*z)
+	else
+		return Vector3.new(-x*z, -y*z, x*x + y*y)
+	end
+end
+*/
+fn choose_perpendicular_direction(d:Planar64Vec3)->Planar64Vec3{
+	let x=d.x.abs();
+	let y=d.y.abs();
+	let z=d.z.abs();
+	if x<y&&x<z{
+		Vector3::new([Fixed::ZERO,-d.z,d.y])
+	}else if y<z{
+		Vector3::new([d.z,Fixed::ZERO,-d.x])
+	}else{
+		Vector3::new([-d.y,d.x,Fixed::ZERO])
+	}
+}
+
+const fn choose_any_direction()->Planar64Vec3{
+	vec3::X
+}
+
+fn narrow_dir2(dir:Vector3<Fixed<2,64>>)->Planar64Vec3{
+	if dir==vec3::zero(){
+		return dir.narrow_1().unwrap();
+	}
+	let x=dir.x.as_bits().unsigned_abs().bits();
+	let y=dir.y.as_bits().unsigned_abs().bits();
+	let z=dir.z.as_bits().unsigned_abs().bits();
+	let big=x.max(y).max(z);
+	const MAX_BITS:u32=64+31;
+	if MAX_BITS<big{
+		dir>>(big-MAX_BITS)
+	}else{
+		dir
+	}.narrow_1().unwrap()
+}
+fn narrow_dir3(dir:Vector3<Fixed<3,96>>)->Planar64Vec3{
+	if dir==vec3::zero(){
+		return dir.narrow_1().unwrap();
+	}
+	let x=dir.x.as_bits().unsigned_abs().bits();
+	let y=dir.y.as_bits().unsigned_abs().bits();
+	let z=dir.z.as_bits().unsigned_abs().bits();
+	let big=x.max(y).max(z);
+	const MAX_BITS:u32=96+31;
+	if MAX_BITS<big{
+		dir>>(big-MAX_BITS)
+	}else{
+		dir
+	}.narrow_1().unwrap()
+}
+
+fn reduce1<M:MeshQuery>(
+	[v0]:Simplex<1,M::Vert>,
+	mesh:&M,
+	point:Planar64Vec3,
+)->Reduced<M::Vert>{
+	// --debug.profilebegin("reduceSimplex0")
+	// local a = a1 - a0
+	let p0=mesh.vert(v0);
+
+	// local p = -a
+	let p=-(p0+point);
+
+	// local direction = p
+	let mut dir=p;
+
+	// if direction.magnitude == 0 then
+	// 	direction = chooseAnyDirection()
+	if dir==vec3::zero(){
+		dir=choose_any_direction();
+	}
+
+	// return direction, a0, a1
+	Reduced{
+		dir,
+		simplex:Simplex1_3::Simplex1([v0]),
+	}
+}
+
+// local function reduceSimplex1(a0, a1, b0, b1)
+fn reduce2<M:MeshQuery>(
+	[v0,v1]:Simplex<2,M::Vert>,
+	mesh:&M,
+	point:Planar64Vec3,
+)->Reduced<M::Vert>{
+	// --debug.profilebegin("reduceSimplex1")
+	// local a = a1 - a0
+	// local b = b1 - b0
+	let p0=mesh.vert(v0);
+	let p1=mesh.vert(v1);
+
+	// local p = -a
+	// local u = b - a
+	let p=-(p0+point);
+	let u=p1-p0;
+
+	// -- modify to take into account the radiuses
+	// local p_u = p:Dot(u)
+	let p_u=p.dot(u);
+
+	// if p_u >= 0 then
+	if !p_u.is_negative(){
+		// local direction = u:Cross(p):Cross(u)
+		let direction=u.cross(p).cross(u);
+
+		// if direction.magnitude == 0 then
+		if direction==vec3::zero(){
+			return Reduced{
+				dir:choose_perpendicular_direction(u),
+				simplex:Simplex1_3::Simplex2([v0,v1]),
+			};
+		}
+
+		// -- modify the direction to take into account a0R and b0R
+		// return direction, a0, a1, b0, b1
+		return Reduced{
+			dir:narrow_dir3(direction),
+			simplex:Simplex1_3::Simplex2([v0,v1]),
+		};
+	}
+
+	// local direction = p
+	let mut dir=p;
+
+	// if direction.magnitude == 0 then
+	if dir==vec3::zero(){
+		dir=choose_perpendicular_direction(u);
+	}
+
+	// return direction, a0, a1
+	Reduced{
+		dir,
+		simplex:Simplex1_3::Simplex1([v0]),
+	}
+}
+
+// local function reduceSimplex2(a0, a1, b0, b1, c0, c1)
+fn reduce3<M:MeshQuery>(
+	[v0,mut v1,v2]:Simplex<3,M::Vert>,
+	mesh:&M,
+	point:Planar64Vec3,
+)->Reduced<M::Vert>{
+	// --debug.profilebegin("reduceSimplex2")
+	// local a = a1 - a0
+	// local b = b1 - b0
+	// local c = c1 - c0
+	let p0=mesh.vert(v0);
+	let p1=mesh.vert(v1);
+	let p2=mesh.vert(v2);
+
+	// local p = -a
+	// local u = b - a
+	// local v = c - a
+	let p=-(p0+point);
+	let mut u=p1-p0;
+	let v=p2-p0;
+
+	// local uv = u:Cross(v)
+	// local up = u:Cross(p)
+	// local pv = p:Cross(v)
+	// local uv_up = uv:Dot(up)
+	// local uv_pv = uv:Dot(pv)
+	let mut uv=u.cross(v);
+	let mut up=u.cross(p);
+	let pv=p.cross(v);
+	let uv_up=uv.dot(up);
+	let uv_pv=uv.dot(pv);
+
+	// if uv_up >= 0 and uv_pv >= 0 then
+	if !uv_up.is_negative()&&!uv_pv.is_negative(){
+		// local uvp = uv:Dot(p)
+		let uvp=uv.dot(p);
+
+		// local direction = uvp < 0 and -uv or uv
+		let direction=if uvp.is_negative(){
+			-uv
+		}else{
+			uv
+		};
+
+		// return direction, a0, a1, b0, b1, c0, c1
+		return Reduced{
+			dir:narrow_dir2(direction),
+			simplex:Simplex1_3::Simplex3([v0,v1,v2]),
+		};
+	}
+
+	// local u_u = u:Dot(u)
+	// local v_v = v:Dot(v)
+	// local uDist = uv_up/(u_u*v.magnitude)
+	// local vDist = uv_pv/(v_v*u.magnitude)
+	// local minDist2 = math.min(uDist, vDist)
+	let u_dist=uv_up*v.length();
+	let v_dist=uv_pv*u.length();
+
+	// if vDist == minDist2 then
+	if v_dist<u_dist{
+		u=v;
+		up=-pv;
+		uv=-uv;
+		// b0 = c0
+		// b1 = c1
+		v1=v2;
+	}
+
+	// local p_u = p:Dot(u)
+	let p_u=p.dot(u);
+
+	// if p_u >= 0 then
+	if !p_u.is_negative(){
+		// local direction = up:Cross(u)
+		let direction=up.cross(u);
+		// if direction.magnitude == 0 then
+		if direction==vec3::zero(){
+			// direction = uv
+			return Reduced{
+				dir:narrow_dir2(uv),
+				simplex:Simplex1_3::Simplex2([v0,v1]),
+			};
+		}
+
+		// return direction, a0, a1, b0, b1
+		return Reduced{
+			dir:narrow_dir3(direction),
+			simplex:Simplex1_3::Simplex2([v0,v1]),
+		};
+	}
+
+	// local direction = p
+	let dir=p;
+	// if direction.magnitude == 0 then
+	if dir==vec3::zero(){
+		// direction = uv
+		return Reduced{
+			dir:narrow_dir2(uv),
+			simplex:Simplex1_3::Simplex1([v0]),
+		};
+	}
+	// return direction, a0, a0
+	Reduced{
+		dir,
+		simplex:Simplex1_3::Simplex1([v0]),
+	}
+}
+
+// local function reduceSimplex3(a0, a1, b0, b1, c0, c1, d0, d1)
+fn reduce4<M:MeshQuery>(
+	[v0,mut v1,mut v2,v3]:Simplex<4,M::Vert>,
+	mesh:&M,
+	point:Planar64Vec3,
+)->Reduce<M::Vert>{
+	// --debug.profilebegin("reduceSimplex3")
+	// local a = a1 - a0
+	// local b = b1 - b0
+	// local c = c1 - c0
+	// local d = d1 - d0
+	let p0=mesh.vert(v0);
+	let p1=mesh.vert(v1);
+	let p2=mesh.vert(v2);
+	let p3=mesh.vert(v3);
+
+	// local p = -a
+	// local u = b - a
+	// local v = c - a
+	// local w = d - a
+	let p=-(p0+point);
+	let mut u=p1-p0;
+	let mut v=p2-p0;
+	let w=p3-p0;
+
+	// local uv = u:Cross(v)
+	// local vw = v:Cross(w)
+	// local wu = w:Cross(u)
+	// local uvw = uv:Dot(w)
+	// local pvw = vw:Dot(p)
+	// local upw = wu:Dot(p)
+	// local uvp = uv:Dot(p)
+	let mut uv=u.cross(v);
+	let vw=v.cross(w);
+	let wu=w.cross(u);
+	let uv_w=uv.dot(w);
+	let pv_w=vw.dot(p);
+	let up_w=wu.dot(p);
+	let uv_p=uv.dot(p);
+
+	// if pvw/uvw >= 0 and upw/uvw >= 0 and uvp/uvw >= 0 then
+	if !pv_w.div_sign(uv_w).is_negative()
+	 ||!up_w.div_sign(uv_w).is_negative()
+	 ||!uv_p.div_sign(uv_w).is_negative(){
+		// origin is contained, this is a positive detection
+		// local direction = Vector3.new(0, 0, 0)
+		// return direction, a0, a1, b0, b1, c0, c1, d0, d1
+		return Reduce::Escape([v0,v1,v2,v3]);
+	}
+
+	// local uvwSign = uvw < 0 and -1 or uvw > 0 and 1 or 0
+	// local uvDist = uvp*uvwSign/uv.magnitude
+	// local vwDist = pvw*uvwSign/vw.magnitude
+	// local wuDist = upw*uvwSign/wu.magnitude
+	// local minDist3 = math.min(uvDist, vwDist, wuDist)
+	let uv_dist=uv_p.mul_sign(uv_w);
+	let vw_dist=pv_w.mul_sign(uv_w);
+	let wu_dist=up_w.mul_sign(uv_w);
+	let wu_len=wu.length();
+	let uv_len=uv.length();
+	let vw_len=vw.length();
+
+	if vw_dist*wu_len<wu_dist*vw_len{
+		// if vwDist == minDist3 then
+		if vw_dist*uv_len<uv_dist*vw_len{
+			(u,v)=(v,w);
+			uv=vw;
+			// uv_p=pv_w; // unused
+			// b0, c0 = c0, d0
+			// b1, c1 = c1, d1
+			(v1,v2)=(v2,v3);
+		}else{
+			v2=v3;
+		}
+	}else{
+		// elseif wuDist == minDist3 then
+		if wu_dist*uv_len<uv_dist*wu_len{
+			(u,v)=(w,u);
+			uv=wu;
+			// uv_p=up_w; // unused
+			// b0, c0 = d0, b0
+			// b1, c1 = d1, b1
+			// before [a,b,c,d]
+			(v1,v2)=(v3,v1);
+			// after [a,d,b]
+		}else{
+			v2=v3;
+		}
+	}
+
+	// local up = u:Cross(p)
+	// local pv = p:Cross(v)
+	// local uv_up = uv:Dot(up)
+	// local uv_pv = uv:Dot(pv)
+	let mut up=u.cross(p);
+	let pv=p.cross(v);
+	let uv_up=uv.dot(up);
+	let uv_pv=uv.dot(pv);
+
+	// if uv_up >= 0 and uv_pv >= 0 then
+	if !uv_up.is_negative()&&!uv_pv.is_negative(){
+		// local direction = uvw < 0 and uv or -uv
+		// return direction, a0, a1, b0, b1, c0, c1
+		if uv_w.is_negative(){
+			return Reduce::Reduced(Reduced{
+				dir:narrow_dir2(uv),
+				simplex:Simplex1_3::Simplex3([v0,v1,v2]),
+			});
+		}else{
+			return Reduce::Reduced(Reduced{
+				dir:narrow_dir2(-uv),
+				simplex:Simplex1_3::Simplex3([v0,v1,v2]),
+			});
+		}
+	}
+
+	// local u_u = u:Dot(u)
+	// local v_v = v:Dot(v)
+	// local uDist = uv_up/(u_u*v.magnitude)
+	// local vDist = uv_pv/(v_v*u.magnitude)
+	// local minDist2 = math.min(uDist, vDist)
+	let u_dist=uv_up*v.length();
+	let v_dist=uv_pv*u.length();
+
+	// if vDist == minDist2 then
+	if v_dist<u_dist{
+		u=v;
+		up=-pv;
+		uv=-uv;
+		// b0 = c0
+		// b1 = c1
+		v1=v2;
+	}
+
+	// local p_u = p:Dot(u)
+	let p_u=p.dot(u);
+
+	// if p_u >= 0 then
+	if !p_u.is_negative(){
+		// local direction = up:Cross(u)
+		let direction=up.cross(u);
+		// if direction.magnitude == 0 then
+		if direction==vec3::zero(){
+			// direction = uvw < 0 and uv or -uv
+			// return direction, a0, a1, b0, b1
+			if uv_w.is_negative(){
+				return Reduce::Reduced(Reduced{
+					dir:narrow_dir2(uv),
+					simplex:Simplex1_3::Simplex2([v0,v1]),
+				});
+			}else{
+				return Reduce::Reduced(Reduced{
+					dir:narrow_dir2(-uv),
+					simplex:Simplex1_3::Simplex2([v0,v1]),
+				});
+			}
+		}
+
+		// return direction, a0, a1, b0, b1
+		return Reduce::Reduced(Reduced{
+			dir:narrow_dir3(direction),
+			simplex:Simplex1_3::Simplex2([v0,v1]),
+		});
+	}
+
+	// local direction = p
+	let dir=p;
+	// if direction.magnitude == 0 then
+	if dir==vec3::zero(){
+		// direction = uvw < 0 and uv or -uv
+		if uv_w.is_negative(){
+			return Reduce::Reduced(Reduced{
+				dir:narrow_dir2(uv),
+				simplex:Simplex1_3::Simplex1([v0]),
+			});
+		}else{
+			return Reduce::Reduced(Reduced{
+				dir:narrow_dir2(-uv),
+				simplex:Simplex1_3::Simplex1([v0]),
+			});
+		}
+	}
+
+	// return direction, a0, a1
+	Reduce::Reduced(Reduced{
+		dir,
+		simplex:Simplex1_3::Simplex1([v0]),
+	})
+}
+
+struct Reduced<Vert>{
+	dir:Planar64Vec3,
+	simplex:Simplex1_3<Vert>,
+}
+
+enum Reduce<Vert>{
+	Escape(Simplex<4,Vert>),
+	Reduced(Reduced<Vert>),
+}
+
+impl<Vert> Simplex2_4<Vert>{
+	fn reduce<M:MeshQuery<Vert=Vert>>(self,mesh:&M,point:Planar64Vec3)->Reduce<Vert>{
+		match self{
+			Self::Simplex2(simplex)=>Reduce::Reduced(reduce2(simplex,mesh,point)),
+			Self::Simplex3(simplex)=>Reduce::Reduced(reduce3(simplex,mesh,point)),
+			Self::Simplex4(simplex)=>reduce4(simplex,mesh,point),
+		}
+	}
+}
+
+pub fn contains_point(mesh:&MinkowskiMesh<'_>,point:Planar64Vec3)->bool{
+	const ENABLE_FAST_FAIL:bool=true;
+	// TODO: remove mesh negation
+	minimum_difference::<ENABLE_FAST_FAIL,_,_>(&-mesh,point,
+		// on_exact
+		|is_intersecting,_simplex|{
+			is_intersecting
+		},
+		// on_escape
+		|_simplex|{
+			// intersection is guaranteed at this point
+			true
+		},
+		// fast_fail value
+		||false
+	)
+}
+
+//infinity fev algorithm state transition
+#[derive(Debug)]
+enum Transition<Vert>{
+	Done,//found closest vert, no edges are better
+	Vert(Vert),//transition to vert
+}
+enum EV<M:MeshQuery>{
+	Vert(M::Vert),
+	Edge(<M::Edge as DirectedEdge>::UndirectedEdge),
+}
+impl<M:MeshQuery> From<EV<M>> for FEV<M>{
+	fn from(value:EV<M>)->Self{
+		match value{
+			EV::Vert(minkowski_vert)=>FEV::Vert(minkowski_vert),
+			EV::Edge(minkowski_edge)=>FEV::Edge(minkowski_edge),
+		}
+	}
+}
+
+trait Contains{
+	fn contains(&self,point:Planar64Vec3)->bool;
+}
+
+// convenience type to check if a point is within some threshold of a plane.
+struct ThickPlane{
+	point:Planar64Vec3,
+	normal:Vector3<Fixed<2,64>>,
+	epsilon:Fixed<3,96>,
+}
+impl ThickPlane{
+	fn new<M:MeshQuery>(mesh:&M,[v0,v1,v2]:Simplex<3,M::Vert>)->Self{
+		let p0=mesh.vert(v0);
+		let p1=mesh.vert(v1);
+		let p2=mesh.vert(v2);
+		let point=p0;
+		let normal=(p1-p0).cross(p2-p0);
+		// Allow ~ 2*sqrt(3) units of thickness on the plane
+		// This is to account for the variance of two voxels across the longest diagonal
+		let epsilon=(normal.length()*(Planar64::EPSILON*3)).wrap_3();
+		Self{point,normal,epsilon}
+	}
+}
+impl Contains for ThickPlane{
+	fn contains(&self,point:Planar64Vec3)->bool{
+		(point-self.point).dot(self.normal).abs()<=self.epsilon
+	}
+}
+
+struct ThickLine{
+	point:Planar64Vec3,
+	dir:Planar64Vec3,
+	epsilon:Fixed<4,128>,
+}
+impl ThickLine{
+	fn new<M:MeshQuery>(mesh:&M,[v0,v1]:Simplex<2,M::Vert>)->Self{
+		let p0=mesh.vert(v0);
+		let p1=mesh.vert(v1);
+		let point=p0;
+		let dir=p1-p0;
+		// Allow ~ 2*sqrt(3) units of thickness on the plane
+		// This is to account for the variance of two voxels across the longest diagonal
+		let epsilon=(dir.length_squared()*(Planar64::EPSILON*3)).widen_4();
+		Self{point,dir,epsilon}
+	}
+}
+impl Contains for ThickLine{
+	fn contains(&self,point:Planar64Vec3)->bool{
+		(point-self.point).cross(self.dir).length_squared()<=self.epsilon
+	}
+}
+
+struct EVFinder<'a,M,C>{
+	mesh:&'a M,
+	constraint:C,
+	best_distance_squared:Fixed<2,64>,
+}
+
+impl<M:MeshQuery,C:Contains> EVFinder<'_,M,C>{
+	fn next_transition_vert(&mut self,vert_id:M::Vert,point:Planar64Vec3)->Transition<M::Vert>{
+		let mut best_transition=Transition::Done;
+		for &directed_edge_id in self.mesh.vert_edges(vert_id).as_ref(){
+			//test if this edge's opposite vertex closer
+			let edge_verts=self.mesh.edge_verts(directed_edge_id.as_undirected());
+			//select opposite vertex
+			let test_vert_id=edge_verts.as_ref()[directed_edge_id.parity() as usize];
+			let test_pos=self.mesh.vert(test_vert_id);
+			let diff=point-test_pos;
+			let distance_squared=diff.dot(diff);
+			// ensure test_vert_id is coplanar to simplex
+			if distance_squared<self.best_distance_squared&&self.constraint.contains(test_pos){
+				best_transition=Transition::Vert(test_vert_id);
+				self.best_distance_squared=distance_squared;
+			}
+		}
+		best_transition
+	}
+	fn final_ev(&mut self,vert_id:M::Vert,point:Planar64Vec3)->EV<M>{
+		let mut best_transition=EV::Vert(vert_id);
+		let vert_pos=self.mesh.vert(vert_id);
+		let diff=point-vert_pos;
+		for &directed_edge_id in self.mesh.vert_edges(vert_id).as_ref(){
+			//test if this edge is closer
+			let edge_verts=self.mesh.edge_verts(directed_edge_id.as_undirected());
+			let test_vert_id=edge_verts.as_ref()[directed_edge_id.parity() as usize];
+			let test_pos=self.mesh.vert(test_vert_id);
+			let edge_n=test_pos-vert_pos;
+			let d=edge_n.dot(diff);
+			//test the edge
+			let edge_nn=edge_n.dot(edge_n);
+			// ensure edge contains closest point and directed_edge_id is coplanar to simplex
+			if !d.is_negative()&&d<=edge_nn&&self.constraint.contains(test_pos){
+				let distance_squared={
+					let c=diff.cross(edge_n);
+					//wrap for speed
+					(c.dot(c)/edge_nn).divide().wrap_2()
+				};
+				if distance_squared<=self.best_distance_squared{
+					best_transition=EV::Edge(directed_edge_id.as_undirected());
+					self.best_distance_squared=distance_squared;
+				}
+			}
+		}
+		best_transition
+	}
+	fn crawl_boundaries(&mut self,mut vert_id:M::Vert,point:Planar64Vec3)->EV<M>{
+		loop{
+			match self.next_transition_vert(vert_id,point){
+				Transition::Done=>return self.final_ev(vert_id,point),
+				Transition::Vert(new_vert_id)=>vert_id=new_vert_id,
+			}
+		}
+	}
+}
+/// This function drops a vertex down to an edge or a face if the path from infinity did not cross any vertex-edge boundaries but the point is supposed to have already crossed a boundary down from a vertex
+fn crawl_to_closest_ev<M:MeshQuery>(mesh:&M,simplex:Simplex<2,M::Vert>,point:Planar64Vec3)->EV<M>{
+	// naively start at the closest vertex
+	// the closest vertex is not necessarily the one with the fewest boundary hops
+	// but it doesn't matter, we will get there regardless.
+	let (vert_id,best_distance_squared)=simplex.into_iter().map(|vert_id|{
+		let diff=point-mesh.vert(vert_id);
+		(vert_id,diff.dot(diff))
+	}).min_by_key(|&(_,d)|d).unwrap();
+
+	let constraint=ThickLine::new(mesh,simplex);
+	let mut finder=EVFinder{constraint,mesh,best_distance_squared};
+	//start on any vertex
+	//cross uncrossable vertex-edge boundaries until you find the closest vertex or edge
+	//cross edge-face boundary if it's uncrossable
+	finder.crawl_boundaries(vert_id,point)
+}
+
+/// This function drops a vertex down to an edge or a face if the path from infinity did not cross any vertex-edge boundaries but the point is supposed to have already crossed a boundary down from a vertex
+fn crawl_to_closest_fev<'a>(mesh:&MinkowskiMesh<'a>,simplex:Simplex<3,MinkowskiVert>,point:Planar64Vec3)->FEV::<MinkowskiMesh<'a>>{
+	// naively start at the closest vertex
+	// the closest vertex is not necessarily the one with the fewest boundary hops
+	// but it doesn't matter, we will get there regardless.
+	let (vert_id,best_distance_squared)=simplex.into_iter().map(|vert_id|{
+		let diff=point-mesh.vert(vert_id);
+		(vert_id,diff.dot(diff))
+	}).min_by_key(|&(_,d)|d).unwrap();
+
+	let constraint=ThickPlane::new(mesh,simplex);
+	let mut finder=EVFinder{constraint,mesh,best_distance_squared};
+	//start on any vertex
+	//cross uncrossable vertex-edge boundaries until you find the closest vertex or edge
+	//cross edge-face boundary if it's uncrossable
+	match finder.crawl_boundaries(vert_id,point){
+		//if a vert is returned, it is the closest point to the infinity point
+		EV::Vert(vert_id)=>FEV::Vert(vert_id),
+		EV::Edge(edge_id)=>{
+			//cross to face if we are on the wrong side
+			let edge_n=mesh.edge_n(edge_id);
+			// point is multiplied by two because vert_sum sums two vertices.
+			let delta_pos=point*2-{
+				let &[v0,v1]=mesh.edge_verts(edge_id).as_ref();
+				mesh.vert(v0)+mesh.vert(v1)
+			};
+			for (i,&face_id) in mesh.edge_faces(edge_id).as_ref().iter().enumerate(){
+				//test if this face is closer
+				let (face_n,d)=mesh.face_nd(face_id);
+				//if test point is behind face, the face is invalid
+				// TODO: find out why I thought of this backwards
+				if !(face_n.dot(point)-d).is_positive(){
+					continue;
+				}
+				//edge-face boundary nd, n facing out of the face towards the edge
+				let boundary_n=face_n.cross(edge_n)*(i as i64*2-1);
+				let boundary_d=boundary_n.dot(delta_pos);
+				//is test point behind edge, i.e. contained in the face
+				if !boundary_d.is_positive(){
+					//both faces cannot pass this condition, return early if one does.
+					return FEV::Face(face_id);
+				}
+			}
+			FEV::Edge(edge_id)
+		},
+	}
+}
+
+pub fn closest_fev_not_inside<'a>(mesh:&MinkowskiMesh<'a>,point:Planar64Vec3)->Option<FEV<MinkowskiMesh<'a>>>{
+	const ENABLE_FAST_FAIL:bool=false;
+	// TODO: remove mesh negation
+	minimum_difference::<ENABLE_FAST_FAIL,_,_>(&-mesh,point,
+		// on_exact
+		|is_intersecting,simplex|{
+			if is_intersecting{
+				return None;
+			}
+			// Convert simplex to FEV
+			// Vertices must be inverted since the mesh is inverted
+			Some(match simplex{
+				Simplex1_3::Simplex1([v0])=>FEV::Vert(-v0),
+				Simplex1_3::Simplex2([v0,v1])=>{
+					// invert
+					let (v0,v1)=(-v0,-v1);
+					let ev=crawl_to_closest_ev(mesh,[v0,v1],point);
+					if !matches!(ev,EV::Edge(_)){
+						println!("I can't believe it's not an edge!");
+					}
+					ev.into()
+				},
+				Simplex1_3::Simplex3([v0,v1,v2])=>{
+					// invert
+					let (v0,v1,v2)=(-v0,-v1,-v2);
+					// Shimmy to the side until you find a face that contains the closest point
+					// it's ALWAYS representable as a face, but this algorithm may
+					// return E or V in edge cases but I don't think that will break the face crawler
+					let fev=crawl_to_closest_fev(mesh,[v0,v1,v2],point);
+					if !matches!(fev,FEV::Face(_)){
+						println!("I can't believe it's not a face!");
+					}
+					fev
+				},
+			})
+		},
+		// on_escape
+		|_simplex|{
+			// intersection is guaranteed at this point
+			// local norm, dist, u0, u1, v0, v1, w0, w1 = expand(queryP, queryQ, a0, a1, b0, b1, c0, c1, d0, d1, 1e-5)
+			// let simplex=refine_to_exact(mesh,simplex);
+			None
+		},
+		// fast_fail value is irrelevant and will never be returned!
+		||unreachable!()
+	)
+}
+
+// local function minimumDifference(
+// 	queryP, radiusP,
+// 	queryQ, radiusQ,
+// 	exitRadius, testIntersection
+// )
+fn minimum_difference<const ENABLE_FAST_FAIL:bool,T,M:MeshQuery>(
+	mesh:&M,
+	point:Planar64Vec3,
+	on_exact:impl FnOnce(bool,Simplex1_3<M::Vert>)->T,
+	on_escape:impl FnOnce(Simplex<4,M::Vert>)->T,
+	on_fast_fail:impl FnOnce()->T,
+)->T{
+	// local initialAxis = queryQ() - queryP()
+	// local new_point_p = queryP(initialAxis)
+	// local new_point_q = queryQ(-initialAxis)
+	// local direction, a0, a1, b0, b1, c0, c1, d0, d1
+	let mut initial_axis=mesh.hint_point()+point;
+	// degenerate case
+	if initial_axis==vec3::zero(){
+		initial_axis=choose_any_direction();
+	}
+	let last_point=mesh.farthest_vert(-initial_axis);
+	// this represents the 'a' value in the commented code
+	let mut last_pos=mesh.vert(last_point);
+	let Reduced{dir:mut direction,simplex:mut simplex_small}=reduce1([last_point],mesh,point);
+
+	// exitRadius = testIntersection and 0 or exitRadius or 1/0
+	// for _ = 1, 100 do
+	loop{
+		// new_point_p = queryP(-direction)
+		// new_point_q = queryQ(direction)
+		// local next_point = new_point_q - new_point_p
+		let next_point=mesh.farthest_vert(direction);
+		let next_pos=mesh.vert(next_point);
+
+		// if -direction:Dot(next_point) > (exitRadius + radiusP + radiusQ)*direction.magnitude then
+		if ENABLE_FAST_FAIL&&direction.dot(next_pos+point).is_negative(){
+			return on_fast_fail();
+		}
+
+		let simplex_big=simplex_small.push_front(next_point);
+
+		// if
+		// 	direction:Dot(next_point - a) <= 0 or
+		// 	absDet(next_point, a, b, c) < 1e-6
+		if !direction.dot(next_pos-last_pos).is_positive()
+		||simplex_big.det_is_zero(mesh){
+			// Found enough information to compute the exact closest point.
+			// local norm = direction.unit
+			// local dist = a:Dot(norm)
+			// local hits = -dist < radiusP + radiusQ
+			let is_intersecting=(last_pos+point).dot(direction).is_positive();
+			return on_exact(is_intersecting,simplex_small);
+		}
+
+		// direction, a0, a1, b0, b1, c0, c1, d0, d1 = reduceSimplex(new_point_p, new_point_q, a0, a1, b0, b1, c0, c1)
+		match simplex_big.reduce(mesh,point){
+			// if a and b and c and d then
+			Reduce::Escape(simplex)=>{
+				// Enough information to conclude that the meshes are intersecting.
+				// Topology information is computed if needed.
+				return on_escape(simplex);
+			},
+			Reduce::Reduced(reduced)=>{
+				direction=reduced.dir;
+				simplex_small=reduced.simplex;
+			},
+		}
+
+		// next loop this will be a
+		last_pos=next_pos;
+	}
+}
+
+#[cfg(test)]
+mod test{
+	use super::*;
+    use crate::model::{PhysicsMesh,PhysicsMeshView};
+
+    fn mesh_contains_point(mesh:PhysicsMeshView<'_>,point:Planar64Vec3)->bool{
+		const ENABLE_FAST_FAIL:bool=true;
+		// TODO: remove mesh negation
+		minimum_difference::<ENABLE_FAST_FAIL,_,_>(&mesh,point,
+			// on_exact
+			|is_intersecting,_simplex|{
+				is_intersecting
+			},
+			// on_escape
+			|_simplex|{
+				// intersection is guaranteed at this point
+				true
+			},
+			// fast_fail value
+			||false
+		)
+	}
+
+	#[test]
+	fn test_cube_points(){
+		let mesh=PhysicsMesh::unit_cube();
+		let mesh_view=mesh.complete_mesh_view();
+		for x in -2..=2{
+			for y in -2..=2{
+				for z in -2..=2{
+					let point=vec3::int(x,y,z)>>1;
+					assert!(mesh_contains_point(mesh_view,point),"Mesh did not contain point {point}");
+				}
+			}
+		}
+	}
+}

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),
@@ -90,16 +90,10 @@ pub trait MeshQuery{
 		let &[v0,v1]=self.edge_verts(directed_edge_id.as_undirected()).as_ref();
 		(self.vert(v1)-self.vert(v0))*((directed_edge_id.parity() as i64)*2-1)
 	}
-	/// Returns an iterator over the vertices in the direction of the directed edges.
-	/// Intended to be used to find adjacent vertices:
-	/// `self.directed_verts(self.vert_edges(vert_id).as_ref())`
-	/// TODO: rewrite this function as `adjacent_vertices`
-	fn directed_verts(&self,edges:&[Self::Edge])->impl Iterator<Item=Self::Vert>{
-		edges.iter().map(|e|{
-			let edge_verts=self.edge_verts(e.as_undirected());
-			edge_verts.as_ref()[e.parity() as usize]
-		})
-	}
+	/// This must return a point inside the mesh.
+	#[expect(dead_code)]
+	fn hint_point(&self)->Planar64Vec3;
+	fn farthest_vert(&self,dir:Planar64Vec3)->Self::Vert;
 	fn vert(&self,vert_id:Self::Vert)->Planar64Vec3;
 	fn face_nd(&self,face_id:Self::Face)->(Self::Normal,Self::Offset);
 	fn face_edges(&self,face_id:Self::Face)->impl AsRef<[Self::Edge]>;
@@ -110,17 +104,21 @@ pub trait MeshQuery{
 }
 #[derive(Debug)]
 struct FaceRefs{
+	// I didn't write it down, but I assume the edges are directed
+	// clockwise when looking towards the face normal, i.e. right hand rule.
 	edges:Vec<SubmeshDirectedEdgeId>,
+	//verts are redundant, use edge[i].verts[0]
 	//verts:Vec<VertId>,
 }
 #[derive(Debug)]
 struct EdgeRefs{
 	faces:[SubmeshFaceId;2],//left, right
-	verts:[SubmeshVertId;2],//bottom, top
+	verts:[SubmeshVertId;2],//start, end
 }
 #[derive(Debug)]
 struct VertRefs{
 	faces:Vec<SubmeshFaceId>,
+	// edges are always directed away from the vert
 	edges:Vec<SubmeshDirectedEdgeId>,
 }
 #[derive(Debug)]
@@ -442,11 +440,29 @@ impl TryFrom<&model::Mesh> for PhysicsMesh{
 	}
 }
 
-#[derive(Debug)]
+#[derive(Debug,Clone,Copy)]
 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;
@@ -457,13 +473,27 @@ impl MeshQuery for PhysicsMeshView<'_>{
 		let face_idx=self.topology.faces[face_id.get() as usize].get() as usize;
 		(self.data.faces[face_idx].normal,self.data.faces[face_idx].dot)
 	}
+	fn hint_point(&self)->Planar64Vec3{
+		// invariant: meshes always encompass the origin
+		vec3::zero()
+	}
+	fn farthest_vert(&self,dir:Planar64Vec3)->SubmeshVertId{
+		//this happens to be well-defined.  there are no virtual virtices
+		SubmeshVertId::new(
+			self.topology.verts.iter()
+			.enumerate()
+			.max_by_key(|&(_,&vert_id)|
+				dir.dot(self.data.verts[vert_id.get() as usize].0)
+			)
+			//assume there is more than zero vertices.
+			.unwrap().0 as u32
+		)
+	}
 	//ideally I never calculate the vertex position, but I have to for the graphical meshes...
 	fn vert(&self,vert_id:SubmeshVertId)->Planar64Vec3{
 		let vert_idx=self.topology.verts[vert_id.get() as usize].get() as usize;
 		self.data.verts[vert_idx].0
 	}
-	/// Directed edges going clockwise when looking in the direction of the face normal.
-	/// (Edit this documentation if this is wrong!)
 	fn face_edges(&self,face_id:SubmeshFaceId)->impl AsRef<[SubmeshDirectedEdgeId]>{
 		self.topology.face_topology[face_id.get() as usize].edges.as_slice()
 	}
@@ -497,7 +527,7 @@ impl PhysicsMeshTransform{
 	}
 }
 
-#[derive(Debug)]
+#[derive(Debug,Clone,Copy)]
 pub struct TransformedMesh<'a>{
 	view:PhysicsMeshView<'a>,
 	transform:&'a PhysicsMeshTransform,
@@ -515,17 +545,8 @@ 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))
 	}
-	fn farthest_vert(&self,dir:Planar64Vec3)->SubmeshVertId{
-		//this happens to be well-defined.  there are no virtual virtices
-		SubmeshVertId::new(
-			self.view.topology.verts.iter()
-			.enumerate()
-			.max_by_key(|&(_,&vert_id)|
-				dir.dot(self.transform.vertex.transform_point3(self.view.data.verts[vert_id.get() as usize].0))
-			)
-			//assume there is more than zero vertices.
-			.unwrap().0 as u32
-		)
+	pub fn faces(&self)->impl Iterator<Item=SubmeshFaceId>{
+		(0..self.view.topology.faces.len() as u32).map(SubmeshFaceId::new)
 	}
 }
 impl MeshQuery for TransformedMesh<'_>{
@@ -544,6 +565,21 @@ impl MeshQuery for TransformedMesh<'_>{
 		// wrap for speed
 		self.transform.vertex.transform_point3(self.view.vert(vert_id)).wrap_1()
 	}
+	fn hint_point(&self)->Planar64Vec3{
+		self.transform.vertex.translation
+	}
+	fn farthest_vert(&self,dir:Planar64Vec3)->SubmeshVertId{
+		//this happens to be well-defined.  there are no virtual virtices
+		SubmeshVertId::new(
+			self.view.topology.verts.iter()
+			.enumerate()
+			.max_by_key(|&(_,&vert_id)|
+				dir.dot(self.transform.vertex.transform_point3(self.view.data.verts[vert_id.get() as usize].0))
+			)
+			//assume there is more than zero vertices.
+			.unwrap().0 as u32
+		)
+	}
 	#[inline]
 	fn face_edges(&self,face_id:SubmeshFaceId)->impl AsRef<[SubmeshDirectedEdgeId]>{
 		self.view.face_edges(face_id)
@@ -570,11 +606,20 @@ impl MeshQuery for TransformedMesh<'_>{
 //(face,vertex)
 //(edge,edge)
 //(vertex,face)
-#[derive(Clone,Copy,Debug)]
+#[derive(Clone,Copy,Debug,Eq,PartialEq)]
 pub enum MinkowskiVert{
 	VertVert(SubmeshVertId,SubmeshVertId),
 }
-#[derive(Clone,Copy,Debug)]
+// TODO: remove this
+impl core::ops::Neg for MinkowskiVert{
+	type Output=Self;
+	fn neg(self)->Self::Output{
+		match self{
+			MinkowskiVert::VertVert(v0,v1)=>MinkowskiVert::VertVert(v1,v0),
+		}
+	}
+}
+#[derive(Clone,Copy,Debug,Eq,PartialEq)]
 pub enum MinkowskiEdge{
 	VertEdge(SubmeshVertId,SubmeshEdgeId),
 	EdgeVert(SubmeshEdgeId,SubmeshVertId),
@@ -589,7 +634,7 @@ impl UndirectedEdge for MinkowskiEdge{
 		}
 	}
 }
-#[derive(Clone,Copy,Debug)]
+#[derive(Clone,Copy,Debug,Eq,PartialEq)]
 pub enum MinkowskiDirectedEdge{
 	VertEdge(SubmeshVertId,SubmeshDirectedEdgeId),
 	EdgeVert(SubmeshDirectedEdgeId,SubmeshVertId),
@@ -610,7 +655,7 @@ impl DirectedEdge for MinkowskiDirectedEdge{
 		}
 	}
 }
-#[derive(Clone,Copy,Debug,Hash,Eq,PartialEq)]
+#[derive(Clone,Copy,Debug,Hash)]
 pub enum MinkowskiFace{
 	VertFace(SubmeshVertId,SubmeshFaceId),
 	EdgeEdge(SubmeshEdgeId,SubmeshEdgeId,bool),
@@ -632,6 +677,14 @@ pub fn into_giga_time(time:Time,relative_to:Time)->GigaTime{
 	Ratio::new(r.num.widen_4(),r.den.widen_4())
 }
 
+// TODO: remove this
+impl<'a> core::ops::Neg for &MinkowskiMesh<'a>{
+	type Output=MinkowskiMesh<'a>;
+	fn neg(self)->Self::Output{
+		MinkowskiMesh::minkowski_sum(self.mesh1,self.mesh0)
+	}
+}
+
 impl MinkowskiMesh<'_>{
 	pub fn minkowski_sum<'a>(mesh0:TransformedMesh<'a>,mesh1:TransformedMesh<'a>)->MinkowskiMesh<'a>{
 		MinkowskiMesh{
@@ -639,212 +692,27 @@ impl MinkowskiMesh<'_>{
 			mesh1,
 		}
 	}
-	fn farthest_vert(&self,dir:Planar64Vec3)->MinkowskiVert{
-		MinkowskiVert::VertVert(self.mesh0.farthest_vert(dir),self.mesh1.farthest_vert(-dir))
-	}
-	fn closest_fev_not_inside(&self,relative_position:Planar64Vec3)->Option<FEV<MinkowskiMesh<'_>>>{
-		// Make a fast guess as to what the closest point will be.
-		let MinkowskiVert::VertVert(mut v0,mut v1)=self.farthest_vert(relative_position);
-		// TODO: alternate naive vertex searches to improve the robustness
-		// in the "tall bipyramid" failure case
-		let mut m0v=self.mesh0.vert(v0);
-		let mut m1v=self.mesh1.vert(v1);
-		let mut best_distance_squared={
-			let diff=relative_position+m1v-m0v;
-			diff.dot(diff)
-		};
-		let mut v0e=self.mesh0.vert_edges(v0);
-		let mut v1e=self.mesh1.vert_edges(v1);
-		let mut v0e_ref=v0e.as_ref();
-		let mut v1e_ref=v1e.as_ref();
-		// repeatedly check adjacent vertex permutations to see if they are closer
-		loop{
-			let mut best_v0=None;
-			let mut best_v1=None;
-
-			// check vertices adjacent to v1 against v0
-			for m1_test_vert in self.mesh1.directed_verts(v1e_ref){
-				let m1v_test=self.mesh1.vert(m1_test_vert);
-				let diff=relative_position+m1v_test-m0v;
-				let d=diff.dot(diff);
-				if d<best_distance_squared{
-					best_distance_squared=d;
-					best_v0=None;
-					best_v1=Some(m1_test_vert);
-				}
-			}
-
-			// check vertices adjacent to v0 against v1
-			for m0_test_vert in self.mesh0.directed_verts(v0e_ref){
-				let m0v_test=self.mesh0.vert(m0_test_vert);
-				let diff=relative_position+m1v-m0v_test;
-				let d=diff.dot(diff);
-				if d<best_distance_squared{
-					best_distance_squared=d;
-					best_v0=Some(m0_test_vert);
-					best_v1=None;
-				}
-			}
-
-			// check permutations of adjacent vertices
-			for m0_test_vert in self.mesh0.directed_verts(v0e_ref){
-				let m0v_test=self.mesh0.vert(m0_test_vert);
-				for m1_test_vert in self.mesh1.directed_verts(v1e_ref){
-					let m1v_test=self.mesh1.vert(m1_test_vert);
-					let diff=relative_position+m1v_test-m0v_test;
-					let d=diff.dot(diff);
-					if d<best_distance_squared{
-						best_distance_squared=d;
-						best_v0=Some(m0_test_vert);
-						best_v1=Some(m1_test_vert);
-					}
-				}
-			}
-
-			// end condition
-			let v0_changed=match best_v0{
-				Some(new_v0)=>{
-					v0=new_v0;
-					m0v=self.mesh0.vert(v0);
-					v0e=self.mesh0.vert_edges(v0);
-					v0e_ref=v0e.as_ref();
-					true
-				},
-				None=>false,
-			};
-			let v1_changed=match best_v1{
-				Some(new_v1)=>{
-					v1=new_v1;
-					m1v=self.mesh1.vert(v1);
-					v1e=self.mesh1.vert_edges(v1);
-					v1e_ref=v1e.as_ref();
-					true
-				},
-				None=>false,
-			};
-			// if neither vertex changes, we found the closest two vertices.
-			if !(v0_changed&&v1_changed){
-				break;
-			}
-		}
-		// now you have the two closest vertices
-		let mut best_fev=FEV::Vert(MinkowskiVert::VertVert(v0,v1));
-		// ==== FEV::Edge ====
-		// test VertEdges
-		for &e1 in v1e_ref{
-			let edge_verts=self.mesh1.edge_verts(e1.as_undirected());
-			let &[ev0_id,ev1_id]=edge_verts.as_ref();
-			let (ev0,ev1)=(self.mesh1.vert(ev0_id),self.mesh1.vert(ev1_id));
-			let edge_n=ev1-ev0;
-			// use relative coordinates to make including relative_position easier
-			let diff=relative_position+m0v-ev0;
-			let d=diff.dot(edge_n);
-			// is test point between edge vertices
-			let edge_nn=edge_n.dot(edge_n);
-			if d.is_positive()&&d<edge_nn{
-				let distance_squared={
-					let c=diff.cross(edge_n);
-					//wrap for speed
-					(c.dot(c)/edge_nn).divide().wrap_2()
-				};
-				if distance_squared<best_distance_squared{
-					best_distance_squared=distance_squared;
-					best_fev=FEV::Edge(MinkowskiEdge::VertEdge(v0,e1.as_undirected()))
-				}
-			}
-		}
-		// test EdgeVerts
-		for &e0 in v0e_ref{
-			let edge_verts=self.mesh0.edge_verts(e0.as_undirected());
-			let &[ev0_id,ev1_id]=edge_verts.as_ref();
-			let (ev0,ev1)=(self.mesh0.vert(ev0_id),self.mesh0.vert(ev1_id));
-			let edge_n=ev1-ev0;
-			// use relative coordinates to make including relative_position easier
-			let diff=m1v-relative_position-ev0;
-			let d=diff.dot(edge_n);
-			// is test point between edge vertices
-			let edge_nn=edge_n.dot(edge_n);
-			if d.is_positive()&&d<edge_nn{
-				let distance_squared={
-					let c=diff.cross(edge_n);
-					//wrap for speed
-					(c.dot(c)/edge_nn).divide().wrap_2()
-				};
-				if distance_squared<best_distance_squared{
-					best_distance_squared=distance_squared;
-					best_fev=FEV::Edge(MinkowskiEdge::EdgeVert(e0.as_undirected(),v1))
-				}
-			}
-		}
-
-		// ==== FEV::Face ====
-		// test VertFaces
-		'outer: for &f1 in self.mesh1.vert_faces(v1).as_ref(){
-			let (n,d)=self.mesh1.face_nd(f1);
-			// Test the face's voronoi column
-			for &e1 in self.mesh1.face_edges(f1).as_ref(){
-				let edge_n=self.mesh1.directed_edge_n(e1);
-				let boundary_n=n.cross(edge_n);
-				let &[ev0_id,ev1_id]=self.mesh1.edge_verts(e1.as_undirected()).as_ref();
-				let (ev0,ev1)=(self.mesh1.vert(ev0_id),self.mesh1.vert(ev1_id));
-				let diff=(relative_position+m0v)*2-(ev0+ev1);
-				if boundary_n.dot(diff).is_negative(){
-					// The test point is outside the face's voronoi column.
-					continue 'outer;
-				}
-			}
-			// Calculate distance
-			let d=n.dot(relative_position+m0v)-d;
-			// Wrap for speed
-			let distance_squared=(d*d).wrap_2();
-			if distance_squared<best_distance_squared{
-				best_distance_squared=distance_squared;
-				best_fev=FEV::Face(MinkowskiFace::VertFace(v0,f1));
-			}
-		}
-		// test FaceVerts
-		'outer: for &f0 in self.mesh0.vert_faces(v0).as_ref(){
-			let (n,d)=self.mesh0.face_nd(f0);
-			// Test the face's voronoi column
-			for &e0 in self.mesh0.face_edges(f0).as_ref(){
-				let edge_n=self.mesh0.directed_edge_n(e0);
-				let boundary_n=n.cross(edge_n);
-				let &[ev0_id,ev1_id]=self.mesh0.edge_verts(e0.as_undirected()).as_ref();
-				let (ev0,ev1)=(self.mesh0.vert(ev0_id),self.mesh0.vert(ev1_id));
-				let diff=(m1v-relative_position)*2-(ev0+ev1);
-				if boundary_n.dot(diff).is_negative(){
-					// The test point is outside the face's voronoi column.
-					continue 'outer;
-				}
-			}
-			// Calculate distance
-			let d=n.dot(relative_position+m0v)-d;
-			// Wrap for speed
-			let distance_squared=(d*d).wrap_2();
-			if distance_squared<best_distance_squared{
-				best_distance_squared=distance_squared;
-				best_fev=FEV::Face(MinkowskiFace::FaceVert(f0,v1));
-			}
-		}
-		// test EdgeEdges
-		Some(best_fev)
-	}
 	pub fn predict_collision_in(&self,relative_body:&Body,range:impl RangeBounds<Time>)->Option<(MinkowskiFace,GigaTime)>{
-		self.closest_fev_not_inside(relative_body.position).and_then(|fev|{
-			//continue forwards along the body parabola
-			fev.crawl(self,relative_body,range.start_bound(),range.end_bound()).hit()
-		})
+		let fev=crate::minimum_difference::closest_fev_not_inside(self,relative_body.position)?;
+		//continue forwards along the body parabola
+		fev.crawl(self,relative_body,range.start_bound(),range.end_bound()).hit()
 	}
 	pub fn predict_collision_out(&self,relative_body:&Body,range:impl RangeBounds<Time>)->Option<(MinkowskiFace,GigaTime)>{
 		let (lower_bound,upper_bound)=(range.start_bound(),range.end_bound());
+		// TODO: handle unbounded collision using infinity fev
+		let time=match upper_bound{
+			Bound::Included(&time)=>time,
+			Bound::Excluded(&time)=>time,
+			Bound::Unbounded=>unimplemented!("unbounded collision out"),
+		};
+		let fev=crate::minimum_difference::closest_fev_not_inside(self,relative_body.extrapolated_position(time))?;
 		// swap and negate bounds to do a time inversion
 		let (lower_bound,upper_bound)=(upper_bound.map(|&t|-t),lower_bound.map(|&t|-t));
-		self.closest_fev_not_inside(relative_body.position).and_then(|fev|{
-			//continue backwards along the body parabola
-			fev.crawl(self,&-relative_body,lower_bound.as_ref(),upper_bound.as_ref()).hit()
-				//no need to test -time<time_limit because of the first step
-				.map(|(face,time)|(face,-time))
-		})
+		let infinity_body=-relative_body;
+		//continue backwards along the body parabola
+		fev.crawl(self,&infinity_body,lower_bound.as_ref(),upper_bound.as_ref()).hit()
+			//no need to test -time<time_limit because of the first step
+			.map(|(face,time)|(face,-time))
 	}
 	pub fn predict_collision_face_out(&self,relative_body:&Body,range:impl RangeBounds<Time>,contact_face_id:MinkowskiFace)->Option<(MinkowskiDirectedEdge,GigaTime)>{
 		// TODO: make better
@@ -874,10 +742,8 @@ impl MinkowskiMesh<'_>{
 		}
 		best_edge
 	}
-	pub fn is_point_in_mesh(&self,point:Planar64Vec3)->bool{
-		// TODO
-		println!("Unimplemented is_point_in_mesh called! {point}");
-		false
+	pub fn contains_point(&self,point:Planar64Vec3)->bool{
+		crate::minimum_difference::contains_point(self,point)
 	}
 }
 impl MeshQuery for MinkowskiMesh<'_>{
@@ -916,6 +782,12 @@ impl MeshQuery for MinkowskiMesh<'_>{
 			},
 		}
 	}
+	fn hint_point(&self)->Planar64Vec3{
+		self.mesh0.transform.vertex.translation-self.mesh1.transform.vertex.translation
+	}
+	fn farthest_vert(&self,dir:Planar64Vec3)->MinkowskiVert{
+		MinkowskiVert::VertVert(self.mesh0.farthest_vert(dir),self.mesh1.farthest_vert(-dir))
+	}
 	fn face_edges(&self,face_id:MinkowskiFace)->impl AsRef<[MinkowskiDirectedEdge]>{
 		match face_id{
 			MinkowskiFace::VertFace(v0,f1)=>{

engine/physics/src/push_solve.rs

@@ -39,20 +39,18 @@ impl Contact{
 
 //note that this is horrible with fixed point arithmetic
 fn solve1(c0:&Contact)->Option<Ratio<Vector3<Fixed<3,96>>,Fixed<2,64>>>{
-	const EPSILON:Fixed<2,64>=Fixed::from_bits(Fixed::<2,64>::ONE.to_bits().shr(10));
 	let det=c0.normal.dot(c0.velocity);
-	if det.abs()<EPSILON{
+	if det.abs()==Fixed::ZERO{
 		return None;
 	}
 	let d0=c0.normal.dot(c0.position);
 	Some(c0.normal*d0/det)
 }
 fn solve2(c0:&Contact,c1:&Contact)->Option<Ratio<Vector3<Fixed<5,160>>,Fixed<4,128>>>{
-	const EPSILON:Fixed<4,128>=Fixed::from_bits(Fixed::<4,128>::ONE.to_bits().shr(10));
 	let u0_u1=c0.velocity.cross(c1.velocity);
 	let n0_n1=c0.normal.cross(c1.normal);
 	let det=u0_u1.dot(n0_n1);
-	if det.abs()<EPSILON{
+	if det.abs()==Fixed::ZERO{
 		return None;
 	}
 	let d0=c0.normal.dot(c0.position);
@@ -60,10 +58,9 @@ fn solve2(c0:&Contact,c1:&Contact)->Option<Ratio<Vector3<Fixed<5,160>>,Fixed<4,1
 	Some((c1.normal.cross(u0_u1)*d0+u0_u1.cross(c0.normal)*d1)/det)
 }
 fn solve3(c0:&Contact,c1:&Contact,c2:&Contact)->Option<Ratio<Vector3<Fixed<4,128>>,Fixed<3,96>>>{
-	const EPSILON:Fixed<3,96>=Fixed::from_bits(Fixed::<3,96>::ONE.to_bits().shr(10));
 	let n0_n1=c0.normal.cross(c1.normal);
 	let det=c2.normal.dot(n0_n1);
-	if det.abs()<EPSILON{
+	if det.abs()==Fixed::ZERO{
 		return None;
 	}
 	let d0=c0.normal.dot(c0.position);
@@ -149,7 +146,7 @@ fn is_space_enclosed_4(
 }
 
 const fn get_push_ray_0(point:Planar64Vec3)->Ray{
-	Ray{origin:point,direction:vec3::ZERO}
+	Ray{origin:point,direction:vec3::zero()}
 }
 fn get_push_ray_1(point:Planar64Vec3,c0:&Contact)->Option<Ray>{
 	//wrap for speed
@@ -321,13 +318,13 @@ mod tests{
 	fn test_push_solve(){
 		let contacts=vec![
 			Contact{
-				position:vec3::ZERO,
+				position:vec3::zero(),
 				velocity:vec3::Y,
 				normal:vec3::Y,
 			}
 		];
 		assert_eq!(
-			vec3::ZERO,
+			vec3::zero(),
 			push_solve(&contacts,vec3::NEG_Y)
 		);
 	}

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
 	}

integration-testing/src/test_scenes.rs

@@ -77,7 +77,7 @@ fn simultaneous_collision(){
 		Time::ZERO,
 	);
 	let mut physics=PhysicsState::new_with_body(body);
-	physics.style_mut().gravity=vec3::ZERO;
+	physics.style_mut().gravity=vec3::zero();
 	let mut phys_iter=PhysicsContext::iter_internal(&mut physics,&physics_data,Time::from_secs(2))
 		.filter(|ins|!matches!(ins.instruction,InternalInstruction::StrafeTick));
 	// the order that they hit does matter, but we aren't currently worrying about that.
@@ -101,7 +101,7 @@ fn bug_3(){
 		Time::ZERO,
 	);
 	let mut physics=PhysicsState::new_with_body(body);
-	physics.style_mut().gravity=vec3::ZERO;
+	physics.style_mut().gravity=vec3::zero();
 	let mut phys_iter=PhysicsContext::iter_internal(&mut physics,&physics_data,Time::from_secs(3))
 		.filter(|ins|!matches!(ins.instruction,InternalInstruction::StrafeTick));
 	// touch side of part at 0,0,0

lib/bsp_loader/src/brush.rs

@@ -210,7 +210,7 @@ pub fn faces_to_mesh(faces:Vec<Vec<integer::Planar64Vec3>>)->model::Mesh{
 	let color=mb.acquire_color_id(glam::Vec4::ONE);
 	let tex=mb.acquire_tex_id(glam::Vec2::ZERO);
 	// normals are ignored by physics
-	let normal=mb.acquire_normal_id(integer::vec3::ZERO);
+	let normal=mb.acquire_normal_id(integer::vec3::zero());
 
 	let polygon_list=faces.into_iter().map(|face|{
 		face.into_iter().map(|pos|{

lib/bsp_loader/src/bsp.rs

@@ -105,7 +105,7 @@ pub fn convert<'a>(
 					water:Some(attr::IntersectingWater{
 						viscosity:integer::Planar64::ONE,
 						density:integer::Planar64::ONE,
-						velocity:integer::vec3::ZERO,
+						velocity:integer::vec3::zero(),
 					}),
 				},
 				general:attr::GeneralAttributes::default(),
@@ -295,7 +295,7 @@ pub fn convert<'a>(
 					attributes,
 					transform:integer::Planar64Affine3::new(
 						integer::mat3::identity(),
-						integer::vec3::ZERO,
+						integer::vec3::zero(),
 					),
 					color:glam::Vec4::ONE,
 				});

lib/common/src/gameplay_style.rs

@@ -319,7 +319,7 @@ impl WalkSettings{
 		self.accelerate.accel.min((-gravity.y*friction).clamp_1())
 	}
 	pub fn get_walk_target_velocity(&self,control_dir:Planar64Vec3,normal:Planar64Vec3)->Planar64Vec3{
-		if control_dir==crate::integer::vec3::ZERO{
+		if control_dir==crate::integer::vec3::zero(){
 			return control_dir;
 		}
 		let nn=normal.length_squared();
@@ -329,13 +329,13 @@ impl WalkSettings{
 		let dd=d*d;
 		if dd<nnmm{
 			let cr=normal.cross(control_dir);
-			if cr==crate::integer::vec3::ZERO_2{
-				crate::integer::vec3::ZERO
+			if cr==crate::integer::vec3::zero(){
+				crate::integer::vec3::zero()
 			}else{
 				(cr.cross(normal)*self.accelerate.topspeed/((nn*(nnmm-dd)).sqrt())).divide().clamp_1()
 			}
 		}else{
-			crate::integer::vec3::ZERO
+			crate::integer::vec3::zero()
 		}
 	}
 	pub fn is_slope_walkable(&self,normal:Planar64Vec3,up:Planar64Vec3)->bool{
@@ -360,7 +360,7 @@ impl LadderSettings{
 		self.accelerate.accel
 	}
 	pub fn get_ladder_target_velocity(&self,mut control_dir:Planar64Vec3,normal:Planar64Vec3)->Planar64Vec3{
-		if control_dir==crate::integer::vec3::ZERO{
+		if control_dir==crate::integer::vec3::zero(){
 			return control_dir;
 		}
 		let nn=normal.length_squared();
@@ -382,13 +382,13 @@ impl LadderSettings{
 		//- fix the underlying issue
 		if dd<nnmm{
 			let cr=normal.cross(control_dir);
-			if cr==crate::integer::vec3::ZERO_2{
-				crate::integer::vec3::ZERO
+			if cr==crate::integer::vec3::zero(){
+				crate::integer::vec3::zero()
 			}else{
 				(cr.cross(normal)*self.accelerate.topspeed/((nn*(nnmm-dd)).sqrt())).divide().clamp_1()
 			}
 		}else{
-			crate::integer::vec3::ZERO
+			crate::integer::vec3::zero()
 		}
 	}
 }

lib/common/src/integer.rs

@@ -561,12 +561,6 @@ pub mod vec3{
 	pub use linear_ops::types::Vector3;
 	pub const MIN:Planar64Vec3=Planar64Vec3::new([Planar64::MIN;3]);
 	pub const MAX:Planar64Vec3=Planar64Vec3::new([Planar64::MAX;3]);
-	pub const ZERO:Planar64Vec3=Planar64Vec3::new([Planar64::ZERO;3]);
-	pub const ZERO_2:Vector3<Fixed::<2,64>>=Vector3::new([Fixed::<2,64>::ZERO;3]);
-	pub const ZERO_3:Vector3<Fixed::<3,96>>=Vector3::new([Fixed::<3,96>::ZERO;3]);
-	pub const ZERO_4:Vector3<Fixed::<4,128>>=Vector3::new([Fixed::<4,128>::ZERO;3]);
-	pub const ZERO_5:Vector3<Fixed::<5,160>>=Vector3::new([Fixed::<5,160>::ZERO;3]);
-	pub const ZERO_6:Vector3<Fixed::<6,192>>=Vector3::new([Fixed::<6,192>::ZERO;3]);
 	pub const X:Planar64Vec3=Planar64Vec3::new([Planar64::ONE,Planar64::ZERO,Planar64::ZERO]);
 	pub const Y:Planar64Vec3=Planar64Vec3::new([Planar64::ZERO,Planar64::ONE,Planar64::ZERO]);
 	pub const Z:Planar64Vec3=Planar64Vec3::new([Planar64::ZERO,Planar64::ZERO,Planar64::ONE]);
@@ -575,6 +569,10 @@ pub mod vec3{
 	pub const NEG_Y:Planar64Vec3=Planar64Vec3::new([Planar64::ZERO,Planar64::NEG_ONE,Planar64::ZERO]);
 	pub const NEG_Z:Planar64Vec3=Planar64Vec3::new([Planar64::ZERO,Planar64::ZERO,Planar64::NEG_ONE]);
 	pub const NEG_ONE:Planar64Vec3=Planar64Vec3::new([Planar64::NEG_ONE,Planar64::NEG_ONE,Planar64::NEG_ONE]);
+	// TODO: use #![feature(generic_const_items)] when stabilized https://github.com/rust-lang/rust/issues/113521
+	pub const fn zero<const N:usize,const F:usize>()->Vector3<Fixed<N,F>>{
+		Vector3::new([Fixed::ZERO;3])
+	}
 	#[inline]
 	pub const fn int(x:i32,y:i32,z:i32)->Planar64Vec3{
 		Planar64Vec3::new([Planar64::raw((x as i64)<<32),Planar64::raw((y as i64)<<32),Planar64::raw((z as i64)<<32)])
@@ -663,7 +661,7 @@ pub struct Planar64Affine3{
 	pub translation:Planar64Vec3,
 }
 impl Planar64Affine3{
-	pub const IDENTITY:Self=Self::new(mat3::identity(),vec3::ZERO);
+	pub const IDENTITY:Self=Self::new(mat3::identity(),vec3::zero());
 	#[inline]
 	pub const fn new(matrix3:Planar64Mat3,translation:Planar64Vec3)->Self{
 		Self{matrix3,translation}

lib/fixed_wide/src/fixed.rs

@@ -70,6 +70,34 @@ impl<const N:usize,const F:usize> Fixed<N,F>{
 	pub const fn midpoint(self,other:Self)->Self{
 		Self::from_bits(self.bits.midpoint(other.bits))
 	}
+	#[inline]
+	pub const fn min(self,other:Self)->Self{
+		Self::from_bits(self.bits.min(other.bits))
+	}
+	#[inline]
+	pub const fn max(self,other:Self)->Self{
+		Self::from_bits(self.bits.max(other.bits))
+	}
+	/// return the result of self*sign(other)
+	#[inline]
+	pub const fn mul_sign<const N1:usize,const F1:usize>(self,other:Fixed<N1,F1>)->Self{
+		if other.is_negative(){
+			Self::from_bits(self.bits.neg())
+		}else if other.is_zero(){
+			Fixed::ZERO
+		}else{
+			self
+		}
+	}
+	/// return the result of self/sign(other) (divide by zero does not change the sign)
+	#[inline]
+	pub const fn div_sign<const N1:usize,const F1:usize>(self,other:Fixed<N1,F1>)->Self{
+		if other.is_negative(){
+			Self::from_bits(self.bits.neg())
+		}else{
+			self
+		}
+	}
 }
 impl<const F:usize> Fixed<1,F>{
 	/// My old code called this function everywhere so let's provide it
@@ -101,28 +129,6 @@ impl_from!(
 	i8,i16,i32,i64,i128,isize
 );
 
-impl<const N:usize,const F:usize,T> PartialEq<T> for Fixed<N,F>
-where
-		T:Copy,
-		BInt::<N>:From<T>,
-{
-	#[inline]
-	fn eq(&self,&other:&T)->bool{
-		self.bits.eq(&other.into())
-	}
-}
-
-impl<const N:usize,const F:usize,T> PartialOrd<T> for Fixed<N,F>
-	where
-		T:Copy,
-		BInt::<N>:From<T>,
-{
-	#[inline]
-	fn partial_cmp(&self,&other:&T)->Option<std::cmp::Ordering>{
-		self.bits.partial_cmp(&other.into())
-	}
-}
-
 impl<const N:usize,const F:usize> std::ops::Neg for Fixed<N,F>{
 	type Output=Self;
 	#[inline]
@@ -328,16 +334,6 @@ macro_rules! impl_additive_operator {
 				self.$method(other)
 			}
 		}
-		impl<const N:usize,const F:usize,U> core::ops::$trait<U> for $struct<N,F>
-			where
-				BInt::<N>:From<U>,
-		{
-			type Output = $output;
-			#[inline]
-			fn $method(self, other: U) -> Self::Output {
-				Self::from_bits(self.bits.$method(BInt::<N>::from(other).shl(F as u32)))
-			}
-		}
 	};
 }
 macro_rules! impl_additive_assign_operator {
@@ -348,15 +344,6 @@ macro_rules! impl_additive_assign_operator {
 				self.bits.$method(other.bits);
 			}
 		}
-		impl<const N:usize,const F:usize,U> core::ops::$trait<U> for $struct<N,F>
-		where
-				BInt::<N>:From<U>,
-		{
-			#[inline]
-			fn $method(&mut self, other: U) {
-				self.bits.$method(BInt::<N>::from(other).shl(F as u32));
-			}
-		}
 	};
 }
 
@@ -379,7 +366,7 @@ impl_additive_operator!( Fixed, BitXor, bitxor, Self );
 // non-wide operators.  The result is the same width as the inputs.
 
 // This macro is not used in the default configuration.
-#[allow(unused_macros)]
+#[expect(unused_macros)]
 macro_rules! impl_multiplicative_operator_not_const_generic {
 	( ($struct: ident, $trait: ident, $method: ident, $output: ty ), $width:expr ) => {
 		impl<const F:usize> core::ops::$trait for $struct<$width,F>{
@@ -558,7 +545,7 @@ impl_shift_operator!( Fixed, Shr, shr, Self );
 
 // wide operators.  The result width is the sum of the input widths, i.e. none of the multiplication
 
-#[allow(unused_macros)]
+#[expect(unused_macros)]
 macro_rules! impl_wide_operators{
 	($lhs:expr,$rhs:expr)=>{
 		impl core::ops::Mul<Fixed<$rhs,{$rhs*32}>> for Fixed<$lhs,{$lhs*32}>{

lib/rbx_loader/Cargo.toml

@@ -13,16 +13,16 @@ authors = ["Rhys Lloyd <krakow20@gmail.com>"]
 bytemuck = "1.14.3"
 glam = "0.30.0"
 regex = { version = "1.11.3", default-features = false }
-rbx_binary = { version = "1.0.1-sn5", registry = "strafesnet" }
-rbx_dom_weak =  { version = "3.0.1-sn5", registry = "strafesnet" }
 rbx_mesh = "0.5.0"
-rbx_reflection = "5.0.0"
-rbx_reflection_database = "1.0.0"
-rbx_xml = { version = "1.0.1-sn5", registry = "strafesnet" }
 rbxassetid = { version = "0.1.0", path = "../rbxassetid", registry = "strafesnet" }
 roblox_emulator = { version = "0.5.1", path = "../roblox_emulator", default-features = false, registry = "strafesnet" }
 strafesnet_common = { version = "0.7.0", path = "../common", registry = "strafesnet" }
 strafesnet_deferred_loader = { version = "0.5.1", path = "../deferred_loader", registry = "strafesnet" }
+rbx_binary = "2.0.1"
+rbx_dom_weak = "4.1.0"
+rbx_reflection = "6.1.0"
+rbx_reflection_database = "2.0.2"
+rbx_xml = "2.0.1"
 
 [lints]
 workspace = true

lib/rbx_loader/src/rbx.rs

@@ -251,7 +251,7 @@ fn get_attributes(name:&str,can_collide:bool,velocity:Planar64Vec3,model_id:mode
 		}
 	}
 	//need some way to skip this
-	if allow_booster&&velocity!=vec3::ZERO{
+	if allow_booster&&velocity!=vec3::zero(){
 		general.booster=Some(attr::Booster::Velocity(velocity));
 	}
 	Ok(match force_can_collide{
@@ -559,7 +559,7 @@ pub fn convert<'a>(
 	//just going to leave it like this for now instead of reworking the data structures for this whole thing
 	let textureless_render_group=render_config_deferred_loader.acquire_render_config_id(None);
 
-	let db=rbx_reflection_database::get();
+	let db=rbx_reflection_database::get().unwrap();
 	let basepart=&db.classes["BasePart"];
 	let baseparts=dom.descendants().filter(|&instance|
 		db.classes.get(instance.class.as_str()).is_some_and(|class|

lib/rbx_loader/src/union.rs

@@ -250,7 +250,7 @@ pub fn convert(
 		// generate a unit cube as default physics
 		let pos_list=CUBE_DEFAULT_VERTICES.map(|pos|mb.acquire_pos_id(pos>>1));
 		let tex=mb.acquire_tex_id(glam::Vec2::ZERO);
-		let normal=mb.acquire_normal_id(vec3::ZERO);
+		let normal=mb.acquire_normal_id(vec3::zero());
 		let color=mb.acquire_color_id(glam::Vec4::ONE);
 		let polygon_group=PolygonGroup::PolygonList(PolygonList::new(CUBE_DEFAULT_POLYS.map(|poly|poly.map(|[pos_id,_]|
 			mb.acquire_vertex_id(IndexedVertex{pos:pos_list[pos_id as usize],tex,normal,color})

lib/roblox_emulator/Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "roblox_emulator"
-version = "0.5.1"
+version = "0.5.2"
 edition = "2024"
 repository = "https://git.itzana.me/StrafesNET/strafe-project"
 license = "MIT OR Apache-2.0"
@@ -15,10 +15,10 @@ run-service=[]
 glam = "0.30.0"
 mlua = { version = "0.11.3", features = ["luau"] }
 phf = { version = "0.13.1", features = ["macros"] }
-rbx_dom_weak = { version = "3.0.1-sn5", registry = "strafesnet" }
-rbx_reflection = "5.0.0"
-rbx_reflection_database = "1.0.0"
-rbx_types = "2.0.0"
+rbx_dom_weak = "4.1.0"
+rbx_reflection = "6.1.0"
+rbx_reflection_database = "2.0.2"
+rbx_types = "3.1.0"
 
 [lints]
 workspace = true

lib/roblox_emulator/src/context.rs

@@ -52,7 +52,7 @@ impl Context{
 	}
 	/// Creates an iterator over all items of a particular class.
 	pub fn superclass_iter<'a>(&'a self,superclass:&'a str)->impl Iterator<Item=Ref>+'a{
-		let db=rbx_reflection_database::get();
+		let db=rbx_reflection_database::get().unwrap();
 		let Some(superclass)=db.classes.get(superclass)else{
 			panic!("Invalid class");
 		};

lib/roblox_emulator/src/runner/enum.rs

@@ -37,7 +37,7 @@ impl PartialEq for EnumItem<'_>{
 pub struct Enums;
 impl Enums{
 	pub fn get(&self,index:&str)->Option<EnumItems<'static>>{
-		let db=rbx_reflection_database::get();
+		let db=rbx_reflection_database::get().unwrap();
 		db.enums.get(index).map(|ed|EnumItems{ed})
 	}
 }

lib/roblox_emulator/src/runner/instance/instance.rs

@@ -37,7 +37,7 @@ pub fn dom_mut<T>(lua:&mlua::Lua,mut f:impl FnMut(&mut WeakDom)->mlua::Result<T>
 }
 
 pub fn class_is_a(class:&str,superclass:&str)->bool{
-	let db=rbx_reflection_database::get();
+	let db=rbx_reflection_database::get().unwrap();
 	let (Some(class),Some(superclass))=(db.classes.get(class),db.classes.get(superclass))else{
 		return false;
 	};
@@ -80,14 +80,14 @@ pub fn find_first_descendant_of_class<'a>(dom:&'a WeakDom,instance:&rbx_dom_weak
 }
 
 pub fn find_first_child_which_is_a<'a>(dom:&'a WeakDom,instance:&rbx_dom_weak::Instance,superclass:&str)->Option<&'a rbx_dom_weak::Instance>{
-	let db=rbx_reflection_database::get();
+	let db=rbx_reflection_database::get().unwrap();
 	let superclass_descriptor=db.classes.get(superclass)?;
 	instance.children().iter().filter_map(|&r|dom.get_by_ref(r)).find(|inst|{
 		db.classes.get(inst.class.as_str()).is_some_and(|descriptor|db.has_superclass(descriptor,superclass_descriptor))
 	})
 }
 pub fn find_first_descendant_which_is_a<'a>(dom:&'a WeakDom,instance:&rbx_dom_weak::Instance,superclass:&str)->Option<&'a rbx_dom_weak::Instance>{
-	let db=rbx_reflection_database::get();
+	let db=rbx_reflection_database::get().unwrap();
 	let superclass_descriptor=db.classes.get(superclass)?;
 	dom.descendants_of(instance.referent()).find(|inst|{
 		db.classes.get(inst.class.as_str()).is_some_and(|descriptor|db.has_superclass(descriptor,superclass_descriptor))
@@ -282,7 +282,7 @@ impl mlua::UserData for Instance{
 			dom_mut(lua,|dom|{
 				let instance=this.get(dom)?;
 				//println!("__index t={} i={index:?}",instance.name);
-				let db=rbx_reflection_database::get();
+				let db=rbx_reflection_database::get().unwrap();
 				let class=db.classes.get(instance.class.as_str()).ok_or_else(||mlua::Error::runtime("Class missing"))?;
 				// Find existing property
 				// Interestingly, ustr can know ahead of time if
@@ -344,7 +344,7 @@ impl mlua::UserData for Instance{
 			let index_str=&*index.to_str()?;
 			dom_mut(lua,|dom|{
 				let instance=this.get_mut(dom)?;
-				let db=rbx_reflection_database::get();
+				let db=rbx_reflection_database::get().unwrap();
 				let class=db.classes.get(instance.class.as_str()).ok_or_else(||mlua::Error::runtime("Class missing"))?;
 				let property=db.superclasses_iter(class).find_map(|cls|
 					cls.properties.get(index_str)

map-tool/Cargo.toml

@@ -13,10 +13,10 @@ futures = "0.3.31"
 image = "0.25.2"
 image_dds = "0.7.1"
 rbx_asset = { version = "0.5.0", registry = "strafesnet" }
-rbx_binary  = { version = "1.0.1-sn5", registry = "strafesnet" }
-rbx_dom_weak =  { version = "3.0.1-sn5", registry = "strafesnet" }
-rbx_reflection_database = "1.0.0"
-rbx_xml = { version = "1.0.1-sn5", registry = "strafesnet" }
+rbx_binary = "2.0.1"
+rbx_dom_weak = "4.1.0"
+rbx_reflection_database = "2.0.2"
+rbx_xml = "2.0.1"
 rbxassetid = { version = "0.1.0", registry = "strafesnet" }
 strafesnet_bsp_loader = { version = "0.3.1", path = "../lib/bsp_loader", registry = "strafesnet" }
 strafesnet_deferred_loader = { version = "0.5.1", path = "../lib/deferred_loader", registry = "strafesnet" }

strafe-client/Cargo.toml

@@ -28,7 +28,7 @@ strafesnet_rbx_loader = { path = "../lib/rbx_loader", registry = "strafesnet", o
 strafesnet_session = { path = "../engine/session", registry = "strafesnet" }
 strafesnet_settings = { path = "../engine/settings", registry = "strafesnet" }
 strafesnet_snf = { path = "../lib/snf", registry = "strafesnet", optional = true }
-wgpu = "27.0.0"
+wgpu = "28.0.0"
 winit = "0.30.7"
 
 [profile.dev]

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{
@@ -52,7 +52,7 @@ impl<'a> SetupContextPartial2<'a>{
 		let required_features=required_features();
 
 		//no helper function smh gotta write it myself
-		let adapters=self.instance.enumerate_adapters(self.backends);
+		let adapters=pollster::block_on(self.instance.enumerate_adapters(self.backends));
 
 		let mut chosen_adapter=None;
 		let mut chosen_adapter_score=0;
@@ -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;
+}