diff --git a/blade-graphics/src/gles/pipeline.rs b/blade-graphics/src/gles/pipeline.rs index fb3a55d5..65caa193 100644 --- a/blade-graphics/src/gles/pipeline.rs +++ b/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 { + Vec::new() + } + fn destroy_compute_pipeline(&self, pipeline: &mut super::ComputePipeline) { unsafe { self.destroy_pipeline(&mut pipeline.inner); diff --git a/blade-graphics/src/lib.rs b/blade-graphics/src/lib.rs index 771d422c..bddda544 100644 --- a/blade-graphics/src/lib.rs +++ b/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, +} + /// Primitive type the input mesh is composed of. #[derive(Copy, Clone, Debug, Default, Hash, Eq, PartialEq)] pub enum PrimitiveTopology { diff --git a/blade-graphics/src/metal/pipeline.rs b/blade-graphics/src/metal/pipeline.rs index 92bacddf..409d82fd 100644 --- a/blade-graphics/src/metal/pipeline.rs +++ b/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 { + 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? } diff --git a/blade-graphics/src/traits.rs b/blade-graphics/src/traits.rs index 8c721de7..e9d8086b 100644 --- a/blade-graphics/src/traits.rs +++ b/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; } pub trait CommandDevice { diff --git a/blade-graphics/src/vulkan/init.rs b/blade-graphics/src/vulkan/init.rs index d71fc24d..901229cb 100644 --- a/blade-graphics/src/vulkan/init.rs +++ b/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, diff --git a/blade-graphics/src/vulkan/mod.rs b/blade-graphics/src/vulkan/mod.rs index b9d88020..323e9cc4 100644 --- a/blade-graphics/src/vulkan/mod.rs +++ b/blade-graphics/src/vulkan/mod.rs @@ -64,6 +64,7 @@ struct Device { max_inline_uniform_block_size: u32, buffer_marker: Option, shader_info: Option, + pipeline_executable_properties: Option, full_screen_exclusive: Option, #[cfg(target_os = "windows")] external_memory: Option, diff --git a/blade-graphics/src/vulkan/pipeline.rs b/blade-graphics/src/vulkan/pipeline.rs index 54681f2f..93635cfe 100644 --- a/blade-graphics/src/vulkan/pipeline.rs +++ b/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 { + 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 {