Change the shape of cooperative matrix to support AMD

Cargo.toml

@@ -69,6 +69,7 @@ egui-winit = { version = "0.33", default-features = false, features = [
 ] }
 transform-gizmo-egui = "0.8"
 env_logger = "0.11"
+half = { version = "2", features = ["bytemuck"] }
 num_cpus = { workspace = true }
 glam = { workspace = true }
 log = { workspace = true }

blade-graphics/src/gles/mod.rs

@@ -453,7 +453,8 @@ impl Context {
             ray_query: crate::ShaderVisibility::empty(),
             sample_count_mask: 0x1 | 0x4, //TODO: accurate info
             dual_source_blending: false,
-            cooperative_matrix: false,
+            shader_float16: false,
+            cooperative_matrix: crate::CooperativeMatrix::default(),
         }
     }
 

blade-graphics/src/lib.rs

@@ -169,8 +169,8 @@ pub enum NotSupportedError {
 
 impl fmt::Display for NotSupportedError {
     fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
-        match self {
-            Self::Platform(e) => write!(f, "platform error: {}", e),
+        match *self {
+            Self::Platform(ref e) => write!(f, "platform error: {}", e),
             Self::NoSupportedDeviceFound => f.write_str("no supported device found"),
             Self::PlatformNotSupported => f.write_str("platform not supported"),
         }
@@ -196,7 +196,7 @@ pub enum DeviceError {
 
 impl fmt::Display for DeviceError {
     fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
-        match self {
+        match *self {
             Self::DeviceLost => f.write_str("device lost"),
             Self::OutOfMemory => f.write_str("out of memory"),
         }
@@ -216,6 +216,25 @@ pub struct MemoryStats {
     pub usage: u64,
 }
 
+/// Cooperative matrix support information.
+///
+/// Each field is a tile size (8 or 16), or 0 if that configuration
+/// is not supported. Naga supports square tiles only (8×8 and 16×16).
+#[derive(Clone, Copy, Debug, Default, PartialEq)]
+pub struct CooperativeMatrix {
+    /// Tile size for all-f32 operations.
+    pub f32_tile: u32,
+    /// Tile size for f16-input, f32-accumulator operations.
+    pub f16_tile: u32,
+}
+
+impl CooperativeMatrix {
+    /// Returns true if any cooperative matrix configuration is supported.
+    pub fn is_supported(&self) -> bool {
+        self.f32_tile > 0 || self.f16_tile > 0
+    }
+}
+
 #[derive(Clone, Debug, Default, PartialEq)]
 pub struct Capabilities {
     /// Support binding arrays of handles.
@@ -226,8 +245,10 @@ pub struct Capabilities {
     pub sample_count_mask: u32,
     /// Support for dual-source blending.
     pub dual_source_blending: bool,
-    /// Support for cooperative matrix operations.
-    pub cooperative_matrix: bool,
+    /// Support for 16-bit floating-point types in shaders.
+    pub shader_float16: bool,
+    /// Cooperative matrix support.
+    pub cooperative_matrix: CooperativeMatrix,
 }
 
 #[derive(Clone, Debug, Default)]

blade-graphics/src/metal/mod.rs

@@ -538,9 +538,19 @@ impl Context {
                 .filter(|&count| device.supportsTextureSampleCount(count as _))
                 .sum(),
             dual_source_blending: true,
-            cooperative_matrix: device.supportsFamily(metal::MTLGPUFamily::Apple7)
+            // Metal Shading Language supports half-precision floats on all supported devices.
+            shader_float16: true,
+            cooperative_matrix: if device.supportsFamily(metal::MTLGPUFamily::Apple7)
                 || device.supportsFamily(metal::MTLGPUFamily::Mac2)
-                || device.supportsFamily(metal::MTLGPUFamily::Metal3),
+                || device.supportsFamily(metal::MTLGPUFamily::Metal3)
+            {
+                crate::CooperativeMatrix {
+                    f32_tile: 8,
+                    f16_tile: 0,
+                }
+            } else {
+                crate::CooperativeMatrix::default()
+            },
         }
     }
 

blade-graphics/src/shader.rs

@@ -39,9 +39,13 @@ impl super::Context {
             naga::valid::Capabilities::DUAL_SOURCE_BLENDING,
             device_caps.dual_source_blending,
         );
+        caps.set(
+            naga::valid::Capabilities::SHADER_FLOAT16,
+            device_caps.shader_float16,
+        );
         caps.set(
             naga::valid::Capabilities::COOPERATIVE_MATRIX,
-            device_caps.cooperative_matrix,
+            device_caps.cooperative_matrix.is_supported(),
         );
         naga::valid::Validator::new(flags, caps)
             .validate(module)

blade-graphics/src/vulkan/init.rs

@@ -71,7 +71,8 @@ struct AdapterCapabilities {
     external_memory: bool,
     timing: bool,
     dual_source_blending: bool,
-    cooperative_matrix: bool,
+    shader_float16: bool,
+    cooperative_matrix: crate::CooperativeMatrix,
     memory_budget: bool,
     bugs: SystemBugs,
 }
@@ -256,6 +257,8 @@ unsafe fn inspect_adapter(
     let mut ray_query_features = vk::PhysicalDeviceRayQueryFeaturesKHR::default();
     let mut cooperative_matrix_features = vk::PhysicalDeviceCooperativeMatrixFeaturesKHR::default();
     let mut vulkan_memory_model_features = vk::PhysicalDeviceVulkanMemoryModelFeatures::default();
+    let mut float16_int8_features = vk::PhysicalDeviceShaderFloat16Int8Features::default();
+    let mut storage_16bit_features = vk::PhysicalDevice16BitStorageFeatures::default();
     let mut features2_khr = vk::PhysicalDeviceFeatures2::default()
         .push_next(&mut inline_uniform_block_features)
         .push_next(&mut timeline_semaphore_features)
@@ -265,12 +268,15 @@ unsafe fn inspect_adapter(
         .push_next(&mut acceleration_structure_features)
         .push_next(&mut ray_query_features)
         .push_next(&mut cooperative_matrix_features)
-        .push_next(&mut vulkan_memory_model_features);
+        .push_next(&mut vulkan_memory_model_features)
+        .push_next(&mut float16_int8_features)
+        .push_next(&mut storage_16bit_features);
     instance
         .get_physical_device_properties2
         .get_physical_device_features2(phd, &mut features2_khr);
 
     let dual_source_blending = features2_khr.features.dual_src_blend != 0;
+    let shader_float16 = float16_int8_features.shader_float16 != 0;
 
     let has_inline_ub = supported_extensions.contains(&vk::EXT_INLINE_UNIFORM_BLOCK_NAME)
         && inline_uniform_block_properties.max_inline_uniform_block_size
@@ -382,23 +388,67 @@ unsafe fn inspect_adapter(
 
     let cooperative_matrix = if !supported_extensions.contains(&vk::KHR_COOPERATIVE_MATRIX_NAME) {
         log::info!("No cooperative matrix extension support");
-        false
+        crate::CooperativeMatrix::default()
     } else if cooperative_matrix_features.cooperative_matrix == vk::FALSE {
         log::info!(
             "No cooperative matrix feature support. Features = {:?}",
             cooperative_matrix_features
         );
-        false
+        crate::CooperativeMatrix::default()
     } else if vulkan_memory_model_features.vulkan_memory_model == vk::FALSE {
         log::info!(
             "No Vulkan memory model support (required for cooperative matrix). Features = {:?}",
             vulkan_memory_model_features
         );
-        false
+        crate::CooperativeMatrix::default()
     } else {
-        log::info!("Cooperative matrix is supported");
-        true
+        // Query supported cooperative matrix configurations and find
+        // square float configurations (Naga supports 8x8 and 16x16).
+        let coop_props = instance
+            .cooperative_matrix
+            .get_physical_device_cooperative_matrix_properties(phd)
+            .unwrap_or_default();
+        let find_tile = |a_type, b_type, c_type, result_type| {
+            [8u32, 16].into_iter().find(|&size| {
+                coop_props.iter().any(|p| {
+                    p.m_size == size
+                        && p.n_size == size
+                        && p.k_size == size
+                        && p.a_type == a_type
+                        && p.b_type == b_type
+                        && p.c_type == c_type
+                        && p.result_type == result_type
+                        && p.scope == vk::ScopeKHR::SUBGROUP
+                })
+            })
+        };
+        let f32t = vk::ComponentTypeKHR::FLOAT32;
+        let f16t = vk::ComponentTypeKHR::FLOAT16;
+        let f32_tile = find_tile(f32t, f32t, f32t, f32t).unwrap_or(0);
+        let f16_tile = if float16_int8_features.shader_float16 != 0
+            && storage_16bit_features.storage_buffer16_bit_access != 0
+        {
+            find_tile(f16t, f16t, f32t, f32t).unwrap_or(0)
+        } else {
+            0
+        };
+        let cm = crate::CooperativeMatrix { f32_tile, f16_tile };
+        if cm.is_supported() {
+            log::info!(
+                "Cooperative matrix: f32 tile={}, f16 tile={}",
+                cm.f32_tile,
+                cm.f16_tile,
+            );
+        } else {
+            log::info!(
+                "Cooperative matrix extension present but no usable config. Properties: {:?}",
+                coop_props
+            );
+        }
+        cm
     };
+    // Auto-enable shader_float16 when cooperative matrix has f16 support.
+    let shader_float16 = shader_float16 || cooperative_matrix.f16_tile > 0;
 
     let buffer_marker = supported_extensions.contains(&vk::AMD_BUFFER_MARKER_NAME);
     let shader_info = supported_extensions.contains(&vk::AMD_SHADER_INFO_NAME);
@@ -434,6 +484,7 @@ unsafe fn inspect_adapter(
         external_memory,
         timing,
         dual_source_blending,
+        shader_float16,
         cooperative_matrix,
         memory_budget,
         bugs,
@@ -601,7 +652,7 @@ impl super::Context {
                 }
             } else {
                 unsafe { entry.create_instance(&create_info, None) }
-                    .map_err(|e| crate::PlatformError::init(e))?
+                    .map_err(crate::PlatformError::init)?
             }
         };
 
@@ -610,6 +661,7 @@ impl super::Context {
                 _debug_utils: ext::debug_utils::Instance::new(&entry, &core_instance),
                 get_physical_device_properties2:
                     khr::get_physical_device_properties2::Instance::new(&entry, &core_instance),
+                cooperative_matrix: khr::cooperative_matrix::Instance::new(&entry, &core_instance),
                 get_surface_capabilities2: if desc.presentation {
                     Some(khr::get_surface_capabilities2::Instance::new(
                         &entry,
@@ -663,7 +715,7 @@ impl super::Context {
             instance
                 .core
                 .enumerate_physical_devices()
-                .map_err(|e| crate::PlatformError::init(e))?
+                .map_err(crate::PlatformError::init)?
                 .into_iter()
                 .find_map(|phd| {
                     inspect_adapter(
@@ -738,7 +790,7 @@ impl super::Context {
                     vk::KHR_EXTERNAL_MEMORY_FD_NAME
                 });
             }
-            if capabilities.cooperative_matrix {
+            if capabilities.cooperative_matrix.is_supported() {
                 device_extensions.push(vk::KHR_COOPERATIVE_MATRIX_NAME);
                 if capabilities.api_version < vk::API_VERSION_1_2 {
                     device_extensions.push(vk::KHR_VULKAN_MEMORY_MODEL_NAME);
@@ -810,9 +862,27 @@ impl super::Context {
                     .push_next(&mut khr_ray_query);
             }
 
+            let mut khr_float16_int8;
+            let mut storage_16bit;
+            if capabilities.shader_float16 {
+                khr_float16_int8 = vk::PhysicalDeviceShaderFloat16Int8Features {
+                    shader_float16: vk::TRUE,
+                    ..Default::default()
+                };
+                device_create_info = device_create_info.push_next(&mut khr_float16_int8);
+            }
+            if capabilities.cooperative_matrix.f16_tile > 0 {
+                storage_16bit = vk::PhysicalDevice16BitStorageFeatures {
+                    storage_buffer16_bit_access: vk::TRUE,
+                    uniform_and_storage_buffer16_bit_access: vk::TRUE,
+                    ..Default::default()
+                };
+                device_create_info = device_create_info.push_next(&mut storage_16bit);
+            }
+
             let mut khr_cooperative_matrix;
             let mut vulkan_memory_model;
-            if capabilities.cooperative_matrix {
+            if capabilities.cooperative_matrix.is_supported() {
                 khr_cooperative_matrix = vk::PhysicalDeviceCooperativeMatrixFeaturesKHR {
                     cooperative_matrix: vk::TRUE,
                     ..Default::default()
@@ -861,7 +931,7 @@ impl super::Context {
                 instance
                     .core
                     .create_device(physical_device, &device_create_info, None)
-                    .map_err(|e| crate::PlatformError::init(e))?
+                    .map_err(crate::PlatformError::init)?
             }
         };
 
@@ -1123,6 +1193,7 @@ impl super::Context {
                     .limits
                     .framebuffer_depth_sample_counts,
             dual_source_blending: capabilities.dual_source_blending,
+            shader_float16: capabilities.shader_float16,
             cooperative_matrix: capabilities.cooperative_matrix,
             binding_array: capabilities.binding_array,
             memory_budget: capabilities.memory_budget,
@@ -1153,6 +1224,7 @@ impl super::Context {
             },
             sample_count_mask: self.sample_count_flags.as_raw(),
             dual_source_blending: self.dual_source_blending,
+            shader_float16: self.shader_float16,
             cooperative_matrix: self.cooperative_matrix,
         }
     }

blade-graphics/src/vulkan/mod.rs

@@ -19,6 +19,7 @@ struct Instance {
     core: ash::Instance,
     _debug_utils: ash::ext::debug_utils::Instance,
     get_physical_device_properties2: khr::get_physical_device_properties2::Instance,
+    cooperative_matrix: khr::cooperative_matrix::Instance,
     get_surface_capabilities2: Option<khr::get_surface_capabilities2::Instance>,
     surface: Option<khr::surface::Instance>,
 }
@@ -265,7 +266,8 @@ pub struct Context {
     min_uniform_buffer_offset_alignment: u64,
     sample_count_flags: vk::SampleCountFlags,
     dual_source_blending: bool,
-    cooperative_matrix: bool,
+    shader_float16: bool,
+    cooperative_matrix: crate::CooperativeMatrix,
     binding_array: bool,
     memory_budget: bool,
     instance: Instance,

blade-graphics/src/vulkan/surface.rs

@@ -198,7 +198,7 @@ impl super::Context {
                 window.window_handle().unwrap().as_raw(),
                 None,
             )
-            .map_err(|e| crate::PlatformError::init(e))?
+            .map_err(crate::PlatformError::init)?
         };
 
         let khr_surface = self

blade-particle/src/system.rs

@@ -215,12 +215,12 @@ impl ParticleSystem {
             EmitterShape::Sphere { radius } => radius,
         };
 
-        let (colors, color_count) = match &self.effect.particle.color {
+        let (colors, color_count) = match self.effect.particle.color {
             ColorConfig::Solid(c) => {
-                let packed = pack_color(*c);
+                let packed = pack_color(c);
                 ([packed, packed, packed, packed], 1u32)
             }
-            ColorConfig::Palette(palette) => {
+            ColorConfig::Palette(ref palette) => {
                 let mut colors = [0u32; 4];
                 let count = palette.len().min(4);
                 for i in 0..count {
@@ -267,7 +267,7 @@ impl ParticleSystem {
             let mut pc = pass.with(&pipeline.reset_pipeline);
             pc.bind(0, &self.main_data());
             let group_size = pipeline.reset_pipeline.get_workgroup_size();
-            let group_count = (self.capacity as u32 + group_size[0] - 1) / group_size[0];
+            let group_count = (self.capacity as u32).div_ceil(group_size[0]);
             pc.dispatch([group_count, 1, 1]);
             self.needs_reset = false;
         }
@@ -298,7 +298,7 @@ impl ParticleSystem {
                 self.emit_accumulator -= emit_count as f32;
                 let params = self.make_emit_params(emit_count, self.origin);
                 let wg_size = pipeline.emit_pipeline.get_workgroup_size()[0];
-                let groups = (emit_count + wg_size - 1) / wg_size;
+                let groups = emit_count.div_ceil(wg_size);
                 let mut pass = encoder.compute("particle emit continuous");
                 let mut pc = pass.with(&pipeline.emit_pipeline);
                 pc.bind(0, &main_data);
@@ -317,7 +317,7 @@ impl ParticleSystem {
         for burst in bursts {
             let params = self.make_emit_params(burst.count, burst.position);
             let wg_size = pipeline.emit_pipeline.get_workgroup_size()[0];
-            let groups = (burst.count + wg_size - 1) / wg_size;
+            let groups = burst.count.div_ceil(wg_size);
             let mut pass = encoder.compute("particle emit burst");
             let mut pc = pass.with(&pipeline.emit_pipeline);
             pc.bind(0, &main_data);