[](https://matrix.to/#/#webgpu-dawn:matrix.org)
[](https://matrix.to/#/#webgpu-dawn:matrix.org)
[VkPhysicalDeviceCooperativeMatrixPropertiesKHR](https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VkPhysicalDeviceCooperativeMatrixPropertiesKHR), but I’ve only ever seen compute supported.
OpTypeCooperativeMatrixKHR represents the abstract type. It parameterized as follows:
* Component type: Matrix element type (must be a numerical scalar) (Float16/32/64, S/UInt8/16/32/64 currently)
* Scope: the scope of operations on the type (Device, QueueFamily, Workgroup, or Subgroup)
* Rows: Number of rows
* Columns: Number of columns
* Use: an enum specifying which kind of matrix this is used as:
* MatrixAKHR: LHS of a multiply (M rows x K columns)
* MatrixBKHR: RHS of a multiply (K rows x N columns)
* MatrixAccumulatorKHR: result of multiply-accumulate (M rows x N columns)
SPIR-V can represent a wide variety of cooperative matrices, but actual devices only support a subset of parameterizations. These are queried from the API by calling [vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR](https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR). It returns a linked list of [valid enumerations](https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VkCooperativeMatrixPropertiesKHR). Preferred variants are listed earlier in the chain. At least one entry in the list must have power of two values for all of MSize, KSize, and NSize. The enumerations indicate whether they support saturating accumulation. Because the type is abstract it can only be stored in the Function and Private storage classes. Special load and store instructions are used to translate to/from backing memory.
The fundamental operation is OpCooperativeMatrixMulAddKHR. It performs a multiply-accumulate operation. It requires consistent values of scope, M, N, and K for all types. The component types may differ, the operations occur using the result type. Signedness of results and operands is indicated using the Cooperative Matrix Operands (as is saturating accumulation). The result is undefined if saturating accumulation is used, but intermediate results overflow (same as integer dot product).
OpCooperativeMatrixLoadKHR and OpCooperativeMatrixStoreKHR share the following operands:
* Pointer: a logical pointer into an array (stride is ignored) (limited to Workgroup, StorageBuffer, and PhysicalStorageBuffer storage classes in Vulkan)
* Memory Layout: an enum for row- or column-majorness
* Stride: an optional operand indicating the stride between elements (aligned to min(16, align(col/row))). The stride counts the number of elements in the pointee type; it is not a byte count. For example, if the pointee type of the Pointer argument is 2xf16 then a stride of 8 translates to a byte stride of 8x (2x2) bytes.
These operations convert to/from the abstract matrix type and memory. The operands **must** be uniform values for the scope.
Additional operations:
* OpCooperativeMatrixLengthKHR: Number of accessible components from an invocation for a given cooperative matrix type.
* Conversions: All standard conversions are allowed (FToU, UToF, FToS, SToF, U, S, F, bitcast).
* Arithmetic: Negate, add, sub, mul, div using the opcode for the appropriate component type. Additionally, OpMatrixTimesScalar is also supported.
* Constants: OpCompositeConstruct and OpConstantComposite can provide a single fill value for a cooperative matrix.
All operations involving a cooperative matrix **must** only be executed in uniform control flow for scope of the operation.
Cooperative matrices **require** the Vulkan memory model in SPIR-V.
OpTypeCooperativeMatrixKHR can only be instantiated by a variable in the Function or Private storage classes.
### HLSL/Direct3D
Microsoft calls them WaveMatrices, and added them to [SM6.8](https://microsoft.github.io/DirectX-Specs/d3d/HLSL_SM_6_8_WaveMatrix.html) as experimental.
Google provided feedback via https://github.com/microsoft/hlsl-specs/issues/72
This experimental feature was later withdrawn.
Jul 14, 2025 [Chris Bieneman](mailto:cbieneman@microsoft.com) posted a revised proposal for HLSL linalg. See https://github.com/microsoft/hlsl-specs/pull/556
### MSL/Metal
Apple calls them simdgroup matrices and support has existed since MSL 2.3. They are limited to relatively few GPUs though (Apple7+ in the [feature set](https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf)). That is iPhone 12+, some newer iPads, and the M1 and later macs.
MSL support a small number of matrix variants:
* simdgroup_float8x8
* simdgroup_half8x8
* simdgroup_bfloat8x8 (MSL 3.1)
The simdgroup_float8x8 and simdgroup_half8x8 types are equivalent to a SPIR-V 8x8 cooperative matrix with Subgroup scope and corresponding component type. Unlike SPIR-V and HLSL, MSL does not specialize the matrices based on the use.
MSL also has more limited supported operations on simdgroup matrices. For load/store/creation:
* simdgroup_matrix<_T_, _M_, _N_>: Create diagonal matrix with single value.
* make_filled_simdgroup_matrix<_T_, _M_, _N_>: Create single value filled matrix.
* simdgroup_load/store: specify stride, layout (via transpose), and an offset (via origin)
Similar to SPIR-V the device or threadgroup pointer points to a matrix component type.
MSL supports multiply and multiply-accumulate (non-saturating) arithmetic operations.
All functions **must** be called from simdgroup uniform control flow. **Presumably** load and store parameter values **must** be simdgroup uniform though I can’t see this called out in the spec.
MSL allows interfaces to be declared as simdgroup_matrix, but no operators exist for them to make them usable. Effectively, this limits them to function storage.
### SYCL
SYCL is a single-source CPU/GPU compute language, and may reflect support among GPU hardware, at least as exposed through OpenCL.
SYCL has a corresponding proposal, using the terminology “joint matrix”. See https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix/sycl_ext_oneapi_matrix.asciidoc
## Proposal
### WGSL
#### Enable Extension
Add a new enable extension, `chromium_experimental_subgroup_matrix`. Enables the declaration of subgroup_matrix types, related built-in functions, and limited use of some scalar types (u8 and i8). Implicitly depends on subgroups. WGSL does not require any subgroups enable, but the API does.
#### Types
Add new types:
* subgroup_matrix_left<_T_, _K_, _M_>: M rows x K columns matrix of T components.
* subgroup_matrix_right<_T_, _N_, _K_>: K rows x N columns matrix of T components
* subgroup_matrix_result<_T_, _N_, _M_>: M rows x N columns matrix of T components
* M, N, and K are override-expressions of positive integer values
* For _initial prototyping and feedback_, we will restrict them to const-expression. We should finalize const vs. override before full standardization.
* Similar to arrays, two subgroup matrices are the same if:
* They are the same base type
* They have the same component type
* The have matching row counts such that:
* Both are equal valued const-expressions, or
* Both resolve to the same override-declaration
* The have matching column counts such that:
* Both are equal valued const-expressions, or
* Both resolve to the same override-declaration
* Note: like arrays, types cannot match if the row or column counts use an override-expression that is not equivalent to override identifier.
* T is the component type and must be one of the entries in the table
* The scalar shader type is the associated type usable in the shader code for scalar operations and data representation in memory
* Can be expanded in the future to support more types (e.g. bfloat16) via new enables.
* The u8 and i8 cases are predeclared types that are not otherwise usable in WGSL. For layout calculations, they are of size 1 byte and have an alignment requirement of 1 byte.
| Type | Extra Enable | Element Stride (bytes) | Shader Scalar Type | Min Value | Max Value |
| ---- | ------------ | ---------------------- | ------------------ | --------- | --------- |
| f32 | | 4 | f32 | | |
| f16 | f16 | 2 | f16 | | |
| u32 | | 4 | u32 | | |
| i32 | | 4 | i32 | | |
| u8 | | 1 | u32 | 0 | 255 |
| i8 | | 1 | i32 | -128 | 127 |
These types are not considered “composite” in the WGSL taxonomy, because they are not decomposible. You can’t reference a sub-vector or a single component. The numeric dimensions must be override-expressions. These types cannot be part of any interface (i.e. they can only be instantiated in Function and Private address spaces). They are plain types (similar to atomics) so that they can be included in composite types. An important use case is to make arrays of these matrices.
It is a pipeline-creation error if any matrix type is not included in a supported `GPUSubgroupMatrixConfig`, `config`, on the device:
* For `subgroup_matrix_left`:
* `M` equals `config.M`
* `K` equals `config.K`
* `T` matches `config.componentType`
* For `subgroup_matrix_right`:
* `K` equals `config.K`
* `N` equals `config.N`
* `T` matches `config.componentType`
* For `subgroup_matrix_result`:
* `M` equals `config.M`
* `N` equals `config.N`
* `T` matches `config.resultType`
Why use a “left” and “right” matrix type, instead of a single matrix type like Metal has?
* D3D and Vulkan both need them.
* There is no free transpose operation. If you had a free transpose, then you can take advantage of: A\* B = transpose(transpose(B)\*transpose(A)). I haven’t seen a transpose operation in APIs or hardware specs.
#### Variables
A variable containing a subgroup matrix can only be instantiated in Function or Private address spaces. (This limitation comes from SPIR-V and Metal). These variables are meant to be used as very temporary scratch space.
#### Loading and storing
Builtin functions are used to load and store subgroup matrix values from variables in workgroup and storage address spaces. The builtins do two things:
* Map matrix row and column indices to external memory locations. This mapping has two parameters: majorness (row-major, or column-major), and an integer Stride. Let Base be the byte address of the start of the external matrix, i.e. where the [0,0]’th element of the matrix is stored. Then
* For row-major:
* Matrix entry [r,c] maps to the sizeof(T) bytes located at Base + Stride\*r\*sizeof(T) + sizeof(T)\*c
* Stride >= number of matrix columns.
* For column-major:
* Matrix entry [r,c] maps to the sizeof(T) bytes located at Base + Stride\*c\*sizeof(T) + sizeof(T)\*r
* Stride >= number of matrix rows.
* Reinterpret data values between the shader scalar type and the external component type T, when those types differ.
For a subgroup_matrix_left/right/result<T, Cols, Rows>, loads and stores are out-of-bounds if the length of the array of the pointer argument is less than Offset + Stride \* Rows\* Cols.
#### Attributes
Add an additional error condition to `workgroup_size`. It is a pipeline-creation error if the x dimension is not a multiple of `GPUAdapterInfo.subgroupMaxSize` and the shader statically accesses any subgroup matrix.
Note: this is where the requirement on the subgroups feature stems from in practice.
#### Expressions
Supported expressions:
* Parenthesization
* Identifier expressions
* Indirection
* Address-of
* Function call expression
* Type expressions
Explicitly not supported:
* Decomposition expressions
Possible future expansion:
* Arithmetic expressions
#### Built-in Values
New built-in value: `subgroup_id`.
Input in compute shaders of type `u32`.
It is the current invocation's subgroup's ID within the workgroup.
_Note: This was not included in the subgroups feature because there was insufficient documentation detailing the mapping in HLSL between workgroups and subgroups. This experiment is not implemented on D3D so the built-in value is available on supported platforms. If guaranteed mapping from Microsoft is provided for D3D, subgroup_id could be added as a language feature on top of the subgroups feature._
#### Built-in Functions
Calls to these functions:
* Must only be used in a compute shader stage.
* Trigger a `chromium_experimental.subgroup_matrix_uniformity` diagnostic if uniformity analysis cannot prove the call is in subgroup uniform control flow.
##### Value constructors
**Overloads**:
```rust
@must_use fn T(value : S) -> T
@must_use fn T() -> T
```
**Preconditions**: