diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 1e3f0d0585..6e8387c005 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -727,7 +727,11 @@ impl<'a, W: Write> Writer<'a, W> { self.collect_reflection_info() } - fn write_array_size(&mut self, size: crate::ArraySize) -> BackendResult { + fn write_array_size( + &mut self, + base: Handle, + size: crate::ArraySize, + ) -> BackendResult { write!(self.out, "[")?; // Write the array size @@ -751,6 +755,16 @@ impl<'a, W: Write> Writer<'a, W> { } write!(self.out, "]")?; + + if let TypeInner::Array { + base: next_base, + size: next_size, + .. + } = self.module.types[base].inner + { + self.write_array_size(next_base, next_size)?; + } + Ok(()) } @@ -807,7 +821,7 @@ impl<'a, W: Write> Writer<'a, W> { // GLSL arrays are written as `type name[size]` // Current code is written arrays only as `[size]` // Base `type` and `name` should be written outside - TypeInner::Array { size, .. } => self.write_array_size(size)?, + TypeInner::Array { base, size, .. } => self.write_array_size(base, size)?, // Panic if either Image, Sampler, Pointer, or a Struct is being written // // Write all variants instead of `_` so that if new variants are added a @@ -994,8 +1008,8 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, " ")?; self.write_global_name(handle, global)?; - if let TypeInner::Array { size, .. } = self.module.types[global.ty].inner { - self.write_array_size(size)?; + if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner { + self.write_array_size(base, size)?; } if global.space.initializable() && is_value_init_supported(self.module, global.ty) { @@ -1298,6 +1312,11 @@ impl<'a, W: Write> Writer<'a, W> { // The leading space is important write!(this.out, " {}", &this.names[&ctx.argument_key(i as u32)])?; + // Write array size + if let TypeInner::Array { base, size, .. } = this.module.types[arg.ty].inner { + this.write_array_size(base, size)?; + } + Ok(()) })?; @@ -1356,8 +1375,8 @@ impl<'a, W: Write> Writer<'a, W> { // The leading space is important write!(self.out, " {}", self.names[&ctx.name_key(handle)])?; // Write size for array type - if let TypeInner::Array { size, .. } = self.module.types[local.ty].inner { - self.write_array_size(size)?; + if let TypeInner::Array { base, size, .. } = self.module.types[local.ty].inner { + self.write_array_size(base, size)?; } // Write the local initializer if needed if let Some(init) = local.init { @@ -1448,8 +1467,8 @@ impl<'a, W: Write> Writer<'a, W> { // `type(components)` where `components` is a comma separated list of constants crate::ConstantInner::Composite { ty, ref components } => { self.write_type(ty)?; - if let TypeInner::Array { size, .. } = self.module.types[ty].inner { - self.write_array_size(size)?; + if let TypeInner::Array { base, size, .. } = self.module.types[ty].inner { + self.write_array_size(base, size)?; } write!(self.out, "(")?; @@ -1525,7 +1544,7 @@ impl<'a, W: Write> Writer<'a, W> { &self.names[&NameKey::StructMember(handle, idx as u32)] )?; // Write [size] - self.write_array_size(size)?; + self.write_array_size(base, size)?; // Newline is important writeln!(self.out, ";")?; } @@ -2056,8 +2075,8 @@ impl<'a, W: Write> Writer<'a, W> { self.write_type(ty)?; let resolved = ctx.info[expr].ty.inner_with(&self.module.types); - if let TypeInner::Array { size, .. } = *resolved { - self.write_array_size(size)?; + if let TypeInner::Array { base, size, .. } = *resolved { + self.write_array_size(base, size)?; } write!(self.out, "(")?; @@ -2909,8 +2928,8 @@ impl<'a, W: Write> Writer<'a, W> { let resolved = base_ty_res.inner_with(&self.module.types); write!(self.out, " {}", name)?; - if let TypeInner::Array { size, .. } = *resolved { - self.write_array_size(size)?; + if let TypeInner::Array { base, size, .. } = *resolved { + self.write_array_size(base, size)?; } write!(self.out, " = ")?; self.write_expr(handle, ctx)?; @@ -2948,7 +2967,7 @@ impl<'a, W: Write> Writer<'a, W> { proc::IndexableLength::Dynamic => return Ok(()), }; self.write_type(base)?; - self.write_array_size(size)?; + self.write_array_size(base, size)?; write!(self.out, "(")?; for _ in 1..count { self.write_zero_init_value(base)?; diff --git a/src/back/hlsl/conv.rs b/src/back/hlsl/conv.rs index 978b073880..547b7632f0 100644 --- a/src/back/hlsl/conv.rs +++ b/src/back/hlsl/conv.rs @@ -1,3 +1,5 @@ +use std::borrow::Cow; + use super::Error; impl crate::ScalarKind { @@ -35,6 +37,79 @@ impl crate::TypeInner { _ => false, } } + + pub(super) fn try_size_hlsl( + &self, + types: &crate::UniqueArena, + constants: &crate::Arena, + ) -> Result { + Ok(match *self { + Self::Matrix { + columns, + rows, + width, + } => { + let aligned_rows = if rows > crate::VectorSize::Bi { 4 } else { 2 }; + let stride = aligned_rows * width as u32; + let last_row_size = rows as u32 * width as u32; + ((columns as u32 - 1) * stride) + last_row_size + } + Self::Array { base, size, stride } => { + let count = match size { + crate::ArraySize::Constant(handle) => { + let constant = constants.try_get(handle)?; + constant.to_array_length().unwrap_or(1) + } + // A dynamically-sized array has to have at least one element + crate::ArraySize::Dynamic => 1, + }; + let last_el_size = types[base].inner.try_size_hlsl(types, constants)?; + ((count - 1) * stride) + last_el_size + } + _ => self.try_size(constants)?, + }) + } + + /// Used to generate the name of the wrapped type constructor + pub(super) fn hlsl_type_id<'a>( + &self, + base: crate::Handle, + types: &crate::UniqueArena, + constants: &crate::Arena, + names: &'a crate::FastHashMap, + ) -> Result, Error> { + Ok(match types[base].inner { + crate::TypeInner::Scalar { kind, width } => Cow::Borrowed(kind.to_hlsl_str(width)?), + crate::TypeInner::Vector { size, kind, width } => Cow::Owned(format!( + "{}{}", + kind.to_hlsl_str(width)?, + crate::back::vector_size_str(size) + )), + crate::TypeInner::Matrix { + columns, + rows, + width, + } => Cow::Owned(format!( + "{}{}x{}", + crate::ScalarKind::Float.to_hlsl_str(width)?, + crate::back::vector_size_str(columns), + crate::back::vector_size_str(rows), + )), + crate::TypeInner::Array { + base, + size: crate::ArraySize::Constant(size), + .. + } => Cow::Owned(format!( + "array{}_{}_", + constants[size].to_array_length().unwrap(), + self.hlsl_type_id(base, types, constants, names)? + )), + crate::TypeInner::Struct { .. } => { + Cow::Borrowed(&names[&crate::proc::NameKey::Type(base)]) + } + _ => unreachable!(), + }) + } } impl crate::StorageFormat { diff --git a/src/back/hlsl/help.rs b/src/back/hlsl/help.rs index c652a1d6e3..78e87cbd49 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -26,7 +26,7 @@ int dim_1d = NagaDimensions1D(image_1d); ``` */ -use super::{super::FunctionCtx, BackendResult, Error}; +use super::{super::FunctionCtx, BackendResult}; use crate::{arena::Handle, proc::NameKey}; use std::fmt::Write; @@ -350,9 +350,15 @@ impl<'a, W: Write> super::Writer<'a, W> { pub(super) fn write_wrapped_constructor_function_name( &mut self, + module: &crate::Module, constructor: WrappedConstructor, ) -> BackendResult { - let name = &self.names[&NameKey::Type(constructor.ty)]; + let name = module.types[constructor.ty].inner.hlsl_type_id( + constructor.ty, + &module.types, + &module.constants, + &self.names, + )?; write!(self.out, "Construct{}", name)?; Ok(()) } @@ -369,69 +375,114 @@ impl<'a, W: Write> super::Writer<'a, W> { const RETURN_VARIABLE_NAME: &str = "ret"; // Write function return type and name - let struct_name = &self.names[&NameKey::Type(constructor.ty)]; - write!(self.out, "{} ", struct_name)?; - self.write_wrapped_constructor_function_name(constructor)?; + self.write_type(module, constructor.ty)?; + write!(self.out, " ")?; + self.write_wrapped_constructor_function_name(module, constructor)?; // Write function parameters write!(self.out, "(")?; - let members = match module.types[constructor.ty].inner { - crate::TypeInner::Struct { ref members, .. } => members, - _ => return Err(Error::Unimplemented("non-struct constructor".to_string())), - }; - for (i, member) in members.iter().enumerate() { + + let mut write_arg = |i, ty| -> BackendResult { if i != 0 { write!(self.out, ", ")?; } - self.write_type(module, member.ty)?; + self.write_type(module, ty)?; write!(self.out, " {}{}", ARGUMENT_VARIABLE_NAME, i)?; - if let crate::TypeInner::Array { size, .. } = module.types[member.ty].inner { - self.write_array_size(module, size)?; + if let crate::TypeInner::Array { base, size, .. } = module.types[ty].inner { + self.write_array_size(module, base, size)?; } + Ok(()) + }; + + match module.types[constructor.ty].inner { + crate::TypeInner::Struct { ref members, .. } => { + for (i, member) in members.iter().enumerate() { + write_arg(i, member.ty)?; + } + } + crate::TypeInner::Array { + base, + size: crate::ArraySize::Constant(size), + .. + } => { + let count = module.constants[size].to_array_length().unwrap(); + for i in 0..count as usize { + write_arg(i, base)?; + } + } + _ => unreachable!(), + }; + + write!(self.out, ")")?; + + if let crate::TypeInner::Array { base, size, .. } = module.types[constructor.ty].inner { + self.write_array_size(module, base, size)?; } - // Write function body - writeln!(self.out, ") {{")?; - let struct_name = &self.names[&NameKey::Type(constructor.ty)]; - writeln!( - self.out, - "{}{} {};", - INDENT, struct_name, RETURN_VARIABLE_NAME - )?; - for i in 0..members.len() as u32 { - let member = &members[i as usize]; - - let field_name = &self.names[&NameKey::StructMember(constructor.ty, i)]; - - match module.types[member.ty].inner { - crate::TypeInner::Matrix { - columns, - rows: crate::VectorSize::Bi, - .. - } if member.binding.is_none() => { - for j in 0..columns as u8 { - writeln!( - self.out, - "{}{}.{}_{} = {}{}[{}];", - INDENT, - RETURN_VARIABLE_NAME, - field_name, - j, - ARGUMENT_VARIABLE_NAME, - i, - j - )?; + // Write function body + writeln!(self.out, " {{")?; + + match module.types[constructor.ty].inner { + crate::TypeInner::Struct { ref members, .. } => { + let struct_name = &self.names[&NameKey::Type(constructor.ty)]; + writeln!( + self.out, + "{}{} {};", + INDENT, struct_name, RETURN_VARIABLE_NAME + )?; + for (i, member) in members.iter().enumerate() { + let field_name = &self.names[&NameKey::StructMember(constructor.ty, i as u32)]; + + match module.types[member.ty].inner { + crate::TypeInner::Matrix { + columns, + rows: crate::VectorSize::Bi, + .. + } if member.binding.is_none() => { + for j in 0..columns as u8 { + writeln!( + self.out, + "{}{}.{}_{} = {}{}[{}];", + INDENT, + RETURN_VARIABLE_NAME, + field_name, + j, + ARGUMENT_VARIABLE_NAME, + i, + j + )?; + } + } + _ => { + writeln!( + self.out, + "{}{}.{} = {}{};", + INDENT, RETURN_VARIABLE_NAME, field_name, ARGUMENT_VARIABLE_NAME, i, + )?; + } } } - _ => { - //TODO: handle arrays? - writeln!( - self.out, - "{}{}.{} = {}{};", - INDENT, RETURN_VARIABLE_NAME, field_name, ARGUMENT_VARIABLE_NAME, i, - )?; + } + crate::TypeInner::Array { + base, + size: crate::ArraySize::Constant(size), + .. + } => { + write!(self.out, "{}", INDENT)?; + self.write_type(module, base)?; + write!(self.out, " {}", RETURN_VARIABLE_NAME)?; + self.write_array_size(module, base, crate::ArraySize::Constant(size))?; + write!(self.out, " = {{ ")?; + let count = module.constants[size].to_array_length().unwrap(); + for i in 0..count { + if i != 0 { + write!(self.out, ", ")?; + } + write!(self.out, "{}{}", ARGUMENT_VARIABLE_NAME, i)?; } + writeln!(self.out, " }};",)?; } + _ => unreachable!(), } // Write return value @@ -831,7 +882,9 @@ impl<'a, W: Write> super::Writer<'a, W> { } crate::Expression::Compose { ty, components: _ } => { let constructor = match module.types[ty].inner { - crate::TypeInner::Struct { .. } => WrappedConstructor { ty }, + crate::TypeInner::Struct { .. } | crate::TypeInner::Array { .. } => { + WrappedConstructor { ty } + } _ => continue, }; if !self.wrapped.constructors.contains(&constructor) { @@ -887,6 +940,33 @@ impl<'a, W: Write> super::Writer<'a, W> { Ok(()) } + pub(super) fn write_wrapped_constructor_function_for_constant( + &mut self, + module: &crate::Module, + constant: &crate::Constant, + ) -> BackendResult { + if let crate::ConstantInner::Composite { ty, ref components } = constant.inner { + match module.types[ty].inner { + crate::TypeInner::Struct { .. } | crate::TypeInner::Array { .. } => { + let constructor = WrappedConstructor { ty }; + if !self.wrapped.constructors.contains(&constructor) { + self.write_wrapped_constructor_function(module, constructor)?; + self.wrapped.constructors.insert(constructor); + } + } + _ => {} + } + for constant in components { + self.write_wrapped_constructor_function_for_constant( + module, + &module.constants[*constant], + )?; + } + } + + Ok(()) + } + pub(super) fn write_texture_coordinates( &mut self, kind: &str, diff --git a/src/back/hlsl/storage.rs b/src/back/hlsl/storage.rs index cab351aa59..cea989deec 100644 --- a/src/back/hlsl/storage.rs +++ b/src/back/hlsl/storage.rs @@ -307,7 +307,7 @@ impl super::Writer<'_, W> { self.write_value_type(module, &module.types[base].inner)?; let depth = level.next().0; write!(self.out, " {}{}", STORE_TEMP_NAME, depth)?; - self.write_array_size(module, crate::ArraySize::Constant(const_handle))?; + self.write_array_size(module, base, crate::ArraySize::Constant(const_handle))?; write!(self.out, " = ")?; self.write_store_value(module, &value, func_ctx)?; writeln!(self.out, ";")?; diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index b514d11557..b3e61e7f53 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -152,7 +152,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Write all structs for (handle, ty) in module.types.iter() { - if let TypeInner::Struct { ref members, .. } = ty.inner { + if let TypeInner::Struct { ref members, span } = ty.inner { if let Some(member) = members.last() { if let TypeInner::Array { size: crate::ArraySize::Dynamic, @@ -176,12 +176,18 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { module, handle, members, + span, ep_result.map(|r| (r.0, Io::Output)), )?; writeln!(self.out)?; } } + // Write wrapped constructor functions used in constants + for (_, constant) in module.constants.iter() { + self.write_wrapped_constructor_function_for_constant(module, constant)?; + } + // Write all globals for (ty, _) in module.global_variables.iter() { self.write_global(module, ty)?; @@ -525,8 +531,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { let arg_name = &self.names[&NameKey::EntryPointArgument(ep_index, arg_index as u32)]; write!(self.out, " {}", arg_name)?; match module.types[arg.ty].inner { - TypeInner::Array { size, .. } => { - self.write_array_size(module, size)?; + TypeInner::Array { base, size, .. } => { + self.write_array_size(module, base, size)?; let fake_member = fake_iter.next().unwrap(); writeln!(self.out, " = {}.{};", ep_input.arg_name, fake_member.name)?; } @@ -631,8 +637,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ")")?; } else { // need to write the array size if the type was emitted with `write_type` - if let TypeInner::Array { size, .. } = module.types[global.ty].inner { - self.write_array_size(module, size)?; + if let TypeInner::Array { base, size, .. } = module.types[global.ty].inner { + self.write_array_size(module, base, size)?; } if global.space == crate::AddressSpace::Private { write!(self.out, " = ")?; @@ -650,8 +656,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { let sub_name = &self.names[&NameKey::GlobalVariable(handle)]; write!(self.out, " {}", sub_name)?; // need to write the array size if the type was emitted with `write_type` - if let TypeInner::Array { size, .. } = module.types[global.ty].inner { - self.write_array_size(module, size)?; + if let TypeInner::Array { base, size, .. } = module.types[global.ty].inner { + self.write_array_size(module, base, size)?; } writeln!(self.out, "; }}")?; } else { @@ -713,6 +719,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { pub(super) fn write_array_size( &mut self, module: &Module, + base: Handle, size: crate::ArraySize, ) -> BackendResult { write!(self.out, "[")?; @@ -729,6 +736,16 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } write!(self.out, "]")?; + + if let TypeInner::Array { + base: next_base, + size: next_size, + .. + } = module.types[base].inner + { + self.write_array_size(module, next_base, next_size)?; + } + Ok(()) } @@ -741,6 +758,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { module: &Module, handle: Handle, members: &[crate::StructMember], + span: u32, shader_stage: Option<(ShaderStage, Io)>, ) -> BackendResult { // Write struct name @@ -759,7 +777,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } } let ty_inner = &module.types[member.ty].inner; - last_offset = member.offset + ty_inner.size(&module.constants); + last_offset = member.offset + + ty_inner + .try_size_hlsl(&module.types, &module.constants) + .unwrap(); // The indentation is only for readability write!(self.out, "{}", back::INDENT)?; @@ -782,7 +803,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { &self.names[&NameKey::StructMember(handle, index as u32)] )?; // Write [size] - self.write_array_size(module, size)?; + self.write_array_size(module, base, size)?; } // We treat matrices of the form `matCx2` as a sequence of C `vec2`s // (see top level module docs for details). @@ -832,6 +853,14 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, ";")?; } + // add padding at the end since sizes of types don't get rounded up to their alignment in HLSL + if members.last().unwrap().binding.is_none() && span > last_offset { + let padding = (span - last_offset) / 4; + for i in 0..padding { + writeln!(self.out, "{}int _end_pad_{};", back::INDENT, i)?; + } + } + writeln!(self.out, "}};")?; Ok(()) } @@ -904,8 +933,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // HLSL arrays are written as `type name[size]` // Current code is written arrays only as `[size]` // Base `type` and `name` should be written outside - TypeInner::Array { size, .. } => { - self.write_array_size(module, size)?; + TypeInner::Array { base, size, .. } => { + self.write_array_size(module, base, size)?; } _ => { return Err(Error::Unimplemented(format!( @@ -989,8 +1018,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Write argument name. Space is important. write!(self.out, " {}", argument_name)?; - if let TypeInner::Array { size, .. } = module.types[arg.ty].inner { - self.write_array_size(module, size)?; + if let TypeInner::Array { base, size, .. } = module.types[arg.ty].inner { + self.write_array_size(module, base, size)?; } } } @@ -1009,8 +1038,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { &self.names[&NameKey::EntryPointArgument(ep_index, index as u32)]; write!(self.out, " {}", argument_name)?; - if let TypeInner::Array { size, .. } = module.types[arg.ty].inner { - self.write_array_size(module, size)?; + if let TypeInner::Array { base, size, .. } = module.types[arg.ty].inner { + self.write_array_size(module, base, size)?; } if let Some(ref binding) = arg.binding { @@ -1053,8 +1082,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.write_type(module, local.ty)?; write!(self.out, " {}", self.names[&func_ctx.name_key(handle)])?; // Write size for array type - if let TypeInner::Array { size, .. } = module.types[local.ty].inner { - self.write_array_size(module, size)?; + if let TypeInner::Array { base, size, .. } = module.types[local.ty].inner { + self.write_array_size(module, base, size)?; } write!(self.out, " = ")?; @@ -1706,19 +1735,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { match *expression { Expression::Constant(constant) => self.write_constant(module, constant)?, Expression::Compose { ty, ref components } => { - let (brace_open, brace_close) = match module.types[ty].inner { - TypeInner::Struct { .. } => { - self.write_wrapped_constructor_function_name(WrappedConstructor { ty })?; - ("(", ")") + match module.types[ty].inner { + TypeInner::Struct { .. } | TypeInner::Array { .. } => { + self.write_wrapped_constructor_function_name( + module, + WrappedConstructor { ty }, + )?; } - TypeInner::Array { .. } => ("{ ", " }"), _ => { self.write_type(module, ty)?; - ("(", ")") } }; - write!(self.out, "{}", brace_open)?; + write!(self.out, "(")?; for (index, &component) in components.iter().enumerate() { if index != 0 { @@ -1728,7 +1757,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.write_expr(module, component, func_ctx)?; } - write!(self.out, "{}", brace_close)?; + write!(self.out, ")")?; } // All of the multiplication can be expressed as `mul`, // except vector * vector, which needs to use the "*" operator. @@ -2324,15 +2353,15 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { ty: Handle, components: &[Handle], ) -> BackendResult { - let (open_b, close_b) = match module.types[ty].inner { - TypeInner::Array { .. } | TypeInner::Struct { .. } => ("{ ", " }"), + match module.types[ty].inner { + TypeInner::Struct { .. } | TypeInner::Array { .. } => { + self.write_wrapped_constructor_function_name(module, WrappedConstructor { ty })?; + } _ => { - // We should write type only for non struct/array constants self.write_type(module, ty)?; - ("(", ")") } }; - write!(self.out, "{}", open_b)?; + write!(self.out, "(")?; for (index, constant) in components.iter().enumerate() { self.write_constant(module, *constant)?; // Only write a comma if isn't the last element @@ -2341,7 +2370,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ", ")?; } } - write!(self.out, "{}", close_b)?; + write!(self.out, ")")?; Ok(()) } @@ -2392,8 +2421,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, " {}", name)?; // If rhs is a array type, we should write array size - if let TypeInner::Array { size, .. } = *resolved { - self.write_array_size(module, size)?; + if let TypeInner::Array { base, size, .. } = *resolved { + self.write_array_size(module, base, size)?; } write!(self.out, " = ")?; self.write_expr(module, handle, ctx)?; diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index a034fdda7d..e64ef2152d 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -54,6 +54,10 @@ fn read_from_private(foo: ptr) -> f32 { return *foo; } +fn test_arr_as_arg(a: array, 5>) -> f32 { + return a[4][9]; +} + @stage(vertex) fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0.0; @@ -79,6 +83,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { c[vi + 1u] = 42; let value = c[vi]; + var _ = test_arr_as_arg(array, 5>()); + return vec4(matrix * vec4(vec4(value)), 2.0); } diff --git a/tests/in/padding.param.ron b/tests/in/padding.param.ron new file mode 100644 index 0000000000..9e20176bac --- /dev/null +++ b/tests/in/padding.param.ron @@ -0,0 +1,22 @@ +( + spv: ( + version: (1, 1), + debug: true, + adjust_coordinate_space: false, + ), + msl: ( + lang_version: (2, 0), + per_stage_map: ( + vs: ( + resources: { + (group: 0, binding: 0): (buffer: Some(0), mutable: false), + (group: 0, binding: 1): (buffer: Some(1), mutable: false), + (group: 0, binding: 2): (buffer: Some(2), mutable: false), + }, + ), + ), + inline_samplers: [], + spirv_cross_compatibility: false, + fake_missing_bindings: false, + ), +) diff --git a/tests/in/padding.wgsl b/tests/in/padding.wgsl new file mode 100644 index 0000000000..0ff60f3932 --- /dev/null +++ b/tests/in/padding.wgsl @@ -0,0 +1,33 @@ +struct S { + a: vec3, +} + +struct Test { + a: S, + b: f32, // offset: 16 +} + +struct Test2 { + a: array, 2>, + b: f32, // offset: 32 +} + +struct Test3 { + a: mat4x3, + b: f32, // offset: 64 +} + +@group(0) @binding(0) +var input: Test; + +@group(0) @binding(1) +var input2: Test2; + +@group(0) @binding(2) +var input3: Test3; + + +@stage(vertex) +fn vertex() -> @builtin(position) vec4 { + return vec4(1.0) * input.b * input2.b * input3.b; +} \ No newline at end of file diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index ff67b1b988..2b92925a0b 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -25,6 +25,10 @@ float read_from_private(inout float foo_1) { return _e3; } +float test_arr_as_arg(float a[5][10]) { + return a[4][9]; +} + void main() { int tmp = 0; int value = _group_0_binding_0_cs.atom; diff --git a/tests/out/glsl/access.foo_frag.Fragment.glsl b/tests/out/glsl/access.foo_frag.Fragment.glsl index b7c62d40e6..e07dad7a00 100644 --- a/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -24,6 +24,10 @@ float read_from_private(inout float foo_1) { return _e3; } +float test_arr_as_arg(float a[5][10]) { + return a[4][9]; +} + void main() { _group_0_binding_0_fs.matrix[1][2] = 1.0; _group_0_binding_0_fs.matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0)); diff --git a/tests/out/glsl/access.foo_vert.Vertex.glsl b/tests/out/glsl/access.foo_vert.Vertex.glsl index b01479c243..c545e44286 100644 --- a/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -22,35 +22,35 @@ uniform Baz_block_1Vertex { Baz _group_0_binding_1_vs; }; void test_matrix_within_struct_accesses() { int idx = 9; - mat3x2 unnamed = mat3x2(0.0); - vec2 unnamed_1 = vec2(0.0); + mat3x2 unnamed_1 = mat3x2(0.0); vec2 unnamed_2 = vec2(0.0); - float unnamed_3 = 0.0; + vec2 unnamed_3 = vec2(0.0); float unnamed_4 = 0.0; float unnamed_5 = 0.0; float unnamed_6 = 0.0; + float unnamed_7 = 0.0; Baz t = Baz(mat3x2(0.0)); int _e4 = idx; idx = (_e4 - 1); mat3x2 _e8 = _group_0_binding_1_vs.m; - unnamed = _e8; + unnamed_1 = _e8; vec2 _e13 = _group_0_binding_1_vs.m[0]; - unnamed_1 = _e13; + unnamed_2 = _e13; int _e16 = idx; vec2 _e18 = _group_0_binding_1_vs.m[_e16]; - unnamed_2 = _e18; + unnamed_3 = _e18; float _e25 = _group_0_binding_1_vs.m[0][1]; - unnamed_3 = _e25; + unnamed_4 = _e25; int _e30 = idx; float _e32 = _group_0_binding_1_vs.m[0][_e30]; - unnamed_4 = _e32; + unnamed_5 = _e32; int _e35 = idx; float _e39 = _group_0_binding_1_vs.m[_e35][1]; - unnamed_5 = _e39; + unnamed_6 = _e39; int _e42 = idx; int _e44 = idx; float _e46 = _group_0_binding_1_vs.m[_e42][_e44]; - unnamed_6 = _e46; + unnamed_7 = _e46; t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0))); int _e57 = idx; idx = (_e57 + 1); @@ -74,21 +74,28 @@ float read_from_private(inout float foo_1) { return _e3; } +float test_arr_as_arg(float a[5][10]) { + return a[4][9]; +} + void main() { uint vi = uint(gl_VertexID); float foo = 0.0; int c[5] = int[5](0, 0, 0, 0, 0); + float unnamed = 0.0; float baz_1 = foo; foo = 1.0; test_matrix_within_struct_accesses(); mat4x3 matrix = _group_0_binding_0_vs.matrix; uvec2 arr[2] = _group_0_binding_0_vs.arr; float b = _group_0_binding_0_vs.matrix[3][0]; - int a = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; + int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; float _e28 = read_from_private(foo); - c = int[5](a, int(b), 3, 4, 5); + c = int[5](a_1, int(b), 3, 4, 5); c[(vi + 1u)] = 42; int value = c[vi]; + float _e42 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0))); + unnamed = _e42; gl_Position = vec4((matrix * vec4(ivec4(value))), 2.0); gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w); return; diff --git a/tests/out/glsl/padding.vertex.Vertex.glsl b/tests/out/glsl/padding.vertex.Vertex.glsl new file mode 100644 index 0000000000..21a231a37b --- /dev/null +++ b/tests/out/glsl/padding.vertex.Vertex.glsl @@ -0,0 +1,36 @@ +#version 310 es + +precision highp float; +precision highp int; + +struct S { + vec3 a; +}; +struct Test { + S a; + float b; +}; +struct Test2_ { + vec3 a[2]; + float b; +}; +struct Test3_ { + mat4x3 a; + float b; +}; +uniform Test_block_0Vertex { Test _group_0_binding_0_vs; }; + +uniform Test2__block_1Vertex { Test2_ _group_0_binding_1_vs; }; + +uniform Test3__block_2Vertex { Test3_ _group_0_binding_2_vs; }; + + +void main() { + float _e6 = _group_0_binding_0_vs.b; + float _e9 = _group_0_binding_1_vs.b; + float _e12 = _group_0_binding_2_vs.b; + gl_Position = (((vec4(1.0) * _e6) * _e9) * _e12); + gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w); + return; +} + diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 7f63c78d87..4a807a49f7 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -1,12 +1,23 @@ struct AlignedWrapper { int value; + int _end_pad_0; }; struct Baz { float2 m_0; float2 m_1; float2 m_2; }; +float Constructarray10_float_(float arg0, float arg1, float arg2, float arg3, float arg4, float arg5, float arg6, float arg7, float arg8, float arg9)[10] { + float ret[10] = { arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7, arg8, arg9 }; + return ret; +} + +float Constructarray5_array10_float__(float arg0[10], float arg1[10], float arg2[10], float arg3[10], float arg4[10])[5][10] { + float ret[5][10] = { arg0, arg1, arg2, arg3, arg4 }; + return ret; +} + RWByteAddressBuffer bar : register(u0); cbuffer baz : register(b1) { Baz baz; } @@ -47,36 +58,36 @@ Baz ConstructBaz(float3x2 arg0) { void test_matrix_within_struct_accesses() { int idx = 9; - float3x2 unnamed = (float3x2)0; - float2 unnamed_1 = (float2)0; + float3x2 unnamed_1 = (float3x2)0; float2 unnamed_2 = (float2)0; - float unnamed_3 = (float)0; + float2 unnamed_3 = (float2)0; float unnamed_4 = (float)0; float unnamed_5 = (float)0; float unnamed_6 = (float)0; + float unnamed_7 = (float)0; Baz t = (Baz)0; int _expr4 = idx; idx = (_expr4 - 1); float3x2 _expr8 = GetMatmOnBaz(baz); - unnamed = _expr8; + unnamed_1 = _expr8; float2 _expr13 = GetMatmOnBaz(baz)[0]; - unnamed_1 = _expr13; + unnamed_2 = _expr13; int _expr16 = idx; float2 _expr18 = GetMatmOnBaz(baz)[_expr16]; - unnamed_2 = _expr18; + unnamed_3 = _expr18; float _expr25 = GetMatmOnBaz(baz)[0][1]; - unnamed_3 = _expr25; + unnamed_4 = _expr25; int _expr30 = idx; float _expr32 = GetMatmOnBaz(baz)[0][_expr30]; - unnamed_4 = _expr32; + unnamed_5 = _expr32; int _expr35 = idx; float _expr39 = GetMatmOnBaz(baz)[_expr35][1]; - unnamed_5 = _expr39; + unnamed_6 = _expr39; int _expr42 = idx; int _expr44 = idx; float _expr46 = GetMatmOnBaz(baz)[_expr42][_expr44]; - unnamed_6 = _expr46; + unnamed_7 = _expr46; t = ConstructBaz(float3x2(float2(1.0.xx), float2(2.0.xx), float2(3.0.xx))); int _expr57 = idx; idx = (_expr57 + 1); @@ -101,6 +112,11 @@ float read_from_private(inout float foo_1) return _expr3; } +float test_arr_as_arg(float a[5][10]) +{ + return a[4][9]; +} + uint NagaBufferLengthRW(RWByteAddressBuffer buffer) { uint ret; @@ -108,10 +124,16 @@ uint NagaBufferLengthRW(RWByteAddressBuffer buffer) return ret; } +int Constructarray5_int_(int arg0, int arg1, int arg2, int arg3, int arg4)[5] { + int ret[5] = { arg0, arg1, arg2, arg3, arg4 }; + return ret; +} + float4 foo_vert(uint vi : SV_VertexID) : SV_Position { float foo = 0.0; int c[5] = {(int)0,(int)0,(int)0,(int)0,(int)0}; + float unnamed = (float)0; float baz_1 = foo; foo = 1.0; @@ -119,17 +141,24 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position float4x3 matrix_ = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48))); uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))}; float b = asfloat(bar.Load(0+48+0)); - int a = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120)); + int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120)); const float _e28 = read_from_private(foo); { - int _result[5]={ a, int(b), 3, 4, 5 }; + int _result[5]=Constructarray5_int_(a_1, int(b), 3, 4, 5); for(int _i=0; _i<5; ++_i) c[_i] = _result[_i]; } c[(vi + 1u)] = 42; int value = c[vi]; + const float _e42 = test_arr_as_arg(Constructarray5_array10_float__(Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0))); + unnamed = _e42; return float4(mul(float4(int4(value.xxxx)), matrix_), 2.0); } +uint2 Constructarray2_uint2_(uint2 arg0, uint2 arg1)[2] { + uint2 ret[2] = { arg0, arg1 }; + return ret; +} + float4 foo_frag() : SV_Target0 { bar.Store(8+16+0, asuint(1.0)); @@ -141,7 +170,7 @@ float4 foo_frag() : SV_Target0 bar.Store3(0+48, asuint(_value2[3])); } { - uint2 _value2[2] = { uint2(0u.xx), uint2(1u.xx) }; + uint2 _value2[2] = Constructarray2_uint2_(uint2(0u.xx), uint2(1u.xx)); bar.Store2(104+0, asuint(_value2[0])); bar.Store2(104+8, asuint(_value2[1])); } diff --git a/tests/out/hlsl/operators.hlsl b/tests/out/hlsl/operators.hlsl index 08aed4f934..984d519f92 100644 --- a/tests/out/hlsl/operators.hlsl +++ b/tests/out/hlsl/operators.hlsl @@ -6,8 +6,23 @@ static const int4 v_i32_one = int4(1, 1, 1, 1); struct Foo { float4 a; int b; + int _end_pad_0; + int _end_pad_1; + int _end_pad_2; }; +Foo ConstructFoo(float4 arg0, int arg1) { + Foo ret; + ret.a = arg0; + ret.b = arg1; + return ret; +} + +Foo Constructarray3_Foo_(Foo arg0, Foo arg1, Foo arg2)[3] { + Foo ret[3] = { arg0, arg1, arg2 }; + return ret; +} + float4 builtins() { int s1_ = (true ? 1 : 0); @@ -43,10 +58,8 @@ float3 bool_cast(float3 x) return float3(y); } -Foo ConstructFoo(float4 arg0, int arg1) { - Foo ret; - ret.a = arg0; - ret.b = arg1; +int Constructarray4_int_(int arg0, int arg1, int arg2, int arg3)[4] { + int ret[4] = { arg0, arg1, arg2, arg3 }; return ret; } @@ -59,8 +72,8 @@ float constructors() float unnamed_3 = 0.0; uint2 unnamed_4 = uint2(0u, 0u); float2x2 unnamed_5 = float2x2(float2(0.0, 0.0), float2(0.0, 0.0)); - Foo unnamed_6[3] = { { float4(0.0, 0.0, 0.0, 0.0), 0 }, { float4(0.0, 0.0, 0.0, 0.0), 0 }, { float4(0.0, 0.0, 0.0, 0.0), 0 } }; - Foo unnamed_7 = { float4(0.0, 0.0, 0.0, 0.0), 0 }; + Foo unnamed_6[3] = Constructarray3_Foo_(ConstructFoo(float4(0.0, 0.0, 0.0, 0.0), 0), ConstructFoo(float4(0.0, 0.0, 0.0, 0.0), 0), ConstructFoo(float4(0.0, 0.0, 0.0, 0.0), 0)); + Foo unnamed_7 = ConstructFoo(float4(0.0, 0.0, 0.0, 0.0), 0); uint2 unnamed_8 = (uint2)0; float2x2 unnamed_9 = (float2x2)0; int unnamed_10[4] = {(int)0,(int)0,(int)0,(int)0}; @@ -71,7 +84,7 @@ float constructors() unnamed_8 = uint2(0u.xx); unnamed_9 = float2x2(float2(0.0.xx), float2(0.0.xx)); { - int _result[4]={ 0, 1, 2, 3 }; + int _result[4]=Constructarray4_int_(0, 1, 2, 3); for(int _i=0; _i<4; ++_i) unnamed_10[_i] = _result[_i]; } float _expr70 = foo.a.x; diff --git a/tests/out/hlsl/padding.hlsl b/tests/out/hlsl/padding.hlsl new file mode 100644 index 0000000000..2bf074c067 --- /dev/null +++ b/tests/out/hlsl/padding.hlsl @@ -0,0 +1,43 @@ + +struct S { + float3 a; + int _end_pad_0; +}; + +struct Test { + S a; + float b; + int _end_pad_0; + int _end_pad_1; + int _end_pad_2; +}; + +struct Test2_ { + float3 a[2]; + int _pad1_0; + float b; + int _end_pad_0; + int _end_pad_1; + int _end_pad_2; +}; + +struct Test3_ { + row_major float4x3 a; + int _pad1_0; + float b; + int _end_pad_0; + int _end_pad_1; + int _end_pad_2; +}; + +cbuffer input : register(b0) { Test input; } +cbuffer input2_ : register(b1) { Test2_ input2_; } +cbuffer input3_ : register(b2) { Test3_ input3_; } + +float4 vertex() : SV_Position +{ + float _expr6 = input.b; + float _expr9 = input2_.b; + float _expr12 = input3_.b; + return (((float4(1.0.xxxx) * _expr6) * _expr9) * _expr12); +} diff --git a/tests/out/hlsl/padding.hlsl.config b/tests/out/hlsl/padding.hlsl.config new file mode 100644 index 0000000000..875a4fa4a3 --- /dev/null +++ b/tests/out/hlsl/padding.hlsl.config @@ -0,0 +1,3 @@ +vertex=(vertex:vs_5_1 ) +fragment=() +compute=() diff --git a/tests/out/hlsl/quad-vert.hlsl b/tests/out/hlsl/quad-vert.hlsl index a02dde77f0..9f5e082009 100644 --- a/tests/out/hlsl/quad-vert.hlsl +++ b/tests/out/hlsl/quad-vert.hlsl @@ -11,9 +11,23 @@ struct type_9 { float4 gl_Position : SV_Position; }; +float Constructarray1_float_(float arg0)[1] { + float ret[1] = { arg0 }; + return ret; +} + +gl_PerVertex Constructgl_PerVertex(float4 arg0, float arg1, float arg2[1], float arg3[1]) { + gl_PerVertex ret; + ret.gl_Position = arg0; + ret.gl_PointSize = arg1; + ret.gl_ClipDistance = arg2; + ret.gl_CullDistance = arg3; + return ret; +} + static float2 v_uv = (float2)0; static float2 a_uv_1 = (float2)0; -static gl_PerVertex perVertexStruct = { float4(0.0, 0.0, 0.0, 1.0), 1.0, { 0.0 }, { 0.0 } }; +static gl_PerVertex perVertexStruct = Constructgl_PerVertex(float4(0.0, 0.0, 0.0, 1.0), 1.0, Constructarray1_float_(0.0), Constructarray1_float_(0.0)); static float2 a_pos_1 = (float2)0; struct VertexOutput_main { diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index 0f3b5cb122..a791ee6ab2 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -29,43 +29,51 @@ struct Bar { struct Baz { metal::float3x2 m; }; -struct type_15 { +struct type_12 { + float inner[10]; +}; +struct type_13 { + type_12 inner[5]; +}; +struct type_17 { int inner[5]; }; +constant type_12 const_type_12_ = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0}; +constant type_13 const_type_13_ = {const_type_12_, const_type_12_, const_type_12_, const_type_12_, const_type_12_}; void test_matrix_within_struct_accesses( constant Baz& baz ) { int idx = 9; - metal::float3x2 unnamed; - metal::float2 unnamed_1; + metal::float3x2 unnamed_1; metal::float2 unnamed_2; - float unnamed_3; + metal::float2 unnamed_3; float unnamed_4; float unnamed_5; float unnamed_6; + float unnamed_7; Baz t; int _e4 = idx; idx = _e4 - 1; metal::float3x2 _e8 = baz.m; - unnamed = _e8; + unnamed_1 = _e8; metal::float2 _e13 = baz.m[0]; - unnamed_1 = _e13; + unnamed_2 = _e13; int _e16 = idx; metal::float2 _e18 = baz.m[_e16]; - unnamed_2 = _e18; + unnamed_3 = _e18; float _e25 = baz.m[0].y; - unnamed_3 = _e25; + unnamed_4 = _e25; int _e30 = idx; float _e32 = baz.m[0][_e30]; - unnamed_4 = _e32; + unnamed_5 = _e32; int _e35 = idx; float _e39 = baz.m[_e35].y; - unnamed_5 = _e39; + unnamed_6 = _e39; int _e42 = idx; int _e44 = idx; float _e46 = baz.m[_e42][_e44]; - unnamed_6 = _e46; + unnamed_7 = _e46; t = Baz {metal::float3x2(metal::float2(1.0), metal::float2(2.0), metal::float2(3.0))}; int _e57 = idx; idx = _e57 + 1; @@ -91,6 +99,12 @@ float read_from_private( return _e3; } +float test_arr_as_arg( + type_13 a +) { + return a.inner[4].inner[9]; +} + struct foo_vertInput { }; struct foo_vertOutput { @@ -103,18 +117,21 @@ vertex foo_vertOutput foo_vert( , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { float foo = 0.0; - type_15 c; + type_17 c; + float unnamed; float baz_1 = foo; foo = 1.0; test_matrix_within_struct_accesses(baz); metal::float4x3 matrix = bar.matrix; type_6 arr = bar.arr; float b = bar.matrix[3].x; - int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value; + int a_1 = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value; float _e28 = read_from_private(foo); - for(int _i=0; _i<5; ++_i) c.inner[_i] = type_15 {a, static_cast(b), 3, 4, 5}.inner[_i]; + for(int _i=0; _i<5; ++_i) c.inner[_i] = type_17 {a_1, static_cast(b), 3, 4, 5}.inner[_i]; c.inner[vi + 1u] = 42; int value = c.inner[vi]; + float _e42 = test_arr_as_arg(const_type_13_); + unnamed = _e42; return foo_vertOutput { metal::float4(matrix * static_cast(metal::int4(value)), 2.0) }; } diff --git a/tests/out/msl/padding.msl b/tests/out/msl/padding.msl new file mode 100644 index 0000000000..23aef35813 --- /dev/null +++ b/tests/out/msl/padding.msl @@ -0,0 +1,38 @@ +// language: metal2.0 +#include +#include + +using metal::uint; + +struct S { + metal::float3 a; +}; +struct Test { + S a; + float b; +}; +struct type_2 { + metal::float3 inner[2]; +}; +struct Test2_ { + type_2 a; + float b; +}; +struct Test3_ { + metal::float4x3 a; + float b; +}; + +struct vertex_Output { + metal::float4 member [[position]]; +}; +vertex vertex_Output vertex_( + constant Test& input [[buffer(0)]] +, constant Test2_& input2_ [[buffer(1)]] +, constant Test3_& input3_ [[buffer(2)]] +) { + float _e6 = input.b; + float _e9 = input2_.b; + float _e12 = input3_.b; + return vertex_Output { ((metal::float4(1.0) * _e6) * _e9) * _e12 }; +} diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index 0f89593901..41d9c1eda3 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,75 +1,80 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 213 +; Bound: 227 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %137 "foo_vert" %132 %135 -OpEntryPoint Fragment %173 "foo_frag" %172 -OpEntryPoint GLCompute %190 "atomics" -OpExecutionMode %173 OriginUpperLeft -OpExecutionMode %190 LocalSize 1 1 1 +OpEntryPoint Vertex %150 "foo_vert" %145 %148 +OpEntryPoint Fragment %187 "foo_frag" %186 +OpEntryPoint GLCompute %204 "atomics" +OpExecutionMode %187 OriginUpperLeft +OpExecutionMode %204 LocalSize 1 1 1 OpSource GLSL 450 -OpMemberName %31 0 "value" -OpName %31 "AlignedWrapper" -OpMemberName %40 0 "matrix" -OpMemberName %40 1 "matrix_array" -OpMemberName %40 2 "atom" -OpMemberName %40 3 "arr" -OpMemberName %40 4 "data" -OpName %40 "Bar" -OpMemberName %42 0 "m" -OpName %42 "Baz" -OpName %47 "bar" -OpName %49 "baz" -OpName %52 "idx" -OpName %54 "_" -OpName %56 "_" -OpName %58 "_" +OpMemberName %32 0 "value" +OpName %32 "AlignedWrapper" +OpMemberName %41 0 "matrix" +OpMemberName %41 1 "matrix_array" +OpMemberName %41 2 "atom" +OpMemberName %41 3 "arr" +OpMemberName %41 4 "data" +OpName %41 "Bar" +OpMemberName %43 0 "m" +OpName %43 "Baz" +OpName %52 "bar" +OpName %54 "baz" +OpName %57 "idx" OpName %59 "_" -OpName %60 "_" OpName %61 "_" -OpName %62 "_" -OpName %63 "t" -OpName %66 "test_matrix_within_struct_accesses" -OpName %123 "foo" -OpName %124 "read_from_private" +OpName %63 "_" +OpName %64 "_" +OpName %65 "_" +OpName %66 "_" +OpName %67 "_" +OpName %68 "t" +OpName %71 "test_matrix_within_struct_accesses" OpName %128 "foo" -OpName %129 "c" -OpName %132 "vi" -OpName %137 "foo_vert" -OpName %173 "foo_frag" -OpName %188 "tmp" -OpName %190 "atomics" -OpMemberDecorate %31 0 Offset 0 -OpDecorate %36 ArrayStride 16 -OpDecorate %38 ArrayStride 8 +OpName %129 "read_from_private" +OpName %134 "a" +OpName %135 "test_arr_as_arg" +OpName %140 "foo" +OpName %141 "c" +OpName %143 "_" +OpName %145 "vi" +OpName %150 "foo_vert" +OpName %187 "foo_frag" +OpName %202 "tmp" +OpName %204 "atomics" +OpMemberDecorate %32 0 Offset 0 +OpDecorate %37 ArrayStride 16 OpDecorate %39 ArrayStride 8 -OpMemberDecorate %40 0 Offset 0 -OpMemberDecorate %40 0 ColMajor -OpMemberDecorate %40 0 MatrixStride 16 -OpMemberDecorate %40 1 Offset 64 -OpMemberDecorate %40 1 ColMajor -OpMemberDecorate %40 1 MatrixStride 8 -OpMemberDecorate %40 2 Offset 96 -OpMemberDecorate %40 3 Offset 104 -OpMemberDecorate %40 4 Offset 120 -OpMemberDecorate %42 0 Offset 0 -OpMemberDecorate %42 0 ColMajor -OpMemberDecorate %42 0 MatrixStride 8 -OpDecorate %46 ArrayStride 4 -OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 0 -OpDecorate %40 Block -OpDecorate %49 DescriptorSet 0 -OpDecorate %49 Binding 1 -OpDecorate %50 Block -OpMemberDecorate %50 0 Offset 0 -OpDecorate %132 BuiltIn VertexIndex -OpDecorate %135 BuiltIn Position -OpDecorate %172 Location 0 +OpDecorate %40 ArrayStride 8 +OpMemberDecorate %41 0 Offset 0 +OpMemberDecorate %41 0 ColMajor +OpMemberDecorate %41 0 MatrixStride 16 +OpMemberDecorate %41 1 Offset 64 +OpMemberDecorate %41 1 ColMajor +OpMemberDecorate %41 1 MatrixStride 8 +OpMemberDecorate %41 2 Offset 96 +OpMemberDecorate %41 3 Offset 104 +OpMemberDecorate %41 4 Offset 120 +OpMemberDecorate %43 0 Offset 0 +OpMemberDecorate %43 0 ColMajor +OpMemberDecorate %43 0 MatrixStride 8 +OpDecorate %45 ArrayStride 4 +OpDecorate %46 ArrayStride 40 +OpDecorate %49 ArrayStride 4 +OpDecorate %52 DescriptorSet 0 +OpDecorate %52 Binding 0 +OpDecorate %41 Block +OpDecorate %54 DescriptorSet 0 +OpDecorate %54 Binding 1 +OpDecorate %55 Block +OpMemberDecorate %55 0 Offset 0 +OpDecorate %145 BuiltIn VertexIndex +OpDecorate %148 BuiltIn Position +OpDecorate %186 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -89,245 +94,263 @@ OpDecorate %172 Location 0 %18 = OpConstant %9 20.0 %19 = OpConstant %9 30.0 %20 = OpConstant %9 40.0 -%21 = OpConstant %9 0.0 -%23 = OpTypeInt 32 0 -%22 = OpConstant %23 3 -%24 = OpConstant %23 2 -%25 = OpConstant %4 5 -%26 = OpConstant %4 3 -%27 = OpConstant %4 4 -%28 = OpConstant %23 1 -%29 = OpConstant %4 42 -%30 = OpConstant %23 0 -%31 = OpTypeStruct %4 -%33 = OpTypeVector %9 3 -%32 = OpTypeMatrix %33 4 -%35 = OpTypeVector %9 2 -%34 = OpTypeMatrix %35 2 -%36 = OpTypeArray %34 %3 -%37 = OpTypeVector %23 2 -%38 = OpTypeArray %37 %3 -%39 = OpTypeRuntimeArray %31 -%40 = OpTypeStruct %32 %36 %4 %38 %39 -%41 = OpTypeMatrix %35 3 -%42 = OpTypeStruct %41 -%43 = OpTypePointer Function %9 -%44 = OpTypeVector %9 4 -%45 = OpTypePointer StorageBuffer %4 -%46 = OpTypeArray %4 %25 -%48 = OpTypePointer StorageBuffer %40 -%47 = OpVariable %48 StorageBuffer -%50 = OpTypeStruct %42 -%51 = OpTypePointer Uniform %50 -%49 = OpVariable %51 Uniform -%53 = OpTypePointer Function %4 -%55 = OpTypePointer Function %41 -%57 = OpTypePointer Function %35 -%64 = OpTypePointer Function %42 -%67 = OpTypeFunction %2 -%68 = OpTypePointer Uniform %42 -%73 = OpTypePointer Uniform %41 -%76 = OpTypePointer Uniform %35 -%82 = OpTypePointer Uniform %9 -%107 = OpTypePointer Function %35 -%113 = OpTypePointer Function %9 -%125 = OpTypeFunction %9 %43 -%130 = OpTypePointer Function %46 -%133 = OpTypePointer Input %23 -%132 = OpVariable %133 Input -%136 = OpTypePointer Output %44 -%135 = OpVariable %136 Output -%142 = OpTypePointer StorageBuffer %32 -%145 = OpTypePointer StorageBuffer %38 -%148 = OpTypePointer StorageBuffer %33 -%149 = OpTypePointer StorageBuffer %9 -%152 = OpTypePointer StorageBuffer %39 -%155 = OpTypePointer StorageBuffer %31 -%156 = OpConstant %23 4 -%166 = OpTypeVector %4 4 -%172 = OpVariable %136 Output -%192 = OpTypePointer StorageBuffer %4 -%195 = OpConstant %23 64 -%66 = OpFunction %2 None %67 -%65 = OpLabel -%61 = OpVariable %43 Function -%58 = OpVariable %57 Function -%52 = OpVariable %53 Function %5 -%62 = OpVariable %43 Function -%59 = OpVariable %43 Function -%54 = OpVariable %55 Function -%63 = OpVariable %64 Function -%60 = OpVariable %43 Function -%56 = OpVariable %57 Function -%69 = OpAccessChain %68 %49 %30 -OpBranch %70 +%21 = OpConstant %4 10 +%22 = OpConstant %4 5 +%23 = OpConstant %4 4 +%24 = OpConstant %9 0.0 +%26 = OpTypeInt 32 0 +%25 = OpConstant %26 3 +%27 = OpConstant %26 2 +%28 = OpConstant %4 3 +%29 = OpConstant %26 1 +%30 = OpConstant %4 42 +%31 = OpConstant %26 0 +%32 = OpTypeStruct %4 +%34 = OpTypeVector %9 3 +%33 = OpTypeMatrix %34 4 +%36 = OpTypeVector %9 2 +%35 = OpTypeMatrix %36 2 +%37 = OpTypeArray %35 %3 +%38 = OpTypeVector %26 2 +%39 = OpTypeArray %38 %3 +%40 = OpTypeRuntimeArray %32 +%41 = OpTypeStruct %33 %37 %4 %39 %40 +%42 = OpTypeMatrix %36 3 +%43 = OpTypeStruct %42 +%44 = OpTypePointer Function %9 +%45 = OpTypeArray %9 %21 +%46 = OpTypeArray %45 %22 +%47 = OpTypeVector %9 4 +%48 = OpTypePointer StorageBuffer %4 +%49 = OpTypeArray %4 %22 +%50 = OpConstantComposite %45 %24 %24 %24 %24 %24 %24 %24 %24 %24 %24 +%51 = OpConstantComposite %46 %50 %50 %50 %50 %50 +%53 = OpTypePointer StorageBuffer %41 +%52 = OpVariable %53 StorageBuffer +%55 = OpTypeStruct %43 +%56 = OpTypePointer Uniform %55 +%54 = OpVariable %56 Uniform +%58 = OpTypePointer Function %4 +%60 = OpTypePointer Function %42 +%62 = OpTypePointer Function %36 +%69 = OpTypePointer Function %43 +%72 = OpTypeFunction %2 +%73 = OpTypePointer Uniform %43 +%78 = OpTypePointer Uniform %42 +%81 = OpTypePointer Uniform %36 +%87 = OpTypePointer Uniform %9 +%112 = OpTypePointer Function %36 +%118 = OpTypePointer Function %9 +%130 = OpTypeFunction %9 %44 +%136 = OpTypeFunction %9 %46 +%142 = OpTypePointer Function %49 +%146 = OpTypePointer Input %26 +%145 = OpVariable %146 Input +%149 = OpTypePointer Output %47 +%148 = OpVariable %149 Output +%155 = OpTypePointer StorageBuffer %33 +%158 = OpTypePointer StorageBuffer %39 +%161 = OpTypePointer StorageBuffer %34 +%162 = OpTypePointer StorageBuffer %9 +%165 = OpTypePointer StorageBuffer %40 +%168 = OpTypePointer StorageBuffer %32 +%169 = OpConstant %26 4 +%180 = OpTypeVector %4 4 +%186 = OpVariable %149 Output +%206 = OpTypePointer StorageBuffer %4 +%209 = OpConstant %26 64 +%71 = OpFunction %2 None %72 %70 = OpLabel -%71 = OpLoad %4 %52 -%72 = OpISub %4 %71 %6 -OpStore %52 %72 -%74 = OpAccessChain %73 %69 %30 -%75 = OpLoad %41 %74 -OpStore %54 %75 -%77 = OpAccessChain %76 %69 %30 %30 -%78 = OpLoad %35 %77 -OpStore %56 %78 -%79 = OpLoad %4 %52 -%80 = OpAccessChain %76 %69 %30 %79 -%81 = OpLoad %35 %80 -OpStore %58 %81 -%83 = OpAccessChain %82 %69 %30 %30 %28 -%84 = OpLoad %9 %83 -OpStore %59 %84 -%85 = OpLoad %4 %52 -%86 = OpAccessChain %82 %69 %30 %30 %85 -%87 = OpLoad %9 %86 -OpStore %60 %87 -%88 = OpLoad %4 %52 -%89 = OpAccessChain %82 %69 %30 %88 %28 -%90 = OpLoad %9 %89 -OpStore %61 %90 -%91 = OpLoad %4 %52 -%92 = OpLoad %4 %52 -%93 = OpAccessChain %82 %69 %30 %91 %92 -%94 = OpLoad %9 %93 -OpStore %62 %94 -%95 = OpCompositeConstruct %35 %8 %8 -%96 = OpCompositeConstruct %35 %10 %10 -%97 = OpCompositeConstruct %35 %11 %11 -%98 = OpCompositeConstruct %41 %95 %96 %97 -%99 = OpCompositeConstruct %42 %98 -OpStore %63 %99 -%100 = OpLoad %4 %52 -%101 = OpIAdd %4 %100 %6 -OpStore %52 %101 -%102 = OpCompositeConstruct %35 %12 %12 -%103 = OpCompositeConstruct %35 %13 %13 -%104 = OpCompositeConstruct %35 %14 %14 -%105 = OpCompositeConstruct %41 %102 %103 %104 -%106 = OpAccessChain %55 %63 %30 -OpStore %106 %105 -%108 = OpCompositeConstruct %35 %15 %15 -%109 = OpAccessChain %107 %63 %30 %30 -OpStore %109 %108 -%110 = OpLoad %4 %52 -%111 = OpCompositeConstruct %35 %16 %16 -%112 = OpAccessChain %107 %63 %30 %110 -OpStore %112 %111 -%114 = OpAccessChain %113 %63 %30 %30 %28 -OpStore %114 %17 -%115 = OpLoad %4 %52 -%116 = OpAccessChain %113 %63 %30 %30 %115 -OpStore %116 %18 -%117 = OpLoad %4 %52 -%118 = OpAccessChain %113 %63 %30 %117 %28 -OpStore %118 %19 -%119 = OpLoad %4 %52 -%120 = OpLoad %4 %52 -%121 = OpAccessChain %113 %63 %30 %119 %120 -OpStore %121 %20 +%66 = OpVariable %44 Function +%63 = OpVariable %62 Function +%57 = OpVariable %58 Function %5 +%67 = OpVariable %44 Function +%64 = OpVariable %44 Function +%59 = OpVariable %60 Function +%68 = OpVariable %69 Function +%65 = OpVariable %44 Function +%61 = OpVariable %62 Function +%74 = OpAccessChain %73 %54 %31 +OpBranch %75 +%75 = OpLabel +%76 = OpLoad %4 %57 +%77 = OpISub %4 %76 %6 +OpStore %57 %77 +%79 = OpAccessChain %78 %74 %31 +%80 = OpLoad %42 %79 +OpStore %59 %80 +%82 = OpAccessChain %81 %74 %31 %31 +%83 = OpLoad %36 %82 +OpStore %61 %83 +%84 = OpLoad %4 %57 +%85 = OpAccessChain %81 %74 %31 %84 +%86 = OpLoad %36 %85 +OpStore %63 %86 +%88 = OpAccessChain %87 %74 %31 %31 %29 +%89 = OpLoad %9 %88 +OpStore %64 %89 +%90 = OpLoad %4 %57 +%91 = OpAccessChain %87 %74 %31 %31 %90 +%92 = OpLoad %9 %91 +OpStore %65 %92 +%93 = OpLoad %4 %57 +%94 = OpAccessChain %87 %74 %31 %93 %29 +%95 = OpLoad %9 %94 +OpStore %66 %95 +%96 = OpLoad %4 %57 +%97 = OpLoad %4 %57 +%98 = OpAccessChain %87 %74 %31 %96 %97 +%99 = OpLoad %9 %98 +OpStore %67 %99 +%100 = OpCompositeConstruct %36 %8 %8 +%101 = OpCompositeConstruct %36 %10 %10 +%102 = OpCompositeConstruct %36 %11 %11 +%103 = OpCompositeConstruct %42 %100 %101 %102 +%104 = OpCompositeConstruct %43 %103 +OpStore %68 %104 +%105 = OpLoad %4 %57 +%106 = OpIAdd %4 %105 %6 +OpStore %57 %106 +%107 = OpCompositeConstruct %36 %12 %12 +%108 = OpCompositeConstruct %36 %13 %13 +%109 = OpCompositeConstruct %36 %14 %14 +%110 = OpCompositeConstruct %42 %107 %108 %109 +%111 = OpAccessChain %60 %68 %31 +OpStore %111 %110 +%113 = OpCompositeConstruct %36 %15 %15 +%114 = OpAccessChain %112 %68 %31 %31 +OpStore %114 %113 +%115 = OpLoad %4 %57 +%116 = OpCompositeConstruct %36 %16 %16 +%117 = OpAccessChain %112 %68 %31 %115 +OpStore %117 %116 +%119 = OpAccessChain %118 %68 %31 %31 %29 +OpStore %119 %17 +%120 = OpLoad %4 %57 +%121 = OpAccessChain %118 %68 %31 %31 %120 +OpStore %121 %18 +%122 = OpLoad %4 %57 +%123 = OpAccessChain %118 %68 %31 %122 %29 +OpStore %123 %19 +%124 = OpLoad %4 %57 +%125 = OpLoad %4 %57 +%126 = OpAccessChain %118 %68 %31 %124 %125 +OpStore %126 %20 OpReturn OpFunctionEnd -%124 = OpFunction %9 None %125 -%123 = OpFunctionParameter %43 -%122 = OpLabel -OpBranch %126 -%126 = OpLabel -%127 = OpLoad %9 %123 -OpReturnValue %127 -OpFunctionEnd -%137 = OpFunction %2 None %67 +%129 = OpFunction %9 None %130 +%128 = OpFunctionParameter %44 +%127 = OpLabel +OpBranch %131 %131 = OpLabel -%128 = OpVariable %43 Function %21 -%129 = OpVariable %130 Function -%134 = OpLoad %23 %132 -%138 = OpAccessChain %68 %49 %30 -OpBranch %139 -%139 = OpLabel -%140 = OpLoad %9 %128 -OpStore %128 %8 -%141 = OpFunctionCall %2 %66 -%143 = OpAccessChain %142 %47 %30 -%144 = OpLoad %32 %143 -%146 = OpAccessChain %145 %47 %22 -%147 = OpLoad %38 %146 -%150 = OpAccessChain %149 %47 %30 %22 %30 -%151 = OpLoad %9 %150 -%153 = OpArrayLength %23 %47 4 -%154 = OpISub %23 %153 %24 -%157 = OpAccessChain %45 %47 %156 %154 %30 -%158 = OpLoad %4 %157 -%159 = OpFunctionCall %9 %124 %128 -%160 = OpConvertFToS %4 %151 -%161 = OpCompositeConstruct %46 %158 %160 %26 %27 %25 -OpStore %129 %161 -%162 = OpIAdd %23 %134 %28 -%163 = OpAccessChain %53 %129 %162 -OpStore %163 %29 -%164 = OpAccessChain %53 %129 %134 -%165 = OpLoad %4 %164 -%167 = OpCompositeConstruct %166 %165 %165 %165 %165 -%168 = OpConvertSToF %44 %167 -%169 = OpMatrixTimesVector %33 %144 %168 -%170 = OpCompositeConstruct %44 %169 %10 -OpStore %135 %170 +%132 = OpLoad %9 %128 +OpReturnValue %132 +OpFunctionEnd +%135 = OpFunction %9 None %136 +%134 = OpFunctionParameter %46 +%133 = OpLabel +OpBranch %137 +%137 = OpLabel +%138 = OpCompositeExtract %45 %134 4 +%139 = OpCompositeExtract %9 %138 9 +OpReturnValue %139 +OpFunctionEnd +%150 = OpFunction %2 None %72 +%144 = OpLabel +%140 = OpVariable %44 Function %24 +%141 = OpVariable %142 Function +%143 = OpVariable %44 Function +%147 = OpLoad %26 %145 +%151 = OpAccessChain %73 %54 %31 +OpBranch %152 +%152 = OpLabel +%153 = OpLoad %9 %140 +OpStore %140 %8 +%154 = OpFunctionCall %2 %71 +%156 = OpAccessChain %155 %52 %31 +%157 = OpLoad %33 %156 +%159 = OpAccessChain %158 %52 %25 +%160 = OpLoad %39 %159 +%163 = OpAccessChain %162 %52 %31 %25 %31 +%164 = OpLoad %9 %163 +%166 = OpArrayLength %26 %52 4 +%167 = OpISub %26 %166 %27 +%170 = OpAccessChain %48 %52 %169 %167 %31 +%171 = OpLoad %4 %170 +%172 = OpFunctionCall %9 %129 %140 +%173 = OpConvertFToS %4 %164 +%174 = OpCompositeConstruct %49 %171 %173 %28 %23 %22 +OpStore %141 %174 +%175 = OpIAdd %26 %147 %29 +%176 = OpAccessChain %58 %141 %175 +OpStore %176 %30 +%177 = OpAccessChain %58 %141 %147 +%178 = OpLoad %4 %177 +%179 = OpFunctionCall %9 %135 %51 +OpStore %143 %179 +%181 = OpCompositeConstruct %180 %178 %178 %178 %178 +%182 = OpConvertSToF %47 %181 +%183 = OpMatrixTimesVector %34 %157 %182 +%184 = OpCompositeConstruct %47 %183 %10 +OpStore %148 %184 OpReturn OpFunctionEnd -%173 = OpFunction %2 None %67 -%171 = OpLabel -OpBranch %174 -%174 = OpLabel -%175 = OpAccessChain %149 %47 %30 %28 %24 -OpStore %175 %8 -%176 = OpCompositeConstruct %33 %21 %21 %21 -%177 = OpCompositeConstruct %33 %8 %8 %8 -%178 = OpCompositeConstruct %33 %10 %10 %10 -%179 = OpCompositeConstruct %33 %11 %11 %11 -%180 = OpCompositeConstruct %32 %176 %177 %178 %179 -%181 = OpAccessChain %142 %47 %30 -OpStore %181 %180 -%182 = OpCompositeConstruct %37 %30 %30 -%183 = OpCompositeConstruct %37 %28 %28 -%184 = OpCompositeConstruct %38 %182 %183 -%185 = OpAccessChain %145 %47 %22 -OpStore %185 %184 -%186 = OpAccessChain %45 %47 %156 %28 %30 -OpStore %186 %6 -%187 = OpCompositeConstruct %44 %21 %21 %21 %21 -OpStore %172 %187 +%187 = OpFunction %2 None %72 +%185 = OpLabel +OpBranch %188 +%188 = OpLabel +%189 = OpAccessChain %162 %52 %31 %29 %27 +OpStore %189 %8 +%190 = OpCompositeConstruct %34 %24 %24 %24 +%191 = OpCompositeConstruct %34 %8 %8 %8 +%192 = OpCompositeConstruct %34 %10 %10 %10 +%193 = OpCompositeConstruct %34 %11 %11 %11 +%194 = OpCompositeConstruct %33 %190 %191 %192 %193 +%195 = OpAccessChain %155 %52 %31 +OpStore %195 %194 +%196 = OpCompositeConstruct %38 %31 %31 +%197 = OpCompositeConstruct %38 %29 %29 +%198 = OpCompositeConstruct %39 %196 %197 +%199 = OpAccessChain %158 %52 %25 +OpStore %199 %198 +%200 = OpAccessChain %48 %52 %169 %29 %31 +OpStore %200 %6 +%201 = OpCompositeConstruct %47 %24 %24 %24 %24 +OpStore %186 %201 OpReturn OpFunctionEnd -%190 = OpFunction %2 None %67 -%189 = OpLabel -%188 = OpVariable %53 Function -OpBranch %191 -%191 = OpLabel -%193 = OpAccessChain %192 %47 %24 -%194 = OpAtomicLoad %4 %193 %6 %195 -%197 = OpAccessChain %192 %47 %24 -%196 = OpAtomicIAdd %4 %197 %6 %195 %25 -OpStore %188 %196 -%199 = OpAccessChain %192 %47 %24 -%198 = OpAtomicISub %4 %199 %6 %195 %25 -OpStore %188 %198 -%201 = OpAccessChain %192 %47 %24 -%200 = OpAtomicAnd %4 %201 %6 %195 %25 -OpStore %188 %200 -%203 = OpAccessChain %192 %47 %24 -%202 = OpAtomicOr %4 %203 %6 %195 %25 -OpStore %188 %202 -%205 = OpAccessChain %192 %47 %24 -%204 = OpAtomicXor %4 %205 %6 %195 %25 -OpStore %188 %204 -%207 = OpAccessChain %192 %47 %24 -%206 = OpAtomicSMin %4 %207 %6 %195 %25 -OpStore %188 %206 -%209 = OpAccessChain %192 %47 %24 -%208 = OpAtomicSMax %4 %209 %6 %195 %25 -OpStore %188 %208 -%211 = OpAccessChain %192 %47 %24 -%210 = OpAtomicExchange %4 %211 %6 %195 %25 -OpStore %188 %210 -%212 = OpAccessChain %192 %47 %24 -OpAtomicStore %212 %6 %195 %194 +%204 = OpFunction %2 None %72 +%203 = OpLabel +%202 = OpVariable %58 Function +OpBranch %205 +%205 = OpLabel +%207 = OpAccessChain %206 %52 %27 +%208 = OpAtomicLoad %4 %207 %6 %209 +%211 = OpAccessChain %206 %52 %27 +%210 = OpAtomicIAdd %4 %211 %6 %209 %22 +OpStore %202 %210 +%213 = OpAccessChain %206 %52 %27 +%212 = OpAtomicISub %4 %213 %6 %209 %22 +OpStore %202 %212 +%215 = OpAccessChain %206 %52 %27 +%214 = OpAtomicAnd %4 %215 %6 %209 %22 +OpStore %202 %214 +%217 = OpAccessChain %206 %52 %27 +%216 = OpAtomicOr %4 %217 %6 %209 %22 +OpStore %202 %216 +%219 = OpAccessChain %206 %52 %27 +%218 = OpAtomicXor %4 %219 %6 %209 %22 +OpStore %202 %218 +%221 = OpAccessChain %206 %52 %27 +%220 = OpAtomicSMin %4 %221 %6 %209 %22 +OpStore %202 %220 +%223 = OpAccessChain %206 %52 %27 +%222 = OpAtomicSMax %4 %223 %6 %209 %22 +OpStore %202 %222 +%225 = OpAccessChain %206 %52 %27 +%224 = OpAtomicExchange %4 %225 %6 %209 %22 +OpStore %202 %224 +%226 = OpAccessChain %206 %52 %27 +OpAtomicStore %226 %6 %209 %208 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/padding.spvasm b/tests/out/spv/padding.spvasm new file mode 100644 index 0000000000..d262f4b08a --- /dev/null +++ b/tests/out/spv/padding.spvasm @@ -0,0 +1,99 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 50 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint Vertex %27 "vertex" %25 +OpSource GLSL 450 +OpMemberName %8 0 "a" +OpName %8 "S" +OpMemberName %9 0 "a" +OpMemberName %9 1 "b" +OpName %9 "Test" +OpMemberName %11 0 "a" +OpMemberName %11 1 "b" +OpName %11 "Test2" +OpMemberName %13 0 "a" +OpMemberName %13 1 "b" +OpName %13 "Test3" +OpName %15 "input" +OpName %18 "input2" +OpName %21 "input3" +OpName %27 "vertex" +OpMemberDecorate %8 0 Offset 0 +OpMemberDecorate %9 0 Offset 0 +OpMemberDecorate %9 1 Offset 16 +OpDecorate %10 ArrayStride 16 +OpMemberDecorate %11 0 Offset 0 +OpMemberDecorate %11 1 Offset 32 +OpMemberDecorate %13 0 Offset 0 +OpMemberDecorate %13 0 ColMajor +OpMemberDecorate %13 0 MatrixStride 16 +OpMemberDecorate %13 1 Offset 64 +OpDecorate %15 DescriptorSet 0 +OpDecorate %15 Binding 0 +OpDecorate %16 Block +OpMemberDecorate %16 0 Offset 0 +OpDecorate %18 DescriptorSet 0 +OpDecorate %18 Binding 1 +OpDecorate %19 Block +OpMemberDecorate %19 0 Offset 0 +OpDecorate %21 DescriptorSet 0 +OpDecorate %21 Binding 2 +OpDecorate %22 Block +OpMemberDecorate %22 0 Offset 0 +OpDecorate %25 BuiltIn Position +%2 = OpTypeVoid +%4 = OpTypeInt 32 1 +%3 = OpConstant %4 2 +%6 = OpTypeFloat 32 +%5 = OpConstant %6 1.0 +%7 = OpTypeVector %6 3 +%8 = OpTypeStruct %7 +%9 = OpTypeStruct %8 %6 +%10 = OpTypeArray %7 %3 +%11 = OpTypeStruct %10 %6 +%12 = OpTypeMatrix %7 4 +%13 = OpTypeStruct %12 %6 +%14 = OpTypeVector %6 4 +%16 = OpTypeStruct %9 +%17 = OpTypePointer Uniform %16 +%15 = OpVariable %17 Uniform +%19 = OpTypeStruct %11 +%20 = OpTypePointer Uniform %19 +%18 = OpVariable %20 Uniform +%22 = OpTypeStruct %13 +%23 = OpTypePointer Uniform %22 +%21 = OpVariable %23 Uniform +%26 = OpTypePointer Output %14 +%25 = OpVariable %26 Output +%28 = OpTypeFunction %2 +%29 = OpTypePointer Uniform %9 +%31 = OpTypeInt 32 0 +%30 = OpConstant %31 0 +%33 = OpTypePointer Uniform %11 +%35 = OpTypePointer Uniform %13 +%39 = OpTypePointer Uniform %6 +%40 = OpConstant %31 1 +%27 = OpFunction %2 None %28 +%24 = OpLabel +%32 = OpAccessChain %29 %15 %30 +%34 = OpAccessChain %33 %18 %30 +%36 = OpAccessChain %35 %21 %30 +OpBranch %37 +%37 = OpLabel +%38 = OpCompositeConstruct %14 %5 %5 %5 %5 +%41 = OpAccessChain %39 %32 %40 +%42 = OpLoad %6 %41 +%43 = OpVectorTimesScalar %14 %38 %42 +%44 = OpAccessChain %39 %34 %40 +%45 = OpLoad %6 %44 +%46 = OpVectorTimesScalar %14 %43 %45 +%47 = OpAccessChain %39 %36 %40 +%48 = OpLoad %6 %47 +%49 = OpVectorTimesScalar %14 %46 %48 +OpStore %25 %49 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 53cd3c4ac2..b6fc1921fc 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -21,36 +21,36 @@ var baz: Baz; fn test_matrix_within_struct_accesses() { var idx: i32 = 9; - var unnamed: mat3x2; - var unnamed_1: vec2; + var unnamed_1: mat3x2; var unnamed_2: vec2; - var unnamed_3: f32; + var unnamed_3: vec2; var unnamed_4: f32; var unnamed_5: f32; var unnamed_6: f32; + var unnamed_7: f32; var t: Baz; let _e4 = idx; idx = (_e4 - 1); let _e8 = baz.m; - unnamed = _e8; + unnamed_1 = _e8; let _e13 = baz.m[0]; - unnamed_1 = _e13; + unnamed_2 = _e13; let _e16 = idx; let _e18 = baz.m[_e16]; - unnamed_2 = _e18; + unnamed_3 = _e18; let _e25 = baz.m[0][1]; - unnamed_3 = _e25; + unnamed_4 = _e25; let _e30 = idx; let _e32 = baz.m[0][_e30]; - unnamed_4 = _e32; + unnamed_5 = _e32; let _e35 = idx; let _e39 = baz.m[_e35][1]; - unnamed_5 = _e39; + unnamed_6 = _e39; let _e42 = idx; let _e44 = idx; let _e46 = baz.m[_e42][_e44]; - unnamed_6 = _e46; + unnamed_7 = _e46; t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0))); let _e57 = idx; idx = (_e57 + 1); @@ -74,10 +74,15 @@ fn read_from_private(foo_1: ptr) -> f32 { return _e3; } +fn test_arr_as_arg(a: array,5>) -> f32 { + return a[4][9]; +} + @stage(vertex) fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0.0; var c: array; + var unnamed: f32; let baz_1 = foo; foo = 1.0; @@ -85,12 +90,14 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let matrix = bar.matrix; let arr = bar.arr; let b = bar.matrix[3][0]; - let a = bar.data[(arrayLength((&bar.data)) - 2u)].value; + let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value; let data_pointer = (&bar.data[0].value); let _e28 = read_from_private((&foo)); - c = array(a, i32(b), 3, 4, 5); + c = array(a_1, i32(b), 3, 4, 5); c[(vi + 1u)] = 42; let value = c[vi]; + let _e42 = test_arr_as_arg(array,5>(array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0))); + unnamed = _e42; return vec4((matrix * vec4(vec4(value))), 2.0); } diff --git a/tests/out/wgsl/padding.wgsl b/tests/out/wgsl/padding.wgsl new file mode 100644 index 0000000000..1d3495477e --- /dev/null +++ b/tests/out/wgsl/padding.wgsl @@ -0,0 +1,33 @@ +struct S { + a: vec3, +} + +struct Test { + a: S, + b: f32, +} + +struct Test2_ { + a: array,2>, + b: f32, +} + +struct Test3_ { + a: mat4x3, + b: f32, +} + +@group(0) @binding(0) +var input: Test; +@group(0) @binding(1) +var input2_: Test2_; +@group(0) @binding(2) +var input3_: Test3_; + +@stage(vertex) +fn vertex() -> @builtin(position) vec4 { + let _e6 = input.b; + let _e9 = input2_.b; + let _e12 = input3_.b; + return (((vec4(1.0) * _e6) * _e9) * _e12); +} diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 279a79c70d..0726a4846f 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -458,6 +458,10 @@ fn convert_wgsl() { "access", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ), + ( + "padding", + Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, + ), ("pointers", Targets::SPIRV | Targets::WGSL), ( "control-flow",