From 5def021b6ee3d5660f196825d8fcffc940c18974 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Mon, 1 Feb 2021 09:33:28 -0500 Subject: [PATCH] [spv] replace the internal type inferrence with the typifier (#403) Also, fix the Typifier's handling of logical operations. See https://github.com/gfx-rs/naga/pull/404 for v0.3 branch. --- src/back/spv/writer.rs | 865 ++++++------------ src/proc/typifier.rs | 24 +- tests/snapshots/snapshots__boids.msl.snap | 2 + tests/snapshots/snapshots__boids.spvasm.snap | 298 +++--- .../snapshots/snapshots__collatz.spvasm.snap | 60 +- tests/snapshots/snapshots__quad.spvasm.snap | 62 +- tests/snapshots/snapshots__shadow.msl.snap | 2 + tests/snapshots/snapshots__shadow.spvasm.snap | 407 ++++---- tests/snapshots/snapshots__skybox.spvasm.snap | 388 ++++---- 9 files changed, 901 insertions(+), 1207 deletions(-) diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 3feb3727d4..b5cfe69d70 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -1,8 +1,8 @@ /*! Standard Portable Intermediate Representation (SPIR-V) backend !*/ use super::{Instruction, LogicalLayout, PhysicalLayout, WriterFlags}; -use crate::proc::Layouter; +use crate::proc::{Layouter, ResolveContext, ResolveError, Typifier}; use spirv::Word; -use std::{collections::hash_map::Entry, ops}; +use std::collections::hash_map::Entry; use thiserror::Error; const BITS_PER_BYTE: crate::Bytes = 8; @@ -13,6 +13,8 @@ pub enum Error { MissingCapabilities(Vec), #[error("unimplemented {0:}")] FeatureNotImplemented(&'static str), + #[error(transparent)] + Resolve(#[from] ResolveError), } struct Block { @@ -77,7 +79,6 @@ impl Function { #[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)] enum LocalType { - Void, Scalar { kind: crate::ScalarKind, width: crate::Bytes, @@ -101,6 +102,28 @@ enum LocalType { }, } +impl LocalType { + fn from_inner(inner: &crate::TypeInner) -> Option { + Some(match *inner { + crate::TypeInner::Scalar { kind, width } => LocalType::Scalar { kind, width }, + crate::TypeInner::Vector { size, kind, width } => { + LocalType::Vector { size, kind, width } + } + crate::TypeInner::Matrix { + columns, + rows, + width, + } => LocalType::Matrix { + columns, + rows, + width, + }, + crate::TypeInner::Pointer { base, class } => LocalType::Pointer { base, class }, + _ => return None, + }) + } +} + #[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)] enum LookupType { Handle(crate::Handle), @@ -128,21 +151,6 @@ struct LookupFunctionType { return_type_id: Word, } -enum MaybeOwned<'a, T> { - Owned(T), - Borrowed(&'a T), -} - -impl<'a, T> ops::Deref for MaybeOwned<'a, T> { - type Target = T; - fn deref(&self) -> &T { - match *self { - MaybeOwned::Owned(ref value) => value, - MaybeOwned::Borrowed(reference) => reference, - } - } -} - #[derive(Debug)] enum Dimension { Scalar, @@ -150,8 +158,8 @@ enum Dimension { Matrix, } -fn get_dimension(ty_inner: &crate::TypeInner) -> Dimension { - match *ty_inner { +fn get_dimension(type_inner: &crate::TypeInner) -> Dimension { + match *type_inner { crate::TypeInner::Scalar { .. } => Dimension::Scalar, crate::TypeInner::Vector { .. } => Dimension::Vector, crate::TypeInner::Matrix { .. } => Dimension::Matrix, @@ -185,11 +193,12 @@ pub struct Writer { struct_type_handles: crate::FastHashMap, crate::StorageAccess>, gl450_ext_inst_id: Word, layouter: Layouter, + typifier: Typifier, } // type alias, for success return of write_expression -type WriteExpressionOutput = (Word, LookupType); -type WritePointerExpressionOutput = (Word, LookupType, spirv::StorageClass); +type ExpressionId = Word; +type PointerExpressionId = (Word, spirv::StorageClass); impl Writer { pub fn new( @@ -214,6 +223,7 @@ impl Writer { struct_type_handles: crate::FastHashMap::default(), gl450_ext_inst_id: 0, layouter: Layouter::default(), + typifier: Typifier::new(), } } @@ -250,6 +260,18 @@ impl Writer { arena, LookupType::Local(LocalType::Vector { size, kind, width }), ), + crate::TypeInner::Matrix { + columns, + rows, + width, + } => self.get_type_id( + arena, + LookupType::Local(LocalType::Matrix { + columns, + rows, + width, + }), + ), _ => self.write_type_declaration_arena(arena, handle), }, LookupType::Local(local_ty) => self.write_type_declaration_local(arena, local_ty), @@ -332,15 +354,13 @@ impl Writer { fn create_pointer_type( &mut self, - lookup_type: LookupType, + type_id: spirv::Word, class: spirv::StorageClass, - type_arena: &crate::Arena, - ) -> Result<(Word, LookupType), Error> { - let type_id = self.get_type_id(type_arena, lookup_type)?; + ) -> Result { let id = self.generate_id(); let instruction = super::instructions::instruction_type_pointer(id, class, type_id); instruction.to_words(&mut self.logical_layout.declarations); - Ok((id, lookup_type)) + Ok(id) } fn create_constant(&mut self, type_id: Word, value: &[Word]) -> Word { @@ -357,6 +377,18 @@ impl Writer { ) -> Result { let mut function = Function::default(); + self.typifier.resolve_all( + &ir_function.expressions, + &ir_module.types, + &ResolveContext { + constants: &ir_module.constants, + global_vars: &ir_module.global_variables, + local_vars: &ir_function.local_variables, + functions: &ir_module.functions, + arguments: &ir_function.arguments, + }, + )?; + for (handle, variable) in ir_function.local_variables.iter() { let id = self.generate_id(); @@ -537,7 +569,6 @@ impl Writer { ) -> Result { let id = self.generate_id(); let instruction = match local_ty { - LocalType::Void => unreachable!(), LocalType::Scalar { kind, width } => self.write_scalar(id, kind, width), LocalType::Vector { size, kind, width } => { let scalar_id = @@ -1048,37 +1079,6 @@ impl Writer { id } - fn get_type_inner<'a>( - &self, - ty_arena: &'a crate::Arena, - lookup_ty: LookupType, - ) -> MaybeOwned<'a, crate::TypeInner> { - match lookup_ty { - LookupType::Handle(handle) => MaybeOwned::Borrowed(&ty_arena[handle].inner), - LookupType::Local(local_ty) => match local_ty { - LocalType::Scalar { kind, width } => { - MaybeOwned::Owned(crate::TypeInner::Scalar { kind, width }) - } - LocalType::Vector { size, kind, width } => { - MaybeOwned::Owned(crate::TypeInner::Vector { size, kind, width }) - } - LocalType::Matrix { - columns, - rows, - width, - } => MaybeOwned::Owned(crate::TypeInner::Matrix { - columns, - rows, - width, - }), - LocalType::Pointer { base, class } => { - MaybeOwned::Owned(crate::TypeInner::Pointer { base, class }) - } - LocalType::Void | LocalType::SampledImage { .. } => unreachable!(), - }, - } - } - /// Write an expression and return a value ID. fn write_expression<'a>( &mut self, @@ -1087,18 +1087,20 @@ impl Writer { handle: crate::Handle, block: &mut Block, function: &mut Function, - ) -> Result { - let (raw_expression, lookup_ty) = + ) -> Result { + let (raw_expression, result_type_id) = self.write_expression_raw(ir_module, ir_function, handle, block, function)?; Ok(match raw_expression { - RawExpression::Value(id) => (id, lookup_ty), + RawExpression::Value(id) => id, RawExpression::Pointer(id, _) => { let load_id = self.generate_id(); - let type_id = self.get_type_id(&ir_module.types, lookup_ty)?; block.body.push(super::instructions::instruction_load( - type_id, load_id, id, None, + result_type_id, + load_id, + id, + None, )); - (load_id, lookup_ty) + load_id } }) } @@ -1111,15 +1113,15 @@ impl Writer { handle: crate::Handle, block: &mut Block, function: &mut Function, - ) -> Result { - let (raw_expression, lookup_ty) = + ) -> Result { + let (raw_expression, _) = self.write_expression_raw(ir_module, ir_function, handle, block, function)?; Ok(match raw_expression { RawExpression::Value(_id) => unimplemented!( "Expression {:?} is not a pointer", ir_function.expressions[handle] ), - RawExpression::Pointer(id, class) => (id, lookup_ty, class), + RawExpression::Pointer(id, class) => (id, class), }) } @@ -1131,40 +1133,30 @@ impl Writer { expr_handle: crate::Handle, block: &mut Block, function: &mut Function, - ) -> Result<(RawExpression, LookupType), Error> { - match ir_function.expressions[expr_handle] { + ) -> Result<(RawExpression, ExpressionId), Error> { + let result_lookup_ty = match self.typifier.get_handle(expr_handle) { + Ok(ty_handle) => LookupType::Handle(ty_handle), + Err(inner) => LookupType::Local(LocalType::from_inner(inner).unwrap()), + }; + let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?; + + let raw_expr = match ir_function.expressions[expr_handle] { crate::Expression::Access { base, index } => { let id = self.generate_id(); - let (raw_base_expression, base_lookup_ty) = + let (raw_base_expression, _) = self.write_expression_raw(ir_module, ir_function, base, block, function)?; - let base_ty_inner = self.get_type_inner(&ir_module.types, base_lookup_ty); - let (index_id, _) = + let index_id = self.write_expression(ir_module, ir_function, index, block, function)?; + let base_inner = self.typifier.get(base, &ir_module.types); - let lookup_ty = match *base_ty_inner { - crate::TypeInner::Vector { - size: _, - kind, - width, - } => LookupType::Local(LocalType::Scalar { kind, width }), - crate::TypeInner::Array { base, .. } => LookupType::Handle(base), - ref other => { - log::error!("Unable to index {:?}", other); - return Err(Error::FeatureNotImplemented( - "accessing index of non vector or array", - )); - } - }; - - Ok(match raw_base_expression { + match raw_base_expression { RawExpression::Value(base_id) => { - if let crate::TypeInner::Array { .. } = *base_ty_inner { + if let crate::TypeInner::Array { .. } = *base_inner { return Err(Error::FeatureNotImplemented( "accessing index of a value array", )); } - let result_type_id = self.get_type_id(&ir_module.types, lookup_ty)?; block .body .push(super::instructions::instruction_vector_extract_dynamic( @@ -1174,11 +1166,10 @@ impl Writer { index_id, )); - (RawExpression::Value(id), lookup_ty) + RawExpression::Value(id) } RawExpression::Pointer(base_id, class) => { - let (pointer_type_id, pointer_lookup_ty) = - self.create_pointer_type(lookup_ty, class, &ir_module.types)?; + let pointer_type_id = self.create_pointer_type(result_type_id, class)?; block .body @@ -1189,46 +1180,17 @@ impl Writer { &[index_id], )); - (RawExpression::Pointer(id, class), pointer_lookup_ty) + RawExpression::Pointer(id, class) } - }) + } } crate::Expression::AccessIndex { base, index } => { let id = self.generate_id(); - let (raw_base_expression, base_lookup_ty) = + let (raw_base_expression, _) = self.write_expression_raw(ir_module, ir_function, base, block, function)?; - let base_ty_inner = self.get_type_inner(&ir_module.types, base_lookup_ty); - let lookup_ty = match *base_ty_inner { - crate::TypeInner::Vector { - size: _, - kind, - width, - } => LookupType::Local(LocalType::Scalar { kind, width }), - crate::TypeInner::Matrix { - columns: _, - rows, - width, - } => LookupType::Local(LocalType::Vector { - size: rows, - kind: crate::ScalarKind::Float, - width, - }), - crate::TypeInner::Struct { - block: _, - ref members, - } => LookupType::Handle(members[index as usize].ty), - ref other => { - log::error!("Unable to access index {:?}", other); - return Err(Error::FeatureNotImplemented( - "accessing index of non vector or struct", - )); - } - }; - - Ok(match raw_base_expression { + match raw_base_expression { RawExpression::Value(base_id) => { - let result_type_id = self.get_type_id(&ir_module.types, lookup_ty)?; block .body .push(super::instructions::instruction_composite_extract( @@ -1238,9 +1200,11 @@ impl Writer { &[index], )); - (RawExpression::Value(id), lookup_ty) + RawExpression::Value(id) } RawExpression::Pointer(base_id, class) => { + let pointer_type_id = self.create_pointer_type(result_type_id, class)?; + let const_ty_id = self.get_type_id( &ir_module.types, LookupType::Local(LocalType::Scalar { @@ -1249,8 +1213,6 @@ impl Writer { }), )?; let const_id = self.create_constant(const_ty_id, &[index]); - let (pointer_type_id, pointer_lookup_ty) = - self.create_pointer_type(lookup_ty, class, &ir_module.types)?; block .body @@ -1261,39 +1223,24 @@ impl Writer { &[const_id], )); - (RawExpression::Pointer(id, class), pointer_lookup_ty) + RawExpression::Pointer(id, class) } - }) + } } crate::Expression::GlobalVariable(handle) => { - let var = &ir_module.global_variables[handle]; let (id, class) = self.get_global_variable_id(&ir_module, handle)?; - - Ok(( - RawExpression::Pointer(id, class), - LookupType::Handle(var.ty), - )) + RawExpression::Pointer(id, class) } crate::Expression::Constant(handle) => { - let var = &ir_module.constants[handle]; let id = self.lookup_constant[&handle]; - let lookup_type = match var.inner { - crate::ConstantInner::Scalar { width, ref value } => { - LookupType::Local(LocalType::Scalar { - kind: value.scalar_kind(), - width, - }) - } - crate::ConstantInner::Composite { ty, components: _ } => LookupType::Handle(ty), - }; - Ok((RawExpression::Value(id), lookup_type)) + RawExpression::Value(id) } crate::Expression::Compose { ty, ref components } => { let base_type_id = self.get_type_id(&ir_module.types, LookupType::Handle(ty))?; let mut constituent_ids = Vec::with_capacity(components.len()); for component in components { - let (component_id, _) = self.write_expression( + let component_id = self.write_expression( ir_module, &ir_function, *component, @@ -1304,14 +1251,13 @@ impl Writer { } let id = self.write_composite_construct(base_type_id, &constituent_ids, block); - Ok((RawExpression::Value(id), LookupType::Handle(ty))) + RawExpression::Value(id) } crate::Expression::Unary { op, expr } => { let id = self.generate_id(); - let (expr_id, expr_lookup_ty) = + let expr_id = self.write_expression(ir_module, ir_function, expr, block, function)?; - let expr_ty_inner = self.get_type_inner(&ir_module.types, expr_lookup_ty); - let result_type_id = self.get_type_id(&ir_module.types, expr_lookup_ty)?; + let expr_ty_inner = self.typifier.get(expr, &ir_module.types); let spirv_op = match op { crate::UnaryOperator::Negate => match expr_ty_inner.scalar_kind() { @@ -1319,7 +1265,7 @@ impl Writer { Some(crate::ScalarKind::Sint) => spirv::Op::SNegate, Some(crate::ScalarKind::Bool) => spirv::Op::LogicalNot, Some(crate::ScalarKind::Uint) | None => { - log::error!("Unable to negate {:?}", &*expr_ty_inner); + log::error!("Unable to negate {:?}", expr_ty_inner); return Err(Error::FeatureNotImplemented("negation")); } }, @@ -1332,25 +1278,21 @@ impl Writer { id, expr_id, )); - Ok((RawExpression::Value(id), expr_lookup_ty)) + RawExpression::Value(id) } crate::Expression::Binary { op, left, right } => { let id = self.generate_id(); - let (left_id, left_lookup_ty) = + let left_id = self.write_expression(ir_module, ir_function, left, block, function)?; - let (right_id, right_lookup_ty) = + let right_id = self.write_expression(ir_module, ir_function, right, block, function)?; - let left_ty_inner = self.get_type_inner(&ir_module.types, left_lookup_ty); - let right_ty_inner = self.get_type_inner(&ir_module.types, right_lookup_ty); + let left_ty_inner = self.typifier.get(left, &ir_module.types); + let right_ty_inner = self.typifier.get(right, &ir_module.types); - let left_result_type_id = self.get_type_id(&ir_module.types, left_lookup_ty)?; - let right_result_type_id = self.get_type_id(&ir_module.types, right_lookup_ty)?; + let left_dimension = get_dimension(left_ty_inner); + let right_dimension = get_dimension(right_ty_inner); - let left_dimension = get_dimension(&left_ty_inner); - let right_dimension = get_dimension(&right_ty_inner); - - let mut result_side_left = true; let mut preserve_order = true; let spirv_op = match op { @@ -1370,38 +1312,30 @@ impl Writer { }, _ => unimplemented!(), }, - crate::BinaryOperator::Multiply => { - // whenever there is a vector on the right, - // the result type is a vector. - if let Dimension::Vector = right_dimension { - result_side_left = false; + crate::BinaryOperator::Multiply => match (left_dimension, right_dimension) { + (Dimension::Scalar, Dimension::Vector { .. }) => { + preserve_order = false; + spirv::Op::VectorTimesScalar } - match (left_dimension, right_dimension) { - (Dimension::Scalar, Dimension::Vector { .. }) => { - preserve_order = false; - spirv::Op::VectorTimesScalar - } - (Dimension::Vector, Dimension::Scalar { .. }) => { - spirv::Op::VectorTimesScalar - } - (Dimension::Vector, Dimension::Matrix) => spirv::Op::VectorTimesMatrix, - (Dimension::Matrix, Dimension::Scalar { .. }) => { - spirv::Op::MatrixTimesScalar - } - (Dimension::Matrix, Dimension::Vector) => spirv::Op::MatrixTimesVector, - (Dimension::Matrix, Dimension::Matrix) => spirv::Op::MatrixTimesMatrix, - (Dimension::Vector, Dimension::Vector) - | (Dimension::Scalar, Dimension::Scalar) - if left_ty_inner.scalar_kind() - == Some(crate::ScalarKind::Float) => - { - spirv::Op::FMul - } - (Dimension::Vector, Dimension::Vector) - | (Dimension::Scalar, Dimension::Scalar) => spirv::Op::IMul, - other => unimplemented!("Mul {:?}", other), + (Dimension::Vector, Dimension::Scalar { .. }) => { + spirv::Op::VectorTimesScalar } - } + (Dimension::Vector, Dimension::Matrix) => spirv::Op::VectorTimesMatrix, + (Dimension::Matrix, Dimension::Scalar { .. }) => { + spirv::Op::MatrixTimesScalar + } + (Dimension::Matrix, Dimension::Vector) => spirv::Op::MatrixTimesVector, + (Dimension::Matrix, Dimension::Matrix) => spirv::Op::MatrixTimesMatrix, + (Dimension::Vector, Dimension::Vector) + | (Dimension::Scalar, Dimension::Scalar) + if left_ty_inner.scalar_kind() == Some(crate::ScalarKind::Float) => + { + spirv::Op::FMul + } + (Dimension::Vector, Dimension::Vector) + | (Dimension::Scalar, Dimension::Scalar) => spirv::Op::IMul, + other => unimplemented!("Mul {:?}", other), + }, crate::BinaryOperator::Divide => match left_ty_inner.scalar_kind() { Some(crate::ScalarKind::Sint) => spirv::Op::SDiv, Some(crate::ScalarKind::Uint) => spirv::Op::UDiv, @@ -1467,29 +1401,6 @@ impl Writer { }, }; - let is_comparison = match op { - crate::BinaryOperator::Equal - | crate::BinaryOperator::NotEqual - | crate::BinaryOperator::Less - | crate::BinaryOperator::LessEqual - | crate::BinaryOperator::Greater - | crate::BinaryOperator::GreaterEqual => true, - _ => false, - }; - - let (result_type_id, result_lookup_ty) = if is_comparison { - let local_ty = LookupType::Local(LocalType::Scalar { - kind: crate::ScalarKind::Bool, - width: 1, - }); - let result_ty_id = self.get_type_id(&ir_module.types, local_ty)?; - (result_ty_id, local_ty) - } else if result_side_left { - (left_result_type_id, left_lookup_ty) - } else { - (right_result_type_id, right_lookup_ty) - }; - block.body.push(super::instructions::instruction_binary( spirv_op, result_type_id, @@ -1497,7 +1408,7 @@ impl Writer { if preserve_order { left_id } else { right_id }, if preserve_order { right_id } else { left_id }, )); - Ok((RawExpression::Value(id), result_lookup_ty)) + RawExpression::Value(id) } crate::Expression::Math { fun, @@ -1507,28 +1418,22 @@ impl Writer { } => { use crate::MathFunction as Mf; enum MathOp { - Single(spirv::GLOp), - Double(spirv::GLOp), - Triple(spirv::GLOp), - Other(super::Instruction, LookupType), + Ext(spirv::GLOp), + Custom(super::Instruction), } - let (arg0_id, arg0_lookup_ty) = + let arg0_id = self.write_expression(ir_module, ir_function, arg, block, function)?; - let arg0_type_id = self.get_type_id(&ir_module.types, arg0_lookup_ty)?; + let arg_scalar_kind = self.typifier.get(arg, &ir_module.types).scalar_kind(); let arg1_id = match arg1 { Some(id) => { - let (arg1_id, _) = - self.write_expression(ir_module, ir_function, id, block, function)?; - arg1_id + self.write_expression(ir_module, ir_function, id, block, function)? } None => 0, }; let arg2_id = match arg2 { Some(id) => { - let (arg2_id, _) = - self.write_expression(ir_module, ir_function, id, block, function)?; - arg2_id + self.write_expression(ir_module, ir_function, id, block, function)? } None => 0, }; @@ -1537,339 +1442,147 @@ impl Writer { let math_op = match fun { // comparison Mf::Abs => { - let inst = match self - .get_type_inner(&ir_module.types, arg0_lookup_ty) - .scalar_kind() - { - Some(crate::ScalarKind::Float) => { - super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - spirv::GLOp::FAbs, - arg0_type_id, - id, - &[arg0_id], - ) - } - Some(crate::ScalarKind::Sint) => { - super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - spirv::GLOp::SAbs, - arg0_type_id, - id, - &[arg0_id], - ) - } + match arg_scalar_kind { + Some(crate::ScalarKind::Float) => MathOp::Ext(spirv::GLOp::FAbs), + Some(crate::ScalarKind::Sint) => MathOp::Ext(spirv::GLOp::SAbs), Some(crate::ScalarKind::Uint) => { - super::instructions::instruction_unary( + MathOp::Custom(super::instructions::instruction_unary( spirv::Op::CopyObject, // do nothing - arg0_type_id, + result_type_id, id, arg0_id, - ) + )) } other => unimplemented!("Unexpected abs({:?})", other), - }; - MathOp::Other(inst, arg0_lookup_ty) - } - Mf::Min => { - let op = match self - .get_type_inner(&ir_module.types, arg0_lookup_ty) - .scalar_kind() - { - Some(crate::ScalarKind::Float) => spirv::GLOp::FMin, - Some(crate::ScalarKind::Sint) => spirv::GLOp::SMin, - Some(crate::ScalarKind::Uint) => spirv::GLOp::UMin, - other => unimplemented!("Unexpected min({:?})", other), - }; - MathOp::Double(op) - } - Mf::Max => { - let op = match self - .get_type_inner(&ir_module.types, arg0_lookup_ty) - .scalar_kind() - { - Some(crate::ScalarKind::Float) => spirv::GLOp::FMax, - Some(crate::ScalarKind::Sint) => spirv::GLOp::SMax, - Some(crate::ScalarKind::Uint) => spirv::GLOp::UMax, - other => unimplemented!("Unexpected max({:?})", other), - }; - MathOp::Double(op) - } - Mf::Clamp => { - let op = match self - .get_type_inner(&ir_module.types, arg0_lookup_ty) - .scalar_kind() - { - Some(crate::ScalarKind::Float) => spirv::GLOp::FClamp, - Some(crate::ScalarKind::Sint) => spirv::GLOp::SClamp, - Some(crate::ScalarKind::Uint) => spirv::GLOp::UClamp, - other => unimplemented!("Unexpected max({:?})", other), - }; - MathOp::Triple(op) + } } + Mf::Min => MathOp::Ext(match arg_scalar_kind { + Some(crate::ScalarKind::Float) => spirv::GLOp::FMin, + Some(crate::ScalarKind::Sint) => spirv::GLOp::SMin, + Some(crate::ScalarKind::Uint) => spirv::GLOp::UMin, + other => unimplemented!("Unexpected min({:?})", other), + }), + Mf::Max => MathOp::Ext(match arg_scalar_kind { + Some(crate::ScalarKind::Float) => spirv::GLOp::FMax, + Some(crate::ScalarKind::Sint) => spirv::GLOp::SMax, + Some(crate::ScalarKind::Uint) => spirv::GLOp::UMax, + other => unimplemented!("Unexpected max({:?})", other), + }), + Mf::Clamp => MathOp::Ext(match arg_scalar_kind { + Some(crate::ScalarKind::Float) => spirv::GLOp::FClamp, + Some(crate::ScalarKind::Sint) => spirv::GLOp::SClamp, + Some(crate::ScalarKind::Uint) => spirv::GLOp::UClamp, + other => unimplemented!("Unexpected max({:?})", other), + }), // trigonometry - Mf::Sin => MathOp::Single(spirv::GLOp::Sin), - Mf::Sinh => MathOp::Single(spirv::GLOp::Sinh), - Mf::Asin => MathOp::Single(spirv::GLOp::Asin), - Mf::Cos => MathOp::Single(spirv::GLOp::Cos), - Mf::Cosh => MathOp::Single(spirv::GLOp::Cosh), - Mf::Acos => MathOp::Single(spirv::GLOp::Acos), - Mf::Tan => MathOp::Single(spirv::GLOp::Tan), - Mf::Tanh => MathOp::Single(spirv::GLOp::Tanh), - Mf::Atan => MathOp::Single(spirv::GLOp::Atan), - Mf::Atan2 => MathOp::Double(spirv::GLOp::Atan2), + Mf::Sin => MathOp::Ext(spirv::GLOp::Sin), + Mf::Sinh => MathOp::Ext(spirv::GLOp::Sinh), + Mf::Asin => MathOp::Ext(spirv::GLOp::Asin), + Mf::Cos => MathOp::Ext(spirv::GLOp::Cos), + Mf::Cosh => MathOp::Ext(spirv::GLOp::Cosh), + Mf::Acos => MathOp::Ext(spirv::GLOp::Acos), + Mf::Tan => MathOp::Ext(spirv::GLOp::Tan), + Mf::Tanh => MathOp::Ext(spirv::GLOp::Tanh), + Mf::Atan => MathOp::Ext(spirv::GLOp::Atan), + Mf::Atan2 => MathOp::Ext(spirv::GLOp::Atan2), // decomposition - Mf::Ceil => MathOp::Single(spirv::GLOp::Ceil), - Mf::Round => MathOp::Single(spirv::GLOp::Round), - Mf::Floor => MathOp::Single(spirv::GLOp::Floor), - Mf::Fract => MathOp::Single(spirv::GLOp::Fract), - Mf::Trunc => MathOp::Single(spirv::GLOp::Trunc), - Mf::Modf => MathOp::Double(spirv::GLOp::Modf), - Mf::Frexp => MathOp::Double(spirv::GLOp::Frexp), - Mf::Ldexp => MathOp::Double(spirv::GLOp::Ldexp), + Mf::Ceil => MathOp::Ext(spirv::GLOp::Ceil), + Mf::Round => MathOp::Ext(spirv::GLOp::Round), + Mf::Floor => MathOp::Ext(spirv::GLOp::Floor), + Mf::Fract => MathOp::Ext(spirv::GLOp::Fract), + Mf::Trunc => MathOp::Ext(spirv::GLOp::Trunc), + Mf::Modf => MathOp::Ext(spirv::GLOp::Modf), + Mf::Frexp => MathOp::Ext(spirv::GLOp::Frexp), + Mf::Ldexp => MathOp::Ext(spirv::GLOp::Ldexp), // geometry - Mf::Dot => { - let result_lookup_ty = - match *self.get_type_inner(&ir_module.types, arg0_lookup_ty) { - crate::TypeInner::Scalar { kind, width } - | crate::TypeInner::Vector { - size: _, - kind, - width, - } => LookupType::Local(LocalType::Scalar { kind, width }), - _ => unreachable!(), - }; - let result_type_id = - self.get_type_id(&ir_module.types, result_lookup_ty)?; - - let inst = super::instructions::instruction_binary( - spirv::Op::Dot, - result_type_id, - id, - arg0_id, - arg1_id, - ); - MathOp::Other(inst, result_lookup_ty) - } - Mf::Cross => MathOp::Double(spirv::GLOp::Cross), - Mf::Distance => { - let result_lookup_ty = - match *self.get_type_inner(&ir_module.types, arg0_lookup_ty) { - crate::TypeInner::Scalar { kind, width } - | crate::TypeInner::Vector { - size: _, - kind, - width, - } => LookupType::Local(LocalType::Scalar { kind, width }), - _ => unreachable!(), - }; - let result_type_id = - self.get_type_id(&ir_module.types, result_lookup_ty)?; - - let inst = super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - spirv::GLOp::Distance, - result_type_id, - id, - &[arg0_id, arg1_id], - ); - MathOp::Other(inst, result_lookup_ty) - } - Mf::Length => { - let result_lookup_ty = - match *self.get_type_inner(&ir_module.types, arg0_lookup_ty) { - crate::TypeInner::Scalar { kind, width } - | crate::TypeInner::Vector { - size: _, - kind, - width, - } => LookupType::Local(LocalType::Scalar { kind, width }), - _ => unreachable!(), - }; - let result_type_id = - self.get_type_id(&ir_module.types, result_lookup_ty)?; - - let inst = super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - spirv::GLOp::Length, - result_type_id, - id, - &[arg0_id], - ); - MathOp::Other(inst, result_lookup_ty) - } - Mf::Normalize => MathOp::Single(spirv::GLOp::Normalize), - Mf::FaceForward => MathOp::Triple(spirv::GLOp::FaceForward), - Mf::Reflect => MathOp::Double(spirv::GLOp::Reflect), + Mf::Dot => MathOp::Custom(super::instructions::instruction_binary( + spirv::Op::Dot, + result_type_id, + id, + arg0_id, + arg1_id, + )), + Mf::Cross => MathOp::Ext(spirv::GLOp::Cross), + Mf::Distance => MathOp::Ext(spirv::GLOp::Distance), + Mf::Length => MathOp::Ext(spirv::GLOp::Length), + Mf::Normalize => MathOp::Ext(spirv::GLOp::Normalize), + Mf::FaceForward => MathOp::Ext(spirv::GLOp::FaceForward), + Mf::Reflect => MathOp::Ext(spirv::GLOp::Reflect), // exponent - Mf::Exp => MathOp::Single(spirv::GLOp::Exp), - Mf::Exp2 => MathOp::Single(spirv::GLOp::Exp2), - Mf::Log => MathOp::Single(spirv::GLOp::Log), - Mf::Log2 => MathOp::Single(spirv::GLOp::Log2), - Mf::Pow => MathOp::Double(spirv::GLOp::Pow), + Mf::Exp => MathOp::Ext(spirv::GLOp::Exp), + Mf::Exp2 => MathOp::Ext(spirv::GLOp::Exp2), + Mf::Log => MathOp::Ext(spirv::GLOp::Log), + Mf::Log2 => MathOp::Ext(spirv::GLOp::Log2), + Mf::Pow => MathOp::Ext(spirv::GLOp::Pow), // computational - Mf::Sign => { - let op = match self - .get_type_inner(&ir_module.types, arg0_lookup_ty) - .scalar_kind() - { - Some(crate::ScalarKind::Float) => spirv::GLOp::FSign, - Some(crate::ScalarKind::Sint) => spirv::GLOp::SSign, - other => unimplemented!("Unexpected sign({:?})", other), - }; - MathOp::Single(op) - } - Mf::Fma => MathOp::Triple(spirv::GLOp::Fma), - Mf::Mix => { - let op = match self - .get_type_inner(&ir_module.types, arg0_lookup_ty) - .scalar_kind() - { - Some(crate::ScalarKind::Float) => spirv::GLOp::FMix, - other => unimplemented!("Unexpected mix({:?})", other), - }; - MathOp::Triple(op) - } - Mf::Step => MathOp::Triple(spirv::GLOp::Step), - Mf::SmoothStep => MathOp::Triple(spirv::GLOp::SmoothStep), - Mf::Sqrt => MathOp::Single(spirv::GLOp::Sqrt), - Mf::InverseSqrt => MathOp::Single(spirv::GLOp::InverseSqrt), - Mf::Transpose => { - let result_lookup_ty = - match *self.get_type_inner(&ir_module.types, arg0_lookup_ty) { - crate::TypeInner::Matrix { - columns, - rows, - width, - } => LookupType::Local(LocalType::Matrix { - columns: rows, - rows: columns, - width, - }), - _ => unreachable!(), - }; - let result_type_id = - self.get_type_id(&ir_module.types, result_lookup_ty)?; - - let inst = super::instructions::instruction_unary( - spirv::Op::Transpose, - result_type_id, - id, - arg0_id, - ); - MathOp::Other(inst, result_lookup_ty) - } - Mf::Determinant => { - let result_lookup_ty = - match *self.get_type_inner(&ir_module.types, arg0_lookup_ty) { - crate::TypeInner::Matrix { width, .. } => { - LookupType::Local(LocalType::Scalar { - kind: crate::ScalarKind::Float, - width, - }) - } - _ => unreachable!(), - }; - let result_type_id = - self.get_type_id(&ir_module.types, result_lookup_ty)?; - - let inst = super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - spirv::GLOp::Determinant, - result_type_id, - id, - &[arg0_id], - ); - MathOp::Other(inst, result_lookup_ty) - } - + Mf::Sign => MathOp::Ext(match arg_scalar_kind { + Some(crate::ScalarKind::Float) => spirv::GLOp::FSign, + Some(crate::ScalarKind::Sint) => spirv::GLOp::SSign, + other => unimplemented!("Unexpected sign({:?})", other), + }), + Mf::Fma => MathOp::Ext(spirv::GLOp::Fma), + Mf::Mix => MathOp::Ext(spirv::GLOp::FMix), + Mf::Step => MathOp::Ext(spirv::GLOp::Step), + 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( + spirv::Op::Transpose, + result_type_id, + id, + arg0_id, + )), + Mf::Determinant => MathOp::Ext(spirv::GLOp::Determinant), Mf::Outer | Mf::ReverseBits | Mf::CountOneBits => { log::error!("unimplemented math function {:?}", fun); return Err(Error::FeatureNotImplemented("math function")); } }; - let (instruction, result_lookup_ty) = match math_op { - MathOp::Single(op) => { - let inst = super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - op, - arg0_type_id, - id, - &[arg0_id], - ); - (inst, arg0_lookup_ty) - } - MathOp::Double(op) => { - let inst = super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - op, - arg0_type_id, - id, - &[arg0_id, arg1_id], - ); - (inst, arg0_lookup_ty) - } - MathOp::Triple(op) => { - let inst = super::instructions::instruction_ext_inst( - self.gl450_ext_inst_id, - op, - arg0_type_id, - id, - &[arg0_id, arg1_id, arg2_id], - ); - (inst, arg0_lookup_ty) - } - MathOp::Other(inst, result_lookup_ty) => (inst, result_lookup_ty), - }; - - block.body.push(instruction); - Ok((RawExpression::Value(id), result_lookup_ty)) + block.body.push(match math_op { + MathOp::Ext(op) => super::instructions::instruction_ext_inst( + self.gl450_ext_inst_id, + op, + result_type_id, + id, + &[arg0_id, arg1_id, arg2_id][..fun.argument_count()], + ), + MathOp::Custom(inst) => inst, + }); + RawExpression::Value(id) } crate::Expression::LocalVariable(variable) => { - let var = &ir_function.local_variables[variable]; let local_var = &function.variables[&variable]; - Ok(( - RawExpression::Pointer(local_var.id, spirv::StorageClass::Function), - LookupType::Handle(var.ty), - )) + RawExpression::Pointer(local_var.id, spirv::StorageClass::Function) } crate::Expression::FunctionArgument(index) => { - let ty_handle = ir_function.arguments[index as usize].ty; let id = function.parameters[index as usize].result_id.unwrap(); - Ok((RawExpression::Value(id), LookupType::Handle(ty_handle))) + RawExpression::Value(id) } crate::Expression::Call { function: local_function, ref arguments, } => { - let target_function = &ir_module.functions[local_function]; let id = self.generate_id(); + //TODO: avoid heap allocation let mut argument_ids = vec![]; for argument in arguments { - let (arg_id, _) = + let arg_id = self.write_expression(ir_module, ir_function, *argument, block, function)?; argument_ids.push(arg_id); } - let return_type_id = - self.get_function_return_type(target_function.return_type, &ir_module.types)?; - block .body .push(super::instructions::instruction_function_call( - return_type_id, + result_type_id, id, *self.lookup_function.get(&local_function).unwrap(), argument_ids.as_slice(), )); - let result_type = match target_function.return_type { - Some(ty_handle) => LookupType::Handle(ty_handle), - None => LookupType::Local(LocalType::Void), - }; - Ok((RawExpression::Value(id), result_type)) + RawExpression::Value(id) } crate::Expression::As { expr, @@ -1880,25 +1593,14 @@ impl Writer { return Err(Error::FeatureNotImplemented("bitcast")); } - let (expr_id, expr_type) = + let expr_id = self.write_expression(ir_module, ir_function, expr, block, function)?; + let expr_kind = self + .typifier + .get(expr, &ir_module.types) + .scalar_kind() + .unwrap(); - let expr_type_inner = self.get_type_inner(&ir_module.types, expr_type); - - let (expr_kind, local_type) = match *expr_type_inner { - crate::TypeInner::Scalar { - kind: expr_kind, - width, - } => (expr_kind, LocalType::Scalar { kind, width }), - crate::TypeInner::Vector { - size, - kind: expr_kind, - width, - } => (expr_kind, LocalType::Vector { size, kind, width }), - _ => unreachable!(), - }; - - let lookup_type = LookupType::Local(local_type); let op = match (expr_kind, kind) { _ if !convert => spirv::Op::Bitcast, (crate::ScalarKind::Float, crate::ScalarKind::Uint) => spirv::Op::ConvertFToU, @@ -1906,17 +1608,15 @@ impl Writer { (crate::ScalarKind::Sint, crate::ScalarKind::Float) => spirv::Op::ConvertSToF, (crate::ScalarKind::Uint, crate::ScalarKind::Float) => spirv::Op::ConvertUToF, // We assume it's either an identity cast, or int-uint. - // In both cases no SPIR-V instructions need to be generated. - _ => return Ok((RawExpression::Value(expr_id), lookup_type)), + _ => spirv::Op::Bitcast, }; let id = self.generate_id(); - let kind_type_id = self.get_type_id(&ir_module.types, lookup_type)?; let instruction = - super::instructions::instruction_unary(op, kind_type_id, id, expr_id); + super::instructions::instruction_unary(op, result_type_id, id, expr_id); block.body.push(instruction); - Ok((RawExpression::Value(id), lookup_type)) + RawExpression::Value(id) } crate::Expression::ImageSample { image, @@ -1929,28 +1629,22 @@ impl Writer { } => { use super::instructions::SampleLod; // image - let (image_id, image_lookup_ty) = + let image_id = self.write_expression(ir_module, ir_function, image, block, function)?; - - let image_ty = match image_lookup_ty { - LookupType::Handle(handle) => handle, - LookupType::Local(_) => unreachable!(), - }; + let image_type = self.typifier.get_handle(image).unwrap(); // OpTypeSampledImage let sampled_image_type_id = self.get_type_id( &ir_module.types, - LookupType::Local(LocalType::SampledImage { - image_type: image_ty, - }), + LookupType::Local(LocalType::SampledImage { image_type }), )?; // sampler - let (sampler_id, _) = + let sampler_id = self.write_expression(ir_module, ir_function, sampler, block, function)?; // coordinate - let (mut coordinate_id, coordinate_lookup_ty) = + let mut coordinate_id = self.write_expression(ir_module, ir_function, coordinate, block, function)?; if let Some(array_index) = array_index { @@ -1963,7 +1657,7 @@ impl Writer { )?; let mut constituent_ids = [0u32; 4]; - let size = match *self.get_type_inner(&ir_module.types, coordinate_lookup_ty) { + let size = match *self.typifier.get(coordinate, &ir_module.types) { crate::TypeInner::Scalar { .. } => { constituent_ids[0] = coordinate_id; crate::VectorSize::Bi @@ -1995,7 +1689,7 @@ impl Writer { let array_index_f32_id = self.generate_id(); constituent_ids[size as usize - 1] = array_index_f32_id; - let (array_index_u32_id, _) = self.write_expression( + let array_index_u32_id = self.write_expression( ir_module, ir_function, array_index, @@ -2026,29 +1720,6 @@ impl Writer { ); } - // component kind - let image_type = &ir_module.types[image_ty]; - let image_sample_result_type = match image_type.inner { - crate::TypeInner::Image { class, .. } => { - let width = 4; - LookupType::Local(match class { - crate::ImageClass::Sampled { kind, multi: _ } => LocalType::Vector { - kind, - width, - size: crate::VectorSize::Quad, - }, - crate::ImageClass::Depth => LocalType::Scalar { - kind: crate::ScalarKind::Float, - width, - }, - crate::ImageClass::Storage(_) => { - unimplemented!("Unexpected storage image being sampled") - } - }) - } - ref other => unimplemented!("Unexpected image type {:?}", other), - }; - let sampled_image_id = self.generate_id(); block .body @@ -2059,12 +1730,10 @@ impl Writer { sampler_id, )); let id = self.generate_id(); - let image_sample_result_type_id = - self.get_type_id(&ir_module.types, image_sample_result_type)?; let depth_id = match depth_ref { Some(handle) => { - let (expr_id, _) = + let expr_id = self.write_expression(ir_module, ir_function, handle, block, function)?; Some(expr_id) } @@ -2074,7 +1743,7 @@ impl Writer { let mut main_instruction = match level { crate::SampleLevel::Zero => { let mut inst = super::instructions::instruction_image_sample( - image_sample_result_type_id, + result_type_id, id, SampleLod::Explicit, sampled_image_id, @@ -2095,7 +1764,7 @@ impl Writer { inst } crate::SampleLevel::Auto => super::instructions::instruction_image_sample( - image_sample_result_type_id, + result_type_id, id, SampleLod::Implicit, sampled_image_id, @@ -2104,7 +1773,7 @@ impl Writer { ), crate::SampleLevel::Exact(lod_handle) => { let mut inst = super::instructions::instruction_image_sample( - image_sample_result_type_id, + result_type_id, id, SampleLod::Explicit, sampled_image_id, @@ -2112,7 +1781,7 @@ impl Writer { depth_id, ); - let (lod_id, _) = self.write_expression( + let lod_id = self.write_expression( ir_module, ir_function, lod_handle, @@ -2126,7 +1795,7 @@ impl Writer { } crate::SampleLevel::Bias(bias_handle) => { let mut inst = super::instructions::instruction_image_sample( - image_sample_result_type_id, + result_type_id, id, SampleLod::Implicit, sampled_image_id, @@ -2134,7 +1803,7 @@ impl Writer { depth_id, ); - let (bias_id, _) = self.write_expression( + let bias_id = self.write_expression( ir_module, ir_function, bias_handle, @@ -2148,7 +1817,7 @@ impl Writer { } crate::SampleLevel::Gradient { x, y } => { let mut inst = super::instructions::instruction_image_sample( - image_sample_result_type_id, + result_type_id, id, SampleLod::Explicit, sampled_image_id, @@ -2156,9 +1825,9 @@ impl Writer { depth_id, ); - let (x_id, _) = + let x_id = self.write_expression(ir_module, ir_function, x, block, function)?; - let (y_id, _) = + let y_id = self.write_expression(ir_module, ir_function, y, block, function)?; inst.add_operand(spirv::ImageOperands::GRAD.bits()); inst.add_operand(x_id); @@ -2175,13 +1844,15 @@ impl Writer { } block.body.push(main_instruction); - Ok((RawExpression::Value(id), image_sample_result_type)) + RawExpression::Value(id) } ref other => { log::error!("unimplemented {:?}", other); - Err(Error::FeatureNotImplemented("expression")) + return Err(Error::FeatureNotImplemented("expression")); } - } + }; + + Ok((raw_expr, result_type_id)) } #[allow(clippy::too_many_arguments)] @@ -2224,7 +1895,7 @@ impl Writer { ref accept, ref reject, } => { - let (condition_id, _) = self.write_expression( + let condition_id = self.write_expression( ir_module, ir_function, *condition, @@ -2332,7 +2003,7 @@ impl Writer { )); } crate::Statement::Return { value: Some(value) } => { - let (id, _) = + let id = self.write_expression(ir_module, ir_function, value, &mut block, function)?; block.termination = Some(super::instructions::instruction_return_value(id)); } @@ -2343,14 +2014,14 @@ impl Writer { block.termination = Some(super::instructions::instruction_kill()); } crate::Statement::Store { pointer, value } => { - let (pointer_id, _, _) = self.write_expression_pointer( + let (pointer_id, _) = self.write_expression_pointer( ir_module, ir_function, pointer, &mut block, function, )?; - let (value_id, _) = + let value_id = self.write_expression(ir_module, ir_function, value, &mut block, function)?; block.body.push(super::instructions::instruction_store( diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index d41b17f891..b367713c9a 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -95,10 +95,10 @@ impl Typifier { pub fn get_handle( &self, expr_handle: Handle, - ) -> Option> { + ) -> Result, &crate::TypeInner> { match self.resolutions[expr_handle.index()] { - Resolution::Handle(ty_handle) => Some(ty_handle), - Resolution::Value(_) => None, + Resolution::Handle(ty_handle) => Ok(ty_handle), + Resolution::Value(ref inner) => Err(inner), } } @@ -258,7 +258,23 @@ impl Typifier { | crate::BinaryOperator::Greater | crate::BinaryOperator::GreaterEqual | crate::BinaryOperator::LogicalAnd - | crate::BinaryOperator::LogicalOr => self.resolutions[left.index()].clone(), + | crate::BinaryOperator::LogicalOr => { + let kind = crate::ScalarKind::Bool; + let width = 1; + let inner = match *self.get(left, types) { + crate::TypeInner::Scalar { .. } => crate::TypeInner::Scalar { kind, width }, + crate::TypeInner::Vector { size, .. } => { + crate::TypeInner::Vector { size, kind, width } + } + ref other => { + return Err(ResolveError::IncompatibleOperand { + op: "logical".to_string(), + operand: format!("{:?}", other), + }) + } + }; + Resolution::Value(inner) + } crate::BinaryOperator::And | crate::BinaryOperator::ExclusiveOr | crate::BinaryOperator::InclusiveOr diff --git a/tests/snapshots/snapshots__boids.msl.snap b/tests/snapshots/snapshots__boids.msl.snap index 4cc30a5530..47cdc704cf 100644 --- a/tests/snapshots/snapshots__boids.msl.snap +++ b/tests/snapshots/snapshots__boids.msl.snap @@ -36,6 +36,8 @@ typedef metal::uint3 type4; typedef int type5; +typedef bool type6; + kernel void main1( constant SimParams& params [[buffer(0)]], constant Particles& particlesSrc [[buffer(1)]], diff --git a/tests/snapshots/snapshots__boids.spvasm.snap b/tests/snapshots/snapshots__boids.spvasm.snap index d98c8e7433..fac29f9aa5 100644 --- a/tests/snapshots/snapshots__boids.spvasm.snap +++ b/tests/snapshots/snapshots__boids.spvasm.snap @@ -9,26 +9,26 @@ expression: dis OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %29 "main" %34 +OpEntryPoint GLCompute %29 "main" %36 OpExecutionMode %29 LocalSize 64 1 1 -OpDecorate %34 BuiltIn GlobalInvocationId -OpDecorate %48 BufferBlock -OpMemberDecorate %48 0 Offset 0 -OpDecorate %49 ArrayStride 16 -OpMemberDecorate %50 0 Offset 0 -OpMemberDecorate %50 1 Offset 8 -OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 1 -OpDecorate %122 Block -OpMemberDecorate %122 0 Offset 0 -OpMemberDecorate %122 1 Offset 4 -OpMemberDecorate %122 2 Offset 8 -OpMemberDecorate %122 3 Offset 12 -OpMemberDecorate %122 4 Offset 16 -OpMemberDecorate %122 5 Offset 20 -OpMemberDecorate %122 6 Offset 24 -OpDecorate %121 DescriptorSet 0 -OpDecorate %121 Binding 0 +OpDecorate %36 BuiltIn GlobalInvocationId +OpMemberDecorate %45 0 Offset 0 +OpMemberDecorate %45 1 Offset 8 +OpDecorate %47 ArrayStride 16 +OpDecorate %49 BufferBlock +OpMemberDecorate %49 0 Offset 0 +OpDecorate %50 DescriptorSet 0 +OpDecorate %50 Binding 1 +OpDecorate %121 Block +OpMemberDecorate %121 0 Offset 0 +OpMemberDecorate %121 1 Offset 4 +OpMemberDecorate %121 2 Offset 8 +OpMemberDecorate %121 3 Offset 12 +OpMemberDecorate %121 4 Offset 16 +OpMemberDecorate %121 5 Offset 20 +OpMemberDecorate %121 6 Offset 24 +OpDecorate %122 DescriptorSet 0 +OpDecorate %122 Binding 0 OpDecorate %274 DescriptorSet 0 OpDecorate %274 Binding 2 %3 = OpTypeInt 32 1 @@ -49,91 +49,91 @@ OpDecorate %274 Binding 2 %27 = OpTypePointer Function %8 %28 = OpTypeVoid %30 = OpTypeFunction %28 +%32 = OpTypeBool %35 = OpTypeVector %8 3 -%36 = OpTypePointer Input %35 -%34 = OpVariable %36 Input -%37 = OpConstant %3 0 +%37 = OpTypePointer Input %35 +%36 = OpVariable %37 Input %38 = OpTypePointer Input %8 -%40 = OpTypeBool -%50 = OpTypeStruct %15 %15 -%49 = OpTypeRuntimeArray %50 -%48 = OpTypeStruct %49 -%51 = OpTypePointer Uniform %48 -%47 = OpVariable %51 Uniform -%52 = OpConstant %3 0 -%53 = OpTypePointer Uniform %49 -%55 = OpConstant %3 0 -%56 = OpTypePointer Input %8 -%58 = OpTypePointer Uniform %50 -%59 = OpConstant %3 0 -%60 = OpTypePointer Uniform %15 -%65 = OpConstant %3 0 -%66 = OpTypePointer Uniform %49 -%68 = OpConstant %3 0 -%69 = OpTypePointer Input %8 -%71 = OpTypePointer Uniform %50 -%72 = OpConstant %3 1 -%73 = OpTypePointer Uniform %15 -%90 = OpConstant %3 0 -%91 = OpTypePointer Input %8 -%99 = OpConstant %3 0 -%100 = OpTypePointer Uniform %49 -%102 = OpTypePointer Uniform %50 -%103 = OpConstant %3 0 -%104 = OpTypePointer Uniform %15 -%109 = OpConstant %3 0 -%110 = OpTypePointer Uniform %49 -%112 = OpTypePointer Uniform %50 -%113 = OpConstant %3 1 -%114 = OpTypePointer Uniform %15 -%122 = OpTypeStruct %5 %5 %5 %5 %5 %5 %5 -%123 = OpTypePointer Uniform %122 -%121 = OpVariable %123 Uniform -%124 = OpConstant %3 1 -%125 = OpTypePointer Uniform %5 -%140 = OpConstant %3 2 -%141 = OpTypePointer Uniform %5 -%156 = OpConstant %3 3 -%157 = OpTypePointer Uniform %5 -%198 = OpConstant %3 4 -%199 = OpTypePointer Uniform %5 -%204 = OpConstant %3 5 -%205 = OpTypePointer Uniform %5 -%210 = OpConstant %3 6 -%211 = OpTypePointer Uniform %5 -%224 = OpConstant %3 0 -%225 = OpTypePointer Uniform %5 -%229 = OpConstant %3 0 -%230 = OpTypePointer Function %5 -%236 = OpConstant %3 0 -%237 = OpTypePointer Function %5 -%240 = OpConstant %3 0 -%241 = OpTypePointer Function %5 -%247 = OpConstant %3 0 -%248 = OpTypePointer Function %5 -%251 = OpConstant %3 1 -%252 = OpTypePointer Function %5 -%258 = OpConstant %3 1 -%259 = OpTypePointer Function %5 -%262 = OpConstant %3 1 -%263 = OpTypePointer Function %5 -%269 = OpConstant %3 1 -%270 = OpTypePointer Function %5 +%39 = OpConstant %3 0 +%45 = OpTypeStruct %15 %15 +%47 = OpTypeRuntimeArray %45 +%49 = OpTypeStruct %47 +%51 = OpTypePointer Uniform %49 +%50 = OpVariable %51 Uniform +%52 = OpTypePointer Uniform %47 +%53 = OpConstant %3 0 +%55 = OpTypePointer Input %8 +%56 = OpConstant %3 0 +%58 = OpTypePointer Uniform %45 +%59 = OpTypePointer Uniform %15 +%60 = OpConstant %3 0 +%65 = OpTypePointer Uniform %47 +%66 = OpConstant %3 0 +%68 = OpTypePointer Input %8 +%69 = OpConstant %3 0 +%71 = OpTypePointer Uniform %45 +%72 = OpTypePointer Uniform %15 +%73 = OpConstant %3 1 +%90 = OpTypePointer Input %8 +%91 = OpConstant %3 0 +%99 = OpTypePointer Uniform %47 +%100 = OpConstant %3 0 +%102 = OpTypePointer Uniform %45 +%103 = OpTypePointer Uniform %15 +%104 = OpConstant %3 0 +%109 = OpTypePointer Uniform %47 +%110 = OpConstant %3 0 +%112 = OpTypePointer Uniform %45 +%113 = OpTypePointer Uniform %15 +%114 = OpConstant %3 1 +%121 = OpTypeStruct %5 %5 %5 %5 %5 %5 %5 +%123 = OpTypePointer Uniform %121 +%122 = OpVariable %123 Uniform +%124 = OpTypePointer Uniform %5 +%125 = OpConstant %3 1 +%140 = OpTypePointer Uniform %5 +%141 = OpConstant %3 2 +%156 = OpTypePointer Uniform %5 +%157 = OpConstant %3 3 +%198 = OpTypePointer Uniform %5 +%199 = OpConstant %3 4 +%204 = OpTypePointer Uniform %5 +%205 = OpConstant %3 5 +%210 = OpTypePointer Uniform %5 +%211 = OpConstant %3 6 +%224 = OpTypePointer Uniform %5 +%225 = OpConstant %3 0 +%229 = OpTypePointer Function %5 +%230 = OpConstant %3 0 +%236 = OpTypePointer Function %5 +%237 = OpConstant %3 0 +%240 = OpTypePointer Function %5 +%241 = OpConstant %3 0 +%247 = OpTypePointer Function %5 +%248 = OpConstant %3 0 +%251 = OpTypePointer Function %5 +%252 = OpConstant %3 1 +%258 = OpTypePointer Function %5 +%259 = OpConstant %3 1 +%262 = OpTypePointer Function %5 +%263 = OpConstant %3 1 +%269 = OpTypePointer Function %5 +%270 = OpConstant %3 1 %274 = OpVariable %51 Uniform -%275 = OpConstant %3 0 -%276 = OpTypePointer Uniform %49 -%278 = OpConstant %3 0 -%279 = OpTypePointer Input %8 -%281 = OpTypePointer Uniform %50 -%282 = OpConstant %3 0 -%283 = OpTypePointer Uniform %15 -%288 = OpConstant %3 0 -%289 = OpTypePointer Uniform %49 -%291 = OpConstant %3 0 -%292 = OpTypePointer Input %8 -%294 = OpTypePointer Uniform %50 -%295 = OpConstant %3 1 -%296 = OpTypePointer Uniform %15 +%275 = OpTypePointer Uniform %47 +%276 = OpConstant %3 0 +%278 = OpTypePointer Input %8 +%279 = OpConstant %3 0 +%281 = OpTypePointer Uniform %45 +%282 = OpTypePointer Uniform %15 +%283 = OpConstant %3 0 +%288 = OpTypePointer Uniform %47 +%289 = OpConstant %3 0 +%291 = OpTypePointer Input %8 +%292 = OpConstant %3 0 +%294 = OpTypePointer Uniform %45 +%295 = OpTypePointer Uniform %15 +%296 = OpConstant %3 1 %29 = OpFunction %28 None %30 %31 = OpLabel %26 = OpVariable %27 Function %7 @@ -146,28 +146,28 @@ OpDecorate %274 Binding 2 %25 = OpVariable %16 Function %21 = OpVariable %22 Function %6 %18 = OpVariable %16 Function -%33 = OpAccessChain %38 %34 %37 -%39 = OpLoad %8 %33 -%32 = OpUGreaterThanEqual %40 %39 %2 +%34 = OpAccessChain %38 %36 %39 +%40 = OpLoad %8 %34 +%33 = OpUGreaterThanEqual %32 %40 %2 OpSelectionMerge %41 None -OpBranchConditional %32 %42 %43 +OpBranchConditional %33 %42 %43 %42 = OpLabel OpReturn %43 = OpLabel OpBranch %41 %41 = OpLabel -%46 = OpAccessChain %53 %47 %52 -%54 = OpAccessChain %56 %34 %55 +%48 = OpAccessChain %52 %50 %53 +%54 = OpAccessChain %55 %36 %56 %57 = OpLoad %8 %54 -%45 = OpAccessChain %58 %46 %57 -%44 = OpAccessChain %60 %45 %59 +%46 = OpAccessChain %58 %48 %57 +%44 = OpAccessChain %59 %46 %60 %61 = OpLoad %15 %44 OpStore %14 %61 -%64 = OpAccessChain %66 %47 %65 -%67 = OpAccessChain %69 %34 %68 +%64 = OpAccessChain %65 %50 %66 +%67 = OpAccessChain %68 %36 %69 %70 = OpLoad %8 %67 %63 = OpAccessChain %71 %64 %70 -%62 = OpAccessChain %73 %63 %72 +%62 = OpAccessChain %72 %63 %73 %74 = OpLoad %15 %62 OpStore %17 %74 %75 = OpCompositeConstruct %15 %4 %4 @@ -182,7 +182,7 @@ OpLoopMerge %79 %81 None OpBranch %80 %80 = OpLabel %83 = OpLoad %8 %26 -%82 = OpUGreaterThanEqual %40 %83 %2 +%82 = OpUGreaterThanEqual %32 %83 %2 OpSelectionMerge %84 None OpBranchConditional %82 %85 %86 %85 = OpLabel @@ -191,9 +191,9 @@ OpBranch %79 OpBranch %84 %84 = OpLabel %88 = OpLoad %8 %26 -%89 = OpAccessChain %91 %34 %90 +%89 = OpAccessChain %90 %36 %91 %92 = OpLoad %8 %89 -%87 = OpIEqual %40 %88 %92 +%87 = OpIEqual %32 %88 %92 OpSelectionMerge %93 None OpBranchConditional %87 %94 %95 %94 = OpLabel @@ -201,24 +201,24 @@ OpBranch %81 %95 = OpLabel OpBranch %93 %93 = OpLabel -%98 = OpAccessChain %100 %47 %99 +%98 = OpAccessChain %99 %50 %100 %101 = OpLoad %8 %26 %97 = OpAccessChain %102 %98 %101 -%96 = OpAccessChain %104 %97 %103 +%96 = OpAccessChain %103 %97 %104 %105 = OpLoad %15 %96 OpStore %24 %105 -%108 = OpAccessChain %110 %47 %109 +%108 = OpAccessChain %109 %50 %110 %111 = OpLoad %8 %26 %107 = OpAccessChain %112 %108 %111 -%106 = OpAccessChain %114 %107 %113 +%106 = OpAccessChain %113 %107 %114 %115 = OpLoad %15 %106 OpStore %25 %115 %117 = OpLoad %15 %24 %118 = OpLoad %15 %14 %119 = OpExtInst %5 %1 Distance %117 %118 -%120 = OpAccessChain %125 %121 %124 +%120 = OpAccessChain %124 %122 %125 %126 = OpLoad %5 %120 -%116 = OpFOrdLessThan %40 %119 %126 +%116 = OpFOrdLessThan %32 %119 %126 OpSelectionMerge %127 None OpBranchConditional %116 %128 %129 %128 = OpLabel @@ -236,9 +236,9 @@ OpBranch %127 %136 = OpLoad %15 %24 %137 = OpLoad %15 %14 %138 = OpExtInst %5 %1 Distance %136 %137 -%139 = OpAccessChain %141 %121 %140 +%139 = OpAccessChain %140 %122 %141 %142 = OpLoad %5 %139 -%135 = OpFOrdLessThan %40 %138 %142 +%135 = OpFOrdLessThan %32 %138 %142 OpSelectionMerge %143 None OpBranchConditional %135 %144 %145 %144 = OpLabel @@ -255,9 +255,9 @@ OpBranch %143 %152 = OpLoad %15 %24 %153 = OpLoad %15 %14 %154 = OpExtInst %5 %1 Distance %152 %153 -%155 = OpAccessChain %157 %121 %156 +%155 = OpAccessChain %156 %122 %157 %158 = OpLoad %5 %155 -%151 = OpFOrdLessThan %40 %154 %158 +%151 = OpFOrdLessThan %32 %154 %158 OpSelectionMerge %159 None OpBranchConditional %151 %160 %161 %160 = OpLabel @@ -280,7 +280,7 @@ OpStore %26 %167 OpBranch %78 %79 = OpLabel %170 = OpLoad %3 %21 -%169 = OpSGreaterThan %40 %170 %6 +%169 = OpSGreaterThan %32 %170 %6 OpSelectionMerge %171 None OpBranchConditional %169 %172 %173 %172 = OpLabel @@ -297,7 +297,7 @@ OpBranch %171 OpBranch %171 %171 = OpLabel %182 = OpLoad %3 %23 -%181 = OpSGreaterThan %40 %182 %6 +%181 = OpSGreaterThan %32 %182 %6 OpSelectionMerge %183 None OpBranchConditional %181 %184 %185 %184 = OpLabel @@ -313,17 +313,17 @@ OpBranch %183 %183 = OpLabel %194 = OpLoad %15 %17 %196 = OpLoad %15 %18 -%197 = OpAccessChain %199 %121 %198 +%197 = OpAccessChain %198 %122 %199 %200 = OpLoad %5 %197 %195 = OpVectorTimesScalar %15 %196 %200 %193 = OpFAdd %15 %194 %195 %202 = OpLoad %15 %20 -%203 = OpAccessChain %205 %121 %204 +%203 = OpAccessChain %204 %122 %205 %206 = OpLoad %5 %203 %201 = OpVectorTimesScalar %15 %202 %206 %192 = OpFAdd %15 %193 %201 %208 = OpLoad %15 %19 -%209 = OpAccessChain %211 %121 %210 +%209 = OpAccessChain %210 %122 %211 %212 = OpLoad %5 %209 %207 = OpVectorTimesScalar %15 %208 %212 %191 = OpFAdd %15 %192 %207 @@ -337,71 +337,71 @@ OpStore %17 %191 OpStore %17 %213 %220 = OpLoad %15 %14 %222 = OpLoad %15 %17 -%223 = OpAccessChain %225 %121 %224 +%223 = OpAccessChain %224 %122 %225 %226 = OpLoad %5 %223 %221 = OpVectorTimesScalar %15 %222 %226 %219 = OpFAdd %15 %220 %221 OpStore %14 %219 -%228 = OpAccessChain %230 %14 %229 +%228 = OpAccessChain %229 %14 %230 %231 = OpLoad %5 %228 -%227 = OpFOrdLessThan %40 %231 %13 +%227 = OpFOrdLessThan %32 %231 %13 OpSelectionMerge %232 None OpBranchConditional %227 %233 %234 %233 = OpLabel -%235 = OpAccessChain %237 %14 %236 +%235 = OpAccessChain %236 %14 %237 OpStore %235 %11 OpBranch %232 %234 = OpLabel OpBranch %232 %232 = OpLabel -%239 = OpAccessChain %241 %14 %240 +%239 = OpAccessChain %240 %14 %241 %242 = OpLoad %5 %239 -%238 = OpFOrdGreaterThan %40 %242 %11 +%238 = OpFOrdGreaterThan %32 %242 %11 OpSelectionMerge %243 None OpBranchConditional %238 %244 %245 %244 = OpLabel -%246 = OpAccessChain %248 %14 %247 +%246 = OpAccessChain %247 %14 %248 OpStore %246 %13 OpBranch %243 %245 = OpLabel OpBranch %243 %243 = OpLabel -%250 = OpAccessChain %252 %14 %251 +%250 = OpAccessChain %251 %14 %252 %253 = OpLoad %5 %250 -%249 = OpFOrdLessThan %40 %253 %13 +%249 = OpFOrdLessThan %32 %253 %13 OpSelectionMerge %254 None OpBranchConditional %249 %255 %256 %255 = OpLabel -%257 = OpAccessChain %259 %14 %258 +%257 = OpAccessChain %258 %14 %259 OpStore %257 %11 OpBranch %254 %256 = OpLabel OpBranch %254 %254 = OpLabel -%261 = OpAccessChain %263 %14 %262 +%261 = OpAccessChain %262 %14 %263 %264 = OpLoad %5 %261 -%260 = OpFOrdGreaterThan %40 %264 %11 +%260 = OpFOrdGreaterThan %32 %264 %11 OpSelectionMerge %265 None OpBranchConditional %260 %266 %267 %266 = OpLabel -%268 = OpAccessChain %270 %14 %269 +%268 = OpAccessChain %269 %14 %270 OpStore %268 %13 OpBranch %265 %267 = OpLabel OpBranch %265 %265 = OpLabel -%273 = OpAccessChain %276 %274 %275 -%277 = OpAccessChain %279 %34 %278 +%273 = OpAccessChain %275 %274 %276 +%277 = OpAccessChain %278 %36 %279 %280 = OpLoad %8 %277 %272 = OpAccessChain %281 %273 %280 -%271 = OpAccessChain %283 %272 %282 +%271 = OpAccessChain %282 %272 %283 %284 = OpLoad %15 %14 OpStore %271 %284 -%287 = OpAccessChain %289 %274 %288 -%290 = OpAccessChain %292 %34 %291 +%287 = OpAccessChain %288 %274 %289 +%290 = OpAccessChain %291 %36 %292 %293 = OpLoad %8 %290 %286 = OpAccessChain %294 %287 %293 -%285 = OpAccessChain %296 %286 %295 +%285 = OpAccessChain %295 %286 %296 %297 = OpLoad %15 %17 OpStore %285 %297 OpReturn diff --git a/tests/snapshots/snapshots__collatz.spvasm.snap b/tests/snapshots/snapshots__collatz.spvasm.snap index b2d30f32d4..89a571a9cd 100644 --- a/tests/snapshots/snapshots__collatz.spvasm.snap +++ b/tests/snapshots/snapshots__collatz.spvasm.snap @@ -9,14 +9,14 @@ expression: dis OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %39 "main" %52 +OpEntryPoint GLCompute %39 "main" %53 OpExecutionMode %39 LocalSize 1 1 1 +OpDecorate %43 ArrayStride 4 OpDecorate %45 BufferBlock OpMemberDecorate %45 0 Offset 0 -OpDecorate %46 ArrayStride 4 -OpDecorate %44 DescriptorSet 0 -OpDecorate %44 Binding 0 -OpDecorate %52 BuiltIn GlobalInvocationId +OpDecorate %46 DescriptorSet 0 +OpDecorate %46 Binding 0 +OpDecorate %53 BuiltIn GlobalInvocationId %3 = OpTypeInt 32 0 %2 = OpConstant %3 0 %4 = OpConstant %3 1 @@ -24,26 +24,26 @@ OpDecorate %52 BuiltIn GlobalInvocationId %6 = OpConstant %3 3 %8 = OpTypePointer Function %3 %12 = OpTypeFunction %3 %3 -%20 = OpTypeBool +%18 = OpTypeBool %38 = OpTypeVoid %40 = OpTypeFunction %38 -%46 = OpTypeRuntimeArray %3 -%45 = OpTypeStruct %46 +%43 = OpTypeRuntimeArray %3 +%45 = OpTypeStruct %43 %47 = OpTypePointer Uniform %45 -%44 = OpVariable %47 Uniform -%48 = OpTypeInt 32 1 -%49 = OpConstant %48 0 -%50 = OpTypePointer Uniform %46 -%53 = OpTypeVector %3 3 -%54 = OpTypePointer Input %53 -%52 = OpVariable %54 Input -%55 = OpConstant %48 0 -%56 = OpTypePointer Input %3 +%46 = OpVariable %47 Uniform +%48 = OpTypePointer Uniform %43 +%49 = OpTypeInt 32 1 +%50 = OpConstant %49 0 +%52 = OpTypeVector %3 3 +%54 = OpTypePointer Input %52 +%53 = OpVariable %54 Input +%55 = OpTypePointer Input %3 +%56 = OpConstant %49 0 %58 = OpTypePointer Uniform %3 -%62 = OpConstant %48 0 -%63 = OpTypePointer Uniform %46 -%65 = OpConstant %48 0 -%66 = OpTypePointer Input %3 +%62 = OpTypePointer Uniform %43 +%63 = OpConstant %49 0 +%65 = OpTypePointer Input %3 +%66 = OpConstant %49 0 %68 = OpTypePointer Uniform %3 %11 = OpFunction %3 None %12 %10 = OpFunctionParameter %3 @@ -56,10 +56,10 @@ OpBranch %14 OpLoopMerge %15 %17 None OpBranch %16 %16 = OpLabel -%19 = OpLoad %3 %7 -%18 = OpULessThanEqual %20 %19 %4 +%20 = OpLoad %3 %7 +%19 = OpULessThanEqual %18 %20 %4 OpSelectionMerge %21 None -OpBranchConditional %18 %22 %23 +OpBranchConditional %19 %22 %23 %22 = OpLabel OpBranch %15 %23 = OpLabel @@ -67,7 +67,7 @@ OpBranch %21 %21 = OpLabel %26 = OpLoad %3 %7 %25 = OpUMod %3 %26 %5 -%24 = OpIEqual %20 %25 %2 +%24 = OpIEqual %18 %25 %2 OpSelectionMerge %27 None OpBranchConditional %24 %28 %29 %28 = OpLabel @@ -94,12 +94,12 @@ OpReturnValue %37 OpFunctionEnd %39 = OpFunction %38 None %40 %41 = OpLabel -%43 = OpAccessChain %50 %44 %49 -%51 = OpAccessChain %56 %52 %55 +%44 = OpAccessChain %48 %46 %50 +%51 = OpAccessChain %55 %53 %56 %57 = OpLoad %3 %51 -%42 = OpAccessChain %58 %43 %57 -%61 = OpAccessChain %63 %44 %62 -%64 = OpAccessChain %66 %52 %65 +%42 = OpAccessChain %58 %44 %57 +%61 = OpAccessChain %62 %46 %63 +%64 = OpAccessChain %65 %53 %66 %67 = OpLoad %3 %64 %60 = OpAccessChain %68 %61 %67 %69 = OpLoad %3 %60 diff --git a/tests/snapshots/snapshots__quad.spvasm.snap b/tests/snapshots/snapshots__quad.spvasm.snap index 145c7399a7..1335b66ec9 100644 --- a/tests/snapshots/snapshots__quad.spvasm.snap +++ b/tests/snapshots/snapshots__quad.spvasm.snap @@ -9,18 +9,18 @@ expression: dis OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %7 "main" %20 %13 %10 %16 +OpEntryPoint Vertex %7 "main" %20 %13 %11 %17 OpEntryPoint Fragment %23 "main" %35 %25 OpExecutionMode %23 OriginUpperLeft -OpDecorate %10 Location 0 +OpDecorate %11 Location 0 OpDecorate %13 Location 1 -OpDecorate %16 BuiltIn Position +OpDecorate %17 BuiltIn Position OpDecorate %20 Location 0 OpDecorate %25 Location 0 -OpDecorate %26 DescriptorSet 0 -OpDecorate %26 Binding 0 -OpDecorate %31 DescriptorSet 0 -OpDecorate %31 Binding 1 +OpDecorate %27 DescriptorSet 0 +OpDecorate %27 Binding 0 +OpDecorate %32 DescriptorSet 0 +OpDecorate %32 Binding 1 OpDecorate %35 Location 0 %3 = OpTypeFloat 32 %2 = OpConstant %3 1.2 @@ -28,41 +28,41 @@ OpDecorate %35 Location 0 %5 = OpConstant %3 1.0 %6 = OpTypeVoid %8 = OpTypeFunction %6 -%11 = OpTypeVector %3 2 -%12 = OpTypePointer Output %11 -%10 = OpVariable %12 Output -%14 = OpTypePointer Input %11 +%10 = OpTypeVector %3 2 +%12 = OpTypePointer Output %10 +%11 = OpVariable %12 Output +%14 = OpTypePointer Input %10 %13 = OpVariable %14 Input -%17 = OpTypeVector %3 4 -%18 = OpTypePointer Output %17 -%16 = OpVariable %18 Output +%16 = OpTypeVector %3 4 +%18 = OpTypePointer Output %16 +%17 = OpVariable %18 Output %20 = OpVariable %14 Input %25 = OpVariable %18 Output -%27 = OpTypeImage %3 2D 0 0 0 1 Unknown -%28 = OpTypePointer UniformConstant %27 -%26 = OpVariable %28 UniformConstant -%30 = OpTypeSampledImage %27 -%32 = OpTypeSampler -%33 = OpTypePointer UniformConstant %32 -%31 = OpVariable %33 UniformConstant +%26 = OpTypeImage %3 2D 0 0 0 1 Unknown +%28 = OpTypePointer UniformConstant %26 +%27 = OpVariable %28 UniformConstant +%30 = OpTypeSampledImage %26 +%31 = OpTypeSampler +%33 = OpTypePointer UniformConstant %31 +%32 = OpVariable %33 UniformConstant %35 = OpVariable %14 Input %7 = OpFunction %6 None %8 %9 = OpLabel -%15 = OpLoad %11 %13 -OpStore %10 %15 -%21 = OpLoad %11 %20 -%19 = OpVectorTimesScalar %11 %21 %2 -%22 = OpCompositeConstruct %17 %19 %4 %5 -OpStore %16 %22 +%15 = OpLoad %10 %13 +OpStore %11 %15 +%21 = OpLoad %10 %20 +%19 = OpVectorTimesScalar %10 %21 %2 +%22 = OpCompositeConstruct %16 %19 %4 %5 +OpStore %17 %22 OpReturn OpFunctionEnd %23 = OpFunction %6 None %8 %24 = OpLabel -%29 = OpLoad %27 %26 -%34 = OpLoad %32 %31 -%36 = OpLoad %11 %35 +%29 = OpLoad %26 %27 +%34 = OpLoad %31 %32 +%36 = OpLoad %10 %35 %37 = OpSampledImage %30 %29 %34 -%38 = OpImageSampleImplicitLod %17 %37 %36 +%38 = OpImageSampleImplicitLod %16 %37 %36 OpStore %25 %38 OpReturn OpFunctionEnd diff --git a/tests/snapshots/snapshots__shadow.msl.snap b/tests/snapshots/snapshots__shadow.msl.snap index b8ad9bf198..0d62314ac0 100644 --- a/tests/snapshots/snapshots__shadow.msl.snap +++ b/tests/snapshots/snapshots__shadow.msl.snap @@ -41,6 +41,8 @@ typedef int type9; typedef metal::float3 type10; +typedef bool type11; + type7 fetch_shadow( type6 light_id, type2 homogeneous_coords diff --git a/tests/snapshots/snapshots__shadow.spvasm.snap b/tests/snapshots/snapshots__shadow.spvasm.snap index c0f1f29a2f..8d13d12835 100644 --- a/tests/snapshots/snapshots__shadow.spvasm.snap +++ b/tests/snapshots/snapshots__shadow.spvasm.snap @@ -5,34 +5,34 @@ expression: dis ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 220 +; Bound: 221 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %62 "fs_main" %113 %110 %216 -OpExecutionMode %62 OriginUpperLeft -OpDecorate %26 DescriptorSet 0 -OpDecorate %26 Binding 2 -OpDecorate %31 DescriptorSet 0 -OpDecorate %31 Binding 3 -OpDecorate %74 Block -OpMemberDecorate %74 0 Offset 0 -OpDecorate %73 DescriptorSet 0 -OpDecorate %73 Binding 0 -OpDecorate %98 BufferBlock -OpMemberDecorate %98 0 Offset 0 -OpMemberDecorate %98 0 NonWritable +OpEntryPoint Fragment %64 "fs_main" %114 %111 %217 +OpExecutionMode %64 OriginUpperLeft +OpDecorate %27 DescriptorSet 0 +OpDecorate %27 Binding 2 +OpDecorate %32 DescriptorSet 0 +OpDecorate %32 Binding 3 +OpDecorate %76 Block +OpMemberDecorate %76 0 Offset 0 +OpDecorate %77 DescriptorSet 0 +OpDecorate %77 Binding 0 +OpMemberDecorate %97 0 Offset 0 +OpMemberDecorate %97 0 ColMajor +OpMemberDecorate %97 0 MatrixStride 16 +OpMemberDecorate %97 1 Offset 64 +OpMemberDecorate %97 2 Offset 80 OpDecorate %99 ArrayStride 96 -OpMemberDecorate %100 0 Offset 0 -OpMemberDecorate %100 0 ColMajor -OpMemberDecorate %100 0 MatrixStride 16 -OpMemberDecorate %100 1 Offset 64 -OpMemberDecorate %100 2 Offset 80 -OpDecorate %97 DescriptorSet 0 -OpDecorate %97 Binding 1 -OpDecorate %110 Location 1 -OpDecorate %113 Location 0 -OpDecorate %216 Location 0 +OpDecorate %101 BufferBlock +OpMemberDecorate %101 0 Offset 0 +OpMemberDecorate %101 0 NonWritable +OpDecorate %102 DescriptorSet 0 +OpDecorate %102 Binding 1 +OpDecorate %111 Location 1 +OpDecorate %114 Location 0 +OpDecorate %217 Location 0 %3 = OpTypeFloat 32 %2 = OpConstant %3 0.0 %4 = OpConstant %3 1.0 @@ -47,227 +47,228 @@ OpDecorate %216 Location 0 %13 = OpConstant %11 1 %16 = OpTypeVector %3 4 %18 = OpTypeFunction %3 %11 %16 -%22 = OpTypeBool -%27 = OpTypeImage %3 2D 1 1 0 1 Unknown -%28 = OpTypePointer UniformConstant %27 -%26 = OpVariable %28 UniformConstant -%30 = OpTypeSampledImage %27 -%32 = OpTypeSampler -%33 = OpTypePointer UniformConstant %32 -%31 = OpVariable %33 UniformConstant -%38 = OpTypeVector %3 2 -%56 = OpConstant %3 0.0 -%58 = OpTypePointer Function %9 -%60 = OpTypePointer Function %11 -%61 = OpTypeVoid -%63 = OpTypeFunction %61 -%75 = OpTypeVector %11 4 -%74 = OpTypeStruct %75 -%76 = OpTypePointer Uniform %74 -%73 = OpVariable %76 Uniform -%77 = OpTypeInt 32 1 -%78 = OpConstant %77 0 -%79 = OpTypePointer Uniform %75 -%80 = OpConstant %77 0 +%20 = OpTypeBool +%26 = OpTypeImage %3 2D 1 1 0 1 Unknown +%28 = OpTypePointer UniformConstant %26 +%27 = OpVariable %28 UniformConstant +%30 = OpTypeSampledImage %26 +%31 = OpTypeSampler +%33 = OpTypePointer UniformConstant %31 +%32 = OpVariable %33 UniformConstant +%35 = OpTypeVector %3 2 +%49 = OpTypeInt 32 1 +%58 = OpConstant %3 0.0 +%60 = OpTypePointer Function %9 +%62 = OpTypePointer Function %11 +%63 = OpTypeVoid +%65 = OpTypeFunction %63 +%74 = OpTypeVector %11 4 +%76 = OpTypeStruct %74 +%78 = OpTypePointer Uniform %76 +%77 = OpVariable %78 Uniform +%79 = OpTypePointer Uniform %74 +%80 = OpConstant %49 0 %81 = OpTypePointer Uniform %11 -%101 = OpTypeMatrix %16 4 -%100 = OpTypeStruct %101 %16 %16 -%99 = OpTypeRuntimeArray %100 -%98 = OpTypeStruct %99 -%102 = OpTypePointer Uniform %98 -%97 = OpVariable %102 Uniform -%103 = OpConstant %77 0 +%82 = OpConstant %49 0 +%95 = OpTypeMatrix %16 4 +%97 = OpTypeStruct %95 %16 %16 +%99 = OpTypeRuntimeArray %97 +%101 = OpTypeStruct %99 +%103 = OpTypePointer Uniform %101 +%102 = OpVariable %103 Uniform %104 = OpTypePointer Uniform %99 -%106 = OpTypePointer Uniform %100 -%107 = OpConstant %77 0 -%108 = OpTypePointer Uniform %101 -%111 = OpTypePointer Input %16 -%110 = OpVariable %111 Input -%114 = OpTypePointer Input %9 -%113 = OpVariable %114 Input -%122 = OpConstant %77 0 +%105 = OpConstant %49 0 +%107 = OpTypePointer Uniform %97 +%108 = OpTypePointer Uniform %95 +%109 = OpConstant %49 0 +%112 = OpTypePointer Input %16 +%111 = OpVariable %112 Input +%115 = OpTypePointer Input %9 +%114 = OpVariable %115 Input %123 = OpTypePointer Uniform %99 -%125 = OpTypePointer Uniform %100 -%126 = OpConstant %77 1 +%124 = OpConstant %49 0 +%126 = OpTypePointer Uniform %97 %127 = OpTypePointer Uniform %16 -%128 = OpConstant %77 0 +%128 = OpConstant %49 1 %129 = OpTypePointer Uniform %3 -%135 = OpConstant %77 0 +%130 = OpConstant %49 0 %136 = OpTypePointer Uniform %99 -%138 = OpTypePointer Uniform %100 -%139 = OpConstant %77 1 +%137 = OpConstant %49 0 +%139 = OpTypePointer Uniform %97 %140 = OpTypePointer Uniform %16 -%141 = OpConstant %77 1 +%141 = OpConstant %49 1 %142 = OpTypePointer Uniform %3 -%148 = OpConstant %77 0 +%143 = OpConstant %49 1 %149 = OpTypePointer Uniform %99 -%151 = OpTypePointer Uniform %100 -%152 = OpConstant %77 1 +%150 = OpConstant %49 0 +%152 = OpTypePointer Uniform %97 %153 = OpTypePointer Uniform %16 -%154 = OpConstant %77 2 +%154 = OpConstant %49 1 %155 = OpTypePointer Uniform %3 -%159 = OpConstant %77 0 +%156 = OpConstant %49 2 %160 = OpTypePointer Input %3 -%163 = OpConstant %77 1 +%161 = OpConstant %49 0 %164 = OpTypePointer Input %3 -%167 = OpConstant %77 2 +%165 = OpConstant %49 1 %168 = OpTypePointer Input %3 -%178 = OpConstant %77 0 +%169 = OpConstant %49 2 %179 = OpTypePointer Uniform %99 -%181 = OpTypePointer Uniform %100 -%182 = OpConstant %77 2 +%180 = OpConstant %49 0 +%182 = OpTypePointer Uniform %97 %183 = OpTypePointer Uniform %16 -%184 = OpConstant %77 0 +%184 = OpConstant %49 2 %185 = OpTypePointer Uniform %3 -%191 = OpConstant %77 0 +%186 = OpConstant %49 0 %192 = OpTypePointer Uniform %99 -%194 = OpTypePointer Uniform %100 -%195 = OpConstant %77 2 +%193 = OpConstant %49 0 +%195 = OpTypePointer Uniform %97 %196 = OpTypePointer Uniform %16 -%197 = OpConstant %77 1 +%197 = OpConstant %49 2 %198 = OpTypePointer Uniform %3 -%204 = OpConstant %77 0 +%199 = OpConstant %49 1 %205 = OpTypePointer Uniform %99 -%207 = OpTypePointer Uniform %100 -%208 = OpConstant %77 2 +%206 = OpConstant %49 0 +%208 = OpTypePointer Uniform %97 %209 = OpTypePointer Uniform %16 -%210 = OpConstant %77 2 +%210 = OpConstant %49 2 %211 = OpTypePointer Uniform %3 -%217 = OpTypePointer Output %16 -%216 = OpVariable %217 Output +%212 = OpConstant %49 2 +%218 = OpTypePointer Output %16 +%217 = OpVariable %218 Output %17 = OpFunction %3 None %18 %14 = OpFunctionParameter %11 %15 = OpFunctionParameter %16 %19 = OpLabel -%21 = OpCompositeExtract %3 %15 3 -%20 = OpFOrdLessThanEqual %22 %21 %2 +%22 = OpCompositeExtract %3 %15 3 +%21 = OpFOrdLessThanEqual %20 %22 %2 OpSelectionMerge %23 None -OpBranchConditional %20 %24 %25 +OpBranchConditional %21 %24 %25 %24 = OpLabel OpReturnValue %4 %25 = OpLabel OpBranch %23 %23 = OpLabel -%29 = OpLoad %27 %26 -%34 = OpLoad %32 %31 +%29 = OpLoad %26 %27 +%34 = OpLoad %31 %32 %39 = OpCompositeExtract %3 %15 0 %40 = OpCompositeExtract %3 %15 1 -%41 = OpCompositeConstruct %38 %39 %40 -%42 = OpCompositeConstruct %38 %5 %6 -%37 = OpFMul %38 %41 %42 +%41 = OpCompositeConstruct %35 %39 %40 +%42 = OpCompositeConstruct %35 %5 %6 +%38 = OpFMul %35 %41 %42 %44 = OpCompositeExtract %3 %15 3 %43 = OpFDiv %3 %4 %44 -%36 = OpVectorTimesScalar %38 %37 %43 -%45 = OpCompositeConstruct %38 %5 %5 -%35 = OpFAdd %38 %36 %45 -%46 = OpCompositeExtract %3 %35 0 -%47 = OpCompositeExtract %3 %35 1 -%48 = OpConvertUToF %3 %14 -%49 = OpCompositeConstruct %9 %46 %47 %48 -%50 = OpSampledImage %30 %29 %34 -%53 = OpCompositeExtract %3 %15 2 -%55 = OpCompositeExtract %3 %15 3 -%54 = OpFDiv %3 %4 %55 -%52 = OpFMul %3 %53 %54 -%51 = OpImageSampleDrefExplicitLod %3 %50 %49 %52 Lod %56 -OpReturnValue %51 +%37 = OpVectorTimesScalar %35 %38 %43 +%45 = OpCompositeConstruct %35 %5 %5 +%36 = OpFAdd %35 %37 %45 +%46 = OpCompositeExtract %3 %36 0 +%47 = OpCompositeExtract %3 %36 1 +%50 = OpBitcast %49 %14 +%48 = OpConvertUToF %3 %50 +%51 = OpCompositeConstruct %9 %46 %47 %48 +%52 = OpSampledImage %30 %29 %34 +%55 = OpCompositeExtract %3 %15 2 +%57 = OpCompositeExtract %3 %15 3 +%56 = OpFDiv %3 %4 %57 +%54 = OpFMul %3 %55 %56 +%53 = OpImageSampleDrefExplicitLod %3 %52 %51 %54 Lod %58 +OpReturnValue %53 OpFunctionEnd -%62 = OpFunction %61 None %63 -%64 = OpLabel -%57 = OpVariable %58 Function %8 -%59 = OpVariable %60 Function %12 -OpBranch %65 -%65 = OpLabel -OpLoopMerge %66 %68 None +%64 = OpFunction %63 None %65 +%66 = OpLabel +%59 = OpVariable %60 Function %8 +%61 = OpVariable %62 Function %12 OpBranch %67 %67 = OpLabel -%70 = OpLoad %11 %59 -%72 = OpAccessChain %79 %73 %78 -%71 = OpAccessChain %81 %72 %80 -%82 = OpLoad %11 %71 -%83 = OpExtInst %11 %1 UMin %82 %10 -%69 = OpUGreaterThanEqual %22 %70 %83 -OpSelectionMerge %84 None -OpBranchConditional %69 %85 %86 -%85 = OpLabel -OpBranch %66 +OpLoopMerge %68 %70 None +OpBranch %69 +%69 = OpLabel +%72 = OpLoad %11 %61 +%75 = OpAccessChain %79 %77 %80 +%73 = OpAccessChain %81 %75 %82 +%83 = OpLoad %11 %73 +%84 = OpExtInst %11 %1 UMin %83 %10 +%71 = OpUGreaterThanEqual %20 %72 %84 +OpSelectionMerge %85 None +OpBranchConditional %71 %86 %87 %86 = OpLabel -OpBranch %84 -%84 = OpLabel -%88 = OpLoad %9 %57 -%92 = OpLoad %11 %59 -%96 = OpAccessChain %104 %97 %103 -%105 = OpLoad %11 %59 -%95 = OpAccessChain %106 %96 %105 -%94 = OpAccessChain %108 %95 %107 -%109 = OpLoad %101 %94 -%112 = OpLoad %16 %110 -%93 = OpMatrixTimesVector %16 %109 %112 -%91 = OpFunctionCall %3 %17 %92 %93 -%115 = OpLoad %9 %113 -%116 = OpExtInst %9 %1 Normalize %115 -%121 = OpAccessChain %123 %97 %122 -%124 = OpLoad %11 %59 -%120 = OpAccessChain %125 %121 %124 -%119 = OpAccessChain %127 %120 %126 -%118 = OpAccessChain %129 %119 %128 -%130 = OpLoad %3 %118 -%134 = OpAccessChain %136 %97 %135 -%137 = OpLoad %11 %59 -%133 = OpAccessChain %138 %134 %137 -%132 = OpAccessChain %140 %133 %139 -%131 = OpAccessChain %142 %132 %141 -%143 = OpLoad %3 %131 -%147 = OpAccessChain %149 %97 %148 -%150 = OpLoad %11 %59 -%146 = OpAccessChain %151 %147 %150 -%145 = OpAccessChain %153 %146 %152 -%144 = OpAccessChain %155 %145 %154 -%156 = OpLoad %3 %144 -%157 = OpCompositeConstruct %9 %130 %143 %156 -%158 = OpAccessChain %160 %110 %159 -%161 = OpLoad %3 %158 -%162 = OpAccessChain %164 %110 %163 -%165 = OpLoad %3 %162 -%166 = OpAccessChain %168 %110 %167 -%169 = OpLoad %3 %166 -%170 = OpCompositeConstruct %9 %161 %165 %169 -%117 = OpFSub %9 %157 %170 -%171 = OpExtInst %9 %1 Normalize %117 -%172 = OpDot %3 %116 %171 -%173 = OpExtInst %3 %1 FMax %2 %172 -%90 = OpFMul %3 %91 %173 -%177 = OpAccessChain %179 %97 %178 -%180 = OpLoad %11 %59 -%176 = OpAccessChain %181 %177 %180 -%175 = OpAccessChain %183 %176 %182 -%174 = OpAccessChain %185 %175 %184 -%186 = OpLoad %3 %174 -%190 = OpAccessChain %192 %97 %191 -%193 = OpLoad %11 %59 -%189 = OpAccessChain %194 %190 %193 -%188 = OpAccessChain %196 %189 %195 -%187 = OpAccessChain %198 %188 %197 -%199 = OpLoad %3 %187 -%203 = OpAccessChain %205 %97 %204 -%206 = OpLoad %11 %59 -%202 = OpAccessChain %207 %203 %206 -%201 = OpAccessChain %209 %202 %208 -%200 = OpAccessChain %211 %201 %210 -%212 = OpLoad %3 %200 -%213 = OpCompositeConstruct %9 %186 %199 %212 -%89 = OpVectorTimesScalar %9 %213 %90 -%87 = OpFAdd %9 %88 %89 -OpStore %57 %87 OpBranch %68 +%87 = OpLabel +OpBranch %85 +%85 = OpLabel +%89 = OpLoad %9 %59 +%93 = OpLoad %11 %61 +%100 = OpAccessChain %104 %102 %105 +%106 = OpLoad %11 %61 +%98 = OpAccessChain %107 %100 %106 +%96 = OpAccessChain %108 %98 %109 +%110 = OpLoad %95 %96 +%113 = OpLoad %16 %111 +%94 = OpMatrixTimesVector %16 %110 %113 +%92 = OpFunctionCall %3 %17 %93 %94 +%116 = OpLoad %9 %114 +%117 = OpExtInst %9 %1 Normalize %116 +%122 = OpAccessChain %123 %102 %124 +%125 = OpLoad %11 %61 +%121 = OpAccessChain %126 %122 %125 +%120 = OpAccessChain %127 %121 %128 +%119 = OpAccessChain %129 %120 %130 +%131 = OpLoad %3 %119 +%135 = OpAccessChain %136 %102 %137 +%138 = OpLoad %11 %61 +%134 = OpAccessChain %139 %135 %138 +%133 = OpAccessChain %140 %134 %141 +%132 = OpAccessChain %142 %133 %143 +%144 = OpLoad %3 %132 +%148 = OpAccessChain %149 %102 %150 +%151 = OpLoad %11 %61 +%147 = OpAccessChain %152 %148 %151 +%146 = OpAccessChain %153 %147 %154 +%145 = OpAccessChain %155 %146 %156 +%157 = OpLoad %3 %145 +%158 = OpCompositeConstruct %9 %131 %144 %157 +%159 = OpAccessChain %160 %111 %161 +%162 = OpLoad %3 %159 +%163 = OpAccessChain %164 %111 %165 +%166 = OpLoad %3 %163 +%167 = OpAccessChain %168 %111 %169 +%170 = OpLoad %3 %167 +%171 = OpCompositeConstruct %9 %162 %166 %170 +%118 = OpFSub %9 %158 %171 +%172 = OpExtInst %9 %1 Normalize %118 +%173 = OpDot %3 %117 %172 +%174 = OpExtInst %3 %1 FMax %2 %173 +%91 = OpFMul %3 %92 %174 +%178 = OpAccessChain %179 %102 %180 +%181 = OpLoad %11 %61 +%177 = OpAccessChain %182 %178 %181 +%176 = OpAccessChain %183 %177 %184 +%175 = OpAccessChain %185 %176 %186 +%187 = OpLoad %3 %175 +%191 = OpAccessChain %192 %102 %193 +%194 = OpLoad %11 %61 +%190 = OpAccessChain %195 %191 %194 +%189 = OpAccessChain %196 %190 %197 +%188 = OpAccessChain %198 %189 %199 +%200 = OpLoad %3 %188 +%204 = OpAccessChain %205 %102 %206 +%207 = OpLoad %11 %61 +%203 = OpAccessChain %208 %204 %207 +%202 = OpAccessChain %209 %203 %210 +%201 = OpAccessChain %211 %202 %212 +%213 = OpLoad %3 %201 +%214 = OpCompositeConstruct %9 %187 %200 %213 +%90 = OpVectorTimesScalar %9 %214 %91 +%88 = OpFAdd %9 %89 %90 +OpStore %59 %88 +OpBranch %70 +%70 = OpLabel +%216 = OpLoad %11 %61 +%215 = OpIAdd %11 %216 %13 +OpStore %61 %215 +OpBranch %67 %68 = OpLabel -%215 = OpLoad %11 %59 -%214 = OpIAdd %11 %215 %13 -OpStore %59 %214 -OpBranch %65 -%66 = OpLabel -%218 = OpLoad %9 %57 -%219 = OpCompositeConstruct %16 %218 %4 -OpStore %216 %219 +%219 = OpLoad %9 %59 +%220 = OpCompositeConstruct %16 %219 %4 +OpStore %217 %220 OpReturn OpFunctionEnd diff --git a/tests/snapshots/snapshots__skybox.spvasm.snap b/tests/snapshots/snapshots__skybox.spvasm.snap index 6c23904ab4..d0577615be 100644 --- a/tests/snapshots/snapshots__skybox.spvasm.snap +++ b/tests/snapshots/snapshots__skybox.spvasm.snap @@ -5,31 +5,31 @@ expression: dis ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 185 +; Bound: 187 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %16 "vs_main" %157 %44 %20 -OpEntryPoint Fragment %168 "fs_main" %180 %170 -OpExecutionMode %168 OriginUpperLeft -OpDecorate %20 BuiltIn VertexIndex -OpDecorate %29 Block -OpMemberDecorate %29 0 Offset 0 -OpMemberDecorate %29 0 ColMajor -OpMemberDecorate %29 0 MatrixStride 16 -OpMemberDecorate %29 1 Offset 64 -OpMemberDecorate %29 1 ColMajor -OpMemberDecorate %29 1 MatrixStride 16 -OpDecorate %28 DescriptorSet 0 -OpDecorate %28 Binding 0 -OpDecorate %44 Location 0 -OpDecorate %157 BuiltIn Position -OpDecorate %170 Location 0 -OpDecorate %171 DescriptorSet 0 -OpDecorate %171 Binding 1 -OpDecorate %176 DescriptorSet 0 -OpDecorate %176 Binding 2 -OpDecorate %180 Location 0 +OpEntryPoint Vertex %16 "vs_main" %159 %47 %21 +OpEntryPoint Fragment %170 "fs_main" %182 %172 +OpExecutionMode %170 OriginUpperLeft +OpDecorate %21 BuiltIn VertexIndex +OpDecorate %31 Block +OpMemberDecorate %31 0 Offset 0 +OpMemberDecorate %31 0 ColMajor +OpMemberDecorate %31 0 MatrixStride 16 +OpMemberDecorate %31 1 Offset 64 +OpMemberDecorate %31 1 ColMajor +OpMemberDecorate %31 1 MatrixStride 16 +OpDecorate %32 DescriptorSet 0 +OpDecorate %32 Binding 0 +OpDecorate %47 Location 0 +OpDecorate %159 BuiltIn Position +OpDecorate %172 Location 0 +OpDecorate %174 DescriptorSet 0 +OpDecorate %174 Binding 1 +OpDecorate %179 DescriptorSet 0 +OpDecorate %179 Binding 2 +OpDecorate %182 Location 0 %3 = OpTypeInt 32 1 %2 = OpConstant %3 2 %4 = OpConstant %3 1 @@ -42,184 +42,186 @@ OpDecorate %180 Location 0 %14 = OpTypePointer Function %13 %15 = OpTypeVoid %17 = OpTypeFunction %15 -%21 = OpTypeInt 32 0 -%22 = OpTypePointer Input %21 -%20 = OpVariable %22 Input -%30 = OpTypeMatrix %13 4 -%29 = OpTypeStruct %30 %30 -%31 = OpTypePointer Uniform %29 -%28 = OpVariable %31 Uniform -%32 = OpConstant %3 0 -%33 = OpTypePointer Uniform %30 -%45 = OpTypeVector %6 3 -%46 = OpTypePointer Output %45 -%44 = OpVariable %46 Output -%48 = OpTypeMatrix %45 3 -%52 = OpConstant %3 1 -%53 = OpTypePointer Uniform %30 -%54 = OpConstant %3 0 -%55 = OpTypePointer Uniform %13 -%56 = OpConstant %3 0 -%57 = OpTypePointer Uniform %6 -%62 = OpConstant %3 1 -%63 = OpTypePointer Uniform %30 -%64 = OpConstant %3 0 -%65 = OpTypePointer Uniform %13 -%66 = OpConstant %3 1 -%67 = OpTypePointer Uniform %6 -%72 = OpConstant %3 1 -%73 = OpTypePointer Uniform %30 -%74 = OpConstant %3 0 -%75 = OpTypePointer Uniform %13 -%76 = OpConstant %3 2 -%77 = OpTypePointer Uniform %6 -%83 = OpConstant %3 1 -%84 = OpTypePointer Uniform %30 -%85 = OpConstant %3 1 -%86 = OpTypePointer Uniform %13 -%87 = OpConstant %3 0 -%88 = OpTypePointer Uniform %6 -%93 = OpConstant %3 1 -%94 = OpTypePointer Uniform %30 -%95 = OpConstant %3 1 -%96 = OpTypePointer Uniform %13 -%97 = OpConstant %3 1 -%98 = OpTypePointer Uniform %6 -%103 = OpConstant %3 1 -%104 = OpTypePointer Uniform %30 -%105 = OpConstant %3 1 -%106 = OpTypePointer Uniform %13 -%107 = OpConstant %3 2 -%108 = OpTypePointer Uniform %6 -%114 = OpConstant %3 1 -%115 = OpTypePointer Uniform %30 -%116 = OpConstant %3 2 -%117 = OpTypePointer Uniform %13 -%118 = OpConstant %3 0 -%119 = OpTypePointer Uniform %6 -%124 = OpConstant %3 1 -%125 = OpTypePointer Uniform %30 -%126 = OpConstant %3 2 -%127 = OpTypePointer Uniform %13 -%128 = OpConstant %3 1 -%129 = OpTypePointer Uniform %6 -%134 = OpConstant %3 1 -%135 = OpTypePointer Uniform %30 -%136 = OpConstant %3 2 -%137 = OpTypePointer Uniform %13 -%138 = OpConstant %3 2 -%139 = OpTypePointer Uniform %6 -%145 = OpConstant %3 0 -%146 = OpTypePointer Function %6 -%149 = OpConstant %3 1 -%150 = OpTypePointer Function %6 -%153 = OpConstant %3 2 -%154 = OpTypePointer Function %6 -%158 = OpTypePointer Output %13 -%157 = OpVariable %158 Output -%170 = OpVariable %158 Output -%172 = OpTypeImage %6 Cube 0 0 0 1 Unknown -%173 = OpTypePointer UniformConstant %172 -%171 = OpVariable %173 UniformConstant -%175 = OpTypeSampledImage %172 -%177 = OpTypeSampler -%178 = OpTypePointer UniformConstant %177 -%176 = OpVariable %178 UniformConstant -%181 = OpTypePointer Input %45 -%180 = OpVariable %181 Input +%20 = OpTypeInt 32 0 +%22 = OpTypePointer Input %20 +%21 = OpVariable %22 Input +%29 = OpTypeMatrix %13 4 +%31 = OpTypeStruct %29 %29 +%33 = OpTypePointer Uniform %31 +%32 = OpVariable %33 Uniform +%34 = OpTypePointer Uniform %29 +%35 = OpConstant %3 0 +%46 = OpTypeVector %6 3 +%48 = OpTypePointer Output %46 +%47 = OpVariable %48 Output +%50 = OpTypeMatrix %46 3 +%54 = OpTypePointer Uniform %29 +%55 = OpConstant %3 1 +%56 = OpTypePointer Uniform %13 +%57 = OpConstant %3 0 +%58 = OpTypePointer Uniform %6 +%59 = OpConstant %3 0 +%64 = OpTypePointer Uniform %29 +%65 = OpConstant %3 1 +%66 = OpTypePointer Uniform %13 +%67 = OpConstant %3 0 +%68 = OpTypePointer Uniform %6 +%69 = OpConstant %3 1 +%74 = OpTypePointer Uniform %29 +%75 = OpConstant %3 1 +%76 = OpTypePointer Uniform %13 +%77 = OpConstant %3 0 +%78 = OpTypePointer Uniform %6 +%79 = OpConstant %3 2 +%85 = OpTypePointer Uniform %29 +%86 = OpConstant %3 1 +%87 = OpTypePointer Uniform %13 +%88 = OpConstant %3 1 +%89 = OpTypePointer Uniform %6 +%90 = OpConstant %3 0 +%95 = OpTypePointer Uniform %29 +%96 = OpConstant %3 1 +%97 = OpTypePointer Uniform %13 +%98 = OpConstant %3 1 +%99 = OpTypePointer Uniform %6 +%100 = OpConstant %3 1 +%105 = OpTypePointer Uniform %29 +%106 = OpConstant %3 1 +%107 = OpTypePointer Uniform %13 +%108 = OpConstant %3 1 +%109 = OpTypePointer Uniform %6 +%110 = OpConstant %3 2 +%116 = OpTypePointer Uniform %29 +%117 = OpConstant %3 1 +%118 = OpTypePointer Uniform %13 +%119 = OpConstant %3 2 +%120 = OpTypePointer Uniform %6 +%121 = OpConstant %3 0 +%126 = OpTypePointer Uniform %29 +%127 = OpConstant %3 1 +%128 = OpTypePointer Uniform %13 +%129 = OpConstant %3 2 +%130 = OpTypePointer Uniform %6 +%131 = OpConstant %3 1 +%136 = OpTypePointer Uniform %29 +%137 = OpConstant %3 1 +%138 = OpTypePointer Uniform %13 +%139 = OpConstant %3 2 +%140 = OpTypePointer Uniform %6 +%141 = OpConstant %3 2 +%147 = OpTypePointer Function %6 +%148 = OpConstant %3 0 +%151 = OpTypePointer Function %6 +%152 = OpConstant %3 1 +%155 = OpTypePointer Function %6 +%156 = OpConstant %3 2 +%160 = OpTypePointer Output %13 +%159 = OpVariable %160 Output +%172 = OpVariable %160 Output +%173 = OpTypeImage %6 Cube 0 0 0 1 Unknown +%175 = OpTypePointer UniformConstant %173 +%174 = OpVariable %175 UniformConstant +%177 = OpTypeSampledImage %173 +%178 = OpTypeSampler +%180 = OpTypePointer UniformConstant %178 +%179 = OpVariable %180 UniformConstant +%183 = OpTypePointer Input %46 +%182 = OpVariable %183 Input %16 = OpFunction %15 None %17 %18 = OpLabel %9 = OpVariable %10 Function %11 = OpVariable %10 Function %12 = OpVariable %14 Function -%23 = OpLoad %21 %20 -%19 = OpSDiv %3 %23 %2 +%23 = OpLoad %20 %21 +%24 = OpBitcast %3 %23 +%19 = OpSDiv %3 %24 %2 OpStore %9 %19 -%25 = OpLoad %21 %20 -%24 = OpBitwiseAnd %3 %25 %4 -OpStore %11 %24 -%27 = OpAccessChain %33 %28 %32 -%34 = OpLoad %30 %27 -%37 = OpLoad %3 %9 -%38 = OpConvertSToF %6 %37 -%36 = OpFMul %6 %38 %5 -%35 = OpFSub %6 %36 %7 -%41 = OpLoad %3 %11 -%42 = OpConvertSToF %6 %41 -%40 = OpFMul %6 %42 %5 -%39 = OpFSub %6 %40 %7 -%43 = OpCompositeConstruct %13 %35 %39 %8 %7 -%26 = OpMatrixTimesVector %13 %34 %43 -OpStore %12 %26 -%51 = OpAccessChain %53 %28 %52 -%50 = OpAccessChain %55 %51 %54 -%49 = OpAccessChain %57 %50 %56 -%58 = OpLoad %6 %49 -%61 = OpAccessChain %63 %28 %62 -%60 = OpAccessChain %65 %61 %64 -%59 = OpAccessChain %67 %60 %66 -%68 = OpLoad %6 %59 -%71 = OpAccessChain %73 %28 %72 -%70 = OpAccessChain %75 %71 %74 -%69 = OpAccessChain %77 %70 %76 -%78 = OpLoad %6 %69 -%79 = OpCompositeConstruct %45 %58 %68 %78 -%82 = OpAccessChain %84 %28 %83 -%81 = OpAccessChain %86 %82 %85 -%80 = OpAccessChain %88 %81 %87 -%89 = OpLoad %6 %80 -%92 = OpAccessChain %94 %28 %93 -%91 = OpAccessChain %96 %92 %95 -%90 = OpAccessChain %98 %91 %97 -%99 = OpLoad %6 %90 -%102 = OpAccessChain %104 %28 %103 -%101 = OpAccessChain %106 %102 %105 -%100 = OpAccessChain %108 %101 %107 -%109 = OpLoad %6 %100 -%110 = OpCompositeConstruct %45 %89 %99 %109 -%113 = OpAccessChain %115 %28 %114 -%112 = OpAccessChain %117 %113 %116 -%111 = OpAccessChain %119 %112 %118 -%120 = OpLoad %6 %111 -%123 = OpAccessChain %125 %28 %124 -%122 = OpAccessChain %127 %123 %126 -%121 = OpAccessChain %129 %122 %128 -%130 = OpLoad %6 %121 -%133 = OpAccessChain %135 %28 %134 -%132 = OpAccessChain %137 %133 %136 -%131 = OpAccessChain %139 %132 %138 -%140 = OpLoad %6 %131 -%141 = OpCompositeConstruct %45 %120 %130 %140 -%142 = OpCompositeConstruct %48 %79 %110 %141 -%143 = OpTranspose %48 %142 -%144 = OpAccessChain %146 %12 %145 -%147 = OpLoad %6 %144 -%148 = OpAccessChain %150 %12 %149 -%151 = OpLoad %6 %148 -%152 = OpAccessChain %154 %12 %153 -%155 = OpLoad %6 %152 -%156 = OpCompositeConstruct %45 %147 %151 %155 -%47 = OpMatrixTimesVector %45 %143 %156 -OpStore %44 %47 -%161 = OpLoad %3 %9 -%162 = OpConvertSToF %6 %161 -%160 = OpFMul %6 %162 %5 -%159 = OpFSub %6 %160 %7 -%165 = OpLoad %3 %11 -%166 = OpConvertSToF %6 %165 -%164 = OpFMul %6 %166 %5 -%163 = OpFSub %6 %164 %7 -%167 = OpCompositeConstruct %13 %159 %163 %8 %7 -OpStore %157 %167 +%26 = OpLoad %20 %21 +%27 = OpBitcast %3 %26 +%25 = OpBitwiseAnd %3 %27 %4 +OpStore %11 %25 +%30 = OpAccessChain %34 %32 %35 +%36 = OpLoad %29 %30 +%39 = OpLoad %3 %9 +%40 = OpConvertSToF %6 %39 +%38 = OpFMul %6 %40 %5 +%37 = OpFSub %6 %38 %7 +%43 = OpLoad %3 %11 +%44 = OpConvertSToF %6 %43 +%42 = OpFMul %6 %44 %5 +%41 = OpFSub %6 %42 %7 +%45 = OpCompositeConstruct %13 %37 %41 %8 %7 +%28 = OpMatrixTimesVector %13 %36 %45 +OpStore %12 %28 +%53 = OpAccessChain %54 %32 %55 +%52 = OpAccessChain %56 %53 %57 +%51 = OpAccessChain %58 %52 %59 +%60 = OpLoad %6 %51 +%63 = OpAccessChain %64 %32 %65 +%62 = OpAccessChain %66 %63 %67 +%61 = OpAccessChain %68 %62 %69 +%70 = OpLoad %6 %61 +%73 = OpAccessChain %74 %32 %75 +%72 = OpAccessChain %76 %73 %77 +%71 = OpAccessChain %78 %72 %79 +%80 = OpLoad %6 %71 +%81 = OpCompositeConstruct %46 %60 %70 %80 +%84 = OpAccessChain %85 %32 %86 +%83 = OpAccessChain %87 %84 %88 +%82 = OpAccessChain %89 %83 %90 +%91 = OpLoad %6 %82 +%94 = OpAccessChain %95 %32 %96 +%93 = OpAccessChain %97 %94 %98 +%92 = OpAccessChain %99 %93 %100 +%101 = OpLoad %6 %92 +%104 = OpAccessChain %105 %32 %106 +%103 = OpAccessChain %107 %104 %108 +%102 = OpAccessChain %109 %103 %110 +%111 = OpLoad %6 %102 +%112 = OpCompositeConstruct %46 %91 %101 %111 +%115 = OpAccessChain %116 %32 %117 +%114 = OpAccessChain %118 %115 %119 +%113 = OpAccessChain %120 %114 %121 +%122 = OpLoad %6 %113 +%125 = OpAccessChain %126 %32 %127 +%124 = OpAccessChain %128 %125 %129 +%123 = OpAccessChain %130 %124 %131 +%132 = OpLoad %6 %123 +%135 = OpAccessChain %136 %32 %137 +%134 = OpAccessChain %138 %135 %139 +%133 = OpAccessChain %140 %134 %141 +%142 = OpLoad %6 %133 +%143 = OpCompositeConstruct %46 %122 %132 %142 +%144 = OpCompositeConstruct %50 %81 %112 %143 +%145 = OpTranspose %50 %144 +%146 = OpAccessChain %147 %12 %148 +%149 = OpLoad %6 %146 +%150 = OpAccessChain %151 %12 %152 +%153 = OpLoad %6 %150 +%154 = OpAccessChain %155 %12 %156 +%157 = OpLoad %6 %154 +%158 = OpCompositeConstruct %46 %149 %153 %157 +%49 = OpMatrixTimesVector %46 %145 %158 +OpStore %47 %49 +%163 = OpLoad %3 %9 +%164 = OpConvertSToF %6 %163 +%162 = OpFMul %6 %164 %5 +%161 = OpFSub %6 %162 %7 +%167 = OpLoad %3 %11 +%168 = OpConvertSToF %6 %167 +%166 = OpFMul %6 %168 %5 +%165 = OpFSub %6 %166 %7 +%169 = OpCompositeConstruct %13 %161 %165 %8 %7 +OpStore %159 %169 OpReturn OpFunctionEnd -%168 = OpFunction %15 None %17 -%169 = OpLabel -%174 = OpLoad %172 %171 -%179 = OpLoad %177 %176 -%182 = OpLoad %45 %180 -%183 = OpSampledImage %175 %174 %179 -%184 = OpImageSampleImplicitLod %13 %183 %182 -OpStore %170 %184 +%170 = OpFunction %15 None %17 +%171 = OpLabel +%176 = OpLoad %173 %174 +%181 = OpLoad %178 %179 +%184 = OpLoad %46 %182 +%185 = OpSampledImage %177 %176 %181 +%186 = OpImageSampleImplicitLod %13 %185 %184 +OpStore %172 %186 OpReturn OpFunctionEnd