github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

[hlsl-out] Add more padding when necessary (#1814)

Browse Source 
* [hlsl-out] add padding at the end of structs and after struct members of type matrix and array (when necessary)

* use wrapped constructor fn for constants

* add array as fn arg test

* fix glsl array fn arg

* add wrapped constructor for arrays

* [glsl-out] support multidimensional arrays

* address comments
This commit is contained in:
Teodor Tanasoaia
2022-04-12 05:34:06 +02:00
committed by GitHub
parent 0e77d26241
commit bd62887a13
24 changed files with 1088 additions and 450 deletions
Download Patch File Download Diff File Expand all files Collapse all files

47
src/back/glsl/mod.rs
View File

@@ -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<crate::Type>,
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)?;

75
src/back/hlsl/conv.rs
View File

@@ -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<crate::Type>,
constants: &crate::Arena<crate::Constant>,
) -> Result<u32, crate::arena::BadHandle> {
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<crate::Type>,
types: &crate::UniqueArena<crate::Type>,
constants: &crate::Arena<crate::Constant>,
names: &'a crate::FastHashMap<crate::proc::NameKey, String>,
) -> Result<Cow<'a, str>, 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 {

182
src/back/hlsl/help.rs
View File

@@ -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, ") {{")?;
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];
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)];
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
)?;
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,

2
src/back/hlsl/storage.rs
View File

@@ -307,7 +307,7 @@ impl<W: fmt::Write> 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, ";")?;

95
src/back/hlsl/writer.rs
View File

@@ -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<crate::Type>,
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<crate::Type>,
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<crate::Type>,
components: &[Handle<crate::Constant>],
) -> 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)?;

6
tests/in/access.wgsl
View File

@@ -54,6 +54,10 @@ fn read_from_private(foo: ptr<function, f32>) -> f32 {
return *foo;
}
fn test_arr_as_arg(a: array<array<f32, 10>, 5>) -> f32 {
return a[4][9];
}
@stage(vertex)
fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
var foo: f32 = 0.0;
@@ -79,6 +83,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
c[vi + 1u] = 42;
let value = c[vi];
var _ = test_arr_as_arg(array<array<f32, 10>, 5>());
return vec4<f32>(matrix * vec4<f32>(vec4<i32>(value)), 2.0);
}

22
tests/in/padding.param.ron Normal file
View File

@@ -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,
),
)

33
tests/in/padding.wgsl Normal file
View File

@@ -0,0 +1,33 @@
struct S {
a: vec3<f32>,
}
struct Test {
a: S,
b: f32, // offset: 16
}
struct Test2 {
a: array<vec3<f32>, 2>,
b: f32, // offset: 32
}
struct Test3 {
a: mat4x3<f32>,
b: f32, // offset: 64
}
@group(0) @binding(0)
var<uniform> input: Test;
@group(0) @binding(1)
var<uniform> input2: Test2;
@group(0) @binding(2)
var<uniform> input3: Test3;
@stage(vertex)
fn vertex() -> @builtin(position) vec4<f32> {
return vec4<f32>(1.0) * input.b * input2.b * input3.b;
}

4
tests/out/glsl/access.atomics.Compute.glsl
View File

@@ -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;

4
tests/out/glsl/access.foo_frag.Fragment.glsl
View File

@@ -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));

31
tests/out/glsl/access.foo_vert.Vertex.glsl
View File

@@ -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;

36
tests/out/glsl/padding.vertex.Vertex.glsl Normal file
View File

@@ -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;
}

55
tests/out/hlsl/access.hlsl
View File

@@ -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]));
}

27
tests/out/hlsl/operators.hlsl
View File

@@ -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;

43
tests/out/hlsl/padding.hlsl Normal file
View File

@@ -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);
}

3
tests/out/hlsl/padding.hlsl.config Normal file
View File

@@ -0,0 +1,3 @@
vertex=(vertex:vs_5_1 )
fragment=()
compute=()

16
tests/out/hlsl/quad-vert.hlsl
View File

@@ -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 {

45
tests/out/msl/access.msl
View File

@@ -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<int>(b), 3, 4, 5}.inner[_i];
for(int _i=0; _i<5; ++_i) c.inner[_i] = type_17 {a_1, static_cast<int>(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::float4>(metal::int4(value)), 2.0) };
}

38
tests/out/msl/padding.msl Normal file
View File

@@ -0,0 +1,38 @@
// language: metal2.0
#include <metal_stdlib>
#include <simd/simd.h>
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 };
}

607
tests/out/spv/access.spvasm
View File

@@ -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

99
tests/out/spv/padding.spvasm Normal file
View File

@@ -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

31
tests/out/wgsl/access.wgsl
View File

@@ -21,36 +21,36 @@ var<uniform> baz: Baz;
fn test_matrix_within_struct_accesses() {
var idx: i32 = 9;
var unnamed: mat3x2<f32>;
var unnamed_1: vec2<f32>;
var unnamed_1: mat3x2<f32>;
var unnamed_2: vec2<f32>;
var unnamed_3: f32;
var unnamed_3: vec2<f32>;
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<f32>(vec2<f32>(1.0), vec2<f32>(2.0), vec2<f32>(3.0)));
let _e57 = idx;
idx = (_e57 + 1);
@@ -74,10 +74,15 @@ fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
return _e3;
}
fn test_arr_as_arg(a: array<array<f32,10>,5>) -> f32 {
return a[4][9];
}
@stage(vertex)
fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
var foo: f32 = 0.0;
var c: array<i32,5>;
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<f32> {
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<i32,5>(a, i32(b), 3, 4, 5);
c = array<i32,5>(a_1, i32(b), 3, 4, 5);
c[(vi + 1u)] = 42;
let value = c[vi];
let _e42 = test_arr_as_arg(array<array<f32,10>,5>(array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
unnamed = _e42;
return vec4<f32>((matrix * vec4<f32>(vec4<i32>(value))), 2.0);
}

33
tests/out/wgsl/padding.wgsl Normal file
View File

@@ -0,0 +1,33 @@
struct S {
a: vec3<f32>,
}
struct Test {
a: S,
b: f32,
}
struct Test2_ {
a: array<vec3<f32>,2>,
b: f32,
}
struct Test3_ {
a: mat4x3<f32>,
b: f32,
}
@group(0) @binding(0)
var<uniform> input: Test;
@group(0) @binding(1)
var<uniform> input2_: Test2_;
@group(0) @binding(2)
var<uniform> input3_: Test3_;
@stage(vertex)
fn vertex() -> @builtin(position) vec4<f32> {
let _e6 = input.b;
let _e9 = input2_.b;
let _e12 = input3_.b;
return (((vec4<f32>(1.0) * _e6) * _e9) * _e12);
}

4
tests/snapshots.rs
View File

@@ -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",