Enable builtins.

[dvc94ch]

No description provided.

[eddyb]

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

[dvc94ch]

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

[dvc94ch]

Is there an example of an existing builtin test somewhere?

[expenses]

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)

[dvc94ch]

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.

[hrydgard]

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

[dvc94ch]

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

[eddyb]

Looking at the Vulkan spec (since SPIR-V fails to document it), looks like SubgroupLocalInvocationId does indeed go on a variable.

(Whereas for WorkgroupSize it mentions "constants" - though technically ambiguous/a misuse of "variable")