diff --git a/src/back/spv/instructions.rs b/src/back/spv/instructions.rs index fe03a4d61c..064bfa0014 100644 --- a/src/back/spv/instructions.rs +++ b/src/back/spv/instructions.rs @@ -93,10 +93,14 @@ pub(super) fn instruction_entry_point( pub(super) fn instruction_execution_mode( entry_point_id: Word, execution_mode: spirv::ExecutionMode, + args: &[Word], ) -> Instruction { let mut instruction = Instruction::new(Op::ExecutionMode); instruction.add_operand(entry_point_id); instruction.add_operand(execution_mode as u32); + for arg in args { + instruction.add_operand(*arg); + } instruction } diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index f4d71b82fb..8a5bbbd2ba 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -421,11 +421,20 @@ impl Writer { crate::ShaderStage::Vertex => {} crate::ShaderStage::Fragment => { let execution_mode = spirv::ExecutionMode::OriginUpperLeft; - self.try_add_capabilities(execution_mode.required_capabilities()); - super::instructions::instruction_execution_mode(function_id, execution_mode) + //self.try_add_capabilities(execution_mode.required_capabilities()); + super::instructions::instruction_execution_mode(function_id, execution_mode, &[]) .to_words(&mut self.logical_layout.execution_modes); } - crate::ShaderStage::Compute => {} + crate::ShaderStage::Compute => { + let execution_mode = spirv::ExecutionMode::LocalSize; + //self.try_add_capabilities(execution_mode.required_capabilities()); + super::instructions::instruction_execution_mode( + function_id, + execution_mode, + &entry_point.workgroup_size, + ) + .to_words(&mut self.logical_layout.execution_modes); + } } if self.writer_flags.contains(WriterFlags::DEBUG) {