Add workgroup id and size builtins

src/back/glsl/mod.rs

@@ -1643,6 +1643,7 @@ fn glsl_built_in(built_in: BuiltIn) -> &'static str {
         BuiltIn::LocalInvocationId => "gl_LocalInvocationID",
         BuiltIn::LocalInvocationIndex => "gl_LocalInvocationIndex",
         BuiltIn::WorkGroupId => "gl_WorkGroupID",
+        BuiltIn::WorkGroupSize => "gl_WorkGroupSize",
     }
 }
 

src/back/msl/mod.rs

@@ -173,6 +173,7 @@ impl ResolvedBinding {
                     Bi::LocalInvocationId => "thread_position_in_threadgroup",
                     Bi::LocalInvocationIndex => "thread_index_in_threadgroup",
                     Bi::WorkGroupId => "threadgroup_position_in_grid",
+                    Bi::WorkGroupSize => "dispatch_threads_per_threadgroup",
                 };
                 Ok(write!(out, "{}", name)?)
             }

src/back/spv/writer.rs

@@ -791,6 +791,7 @@ impl Writer {
                     crate::BuiltIn::LocalInvocationId => spirv::BuiltIn::LocalInvocationId,
                     crate::BuiltIn::LocalInvocationIndex => spirv::BuiltIn::LocalInvocationIndex,
                     crate::BuiltIn::WorkGroupId => spirv::BuiltIn::WorkgroupId,
+                    crate::BuiltIn::WorkGroupSize => spirv::BuiltIn::WorkgroupSize,
                 };
 
                 self.annotations

src/front/glsl/variables.rs

@@ -86,6 +86,41 @@ impl Program {
 
                 expression = Some(exp);
             }
+            "gl_InstanceIndex" => {
+                #[cfg(feature = "glsl-validate")]
+                match self.shader_stage {
+                    ShaderStage::Vertex => {}
+                    _ => {
+                        return Err(ErrorKind::VariableNotAvailable(name.into()));
+                    }
+                };
+                let h = self
+                    .module
+                    .global_variables
+                    .fetch_or_append(GlobalVariable {
+                        name: Some(name.into()),
+                        class: StorageClass::Input,
+                        binding: Some(Binding::BuiltIn(BuiltIn::InstanceIndex)),
+                        ty: self.module.types.fetch_or_append(Type {
+                            name: None,
+                            inner: TypeInner::Scalar {
+                                kind: ScalarKind::Uint,
+                                width: 4,
+                            },
+                        }),
+                        init: None,
+                        interpolation: None,
+                        storage_access: StorageAccess::empty(),
+                    });
+                self.lookup_global_variables.insert(name.into(), h);
+                let exp = self
+                    .context
+                    .expressions
+                    .append(Expression::GlobalVariable(h));
+                self.context.lookup_global_var_exps.insert(name.into(), exp);
+
+                expression = Some(exp);
+            }
             _ => {}
         }
 

src/front/spv/convert.rs

@@ -118,6 +118,7 @@ pub fn map_builtin(word: spirv::Word) -> Result<crate::BuiltIn, Error> {
         Some(Bi::LocalInvocationId) => crate::BuiltIn::LocalInvocationId,
         Some(Bi::LocalInvocationIndex) => crate::BuiltIn::LocalInvocationIndex,
         Some(Bi::WorkgroupId) => crate::BuiltIn::WorkGroupId,
+        Some(Bi::WorkgroupSize) => crate::BuiltIn::WorkGroupSize,
         _ => return Err(Error::UnsupportedBuiltIn(word)),
     })
 }

src/front/wgsl/conv.rs

@@ -25,6 +25,8 @@ pub fn map_built_in(word: &str) -> Result<crate::BuiltIn, Error<'_>> {
         "global_invocation_id" => crate::BuiltIn::GlobalInvocationId,
         "local_invocation_id" => crate::BuiltIn::LocalInvocationId,
         "local_invocation_index" => crate::BuiltIn::LocalInvocationIndex,
+        "workgroup_id" => crate::BuiltIn::WorkGroupId,
+        "workgroup_size" => crate::BuiltIn::WorkGroupSize,
         _ => return Err(Error::UnknownBuiltin(word)),
     })
 }

src/lib.rs

@@ -153,6 +153,7 @@ pub enum BuiltIn {
     LocalInvocationId,
     LocalInvocationIndex,
     WorkGroupId,
+    WorkGroupSize,
 }
 
 /// Number of bytes.

src/proc/validator.rs

@@ -154,13 +154,14 @@ impl crate::GlobalVariable {
                         kind: Sk::Bool,
                         width,
                     },
-                    Bi::GlobalInvocationId | Bi::LocalInvocationId | Bi::WorkGroupId => {
-                        Ti::Vector {
-                            size: Vs::Tri,
-                            kind: Sk::Uint,
-                            width,
-                        }
-                    }
+                    Bi::GlobalInvocationId
+                    | Bi::LocalInvocationId
+                    | Bi::WorkGroupId
+                    | Bi::WorkGroupSize => Ti::Vector {
+                        size: Vs::Tri,
+                        kind: Sk::Uint,
+                        width,
+                    },
                 };
                 if types[self.ty].inner != ty_inner {
                     log::warn!("Wrong builtin type: {:?}", types[self.ty]);
@@ -375,7 +376,8 @@ impl Validator {
                             | (crate::ShaderStage::Compute, crate::BuiltIn::GlobalInvocationId)
                             | (crate::ShaderStage::Compute, crate::BuiltIn::LocalInvocationId)
                             | (crate::ShaderStage::Compute, crate::BuiltIn::LocalInvocationIndex)
-                            | (crate::ShaderStage::Compute, crate::BuiltIn::WorkGroupId) => 0,
+                            | (crate::ShaderStage::Compute, crate::BuiltIn::WorkGroupId)
+                            | (crate::ShaderStage::Compute, crate::BuiltIn::WorkGroupSize) => 0,
                             _ => return Err(EntryPointError::InvalidBuiltIn(built_in)),
                         },
                         Some(crate::Binding::Location(loc)) => 1 << loc,