Add MSL support for mixed-precision cooperative matrix multiply-add#9629
Open
seddonm1 wants to merge 4 commits into
Open
Add MSL support for mixed-precision cooperative matrix multiply-add#9629seddonm1 wants to merge 4 commits into
seddonm1 wants to merge 4 commits into
Conversation
simdgroup_multiply_accumulate
inner-daemons
approved these changes
Jun 5, 2026
inner-daemons
left a comment
Collaborator
There was a problem hiding this comment.
Would nominate this for an award if I could. Very simple, self explanatory, with testing to back it up, and a good changelog entry. LGTM!
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
Adds MSL support for
simdgroup_multiply_accumulatewrappers where cooperative matrix A/B operands use a different scalar type from the C accumulator/result.coopMultiplyAdd(a, b, c)resolves to the type ofc, but the MSL wrapper previously used one scalar type for A, B, C, the temporary accumulator, and the return value. The wrapper now tracks the A/B scalar separately from the C/result scalar and emits a signature that matches mixed accumulator use cases.Testing
Added a new
cooperative-matrix-mixed-accumulatorsnapshot instead of extending the existingcooperative-matrixsnapshot so the mixed-scalar accumulator case is isolated.The new input uses f16 A/B cooperative matrices and an f32 C accumulator, with targets for
SPIRV | METAL | WGSL. This produces snapshots for all supported cooperative matrix backends:naga/tests/out/msl/wgsl-cooperative-matrix-mixed-accumulator.metalnaga/tests/out/spv/wgsl-cooperative-matrix-mixed-accumulator.spvasmnaga/tests/out/wgsl/wgsl-cooperative-matrix-mixed-accumulator.wgslChecklist
wgpumay be affected behaviorally.CHANGELOG.mdentries for the user-facing effects of this change are present.