这是indexloc提供的服务,不要输入任何密码
Skip to content

allow workgroup size components to be pipeline-overrideable constants #1442

@dneto0

Description

@dneto0

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

Metadata

Metadata

Assignees

Labels

wgslWebGPU Shading Language Issues

Type

No type

Projects

No projects

Relationships

None yet

Development

No branches or pull requests

Issue actions