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>,
}
#[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.
///
/// 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`
///
/// <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);
}
}
// 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
- 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<help::WrappedArrayLength>,
image_queries: crate::FastHashSet<help::WrappedImageQuery>,
constructors: crate::FastHashSet<help::WrappedConstructor>,
struct_matrix_access: crate::FastHashSet<help::WrappedStructMatrixAccess>,
}
impl Wrapped {
@@ -176,6 +186,7 @@ impl Wrapped {
self.array_lengths.clear();
self.image_queries.clear();
self.constructors.clear();
self.struct_matrix_access.clear();
}
}

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

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

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

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

37
tests/in/access.wgsl
View File

@@ -15,6 +15,41 @@ struct Bar {
@group(0) @binding(0)
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 {
return *foo;
}
@@ -26,6 +61,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let baz: f32 = foo;
foo = 1.0;
test_matrix_within_struct_accesses();
// test storage loads
let matrix = bar.matrix;
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 {
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;
}

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

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

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

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

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

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

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

@@ -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<int>(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<int>(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::float4>(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;
}

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

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

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

@@ -10,12 +10,68 @@ struct Bar {
data: array<AlignedWrapper>,
}
struct Baz {
m: mat3x2<f32>,
}
@group(0) @binding(0)
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 {
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<f32> {
var foo: f32 = 0.0;
var c: array<i32,5>;
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<i32,5>(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;
}