Speaker
Description
Workgroups are a concept which allow multiple SIMD waves of shaders to run together and share some memory. API features such as compute and mesh shaders expect the driver to support workgroups. Modern hardware typically supports it, but... What if you want to run workgroups on hardware that doesn't really support it, for example you want to run mesh shaders on NVidia? Or maybe you want to support larger workgroups on older hardware that has an upper limit on the workgroup size?
The new compiler pass nir_lower_workgroup_size was invented exactly for this purpose. Your driver can now support workgroups where there aren't any, or support larger workgroups when the HW can't.
In this lightning talk I'll present what the challenge is with "emulating" workgroups; how we solved it, and how we've used it to implement NVK mesh shaders, to support larger compute workgroups on old GPUs on RADV, and even to find ray tracing issues.
| Code of Conduct | Yes |
|---|---|
| GSoC, EVoC or Outreachy | No |
| In-person or virtual presentation | In-person |
