Implement more of the API

WIP. Goal is to get collatz running.

Cargo.lock

@@ -875,6 +875,7 @@ dependencies = [
  "ash-window",
  "bitflags",
  "block",
+ "cocoa-foundation",
  "metal",
  "objc",
  "raw-window-handle",

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"

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::<Vec<u32>>();
         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])

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<Buffer>,
+}
 
 impl MtlInstance {
     pub fn new(
-        window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>,
+        _window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>,
     ) -> Result<(MtlInstance, Option<MtlSurface>), Error> {
         Ok((MtlInstance, None))
     }
 
     // TODO might do some enumeration of devices
 
-    pub fn device(&self, surface: Option<&MtlSurface>) -> Result<MtlDevice, Error> {
+    pub fn device(&self, _surface: Option<&MtlSurface>) -> Result<MtlDevice, Error> {
         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<Self::Buffer, Error> {
         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<Self::Image, Error> {
+    unsafe fn create_image2d(&self, width: u32, height: u32) -> Result<Self::Image, Error> {
         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<Self::CmdBuf, Error> {
-        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<Self::QueryPool, Error> {
-        todo!()
+        // TODO
+        Ok(QueryPool)
     }
 
     unsafe fn fetch_query_pool(&self, pool: &Self::QueryPool) -> Result<Vec<f64>, 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<MtlDevice> 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<MtlDevice> for CmdBuf {
 }
 
 impl crate::PipelineBuilder<MtlDevice> 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<Pipeline, Error> {
-        todo!()
+    unsafe fn create_compute_pipeline(
+        self,
+        device: &MtlDevice,
+        code: &str,
+    ) -> Result<Pipeline, Error> {
+        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<MtlDevice> 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<MtlDevice> for DescriptorSetBuilder {
     }
 
     unsafe fn build(self, device: &MtlDevice, pipeline: &Pipeline) -> Result<DescriptorSet, Error> {
-        todo!()
+        Ok(self.0)
     }
 }
 
@@ -330,5 +409,27 @@ impl MtlSwapchain {
         semaphores: &[&Semaphore],
     ) -> Result<bool, Error> {
         todo!()
-    }    
-}
+    }
+}
+
+#[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)
+    }
+}