-
Notifications
You must be signed in to change notification settings - Fork 329
Closed
Labels
wgslWebGPU Shading Language IssuesWebGPU Shading Language Issues
Milestone
Description
We currently have a workgroup_size attribute for compute shader that takes a 3D grid size expressed as 3 integer literals.
It's very common to tune compute shaders based on workgroup sizes, after initial shader compilation:
- in OpenCL you set it at enqueue time: see the local_work_size parameter to clEnqueueNDRangeKernel
- in Metal it's the threadsPerThreadgroup parameter to dispatchThreads
- in Vulkan you can make workgroup sizes from specialization constants, and then override them at pipeline creation time. See use of the WorkgroupSize builtin.
In the OpenCL and Vulkan cases, I know that the late-binding can fail due to workgroup size problems (as it can fail for other reasons too). OpenCL even has an API for asking for an acceptable workgroup size. (See https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/clGetKernelWorkGroupInfo.html )
We have partner applications that strongly rely on the ability to do this late-ish binding of workgroup sizes.
cc: @alan-baker
kvarkstolk
Metadata
Metadata
Assignees
Labels
wgslWebGPU Shading Language IssuesWebGPU Shading Language Issues