diff --git a/src/back/spv/instructions.rs b/src/back/spv/instructions.rs index 982d85bfe2..f4b50afb3c 100644 --- a/src/back/spv/instructions.rs +++ b/src/back/spv/instructions.rs @@ -1,4 +1,4 @@ -use crate::back::spv::{helpers, Instruction}; +use super::helpers; use spirv::{Op, Word}; pub(super) enum Signedness { @@ -6,745 +6,728 @@ pub(super) enum Signedness { Signed = 1, } -// -// Debug Instructions -// - -pub(super) fn instruction_source( - source_language: spirv::SourceLanguage, - version: u32, -) -> Instruction { - let mut instruction = Instruction::new(Op::Source); - instruction.add_operand(source_language as u32); - instruction.add_operands(helpers::bytes_to_words(&version.to_le_bytes())); - instruction -} - -pub(super) fn instruction_name(target_id: Word, name: &str) -> Instruction { - let mut instruction = Instruction::new(Op::Name); - instruction.add_operand(target_id); - instruction.add_operands(helpers::string_to_words(name)); - instruction -} - -pub(super) fn instruction_member_name(target_id: Word, member: Word, name: &str) -> Instruction { - let mut instruction = Instruction::new(Op::MemberName); - instruction.add_operand(target_id); - instruction.add_operand(member); - instruction.add_operands(helpers::string_to_words(name)); - instruction -} - -// -// Annotation Instructions -// - -pub(super) fn instruction_decorate( - target_id: Word, - decoration: spirv::Decoration, - operands: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::Decorate); - instruction.add_operand(target_id); - instruction.add_operand(decoration as u32); - for operand in operands { - instruction.add_operand(*operand) - } - instruction -} - -pub(super) fn instruction_member_decorate( - target_id: Word, - member_index: Word, - decoration: spirv::Decoration, - operands: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::MemberDecorate); - instruction.add_operand(target_id); - instruction.add_operand(member_index); - instruction.add_operand(decoration as u32); - for operand in operands { - instruction.add_operand(*operand) - } - instruction -} - -// -// Extension Instructions -// - -pub(super) fn instruction_ext_inst_import(id: Word, name: &str) -> Instruction { - let mut instruction = Instruction::new(Op::ExtInstImport); - instruction.set_result(id); - instruction.add_operands(helpers::string_to_words(name)); - instruction -} - -pub(super) fn instruction_ext_inst( - set_id: Word, - op: spirv::GLOp, - result_type_id: Word, - id: Word, - operands: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::ExtInst); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(set_id); - instruction.add_operand(op as u32); - for operand in operands { - instruction.add_operand(*operand) - } - instruction -} - -// -// Mode-Setting Instructions -// - -pub(super) fn instruction_memory_model( - addressing_model: spirv::AddressingModel, - memory_model: spirv::MemoryModel, -) -> Instruction { - let mut instruction = Instruction::new(Op::MemoryModel); - instruction.add_operand(addressing_model as u32); - instruction.add_operand(memory_model as u32); - instruction -} - -pub(super) fn instruction_entry_point( - execution_model: spirv::ExecutionModel, - entry_point_id: Word, - name: &str, - interface_ids: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::EntryPoint); - instruction.add_operand(execution_model as u32); - instruction.add_operand(entry_point_id); - instruction.add_operands(helpers::string_to_words(name)); - - for interface_id in interface_ids { - instruction.add_operand(*interface_id); - } - - instruction -} - -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 -} - -pub(super) fn instruction_capability(capability: spirv::Capability) -> Instruction { - let mut instruction = Instruction::new(Op::Capability); - instruction.add_operand(capability as u32); - instruction -} - -// -// Type-Declaration Instructions -// - -pub(super) fn instruction_type_void(id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::TypeVoid); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_type_bool(id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::TypeBool); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_type_int(id: Word, width: Word, signedness: Signedness) -> Instruction { - let mut instruction = Instruction::new(Op::TypeInt); - instruction.set_result(id); - instruction.add_operand(width); - instruction.add_operand(signedness as u32); - instruction -} - -pub(super) fn instruction_type_float(id: Word, width: Word) -> Instruction { - let mut instruction = Instruction::new(Op::TypeFloat); - instruction.set_result(id); - instruction.add_operand(width); - instruction -} - -pub(super) fn instruction_type_vector( - id: Word, - component_type_id: Word, - component_count: crate::VectorSize, -) -> Instruction { - let mut instruction = Instruction::new(Op::TypeVector); - instruction.set_result(id); - instruction.add_operand(component_type_id); - instruction.add_operand(component_count as u32); - instruction -} - -pub(super) fn instruction_type_matrix( - id: Word, - column_type_id: Word, - column_count: crate::VectorSize, -) -> Instruction { - let mut instruction = Instruction::new(Op::TypeMatrix); - instruction.set_result(id); - instruction.add_operand(column_type_id); - instruction.add_operand(column_count as u32); - instruction -} - -pub(super) fn instruction_type_image( - id: Word, - sampled_type_id: Word, - dim: spirv::Dim, - arrayed: bool, - image_class: crate::ImageClass, -) -> Instruction { - let mut instruction = Instruction::new(Op::TypeImage); - instruction.set_result(id); - instruction.add_operand(sampled_type_id); - instruction.add_operand(dim as u32); - - let (depth, multi, sampled) = match image_class { - crate::ImageClass::Sampled { kind: _, multi } => (false, multi, true), - crate::ImageClass::Depth => (true, false, true), - crate::ImageClass::Storage(_) => (false, false, false), - }; - instruction.add_operand(depth as u32); - instruction.add_operand(arrayed as u32); - instruction.add_operand(multi as u32); - instruction.add_operand(sampled as u32); - - let format = match image_class { - crate::ImageClass::Storage(format) => match format { - crate::StorageFormat::R8Unorm => spirv::ImageFormat::R8, - crate::StorageFormat::R8Snorm => spirv::ImageFormat::R8Snorm, - crate::StorageFormat::R8Uint => spirv::ImageFormat::R8ui, - crate::StorageFormat::R8Sint => spirv::ImageFormat::R8i, - crate::StorageFormat::R16Uint => spirv::ImageFormat::R16ui, - crate::StorageFormat::R16Sint => spirv::ImageFormat::R16i, - crate::StorageFormat::R16Float => spirv::ImageFormat::R16f, - crate::StorageFormat::Rg8Unorm => spirv::ImageFormat::Rg8, - crate::StorageFormat::Rg8Snorm => spirv::ImageFormat::Rg8Snorm, - crate::StorageFormat::Rg8Uint => spirv::ImageFormat::Rg8ui, - crate::StorageFormat::Rg8Sint => spirv::ImageFormat::Rg8i, - crate::StorageFormat::R32Uint => spirv::ImageFormat::R32ui, - crate::StorageFormat::R32Sint => spirv::ImageFormat::R32i, - crate::StorageFormat::R32Float => spirv::ImageFormat::R32f, - crate::StorageFormat::Rg16Uint => spirv::ImageFormat::Rg16ui, - crate::StorageFormat::Rg16Sint => spirv::ImageFormat::Rg16i, - crate::StorageFormat::Rg16Float => spirv::ImageFormat::Rg16f, - crate::StorageFormat::Rgba8Unorm => spirv::ImageFormat::Rgba8, - crate::StorageFormat::Rgba8Snorm => spirv::ImageFormat::Rgba8Snorm, - crate::StorageFormat::Rgba8Uint => spirv::ImageFormat::Rgba8ui, - crate::StorageFormat::Rgba8Sint => spirv::ImageFormat::Rgba8i, - crate::StorageFormat::Rgb10a2Unorm => spirv::ImageFormat::Rgb10a2ui, - crate::StorageFormat::Rg11b10Float => spirv::ImageFormat::R11fG11fB10f, - crate::StorageFormat::Rg32Uint => spirv::ImageFormat::Rg32ui, - crate::StorageFormat::Rg32Sint => spirv::ImageFormat::Rg32i, - crate::StorageFormat::Rg32Float => spirv::ImageFormat::Rg32f, - crate::StorageFormat::Rgba16Uint => spirv::ImageFormat::Rgba16ui, - crate::StorageFormat::Rgba16Sint => spirv::ImageFormat::Rgba16i, - crate::StorageFormat::Rgba16Float => spirv::ImageFormat::Rgba16f, - crate::StorageFormat::Rgba32Uint => spirv::ImageFormat::Rgba32ui, - crate::StorageFormat::Rgba32Sint => spirv::ImageFormat::Rgba32i, - crate::StorageFormat::Rgba32Float => spirv::ImageFormat::Rgba32f, - }, - _ => spirv::ImageFormat::Unknown, - }; - - instruction.add_operand(format as u32); - instruction -} - -pub(super) fn instruction_type_sampler(id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::TypeSampler); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_type_sampled_image(id: Word, image_type_id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::TypeSampledImage); - instruction.set_result(id); - instruction.add_operand(image_type_id); - instruction -} - -pub(super) fn instruction_type_array( - id: Word, - element_type_id: Word, - length_id: Word, -) -> Instruction { - let mut instruction = Instruction::new(Op::TypeArray); - instruction.set_result(id); - instruction.add_operand(element_type_id); - instruction.add_operand(length_id); - instruction -} - -pub(super) fn instruction_type_runtime_array(id: Word, element_type_id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::TypeRuntimeArray); - instruction.set_result(id); - instruction.add_operand(element_type_id); - instruction -} - -pub(super) fn instruction_type_struct(id: Word, member_ids: &[Word]) -> Instruction { - let mut instruction = Instruction::new(Op::TypeStruct); - instruction.set_result(id); - - for member_id in member_ids { - instruction.add_operand(*member_id) - } - - instruction -} - -pub(super) fn instruction_type_pointer( - id: Word, - storage_class: spirv::StorageClass, - type_id: Word, -) -> Instruction { - let mut instruction = Instruction::new(Op::TypePointer); - instruction.set_result(id); - instruction.add_operand(storage_class as u32); - instruction.add_operand(type_id); - instruction -} - -pub(super) fn instruction_type_function( - id: Word, - return_type_id: Word, - parameter_ids: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::TypeFunction); - instruction.set_result(id); - instruction.add_operand(return_type_id); - - for parameter_id in parameter_ids { - instruction.add_operand(*parameter_id); - } - - instruction -} - -// -// Constant-Creation Instructions -// - -pub(super) fn instruction_constant_true(result_type_id: Word, id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::ConstantTrue); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_constant_false(result_type_id: Word, id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::ConstantFalse); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_constant(result_type_id: Word, id: Word, values: &[Word]) -> Instruction { - let mut instruction = Instruction::new(Op::Constant); - instruction.set_type(result_type_id); - instruction.set_result(id); - - for value in values { - instruction.add_operand(*value); - } - - instruction -} - -pub(super) fn instruction_constant_composite( - result_type_id: Word, - id: Word, - constituent_ids: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::ConstantComposite); - instruction.set_type(result_type_id); - instruction.set_result(id); - - for constituent_id in constituent_ids { - instruction.add_operand(*constituent_id); - } - - instruction -} - -// -// Memory Instructions -// - -pub(super) fn instruction_variable( - result_type_id: Word, - id: Word, - storage_class: spirv::StorageClass, - initializer_id: Option, -) -> Instruction { - let mut instruction = Instruction::new(Op::Variable); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(storage_class as u32); - - if let Some(initializer_id) = initializer_id { - instruction.add_operand(initializer_id); - } - - instruction -} - -pub(super) fn instruction_load( - result_type_id: Word, - id: Word, - pointer_id: Word, - memory_access: Option, -) -> Instruction { - let mut instruction = Instruction::new(Op::Load); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(pointer_id); - - if let Some(memory_access) = memory_access { - instruction.add_operand(memory_access.bits()); - } - - instruction -} - -pub(super) fn instruction_store( - pointer_type_id: Word, - object_id: Word, - memory_access: Option, -) -> Instruction { - let mut instruction = Instruction::new(Op::Store); - instruction.add_operand(pointer_type_id); - instruction.add_operand(object_id); - - if let Some(memory_access) = memory_access { - instruction.add_operand(memory_access.bits()); - } - - instruction -} - -pub(super) fn instruction_access_chain( - result_type_id: Word, - id: Word, - base_id: Word, - index_ids: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::AccessChain); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(base_id); - - for index_id in index_ids { - instruction.add_operand(*index_id); - } - - instruction -} - -// -// Function Instructions -// - -pub(super) fn instruction_function( - return_type_id: Word, - id: Word, - function_control: spirv::FunctionControl, - function_type_id: Word, -) -> Instruction { - let mut instruction = Instruction::new(Op::Function); - instruction.set_type(return_type_id); - instruction.set_result(id); - instruction.add_operand(function_control.bits()); - instruction.add_operand(function_type_id); - instruction -} - -pub(super) fn instruction_function_parameter(result_type_id: Word, id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::FunctionParameter); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_function_end() -> Instruction { - Instruction::new(Op::FunctionEnd) -} - -pub(super) fn instruction_function_call( - result_type_id: Word, - id: Word, - function_id: Word, - argument_ids: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::FunctionCall); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(function_id); - - for argument_id in argument_ids { - instruction.add_operand(*argument_id); - } - - instruction -} - -// -// Image Instructions -// -pub(super) fn instruction_sampled_image( - result_type_id: Word, - id: Word, - image: Word, - sampler: Word, -) -> Instruction { - let mut instruction = Instruction::new(Op::SampledImage); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(image); - instruction.add_operand(sampler); - instruction -} - pub(super) enum SampleLod { Explicit, Implicit, } -pub(super) fn instruction_image_sample( - result_type_id: Word, - id: Word, - lod: SampleLod, - sampled_image: Word, - coordinates: Word, - depth_ref: Option, -) -> Instruction { - let op = match (lod, depth_ref) { - (SampleLod::Explicit, None) => Op::ImageSampleExplicitLod, - (SampleLod::Implicit, None) => Op::ImageSampleImplicitLod, - (SampleLod::Explicit, Some(_)) => Op::ImageSampleDrefExplicitLod, - (SampleLod::Implicit, Some(_)) => Op::ImageSampleDrefImplicitLod, - }; - - let mut instruction = Instruction::new(op); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(sampled_image); - instruction.add_operand(coordinates); - if let Some(dref) = depth_ref { - instruction.add_operand(dref); - } - - instruction -} - -// -// Conversion Instructions -// -pub(super) fn instruction_unary( - op: Op, - result_type_id: Word, - id: Word, - value: Word, -) -> Instruction { - let mut instruction = Instruction::new(op); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(value); - instruction -} - -// -// Composite Instructions -// - -pub(super) fn instruction_composite_construct( - result_type_id: Word, - id: Word, - constituent_ids: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::CompositeConstruct); - instruction.set_type(result_type_id); - instruction.set_result(id); - - for constituent_id in constituent_ids { - instruction.add_operand(*constituent_id); - } - - instruction -} - -pub(super) fn instruction_composite_extract( - result_type_id: Word, - id: Word, - composite_id: Word, - indices: &[Word], -) -> Instruction { - let mut instruction = Instruction::new(Op::CompositeExtract); - instruction.set_type(result_type_id); - instruction.set_result(id); - - instruction.add_operand(composite_id); - for index in indices { - instruction.add_operand(*index); - } - - instruction -} - -pub(super) fn instruction_vector_extract_dynamic( - result_type_id: Word, - id: Word, - vector_id: Word, - index_id: Word, -) -> Instruction { - let mut instruction = Instruction::new(Op::VectorExtractDynamic); - instruction.set_type(result_type_id); - instruction.set_result(id); - - instruction.add_operand(vector_id); - instruction.add_operand(index_id); - - instruction -} - -// -// Arithmetic Instructions -// -pub(super) fn instruction_binary( - op: Op, - result_type_id: Word, - id: Word, - operand_1: Word, - operand_2: Word, -) -> Instruction { - let mut instruction = Instruction::new(op); - instruction.set_type(result_type_id); - instruction.set_result(id); - instruction.add_operand(operand_1); - instruction.add_operand(operand_2); - instruction -} - -// -// Bit Instructions -// - -// -// Relational and Logical Instructions -// - -// -// Derivative Instructions -// - -// -// Control-Flow Instructions -// - -pub(super) fn instruction_selection_merge( - merge_id: Word, - selection_control: spirv::SelectionControl, -) -> Instruction { - let mut instruction = Instruction::new(Op::SelectionMerge); - instruction.add_operand(merge_id); - instruction.add_operand(selection_control.bits()); - instruction -} - -pub(super) fn instruction_loop_merge( - merge_id: Word, - continuing_id: Word, - selection_control: spirv::SelectionControl, -) -> Instruction { - let mut instruction = Instruction::new(Op::LoopMerge); - instruction.add_operand(merge_id); - instruction.add_operand(continuing_id); - instruction.add_operand(selection_control.bits()); - instruction -} - -pub(super) fn instruction_label(id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::Label); - instruction.set_result(id); - instruction -} - -pub(super) fn instruction_branch(id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::Branch); - instruction.add_operand(id); - instruction -} - -// TODO Branch Weights not implemented. -pub(super) fn instruction_branch_conditional( - condition_id: Word, - true_label: Word, - false_label: Word, -) -> Instruction { - let mut instruction = Instruction::new(Op::BranchConditional); - instruction.add_operand(condition_id); - instruction.add_operand(true_label); - instruction.add_operand(false_label); - instruction -} - -pub struct Case { +pub(super) struct Case { pub value: Word, pub label_id: Word, } -pub(super) fn instruction_switch( - selector_id: Word, - default_id: Word, - cases: &[Case], -) -> Instruction { - let mut instruction = Instruction::new(Op::Switch); - instruction.add_operand(selector_id); - instruction.add_operand(default_id); - for case in cases { - instruction.add_operand(case.value); - instruction.add_operand(case.label_id); +impl super::Instruction { + // + // Debug Instructions + // + + pub(super) fn source(source_language: spirv::SourceLanguage, version: u32) -> Self { + let mut instruction = Self::new(Op::Source); + instruction.add_operand(source_language as u32); + instruction.add_operands(helpers::bytes_to_words(&version.to_le_bytes())); + instruction } - instruction + + pub(super) fn name(target_id: Word, name: &str) -> Self { + let mut instruction = Self::new(Op::Name); + instruction.add_operand(target_id); + instruction.add_operands(helpers::string_to_words(name)); + instruction + } + + pub(super) fn member_name(target_id: Word, member: Word, name: &str) -> Self { + let mut instruction = Self::new(Op::MemberName); + instruction.add_operand(target_id); + instruction.add_operand(member); + instruction.add_operands(helpers::string_to_words(name)); + instruction + } + + // + // Annotation Instructions + // + + pub(super) fn decorate( + target_id: Word, + decoration: spirv::Decoration, + operands: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::Decorate); + instruction.add_operand(target_id); + instruction.add_operand(decoration as u32); + for operand in operands { + instruction.add_operand(*operand) + } + instruction + } + + pub(super) fn member_decorate( + target_id: Word, + member_index: Word, + decoration: spirv::Decoration, + operands: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::MemberDecorate); + instruction.add_operand(target_id); + instruction.add_operand(member_index); + instruction.add_operand(decoration as u32); + for operand in operands { + instruction.add_operand(*operand) + } + instruction + } + + // + // Extension Instructions + // + + pub(super) fn ext_inst_import(id: Word, name: &str) -> Self { + let mut instruction = Self::new(Op::ExtInstImport); + instruction.set_result(id); + instruction.add_operands(helpers::string_to_words(name)); + instruction + } + + pub(super) fn ext_inst( + set_id: Word, + op: spirv::GLOp, + result_type_id: Word, + id: Word, + operands: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::ExtInst); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(set_id); + instruction.add_operand(op as u32); + for operand in operands { + instruction.add_operand(*operand) + } + instruction + } + + // + // Mode-Setting Instructions + // + + pub(super) fn memory_model( + addressing_model: spirv::AddressingModel, + memory_model: spirv::MemoryModel, + ) -> Self { + let mut instruction = Self::new(Op::MemoryModel); + instruction.add_operand(addressing_model as u32); + instruction.add_operand(memory_model as u32); + instruction + } + + pub(super) fn entry_point( + execution_model: spirv::ExecutionModel, + entry_point_id: Word, + name: &str, + interface_ids: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::EntryPoint); + instruction.add_operand(execution_model as u32); + instruction.add_operand(entry_point_id); + instruction.add_operands(helpers::string_to_words(name)); + + for interface_id in interface_ids { + instruction.add_operand(*interface_id); + } + + instruction + } + + pub(super) fn execution_mode( + entry_point_id: Word, + execution_mode: spirv::ExecutionMode, + args: &[Word], + ) -> Self { + let mut instruction = Self::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 + } + + pub(super) fn capability(capability: spirv::Capability) -> Self { + let mut instruction = Self::new(Op::Capability); + instruction.add_operand(capability as u32); + instruction + } + + // + // Type-Declaration Instructions + // + + pub(super) fn type_void(id: Word) -> Self { + let mut instruction = Self::new(Op::TypeVoid); + instruction.set_result(id); + instruction + } + + pub(super) fn type_bool(id: Word) -> Self { + let mut instruction = Self::new(Op::TypeBool); + instruction.set_result(id); + instruction + } + + pub(super) fn type_int(id: Word, width: Word, signedness: Signedness) -> Self { + let mut instruction = Self::new(Op::TypeInt); + instruction.set_result(id); + instruction.add_operand(width); + instruction.add_operand(signedness as u32); + instruction + } + + pub(super) fn type_float(id: Word, width: Word) -> Self { + let mut instruction = Self::new(Op::TypeFloat); + instruction.set_result(id); + instruction.add_operand(width); + instruction + } + + pub(super) fn type_vector( + id: Word, + component_type_id: Word, + component_count: crate::VectorSize, + ) -> Self { + let mut instruction = Self::new(Op::TypeVector); + instruction.set_result(id); + instruction.add_operand(component_type_id); + instruction.add_operand(component_count as u32); + instruction + } + + pub(super) fn type_matrix( + id: Word, + column_type_id: Word, + column_count: crate::VectorSize, + ) -> Self { + let mut instruction = Self::new(Op::TypeMatrix); + instruction.set_result(id); + instruction.add_operand(column_type_id); + instruction.add_operand(column_count as u32); + instruction + } + + pub(super) fn type_image( + id: Word, + sampled_type_id: Word, + dim: spirv::Dim, + arrayed: bool, + image_class: crate::ImageClass, + ) -> Self { + let mut instruction = Self::new(Op::TypeImage); + instruction.set_result(id); + instruction.add_operand(sampled_type_id); + instruction.add_operand(dim as u32); + + let (depth, multi, sampled) = match image_class { + crate::ImageClass::Sampled { kind: _, multi } => (false, multi, true), + crate::ImageClass::Depth => (true, false, true), + crate::ImageClass::Storage(_) => (false, false, false), + }; + instruction.add_operand(depth as u32); + instruction.add_operand(arrayed as u32); + instruction.add_operand(multi as u32); + instruction.add_operand(sampled as u32); + + let format = match image_class { + crate::ImageClass::Storage(format) => match format { + crate::StorageFormat::R8Unorm => spirv::ImageFormat::R8, + crate::StorageFormat::R8Snorm => spirv::ImageFormat::R8Snorm, + crate::StorageFormat::R8Uint => spirv::ImageFormat::R8ui, + crate::StorageFormat::R8Sint => spirv::ImageFormat::R8i, + crate::StorageFormat::R16Uint => spirv::ImageFormat::R16ui, + crate::StorageFormat::R16Sint => spirv::ImageFormat::R16i, + crate::StorageFormat::R16Float => spirv::ImageFormat::R16f, + crate::StorageFormat::Rg8Unorm => spirv::ImageFormat::Rg8, + crate::StorageFormat::Rg8Snorm => spirv::ImageFormat::Rg8Snorm, + crate::StorageFormat::Rg8Uint => spirv::ImageFormat::Rg8ui, + crate::StorageFormat::Rg8Sint => spirv::ImageFormat::Rg8i, + crate::StorageFormat::R32Uint => spirv::ImageFormat::R32ui, + crate::StorageFormat::R32Sint => spirv::ImageFormat::R32i, + crate::StorageFormat::R32Float => spirv::ImageFormat::R32f, + crate::StorageFormat::Rg16Uint => spirv::ImageFormat::Rg16ui, + crate::StorageFormat::Rg16Sint => spirv::ImageFormat::Rg16i, + crate::StorageFormat::Rg16Float => spirv::ImageFormat::Rg16f, + crate::StorageFormat::Rgba8Unorm => spirv::ImageFormat::Rgba8, + crate::StorageFormat::Rgba8Snorm => spirv::ImageFormat::Rgba8Snorm, + crate::StorageFormat::Rgba8Uint => spirv::ImageFormat::Rgba8ui, + crate::StorageFormat::Rgba8Sint => spirv::ImageFormat::Rgba8i, + crate::StorageFormat::Rgb10a2Unorm => spirv::ImageFormat::Rgb10a2ui, + crate::StorageFormat::Rg11b10Float => spirv::ImageFormat::R11fG11fB10f, + crate::StorageFormat::Rg32Uint => spirv::ImageFormat::Rg32ui, + crate::StorageFormat::Rg32Sint => spirv::ImageFormat::Rg32i, + crate::StorageFormat::Rg32Float => spirv::ImageFormat::Rg32f, + crate::StorageFormat::Rgba16Uint => spirv::ImageFormat::Rgba16ui, + crate::StorageFormat::Rgba16Sint => spirv::ImageFormat::Rgba16i, + crate::StorageFormat::Rgba16Float => spirv::ImageFormat::Rgba16f, + crate::StorageFormat::Rgba32Uint => spirv::ImageFormat::Rgba32ui, + crate::StorageFormat::Rgba32Sint => spirv::ImageFormat::Rgba32i, + crate::StorageFormat::Rgba32Float => spirv::ImageFormat::Rgba32f, + }, + _ => spirv::ImageFormat::Unknown, + }; + + instruction.add_operand(format as u32); + instruction + } + + pub(super) fn type_sampler(id: Word) -> Self { + let mut instruction = Self::new(Op::TypeSampler); + instruction.set_result(id); + instruction + } + + pub(super) fn type_sampled_image(id: Word, image_type_id: Word) -> Self { + let mut instruction = Self::new(Op::TypeSampledImage); + instruction.set_result(id); + instruction.add_operand(image_type_id); + instruction + } + + pub(super) fn type_array(id: Word, element_type_id: Word, length_id: Word) -> Self { + let mut instruction = Self::new(Op::TypeArray); + instruction.set_result(id); + instruction.add_operand(element_type_id); + instruction.add_operand(length_id); + instruction + } + + pub(super) fn type_runtime_array(id: Word, element_type_id: Word) -> Self { + let mut instruction = Self::new(Op::TypeRuntimeArray); + instruction.set_result(id); + instruction.add_operand(element_type_id); + instruction + } + + pub(super) fn type_struct(id: Word, member_ids: &[Word]) -> Self { + let mut instruction = Self::new(Op::TypeStruct); + instruction.set_result(id); + + for member_id in member_ids { + instruction.add_operand(*member_id) + } + + instruction + } + + pub(super) fn type_pointer( + id: Word, + storage_class: spirv::StorageClass, + type_id: Word, + ) -> Self { + let mut instruction = Self::new(Op::TypePointer); + instruction.set_result(id); + instruction.add_operand(storage_class as u32); + instruction.add_operand(type_id); + instruction + } + + pub(super) fn type_function(id: Word, return_type_id: Word, parameter_ids: &[Word]) -> Self { + let mut instruction = Self::new(Op::TypeFunction); + instruction.set_result(id); + instruction.add_operand(return_type_id); + + for parameter_id in parameter_ids { + instruction.add_operand(*parameter_id); + } + + instruction + } + + // + // Constant-Creation Instructions + // + + pub(super) fn constant_true(result_type_id: Word, id: Word) -> Self { + let mut instruction = Self::new(Op::ConstantTrue); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction + } + + pub(super) fn constant_false(result_type_id: Word, id: Word) -> Self { + let mut instruction = Self::new(Op::ConstantFalse); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction + } + + pub(super) fn constant(result_type_id: Word, id: Word, values: &[Word]) -> Self { + let mut instruction = Self::new(Op::Constant); + instruction.set_type(result_type_id); + instruction.set_result(id); + + for value in values { + instruction.add_operand(*value); + } + + instruction + } + + pub(super) fn constant_composite( + result_type_id: Word, + id: Word, + constituent_ids: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::ConstantComposite); + instruction.set_type(result_type_id); + instruction.set_result(id); + + for constituent_id in constituent_ids { + instruction.add_operand(*constituent_id); + } + + instruction + } + + // + // Memory Instructions + // + + pub(super) fn variable( + result_type_id: Word, + id: Word, + storage_class: spirv::StorageClass, + initializer_id: Option, + ) -> Self { + let mut instruction = Self::new(Op::Variable); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(storage_class as u32); + + if let Some(initializer_id) = initializer_id { + instruction.add_operand(initializer_id); + } + + instruction + } + + pub(super) fn load( + result_type_id: Word, + id: Word, + pointer_id: Word, + memory_access: Option, + ) -> Self { + let mut instruction = Self::new(Op::Load); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(pointer_id); + + if let Some(memory_access) = memory_access { + instruction.add_operand(memory_access.bits()); + } + + instruction + } + + pub(super) fn store( + pointer_type_id: Word, + object_id: Word, + memory_access: Option, + ) -> Self { + let mut instruction = Self::new(Op::Store); + instruction.add_operand(pointer_type_id); + instruction.add_operand(object_id); + + if let Some(memory_access) = memory_access { + instruction.add_operand(memory_access.bits()); + } + + instruction + } + + pub(super) fn access_chain( + result_type_id: Word, + id: Word, + base_id: Word, + index_ids: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::AccessChain); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(base_id); + + for index_id in index_ids { + instruction.add_operand(*index_id); + } + + instruction + } + + // + // Function Instructions + // + + pub(super) fn function( + return_type_id: Word, + id: Word, + function_control: spirv::FunctionControl, + function_type_id: Word, + ) -> Self { + let mut instruction = Self::new(Op::Function); + instruction.set_type(return_type_id); + instruction.set_result(id); + instruction.add_operand(function_control.bits()); + instruction.add_operand(function_type_id); + instruction + } + + pub(super) fn function_parameter(result_type_id: Word, id: Word) -> Self { + let mut instruction = Self::new(Op::FunctionParameter); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction + } + + pub(super) fn function_end() -> Self { + Self::new(Op::FunctionEnd) + } + + pub(super) fn function_call( + result_type_id: Word, + id: Word, + function_id: Word, + argument_ids: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::FunctionCall); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(function_id); + + for argument_id in argument_ids { + instruction.add_operand(*argument_id); + } + + instruction + } + + // + // Image Instructions + // + + pub(super) fn sampled_image( + result_type_id: Word, + id: Word, + image: Word, + sampler: Word, + ) -> Self { + let mut instruction = Self::new(Op::SampledImage); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(image); + instruction.add_operand(sampler); + instruction + } + + pub(super) fn image_sample( + result_type_id: Word, + id: Word, + lod: SampleLod, + sampled_image: Word, + coordinates: Word, + depth_ref: Option, + ) -> Self { + let op = match (lod, depth_ref) { + (SampleLod::Explicit, None) => Op::ImageSampleExplicitLod, + (SampleLod::Implicit, None) => Op::ImageSampleImplicitLod, + (SampleLod::Explicit, Some(_)) => Op::ImageSampleDrefExplicitLod, + (SampleLod::Implicit, Some(_)) => Op::ImageSampleDrefImplicitLod, + }; + + let mut instruction = Self::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(sampled_image); + instruction.add_operand(coordinates); + if let Some(dref) = depth_ref { + instruction.add_operand(dref); + } + + instruction + } + + // + // Conversion Instructions + // + pub(super) fn unary(op: Op, result_type_id: Word, id: Word, value: Word) -> Self { + let mut instruction = Self::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(value); + instruction + } + + // + // Composite Instructions + // + + pub(super) fn composite_construct( + result_type_id: Word, + id: Word, + constituent_ids: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::CompositeConstruct); + instruction.set_type(result_type_id); + instruction.set_result(id); + + for constituent_id in constituent_ids { + instruction.add_operand(*constituent_id); + } + + instruction + } + + pub(super) fn composite_extract( + result_type_id: Word, + id: Word, + composite_id: Word, + indices: &[Word], + ) -> Self { + let mut instruction = Self::new(Op::CompositeExtract); + instruction.set_type(result_type_id); + instruction.set_result(id); + + instruction.add_operand(composite_id); + for index in indices { + instruction.add_operand(*index); + } + + instruction + } + + pub(super) fn vector_extract_dynamic( + result_type_id: Word, + id: Word, + vector_id: Word, + index_id: Word, + ) -> Self { + let mut instruction = Self::new(Op::VectorExtractDynamic); + instruction.set_type(result_type_id); + instruction.set_result(id); + + instruction.add_operand(vector_id); + instruction.add_operand(index_id); + + instruction + } + + // + // Arithmetic Instructions + // + pub(super) fn binary( + op: Op, + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, + ) -> Self { + let mut instruction = Self::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(operand_1); + instruction.add_operand(operand_2); + instruction + } + + // + // Bit Instructions + // + + // + // Relational and Logical Instructions + // + + // + // Derivative Instructions + // + + // + // Control-Flow Instructions + // + + pub(super) fn selection_merge( + merge_id: Word, + selection_control: spirv::SelectionControl, + ) -> Self { + let mut instruction = Self::new(Op::SelectionMerge); + instruction.add_operand(merge_id); + instruction.add_operand(selection_control.bits()); + instruction + } + + pub(super) fn loop_merge( + merge_id: Word, + continuing_id: Word, + selection_control: spirv::SelectionControl, + ) -> Self { + let mut instruction = Self::new(Op::LoopMerge); + instruction.add_operand(merge_id); + instruction.add_operand(continuing_id); + instruction.add_operand(selection_control.bits()); + instruction + } + + pub(super) fn label(id: Word) -> Self { + let mut instruction = Self::new(Op::Label); + instruction.set_result(id); + instruction + } + + pub(super) fn branch(id: Word) -> Self { + let mut instruction = Self::new(Op::Branch); + instruction.add_operand(id); + instruction + } + + // TODO Branch Weights not implemented. + pub(super) fn branch_conditional( + condition_id: Word, + true_label: Word, + false_label: Word, + ) -> Self { + let mut instruction = Self::new(Op::BranchConditional); + instruction.add_operand(condition_id); + instruction.add_operand(true_label); + instruction.add_operand(false_label); + instruction + } + + pub(super) fn switch(selector_id: Word, default_id: Word, cases: &[Case]) -> Self { + let mut instruction = Self::new(Op::Switch); + instruction.add_operand(selector_id); + instruction.add_operand(default_id); + for case in cases { + instruction.add_operand(case.value); + instruction.add_operand(case.label_id); + } + instruction + } + + pub(super) fn kill() -> Self { + Self::new(Op::Kill) + } + + pub(super) fn return_void() -> Self { + Self::new(Op::Return) + } + + pub(super) fn return_value(value_id: Word) -> Self { + let mut instruction = Self::new(Op::ReturnValue); + instruction.add_operand(value_id); + instruction + } + + // + // Atomic Instructions + // + + // + // Primitive Instructions + // } - -pub(super) fn instruction_kill() -> Instruction { - Instruction::new(Op::Kill) -} - -pub(super) fn instruction_return() -> Instruction { - Instruction::new(Op::Return) -} - -pub(super) fn instruction_return_value(value_id: Word) -> Instruction { - let mut instruction = Instruction::new(Op::ReturnValue); - instruction.add_operand(value_id); - instruction -} - -// -// Atomic Instructions -// - -// -// Primitive Instructions -// diff --git a/src/back/spv/mod.rs b/src/back/spv/mod.rs index 011645b50e..bbad11f855 100644 --- a/src/back/spv/mod.rs +++ b/src/back/spv/mod.rs @@ -44,7 +44,7 @@ struct LogicalLayout { function_definitions: Vec, } -pub(self) struct Instruction { +struct Instruction { op: spirv::Op, wc: u32, type_id: Option, diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index e9bdda5ee6..1a4f0d24b7 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -58,7 +58,7 @@ impl Function { instruction.to_words(sink); } for (index, block) in self.blocks.iter().enumerate() { - super::instructions::instruction_label(block.label_id).to_words(sink); + Instruction::label(block.label_id).to_words(sink); if index == 0 { for local_var in self.variables.values() { local_var.instruction.to_words(sink); @@ -318,8 +318,7 @@ impl Writer { _ => { let storage_class = self.parse_to_spirv_storage_class(class); let id = self.generate_id(); - let instruction = - super::instructions::instruction_type_pointer(id, storage_class, ty_id); + let instruction = Instruction::type_pointer(id, storage_class, ty_id); instruction.to_words(&mut self.logical_layout.declarations); self.lookup_type.insert( LookupType::Local(LocalType::Pointer { @@ -336,14 +335,14 @@ impl Writer { fn create_pointer_type(&mut self, type_id: spirv::Word, class: spirv::StorageClass) -> Word { let id = self.generate_id(); - let instruction = super::instructions::instruction_type_pointer(id, class, type_id); + let instruction = Instruction::type_pointer(id, class, type_id); instruction.to_words(&mut self.logical_layout.declarations); id } fn create_constant(&mut self, type_id: Word, value: &[Word]) -> Word { let id = self.generate_id(); - let instruction = super::instructions::instruction_constant(type_id, id, value); + let instruction = Instruction::constant(type_id, id, value); instruction.to_words(&mut self.logical_layout.declarations); id } @@ -372,8 +371,7 @@ impl Writer { if self.flags.contains(WriterFlags::DEBUG) { if let Some(ref name) = variable.name { - self.debugs - .push(super::instructions::instruction_name(id, name)); + self.debugs.push(Instruction::name(id, name)); } } @@ -382,7 +380,7 @@ impl Writer { .map(|constant| self.lookup_constant[&constant]); let pointer_type_id = self.get_pointer_id(&ir_module.types, variable.ty, crate::StorageClass::Function)?; - let instruction = super::instructions::instruction_variable( + let instruction = Instruction::variable( pointer_type_id, id, spirv::StorageClass::Function, @@ -406,10 +404,7 @@ impl Writer { parameter_type_ids.push(parameter_type_id); function .parameters - .push(super::instructions::instruction_function_parameter( - parameter_type_id, - id, - )); + .push(Instruction::function_parameter(parameter_type_id, id)); } let lookup_function_type = LookupFunctionType { @@ -419,7 +414,7 @@ impl Writer { let function_id = self.generate_id(); let function_type = self.get_function_type(lookup_function_type); - function.signature = Some(super::instructions::instruction_function( + function.signature = Some(Instruction::function( return_type_id, function_id, spirv::FunctionControl::empty(), @@ -438,8 +433,7 @@ impl Writer { )?; function.to_words(&mut self.logical_layout.function_definitions); - super::instructions::instruction_function_end() - .to_words(&mut self.logical_layout.function_definitions); + Instruction::function_end().to_words(&mut self.logical_layout.function_definitions); Ok(function_id) } @@ -475,14 +469,14 @@ impl Writer { crate::ShaderStage::Fragment => { let execution_mode = spirv::ExecutionMode::OriginUpperLeft; self.check(execution_mode.required_capabilities())?; - super::instructions::instruction_execution_mode(function_id, execution_mode, &[]) + Instruction::execution_mode(function_id, execution_mode, &[]) .to_words(&mut self.logical_layout.execution_modes); spirv::ExecutionModel::Fragment } crate::ShaderStage::Compute => { let execution_mode = spirv::ExecutionMode::LocalSize; self.check(execution_mode.required_capabilities())?; - super::instructions::instruction_execution_mode( + Instruction::execution_mode( function_id, execution_mode, &entry_point.workgroup_size, @@ -494,11 +488,10 @@ impl Writer { self.check(exec_model.required_capabilities())?; if self.flags.contains(WriterFlags::DEBUG) { - self.debugs - .push(super::instructions::instruction_name(function_id, name)); + self.debugs.push(Instruction::name(function_id, name)); } - Ok(super::instructions::instruction_entry_point( + Ok(Instruction::entry_point( exec_model, function_id, name, @@ -509,18 +502,14 @@ impl Writer { fn write_scalar(&self, id: Word, kind: crate::ScalarKind, width: crate::Bytes) -> Instruction { let bits = (width * BITS_PER_BYTE) as u32; match kind { - crate::ScalarKind::Sint => super::instructions::instruction_type_int( - id, - bits, - super::instructions::Signedness::Signed, - ), - crate::ScalarKind::Uint => super::instructions::instruction_type_int( - id, - bits, - super::instructions::Signedness::Unsigned, - ), - crate::ScalarKind::Float => super::instructions::instruction_type_float(id, bits), - crate::ScalarKind::Bool => super::instructions::instruction_type_bool(id), + crate::ScalarKind::Sint => { + Instruction::type_int(id, bits, super::instructions::Signedness::Signed) + } + crate::ScalarKind::Uint => { + Instruction::type_int(id, bits, super::instructions::Signedness::Unsigned) + } + crate::ScalarKind::Float => Instruction::type_float(id, bits), + crate::ScalarKind::Bool => Instruction::type_bool(id), } } @@ -553,7 +542,7 @@ impl Writer { LocalType::Vector { size, kind, width } => { let scalar_id = self.get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width }))?; - super::instructions::instruction_type_vector(id, scalar_id, size) + Instruction::type_vector(id, scalar_id, size) } LocalType::Matrix { columns, @@ -568,14 +557,14 @@ impl Writer { width, }), )?; - super::instructions::instruction_type_matrix(id, vector_id, columns) + Instruction::type_matrix(id, vector_id, columns) } LocalType::Pointer { .. } => { return Err(Error::FeatureNotImplemented("pointer declaration")) } LocalType::SampledImage { image_type } => { let image_type_id = self.get_type_id(arena, LookupType::Handle(image_type))?; - super::instructions::instruction_type_sampled_image(id, image_type_id) + Instruction::type_sampled_image(id, image_type_id) } }; @@ -594,8 +583,7 @@ impl Writer { if self.flags.contains(WriterFlags::DEBUG) { if let Some(ref name) = ty.name { - self.debugs - .push(super::instructions::instruction_name(id, name)); + self.debugs.push(Instruction::name(id, name)); } } @@ -612,7 +600,7 @@ impl Writer { LookupType::Local(LocalType::Vector { size, kind, width }), id, ); - super::instructions::instruction_type_vector(id, scalar_id, size) + Instruction::type_vector(id, scalar_id, size) } crate::TypeInner::Matrix { columns, @@ -635,7 +623,7 @@ impl Writer { }), id, ); - super::instructions::instruction_type_matrix(id, vector_id, columns) + Instruction::type_matrix(id, vector_id, columns) } crate::TypeInner::Image { dim, @@ -659,30 +647,25 @@ impl Writer { let type_id = self.get_type_id(arena, LookupType::Local(local_type))?; let dim = map_dim(dim); self.check(dim.required_capabilities())?; - super::instructions::instruction_type_image(id, type_id, dim, arrayed, class) - } - crate::TypeInner::Sampler { comparison: _ } => { - super::instructions::instruction_type_sampler(id) + Instruction::type_image(id, type_id, dim, arrayed, class) } + crate::TypeInner::Sampler { comparison: _ } => Instruction::type_sampler(id), crate::TypeInner::Array { base, size, stride } => { if let Some(array_stride) = stride { - self.annotations - .push(super::instructions::instruction_decorate( - id, - spirv::Decoration::ArrayStride, - &[array_stride.get()], - )); + self.annotations.push(Instruction::decorate( + id, + spirv::Decoration::ArrayStride, + &[array_stride.get()], + )); } let type_id = self.get_type_id(arena, LookupType::Handle(base))?; match size { crate::ArraySize::Constant(const_handle) => { let length_id = self.lookup_constant[&const_handle]; - super::instructions::instruction_type_array(id, type_id, length_id) - } - crate::ArraySize::Dynamic => { - super::instructions::instruction_type_runtime_array(id, type_id) + Instruction::type_array(id, type_id, length_id) } + crate::ArraySize::Dynamic => Instruction::type_runtime_array(id, type_id), } } crate::TypeInner::Struct { @@ -696,11 +679,7 @@ impl Writer { spirv::Decoration::BufferBlock }; self.annotations - .push(super::instructions::instruction_decorate( - id, - decoration, - &[], - )); + .push(Instruction::decorate(id, decoration, &[])); } let mut current_offset = 0; @@ -708,13 +687,12 @@ impl Writer { for (index, member) in members.iter().enumerate() { let layout = self.layouter.resolve(member.ty); current_offset += layout.pad(current_offset); - self.annotations - .push(super::instructions::instruction_member_decorate( - id, - index as u32, - spirv::Decoration::Offset, - &[current_offset], - )); + self.annotations.push(Instruction::member_decorate( + id, + index as u32, + spirv::Decoration::Offset, + &[current_offset], + )); current_offset += match member.span { Some(span) => span.get(), None => layout.size, @@ -723,11 +701,7 @@ impl Writer { if self.flags.contains(WriterFlags::DEBUG) { if let Some(ref name) = member.name { self.debugs - .push(super::instructions::instruction_member_name( - id, - index as u32, - name, - )); + .push(Instruction::member_name(id, index as u32, name)); } } @@ -741,26 +715,24 @@ impl Writer { crate::VectorSize::Bi => 2 * width, crate::VectorSize::Tri | crate::VectorSize::Quad => 4 * width, }; - self.annotations - .push(super::instructions::instruction_member_decorate( - id, - index as u32, - spirv::Decoration::ColMajor, - &[], - )); - self.annotations - .push(super::instructions::instruction_member_decorate( - id, - index as u32, - spirv::Decoration::MatrixStride, - &[byte_stride as u32], - )); + self.annotations.push(Instruction::member_decorate( + id, + index as u32, + spirv::Decoration::ColMajor, + &[], + )); + self.annotations.push(Instruction::member_decorate( + id, + index as u32, + spirv::Decoration::MatrixStride, + &[byte_stride as u32], + )); } let member_id = self.get_type_id(arena, LookupType::Handle(member.ty))?; member_ids.push(member_id); } - super::instructions::instruction_type_struct(id, member_ids.as_slice()) + Instruction::type_struct(id, member_ids.as_slice()) } crate::TypeInner::Struct { block: false, @@ -771,17 +743,13 @@ impl Writer { let member_id = self.get_type_id(arena, LookupType::Handle(member.ty))?; member_ids.push(member_id); } - super::instructions::instruction_type_struct(id, member_ids.as_slice()) + Instruction::type_struct(id, member_ids.as_slice()) } crate::TypeInner::Pointer { base, class } => { let type_id = self.get_type_id(arena, LookupType::Handle(base))?; self.lookup_type .insert(LookupType::Local(LocalType::Pointer { base, class }), id); - super::instructions::instruction_type_pointer( - id, - self.parse_to_spirv_storage_class(class), - type_id, - ) + Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id) } }; @@ -819,7 +787,7 @@ impl Writer { } _ => unreachable!(), }; - super::instructions::instruction_constant(type_id, id, words) + Instruction::constant(type_id, id, words) } crate::ScalarValue::Uint(val) => { let words = match width { @@ -833,7 +801,7 @@ impl Writer { } _ => unreachable!(), }; - super::instructions::instruction_constant(type_id, id, words) + Instruction::constant(type_id, id, words) } crate::ScalarValue::Float(val) => { let words = match width { @@ -848,14 +816,10 @@ impl Writer { } _ => unreachable!(), }; - super::instructions::instruction_constant(type_id, id, words) - } - crate::ScalarValue::Bool(true) => { - super::instructions::instruction_constant_true(type_id, id) - } - crate::ScalarValue::Bool(false) => { - super::instructions::instruction_constant_false(type_id, id) + Instruction::constant(type_id, id, words) } + crate::ScalarValue::Bool(true) => Instruction::constant_true(type_id, id), + crate::ScalarValue::Bool(false) => Instruction::constant_false(type_id, id), } } crate::ConstantInner::Composite { ty, ref components } => { @@ -866,11 +830,7 @@ impl Writer { } let type_id = self.get_type_id(types, LookupType::Handle(ty))?; - super::instructions::instruction_constant_composite( - type_id, - id, - constituent_ids.as_slice(), - ) + Instruction::constant_composite(type_id, id, constituent_ids.as_slice()) } }; @@ -894,13 +854,11 @@ impl Writer { .map(|constant| self.lookup_constant[&constant]); let pointer_type_id = self.get_pointer_id(&ir_module.types, global_variable.ty, global_variable.class)?; - let instruction = - super::instructions::instruction_variable(pointer_type_id, id, class, init_word); + let instruction = Instruction::variable(pointer_type_id, id, class, init_word); if self.flags.contains(WriterFlags::DEBUG) { if let Some(ref name) = global_variable.name { - self.debugs - .push(super::instructions::instruction_name(id, name)); + self.debugs.push(Instruction::name(id, name)); } } @@ -911,11 +869,7 @@ impl Writer { }; if let Some(decoration) = access_decoration { self.annotations - .push(super::instructions::instruction_decorate( - id, - decoration, - &[], - )); + .push(Instruction::decorate(id, decoration, &[])); } if let Some(interpolation) = global_variable.interpolation { @@ -929,36 +883,29 @@ impl Writer { }; if let Some(decoration) = decoration { self.annotations - .push(super::instructions::instruction_decorate( - id, - decoration, - &[], - )); + .push(Instruction::decorate(id, decoration, &[])); } } match global_variable.binding { Some(crate::Binding::Location(location)) => { - self.annotations - .push(super::instructions::instruction_decorate( - id, - spirv::Decoration::Location, - &[location], - )); + self.annotations.push(Instruction::decorate( + id, + spirv::Decoration::Location, + &[location], + )); } Some(crate::Binding::Resource { group, binding }) => { - self.annotations - .push(super::instructions::instruction_decorate( - id, - spirv::Decoration::DescriptorSet, - &[group], - )); - self.annotations - .push(super::instructions::instruction_decorate( - id, - spirv::Decoration::Binding, - &[binding], - )); + self.annotations.push(Instruction::decorate( + id, + spirv::Decoration::DescriptorSet, + &[group], + )); + self.annotations.push(Instruction::decorate( + id, + spirv::Decoration::Binding, + &[binding], + )); } Some(crate::Binding::BuiltIn(built_in)) => { use crate::BuiltIn as Bi; @@ -985,12 +932,11 @@ impl Writer { Bi::WorkGroupSize => spirv::BuiltIn::WorkgroupSize, }; - self.annotations - .push(super::instructions::instruction_decorate( - id, - spirv::Decoration::BuiltIn, - &[built_in as u32], - )); + self.annotations.push(Instruction::decorate( + id, + spirv::Decoration::BuiltIn, + &[built_in as u32], + )); } None => {} } @@ -1009,7 +955,7 @@ impl Writer { Entry::Occupied(e) => *e.get(), _ => { let id = self.generate_id(); - let instruction = super::instructions::instruction_type_function( + let instruction = Instruction::type_function( id, lookup_function_type.return_type_id, &lookup_function_type.parameter_type_ids, @@ -1028,13 +974,11 @@ impl Writer { block: &mut Block, ) -> Word { let id = self.generate_id(); - block - .body - .push(super::instructions::instruction_composite_construct( - base_type_id, - id, - constituent_ids, - )); + block.body.push(Instruction::composite_construct( + base_type_id, + id, + constituent_ids, + )); id } @@ -1053,12 +997,9 @@ impl Writer { RawExpression::Value(id) => id, RawExpression::Pointer(id, _) => { let load_id = self.generate_id(); - block.body.push(super::instructions::instruction_load( - result_type_id, - load_id, - id, - None, - )); + block + .body + .push(Instruction::load(result_type_id, load_id, id, None)); load_id } }) @@ -1116,28 +1057,24 @@ impl Writer { )); } - block - .body - .push(super::instructions::instruction_vector_extract_dynamic( - result_type_id, - id, - base_id, - index_id, - )); + block.body.push(Instruction::vector_extract_dynamic( + result_type_id, + id, + base_id, + index_id, + )); RawExpression::Value(id) } RawExpression::Pointer(base_id, class) => { let pointer_type_id = self.create_pointer_type(result_type_id, class); - block - .body - .push(super::instructions::instruction_access_chain( - pointer_type_id, - id, - base_id, - &[index_id], - )); + block.body.push(Instruction::access_chain( + pointer_type_id, + id, + base_id, + &[index_id], + )); RawExpression::Pointer(id, class) } @@ -1150,14 +1087,12 @@ impl Writer { match raw_base_expression { RawExpression::Value(base_id) => { - block - .body - .push(super::instructions::instruction_composite_extract( - result_type_id, - id, - base_id, - &[index], - )); + block.body.push(Instruction::composite_extract( + result_type_id, + id, + base_id, + &[index], + )); RawExpression::Value(id) } @@ -1173,14 +1108,12 @@ impl Writer { )?; let const_id = self.create_constant(const_ty_id, &[index]); - block - .body - .push(super::instructions::instruction_access_chain( - pointer_type_id, - id, - base_id, - &[const_id], - )); + block.body.push(Instruction::access_chain( + pointer_type_id, + id, + base_id, + &[const_id], + )); RawExpression::Pointer(id, class) } @@ -1231,12 +1164,9 @@ impl Writer { crate::UnaryOperator::Not => spirv::Op::Not, }; - block.body.push(super::instructions::instruction_unary( - spirv_op, - result_type_id, - id, - expr_id, - )); + block + .body + .push(Instruction::unary(spirv_op, result_type_id, id, expr_id)); RawExpression::Value(id) } crate::Expression::Binary { op, left, right } => { @@ -1360,7 +1290,7 @@ impl Writer { }, }; - block.body.push(super::instructions::instruction_binary( + block.body.push(Instruction::binary( spirv_op, result_type_id, id, @@ -1405,7 +1335,7 @@ impl Writer { Some(crate::ScalarKind::Float) => MathOp::Ext(spirv::GLOp::FAbs), Some(crate::ScalarKind::Sint) => MathOp::Ext(spirv::GLOp::SAbs), Some(crate::ScalarKind::Uint) => { - MathOp::Custom(super::instructions::instruction_unary( + MathOp::Custom(Instruction::unary( spirv::Op::CopyObject, // do nothing result_type_id, id, @@ -1454,7 +1384,7 @@ impl Writer { Mf::Frexp => MathOp::Ext(spirv::GLOp::Frexp), Mf::Ldexp => MathOp::Ext(spirv::GLOp::Ldexp), // geometry - Mf::Dot => MathOp::Custom(super::instructions::instruction_binary( + Mf::Dot => MathOp::Custom(Instruction::binary( spirv::Op::Dot, result_type_id, id, @@ -1485,7 +1415,7 @@ impl Writer { Mf::SmoothStep => MathOp::Ext(spirv::GLOp::SmoothStep), Mf::Sqrt => MathOp::Ext(spirv::GLOp::Sqrt), Mf::InverseSqrt => MathOp::Ext(spirv::GLOp::InverseSqrt), - Mf::Transpose => MathOp::Custom(super::instructions::instruction_unary( + Mf::Transpose => MathOp::Custom(Instruction::unary( spirv::Op::Transpose, result_type_id, id, @@ -1499,7 +1429,7 @@ impl Writer { }; block.body.push(match math_op { - MathOp::Ext(op) => super::instructions::instruction_ext_inst( + MathOp::Ext(op) => Instruction::ext_inst( self.gl450_ext_inst_id, op, result_type_id, @@ -1532,14 +1462,12 @@ impl Writer { argument_ids.push(arg_id); } - block - .body - .push(super::instructions::instruction_function_call( - result_type_id, - id, - *self.lookup_function.get(&local_function).unwrap(), - argument_ids.as_slice(), - )); + block.body.push(Instruction::function_call( + result_type_id, + id, + *self.lookup_function.get(&local_function).unwrap(), + argument_ids.as_slice(), + )); RawExpression::Value(id) } @@ -1571,8 +1499,7 @@ impl Writer { }; let id = self.generate_id(); - let instruction = - super::instructions::instruction_unary(op, result_type_id, id, expr_id); + let instruction = Instruction::unary(op, result_type_id, id, expr_id); block.body.push(instruction); RawExpression::Value(id) @@ -1625,14 +1552,12 @@ impl Writer { for i in 0..size as u32 { let id = self.generate_id(); constituent_ids[i as usize] = id; - block.body.push( - super::instructions::instruction_composite_extract( - coordinate_scalar_type_id, - id, - coordinate_id, - &[i], - ), - ); + block.body.push(Instruction::composite_extract( + coordinate_scalar_type_id, + id, + coordinate_id, + &[i], + )); } match size { crate::VectorSize::Bi => crate::VectorSize::Tri, @@ -1655,7 +1580,7 @@ impl Writer { block, function, )?; - let cast_instruction = super::instructions::instruction_unary( + let cast_instruction = Instruction::unary( spirv::Op::ConvertUToF, coordinate_scalar_type_id, array_index_f32_id, @@ -1680,14 +1605,12 @@ impl Writer { } let sampled_image_id = self.generate_id(); - block - .body - .push(super::instructions::instruction_sampled_image( - sampled_image_type_id, - sampled_image_id, - image_id, - sampler_id, - )); + block.body.push(Instruction::sampled_image( + sampled_image_type_id, + sampled_image_id, + image_id, + sampler_id, + )); let id = self.generate_id(); let depth_id = match depth_ref { @@ -1701,7 +1624,7 @@ impl Writer { let mut main_instruction = match level { crate::SampleLevel::Zero => { - let mut inst = super::instructions::instruction_image_sample( + let mut inst = Instruction::image_sample( result_type_id, id, SampleLod::Explicit, @@ -1722,7 +1645,7 @@ impl Writer { inst } - crate::SampleLevel::Auto => super::instructions::instruction_image_sample( + crate::SampleLevel::Auto => Instruction::image_sample( result_type_id, id, SampleLod::Implicit, @@ -1731,7 +1654,7 @@ impl Writer { depth_id, ), crate::SampleLevel::Exact(lod_handle) => { - let mut inst = super::instructions::instruction_image_sample( + let mut inst = Instruction::image_sample( result_type_id, id, SampleLod::Explicit, @@ -1753,7 +1676,7 @@ impl Writer { inst } crate::SampleLevel::Bias(bias_handle) => { - let mut inst = super::instructions::instruction_image_sample( + let mut inst = Instruction::image_sample( result_type_id, id, SampleLod::Implicit, @@ -1775,7 +1698,7 @@ impl Writer { inst } crate::SampleLevel::Gradient { x, y } => { - let mut inst = super::instructions::instruction_image_sample( + let mut inst = Instruction::image_sample( result_type_id, id, SampleLod::Explicit, @@ -1834,7 +1757,7 @@ impl Writer { match *statement { crate::Statement::Block(ref block_statements) => { let scope_id = self.generate_id(); - function.consume(block, super::instructions::instruction_branch(scope_id)); + function.consume(block, Instruction::branch(scope_id)); let merge_id = self.generate_id(); self.write_block( @@ -1863,22 +1786,16 @@ impl Writer { )?; let merge_id = self.generate_id(); - block - .body - .push(super::instructions::instruction_selection_merge( - merge_id, - spirv::SelectionControl::NONE, - )); + block.body.push(Instruction::selection_merge( + merge_id, + spirv::SelectionControl::NONE, + )); let accept_id = self.generate_id(); let reject_id = self.generate_id(); function.consume( block, - super::instructions::instruction_branch_conditional( - condition_id, - accept_id, - reject_id, - ), + Instruction::branch_conditional(condition_id, accept_id, reject_id), ); self.write_block( @@ -1916,12 +1833,10 @@ impl Writer { )?; let merge_id = self.generate_id(); - block - .body - .push(super::instructions::instruction_selection_merge( - merge_id, - spirv::SelectionControl::NONE, - )); + block.body.push(Instruction::selection_merge( + merge_id, + spirv::SelectionControl::NONE, + )); let default_id = self.generate_id(); let raw_cases = cases @@ -1934,11 +1849,7 @@ impl Writer { function.consume( block, - super::instructions::instruction_switch( - selector_id, - default_id, - &raw_cases, - ), + Instruction::switch(selector_id, default_id, &raw_cases), ); for (i, (case, raw_case)) in cases.iter().zip(raw_cases.iter()).enumerate() { @@ -1978,7 +1889,7 @@ impl Writer { ref continuing, } => { let preamble_id = self.generate_id(); - function.consume(block, super::instructions::instruction_branch(preamble_id)); + function.consume(block, Instruction::branch(preamble_id)); let merge_id = self.generate_id(); let body_id = self.generate_id(); @@ -1987,12 +1898,12 @@ impl Writer { // SPIR-V requires the continuing to the `OpLoopMerge`, // so we have to start a new block with it. block = Block::new(preamble_id); - block.body.push(super::instructions::instruction_loop_merge( + block.body.push(Instruction::loop_merge( merge_id, continuing_id, spirv::SelectionControl::NONE, )); - function.consume(block, super::instructions::instruction_branch(body_id)); + function.consume(block, Instruction::branch(body_id)); self.write_block( body_id, @@ -2023,25 +1934,22 @@ impl Writer { block = Block::new(merge_id); } crate::Statement::Break => { - block.termination = Some(super::instructions::instruction_branch( - loop_context.break_id.unwrap(), - )); + block.termination = Some(Instruction::branch(loop_context.break_id.unwrap())); } crate::Statement::Continue => { - block.termination = Some(super::instructions::instruction_branch( - loop_context.continuing_id.unwrap(), - )); + block.termination = + Some(Instruction::branch(loop_context.continuing_id.unwrap())); } crate::Statement::Return { value: Some(value) } => { let id = self.write_expression(ir_module, ir_function, value, &mut block, function)?; - block.termination = Some(super::instructions::instruction_return_value(id)); + block.termination = Some(Instruction::return_value(id)); } crate::Statement::Return { value: None } => { - block.termination = Some(super::instructions::instruction_return()); + block.termination = Some(Instruction::return_void()); } crate::Statement::Kill => { - block.termination = Some(super::instructions::instruction_kill()); + block.termination = Some(Instruction::kill()); } crate::Statement::Store { pointer, value } => { let (pointer_id, _) = self.write_expression_pointer( @@ -2054,9 +1962,9 @@ impl Writer { let value_id = self.write_expression(ir_module, ir_function, value, &mut block, function)?; - block.body.push(super::instructions::instruction_store( - pointer_id, value_id, None, - )); + block + .body + .push(Instruction::store(pointer_id, value_id, None)); } crate::Statement::Call { function: local_function, @@ -2077,22 +1985,20 @@ impl Writer { argument_ids.push(arg_id); } - block - .body - .push(super::instructions::instruction_function_call( - self.void_type, - id, - *self.lookup_function.get(&local_function).unwrap(), - argument_ids.as_slice(), - )); + block.body.push(Instruction::function_call( + self.void_type, + id, + *self.lookup_function.get(&local_function).unwrap(), + argument_ids.as_slice(), + )); } } } if block.termination.is_none() { block.termination = Some(match exit_id { - Some(id) => super::instructions::instruction_branch(id), - None => super::instructions::instruction_return(), + Some(id) => Instruction::branch(id), + None => Instruction::return_void(), }); } @@ -2105,16 +2011,13 @@ impl Writer { } fn write_logical_layout(&mut self, ir_module: &crate::Module) -> Result<(), Error> { - super::instructions::instruction_type_void(self.void_type) - .to_words(&mut self.logical_layout.declarations); - super::instructions::instruction_ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450") + Instruction::type_void(self.void_type).to_words(&mut self.logical_layout.declarations); + Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450") .to_words(&mut self.logical_layout.ext_inst_imports); if self.flags.contains(WriterFlags::DEBUG) { - self.debugs.push(super::instructions::instruction_source( - spirv::SourceLanguage::GLSL, - 450, - )); + self.debugs + .push(Instruction::source(spirv::SourceLanguage::GLSL, 450)); } for (handle, constant) in ir_module.constants.iter() { @@ -2140,8 +2043,7 @@ impl Writer { } for capability in self.capabilities.iter() { - super::instructions::instruction_capability(*capability) - .to_words(&mut self.logical_layout.capabilities); + Instruction::capability(*capability).to_words(&mut self.logical_layout.capabilities); } let addressing_model = spirv::AddressingModel::Logical; @@ -2149,7 +2051,7 @@ impl Writer { self.check(addressing_model.required_capabilities())?; self.check(memory_model.required_capabilities())?; - super::instructions::instruction_memory_model(addressing_model, memory_model) + Instruction::memory_model(addressing_model, memory_model) .to_words(&mut self.logical_layout.memory_model); if self.flags.contains(WriterFlags::DEBUG) {