From be00a05fa5feee1527569f5b1da143c0f15a7be6 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Mon, 3 May 2021 13:47:51 -0400 Subject: [PATCH] [spv-in] default output builtins --- src/back/msl/writer.rs | 1 + src/front/spv/error.rs | 2 + src/front/spv/function.rs | 55 +++++++------- src/front/spv/mod.rs | 138 +++++++++++++++-------------------- src/front/spv/null.rs | 149 ++++++++++++++++++++++++++++++++++++++ tests/out/quad-vert.msl | 5 +- 6 files changed, 241 insertions(+), 109 deletions(-) create mode 100644 src/front/spv/null.rs diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index d72cd5c61c..e08f0d61dc 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -2027,6 +2027,7 @@ impl Writer { result.binding.as_ref(), )); } + writeln!(self.out, "struct {} {{", stage_out_name)?; for (name, ty, binding) in result_members { let ty_name = TypeContext { diff --git a/src/front/spv/error.rs b/src/front/spv/error.rs index 3a4bc44bee..3f0ae95592 100644 --- a/src/front/spv/error.rs +++ b/src/front/spv/error.rs @@ -99,6 +99,8 @@ pub enum Error { InvalidEdgeClassification, #[error("recursive function call %{0}")] FunctionCallCycle(spirv::Word), + #[error("invalid array size {0:?}")] + InvalidArraySize(Handle), #[error("invalid barrier scope %{0}")] InvalidBarrierScope(spirv::Word), #[error("invalid barrier memory semantics %{0}")] diff --git a/src/front/spv/function.rs b/src/front/spv/function.rs index 856ffc076b..a01ed51d63 100644 --- a/src/front/spv/function.rs +++ b/src/front/spv/function.rs @@ -296,33 +296,34 @@ impl> super::Parser { } } - if self.options.adjust_coordinate_space { - let position_index = members.iter().position(|member| match member.binding { - Some(crate::Binding::BuiltIn(crate::BuiltIn::Position)) => true, - _ => false, - }); - if let Some(component_index) = position_index { - // The IR is Y-up, while SPIR-V is Y-down. - let old_len = function.expressions.len(); - let global_expr = components[component_index]; - let access_expr = function.expressions.append(crate::Expression::AccessIndex { - base: global_expr, - index: 1, - }); - let load_expr = function.expressions.append(crate::Expression::Load { - pointer: access_expr, - }); - let neg_expr = function.expressions.append(crate::Expression::Unary { - op: crate::UnaryOperator::Negate, - expr: load_expr, - }); - function.body.push(crate::Statement::Emit( - function.expressions.range_from(old_len), - )); - function.body.push(crate::Statement::Store { - pointer: access_expr, - value: neg_expr, - }); + for (member_index, member) in members.iter().enumerate() { + match member.binding { + Some(crate::Binding::BuiltIn(crate::BuiltIn::Position)) + if self.options.adjust_coordinate_space => + { + let old_len = function.expressions.len(); + let global_expr = components[member_index]; + let access_expr = + function.expressions.append(crate::Expression::AccessIndex { + base: global_expr, + index: 1, + }); + let load_expr = function.expressions.append(crate::Expression::Load { + pointer: access_expr, + }); + let neg_expr = function.expressions.append(crate::Expression::Unary { + op: crate::UnaryOperator::Negate, + expr: load_expr, + }); + function.body.push(crate::Statement::Emit( + function.expressions.range_from(old_len), + )); + function.body.push(crate::Statement::Store { + pointer: access_expr, + value: neg_expr, + }); + } + _ => {} } } diff --git a/src/front/spv/mod.rs b/src/front/spv/mod.rs index 5603a53c9c..3750bd2383 100644 --- a/src/front/spv/mod.rs +++ b/src/front/spv/mod.rs @@ -31,6 +31,7 @@ mod error; mod flow; mod function; mod image; +mod null; use convert::*; pub use error::Error; @@ -3217,84 +3218,6 @@ impl> Parser { Ok(()) } - fn generate_null_constant( - &mut self, - constants: &mut Arena, - types: &mut Arena, - ty: Handle, - ) -> Result { - fn make_scalar_inner(kind: crate::ScalarKind, width: crate::Bytes) -> crate::ConstantInner { - crate::ConstantInner::Scalar { - width, - value: match kind { - crate::ScalarKind::Uint => crate::ScalarValue::Uint(0), - crate::ScalarKind::Sint => crate::ScalarValue::Sint(0), - crate::ScalarKind::Float => crate::ScalarValue::Float(0.0), - crate::ScalarKind::Bool => crate::ScalarValue::Bool(false), - }, - } - } - - let inner = match types[ty].inner { - crate::TypeInner::Scalar { kind, width } => make_scalar_inner(kind, width), - crate::TypeInner::Vector { size, kind, width } => { - let mut components = Vec::with_capacity(size as usize); - for _ in 0..size as usize { - components.push(constants.fetch_or_append(crate::Constant { - name: None, - specialization: None, - inner: make_scalar_inner(kind, width), - })); - } - crate::ConstantInner::Composite { ty, components } - } - crate::TypeInner::Matrix { - columns, - rows, - width, - } => { - let vector_ty = types.fetch_or_append(crate::Type { - name: None, - inner: crate::TypeInner::Vector { - kind: crate::ScalarKind::Float, - size: rows, - width, - }, - }); - let vector_inner = self.generate_null_constant(constants, types, vector_ty)?; - let vector_handle = constants.fetch_or_append(crate::Constant { - name: None, - specialization: None, - inner: vector_inner, - }); - crate::ConstantInner::Composite { - ty, - components: vec![vector_handle; columns as usize], - } - } - crate::TypeInner::Struct { ref members, .. } => { - let mut components = Vec::with_capacity(members.len()); - // copy out the types to avoid borrowing `members` - let member_tys = members.iter().map(|member| member.ty).collect::>(); - for member_ty in member_tys { - let inner = self.generate_null_constant(constants, types, member_ty)?; - components.push(constants.fetch_or_append(crate::Constant { - name: None, - specialization: None, - inner, - })); - } - crate::ConstantInner::Composite { ty, components } - } - //TODO: arrays - ref other => { - log::warn!("null constant type {:?}", other); - return Err(Error::UnsupportedType(ty)); - } - }; - Ok(inner) - } - fn parse_null_constant( &mut self, inst: Instruction, @@ -3307,7 +3230,7 @@ impl> Parser { let type_lookup = self.lookup_type.lookup(type_id)?; let ty = type_lookup.handle; - let inner = self.generate_null_constant(&mut module.constants, &mut module.types, ty)?; + let inner = null::generate_null_constant(ty, &mut module.types, &mut module.constants)?; self.lookup_constant.insert( id, @@ -3471,14 +3394,67 @@ impl> Parser { (inner, var) } ExtendedClass::Output => { - // For output interface blocks. this would be a structure. + // For output interface blocks, this would be a structure. let binding = dec.io_binding().ok(); + let init = match binding { + Some(crate::Binding::BuiltIn(built_in)) => { + match null::generate_default_built_in( + Some(built_in), + effective_ty, + &mut module.types, + &mut module.constants, + ) { + Ok(handle) => Some(handle), + Err(e) => { + log::warn!("Failed to initialize output built-in: {}", e); + None + } + } + } + Some(crate::Binding::Location { .. }) => None, + None => match module.types[effective_ty].inner { + crate::TypeInner::Struct { ref members, .. } => { + // A temporary to avoid borrowing `module.types` + let pairs = members + .iter() + .map(|member| { + let built_in = match member.binding { + Some(crate::Binding::BuiltIn(built_in)) => Some(built_in), + _ => None, + }; + (built_in, member.ty) + }) + .collect::>(); + + let mut components = Vec::with_capacity(members.len()); + for (built_in, member_ty) in pairs { + let handle = null::generate_default_built_in( + built_in, + member_ty, + &mut module.types, + &mut module.constants, + )?; + components.push(handle); + } + Some(module.constants.append(crate::Constant { + name: None, + specialization: None, + inner: crate::ConstantInner::Composite { + ty: effective_ty, + components, + }, + })) + } + _ => None, + }, + }; + let var = crate::GlobalVariable { name: dec.name, class: crate::StorageClass::Private, binding: None, ty: effective_ty, - init: None, + init, storage_access: crate::StorageAccess::empty(), }; let inner = Variable::Output(crate::FunctionResult { diff --git a/src/front/spv/null.rs b/src/front/spv/null.rs new file mode 100644 index 0000000000..1fc2b30430 --- /dev/null +++ b/src/front/spv/null.rs @@ -0,0 +1,149 @@ +use super::Error; +use crate::arena::{Arena, Handle}; + +fn make_scalar_inner(kind: crate::ScalarKind, width: crate::Bytes) -> crate::ConstantInner { + crate::ConstantInner::Scalar { + width, + value: match kind { + crate::ScalarKind::Uint => crate::ScalarValue::Uint(0), + crate::ScalarKind::Sint => crate::ScalarValue::Sint(0), + crate::ScalarKind::Float => crate::ScalarValue::Float(0.0), + crate::ScalarKind::Bool => crate::ScalarValue::Bool(false), + }, + } +} + +pub fn generate_null_constant( + ty: Handle, + type_arena: &mut Arena, + constant_arena: &mut Arena, +) -> Result { + let inner = match type_arena[ty].inner { + crate::TypeInner::Scalar { kind, width } => make_scalar_inner(kind, width), + crate::TypeInner::Vector { size, kind, width } => { + let mut components = Vec::with_capacity(size as usize); + for _ in 0..size as usize { + components.push(constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner: make_scalar_inner(kind, width), + })); + } + crate::ConstantInner::Composite { ty, components } + } + crate::TypeInner::Matrix { + columns, + rows, + width, + } => { + let vector_ty = type_arena.fetch_or_append(crate::Type { + name: None, + inner: crate::TypeInner::Vector { + kind: crate::ScalarKind::Float, + size: rows, + width, + }, + }); + let vector_inner = generate_null_constant(vector_ty, type_arena, constant_arena)?; + let vector_handle = constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner: vector_inner, + }); + crate::ConstantInner::Composite { + ty, + components: vec![vector_handle; columns as usize], + } + } + crate::TypeInner::Struct { ref members, .. } => { + let mut components = Vec::with_capacity(members.len()); + // copy out the types to avoid borrowing `members` + let member_tys = members.iter().map(|member| member.ty).collect::>(); + for member_ty in member_tys { + let inner = generate_null_constant(member_ty, type_arena, constant_arena)?; + components.push(constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner, + })); + } + crate::ConstantInner::Composite { ty, components } + } + crate::TypeInner::Array { + base, + size: crate::ArraySize::Constant(handle), + .. + } => { + let size = constant_arena[handle] + .to_array_length() + .ok_or(Error::InvalidArraySize(handle))?; + let inner = generate_null_constant(base, type_arena, constant_arena)?; + let value = constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner, + }); + crate::ConstantInner::Composite { + ty, + components: vec![value; size as usize], + } + } + ref other => { + log::warn!("null constant type {:?}", other); + return Err(Error::UnsupportedType(ty)); + } + }; + Ok(inner) +} + +/// Create a default value for an output built-in. +pub fn generate_default_built_in( + built_in: Option, + ty: Handle, + type_arena: &mut Arena, + constant_arena: &mut Arena, +) -> Result, Error> { + let inner = match built_in { + Some(crate::BuiltIn::Position) => { + let zero = constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner: crate::ConstantInner::Scalar { + value: crate::ScalarValue::Float(0.0), + width: 4, + }, + }); + let one = constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner: crate::ConstantInner::Scalar { + value: crate::ScalarValue::Float(1.0), + width: 4, + }, + }); + crate::ConstantInner::Composite { + ty, + components: vec![zero, zero, zero, one], + } + } + Some(crate::BuiltIn::PointSize) => crate::ConstantInner::Scalar { + value: crate::ScalarValue::Float(1.0), + width: 4, + }, + Some(crate::BuiltIn::FragDepth) => crate::ConstantInner::Scalar { + value: crate::ScalarValue::Float(0.0), + width: 4, + }, + Some(crate::BuiltIn::SampleMask) => crate::ConstantInner::Scalar { + value: crate::ScalarValue::Uint(!0), + width: 4, + }, + //Note: `crate::BuiltIn::ClipDistance` is intentionally left for the default path + _ => generate_null_constant(ty, type_arena, constant_arena)?, + }; + Ok(constant_arena.fetch_or_append(crate::Constant { + name: None, + specialization: None, + inner, + })) +} diff --git a/tests/out/quad-vert.msl b/tests/out/quad-vert.msl index d7363aad49..4c99ca960c 100644 --- a/tests/out/quad-vert.msl +++ b/tests/out/quad-vert.msl @@ -16,6 +16,9 @@ struct type10 { float gl_PointSize1; type6 gl_ClipDistance1; }; +constant metal::float4 const_type4_ = {0.0, 0.0, 0.0, 1.0}; +constant type6 const_type6_ = {0.0}; +constant gl_PerVertex const_gl_PerVertex = {const_type4_, 1.0, const_type6_, const_type6_}; void main1( thread metal::float2& v_uv, @@ -44,7 +47,7 @@ vertex main2Output main2( ) { metal::float2 v_uv = {}; metal::float2 a_uv = {}; - gl_PerVertex perVertexStruct = {}; + gl_PerVertex perVertexStruct = const_gl_PerVertex; metal::float2 a_pos = {}; const auto a_uv1 = varyings.a_uv1; const auto a_pos1 = varyings.a_pos1;