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] fix matCx2 translation for uniform buffers (#1802)

Browse Source 
* [hlsl-out] fix matCx2 translation

* fix msl validation (warning: unused variable)

* fix msl buffer declaration

* address comments
This commit is contained in:
Teodor Tanasoaia
2022-04-11 07:29:11 +02:00
committed by GitHub
parent 5ba2d4d079
commit 7aaac25fbf
12 changed files with 1344 additions and 291 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -48,6 +48,12 @@ pub(super) struct WrappedConstructor {
pub(super) ty: Handle<crate::Type>, pub(super) ty: Handle<crate::Type>,
} }
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) struct WrappedStructMatrixAccess {
pub(super) ty: Handle<crate::Type>,
pub(super) index: u32,
}
/// HLSL backend requires its own `ImageQuery` enum. /// HLSL backend requires its own `ImageQuery` enum.
/// ///
/// It is used inside `WrappedImageQuery` and should be unique per ImageQuery function. /// 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 INDENT, struct_name, RETURN_VARIABLE_NAME
)?; )?;
for i in 0..members.len() as u32 { for i in 0..members.len() as u32 {
let member = &members[i as usize];
let field_name = &self.names[&NameKey::StructMember(constructor.ty, i)]; let field_name = &self.names[&NameKey::StructMember(constructor.ty, i)];
//TODO: handle arrays?
writeln!( match module.types[member.ty].inner {
self.out, crate::TypeInner::Matrix {
"{}{}.{} = {}{};", columns,
INDENT, RETURN_VARIABLE_NAME, field_name, ARGUMENT_VARIABLE_NAME, i, 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 // Write return value
@@ -413,6 +445,335 @@ impl<'a, W: Write> super::Writer<'a, W> {
Ok(()) 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` /// Helper function that write wrapped function for `Expression::ImageQuery` and `Expression::ArrayLength`
/// ///
/// <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions> /// <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-to-getdimensions>
@@ -478,6 +839,47 @@ impl<'a, W: Write> super::Writer<'a, W> {
self.wrapped.constructors.insert(constructor); 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);
}
}
_ => {}
}
}
}
_ => {} _ => {}
}; };
} }

11
src/back/hlsl/mod.rs
View File

@@ -6,6 +6,8 @@ Backend for [HLSL][hlsl] (High-Level Shading Language).
- 5.1 - 5.1
- 6.0 - 6.0
# General Matrix Note
All matrix construction/deconstruction is row based in HLSL. This means that when 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. 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. 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 Finally because all of our matrices are transposed, if you use `mat3x4`, it'll become `float3x4` in HLSL
(HLSL has inverted col/row notation). (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 [hlsl]: https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl
*/ */
@@ -169,6 +178,7 @@ struct Wrapped {
array_lengths: crate::FastHashSet<help::WrappedArrayLength>, array_lengths: crate::FastHashSet<help::WrappedArrayLength>,
image_queries: crate::FastHashSet<help::WrappedImageQuery>, image_queries: crate::FastHashSet<help::WrappedImageQuery>,
constructors: crate::FastHashSet<help::WrappedConstructor>, constructors: crate::FastHashSet<help::WrappedConstructor>,
struct_matrix_access: crate::FastHashSet<help::WrappedStructMatrixAccess>,
} }
impl Wrapped { impl Wrapped {
@@ -176,6 +186,7 @@ impl Wrapped {
self.array_lengths.clear(); self.array_lengths.clear();
self.image_queries.clear(); self.image_queries.clear();
self.constructors.clear(); self.constructors.clear();
self.struct_matrix_access.clear();
} }
} }

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

@@ -1,5 +1,5 @@
use super::{ use super::{
help::{WrappedArrayLength, WrappedConstructor, WrappedImageQuery}, help::{WrappedArrayLength, WrappedConstructor, WrappedImageQuery, WrappedStructMatrixAccess},
storage::StoreValue, storage::StoreValue,
BackendResult, Error, Options, BackendResult, Error, Options,
}; };
@@ -784,6 +784,28 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
// Write [size] // Write [size]
self.write_array_size(module, 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 // Write modifier before type
if let Some(ref binding) = member.binding { 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, "[_i] = _result[_i];")?;
writeln!(self.out, "{}}}", level)?; writeln!(self.out, "{}}}", level)?;
} else { } 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<crate::Expression>,
index: u32,
}
enum Index {
Expression(Handle<crate::Expression>),
Static(u32),
}
let get_members = |expr: Handle<crate::Expression>| {
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)?; write!(self.out, "{}", level)?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, " = ")?; if let Some(MatrixAccess { index, base }) = matrix {
self.write_expr(module, value, func_ctx)?; let base_ty_res = &func_ctx.info[base].ty;
writeln!(self.out, ";")? 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 { Statement::Loop {
@@ -1592,8 +1778,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
{ {
// do nothing, the chain is written on `Load`/`Store` // do nothing, the chain is written on `Load`/`Store`
} else { } else {
self.write_expr(module, base, func_ctx)?;
let base_ty_res = &func_ctx.info[base].ty; let base_ty_res = &func_ctx.info[base].ty;
let mut resolved = base_ty_res.inner_with(&module.types); let mut resolved = base_ty_res.inner_with(&module.types);
let base_ty_handle = match *resolved { let base_ty_handle = match *resolved {
@@ -1604,6 +1788,34 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
_ => base_ty_res.handle(), _ => 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 { match *resolved {
TypeInner::Vector { .. } => { TypeInner::Vector { .. } => {
// Write vector access as a swizzle // Write vector access as a swizzle

1
tests/in/access.param.ron
View File

@@ -10,6 +10,7 @@
vs: ( vs: (
resources: { resources: {
(group: 0, binding: 0): (buffer: Some(0), mutable: false), (group: 0, binding: 0): (buffer: Some(0), mutable: false),
(group: 0, binding: 1): (buffer: Some(1), mutable: false),
}, },
sizes_buffer: Some(24), sizes_buffer: Some(24),
), ),

37
tests/in/access.wgsl
View File

@@ -15,6 +15,41 @@ struct Bar {
@group(0) @binding(0) @group(0) @binding(0)
var<storage,read_write> bar: Bar; var<storage,read_write> bar: Bar;
struct Baz {
m: mat3x2<f32>,
}
@group(0) @binding(1)
var<uniform> 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<f32>(vec2<f32>(1.0), vec2<f32>(2.0), vec2<f32>(3.0)));
idx++;
// stores
t.m = mat3x2<f32>(vec2<f32>(6.0), vec2<f32>(5.0), vec2<f32>(4.0));
t.m[0] = vec2<f32>(9.0);
t.m[idx] = vec2<f32>(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<function, f32>) -> f32 { fn read_from_private(foo: ptr<function, f32>) -> f32 {
return *foo; return *foo;
} }
@@ -26,6 +61,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let baz: f32 = foo; let baz: f32 = foo;
foo = 1.0; foo = 1.0;
test_matrix_within_struct_accesses();
// test storage loads // test storage loads
let matrix = bar.matrix; let matrix = bar.matrix;
let arr = bar.arr; let arr = bar.arr;

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

@@ -8,6 +8,9 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct AlignedWrapper { struct AlignedWrapper {
int value; int value;
}; };
struct Baz {
mat3x2 m;
};
layout(std430) buffer Bar_block_0Compute { layout(std430) buffer Bar_block_0Compute {
mat4x3 matrix; mat4x3 matrix;
mat2x2 matrix_array[2]; mat2x2 matrix_array[2];
@@ -18,29 +21,29 @@ layout(std430) buffer Bar_block_0Compute {
float read_from_private(inout float foo_1) { float read_from_private(inout float foo_1) {
float _e2 = foo_1; float _e3 = foo_1;
return _e2; return _e3;
} }
void main() { void main() {
int tmp = 0; int tmp = 0;
int value = _group_0_binding_0_cs.atom; int value = _group_0_binding_0_cs.atom;
int _e6 = atomicAdd(_group_0_binding_0_cs.atom, 5); int _e7 = atomicAdd(_group_0_binding_0_cs.atom, 5);
tmp = _e6; tmp = _e7;
int _e9 = atomicAdd(_group_0_binding_0_cs.atom, -5); int _e10 = atomicAdd(_group_0_binding_0_cs.atom, -5);
tmp = _e9; tmp = _e10;
int _e12 = atomicAnd(_group_0_binding_0_cs.atom, 5); int _e13 = atomicAnd(_group_0_binding_0_cs.atom, 5);
tmp = _e12; tmp = _e13;
int _e15 = atomicOr(_group_0_binding_0_cs.atom, 5); int _e16 = atomicOr(_group_0_binding_0_cs.atom, 5);
tmp = _e15; tmp = _e16;
int _e18 = atomicXor(_group_0_binding_0_cs.atom, 5); int _e19 = atomicXor(_group_0_binding_0_cs.atom, 5);
tmp = _e18; tmp = _e19;
int _e21 = atomicMin(_group_0_binding_0_cs.atom, 5); int _e22 = atomicMin(_group_0_binding_0_cs.atom, 5);
tmp = _e21; tmp = _e22;
int _e24 = atomicMax(_group_0_binding_0_cs.atom, 5); int _e25 = atomicMax(_group_0_binding_0_cs.atom, 5);
tmp = _e24; tmp = _e25;
int _e27 = atomicExchange(_group_0_binding_0_cs.atom, 5); int _e28 = atomicExchange(_group_0_binding_0_cs.atom, 5);
tmp = _e27; tmp = _e28;
_group_0_binding_0_cs.atom = value; _group_0_binding_0_cs.atom = value;
return; return;
} }

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

@@ -6,6 +6,9 @@ precision highp int;
struct AlignedWrapper { struct AlignedWrapper {
int value; int value;
}; };
struct Baz {
mat3x2 m;
};
layout(std430) buffer Bar_block_0Fragment { layout(std430) buffer Bar_block_0Fragment {
mat4x3 matrix; mat4x3 matrix;
mat2x2 matrix_array[2]; mat2x2 matrix_array[2];
@@ -17,8 +20,8 @@ layout(std430) buffer Bar_block_0Fragment {
layout(location = 0) out vec4 _fs2p_location0; layout(location = 0) out vec4 _fs2p_location0;
float read_from_private(inout float foo_1) { float read_from_private(inout float foo_1) {
float _e2 = foo_1; float _e3 = foo_1;
return _e2; return _e3;
} }
void main() { void main() {

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

@@ -6,6 +6,9 @@ precision highp int;
struct AlignedWrapper { struct AlignedWrapper {
int value; int value;
}; };
struct Baz {
mat3x2 m;
};
layout(std430) buffer Bar_block_0Vertex { layout(std430) buffer Bar_block_0Vertex {
mat4x3 matrix; mat4x3 matrix;
mat2x2 matrix_array[2]; mat2x2 matrix_array[2];
@@ -14,23 +17,75 @@ layout(std430) buffer Bar_block_0Vertex {
AlignedWrapper data[]; AlignedWrapper data[];
} _group_0_binding_0_vs; } _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 read_from_private(inout float foo_1) {
float _e2 = foo_1; float _e3 = foo_1;
return _e2; return _e3;
} }
void main() { void main() {
uint vi = uint(gl_VertexID); uint vi = uint(gl_VertexID);
float foo = 0.0; float foo = 0.0;
int c[5] = int[5](0, 0, 0, 0, 0); int c[5] = int[5](0, 0, 0, 0, 0);
float baz = foo; float baz_1 = foo;
foo = 1.0; foo = 1.0;
test_matrix_within_struct_accesses();
mat4x3 matrix = _group_0_binding_0_vs.matrix; mat4x3 matrix = _group_0_binding_0_vs.matrix;
uvec2 arr[2] = _group_0_binding_0_vs.arr; uvec2 arr[2] = _group_0_binding_0_vs.arr;
float b = _group_0_binding_0_vs.matrix[3][0]; 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 = _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 = int[5](a, int(b), 3, 4, 5);
c[(vi + 1u)] = 42; c[(vi + 1u)] = 42;
int value = c[vi]; int value = c[vi];

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

@@ -3,12 +3,102 @@ struct AlignedWrapper {
int value; int value;
}; };
struct Baz {
float2 m_0; float2 m_1; float2 m_2;
};
RWByteAddressBuffer bar : register(u0); 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 read_from_private(inout float foo_1)
{ {
float _expr2 = foo_1; float _expr3 = foo_1;
return _expr2; return _expr3;
} }
uint NagaBufferLengthRW(RWByteAddressBuffer buffer) uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
@@ -23,13 +113,14 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
float foo = 0.0; float foo = 0.0;
int c[5] = {(int)0,(int)0,(int)0,(int)0,(int)0}; int c[5] = {(int)0,(int)0,(int)0,(int)0,(int)0};
float baz = foo; float baz_1 = foo;
foo = 1.0; 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))); 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))}; uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))};
float b = asfloat(bar.Load(0+48+0)); float b = asfloat(bar.Load(0+48+0));
int a = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120)); 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 }; int _result[5]={ a, int(b), 3, 4, 5 };
for(int _i=0; _i<5; ++_i) c[_i] = _result[_i]; for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
@@ -64,22 +155,22 @@ void atomics()
int tmp = (int)0; int tmp = (int)0;
int value_1 = asint(bar.Load(96)); int value_1 = asint(bar.Load(96));
int _e6; bar.InterlockedAdd(96, 5, _e6); int _e7; bar.InterlockedAdd(96, 5, _e7);
tmp = _e6; tmp = _e7;
int _e9; bar.InterlockedAdd(96, -5, _e9); int _e10; bar.InterlockedAdd(96, -5, _e10);
tmp = _e9; tmp = _e10;
int _e12; bar.InterlockedAnd(96, 5, _e12); int _e13; bar.InterlockedAnd(96, 5, _e13);
tmp = _e12; tmp = _e13;
int _e15; bar.InterlockedOr(96, 5, _e15); int _e16; bar.InterlockedOr(96, 5, _e16);
tmp = _e15; tmp = _e16;
int _e18; bar.InterlockedXor(96, 5, _e18); int _e19; bar.InterlockedXor(96, 5, _e19);
tmp = _e18; tmp = _e19;
int _e21; bar.InterlockedMin(96, 5, _e21); int _e22; bar.InterlockedMin(96, 5, _e22);
tmp = _e21; tmp = _e22;
int _e24; bar.InterlockedMax(96, 5, _e24); int _e25; bar.InterlockedMax(96, 5, _e25);
tmp = _e24; tmp = _e25;
int _e27; bar.InterlockedExchange(96, 5, _e27); int _e28; bar.InterlockedExchange(96, 5, _e28);
tmp = _e27; tmp = _e28;
bar.Store(96, asuint(value_1)); bar.Store(96, asuint(value_1));
return; return;
} }

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

@@ -26,15 +26,69 @@ struct Bar {
type_6 arr; type_6 arr;
type_7 data; type_7 data;
}; };
struct type_13 { struct Baz {
metal::float3x2 m;
};
struct type_15 {
int inner[5]; 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( float read_from_private(
thread float& foo_1 thread float& foo_1
) { ) {
float _e2 = foo_1; float _e3 = foo_1;
return _e2; return _e3;
} }
struct foo_vertInput { struct foo_vertInput {
@@ -45,18 +99,20 @@ struct foo_vertOutput {
vertex foo_vertOutput foo_vert( vertex foo_vertOutput foo_vert(
uint vi [[vertex_id]] uint vi [[vertex_id]]
, device Bar const& bar [[buffer(0)]] , device Bar const& bar [[buffer(0)]]
, constant Baz& baz [[buffer(1)]]
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) { ) {
float foo = 0.0; float foo = 0.0;
type_13 c; type_15 c;
float baz = foo; float baz_1 = foo;
foo = 1.0; foo = 1.0;
test_matrix_within_struct_accesses(baz);
metal::float4x3 matrix = bar.matrix; metal::float4x3 matrix = bar.matrix;
type_6 arr = bar.arr; type_6 arr = bar.arr;
float b = bar.matrix[3].x; float b = bar.matrix[3].x;
int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value; int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value;
float _e27 = read_from_private(foo); float _e28 = read_from_private(foo);
for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast<int>(b), 3, 4, 5}.inner[_i]; for(int _i=0; _i<5; ++_i) c.inner[_i] = type_15 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
c.inner[vi + 1u] = 42; c.inner[vi + 1u] = 42;
int value = c.inner[vi]; int value = c.inner[vi];
return foo_vertOutput { metal::float4(matrix * static_cast<metal::float4>(metal::int4(value)), 2.0) }; return foo_vertOutput { metal::float4(matrix * static_cast<metal::float4>(metal::int4(value)), 2.0) };
@@ -84,22 +140,22 @@ kernel void atomics(
) { ) {
int tmp; int tmp;
int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed); 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); int _e7 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e6; tmp = _e7;
int _e9 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e10 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e9; tmp = _e10;
int _e12 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e13 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e12; tmp = _e13;
int _e15 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e16 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e15; tmp = _e16;
int _e18 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e19 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e18; tmp = _e19;
int _e21 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e22 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e21; tmp = _e22;
int _e24 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e25 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e24; tmp = _e25;
int _e27 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); int _e28 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e27; tmp = _e28;
metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed); metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed);
return; return;
} }

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

@@ -1,208 +1,333 @@
; SPIR-V ; SPIR-V
; Version: 1.1 ; Version: 1.1
; Generator: rspirv ; Generator: rspirv
; Bound: 128 ; Bound: 213
OpCapability Shader OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class" OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450" %1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %52 "foo_vert" %47 %50 OpEntryPoint Vertex %137 "foo_vert" %132 %135
OpEntryPoint Fragment %88 "foo_frag" %87 OpEntryPoint Fragment %173 "foo_frag" %172
OpEntryPoint GLCompute %105 "atomics" OpEntryPoint GLCompute %190 "atomics"
OpExecutionMode %88 OriginUpperLeft OpExecutionMode %173 OriginUpperLeft
OpExecutionMode %105 LocalSize 1 1 1 OpExecutionMode %190 LocalSize 1 1 1
OpSource GLSL 450 OpSource GLSL 450
OpMemberName %21 0 "value" OpMemberName %31 0 "value"
OpName %21 "AlignedWrapper" OpName %31 "AlignedWrapper"
OpMemberName %30 0 "matrix" OpMemberName %40 0 "matrix"
OpMemberName %30 1 "matrix_array" OpMemberName %40 1 "matrix_array"
OpMemberName %30 2 "atom" OpMemberName %40 2 "atom"
OpMemberName %30 3 "arr" OpMemberName %40 3 "arr"
OpMemberName %30 4 "data" OpMemberName %40 4 "data"
OpName %30 "Bar" OpName %40 "Bar"
OpName %35 "bar" OpMemberName %42 0 "m"
OpName %38 "foo" OpName %42 "Baz"
OpName %39 "read_from_private" OpName %47 "bar"
OpName %43 "foo" OpName %49 "baz"
OpName %44 "c" OpName %52 "idx"
OpName %47 "vi" OpName %54 "_"
OpName %52 "foo_vert" OpName %56 "_"
OpName %88 "foo_frag" OpName %58 "_"
OpName %103 "tmp" OpName %59 "_"
OpName %105 "atomics" OpName %60 "_"
OpMemberDecorate %21 0 Offset 0 OpName %61 "_"
OpDecorate %26 ArrayStride 16 OpName %62 "_"
OpDecorate %28 ArrayStride 8 OpName %63 "t"
OpDecorate %29 ArrayStride 8 OpName %66 "test_matrix_within_struct_accesses"
OpMemberDecorate %30 0 Offset 0 OpName %123 "foo"
OpMemberDecorate %30 0 ColMajor OpName %124 "read_from_private"
OpMemberDecorate %30 0 MatrixStride 16 OpName %128 "foo"
OpMemberDecorate %30 1 Offset 64 OpName %129 "c"
OpMemberDecorate %30 1 ColMajor OpName %132 "vi"
OpMemberDecorate %30 1 MatrixStride 8 OpName %137 "foo_vert"
OpMemberDecorate %30 2 Offset 96 OpName %173 "foo_frag"
OpMemberDecorate %30 3 Offset 104 OpName %188 "tmp"
OpMemberDecorate %30 4 Offset 120 OpName %190 "atomics"
OpDecorate %34 ArrayStride 4 OpMemberDecorate %31 0 Offset 0
OpDecorate %35 DescriptorSet 0 OpDecorate %36 ArrayStride 16
OpDecorate %35 Binding 0 OpDecorate %38 ArrayStride 8
OpDecorate %30 Block OpDecorate %39 ArrayStride 8
OpDecorate %47 BuiltIn VertexIndex OpMemberDecorate %40 0 Offset 0
OpDecorate %50 BuiltIn Position OpMemberDecorate %40 0 ColMajor
OpDecorate %87 Location 0 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 %2 = OpTypeVoid
%4 = OpTypeInt 32 1 %4 = OpTypeInt 32 1
%3 = OpConstant %4 2 %3 = OpConstant %4 2
%6 = OpTypeFloat 32 %5 = OpConstant %4 9
%5 = OpConstant %6 0.0 %6 = OpConstant %4 1
%7 = OpConstant %6 1.0 %7 = OpConstant %4 0
%9 = OpTypeInt 32 0 %9 = OpTypeFloat 32
%8 = OpConstant %9 3 %8 = OpConstant %9 1.0
%10 = OpConstant %9 2 %10 = OpConstant %9 2.0
%11 = OpConstant %4 0 %11 = OpConstant %9 3.0
%12 = OpConstant %4 5 %12 = OpConstant %9 6.0
%13 = OpConstant %4 3 %13 = OpConstant %9 5.0
%14 = OpConstant %4 4 %14 = OpConstant %9 4.0
%15 = OpConstant %9 1 %15 = OpConstant %9 9.0
%16 = OpConstant %4 42 %16 = OpConstant %9 90.0
%17 = OpConstant %6 2.0 %17 = OpConstant %9 10.0
%18 = OpConstant %4 1 %18 = OpConstant %9 20.0
%19 = OpConstant %6 3.0 %19 = OpConstant %9 30.0
%20 = OpConstant %9 0 %20 = OpConstant %9 40.0
%21 = OpTypeStruct %4 %21 = OpConstant %9 0.0
%23 = OpTypeVector %6 3 %23 = OpTypeInt 32 0
%22 = OpTypeMatrix %23 4 %22 = OpConstant %23 3
%25 = OpTypeVector %6 2 %24 = OpConstant %23 2
%24 = OpTypeMatrix %25 2 %25 = OpConstant %4 5
%26 = OpTypeArray %24 %3 %26 = OpConstant %4 3
%27 = OpTypeVector %9 2 %27 = OpConstant %4 4
%28 = OpTypeArray %27 %3 %28 = OpConstant %23 1
%29 = OpTypeRuntimeArray %21 %29 = OpConstant %4 42
%30 = OpTypeStruct %22 %26 %4 %28 %29 %30 = OpConstant %23 0
%31 = OpTypePointer Function %6 %31 = OpTypeStruct %4
%32 = OpTypeVector %6 4 %33 = OpTypeVector %9 3
%33 = OpTypePointer StorageBuffer %4 %32 = OpTypeMatrix %33 4
%34 = OpTypeArray %4 %12 %35 = OpTypeVector %9 2
%36 = OpTypePointer StorageBuffer %30 %34 = OpTypeMatrix %35 2
%35 = OpVariable %36 StorageBuffer %36 = OpTypeArray %34 %3
%40 = OpTypeFunction %6 %31 %37 = OpTypeVector %23 2
%45 = OpTypePointer Function %34 %38 = OpTypeArray %37 %3
%48 = OpTypePointer Input %9 %39 = OpTypeRuntimeArray %31
%47 = OpVariable %48 Input %40 = OpTypeStruct %32 %36 %4 %38 %39
%51 = OpTypePointer Output %32 %41 = OpTypeMatrix %35 3
%50 = OpVariable %51 Output %42 = OpTypeStruct %41
%53 = OpTypeFunction %2 %43 = OpTypePointer Function %9
%56 = OpTypePointer StorageBuffer %22 %44 = OpTypeVector %9 4
%59 = OpTypePointer StorageBuffer %28 %45 = OpTypePointer StorageBuffer %4
%62 = OpTypePointer StorageBuffer %23 %46 = OpTypeArray %4 %25
%63 = OpTypePointer StorageBuffer %6 %48 = OpTypePointer StorageBuffer %40
%66 = OpTypePointer StorageBuffer %29 %47 = OpVariable %48 StorageBuffer
%69 = OpTypePointer StorageBuffer %21 %50 = OpTypeStruct %42
%70 = OpConstant %9 4 %51 = OpTypePointer Uniform %50
%77 = OpTypePointer Function %4 %49 = OpVariable %51 Uniform
%81 = OpTypeVector %4 4 %53 = OpTypePointer Function %4
%87 = OpVariable %51 Output %55 = OpTypePointer Function %41
%107 = OpTypePointer StorageBuffer %4 %57 = OpTypePointer Function %35
%110 = OpConstant %9 64 %64 = OpTypePointer Function %42
%39 = OpFunction %6 None %40 %67 = OpTypeFunction %2
%38 = OpFunctionParameter %31 %68 = OpTypePointer Uniform %42
%37 = OpLabel %73 = OpTypePointer Uniform %41
OpBranch %41 %76 = OpTypePointer Uniform %35
%41 = OpLabel %82 = OpTypePointer Uniform %9
%42 = OpLoad %6 %38 %107 = OpTypePointer Function %35
OpReturnValue %42 %113 = OpTypePointer Function %9
OpFunctionEnd %125 = OpTypeFunction %9 %43
%52 = OpFunction %2 None %53 %130 = OpTypePointer Function %46
%46 = OpLabel %133 = OpTypePointer Input %23
%43 = OpVariable %31 Function %5 %132 = OpVariable %133 Input
%44 = OpVariable %45 Function %136 = OpTypePointer Output %44
%49 = OpLoad %9 %47 %135 = OpVariable %136 Output
OpBranch %54 %142 = OpTypePointer StorageBuffer %32
%54 = OpLabel %145 = OpTypePointer StorageBuffer %38
%55 = OpLoad %6 %43 %148 = OpTypePointer StorageBuffer %33
OpStore %43 %7 %149 = OpTypePointer StorageBuffer %9
%57 = OpAccessChain %56 %35 %20 %152 = OpTypePointer StorageBuffer %39
%58 = OpLoad %22 %57 %155 = OpTypePointer StorageBuffer %31
%60 = OpAccessChain %59 %35 %8 %156 = OpConstant %23 4
%61 = OpLoad %28 %60 %166 = OpTypeVector %4 4
%64 = OpAccessChain %63 %35 %20 %8 %20 %172 = OpVariable %136 Output
%65 = OpLoad %6 %64 %192 = OpTypePointer StorageBuffer %4
%67 = OpArrayLength %9 %35 4 %195 = OpConstant %23 64
%68 = OpISub %9 %67 %10 %66 = OpFunction %2 None %67
%71 = OpAccessChain %33 %35 %70 %68 %20 %65 = OpLabel
%72 = OpLoad %4 %71 %61 = OpVariable %43 Function
%73 = OpFunctionCall %6 %39 %43 %58 = OpVariable %57 Function
%74 = OpConvertFToS %4 %65 %52 = OpVariable %53 Function %5
%75 = OpCompositeConstruct %34 %72 %74 %13 %14 %12 %62 = OpVariable %43 Function
OpStore %44 %75 %59 = OpVariable %43 Function
%76 = OpIAdd %9 %49 %15 %54 = OpVariable %55 Function
%78 = OpAccessChain %77 %44 %76 %63 = OpVariable %64 Function
OpStore %78 %16 %60 = OpVariable %43 Function
%79 = OpAccessChain %77 %44 %49 %56 = OpVariable %57 Function
%80 = OpLoad %4 %79 %69 = OpAccessChain %68 %49 %30
%82 = OpCompositeConstruct %81 %80 %80 %80 %80 OpBranch %70
%83 = OpConvertSToF %32 %82 %70 = OpLabel
%84 = OpMatrixTimesVector %23 %58 %83 %71 = OpLoad %4 %52
%85 = OpCompositeConstruct %32 %84 %17 %72 = OpISub %4 %71 %6
OpStore %50 %85 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 OpReturn
OpFunctionEnd OpFunctionEnd
%88 = OpFunction %2 None %53 %124 = OpFunction %9 None %125
%86 = OpLabel %123 = OpFunctionParameter %43
OpBranch %89 %122 = OpLabel
%89 = OpLabel OpBranch %126
%90 = OpAccessChain %63 %35 %20 %15 %10 %126 = OpLabel
OpStore %90 %7 %127 = OpLoad %9 %123
%91 = OpCompositeConstruct %23 %5 %5 %5 OpReturnValue %127
%92 = OpCompositeConstruct %23 %7 %7 %7 OpFunctionEnd
%93 = OpCompositeConstruct %23 %17 %17 %17 %137 = OpFunction %2 None %67
%94 = OpCompositeConstruct %23 %19 %19 %19 %131 = OpLabel
%95 = OpCompositeConstruct %22 %91 %92 %93 %94 %128 = OpVariable %43 Function %21
%96 = OpAccessChain %56 %35 %20 %129 = OpVariable %130 Function
OpStore %96 %95 %134 = OpLoad %23 %132
%97 = OpCompositeConstruct %27 %20 %20 %138 = OpAccessChain %68 %49 %30
%98 = OpCompositeConstruct %27 %15 %15 OpBranch %139
%99 = OpCompositeConstruct %28 %97 %98 %139 = OpLabel
%100 = OpAccessChain %59 %35 %8 %140 = OpLoad %9 %128
OpStore %100 %99 OpStore %128 %8
%101 = OpAccessChain %33 %35 %70 %15 %20 %141 = OpFunctionCall %2 %66
OpStore %101 %18 %143 = OpAccessChain %142 %47 %30
%102 = OpCompositeConstruct %32 %5 %5 %5 %5 %144 = OpLoad %32 %143
OpStore %87 %102 %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 OpReturn
OpFunctionEnd OpFunctionEnd
%105 = OpFunction %2 None %53 %173 = OpFunction %2 None %67
%104 = OpLabel %171 = OpLabel
%103 = OpVariable %77 Function OpBranch %174
OpBranch %106 %174 = OpLabel
%106 = OpLabel %175 = OpAccessChain %149 %47 %30 %28 %24
%108 = OpAccessChain %107 %35 %10 OpStore %175 %8
%109 = OpAtomicLoad %4 %108 %18 %110 %176 = OpCompositeConstruct %33 %21 %21 %21
%112 = OpAccessChain %107 %35 %10 %177 = OpCompositeConstruct %33 %8 %8 %8
%111 = OpAtomicIAdd %4 %112 %18 %110 %12 %178 = OpCompositeConstruct %33 %10 %10 %10
OpStore %103 %111 %179 = OpCompositeConstruct %33 %11 %11 %11
%114 = OpAccessChain %107 %35 %10 %180 = OpCompositeConstruct %32 %176 %177 %178 %179
%113 = OpAtomicISub %4 %114 %18 %110 %12 %181 = OpAccessChain %142 %47 %30
OpStore %103 %113 OpStore %181 %180
%116 = OpAccessChain %107 %35 %10 %182 = OpCompositeConstruct %37 %30 %30
%115 = OpAtomicAnd %4 %116 %18 %110 %12 %183 = OpCompositeConstruct %37 %28 %28
OpStore %103 %115 %184 = OpCompositeConstruct %38 %182 %183
%118 = OpAccessChain %107 %35 %10 %185 = OpAccessChain %145 %47 %22
%117 = OpAtomicOr %4 %118 %18 %110 %12 OpStore %185 %184
OpStore %103 %117 %186 = OpAccessChain %45 %47 %156 %28 %30
%120 = OpAccessChain %107 %35 %10 OpStore %186 %6
%119 = OpAtomicXor %4 %120 %18 %110 %12 %187 = OpCompositeConstruct %44 %21 %21 %21 %21
OpStore %103 %119 OpStore %172 %187
%122 = OpAccessChain %107 %35 %10 OpReturn
%121 = OpAtomicSMin %4 %122 %18 %110 %12 OpFunctionEnd
OpStore %103 %121 %190 = OpFunction %2 None %67
%124 = OpAccessChain %107 %35 %10 %189 = OpLabel
%123 = OpAtomicSMax %4 %124 %18 %110 %12 %188 = OpVariable %53 Function
OpStore %103 %123 OpBranch %191
%126 = OpAccessChain %107 %35 %10 %191 = OpLabel
%125 = OpAtomicExchange %4 %126 %18 %110 %12 %193 = OpAccessChain %192 %47 %24
OpStore %103 %125 %194 = OpAtomicLoad %4 %193 %6 %195
%127 = OpAccessChain %107 %35 %10 %197 = OpAccessChain %192 %47 %24
OpAtomicStore %127 %18 %110 %109 %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 OpReturn
OpFunctionEnd OpFunctionEnd

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

@@ -10,12 +10,68 @@ struct Bar {
data: array<AlignedWrapper>, data: array<AlignedWrapper>,
} }
struct Baz {
m: mat3x2<f32>,
}
@group(0) @binding(0) @group(0) @binding(0)
var<storage, read_write> bar: Bar; var<storage, read_write> bar: Bar;
@group(0) @binding(1)
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_2: vec2<f32>;
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<f32>(vec2<f32>(1.0), vec2<f32>(2.0), vec2<f32>(3.0)));
let _e57 = idx;
idx = (_e57 + 1);
t.m = mat3x2<f32>(vec2<f32>(6.0), vec2<f32>(5.0), vec2<f32>(4.0));
t.m[0] = vec2<f32>(9.0);
let _e74 = idx;
t.m[_e74] = vec2<f32>(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<function, f32>) -> f32 { fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
let _e2 = (*foo_1); let _e3 = (*foo_1);
return _e2; return _e3;
} }
@stage(vertex) @stage(vertex)
@@ -23,14 +79,15 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
var foo: f32 = 0.0; var foo: f32 = 0.0;
var c: array<i32,5>; var c: array<i32,5>;
let baz = foo; let baz_1 = foo;
foo = 1.0; foo = 1.0;
test_matrix_within_struct_accesses();
let matrix = bar.matrix; let matrix = bar.matrix;
let arr = bar.arr; let arr = bar.arr;
let b = bar.matrix[3][0]; let b = bar.matrix[3][0];
let a = bar.data[(arrayLength((&bar.data)) - 2u)].value; let a = bar.data[(arrayLength((&bar.data)) - 2u)].value;
let data_pointer = (&bar.data[0].value); let data_pointer = (&bar.data[0].value);
let _e27 = read_from_private((&foo)); let _e28 = read_from_private((&foo));
c = array<i32,5>(a, i32(b), 3, 4, 5); c = array<i32,5>(a, i32(b), 3, 4, 5);
c[(vi + 1u)] = 42; c[(vi + 1u)] = 42;
let value = c[vi]; let value = c[vi];
@@ -51,22 +108,22 @@ fn atomics() {
var tmp: i32; var tmp: i32;
let value_1 = atomicLoad((&bar.atom)); let value_1 = atomicLoad((&bar.atom));
let _e6 = atomicAdd((&bar.atom), 5); let _e7 = atomicAdd((&bar.atom), 5);
tmp = _e6; tmp = _e7;
let _e9 = atomicSub((&bar.atom), 5); let _e10 = atomicSub((&bar.atom), 5);
tmp = _e9; tmp = _e10;
let _e12 = atomicAnd((&bar.atom), 5); let _e13 = atomicAnd((&bar.atom), 5);
tmp = _e12; tmp = _e13;
let _e15 = atomicOr((&bar.atom), 5); let _e16 = atomicOr((&bar.atom), 5);
tmp = _e15; tmp = _e16;
let _e18 = atomicXor((&bar.atom), 5); let _e19 = atomicXor((&bar.atom), 5);
tmp = _e18; tmp = _e19;
let _e21 = atomicMin((&bar.atom), 5); let _e22 = atomicMin((&bar.atom), 5);
tmp = _e21; tmp = _e22;
let _e24 = atomicMax((&bar.atom), 5); let _e25 = atomicMax((&bar.atom), 5);
tmp = _e24; tmp = _e25;
let _e27 = atomicExchange((&bar.atom), 5); let _e28 = atomicExchange((&bar.atom), 5);
tmp = _e27; tmp = _e28;
atomicStore((&bar.atom), value_1); atomicStore((&bar.atom), value_1);
return; return;
} }