diff --git a/src/back/hlsl/help.rs b/src/back/hlsl/help.rs index 91553a508b..c652a1d6e3 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -48,6 +48,12 @@ pub(super) struct WrappedConstructor { pub(super) ty: Handle, } +#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] +pub(super) struct WrappedStructMatrixAccess { + pub(super) ty: Handle, + pub(super) index: u32, +} + /// HLSL backend requires its own `ImageQuery` enum. /// /// It is used inside `WrappedImageQuery` and should be unique per ImageQuery function. @@ -393,13 +399,39 @@ impl<'a, W: Write> super::Writer<'a, W> { 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)]; - //TODO: handle arrays? - writeln!( - self.out, - "{}{}.{} = {}{};", - INDENT, RETURN_VARIABLE_NAME, field_name, ARGUMENT_VARIABLE_NAME, 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 + )?; + } + } + _ => { + //TODO: handle arrays? + writeln!( + self.out, + "{}{}.{} = {}{};", + INDENT, RETURN_VARIABLE_NAME, field_name, ARGUMENT_VARIABLE_NAME, i, + )?; + } + } } // Write return value @@ -413,6 +445,335 @@ impl<'a, W: Write> super::Writer<'a, W> { Ok(()) } + pub(super) fn write_wrapped_struct_matrix_get_function_name( + &mut self, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + let name = &self.names[&NameKey::Type(access.ty)]; + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + write!(self.out, "GetMat{}On{}", field_name, name)?; + Ok(()) + } + + /// Writes a function used to get a matCx2 from within a structure. + pub(super) fn write_wrapped_struct_matrix_get_function( + &mut self, + module: &crate::Module, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + use crate::back::INDENT; + + const STRUCT_ARGUMENT_VARIABLE_NAME: &str = "obj"; + + // Write function return type and name + let member = match module.types[access.ty].inner { + crate::TypeInner::Struct { ref members, .. } => &members[access.index as usize], + _ => unreachable!(), + }; + let ret_ty = &module.types[member.ty].inner; + self.write_value_type(module, ret_ty)?; + write!(self.out, " ")?; + self.write_wrapped_struct_matrix_get_function_name(access)?; + + // Write function parameters + write!(self.out, "(")?; + let struct_name = &self.names[&NameKey::Type(access.ty)]; + write!( + self.out, + "{} {}", + struct_name, STRUCT_ARGUMENT_VARIABLE_NAME + )?; + + // Write function body + writeln!(self.out, ") {{")?; + + // Write return value + write!(self.out, "{}return ", INDENT)?; + self.write_value_type(module, ret_ty)?; + write!(self.out, "(")?; + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + match module.types[member.ty].inner { + crate::TypeInner::Matrix { columns, .. } => { + for i in 0..columns as u8 { + if i != 0 { + write!(self.out, ", ")?; + } + write!( + self.out, + "{}.{}_{}", + STRUCT_ARGUMENT_VARIABLE_NAME, field_name, i + )?; + } + } + _ => unreachable!(), + } + writeln!(self.out, ");")?; + + // End of function body + writeln!(self.out, "}}")?; + // Write extra new line + writeln!(self.out)?; + + Ok(()) + } + + pub(super) fn write_wrapped_struct_matrix_set_function_name( + &mut self, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + let name = &self.names[&NameKey::Type(access.ty)]; + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + write!(self.out, "SetMat{}On{}", field_name, name)?; + Ok(()) + } + + /// Writes a function used to set a matCx2 from within a structure. + pub(super) fn write_wrapped_struct_matrix_set_function( + &mut self, + module: &crate::Module, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + use crate::back::INDENT; + + const STRUCT_ARGUMENT_VARIABLE_NAME: &str = "obj"; + const MATRIX_ARGUMENT_VARIABLE_NAME: &str = "mat"; + + // Write function return type and name + write!(self.out, "void ")?; + self.write_wrapped_struct_matrix_set_function_name(access)?; + + // Write function parameters + write!(self.out, "(")?; + let struct_name = &self.names[&NameKey::Type(access.ty)]; + write!( + self.out, + "{} {}, ", + struct_name, STRUCT_ARGUMENT_VARIABLE_NAME + )?; + let member = match module.types[access.ty].inner { + crate::TypeInner::Struct { ref members, .. } => &members[access.index as usize], + _ => unreachable!(), + }; + self.write_type(module, member.ty)?; + write!(self.out, " {}", MATRIX_ARGUMENT_VARIABLE_NAME)?; + // Write function body + writeln!(self.out, ") {{")?; + + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + + match module.types[member.ty].inner { + crate::TypeInner::Matrix { columns, .. } => { + for i in 0..columns as u8 { + writeln!( + self.out, + "{}{}.{}_{} = {}[{}];", + INDENT, + STRUCT_ARGUMENT_VARIABLE_NAME, + field_name, + i, + MATRIX_ARGUMENT_VARIABLE_NAME, + i + )?; + } + } + _ => unreachable!(), + } + + // End of function body + writeln!(self.out, "}}")?; + // Write extra new line + writeln!(self.out)?; + + Ok(()) + } + + pub(super) fn write_wrapped_struct_matrix_set_vec_function_name( + &mut self, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + let name = &self.names[&NameKey::Type(access.ty)]; + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + write!(self.out, "SetMatVec{}On{}", field_name, name)?; + Ok(()) + } + + /// Writes a function used to set a vec2 on a matCx2 from within a structure. + pub(super) fn write_wrapped_struct_matrix_set_vec_function( + &mut self, + module: &crate::Module, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + use crate::back::INDENT; + + const STRUCT_ARGUMENT_VARIABLE_NAME: &str = "obj"; + const VECTOR_ARGUMENT_VARIABLE_NAME: &str = "vec"; + const MATRIX_INDEX_ARGUMENT_VARIABLE_NAME: &str = "mat_idx"; + + // Write function return type and name + write!(self.out, "void ")?; + self.write_wrapped_struct_matrix_set_vec_function_name(access)?; + + // Write function parameters + write!(self.out, "(")?; + let struct_name = &self.names[&NameKey::Type(access.ty)]; + write!( + self.out, + "{} {}, ", + struct_name, STRUCT_ARGUMENT_VARIABLE_NAME + )?; + let member = match module.types[access.ty].inner { + crate::TypeInner::Struct { ref members, .. } => &members[access.index as usize], + _ => unreachable!(), + }; + let vec_ty = match module.types[member.ty].inner { + crate::TypeInner::Matrix { rows, width, .. } => crate::TypeInner::Vector { + size: rows, + kind: crate::ScalarKind::Float, + width, + }, + _ => unreachable!(), + }; + self.write_value_type(module, &vec_ty)?; + write!( + self.out, + " {}, uint {}", + VECTOR_ARGUMENT_VARIABLE_NAME, MATRIX_INDEX_ARGUMENT_VARIABLE_NAME + )?; + + // Write function body + writeln!(self.out, ") {{")?; + + writeln!( + self.out, + "{}switch({}) {{", + INDENT, MATRIX_INDEX_ARGUMENT_VARIABLE_NAME + )?; + + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + + match module.types[member.ty].inner { + crate::TypeInner::Matrix { columns, .. } => { + for i in 0..columns as u8 { + writeln!( + self.out, + "{}case {}: {}.{}_{} = {};", + INDENT, + i, + STRUCT_ARGUMENT_VARIABLE_NAME, + field_name, + i, + VECTOR_ARGUMENT_VARIABLE_NAME + )?; + } + } + _ => unreachable!(), + } + + writeln!(self.out, "{}}}", INDENT)?; + + // End of function body + writeln!(self.out, "}}")?; + // Write extra new line + writeln!(self.out)?; + + Ok(()) + } + + pub(super) fn write_wrapped_struct_matrix_set_scalar_function_name( + &mut self, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + let name = &self.names[&NameKey::Type(access.ty)]; + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + write!(self.out, "SetMatScalar{}On{}", field_name, name)?; + Ok(()) + } + + /// Writes a function used to set a float on a matCx2 from within a structure. + pub(super) fn write_wrapped_struct_matrix_set_scalar_function( + &mut self, + module: &crate::Module, + access: WrappedStructMatrixAccess, + ) -> BackendResult { + use crate::back::INDENT; + + const STRUCT_ARGUMENT_VARIABLE_NAME: &str = "obj"; + const SCALAR_ARGUMENT_VARIABLE_NAME: &str = "scalar"; + const MATRIX_INDEX_ARGUMENT_VARIABLE_NAME: &str = "mat_idx"; + const VECTOR_INDEX_ARGUMENT_VARIABLE_NAME: &str = "vec_idx"; + + // Write function return type and name + write!(self.out, "void ")?; + self.write_wrapped_struct_matrix_set_scalar_function_name(access)?; + + // Write function parameters + write!(self.out, "(")?; + let struct_name = &self.names[&NameKey::Type(access.ty)]; + write!( + self.out, + "{} {}, ", + struct_name, STRUCT_ARGUMENT_VARIABLE_NAME + )?; + let member = match module.types[access.ty].inner { + crate::TypeInner::Struct { ref members, .. } => &members[access.index as usize], + _ => unreachable!(), + }; + let scalar_ty = match module.types[member.ty].inner { + crate::TypeInner::Matrix { width, .. } => crate::TypeInner::Scalar { + kind: crate::ScalarKind::Float, + width, + }, + _ => unreachable!(), + }; + self.write_value_type(module, &scalar_ty)?; + write!( + self.out, + " {}, uint {}, uint {}", + SCALAR_ARGUMENT_VARIABLE_NAME, + MATRIX_INDEX_ARGUMENT_VARIABLE_NAME, + VECTOR_INDEX_ARGUMENT_VARIABLE_NAME + )?; + + // Write function body + writeln!(self.out, ") {{")?; + + writeln!( + self.out, + "{}switch({}) {{", + INDENT, MATRIX_INDEX_ARGUMENT_VARIABLE_NAME + )?; + + let field_name = &self.names[&NameKey::StructMember(access.ty, access.index)]; + + match module.types[member.ty].inner { + crate::TypeInner::Matrix { columns, .. } => { + for i in 0..columns as u8 { + writeln!( + self.out, + "{}case {}: {}.{}_{}[{}] = {};", + INDENT, + i, + STRUCT_ARGUMENT_VARIABLE_NAME, + field_name, + i, + VECTOR_INDEX_ARGUMENT_VARIABLE_NAME, + SCALAR_ARGUMENT_VARIABLE_NAME + )?; + } + } + _ => unreachable!(), + } + + writeln!(self.out, "{}}}", INDENT)?; + + // End of function body + writeln!(self.out, "}}")?; + // Write extra new line + writeln!(self.out)?; + + Ok(()) + } + /// Helper function that write wrapped function for `Expression::ImageQuery` and `Expression::ArrayLength` /// /// @@ -478,6 +839,47 @@ impl<'a, W: Write> super::Writer<'a, W> { self.wrapped.constructors.insert(constructor); } } + // We treat matrices of the form `matCx2` as a sequence of C `vec2`s + // (see top level module docs for details). + // + // The functions injected here are required to get the matrix accesses working. + crate::Expression::AccessIndex { base, index } => { + let base_ty_res = &func_ctx.info[base].ty; + let mut resolved = base_ty_res.inner_with(&module.types); + let base_ty_handle = match *resolved { + crate::TypeInner::Pointer { base, .. } => { + resolved = &module.types[base].inner; + Some(base) + } + _ => base_ty_res.handle(), + }; + if let crate::TypeInner::Struct { ref members, .. } = *resolved { + let member = &members[index as usize]; + + match module.types[member.ty].inner { + crate::TypeInner::Matrix { + rows: crate::VectorSize::Bi, + .. + } if member.binding.is_none() => { + let ty = base_ty_handle.unwrap(); + let access = WrappedStructMatrixAccess { ty, index }; + + if !self.wrapped.struct_matrix_access.contains(&access) { + self.write_wrapped_struct_matrix_get_function(module, access)?; + self.write_wrapped_struct_matrix_set_function(module, access)?; + self.write_wrapped_struct_matrix_set_vec_function( + module, access, + )?; + self.write_wrapped_struct_matrix_set_scalar_function( + module, access, + )?; + self.wrapped.struct_matrix_access.insert(access); + } + } + _ => {} + } + } + } _ => {} }; } diff --git a/src/back/hlsl/mod.rs b/src/back/hlsl/mod.rs index 8f3134bb1b..3ab88409f8 100644 --- a/src/back/hlsl/mod.rs +++ b/src/back/hlsl/mod.rs @@ -6,6 +6,8 @@ Backend for [HLSL][hlsl] (High-Level Shading Language). - 5.1 - 6.0 +# General Matrix Note + All matrix construction/deconstruction is row based in HLSL. This means that when we construct a matrix from column vectors, our matrix will be implicitly transposed. The inverse transposition happens when we call `[0]` to get the zeroth column vector. @@ -19,6 +21,13 @@ To deal with this, we add `row_major` to all declarations of matrices in Uniform Finally because all of our matrices are transposed, if you use `mat3x4`, it'll become `float3x4` in HLSL (HLSL has inverted col/row notation). +# Matrix struct member of the form `matCx2` Note + +Struct member matrices of the form `matCx2` are translated to a sequence of C `vec2`s due to +differences in alignment between WGSL and HLSL for uniform buffers. + +Accesses to these matrices are handled by injected functions. + [hlsl]: https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl */ @@ -169,6 +178,7 @@ struct Wrapped { array_lengths: crate::FastHashSet, image_queries: crate::FastHashSet, constructors: crate::FastHashSet, + struct_matrix_access: crate::FastHashSet, } impl Wrapped { @@ -176,6 +186,7 @@ impl Wrapped { self.array_lengths.clear(); self.image_queries.clear(); self.constructors.clear(); + self.struct_matrix_access.clear(); } } diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 5516f94a25..b514d11557 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1,5 +1,5 @@ use super::{ - help::{WrappedArrayLength, WrappedConstructor, WrappedImageQuery}, + help::{WrappedArrayLength, WrappedConstructor, WrappedImageQuery, WrappedStructMatrixAccess}, storage::StoreValue, BackendResult, Error, Options, }; @@ -784,6 +784,28 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Write [size] self.write_array_size(module, size)?; } + // We treat matrices of the form `matCx2` as a sequence of C `vec2`s + // (see top level module docs for details). + TypeInner::Matrix { + rows, + columns, + width, + } if member.binding.is_none() && rows == crate::VectorSize::Bi => { + let vec_ty = crate::TypeInner::Vector { + size: rows, + kind: crate::ScalarKind::Float, + width, + }; + let field_name_key = NameKey::StructMember(handle, index as u32); + + for i in 0..columns as u8 { + if i != 0 { + write!(self.out, "; ")?; + } + self.write_value_type(module, &vec_ty)?; + write!(self.out, " {}_{}", &self.names[&field_name_key], i)?; + } + } _ => { // Write modifier before type if let Some(ref binding) = member.binding { @@ -1253,11 +1275,175 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, "[_i] = _result[_i];")?; writeln!(self.out, "{}}}", level)?; } else { + // We treat matrices of the form `matCx2` as a sequence of C `vec2`s + // (see top level module docs for details). + // + // We handle matrix Stores here directly (including sub accesses for Vectors and Scalars). + // Loads are handled by `Expression::AccessIndex` (since sub accesses work fine for Loads). + struct MatrixAccess { + base: Handle, + index: u32, + } + enum Index { + Expression(Handle), + Static(u32), + } + + let get_members = |expr: Handle| { + let base_ty_res = &func_ctx.info[expr].ty; + let resolved = base_ty_res.inner_with(&module.types); + match *resolved { + TypeInner::Pointer { base, .. } => match module.types[base].inner { + TypeInner::Struct { ref members, .. } => Some(members), + _ => None, + }, + _ => None, + } + }; + + let mut matrix = None; + let mut vector = None; + let mut scalar = None; + + let mut current_expr = pointer; + for _ in 0..3 { + let resolved = func_ctx.info[current_expr].ty.inner_with(&module.types); + + match (resolved, &func_ctx.expressions[current_expr]) { + ( + &TypeInner::Pointer { base: ty, .. }, + &crate::Expression::AccessIndex { base, index }, + ) if matches!( + module.types[ty].inner, + TypeInner::Matrix { + rows: crate::VectorSize::Bi, + .. + } + ) && get_members(base) + .map(|members| members[index as usize].binding.is_none()) + == Some(true) => + { + matrix = Some(MatrixAccess { base, index }); + break; + } + ( + &TypeInner::ValuePointer { + size: Some(crate::VectorSize::Bi), + .. + }, + &crate::Expression::Access { base, index }, + ) => { + vector = Some(Index::Expression(index)); + current_expr = base; + } + ( + &TypeInner::ValuePointer { + size: Some(crate::VectorSize::Bi), + .. + }, + &crate::Expression::AccessIndex { base, index }, + ) => { + vector = Some(Index::Static(index)); + current_expr = base; + } + ( + &TypeInner::ValuePointer { size: None, .. }, + &crate::Expression::Access { base, index }, + ) => { + scalar = Some(Index::Expression(index)); + current_expr = base; + } + ( + &TypeInner::ValuePointer { size: None, .. }, + &crate::Expression::AccessIndex { base, index }, + ) => { + scalar = Some(Index::Static(index)); + current_expr = base; + } + _ => break, + } + } + write!(self.out, "{}", level)?; - self.write_expr(module, pointer, func_ctx)?; - write!(self.out, " = ")?; - self.write_expr(module, value, func_ctx)?; - writeln!(self.out, ";")? + + if let Some(MatrixAccess { index, base }) = matrix { + let base_ty_res = &func_ctx.info[base].ty; + let resolved = base_ty_res.inner_with(&module.types); + let ty = match *resolved { + TypeInner::Pointer { base, .. } => base, + _ => base_ty_res.handle().unwrap(), + }; + + if let Some(Index::Static(vec_index)) = vector { + self.write_expr(module, base, func_ctx)?; + write!( + self.out, + ".{}_{}", + &self.names[&NameKey::StructMember(ty, index)], + vec_index + )?; + + if let Some(scalar_index) = scalar { + write!(self.out, "[")?; + match scalar_index { + Index::Static(index) => { + write!(self.out, "{}", index)?; + } + Index::Expression(index) => { + self.write_expr(module, index, func_ctx)?; + } + } + write!(self.out, "]")?; + } + + write!(self.out, " = ")?; + self.write_expr(module, value, func_ctx)?; + writeln!(self.out, ";")?; + } else { + let access = WrappedStructMatrixAccess { ty, index }; + match (&vector, &scalar) { + (&Some(_), &Some(_)) => { + self.write_wrapped_struct_matrix_set_scalar_function_name( + access, + )?; + } + (&Some(_), &None) => { + self.write_wrapped_struct_matrix_set_vec_function_name(access)?; + } + (&None, _) => { + self.write_wrapped_struct_matrix_set_function_name(access)?; + } + } + + write!(self.out, "(")?; + self.write_expr(module, base, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, value, func_ctx)?; + + if let Some(Index::Expression(vec_index)) = vector { + write!(self.out, ", ")?; + self.write_expr(module, vec_index, func_ctx)?; + + if let Some(scalar_index) = scalar { + write!(self.out, ", ")?; + match scalar_index { + Index::Static(index) => { + write!(self.out, "{}", index)?; + } + Index::Expression(index) => { + self.write_expr(module, index, func_ctx)?; + } + } + } + } + writeln!(self.out, ");")?; + } + } else { + self.write_expr(module, pointer, func_ctx)?; + write!(self.out, " = ")?; + self.write_expr(module, value, func_ctx)?; + writeln!(self.out, ";")? + } } } Statement::Loop { @@ -1592,8 +1778,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { { // do nothing, the chain is written on `Load`/`Store` } else { - self.write_expr(module, base, func_ctx)?; - let base_ty_res = &func_ctx.info[base].ty; let mut resolved = base_ty_res.inner_with(&module.types); let base_ty_handle = match *resolved { @@ -1604,6 +1788,34 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { _ => base_ty_res.handle(), }; + // We treat matrices of the form `matCx2` as a sequence of C `vec2`s + // (see top level module docs for details). + // + // We handle matrix reconstruction here for Loads. + // Stores are handled directly by `Statement::Store`. + if let TypeInner::Struct { ref members, .. } = *resolved { + let member = &members[index as usize]; + + match module.types[member.ty].inner { + TypeInner::Matrix { + rows: crate::VectorSize::Bi, + .. + } if member.binding.is_none() => { + let ty = base_ty_handle.unwrap(); + self.write_wrapped_struct_matrix_get_function_name( + WrappedStructMatrixAccess { ty, index }, + )?; + write!(self.out, "(")?; + self.write_expr(module, base, func_ctx)?; + write!(self.out, ")")?; + return Ok(()); + } + _ => {} + } + }; + + self.write_expr(module, base, func_ctx)?; + match *resolved { TypeInner::Vector { .. } => { // Write vector access as a swizzle diff --git a/tests/in/access.param.ron b/tests/in/access.param.ron index 84595754e4..57cb7a632d 100644 --- a/tests/in/access.param.ron +++ b/tests/in/access.param.ron @@ -10,6 +10,7 @@ vs: ( resources: { (group: 0, binding: 0): (buffer: Some(0), mutable: false), + (group: 0, binding: 1): (buffer: Some(1), mutable: false), }, sizes_buffer: Some(24), ), diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index 5ba29d90d5..a034fdda7d 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -15,6 +15,41 @@ struct Bar { @group(0) @binding(0) var bar: Bar; +struct Baz { + m: mat3x2, +} + +@group(0) @binding(1) +var baz: Baz; + +fn test_matrix_within_struct_accesses() { + var idx = 9; + + idx--; + + // loads + var _ = baz.m; + var _ = baz.m[0]; + var _ = baz.m[idx]; + var _ = baz.m[0][1]; + var _ = baz.m[0][idx]; + var _ = baz.m[idx][1]; + var _ = baz.m[idx][idx]; + + var t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0))); + + idx++; + + // stores + t.m = mat3x2(vec2(6.0), vec2(5.0), vec2(4.0)); + t.m[0] = vec2(9.0); + t.m[idx] = vec2(90.0); + t.m[0][1] = 10.0; + t.m[0][idx] = 20.0; + t.m[idx][1] = 30.0; + t.m[idx][idx] = 40.0; +} + fn read_from_private(foo: ptr) -> f32 { return *foo; } @@ -26,6 +61,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let baz: f32 = foo; foo = 1.0; + test_matrix_within_struct_accesses(); + // test storage loads let matrix = bar.matrix; let arr = bar.arr; diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index 5edd19fea5..ff67b1b988 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -8,6 +8,9 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; struct AlignedWrapper { int value; }; +struct Baz { + mat3x2 m; +}; layout(std430) buffer Bar_block_0Compute { mat4x3 matrix; mat2x2 matrix_array[2]; @@ -18,29 +21,29 @@ layout(std430) buffer Bar_block_0Compute { float read_from_private(inout float foo_1) { - float _e2 = foo_1; - return _e2; + float _e3 = foo_1; + return _e3; } void main() { int tmp = 0; int value = _group_0_binding_0_cs.atom; - int _e6 = atomicAdd(_group_0_binding_0_cs.atom, 5); - tmp = _e6; - int _e9 = atomicAdd(_group_0_binding_0_cs.atom, -5); - tmp = _e9; - int _e12 = atomicAnd(_group_0_binding_0_cs.atom, 5); - tmp = _e12; - int _e15 = atomicOr(_group_0_binding_0_cs.atom, 5); - tmp = _e15; - int _e18 = atomicXor(_group_0_binding_0_cs.atom, 5); - tmp = _e18; - int _e21 = atomicMin(_group_0_binding_0_cs.atom, 5); - tmp = _e21; - int _e24 = atomicMax(_group_0_binding_0_cs.atom, 5); - tmp = _e24; - int _e27 = atomicExchange(_group_0_binding_0_cs.atom, 5); - tmp = _e27; + int _e7 = atomicAdd(_group_0_binding_0_cs.atom, 5); + tmp = _e7; + int _e10 = atomicAdd(_group_0_binding_0_cs.atom, -5); + tmp = _e10; + int _e13 = atomicAnd(_group_0_binding_0_cs.atom, 5); + tmp = _e13; + int _e16 = atomicOr(_group_0_binding_0_cs.atom, 5); + tmp = _e16; + int _e19 = atomicXor(_group_0_binding_0_cs.atom, 5); + tmp = _e19; + int _e22 = atomicMin(_group_0_binding_0_cs.atom, 5); + tmp = _e22; + int _e25 = atomicMax(_group_0_binding_0_cs.atom, 5); + tmp = _e25; + int _e28 = atomicExchange(_group_0_binding_0_cs.atom, 5); + tmp = _e28; _group_0_binding_0_cs.atom = value; return; } diff --git a/tests/out/glsl/access.foo_frag.Fragment.glsl b/tests/out/glsl/access.foo_frag.Fragment.glsl index 63b40b7ab1..b7c62d40e6 100644 --- a/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -6,6 +6,9 @@ precision highp int; struct AlignedWrapper { int value; }; +struct Baz { + mat3x2 m; +}; layout(std430) buffer Bar_block_0Fragment { mat4x3 matrix; mat2x2 matrix_array[2]; @@ -17,8 +20,8 @@ layout(std430) buffer Bar_block_0Fragment { layout(location = 0) out vec4 _fs2p_location0; float read_from_private(inout float foo_1) { - float _e2 = foo_1; - return _e2; + float _e3 = foo_1; + return _e3; } void main() { diff --git a/tests/out/glsl/access.foo_vert.Vertex.glsl b/tests/out/glsl/access.foo_vert.Vertex.glsl index f1fb3e09a1..b01479c243 100644 --- a/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -6,6 +6,9 @@ precision highp int; struct AlignedWrapper { int value; }; +struct Baz { + mat3x2 m; +}; layout(std430) buffer Bar_block_0Vertex { mat4x3 matrix; mat2x2 matrix_array[2]; @@ -14,23 +17,75 @@ layout(std430) buffer Bar_block_0Vertex { AlignedWrapper data[]; } _group_0_binding_0_vs; +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); + vec2 unnamed_2 = vec2(0.0); + float unnamed_3 = 0.0; + float unnamed_4 = 0.0; + float unnamed_5 = 0.0; + float unnamed_6 = 0.0; + Baz t = Baz(mat3x2(0.0)); + int _e4 = idx; + idx = (_e4 - 1); + mat3x2 _e8 = _group_0_binding_1_vs.m; + unnamed = _e8; + vec2 _e13 = _group_0_binding_1_vs.m[0]; + unnamed_1 = _e13; + int _e16 = idx; + vec2 _e18 = _group_0_binding_1_vs.m[_e16]; + unnamed_2 = _e18; + float _e25 = _group_0_binding_1_vs.m[0][1]; + unnamed_3 = _e25; + int _e30 = idx; + float _e32 = _group_0_binding_1_vs.m[0][_e30]; + unnamed_4 = _e32; + int _e35 = idx; + float _e39 = _group_0_binding_1_vs.m[_e35][1]; + unnamed_5 = _e39; + int _e42 = idx; + int _e44 = idx; + float _e46 = _group_0_binding_1_vs.m[_e42][_e44]; + unnamed_6 = _e46; + t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0))); + int _e57 = idx; + idx = (_e57 + 1); + t.m = mat3x2(vec2(6.0), vec2(5.0), vec2(4.0)); + t.m[0] = vec2(9.0); + int _e74 = idx; + t.m[_e74] = vec2(90.0); + t.m[0][1] = 10.0; + int _e87 = idx; + t.m[0][_e87] = 20.0; + int _e91 = idx; + t.m[_e91][1] = 30.0; + int _e97 = idx; + int _e99 = idx; + t.m[_e97][_e99] = 40.0; + return; +} float read_from_private(inout float foo_1) { - float _e2 = foo_1; - return _e2; + float _e3 = foo_1; + return _e3; } void main() { uint vi = uint(gl_VertexID); float foo = 0.0; int c[5] = int[5](0, 0, 0, 0, 0); - float baz = foo; + 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; - float _e27 = read_from_private(foo); + float _e28 = read_from_private(foo); c = int[5](a, int(b), 3, 4, 5); c[(vi + 1u)] = 42; int value = c[vi]; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 8af92afdd8..7f63c78d87 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -3,12 +3,102 @@ struct AlignedWrapper { int value; }; +struct Baz { + float2 m_0; float2 m_1; float2 m_2; +}; + RWByteAddressBuffer bar : register(u0); +cbuffer baz : register(b1) { Baz baz; } + +float3x2 GetMatmOnBaz(Baz obj) { + return float3x2(obj.m_0, obj.m_1, obj.m_2); +} + +void SetMatmOnBaz(Baz obj, float3x2 mat) { + obj.m_0 = mat[0]; + obj.m_1 = mat[1]; + obj.m_2 = mat[2]; +} + +void SetMatVecmOnBaz(Baz obj, float2 vec, uint mat_idx) { + switch(mat_idx) { + case 0: obj.m_0 = vec; + case 1: obj.m_1 = vec; + case 2: obj.m_2 = vec; + } +} + +void SetMatScalarmOnBaz(Baz obj, float scalar, uint mat_idx, uint vec_idx) { + switch(mat_idx) { + case 0: obj.m_0[vec_idx] = scalar; + case 1: obj.m_1[vec_idx] = scalar; + case 2: obj.m_2[vec_idx] = scalar; + } +} + +Baz ConstructBaz(float3x2 arg0) { + Baz ret; + ret.m_0 = arg0[0]; + ret.m_1 = arg0[1]; + ret.m_2 = arg0[2]; + return ret; +} + +void test_matrix_within_struct_accesses() +{ + int idx = 9; + float3x2 unnamed = (float3x2)0; + float2 unnamed_1 = (float2)0; + float2 unnamed_2 = (float2)0; + float unnamed_3 = (float)0; + float unnamed_4 = (float)0; + float unnamed_5 = (float)0; + float unnamed_6 = (float)0; + Baz t = (Baz)0; + + int _expr4 = idx; + idx = (_expr4 - 1); + float3x2 _expr8 = GetMatmOnBaz(baz); + unnamed = _expr8; + float2 _expr13 = GetMatmOnBaz(baz)[0]; + unnamed_1 = _expr13; + int _expr16 = idx; + float2 _expr18 = GetMatmOnBaz(baz)[_expr16]; + unnamed_2 = _expr18; + float _expr25 = GetMatmOnBaz(baz)[0][1]; + unnamed_3 = _expr25; + int _expr30 = idx; + float _expr32 = GetMatmOnBaz(baz)[0][_expr30]; + unnamed_4 = _expr32; + int _expr35 = idx; + float _expr39 = GetMatmOnBaz(baz)[_expr35][1]; + unnamed_5 = _expr39; + int _expr42 = idx; + int _expr44 = idx; + float _expr46 = GetMatmOnBaz(baz)[_expr42][_expr44]; + unnamed_6 = _expr46; + t = ConstructBaz(float3x2(float2(1.0.xx), float2(2.0.xx), float2(3.0.xx))); + int _expr57 = idx; + idx = (_expr57 + 1); + SetMatmOnBaz(t, float3x2(float2(6.0.xx), float2(5.0.xx), float2(4.0.xx))); + t.m_0 = float2(9.0.xx); + int _expr74 = idx; + SetMatVecmOnBaz(t, float2(90.0.xx), _expr74); + t.m_0[1] = 10.0; + int _expr87 = idx; + t.m_0[_expr87] = 20.0; + int _expr91 = idx; + SetMatScalarmOnBaz(t, 30.0, _expr91, 1); + int _expr97 = idx; + int _expr99 = idx; + SetMatScalarmOnBaz(t, 40.0, _expr97, _expr99); + return; +} float read_from_private(inout float foo_1) { - float _expr2 = foo_1; - return _expr2; + float _expr3 = foo_1; + return _expr3; } uint NagaBufferLengthRW(RWByteAddressBuffer buffer) @@ -23,13 +113,14 @@ 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 baz = foo; + float baz_1 = foo; foo = 1.0; + test_matrix_within_struct_accesses(); 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)); - const float _e27 = read_from_private(foo); + const float _e28 = read_from_private(foo); { int _result[5]={ a, int(b), 3, 4, 5 }; for(int _i=0; _i<5; ++_i) c[_i] = _result[_i]; @@ -64,22 +155,22 @@ void atomics() int tmp = (int)0; int value_1 = asint(bar.Load(96)); - int _e6; bar.InterlockedAdd(96, 5, _e6); - tmp = _e6; - int _e9; bar.InterlockedAdd(96, -5, _e9); - tmp = _e9; - int _e12; bar.InterlockedAnd(96, 5, _e12); - tmp = _e12; - int _e15; bar.InterlockedOr(96, 5, _e15); - tmp = _e15; - int _e18; bar.InterlockedXor(96, 5, _e18); - tmp = _e18; - int _e21; bar.InterlockedMin(96, 5, _e21); - tmp = _e21; - int _e24; bar.InterlockedMax(96, 5, _e24); - tmp = _e24; - int _e27; bar.InterlockedExchange(96, 5, _e27); - tmp = _e27; + int _e7; bar.InterlockedAdd(96, 5, _e7); + tmp = _e7; + int _e10; bar.InterlockedAdd(96, -5, _e10); + tmp = _e10; + int _e13; bar.InterlockedAnd(96, 5, _e13); + tmp = _e13; + int _e16; bar.InterlockedOr(96, 5, _e16); + tmp = _e16; + int _e19; bar.InterlockedXor(96, 5, _e19); + tmp = _e19; + int _e22; bar.InterlockedMin(96, 5, _e22); + tmp = _e22; + int _e25; bar.InterlockedMax(96, 5, _e25); + tmp = _e25; + int _e28; bar.InterlockedExchange(96, 5, _e28); + tmp = _e28; bar.Store(96, asuint(value_1)); return; } diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index dbdad72e18..0f3b5cb122 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -26,15 +26,69 @@ struct Bar { type_6 arr; type_7 data; }; -struct type_13 { +struct Baz { + metal::float3x2 m; +}; +struct type_15 { int inner[5]; }; +void test_matrix_within_struct_accesses( + constant Baz& baz +) { + int idx = 9; + metal::float3x2 unnamed; + metal::float2 unnamed_1; + metal::float2 unnamed_2; + float unnamed_3; + float unnamed_4; + float unnamed_5; + float unnamed_6; + Baz t; + int _e4 = idx; + idx = _e4 - 1; + metal::float3x2 _e8 = baz.m; + unnamed = _e8; + metal::float2 _e13 = baz.m[0]; + unnamed_1 = _e13; + int _e16 = idx; + metal::float2 _e18 = baz.m[_e16]; + unnamed_2 = _e18; + float _e25 = baz.m[0].y; + unnamed_3 = _e25; + int _e30 = idx; + float _e32 = baz.m[0][_e30]; + unnamed_4 = _e32; + int _e35 = idx; + float _e39 = baz.m[_e35].y; + unnamed_5 = _e39; + int _e42 = idx; + int _e44 = idx; + float _e46 = baz.m[_e42][_e44]; + unnamed_6 = _e46; + t = Baz {metal::float3x2(metal::float2(1.0), metal::float2(2.0), metal::float2(3.0))}; + int _e57 = idx; + idx = _e57 + 1; + t.m = metal::float3x2(metal::float2(6.0), metal::float2(5.0), metal::float2(4.0)); + t.m[0] = metal::float2(9.0); + int _e74 = idx; + t.m[_e74] = metal::float2(90.0); + t.m[0].y = 10.0; + int _e87 = idx; + t.m[0][_e87] = 20.0; + int _e91 = idx; + t.m[_e91].y = 30.0; + int _e97 = idx; + int _e99 = idx; + t.m[_e97][_e99] = 40.0; + return; +} + float read_from_private( thread float& foo_1 ) { - float _e2 = foo_1; - return _e2; + float _e3 = foo_1; + return _e3; } struct foo_vertInput { @@ -45,18 +99,20 @@ struct foo_vertOutput { vertex foo_vertOutput foo_vert( uint vi [[vertex_id]] , device Bar const& bar [[buffer(0)]] +, constant Baz& baz [[buffer(1)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { float foo = 0.0; - type_13 c; - float baz = foo; + type_15 c; + 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; - float _e27 = read_from_private(foo); - for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast(b), 3, 4, 5}.inner[_i]; + 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]; c.inner[vi + 1u] = 42; int value = c.inner[vi]; return foo_vertOutput { metal::float4(matrix * static_cast(metal::int4(value)), 2.0) }; @@ -84,22 +140,22 @@ kernel void atomics( ) { int tmp; int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed); - int _e6 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e6; - int _e9 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e9; - int _e12 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e12; - int _e15 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e15; - int _e18 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e18; - int _e21 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e21; - int _e24 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e24; - int _e27 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); - tmp = _e27; + int _e7 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e7; + int _e10 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e10; + int _e13 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e13; + int _e16 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e16; + int _e19 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e19; + int _e22 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e22; + int _e25 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e25; + int _e28 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e28; metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed); return; } diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index a78c502374..0f89593901 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,208 +1,333 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 128 +; Bound: 213 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %52 "foo_vert" %47 %50 -OpEntryPoint Fragment %88 "foo_frag" %87 -OpEntryPoint GLCompute %105 "atomics" -OpExecutionMode %88 OriginUpperLeft -OpExecutionMode %105 LocalSize 1 1 1 +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 OpSource GLSL 450 -OpMemberName %21 0 "value" -OpName %21 "AlignedWrapper" -OpMemberName %30 0 "matrix" -OpMemberName %30 1 "matrix_array" -OpMemberName %30 2 "atom" -OpMemberName %30 3 "arr" -OpMemberName %30 4 "data" -OpName %30 "Bar" -OpName %35 "bar" -OpName %38 "foo" -OpName %39 "read_from_private" -OpName %43 "foo" -OpName %44 "c" -OpName %47 "vi" -OpName %52 "foo_vert" -OpName %88 "foo_frag" -OpName %103 "tmp" -OpName %105 "atomics" -OpMemberDecorate %21 0 Offset 0 -OpDecorate %26 ArrayStride 16 -OpDecorate %28 ArrayStride 8 -OpDecorate %29 ArrayStride 8 -OpMemberDecorate %30 0 Offset 0 -OpMemberDecorate %30 0 ColMajor -OpMemberDecorate %30 0 MatrixStride 16 -OpMemberDecorate %30 1 Offset 64 -OpMemberDecorate %30 1 ColMajor -OpMemberDecorate %30 1 MatrixStride 8 -OpMemberDecorate %30 2 Offset 96 -OpMemberDecorate %30 3 Offset 104 -OpMemberDecorate %30 4 Offset 120 -OpDecorate %34 ArrayStride 4 -OpDecorate %35 DescriptorSet 0 -OpDecorate %35 Binding 0 -OpDecorate %30 Block -OpDecorate %47 BuiltIn VertexIndex -OpDecorate %50 BuiltIn Position -OpDecorate %87 Location 0 +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 "_" +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 %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 +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 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 -%6 = OpTypeFloat 32 -%5 = OpConstant %6 0.0 -%7 = OpConstant %6 1.0 -%9 = OpTypeInt 32 0 -%8 = OpConstant %9 3 -%10 = OpConstant %9 2 -%11 = OpConstant %4 0 -%12 = OpConstant %4 5 -%13 = OpConstant %4 3 -%14 = OpConstant %4 4 -%15 = OpConstant %9 1 -%16 = OpConstant %4 42 -%17 = OpConstant %6 2.0 -%18 = OpConstant %4 1 -%19 = OpConstant %6 3.0 -%20 = OpConstant %9 0 -%21 = OpTypeStruct %4 -%23 = OpTypeVector %6 3 -%22 = OpTypeMatrix %23 4 -%25 = OpTypeVector %6 2 -%24 = OpTypeMatrix %25 2 -%26 = OpTypeArray %24 %3 -%27 = OpTypeVector %9 2 -%28 = OpTypeArray %27 %3 -%29 = OpTypeRuntimeArray %21 -%30 = OpTypeStruct %22 %26 %4 %28 %29 -%31 = OpTypePointer Function %6 -%32 = OpTypeVector %6 4 -%33 = OpTypePointer StorageBuffer %4 -%34 = OpTypeArray %4 %12 -%36 = OpTypePointer StorageBuffer %30 -%35 = OpVariable %36 StorageBuffer -%40 = OpTypeFunction %6 %31 -%45 = OpTypePointer Function %34 -%48 = OpTypePointer Input %9 -%47 = OpVariable %48 Input -%51 = OpTypePointer Output %32 -%50 = OpVariable %51 Output -%53 = OpTypeFunction %2 -%56 = OpTypePointer StorageBuffer %22 -%59 = OpTypePointer StorageBuffer %28 -%62 = OpTypePointer StorageBuffer %23 -%63 = OpTypePointer StorageBuffer %6 -%66 = OpTypePointer StorageBuffer %29 -%69 = OpTypePointer StorageBuffer %21 -%70 = OpConstant %9 4 -%77 = OpTypePointer Function %4 -%81 = OpTypeVector %4 4 -%87 = OpVariable %51 Output -%107 = OpTypePointer StorageBuffer %4 -%110 = OpConstant %9 64 -%39 = OpFunction %6 None %40 -%38 = OpFunctionParameter %31 -%37 = OpLabel -OpBranch %41 -%41 = OpLabel -%42 = OpLoad %6 %38 -OpReturnValue %42 -OpFunctionEnd -%52 = OpFunction %2 None %53 -%46 = OpLabel -%43 = OpVariable %31 Function %5 -%44 = OpVariable %45 Function -%49 = OpLoad %9 %47 -OpBranch %54 -%54 = OpLabel -%55 = OpLoad %6 %43 -OpStore %43 %7 -%57 = OpAccessChain %56 %35 %20 -%58 = OpLoad %22 %57 -%60 = OpAccessChain %59 %35 %8 -%61 = OpLoad %28 %60 -%64 = OpAccessChain %63 %35 %20 %8 %20 -%65 = OpLoad %6 %64 -%67 = OpArrayLength %9 %35 4 -%68 = OpISub %9 %67 %10 -%71 = OpAccessChain %33 %35 %70 %68 %20 -%72 = OpLoad %4 %71 -%73 = OpFunctionCall %6 %39 %43 -%74 = OpConvertFToS %4 %65 -%75 = OpCompositeConstruct %34 %72 %74 %13 %14 %12 -OpStore %44 %75 -%76 = OpIAdd %9 %49 %15 -%78 = OpAccessChain %77 %44 %76 -OpStore %78 %16 -%79 = OpAccessChain %77 %44 %49 -%80 = OpLoad %4 %79 -%82 = OpCompositeConstruct %81 %80 %80 %80 %80 -%83 = OpConvertSToF %32 %82 -%84 = OpMatrixTimesVector %23 %58 %83 -%85 = OpCompositeConstruct %32 %84 %17 -OpStore %50 %85 +%5 = OpConstant %4 9 +%6 = OpConstant %4 1 +%7 = OpConstant %4 0 +%9 = OpTypeFloat 32 +%8 = OpConstant %9 1.0 +%10 = OpConstant %9 2.0 +%11 = OpConstant %9 3.0 +%12 = OpConstant %9 6.0 +%13 = OpConstant %9 5.0 +%14 = OpConstant %9 4.0 +%15 = OpConstant %9 9.0 +%16 = OpConstant %9 90.0 +%17 = OpConstant %9 10.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 +%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 OpReturn OpFunctionEnd -%88 = OpFunction %2 None %53 -%86 = OpLabel -OpBranch %89 -%89 = OpLabel -%90 = OpAccessChain %63 %35 %20 %15 %10 -OpStore %90 %7 -%91 = OpCompositeConstruct %23 %5 %5 %5 -%92 = OpCompositeConstruct %23 %7 %7 %7 -%93 = OpCompositeConstruct %23 %17 %17 %17 -%94 = OpCompositeConstruct %23 %19 %19 %19 -%95 = OpCompositeConstruct %22 %91 %92 %93 %94 -%96 = OpAccessChain %56 %35 %20 -OpStore %96 %95 -%97 = OpCompositeConstruct %27 %20 %20 -%98 = OpCompositeConstruct %27 %15 %15 -%99 = OpCompositeConstruct %28 %97 %98 -%100 = OpAccessChain %59 %35 %8 -OpStore %100 %99 -%101 = OpAccessChain %33 %35 %70 %15 %20 -OpStore %101 %18 -%102 = OpCompositeConstruct %32 %5 %5 %5 %5 -OpStore %87 %102 +%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 +%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 OpReturn OpFunctionEnd -%105 = OpFunction %2 None %53 -%104 = OpLabel -%103 = OpVariable %77 Function -OpBranch %106 -%106 = OpLabel -%108 = OpAccessChain %107 %35 %10 -%109 = OpAtomicLoad %4 %108 %18 %110 -%112 = OpAccessChain %107 %35 %10 -%111 = OpAtomicIAdd %4 %112 %18 %110 %12 -OpStore %103 %111 -%114 = OpAccessChain %107 %35 %10 -%113 = OpAtomicISub %4 %114 %18 %110 %12 -OpStore %103 %113 -%116 = OpAccessChain %107 %35 %10 -%115 = OpAtomicAnd %4 %116 %18 %110 %12 -OpStore %103 %115 -%118 = OpAccessChain %107 %35 %10 -%117 = OpAtomicOr %4 %118 %18 %110 %12 -OpStore %103 %117 -%120 = OpAccessChain %107 %35 %10 -%119 = OpAtomicXor %4 %120 %18 %110 %12 -OpStore %103 %119 -%122 = OpAccessChain %107 %35 %10 -%121 = OpAtomicSMin %4 %122 %18 %110 %12 -OpStore %103 %121 -%124 = OpAccessChain %107 %35 %10 -%123 = OpAtomicSMax %4 %124 %18 %110 %12 -OpStore %103 %123 -%126 = OpAccessChain %107 %35 %10 -%125 = OpAtomicExchange %4 %126 %18 %110 %12 -OpStore %103 %125 -%127 = OpAccessChain %107 %35 %10 -OpAtomicStore %127 %18 %110 %109 +%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 +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 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 38f0c45120..53cd3c4ac2 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -10,12 +10,68 @@ struct Bar { data: array, } +struct Baz { + m: mat3x2, +} + @group(0) @binding(0) var bar: Bar; +@group(0) @binding(1) +var baz: Baz; + +fn test_matrix_within_struct_accesses() { + var idx: i32 = 9; + var unnamed: mat3x2; + var unnamed_1: vec2; + var unnamed_2: vec2; + var unnamed_3: f32; + var unnamed_4: f32; + var unnamed_5: f32; + var unnamed_6: f32; + var t: Baz; + + let _e4 = idx; + idx = (_e4 - 1); + let _e8 = baz.m; + unnamed = _e8; + let _e13 = baz.m[0]; + unnamed_1 = _e13; + let _e16 = idx; + let _e18 = baz.m[_e16]; + unnamed_2 = _e18; + let _e25 = baz.m[0][1]; + unnamed_3 = _e25; + let _e30 = idx; + let _e32 = baz.m[0][_e30]; + unnamed_4 = _e32; + let _e35 = idx; + let _e39 = baz.m[_e35][1]; + unnamed_5 = _e39; + let _e42 = idx; + let _e44 = idx; + let _e46 = baz.m[_e42][_e44]; + unnamed_6 = _e46; + t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0))); + let _e57 = idx; + idx = (_e57 + 1); + t.m = mat3x2(vec2(6.0), vec2(5.0), vec2(4.0)); + t.m[0] = vec2(9.0); + let _e74 = idx; + t.m[_e74] = vec2(90.0); + t.m[0][1] = 10.0; + let _e87 = idx; + t.m[0][_e87] = 20.0; + let _e91 = idx; + t.m[_e91][1] = 30.0; + let _e97 = idx; + let _e99 = idx; + t.m[_e97][_e99] = 40.0; + return; +} fn read_from_private(foo_1: ptr) -> f32 { - let _e2 = (*foo_1); - return _e2; + let _e3 = (*foo_1); + return _e3; } @stage(vertex) @@ -23,14 +79,15 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0.0; var c: array; - let baz = foo; + let baz_1 = foo; foo = 1.0; + test_matrix_within_struct_accesses(); let matrix = bar.matrix; let arr = bar.arr; let b = bar.matrix[3][0]; let a = bar.data[(arrayLength((&bar.data)) - 2u)].value; let data_pointer = (&bar.data[0].value); - let _e27 = read_from_private((&foo)); + let _e28 = read_from_private((&foo)); c = array(a, i32(b), 3, 4, 5); c[(vi + 1u)] = 42; let value = c[vi]; @@ -51,22 +108,22 @@ fn atomics() { var tmp: i32; let value_1 = atomicLoad((&bar.atom)); - let _e6 = atomicAdd((&bar.atom), 5); - tmp = _e6; - let _e9 = atomicSub((&bar.atom), 5); - tmp = _e9; - let _e12 = atomicAnd((&bar.atom), 5); - tmp = _e12; - let _e15 = atomicOr((&bar.atom), 5); - tmp = _e15; - let _e18 = atomicXor((&bar.atom), 5); - tmp = _e18; - let _e21 = atomicMin((&bar.atom), 5); - tmp = _e21; - let _e24 = atomicMax((&bar.atom), 5); - tmp = _e24; - let _e27 = atomicExchange((&bar.atom), 5); - tmp = _e27; + let _e7 = atomicAdd((&bar.atom), 5); + tmp = _e7; + let _e10 = atomicSub((&bar.atom), 5); + tmp = _e10; + let _e13 = atomicAnd((&bar.atom), 5); + tmp = _e13; + let _e16 = atomicOr((&bar.atom), 5); + tmp = _e16; + let _e19 = atomicXor((&bar.atom), 5); + tmp = _e19; + let _e22 = atomicMin((&bar.atom), 5); + tmp = _e22; + let _e25 = atomicMax((&bar.atom), 5); + tmp = _e25; + let _e28 = atomicExchange((&bar.atom), 5); + tmp = _e28; atomicStore((&bar.atom), value_1); return; }