From edf5d857b1ee4e3431e141beb1cae08facc78a33 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Thu, 10 Dec 2020 00:13:07 -0500 Subject: [PATCH] Add workgroup id and size builtins --- src/back/glsl/mod.rs | 1 + src/back/msl/mod.rs | 1 + src/back/spv/writer.rs | 1 + src/front/glsl/variables.rs | 35 +++++++++++++++++++++++++++++++++++ src/front/spv/convert.rs | 1 + src/front/wgsl/conv.rs | 2 ++ src/lib.rs | 1 + src/proc/validator.rs | 18 ++++++++++-------- 8 files changed, 52 insertions(+), 8 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 3f87d1c760..419452e1cd 100644 --- a/src/back/glsl/mod.rs +++ b/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", } } diff --git a/src/back/msl/mod.rs b/src/back/msl/mod.rs index e96991ba33..cc08869e74 100644 --- a/src/back/msl/mod.rs +++ b/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)?) } diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 9279cce037..db06ba320a 100644 --- a/src/back/spv/writer.rs +++ b/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 diff --git a/src/front/glsl/variables.rs b/src/front/glsl/variables.rs index 9f76335685..50343ba922 100644 --- a/src/front/glsl/variables.rs +++ b/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); + } _ => {} } diff --git a/src/front/spv/convert.rs b/src/front/spv/convert.rs index 6036d08946..167863efa3 100644 --- a/src/front/spv/convert.rs +++ b/src/front/spv/convert.rs @@ -118,6 +118,7 @@ pub fn map_builtin(word: spirv::Word) -> Result { 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)), }) } diff --git a/src/front/wgsl/conv.rs b/src/front/wgsl/conv.rs index 18de009aad..8beada1881 100644 --- a/src/front/wgsl/conv.rs +++ b/src/front/wgsl/conv.rs @@ -25,6 +25,8 @@ pub fn map_built_in(word: &str) -> Result> { "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)), }) } diff --git a/src/lib.rs b/src/lib.rs index 61b835c4b1..dff0233ace 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -153,6 +153,7 @@ pub enum BuiltIn { LocalInvocationId, LocalInvocationIndex, WorkGroupId, + WorkGroupSize, } /// Number of bytes. diff --git a/src/proc/validator.rs b/src/proc/validator.rs index 20061c1c49..f39a5bf85f 100644 --- a/src/proc/validator.rs +++ b/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,