Pipeline statistics API

blade-graphics/src/gles/pipeline.rs

@@ -345,6 +345,13 @@ impl crate::traits::ShaderDevice for super::Context {
         super::ComputePipeline { inner, wg_size }
     }
 
+    fn get_pipeline_statistics(
+        &self,
+        _pipeline: &super::ComputePipeline,
+    ) -> Vec<crate::PipelineExecutableInfo> {
+        Vec::new()
+    }
+
     fn destroy_compute_pipeline(&self, pipeline: &mut super::ComputePipeline) {
         unsafe {
             self.destroy_pipeline(&mut pipeline.inner);

blade-graphics/src/lib.rs

@@ -892,6 +892,26 @@ pub struct ComputePipelineDesc<'a> {
     pub compute: ShaderFunction<'a>,
 }
 
+/// A single statistic reported for a pipeline executable.
+#[derive(Clone, Debug)]
+pub struct PipelineStatistic {
+    /// Name of the statistic (e.g. "numUsedVgprs", "spilledSgprs").
+    pub name: String,
+    /// Human-readable description.
+    pub description: String,
+    /// Numeric value (integer statistics are converted to f64).
+    pub value: f64,
+}
+
+/// Statistics for one executable stage within a pipeline.
+#[derive(Clone, Debug)]
+pub struct PipelineExecutableInfo {
+    /// Name of the executable (e.g. "Compute").
+    pub name: String,
+    /// Individual statistics.
+    pub statistics: Vec<PipelineStatistic>,
+}
+
 /// Primitive type the input mesh is composed of.
 #[derive(Copy, Clone, Debug, Default, Hash, Eq, PartialEq)]
 pub enum PrimitiveTopology {

blade-graphics/src/metal/pipeline.rs

@@ -1,7 +1,7 @@
 use naga::back::msl;
 use objc2::{rc::Retained, runtime::ProtocolObject};
 use objc2_foundation::NSString;
-use objc2_metal::{self as metal, MTLDevice, MTLLibrary};
+use objc2_metal::{self as metal, MTLComputePipelineState, MTLDevice, MTLLibrary};
 
 fn map_blend_factor(factor: crate::BlendFactor) -> metal::MTLBlendFactor {
     use crate::BlendFactor as Bf;
@@ -382,6 +382,32 @@ impl crate::traits::ShaderDevice for super::Context {
         })
     }
 
+    fn get_pipeline_statistics(
+        &self,
+        pipeline: &super::ComputePipeline,
+    ) -> Vec<crate::PipelineExecutableInfo> {
+        vec![crate::PipelineExecutableInfo {
+            name: pipeline.name.clone(),
+            statistics: vec![
+                crate::PipelineStatistic {
+                    name: "maxTotalThreadsPerThreadgroup".to_string(),
+                    description: "Maximum threads per threadgroup for this pipeline".to_string(),
+                    value: pipeline.raw.maxTotalThreadsPerThreadgroup() as f64,
+                },
+                crate::PipelineStatistic {
+                    name: "threadExecutionWidth".to_string(),
+                    description: "SIMD-group width".to_string(),
+                    value: pipeline.raw.threadExecutionWidth() as f64,
+                },
+                crate::PipelineStatistic {
+                    name: "staticThreadgroupMemoryLength".to_string(),
+                    description: "Statically allocated threadgroup memory in bytes".to_string(),
+                    value: pipeline.raw.staticThreadgroupMemoryLength() as f64,
+                },
+            ],
+        }]
+    }
+
     fn destroy_compute_pipeline(&self, _pipeline: &mut super::ComputePipeline) {
         //TODO: is there a way to release?
     }

blade-graphics/src/traits.rs

@@ -39,6 +39,10 @@ pub trait ShaderDevice {
     fn destroy_compute_pipeline(&self, pipeline: &mut Self::ComputePipeline);
     fn create_render_pipeline(&self, desc: super::RenderPipelineDesc) -> Self::RenderPipeline;
     fn destroy_render_pipeline(&self, pipeline: &mut Self::RenderPipeline);
+    fn get_pipeline_statistics(
+        &self,
+        pipeline: &Self::ComputePipeline,
+    ) -> Vec<super::PipelineExecutableInfo>;
 }
 
 pub trait CommandDevice {

blade-graphics/src/vulkan/init.rs

@@ -100,6 +100,7 @@ struct AdapterCapabilities {
     max_inline_uniform_block_size: u32,
     buffer_marker: bool,
     shader_info: bool,
+    pipeline_executable_properties: bool,
     full_screen_exclusive: bool,
     external_memory: bool,
     timing: bool,
@@ -511,6 +512,8 @@ fn inspect_adapter(
 
     let buffer_marker = supported_extensions.contains(&vk::AMD_BUFFER_MARKER_NAME);
     let shader_info = supported_extensions.contains(&vk::AMD_SHADER_INFO_NAME);
+    let pipeline_executable_properties =
+        supported_extensions.contains(&vk::KHR_PIPELINE_EXECUTABLE_PROPERTIES_NAME);
     let full_screen_exclusive = supported_extensions.contains(&vk::EXT_FULL_SCREEN_EXCLUSIVE_NAME);
     let memory_budget = supported_extensions.contains(&vk::EXT_MEMORY_BUDGET_NAME);
 
@@ -541,6 +544,7 @@ fn inspect_adapter(
         max_inline_uniform_block_size,
         buffer_marker,
         shader_info,
+        pipeline_executable_properties,
         full_screen_exclusive,
         external_memory,
         timing,
@@ -933,6 +937,9 @@ impl super::Context {
             if capabilities.shader_info {
                 device_extensions.push(vk::AMD_SHADER_INFO_NAME);
             }
+            if capabilities.pipeline_executable_properties {
+                device_extensions.push(vk::KHR_PIPELINE_EXECUTABLE_PROPERTIES_NAME);
+            }
             if capabilities.full_screen_exclusive {
                 device_extensions.push(vk::EXT_FULL_SCREEN_EXCLUSIVE_NAME);
             }
@@ -1054,6 +1061,17 @@ impl super::Context {
                     .push_next(&mut vulkan_memory_model);
             }
 
+            let mut khr_pipeline_executable_properties;
+            if capabilities.pipeline_executable_properties {
+                khr_pipeline_executable_properties =
+                    vk::PhysicalDevicePipelineExecutablePropertiesFeaturesKHR {
+                        pipeline_executable_info: vk::TRUE,
+                        ..Default::default()
+                    };
+                device_create_info =
+                    device_create_info.push_next(&mut khr_pipeline_executable_properties);
+            }
+
             // TODO: Replace with ash typed struct once available.
             let mut khr_unified_image_layouts;
             if capabilities.unified_image_layouts {
@@ -1140,6 +1158,14 @@ impl super::Context {
             } else {
                 None
             },
+            pipeline_executable_properties: if capabilities.pipeline_executable_properties {
+                Some(khr::pipeline_executable_properties::Device::new(
+                    &instance.core,
+                    &device_core,
+                ))
+            } else {
+                None
+            },
             full_screen_exclusive: if desc.presentation && capabilities.full_screen_exclusive {
                 Some(ext::full_screen_exclusive::Device::new(
                     &instance.core,

blade-graphics/src/vulkan/mod.rs

@@ -64,6 +64,7 @@ struct Device {
     max_inline_uniform_block_size: u32,
     buffer_marker: Option<ash::amd::buffer_marker::Device>,
     shader_info: Option<ash::amd::shader_info::Device>,
+    pipeline_executable_properties: Option<ash::khr::pipeline_executable_properties::Device>,
     full_screen_exclusive: Option<ash::ext::full_screen_exclusive::Device>,
     #[cfg(target_os = "windows")]
     external_memory: Option<ash::khr::external_memory_win32::Device>,

blade-graphics/src/vulkan/pipeline.rs

@@ -399,10 +399,14 @@ impl crate::traits::ShaderDevice for super::Context {
 
         let layout = self.create_pipeline_layout(desc.data_layouts, &group_infos);
 
-        let create_info = vk::ComputePipelineCreateInfo::default()
+        let mut create_info = vk::ComputePipelineCreateInfo::default()
             .layout(layout.raw)
             .stage(cs.create_info);
 
+        if self.device.pipeline_executable_properties.is_some() {
+            create_info.flags |= vk::PipelineCreateFlags::CAPTURE_STATISTICS_KHR;
+        }
+
         let mut raw_vec = unsafe {
             self.device
                 .core
@@ -438,6 +442,73 @@ impl crate::traits::ShaderDevice for super::Context {
         }
     }
 
+    fn get_pipeline_statistics(
+        &self,
+        pipeline: &super::ComputePipeline,
+    ) -> Vec<crate::PipelineExecutableInfo> {
+        let Some(ref ext) = self.device.pipeline_executable_properties else {
+            return Vec::new();
+        };
+
+        let pipeline_info = vk::PipelineInfoKHR::default().pipeline(pipeline.raw);
+        let executables = match unsafe { ext.get_pipeline_executable_properties(&pipeline_info) } {
+            Ok(e) => e,
+            Err(_) => return Vec::new(),
+        };
+
+        executables
+            .iter()
+            .enumerate()
+            .map(|(i, exec)| {
+                let name = exec
+                    .name_as_c_str()
+                    .map(|s| s.to_string_lossy().into_owned())
+                    .unwrap_or_default();
+
+                let exec_info = vk::PipelineExecutableInfoKHR::default()
+                    .pipeline(pipeline.raw)
+                    .executable_index(i as u32);
+
+                let statistics = unsafe { ext.get_pipeline_executable_statistics(&exec_info) }
+                    .unwrap_or_default()
+                    .iter()
+                    .map(|stat| {
+                        let stat_name = stat
+                            .name_as_c_str()
+                            .map(|s| s.to_string_lossy().into_owned())
+                            .unwrap_or_default();
+                        let stat_desc = stat
+                            .description_as_c_str()
+                            .map(|s| s.to_string_lossy().into_owned())
+                            .unwrap_or_default();
+                        let value = unsafe {
+                            match stat.format {
+                                vk::PipelineExecutableStatisticFormatKHR::BOOL32 => {
+                                    if stat.value.b32 != 0 { 1.0 } else { 0.0 }
+                                }
+                                vk::PipelineExecutableStatisticFormatKHR::INT64 => {
+                                    stat.value.i64 as f64
+                                }
+                                vk::PipelineExecutableStatisticFormatKHR::UINT64 => {
+                                    stat.value.u64 as f64
+                                }
+                                vk::PipelineExecutableStatisticFormatKHR::FLOAT64 => stat.value.f64,
+                                _ => 0.0,
+                            }
+                        };
+                        crate::PipelineStatistic {
+                            name: stat_name,
+                            description: stat_desc,
+                            value,
+                        }
+                    })
+                    .collect();
+
+                crate::PipelineExecutableInfo { name, statistics }
+            })
+            .collect()
+    }
+
     fn destroy_compute_pipeline(&self, pipeline: &mut super::ComputePipeline) {
         self.destroy_pipeline_layout(&mut pipeline.layout);
         unsafe {