diff --git a/Cargo.lock b/Cargo.lock index 01a18cf007..a4f241401b 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1917,7 +1917,6 @@ dependencies = [ "arrayvec", "bitflags", "foreign-types", - "fxhash", "log", "metal", "naga", diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs index d7ef19aa08..745220a9e7 100644 --- a/wgpu-core/src/device/mod.rs +++ b/wgpu-core/src/device/mod.rs @@ -2175,6 +2175,7 @@ impl Device { use resource::CreateQuerySetError as Error; match desc.ty { + wgt::QueryType::Occlusion => {} wgt::QueryType::Timestamp => { self.require_features(wgt::Features::TIMESTAMP_QUERY)?; } diff --git a/wgpu-hal/Cargo.toml b/wgpu-hal/Cargo.toml index 6939cd7ff6..ccc14b184c 100644 --- a/wgpu-hal/Cargo.toml +++ b/wgpu-hal/Cargo.toml @@ -18,7 +18,6 @@ metal = ["foreign-types", "mtl", "objc", "parking_lot", "naga/msl-out"] [dependencies] arrayvec = "0.5" bitflags = "1.0" -fxhash = "0.2" log = "0.4" parking_lot = { version = "0.11", optional = true } raw-window-handle = "0.3" diff --git a/wgpu-hal/src/empty.rs b/wgpu-hal/src/empty.rs index a3973da56f..950f72b40e 100644 --- a/wgpu-hal/src/empty.rs +++ b/wgpu-hal/src/empty.rs @@ -174,7 +174,7 @@ impl crate::Device for Context { &self, desc: &crate::ShaderModuleDescriptor, shader: crate::NagaShader, - ) -> Result { + ) -> Result { Ok(Resource) } unsafe fn destroy_shader_module(&self, module: Resource) {} diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs index f3fda88b73..ab151ce556 100644 --- a/wgpu-hal/src/lib.rs +++ b/wgpu-hal/src/lib.rs @@ -88,6 +88,8 @@ pub enum ShaderError { pub enum PipelineError { #[error("linkage failed for stage {0:?}: {1}")] Linkage(wgt::ShaderStage, String), + #[error("entry point for stage {0:?} is invalid")] + EntryPoint(naga::ShaderStage), #[error(transparent)] Device(#[from] DeviceError), } @@ -239,7 +241,7 @@ pub trait Device { &self, desc: &ShaderModuleDescriptor, shader: NagaShader, - ) -> Result; + ) -> Result; unsafe fn destroy_shader_module(&self, module: A::ShaderModule); unsafe fn create_render_pipeline( &self, @@ -738,6 +740,15 @@ pub struct BindGroupEntry { pub resource_index: u32, } +/// BindGroup descriptor. +/// +/// Valid usage: +///. - `entries` has to be sorted by ascending `BindGroupEntry::binding` +///. - `entries` has to have the same set of `BindGroupEntry::binding` as `layout` +///. - each entry has to be compatible with the `layout` +///. - each entry's `BindGroupEntry::resource_index` is within range +/// of the corresponding resource array, selected by the relevant +/// `BindGroupLayoutEntry`. #[derive(Clone, Debug)] pub struct BindGroupDescriptor<'a, A: Api> { pub label: Label<'a>, @@ -754,6 +765,7 @@ pub struct CommandBufferDescriptor<'a> { } /// Naga shader module. +#[derive(Debug)] pub struct NagaShader { /// Shader module IR. pub module: naga::Module, diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs index 6f194a848a..45d745b40d 100644 --- a/wgpu-hal/src/metal/adapter.rs +++ b/wgpu-hal/src/metal/adapter.rs @@ -25,7 +25,6 @@ impl crate::Adapter for super::Adapter { features, }, queue: super::Queue { - raw: raw_device.new_command_queue(), }, }) } @@ -802,10 +801,15 @@ impl super::PrivateCapabilities { supports_binary_archives: family_check && (device.supports_family(MTLGPUFamily::Apple3) || device.supports_family(MTLGPUFamily::Mac1)), + supports_capture_manager: if os_is_mac { + Self::version_at_least(major, minor, 10, 13) + } else { + Self::version_at_least(major, minor, 11, 0) + }, can_set_maximum_drawables_count: os_is_mac || Self::version_at_least(major, minor, 11, 2), can_set_display_sync: os_is_mac && Self::version_at_least(major, minor, 10, 13), - can_set_next_drawable_timeout: if is_mac { + can_set_next_drawable_timeout: if os_is_mac { Self::version_at_least(major, minor, 10, 13) } else { Self::version_at_least(major, minor, 11, 0) diff --git a/wgpu-hal/src/metal/command.rs b/wgpu-hal/src/metal/command.rs index b5d88a0b0b..b705bf8dfb 100644 --- a/wgpu-hal/src/metal/command.rs +++ b/wgpu-hal/src/metal/command.rs @@ -1,19 +1,17 @@ -use super::{Api, Resource}; - use std::ops::Range; -impl crate::CommandBuffer for super::Encoder { +impl crate::CommandBuffer for super::CommandBuffer { unsafe fn finish(&mut self) {} unsafe fn transition_buffers<'a, T>(&mut self, barriers: T) where - T: Iterator>, + T: Iterator>, { } unsafe fn transition_textures<'a, T>(&mut self, barriers: T) where - T: Iterator>, + T: Iterator>, { } @@ -62,7 +60,7 @@ impl crate::CommandBuffer for super::Encoder { // render - unsafe fn begin_render_pass(&mut self, desc: &crate::RenderPassDescriptor) {} + unsafe fn begin_render_pass(&mut self, desc: &crate::RenderPassDescriptor) {} unsafe fn end_render_pass(&mut self) {} unsafe fn set_bind_group( @@ -90,11 +88,11 @@ impl crate::CommandBuffer for super::Encoder { unsafe fn set_index_buffer<'a>( &mut self, - binding: crate::BufferBinding<'a, Api>, + binding: crate::BufferBinding<'a, super::Api>, format: wgt::IndexFormat, ) { } - unsafe fn set_vertex_buffer<'a>(&mut self, index: u32, binding: crate::BufferBinding<'a, Api>) { + unsafe fn set_vertex_buffer<'a>(&mut self, index: u32, binding: crate::BufferBinding<'a, super::Api>) { } unsafe fn set_viewport(&mut self, rect: &crate::Rect, depth_range: Range) {} unsafe fn set_scissor_rect(&mut self, rect: &crate::Rect) {} diff --git a/wgpu-hal/src/metal/conv.rs b/wgpu-hal/src/metal/conv.rs index 95d26fa884..b099edf949 100644 --- a/wgpu-hal/src/metal/conv.rs +++ b/wgpu-hal/src/metal/conv.rs @@ -75,3 +75,180 @@ pub fn map_border_color(border_color: wgt::SamplerBorderColor) -> mtl::MTLSample wgt::SamplerBorderColor::OpaqueWhite => OpaqueWhite, } } + +pub fn map_primitive_topology( + topology: wgt::PrimitiveTopology, +) -> (mtl::MTLPrimitiveTopologyClass, mtl::MTLPrimitiveType) { + use wgt::PrimitiveTopology as Pt; + match topology { + Pt::PointList => ( + mtl::MTLPrimitiveTopologyClass::Point, + mtl::MTLPrimitiveType::Point, + ), + Pt::LineList => ( + mtl::MTLPrimitiveTopologyClass::Line, + mtl::MTLPrimitiveType::Line, + ), + Pt::LineStrip => ( + mtl::MTLPrimitiveTopologyClass::Line, + mtl::MTLPrimitiveType::LineStrip, + ), + Pt::TriangleList => ( + mtl::MTLPrimitiveTopologyClass::Triangle, + mtl::MTLPrimitiveType::Triangle, + ), + Pt::TriangleStrip => ( + mtl::MTLPrimitiveTopologyClass::Triangle, + mtl::MTLPrimitiveType::TriangleStrip, + ), + } +} + +pub fn map_color_write(mask: wgt::ColorWrite) -> mtl::MTLColorWriteMask { + let mut raw_mask = mtl::MTLColorWriteMask::empty(); + + if mask.contains(wgt::ColorWrite::RED) { + raw_mask |= mtl::MTLColorWriteMask::Red; + } + if mask.contains(wgt::ColorWrite::GREEN) { + raw_mask |= mtl::MTLColorWriteMask::Green; + } + if mask.contains(wgt::ColorWrite::BLUE) { + raw_mask |= mtl::MTLColorWriteMask::Blue; + } + if mask.contains(wgt::ColorWrite::ALPHA) { + raw_mask |= mtl::MTLColorWriteMask::Alpha; + } + + raw_mask +} + +pub fn map_blend_factor(factor: wgt::BlendFactor) -> mtl::MTLBlendFactor { + use mtl::MTLBlendFactor::*; + use wgt::BlendFactor as Bf; + + match factor { + Bf::Zero => Zero, + Bf::One => One, + Bf::Src => SourceColor, + Bf::OneMinusSrc => OneMinusSourceColor, + Bf::Dst => DestinationColor, + Bf::OneMinusDst => OneMinusDestinationColor, + Bf::SrcAlpha => SourceAlpha, + Bf::OneMinusSrcAlpha => OneMinusSourceAlpha, + Bf::DstAlpha => DestinationAlpha, + Bf::OneMinusDstAlpha => OneMinusDestinationAlpha, + Bf::Constant => BlendColor, + Bf::OneMinusConstant => OneMinusBlendColor, + //Bf::ConstantAlpha => BlendAlpha, + //Bf::OneMinusConstantAlpha => OneMinusBlendAlpha, + Bf::SrcAlphaSaturated => SourceAlphaSaturated, + //Bf::Src1 => Source1Color, + //Bf::OneMinusSrc1 => OneMinusSource1Color, + //Bf::Src1Alpha => Source1Alpha, + //Bf::OneMinusSrc1Alpha => OneMinusSource1Alpha, + } +} + +pub fn map_blend_op(operation: wgt::BlendOperation) -> mtl::MTLBlendOperation { + use mtl::MTLBlendOperation::*; + use wgt::BlendOperation as Bo; + + match operation { + Bo::Add => Add, + Bo::Subtract => Subtract, + Bo::ReverseSubtract => ReverseSubtract, + Bo::Min => Min, + Bo::Max => Max, + } +} + +pub fn map_blend_component( + component: &wgt::BlendComponent, +) -> ( + mtl::MTLBlendOperation, + mtl::MTLBlendFactor, + mtl::MTLBlendFactor, +) { + ( + map_blend_op(component.operation), + map_blend_factor(component.src_factor), + map_blend_factor(component.dst_factor), + ) +} + +pub fn map_vertex_format(format: wgt::VertexFormat) -> mtl::MTLVertexFormat { + use mtl::MTLVertexFormat::*; + use wgt::VertexFormat as Vf; + + match format { + Vf::Unorm8x2 => UChar2Normalized, + Vf::Snorm8x2 => Char2Normalized, + Vf::Uint8x2 => UChar2, + Vf::Sint8x2 => Char2, + Vf::Unorm8x4 => UChar4Normalized, + Vf::Snorm8x4 => Char4Normalized, + Vf::Uint8x4 => UChar4, + Vf::Sint8x4 => Char4, + Vf::Unorm16x2 => UShort2Normalized, + Vf::Snorm16x2 => Short2Normalized, + Vf::Uint16x2 => UShort2, + Vf::Sint16x2 => Short2, + Vf::Float16x2 => Half2, + Vf::Unorm16x4 => UShort4Normalized, + Vf::Snorm16x4 => Short4Normalized, + Vf::Uint16x4 => UShort4, + Vf::Sint16x4 => Short4, + Vf::Float16x4 => Half4, + Vf::Uint32 => UInt, + Vf::Sint32 => Int, + Vf::Float32 => Float, + Vf::Uint32x2 => UInt2, + Vf::Sint32x2 => Int2, + Vf::Float32x2 => Float2, + Vf::Uint32x3 => UInt3, + Vf::Sint32x3 => Int3, + Vf::Float32x3 => Float3, + Vf::Uint32x4 => UInt4, + Vf::Sint32x4 => Int4, + Vf::Float32x4 => Float4, + } +} + +pub fn map_step_mode(mode: wgt::InputStepMode) -> mtl::MTLVertexStepFunction { + match mode { + wgt::InputStepMode::Vertex => mtl::MTLVertexStepFunction::PerVertex, + wgt::InputStepMode::Instance => mtl::MTLVertexStepFunction::PerInstance, + } +} + +pub fn map_stencil_op(op: wgt::StencilOperation) -> mtl::MTLStencilOperation { + use mtl::MTLStencilOperation::*; + use wgt::StencilOperation as So; + + match op { + So::Keep => Keep, + So::Zero => Zero, + So::Replace => Replace, + So::IncrementClamp => IncrementClamp, + So::IncrementWrap => IncrementWrap, + So::DecrementClamp => DecrementClamp, + So::DecrementWrap => DecrementWrap, + So::Invert => Invert, + } +} + +pub fn map_winding(winding: wgt::FrontFace) -> mtl::MTLWinding { + match winding { + wgt::FrontFace::Cw => mtl::MTLWinding::Clockwise, + wgt::FrontFace::Ccw => mtl::MTLWinding::CounterClockwise, + } +} + +pub fn map_cull_mode(face: Option) -> mtl::MTLCullMode { + match face { + None => mtl::MTLCullMode::None, + Some(wgt::Face::Front) => mtl::MTLCullMode::Front, + Some(wgt::Face::Back) => mtl::MTLCullMode::Back, + } +} diff --git a/wgpu-hal/src/metal/device.rs b/wgpu-hal/src/metal/device.rs index 3274c623f0..975993c25b 100644 --- a/wgpu-hal/src/metal/device.rs +++ b/wgpu-hal/src/metal/device.rs @@ -1,9 +1,136 @@ -use std::{ptr, sync::Arc}; +use std::{ + ptr, + sync::{atomic, Arc}, + thread, time, +}; use super::conv; +use crate::aux::map_naga_stage; type DeviceResult = Result; +struct CompiledShader { + library: mtl::Library, + function: mtl::Function, + wg_size: mtl::MTLSize, + sized_bindings: Vec, +} + +fn create_stencil_desc( + face: &wgt::StencilFaceState, + read_mask: u32, + write_mask: u32, +) -> mtl::StencilDescriptor { + let desc = mtl::StencilDescriptor::new(); + desc.set_stencil_compare_function(conv::map_compare_function(face.compare)); + desc.set_read_mask(read_mask); + desc.set_write_mask(write_mask); + desc.set_stencil_failure_operation(conv::map_stencil_op(face.fail_op)); + desc.set_depth_failure_operation(conv::map_stencil_op(face.depth_fail_op)); + desc.set_depth_stencil_pass_operation(conv::map_stencil_op(face.pass_op)); + desc +} + +fn create_depth_stencil_desc(state: &wgt::DepthStencilState) -> mtl::DepthStencilDescriptor { + let desc = mtl::DepthStencilDescriptor::new(); + desc.set_depth_compare_function(conv::map_compare_function(state.depth_compare)); + desc.set_depth_write_enabled(state.depth_write_enabled); + let s = &state.stencil; + if s.is_enabled() { + let front_desc = create_stencil_desc(&s.front, s.read_mask, s.write_mask); + desc.set_front_face_stencil(Some(&front_desc)); + let back_desc = create_stencil_desc(&s.back, s.read_mask, s.write_mask); + desc.set_front_face_stencil(Some(&back_desc)); + } + desc +} + +impl super::Device { + fn load_shader( + &self, + stage: &crate::ProgrammableStage, + layout: &super::PipelineLayout, + primitive_class: mtl::MTLPrimitiveTopologyClass, + naga_stage: naga::ShaderStage, + ) -> Result { + let stage_bit = map_naga_stage(naga_stage); + let pipeline_options = naga::back::msl::PipelineOptions { + allow_point_size: match primitive_class { + mtl::MTLPrimitiveTopologyClass::Point => true, + _ => false, + }, + }; + + let module = &stage.module.raw.module; + let (source, info) = naga::back::msl::write_string( + module, + &stage.module.raw.info, + &layout.naga_options, + &pipeline_options, + ) + .map_err(|e| crate::PipelineError::Linkage(stage_bit, format!("MSL: {:?}", e)))?; + + let options = mtl::CompileOptions::new(); + options.set_language_version(self.shared.private_caps.msl_version); + + let library = self + .shared + .device + .lock() + .new_library_with_source(source.as_ref(), &options) + .map_err(|err| { + log::warn!("Naga generated shader:\n{}", source); + crate::PipelineError::Linkage(stage_bit, format!("Metal: {}", err)) + })?; + + // collect sizes indices + let mut sized_bindings = Vec::new(); + for (_handle, var) in module.global_variables.iter() { + if let naga::TypeInner::Struct { ref members, .. } = module.types[var.ty].inner { + if let Some(member) = members.last() { + if let naga::TypeInner::Array { + size: naga::ArraySize::Dynamic, + .. + } = module.types[member.ty].inner + { + // Note: unwraps are fine, since the MSL is already generated + let br = var.binding.clone().unwrap(); + sized_bindings.push(br); + } + } + } + } + + let (ep, internal_name) = module + .entry_points + .iter() + .zip(info.entry_point_names) + .find(|&(ep, _)| ep.stage == naga_stage && ep.name == stage.entry_point) + .ok_or(crate::PipelineError::EntryPoint(naga_stage))?; + + let name = internal_name + .as_ref() + .map_err(|e| crate::PipelineError::Linkage(stage_bit, format!("{}", e)))?; + let wg_size = mtl::MTLSize { + width: ep.workgroup_size[0] as _, + height: ep.workgroup_size[1] as _, + depth: ep.workgroup_size[2] as _, + }; + + let function = library.get_function(name, None).map_err(|e| { + log::error!("get_function: {:?}", e); + crate::PipelineError::EntryPoint(naga_stage) + })?; + + Ok(CompiledShader { + library, + function, + wg_size, + sized_bindings, + }) + } +} + impl crate::Device for super::Device { unsafe fn create_buffer(&self, desc: &crate::BufferDescriptor) -> DeviceResult { let map_read = desc.usage.contains(crate::BufferUse::MAP_READ); @@ -15,7 +142,10 @@ impl crate::Device for super::Device { } else { mtl::MTLResourceOptions::StorageModePrivate }; - options.set(mtl::MTLResourceOptions::CPUCacheModeDefaultCache, map_read); + options.set( + mtl::MTLResourceOptions::CPUCacheModeWriteCombined, + map_write, + ); //TODO: HazardTrackingModeUntracked @@ -111,16 +241,16 @@ impl crate::Device for super::Device { texture: &super::Texture, desc: &crate::TextureViewDescriptor, ) -> DeviceResult { - let mtl_format = self.shared.private_caps.map_format(desc.format); + let raw_format = self.shared.private_caps.map_format(desc.format); - let mtl_type = if texture.raw_type == mtl::MTLTextureType::D2Multisample { + let raw_type = if texture.raw_type == mtl::MTLTextureType::D2Multisample { texture.raw_type } else { conv::map_texture_view_dimension(desc.dimension) }; - let raw = if mtl_format == texture.raw_format - && mtl_type == texture.raw_type + let raw = if raw_format == texture.raw_format + && raw_type == texture.raw_type && desc.range == wgt::ImageSubresourceRange::default() { // Some images are marked as framebuffer-only, and we can't create aliases of them. @@ -137,8 +267,8 @@ impl crate::Device for super::Device { }; texture.raw.new_texture_view_from_slice( - mtl_format, - mtl_type, + raw_format, + raw_type, mtl::NSRange { location: desc.range.base_mip_level as _, length: mip_level_count as _, @@ -205,21 +335,19 @@ impl crate::Device for super::Device { unsafe fn create_command_buffer( &self, desc: &crate::CommandBufferDescriptor, - ) -> DeviceResult { - Ok(super::Encoder) + ) -> DeviceResult { + let raw = self.shared.create_command_buffer().to_owned(); + Ok(super::CommandBuffer { raw }) } - unsafe fn destroy_command_buffer(&self, cmd_buf: super::Encoder) {} + unsafe fn destroy_command_buffer(&self, _cmd_buf: super::CommandBuffer) {} unsafe fn create_bind_group_layout( &self, desc: &crate::BindGroupLayoutDescriptor, ) -> DeviceResult { - let map = desc - .entries - .iter() - .cloned() - .map(|entry| (entry.binding, entry)) - .collect(); + let mut map = desc.entries.to_vec(); + map.sort_by_key(|e| e.binding); + Ok(super::BindGroupLayout { entries: Arc::new(map), }) @@ -239,11 +367,6 @@ impl crate::Device for super::Device { sizes_buffer: Option, sizes_count: u8, } - impl StageInfo { - fn stage_bit(&self) -> wgt::ShaderStage { - crate::aux::map_naga_stage(self.stage) - } - } let mut stage_data = super::NAGA_STAGES.map(|&stage| StageInfo { stage, @@ -259,7 +382,7 @@ impl crate::Device for super::Device { // First, place the push constants for info in stage_data.iter_mut() { for pcr in desc.push_constant_ranges { - if pcr.stages.contains(info.stage_bit()) { + if pcr.stages.contains(map_naga_stage(info.stage)) { debug_assert_eq!(pcr.range.end % 4, 0); info.pc_limit = (pcr.range.end / 4).max(info.pc_limit); } @@ -287,7 +410,7 @@ impl crate::Device for super::Device { let mut sized_buffer_bindings = Vec::new(); let base_resource_indices = stage_data.map(|info| info.counters.clone()); - for entry in bgl.entries.values() { + for entry in bgl.entries.iter() { match entry.ty { wgt::BindingType::Buffer { ty, @@ -296,7 +419,7 @@ impl crate::Device for super::Device { } => { if has_dynamic_offset { dynamic_buffers.push(stage_data.map(|info| { - if entry.visibility.contains(info.stage_bit()) { + if entry.visibility.contains(map_naga_stage(info.stage)) { info.counters.buffers } else { !0 @@ -307,7 +430,7 @@ impl crate::Device for super::Device { wgt::BufferBindingType::Storage { .. } => { sized_buffer_bindings.push((entry.binding, entry.visibility)); for info in stage_data.iter_mut() { - if entry.visibility.contains(info.stage_bit()) { + if entry.visibility.contains(map_naga_stage(info.stage)) { info.sizes_count += 1; } } @@ -318,7 +441,7 @@ impl crate::Device for super::Device { } for info in stage_data.iter_mut() { - if !entry.visibility.contains(info.stage_bit()) { + if !entry.visibility.contains(map_naga_stage(info.stage)) { continue; } @@ -418,12 +541,12 @@ impl crate::Device for super::Device { naga_options, bind_group_infos, push_constants_infos: stage_data.map(|info| { - info.pc_buffer - .map(|buffer_index| super::PushConstantsStage { - count: info.pc_limit, - buffer_index, - }) + info.pc_buffer.map(|buffer_index| super::PushConstantsInfo { + count: info.pc_limit, + buffer_index, + }) }), + total_counters: stage_data.map(|info| info.counters.clone()), }) } unsafe fn destroy_pipeline_layout(&self, _pipeline_layout: super::PipelineLayout) {} @@ -432,15 +555,10 @@ impl crate::Device for super::Device { &self, desc: &crate::BindGroupDescriptor, ) -> DeviceResult { - //TODO: avoid heap allocation - let mut entries = desc.entries.to_vec(); - entries.sort_by_key(|e| e.binding); - let mut bg = super::BindGroup::default(); for (&stage, counter) in super::NAGA_STAGES.iter().zip(bg.counters.iter_mut()) { - let stage_bit = crate::aux::map_naga_stage(stage); - for entry in entries.iter() { - let layout = &desc.layout.entries[&entry.binding]; + let stage_bit = map_naga_stage(stage); + for (entry, layout) in desc.entries.iter().zip(desc.layout.entries.iter()) { if !layout.visibility.contains(stage_bit) { continue; } @@ -476,47 +594,310 @@ impl crate::Device for super::Device { &self, desc: &crate::ShaderModuleDescriptor, shader: crate::NagaShader, - ) -> Result { - Ok(Resource) + ) -> Result { + Ok(super::ShaderModule { raw: shader }) } - unsafe fn destroy_shader_module(&self, module: Resource) {} + unsafe fn destroy_shader_module(&self, _module: super::ShaderModule) {} + unsafe fn create_render_pipeline( &self, - desc: &crate::RenderPipelineDescriptor, - ) -> Result { - Ok(Resource) + desc: &crate::RenderPipelineDescriptor, + ) -> Result { + let descriptor = mtl::RenderPipelineDescriptor::new(); + let (primitive_class, raw_primitive_type) = + conv::map_primitive_topology(desc.primitive.topology); + + let vs = self.load_shader( + &desc.vertex_stage, + desc.layout, + primitive_class, + naga::ShaderStage::Vertex, + )?; + + descriptor.set_vertex_function(Some(&vs.function)); + + // Fragment shader + let (fs_lib, fs_sized_bindings) = match desc.fragment_stage { + Some(ref stage) => { + let compiled = self.load_shader( + stage, + desc.layout, + primitive_class, + naga::ShaderStage::Fragment, + )?; + descriptor.set_fragment_function(Some(&compiled.function)); + (Some(compiled.library), compiled.sized_bindings) + } + None => { + // TODO: This is a workaround for what appears to be a Metal validation bug + // A pixel format is required even though no attachments are provided + if desc.color_targets.is_empty() && desc.depth_stencil.is_none() { + descriptor.set_depth_attachment_pixel_format(mtl::MTLPixelFormat::Depth32Float); + } + (None, Vec::new()) + } + }; + + for (i, ct) in desc.color_targets.iter().enumerate() { + let at_descriptor = descriptor.color_attachments().object_at(i as u64).unwrap(); + + let raw_format = self.shared.private_caps.map_format(ct.format); + at_descriptor.set_pixel_format(raw_format); + at_descriptor.set_write_mask(conv::map_color_write(ct.write_mask)); + + if let Some(ref blend) = ct.blend { + at_descriptor.set_blending_enabled(true); + let (color_op, color_src, color_dst) = conv::map_blend_component(&blend.color); + let (alpha_op, alpha_src, alpha_dst) = conv::map_blend_component(&blend.alpha); + + at_descriptor.set_rgb_blend_operation(color_op); + at_descriptor.set_source_rgb_blend_factor(color_src); + at_descriptor.set_destination_rgb_blend_factor(color_dst); + + at_descriptor.set_alpha_blend_operation(alpha_op); + at_descriptor.set_source_alpha_blend_factor(alpha_src); + at_descriptor.set_destination_alpha_blend_factor(alpha_dst); + } + } + + let (raw_depth_stencil, depth_bias) = match desc.depth_stencil { + Some(ref ds) => { + let raw_format = self.shared.private_caps.map_format(ds.format); + let aspects = crate::FormatAspect::from(ds.format); + if aspects.contains(crate::FormatAspect::DEPTH) { + descriptor.set_depth_attachment_pixel_format(raw_format); + } + if aspects.contains(crate::FormatAspect::STENCIL) { + descriptor.set_stencil_attachment_pixel_format(raw_format); + } + + let ds_descriptor = create_depth_stencil_desc(ds); + let raw = self + .shared + .device + .lock() + .new_depth_stencil_state(&ds_descriptor); + (Some(raw), ds.bias) + } + None => (None, wgt::DepthBiasState::default()), + }; + + if desc.layout.total_counters.vs.buffers + (desc.vertex_buffers.len() as u32) + > self.shared.private_caps.max_buffers_per_stage + { + let msg = format!( + "pipeline needs too many buffers in the vertex stage: {} vertex and {} layout", + desc.vertex_buffers.len(), + desc.layout.total_counters.vs.buffers + ); + return Err(crate::PipelineError::Linkage(wgt::ShaderStage::VERTEX, msg)); + } + + if !desc.vertex_buffers.is_empty() { + let vertex_descriptor = mtl::VertexDescriptor::new(); + for (i, vb) in desc.vertex_buffers.iter().enumerate() { + let buffer_index = + self.shared.private_caps.max_buffers_per_stage as u64 - 1 - i as u64; + let buffer_desc = vertex_descriptor.layouts().object_at(buffer_index).unwrap(); + + buffer_desc.set_stride(vb.array_stride); + buffer_desc.set_step_function(conv::map_step_mode(vb.step_mode)); + + for (j, at) in vb.attributes.iter().enumerate() { + let attribute_desc = + vertex_descriptor.attributes().object_at(i as u64).unwrap(); + attribute_desc.set_format(conv::map_vertex_format(at.format)); + attribute_desc.set_buffer_index(buffer_index); + attribute_desc.set_offset(at.offset); + } + } + descriptor.set_vertex_descriptor(Some(&vertex_descriptor)); + } + + if desc.multisample.count != 1 { + //TODO: handle sample mask + descriptor.set_sample_count(desc.multisample.count as u64); + descriptor.set_alpha_to_coverage_enabled(desc.multisample.alpha_to_coverage_enabled); + //descriptor.set_alpha_to_one_enabled(desc.multisample.alpha_to_one_enabled); + } + + if let Some(name) = desc.label { + descriptor.set_label(name); + } + + let raw = self + .shared + .device + .lock() + .new_render_pipeline_state(&descriptor) + .map_err(|e| { + crate::PipelineError::Linkage( + wgt::ShaderStage::VERTEX | wgt::ShaderStage::FRAGMENT, + format!("new_render_pipeline_state: {:?}", e), + ) + })?; + + Ok(super::RenderPipeline { + raw, + vs_lib: vs.library, + fs_lib, + vs_info: super::PipelineStageInfo { + push_constants: desc.layout.push_constants_infos.vs, + sizes_slot: desc.layout.naga_options.per_stage_map.vs.sizes_buffer, + sized_bindings: vs.sized_bindings, + }, + fs_info: super::PipelineStageInfo { + push_constants: desc.layout.push_constants_infos.fs, + sizes_slot: desc.layout.naga_options.per_stage_map.fs.sizes_buffer, + sized_bindings: fs_sized_bindings, + }, + raw_primitive_type, + raw_front_winding: conv::map_winding(desc.primitive.front_face), + raw_cull_mode: conv::map_cull_mode(desc.primitive.cull_mode), + raw_depth_clip_mode: if self.features.contains(wgt::Features::DEPTH_CLAMPING) { + Some(if desc.primitive.clamp_depth { + mtl::MTLDepthClipMode::Clamp + } else { + mtl::MTLDepthClipMode::Clip + }) + } else { + None + }, + raw_depth_stencil, + depth_bias, + }) } - unsafe fn destroy_render_pipeline(&self, pipeline: Resource) {} + unsafe fn destroy_render_pipeline(&self, _pipeline: super::RenderPipeline) {} + unsafe fn create_compute_pipeline( &self, - desc: &crate::ComputePipelineDescriptor, - ) -> Result { - Ok(Resource) - } - unsafe fn destroy_compute_pipeline(&self, pipeline: Resource) {} + desc: &crate::ComputePipelineDescriptor, + ) -> Result { + let descriptor = mtl::ComputePipelineDescriptor::new(); - unsafe fn create_query_set(&self, desc: &wgt::QuerySetDescriptor) -> DeviceResult { - Ok(Resource) + let cs = self.load_shader( + &desc.stage, + desc.layout, + mtl::MTLPrimitiveTopologyClass::Unspecified, + naga::ShaderStage::Compute, + )?; + descriptor.set_compute_function(Some(&cs.function)); + if let Some(name) = desc.label { + descriptor.set_label(name); + } + + let raw = self + .shared + .device + .lock() + .new_compute_pipeline_state(&descriptor) + .map_err(|e| { + crate::PipelineError::Linkage( + wgt::ShaderStage::COMPUTE, + format!("new_compute_pipeline_state: {:?}", e), + ) + })?; + + Ok(super::ComputePipeline { + raw, + cs_info: super::PipelineStageInfo { + push_constants: desc.layout.push_constants_infos.cs, + sizes_slot: desc.layout.naga_options.per_stage_map.cs.sizes_buffer, + sized_bindings: cs.sized_bindings, + }, + cs_lib: cs.library, + work_group_size: cs.wg_size, + }) } - unsafe fn destroy_query_set(&self, set: Resource) {} - unsafe fn create_fence(&self) -> DeviceResult { - Ok(Resource) + unsafe fn destroy_compute_pipeline(&self, _pipeline: super::ComputePipeline) {} + + unsafe fn create_query_set( + &self, + desc: &wgt::QuerySetDescriptor, + ) -> DeviceResult { + match desc.ty { + wgt::QueryType::Occlusion => { + let size = desc.count as u64 * 8; + let options = mtl::MTLResourceOptions::empty(); + //TODO: HazardTrackingModeUntracked + let raw_buffer = self.shared.device.lock().new_buffer(size, options); + raw_buffer.set_label("_QuerySet"); + Ok(super::QuerySet { raw_buffer }) + } + wgt::QueryType::Timestamp | wgt::QueryType::PipelineStatistics(_) => { + Err(crate::DeviceError::OutOfMemory) + } + } } - unsafe fn destroy_fence(&self, fence: Resource) {} - unsafe fn get_fence_value(&self, fence: &Resource) -> DeviceResult { - Ok(0) + unsafe fn destroy_query_set(&self, _set: super::QuerySet) {} + + unsafe fn create_fence(&self) -> DeviceResult { + Ok(super::Fence { + completed_value: atomic::AtomicU64::new(0), + pending_command_buffers: Vec::new(), + }) + } + unsafe fn destroy_fence(&self, _fence: super::Fence) {} + unsafe fn get_fence_value(&self, fence: &super::Fence) -> DeviceResult { + let mut max_value = fence.completed_value.load(atomic::Ordering::Acquire); + for &(value, ref cmd_buf) in fence.pending_command_buffers.iter() { + if cmd_buf.status() == mtl::MTLCommandBufferStatus::Completed { + max_value = value; + } + } + Ok(max_value) } unsafe fn wait( &self, - fence: &Resource, - value: crate::FenceValue, + fence: &super::Fence, + wait_value: crate::FenceValue, timeout_ms: u32, ) -> DeviceResult { - Ok(true) + if wait_value <= fence.completed_value.load(atomic::Ordering::Acquire) { + return Ok(true); + } + + let cmd_buf = match fence + .pending_command_buffers + .iter() + .find(|&&(value, _)| value == wait_value) + { + Some(&(_, ref cmd_buf)) => cmd_buf, + None => { + log::error!("No active command buffers for fence value {}", wait_value); + return Err(crate::DeviceError::Lost); + } + }; + + let start = time::Instant::now(); + loop { + if let mtl::MTLCommandBufferStatus::Completed = cmd_buf.status() { + return Ok(true); + } + if start.elapsed().as_millis() >= timeout_ms as u128 { + return Ok(false); + } + thread::sleep(time::Duration::from_millis(1)); + } } unsafe fn start_capture(&self) -> bool { - false + if !self.shared.private_caps.supports_capture_manager { + return false; + } + let device = self.shared.device.lock(); + let shared_capture_manager = mtl::CaptureManager::shared(); + let default_capture_scope = shared_capture_manager.new_capture_scope_with_device(&device); + shared_capture_manager.set_default_capture_scope(&default_capture_scope); + shared_capture_manager.start_capture_with_scope(&default_capture_scope); + default_capture_scope.begin_scope(); + true + } + unsafe fn stop_capture(&self) { + let shared_capture_manager = mtl::CaptureManager::shared(); + if let Some(default_capture_scope) = shared_capture_manager.default_capture_scope() { + default_capture_scope.end_scope(); + } + shared_capture_manager.stop_capture(); } - unsafe fn stop_capture(&self) {} } diff --git a/wgpu-hal/src/metal/mod.rs b/wgpu-hal/src/metal/mod.rs index 68ec6b2d64..6a41ff2d3d 100644 --- a/wgpu-hal/src/metal/mod.rs +++ b/wgpu-hal/src/metal/mod.rs @@ -4,7 +4,12 @@ mod conv; mod device; mod surface; -use std::{iter, ops, ptr::NonNull, sync::Arc, thread}; +use std::{ + iter, ops, + ptr::NonNull, + sync::{atomic, Arc}, + thread, +}; use arrayvec::ArrayVec; use foreign_types::ForeignTypeRef as _; @@ -12,9 +17,6 @@ use parking_lot::Mutex; #[derive(Clone)] pub struct Api; -pub struct Encoder; -#[derive(Debug)] -pub struct Resource; type ResourceIndex = u32; @@ -25,22 +27,22 @@ impl crate::Api for Api { type Queue = Queue; type Device = Device; - type CommandBuffer = Encoder; + type CommandBuffer = CommandBuffer; type Buffer = Buffer; type Texture = Texture; type SurfaceTexture = SurfaceTexture; type TextureView = TextureView; type Sampler = Sampler; - type QuerySet = Resource; - type Fence = Resource; + type QuerySet = QuerySet; + type Fence = Fence; type BindGroupLayout = BindGroupLayout; type BindGroup = BindGroup; type PipelineLayout = PipelineLayout; - type ShaderModule = Resource; - type RenderPipeline = Resource; - type ComputePipeline = Resource; + type ShaderModule = ShaderModule; + type RenderPipeline = RenderPipeline; + type ComputePipeline = ComputePipeline; } pub struct Instance {} @@ -181,6 +183,7 @@ struct PrivateCapabilities { sample_count_mask: u8, supports_debug_markers: bool, supports_binary_archives: bool, + supports_capture_manager: bool, can_set_maximum_drawables_count: bool, can_set_display_sync: bool, can_set_next_drawable_timeout: bool, @@ -194,10 +197,17 @@ struct PrivateDisabilities { broken_layered_clear_image: bool, } +#[derive(Debug, Default)] +struct Settings { + retain_command_buffer_references: bool, +} + struct AdapterShared { device: Mutex, + queue: Mutex, disabilities: PrivateDisabilities, private_caps: PrivateCapabilities, + settings: Settings, } impl AdapterShared { @@ -208,18 +218,29 @@ impl AdapterShared { Self { disabilities: PrivateDisabilities::new(&device), private_caps: PrivateCapabilities::new(&device), + queue: Mutex::new(device.new_command_queue()), device: Mutex::new(device), + settings: Settings::default(), } } + + fn create_command_buffer(&self) -> &mtl::CommandBufferRef { + let queue = self.queue.lock(); + objc::rc::autoreleasepool(|| { + if self.settings.retain_command_buffer_references { + queue.new_command_buffer() + } else { + queue.new_command_buffer_with_unretained_references() + } + }) + } } struct Adapter { shared: Arc, } -struct Queue { - raw: mtl::CommandQueue, -} +struct Queue {} struct Device { shared: Arc, @@ -256,7 +277,7 @@ impl crate::Queue for Queue { unsafe fn submit( &mut self, command_buffers: I, - signal_fence: Option<(&Resource, crate::FenceValue)>, + signal_fence: Option<(&Fence, crate::FenceValue)>, ) -> Result<(), crate::DeviceError> { Ok(()) } @@ -325,11 +346,10 @@ impl Sampler { } } -type BindingMap = fxhash::FxHashMap; - #[derive(Debug)] pub struct BindGroupLayout { - entries: Arc, + /// Sorted list of BGL entries. + entries: Arc>, } #[derive(Clone, Debug, Default)] @@ -393,7 +413,7 @@ struct BindGroupLayoutInfo { } #[derive(Copy, Clone, Debug, Eq, PartialEq)] -struct PushConstantsStage { +struct PushConstantsInfo { count: u32, buffer_index: ResourceIndex, } @@ -402,7 +422,8 @@ struct PushConstantsStage { pub struct PipelineLayout { naga_options: naga::back::msl::Options, bind_group_infos: ArrayVec<[BindGroupLayoutInfo; crate::MAX_BIND_GROUPS]>, - push_constants_infos: MultiStageData>, + push_constants_infos: MultiStageData>, + total_counters: MultiStageResourceCounters, } trait AsNative { @@ -468,3 +489,72 @@ pub struct BindGroup { unsafe impl Send for BindGroup {} unsafe impl Sync for BindGroup {} + +#[derive(Debug)] +pub struct ShaderModule { + raw: crate::NagaShader, +} + +#[derive(Debug, Default)] +struct PipelineStageInfo { + push_constants: Option, + sizes_slot: Option, + sized_bindings: Vec, +} + +impl PipelineStageInfo { + fn clear(&mut self) { + self.push_constants = None; + self.sizes_slot = None; + self.sized_bindings.clear(); + } + + fn assign_from(&mut self, other: &Self) { + self.push_constants = other.push_constants; + self.sizes_slot = other.sizes_slot; + self.sized_bindings.clear(); + self.sized_bindings.extend_from_slice(&other.sized_bindings); + } +} + +pub struct RenderPipeline { + raw: mtl::RenderPipelineState, + vs_lib: mtl::Library, + fs_lib: Option, + vs_info: PipelineStageInfo, + fs_info: PipelineStageInfo, + raw_primitive_type: mtl::MTLPrimitiveType, + raw_front_winding: mtl::MTLWinding, + raw_cull_mode: mtl::MTLCullMode, + raw_depth_clip_mode: Option, + raw_depth_stencil: Option, + depth_bias: wgt::DepthBiasState, +} + +pub struct ComputePipeline { + raw: mtl::ComputePipelineState, + cs_lib: mtl::Library, + cs_info: PipelineStageInfo, + work_group_size: mtl::MTLSize, +} + +#[derive(Debug)] +pub struct QuerySet { + raw_buffer: mtl::Buffer, +} + +unsafe impl Send for QuerySet {} +unsafe impl Sync for QuerySet {} + +#[derive(Debug)] +pub struct Fence { + completed_value: atomic::AtomicU64, + pending_command_buffers: Vec<(crate::FenceValue, mtl::CommandBuffer)>, +} + +unsafe impl Send for Fence {} +unsafe impl Sync for Fence {} + +pub struct CommandBuffer { + raw: mtl::CommandBuffer, +} diff --git a/wgpu-hal/src/metal/surface.rs b/wgpu-hal/src/metal/surface.rs index 87c32bef9c..6ac31ec0ff 100644 --- a/wgpu-hal/src/metal/surface.rs +++ b/wgpu-hal/src/metal/surface.rs @@ -141,7 +141,7 @@ impl super::Surface { Self::new(None, layer.to_owned()) } - fn dimensions(&self) -> wgt::Extent3d { + pub(super) fn dimensions(&self) -> wgt::Extent3d { let (size, scale): (mtl::CGSize, mtl::CGFloat) = match self.view { Some(view) if !cfg!(target_os = "macos") => unsafe { let bounds: CGRect = msg_send![view.as_ptr(), bounds]; diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index ca96d637ab..3cccf262ee 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -3065,6 +3065,8 @@ pub struct QuerySetDescriptor { #[cfg_attr(feature = "trace", derive(serde::Serialize))] #[cfg_attr(feature = "replay", derive(serde::Deserialize))] pub enum QueryType { + /// Query returns a single 64-bit number, serving as an occlusion boolean. + Occlusion, /// Query returns up to 5 64-bit numbers based on the given flags. /// /// See [`PipelineStatisticsTypes`]'s documentation for more information