Skip to content
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.

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

Cooperative matrix #4195

Open
alan-baker opened this issue Jun 23, 2023 · 2 comments
Open

Cooperative matrix #4195

alan-baker opened this issue Jun 23, 2023 · 2 comments
Labels
feature request A request for a new GPU feature exposed in the API wgsl WebGPU Shading Language Issues
Projects
Milestone

Comments

@alan-baker
Copy link
Contributor

All major platform APIs have now released a similar extensions for cooperative matrix:

  • Metal introduced simdgroup_matrix in MSL 3.1
  • HLSL has support in SM6.8 (currently experimental release)
  • SPIR-V/Vulkan released SPV_KHR_cooperative_matrix

This feature is very useful for accelerating ML operations by performing matrix calculations across multiple threads.

@alan-baker alan-baker added feature request A request for a new GPU feature exposed in the API wgsl WebGPU Shading Language Issues labels Jun 23, 2023
@alan-baker alan-baker added this to the post-V1 milestone Jun 23, 2023
@PWhiddy
Copy link

PWhiddy commented Jan 5, 2024

Would a wgpu extension providing an interface for this make it feasible?

@dneto0
Copy link
Contributor

dneto0 commented Mar 1, 2024

Update on Google's thinking:

  • This is important, with clear use cases
  • Platform support is uneven so far. HLSL's SM 6.8 is still experimental.
  • We think the portability issues for subgroups also affect coopartive matrix multiplies. So we'd like to resolve subgroup portability first.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
feature request A request for a new GPU feature exposed in the API wgsl WebGPU Shading Language Issues
Projects
WGSL
Awaiting triage
Development

No branches or pull requests

3 participants