Can you try adding tests for these? I thought they couldn't work for now.

They seem to work fine. Any reason why you think they shouldn't?

Is there an example of an existing builtin test somewhere?

See: #789 (I followed the git blame)

In discord, it was pointed out that #[spirv(workgroup_size)] is broken - the decoration applies to a constant, not a variable, and it specifies the workgroup size, it's not provided with the workgroup size. (A subsequent discussion also showed that WorkgroupSize in SPIR-V is just plain broken with multi-entry-point shaders - but the LocalSize(Id) execution mode is a replacement for its behavior)

ok workgroup_size might be broken, realized I wasn't actually using it. but I think we went over this in #841, subgroup_local_invocation_id should be supported as it has the following enabling capabilities: Kernel, GroupNonUniform, SubgroupBallotKHR.

I think this should be fine to go in if you just back out the workgroup_size change, then.

Let me know if you want me to drop the example.