Repository navigation
Proposal for subgroup matrix feature - #5335
Conversation
|
Previews, as seen when this build job started (01420ea): |
dneto0
left a comment
There was a problem hiding this comment.
This captures the proposal as intended, and is a good basis for further iteration and discusison.
Key changes from what's in the doc in the Dawn source tree:
- change name of the extension
- pull
subgroup_idout to a separate feature - describe relation to the HLSL
linalgproposal.
I think this is ready to land.
| #### 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. |
There was a problem hiding this comment.
On Apple Silicon, the subgroup-size restriction is likely to be a bottleneck. If we need a bound, could we use subgroupMinSize rather than subgroupMaxSize?
There was a problem hiding this comment.
The better restriction would be that x dimension is a multiple of the selected subgroup size. The issue is that there is no good way to get that information. Metal can query the subgroup size of the pipeline, but Vulkan and D3D cannot. Vulkan and D3D can specify a specific subgroup size, but Metal cannot. So consider a device that reports the subgroup size range [8, 32]. If the compiler selects a subgroup size of 32, a workgroup size of 16 (multiple of min size) would be invalid. Are these devices common? I'm not certain, but it's not clear how to handle it.
We can look into devices more thoroughly compared to supported subgroup sizes, but I'm not certain how restrictive we want to be.
There was a problem hiding this comment.
Metal caps the subgroup size at 32 today, which makes this specification acceptable short-term.
google/dawn@d7d27a6
However, performance usually peaks when subgroup size aligns with the hardware SIMD width. If possible, consider returning the optimal subgroup size (in addition to—or instead of—min/max). For context, NVIDIA warps are 32, while AMD wavefronts are typically 64 (RDNA: 32 or 64), so a single max/min pair can be misleading across vendors.
|
|
||
| ### 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. |
There was a problem hiding this comment.
'relatively few GPUs' is potentially inaccurate, since Apple7 is A14 Bionic which released 5 years ago.
But I agree it is certainly not supported on all Metal devices. I.e., no Intel Mac support, many iPads do not support this, and some iPhones too.
There was a problem hiding this comment.
I'm capturing this feedback in an update to the proposal text in #5346
There's no need to quantify it as many or few, just that it's not all of them.
Minutes from today's WGSL committee meeting Oct 7, 2025
|
|
My unofficial summary: looks good as proposal, let's have further discussion as PRs/issues regarding that document. |
Refs #4195