From 3067733310ea6674fb5f601764352824e5f17420 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Wed, 26 May 2021 22:07:00 -0700 Subject: [PATCH] Implement more of the API WIP. Goal is to get collatz running. --- Cargo.lock | 1 + piet-gpu-hal/Cargo.toml | 3 +- piet-gpu-hal/examples/collatz.rs | 2 +- piet-gpu-hal/src/metal.rs | 179 ++++++++++++++++++++++++------- 4 files changed, 144 insertions(+), 41 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index ee015f0..19fa7db 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -875,6 +875,7 @@ dependencies = [ "ash-window", "bitflags", "block", + "cocoa-foundation", "metal", "objc", "raw-window-handle", diff --git a/piet-gpu-hal/Cargo.toml b/piet-gpu-hal/Cargo.toml index 5fbce83..1826f45 100644 --- a/piet-gpu-hal/Cargo.toml +++ b/piet-gpu-hal/Cargo.toml @@ -24,5 +24,6 @@ wio = "0.2.2" [target.'cfg(target_os="macos")'.dependencies] metal = "0.22" -objc = "0.2.4" +objc = "0.2.5" block = "0.1.6" +cocoa-foundation = "0.1" diff --git a/piet-gpu-hal/examples/collatz.rs b/piet-gpu-hal/examples/collatz.rs index 5dcce21..4220148 100644 --- a/piet-gpu-hal/examples/collatz.rs +++ b/piet-gpu-hal/examples/collatz.rs @@ -10,7 +10,7 @@ fn main() { let usage = BufferUsage::MAP_READ | BufferUsage::STORAGE; let src = (0..256).map(|x| x + 1).collect::>(); let buffer = session.create_buffer_init(&src, usage).unwrap(); - let code = ShaderCode::Spv(include_bytes!("./shader/collatz.spv")); + let code = ShaderCode::Msl(include_str!("./shader/collatz.msl")); let pipeline = session.create_simple_compute_pipeline(code, 1).unwrap(); let descriptor_set = session .create_simple_descriptor_set(&pipeline, &[&buffer]) diff --git a/piet-gpu-hal/src/metal.rs b/piet-gpu-hal/src/metal.rs index f9d67cc..f7a97d6 100644 --- a/piet-gpu-hal/src/metal.rs +++ b/piet-gpu-hal/src/metal.rs @@ -14,18 +14,28 @@ // // Also licensed under MIT license, at your choice. -use crate::{BufferUsage, Error}; +use cocoa_foundation::foundation::NSInteger; +use objc::rc::autoreleasepool; +use objc::runtime::Object; +use objc::{class, msg_send, sel, sel_impl}; + +use metal::MTLFeatureSet; + +use crate::{BufferUsage, Error, GpuInfo}; pub struct MtlInstance; pub struct MtlDevice { device: metal::Device, + cmd_queue: metal::CommandQueue, + gpu_info: GpuInfo, } pub struct MtlSurface; pub struct MtlSwapchain; +#[derive(Clone)] pub struct Buffer { buffer: metal::Buffer, pub(crate) size: u64, @@ -33,34 +43,60 @@ pub struct Buffer { pub struct Image; -pub struct Pipeline; - -pub struct DescriptorSet; - pub struct Fence; pub struct Semaphore; -pub struct CmdBuf; +pub struct CmdBuf { + cmd_buf: metal::CommandBuffer, +} pub struct QueryPool; pub struct PipelineBuilder; -pub struct DescriptorSetBuilder; +pub struct Pipeline(metal::ComputePipelineState); + +#[derive(Default)] +pub struct DescriptorSetBuilder(DescriptorSet); + +#[derive(Default)] +pub struct DescriptorSet { + buffers: Vec, +} impl MtlInstance { pub fn new( - window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>, + _window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>, ) -> Result<(MtlInstance, Option), Error> { Ok((MtlInstance, None)) } // TODO might do some enumeration of devices - pub fn device(&self, surface: Option<&MtlSurface>) -> Result { + pub fn device(&self, _surface: Option<&MtlSurface>) -> Result { if let Some(device) = metal::Device::system_default() { - Ok(MtlDevice { device }) + let cmd_queue = device.new_command_queue(); + let is_mac = device.supports_feature_set(MTLFeatureSet::macOS_GPUFamily1_v1); + let is_ios = device.supports_feature_set(MTLFeatureSet::iOS_GPUFamily1_v1); + let version = NSOperatingSystemVersion::get(); + + let use_staging_buffers = + if (is_mac && version.at_least(10, 15)) || (is_ios && version.at_least(13, 0)) { + !device.has_unified_memory() + } else { + !device.is_low_power() + }; + // TODO: these are conservative; we need to derive these from + // supports_feature_set queries. + let gpu_info = GpuInfo { + has_descriptor_indexing: false, + has_subgroups: false, + subgroup_size: None, + has_memory_model: false, + use_staging_buffers: use_staging_buffers, + }; + Ok(MtlDevice { device, cmd_queue, gpu_info }) } else { Err("can't create system default Metal device".into()) } @@ -103,14 +139,16 @@ impl crate::Device for MtlDevice { type ShaderSource = str; fn query_gpu_info(&self) -> crate::GpuInfo { - todo!() + self.gpu_info.clone() } fn create_buffer(&self, size: u64, usage: BufferUsage) -> Result { let options = if usage.contains(BufferUsage::MAP_READ) { - metal::MTLResourceOptions::StorageModeShared | metal::MTLResourceOptions::CPUCacheModeDefaultCache + metal::MTLResourceOptions::StorageModeShared + | metal::MTLResourceOptions::CPUCacheModeDefaultCache } else if usage.contains(BufferUsage::MAP_WRITE) { - metal::MTLResourceOptions::StorageModeShared | metal::MTLResourceOptions::CPUCacheModeWriteCombined + metal::MTLResourceOptions::StorageModeShared + | metal::MTLResourceOptions::CPUCacheModeWriteCombined } else { metal::MTLResourceOptions::StorageModePrivate }; @@ -119,14 +157,12 @@ impl crate::Device for MtlDevice { } unsafe fn destroy_buffer(&self, buffer: &Self::Buffer) -> Result<(), Error> { - todo!() + // This defers dropping until the buffer object is dropped. We probably need + // to rethink buffer lifetime if descriptor sets can retain references. + Ok(()) } - unsafe fn create_image2d( - &self, - width: u32, - height: u32, - ) -> Result { + unsafe fn create_image2d(&self, width: u32, height: u32) -> Result { todo!() } @@ -135,23 +171,28 @@ impl crate::Device for MtlDevice { } unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder { - todo!() + PipelineBuilder } unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder { - todo!() + DescriptorSetBuilder::default() } fn create_cmd_buf(&self) -> Result { - todo!() + // consider new_command_buffer_with_unretained_references for performance + let cmd_buf = self.cmd_queue.new_command_buffer(); + let cmd_buf = autoreleasepool(|| cmd_buf.to_owned()); + Ok(CmdBuf { cmd_buf }) } fn create_query_pool(&self, n_queries: u32) -> Result { - todo!() + // TODO + Ok(QueryPool) } unsafe fn fetch_query_pool(&self, pool: &Self::QueryPool) -> Result, Error> { - todo!() + // TODO + Ok(Vec::new()) } unsafe fn run_cmd_bufs( @@ -175,7 +216,11 @@ impl crate::Device for MtlDevice { if contents_ptr.is_null() { return Err("probably trying to read from private buffer".into()); } - std::ptr::copy_nonoverlapping((contents_ptr as *const u8).add(offset as usize), dst, size as usize); + std::ptr::copy_nonoverlapping( + (contents_ptr as *const u8).add(offset as usize), + dst, + size as usize, + ); Ok(()) } @@ -190,7 +235,11 @@ impl crate::Device for MtlDevice { if contents_ptr.is_null() { return Err("probably trying to write to private buffer".into()); } - std::ptr::copy_nonoverlapping(contents, (contents_ptr as *mut u8).add(offset as usize), size as usize); + std::ptr::copy_nonoverlapping( + contents, + (contents_ptr as *mut u8).add(offset as usize), + size as usize, + ); Ok(()) } @@ -230,7 +279,28 @@ impl crate::CmdBuf for CmdBuf { descriptor_set: &DescriptorSet, size: (u32, u32, u32), ) { - todo!() + let encoder = self.cmd_buf.new_compute_command_encoder(); + encoder.set_compute_pipeline_state(&pipeline.0); + let mut ix = 0; + for buffer in &descriptor_set.buffers { + encoder.set_buffer(ix, Some(&buffer.buffer), 0); + ix += 1; + } + // TODO: set images + let work_group_count = metal::MTLSize { + width: size.0 as u64, + height: size.1 as u64, + depth: size.2 as u64, + }; + // TODO: we need to pass this in explicitly. In gfx-hal, this is parsed from + // the spv before translation. + let work_group_size = metal::MTLSize { + width: 256, + height: 1, + depth: 1, + }; + encoder.dispatch_thread_groups(work_group_count, work_group_size); + encoder.end_encoding(); } unsafe fn memory_barrier(&mut self) { @@ -280,26 +350,35 @@ impl crate::CmdBuf for CmdBuf { } impl crate::PipelineBuilder for PipelineBuilder { - fn add_buffers(&mut self, n_buffers: u32) { - todo!() + fn add_buffers(&mut self, _n_buffers: u32) { + // My understanding is that Metal infers the pipeline layout from + // the source. } - fn add_images(&mut self, n_images: u32) { - todo!() + fn add_images(&mut self, _n_images: u32) { } - fn add_textures(&mut self, max_textures: u32) { - todo!() + fn add_textures(&mut self, _max_textures: u32) { } - unsafe fn create_compute_pipeline(self, device: &MtlDevice, code: &str) -> Result { - todo!() + unsafe fn create_compute_pipeline( + self, + device: &MtlDevice, + code: &str, + ) -> Result { + let options = metal::CompileOptions::new(); + // Probably want to set MSL version here. + let library = device.device.new_library_with_source(code, &options)?; + // This seems to be the default name from spirv-cross, but we may need to tweak. + let function = library.get_function("main0", None)?; + let pipeline = device.device.new_compute_pipeline_state_with_function(&function)?; + Ok(Pipeline(pipeline)) } } impl crate::DescriptorSetBuilder for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { - todo!() + self.0.buffers.extend(buffers.iter().copied().cloned()); } fn add_images(&mut self, images: &[&Image]) { @@ -311,7 +390,7 @@ impl crate::DescriptorSetBuilder for DescriptorSetBuilder { } unsafe fn build(self, device: &MtlDevice, pipeline: &Pipeline) -> Result { - todo!() + Ok(self.0) } } @@ -330,5 +409,27 @@ impl MtlSwapchain { semaphores: &[&Semaphore], ) -> Result { todo!() - } -} \ No newline at end of file + } +} + +#[repr(C)] +struct NSOperatingSystemVersion { + major: NSInteger, + minor: NSInteger, + patch: NSInteger, +} + +impl NSOperatingSystemVersion { + fn get() -> NSOperatingSystemVersion { + unsafe { + let process_info: *mut Object = msg_send![class!(NSProcessInfo), processInfo]; + msg_send![process_info, operatingSystemVersion] + } + } + + fn at_least(&self, major: u32, minor: u32) -> bool { + let major = major as NSInteger; + let minor = minor as NSInteger; + self.major > major || (self.major == major && self.minor >= minor) + } +}