MSL: add initial cooperative matrix support#2596
MSL: add initial cooperative matrix support#2596HansKristian-Work merged 8 commits intoKhronosGroup:mainfrom
Conversation
|
|
|
Overloaded with other stuff atm, I'll try to get to this next week. |
spirv_msl.cpp
Outdated
| { | ||
| auto is_cooperative_matrix_typed_id = [&](uint32_t typed_id) -> bool { | ||
| auto *type = maybe_get<SPIRType>(typed_id); | ||
| if (!type) |
There was a problem hiding this comment.
A little overzealous, but sure I guess.
| switch (comp.basetype) | ||
| { | ||
| case SPIRType::Float: | ||
| return "simdgroup_float8x8"; |
There was a problem hiding this comment.
I think this would break if you have array of coopmats, but given that MSL coopmat clearly cannot be conformant to Vulkan due to lack of a lot of core features, I guess it's okay to just keep it simple until there is a real need.
|
Test suite is not passing, this needs to be fixed. |
| bfloat data[1]; | ||
| }; | ||
|
|
||
| kernel void main0(device SSBO& ssbo [[buffer(0)]]) |
There was a problem hiding this comment.
Missing a test for shared memory load-store. MSL spec says it's supported.
Also, this is missing tests for load-store with different types. E.g. in SPIR-V you can load a float16_t coopmat via a uint8_t array of data. Just pointer casting and recomputing the row/col stride should work.
This PR adds initial cooperative matrix (
SPV_KHR_cooperative_matrix) support for MSL.Related issue: #2478
Limitations:
For 2, although it is possible to break a larger matrix into several 8x8 matrices or padding a smaller one to 8x8, it requires a lot of effort.
For 3,
Matrix{A,B,C,Result}SignedComponentsKHRMaskis for integer types, whereas simdgroup matrix only has floating-point types. ForSaturatingAccumulationKHRMask, although MSL spec doesn't document, experiment shows that simdgroup_multiply_accumulate doesn't saturate or clamp and cannot be controlled.