Skip to content

[subgroup-matrix] Update proposal for D3D12 #10964

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
alan-baker wants to merge 2 commits into
base: main
Choose a base branch
from
Open

[subgroup-matrix] Update proposal for D3D12 #10964

Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c4a0f11
[subgroup-matrix] Update proposal for D3D12
alan-baker Sep 22, 2026
319cae5
Update proposals/subgroup-matrix.md
alan-baker Sep 22, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Next Next commit
[subgroup-matrix] Update proposal for D3D12 
* Add D3D12 API requirements and mappings
* Add `minSubgroupSize` and `maxSubgroupSize` to
  `GPUSubgroupMatrixConfig`
* Add validation that a specified subgroup size is in range of a config
  • Loading branch information
@alan-baker
alan-baker committed Sep 22, 2026
commit c4a0f116a0215cbe70c7721cb07bb7c146b2e40b
39 changes: 34 additions & 5 deletions proposals/subgroup-matrix.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -97,9 +97,8 @@ as experimental.
This experimental feature was later withdrawn.

Microsoft is now targeting a new feature, linalg::Matrix, for SM6.10.
See https://github.com/microsoft/hlsl-specs/pull/556.
The API side will likely look similar to the WaveMatrix feature, but it is not
included in the proposal yet.
For HLSL, see https://github.com/microsoft/hlsl-specs/blob/main/proposals/0035-linalg-matrix.md.
For the D3D API, see https://microsoft.github.io/DirectX-Specs/d3d/D3D12LinearAlgebraRuntimeFeatureSupport.html.

The HLSL feature relies on SFINAE to provide a templated type for matrices.
The type is templated with the following parameters:
Expand All @@ -116,6 +115,8 @@ Additional functionality includes:
* Conversions: A <-> B, A/B <-> Accumulator
* Coordinates: Can access matrix coordinates when iterating over each thread's values.

The API configurations are additionally limited by WaveSize.

### MSL/Metal

Apple calls them simdgroup matrices and support has existed since MSL 2.3.
Expand Down Expand Up @@ -655,7 +656,7 @@ New GPUFeatureName `subgroup-matrix`
* Vulkan pipelines will need to be compiled with <code>VK_PIPELINE_SHADER_STAGE_CREATE_ALLOW_VARYING_SUBGROUP_SIZE_BIT</code> and <code>VK_PIPELINE_SHADER_STAGE_CREATE_REQUIRE_FULL_SUBGROUPS_BIT</code>, or the SPIR-V module must be version 1.6 or later
* Metal:
* Family is Apple 7+
* D3D: **TODO**
* D3D: `D3D12_LINEAR_ALGEBRA_TIER_1_0` is supported.

New immutable array, <code>subgroupMatrixConfigs</code>, added to <code>GPUAdapterInfo</code>.

Expand All @@ -680,6 +681,8 @@ interface GPUSubgroupMatrixConfig {
readonly attribute unsigned long M;
readonly attribute unsigned long N;
readonly attribute unsigned long K;
readonly attribute unsigned long minSubgroupSize;
readonly attribute unsigned long maxSubgroupSize;
};
```

Expand All @@ -703,6 +706,8 @@ WGSL pipeline-creation checks (repeated for ease of reference):
`GPUSubgroupMatrixConfig`
* The x-dimension of `workgroup_size` is a multiple of
`GPUSupportedLimits::maxSubgroupSize`
* If the shader specifes a `subgroup_size`, it must be in the range
Comment thread
alan-baker marked this conversation as resolved.
Outdated
[minSubgroupSize, maxSubgroupSize] of the `GPUSubgroupMatrixConfig`


### Mapping
Expand Down Expand Up @@ -759,9 +764,30 @@ Filter the list returned from [vkGetPhysicalDeviceCooperativeMatrixPropertiesKHR

`VK_COMPONENT_TYPE_FLOAT16` will need to be filtered out of the device properties if the `shader-f16` feature is not requested.

`minSubgroupSize` and `maxSubgroupSize` for each config can be set to the
adapter's `subgroupMinSize` and `subgroupMaxSize`.

##### D3D12

**TODO**: The feature is still under development.
Filter the list returned from enumerating `D3D12_LINEAR_ALGEBRA_OPERATION_TYPE_WAVE_MATRIX_MULTIPLY` operations:

* MatrixAComponentType matches GPUSubgroupMatrixConfig.componentType
* MatrixBComponentType matches GPUSubgroupMatrixConfig.componentType
* AccumulatorComponentType matches GPUSubgroupMatrixConfig.resultComponentType
* Shape.{M, N, K} matches GPUSubgroupMatrixConfig.{M, N, K}
* `subgroup_size`, if specified, is in range
* Component types match as follows (-> API enum):
* `D3D12_LINEAR_ALGEBRA_DATATYPE_SINT32` -> i32
* `D3D12_LINEAR_ALGEBRA_DATATYPE_UINT32` -> u32
* `D3D12_LINEAR_ALGEBRA_DATATYPE_FLOAT16` -> f16
* `D3D12_LINEAR_ALGEBRA_DATATYPE_FLOAT32` -> f32
* `D3D12_LINEAR_ALGEBRA_DATATYPE_SINT8` -> i8
* `D3D12_LINEAR_ALGEBRA_DATATYPE_UINT8` -> u8

`D3D12_LINEAR_ALGEBRA_DATATYPE_FLOAT16` will need filtered out of the device properties if the `shader-f16` feature is not requested.

**TODO**: confirm D3D does not require a `WaveSize` to be specified (though an
implementation could choose a valid size).

##### Metal

Expand All @@ -774,6 +800,9 @@ Hardcode the following configurations if the feature is supported:

1. Filter out f16 from the device properties if `shader-f16` is not requested on the device.

`minSubgroupSize` and `maxSubgroupSize` for each config can be set to the
adapter's `subgroupMinSize` and `subgroupMaxSize`.

**TODO**: Should we consider using performance primitives if the device supports Metal 4?

## Future Expansion
Expand Down
Loading