archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

[naga hlsl-out] Handle additional cases of Cx2 matrices

Browse Source 
Fixes #4423
This commit is contained in:
Andy Leiserson 2025-03-22 13:01:11 -07:00
parent c868142709
commit bfa7ee8de5
13 changed files with 1492 additions and 238 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
naga/src/back/hlsl/mod.rs
View File

@ -13,11 +13,17 @@ type should be stored in `uniform` and `storage` buffers. The HLSL we
generate must access values in that form, even when it is not what
HLSL would use normally.
The rules described here only apply to WGSL `uniform` variables. WGSL
`storage` buffers are translated as HLSL `ByteAddressBuffers`, for
which we generate `Load` and `Store` method calls with explicit byte
offsets. WGSL pipeline inputs must be scalars or vectors; they cannot
be matrices, which is where the interesting problems arise.
Matching the WGSL memory layout is a concern only for `uniform`
variables. WGSL `storage` buffers are translated as HLSL
`ByteAddressBuffers`, for which we generate `Load` and `Store` method
calls with explicit byte offsets. WGSL pipeline inputs must be scalars
or vectors; they cannot be matrices, which is where the interesting
problems arise. However, when an affected type appears in a struct
definition, the transformations described here are applied without
consideration of where the struct is used.
Access to storage buffers is implemented in `storage.rs`. Access to
uniform buffers is implemented where applicable in `writer.rs`.
## Row- and column-major ordering for matrices
@ -57,10 +63,9 @@ that the columns of a `matKx2<f32>` need only be [aligned as required
for `vec2<f32>`][ilov], which is [eight-byte alignment][8bb].
To compensate for this, any time a `matKx2<f32>` appears in a WGSL
`uniform` variable, whether directly as the variable's type or as part
of a struct/array, we actually emit `K` separate `float2` members, and
assemble/disassemble the matrix from its columns (in WGSL; rows in
HLSL) upon load and store.
`uniform` value or as part of a struct/array, we actually emit `K`
separate `float2` members, and assemble/disassemble the matrix from its
columns (in WGSL; rows in HLSL) upon load and store.
For example, the following WGSL struct type:

144
naga/src/back/hlsl/storage.rs
View File

@ -108,6 +108,13 @@ pub(super) enum StoreValue {
base: Handle<crate::Type>,
member_index: u32,
},
// Access to a single column of a Cx2 matrix within a struct
TempColumnAccess {
depth: usize,
base: Handle<crate::Type>,
member_index: u32,
column: u32,
},
}
impl<W: fmt::Write> super::Writer<'_, W> {
@ -290,6 +297,15 @@ impl<W: fmt::Write> super::Writer<'_, W> {
let name = &self.names[&NameKey::StructMember(base, member_index)];
write!(self.out, "{STORE_TEMP_NAME}{depth}.{name}")?
}
StoreValue::TempColumnAccess {
depth,
base,
member_index,
column,
} => {
let name = &self.names[&NameKey::StructMember(base, member_index)];
write!(self.out, "{STORE_TEMP_NAME}{depth}.{name}_{column}")?
}
}
Ok(())
}
@ -302,6 +318,7 @@ impl<W: fmt::Write> super::Writer<'_, W> {
value: StoreValue,
func_ctx: &FunctionCtx,
level: crate::back::Level,
within_struct: Option<Handle<crate::Type>>,
) -> BackendResult {
let temp_resolution;
let ty_resolution = match value {
@ -325,6 +342,9 @@ impl<W: fmt::Write> super::Writer<'_, W> {
temp_resolution = TypeResolution::Handle(ty_handle);
&temp_resolution
}
StoreValue::TempColumnAccess { .. } => {
unreachable!("attempting write_storage_store for TempColumnAccess");
}
};
match *ty_resolution.inner_with(&module.types) {
crate::TypeInner::Scalar(scalar) => {
@ -372,37 +392,92 @@ impl<W: fmt::Write> super::Writer<'_, W> {
rows,
scalar,
} => {
// first, assign the value to a temporary
writeln!(self.out, "{level}{{")?;
let depth = level.0 + 1;
write!(
self.out,
"{}{}{}x{} {}{} = ",
level.next(),
scalar.to_hlsl_str()?,
columns as u8,
rows as u8,
STORE_TEMP_NAME,
depth,
)?;
self.write_store_value(module, &value, func_ctx)?;
writeln!(self.out, ";")?;
// Note: Matrices containing vec3s, due to padding, act like they contain vec4s.
let row_stride = Alignment::from(rows) * scalar.width as u32;
// then iterate the stores
for i in 0..columns as u32 {
self.temp_access_chain
.push(SubAccess::Offset(i * row_stride));
let ty_inner = crate::TypeInner::Vector { size: rows, scalar };
let sv = StoreValue::TempIndex {
depth,
index: i,
ty: TypeResolution::Value(ty_inner),
};
self.write_storage_store(module, var_handle, sv, func_ctx, level.next())?;
self.temp_access_chain.pop();
writeln!(self.out, "{level}{{")?;
match within_struct {
Some(containing_struct) if rows == crate::VectorSize::Bi => {
// If we are within a struct, then the struct was already assigned to
// a temporary, we don't need to make another.
let mut chain = mem::take(&mut self.temp_access_chain);
for i in 0..columns as u32 {
chain.push(SubAccess::Offset(i * row_stride));
// working around the borrow checker in `self.write_expr`
let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
let StoreValue::TempAccess { member_index, .. } = value else {
unreachable!(
"write_storage_store within_struct but not TempAccess"
);
};
let column_value = StoreValue::TempColumnAccess {
depth: level.0, // note not incrementing, b/c no temp
base: containing_struct,
member_index,
column: i,
};
// See note about DXC and Load/Store in the module's documentation.
if scalar.width == 4 {
write!(
self.out,
"{}{}.Store{}(",
level.next(),
var_name,
rows as u8
)?;
self.write_storage_address(module, &chain, func_ctx)?;
write!(self.out, ", asuint(")?;
self.write_store_value(module, &column_value, func_ctx)?;
writeln!(self.out, "));")?;
} else {
write!(self.out, "{}{var_name}.Store(", level.next())?;
self.write_storage_address(module, &chain, func_ctx)?;
write!(self.out, ", ")?;
self.write_store_value(module, &column_value, func_ctx)?;
writeln!(self.out, ");")?;
}
chain.pop();
}
self.temp_access_chain = chain;
}
_ => {
// first, assign the value to a temporary
let depth = level.0 + 1;
write!(
self.out,
"{}{}{}x{} {}{} = ",
level.next(),
scalar.to_hlsl_str()?,
columns as u8,
rows as u8,
STORE_TEMP_NAME,
depth,
)?;
self.write_store_value(module, &value, func_ctx)?;
writeln!(self.out, ";")?;
// then iterate the stores
for i in 0..columns as u32 {
self.temp_access_chain
.push(SubAccess::Offset(i * row_stride));
let ty_inner = crate::TypeInner::Vector { size: rows, scalar };
let sv = StoreValue::TempIndex {
depth,
index: i,
ty: TypeResolution::Value(ty_inner),
};
self.write_storage_store(
module,
var_handle,
sv,
func_ctx,
level.next(),
None,
)?;
self.temp_access_chain.pop();
}
}
}
// done
writeln!(self.out, "{level}}}")?;
@ -415,7 +490,7 @@ impl<W: fmt::Write> super::Writer<'_, W> {
// first, assign the value to a temporary
writeln!(self.out, "{level}{{")?;
write!(self.out, "{}", level.next())?;
self.write_value_type(module, &module.types[base].inner)?;
self.write_type(module, base)?;
let depth = level.next().0;
write!(self.out, " {STORE_TEMP_NAME}{depth}")?;
self.write_array_size(module, base, crate::ArraySize::Constant(size))?;
@ -430,7 +505,7 @@ impl<W: fmt::Write> super::Writer<'_, W> {
index: i,
ty: TypeResolution::Handle(base),
};
self.write_storage_store(module, var_handle, sv, func_ctx, level.next())?;
self.write_storage_store(module, var_handle, sv, func_ctx, level.next(), None)?;
self.temp_access_chain.pop();
}
// done
@ -461,7 +536,14 @@ impl<W: fmt::Write> super::Writer<'_, W> {
base: struct_ty,
member_index: i as u32,
};
self.write_storage_store(module, var_handle, sv, func_ctx, level.next())?;
self.write_storage_store(
module,
var_handle,
sv,
func_ctx,
level.next(),
Some(struct_ty),
)?;
self.temp_access_chain.pop();
}
// done

71
naga/src/back/hlsl/writer.rs
View File

@ -1945,6 +1945,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
StoreValue::Expression(value),
func_ctx,
level,
None,
)?;
} else {
// We treat matrices of the form `matCx2` as a sequence of C `vec2`s.
@ -2963,6 +2964,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
rows: crate::VectorSize::Bi,
width: 4,
}) = get_inner_matrix_of_struct_array_member(module, base, func_ctx, true)
.or_else(|| get_global_uniform_matrix(module, base, func_ctx))
{
write!(self.out, "__get_col_of_mat{}x2(", columns as u8)?;
self.write_expr(module, base, func_ctx)?;
@ -3075,13 +3077,15 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
{
// do nothing, the chain is written on `Load`/`Store`
} else {
// We write the matrix column access in a special way since
// the type of `base` is our special __matCx2 struct.
// See if we need to write the matrix column access in a
// special way since the type of `base` is our special
// __matCx2 struct.
if let Some(MatrixType {
rows: crate::VectorSize::Bi,
width: 4,
..
}) = get_inner_matrix_of_struct_array_member(module, base, func_ctx, true)
.or_else(|| get_global_uniform_matrix(module, base, func_ctx))
{
self.write_expr(module, base, func_ctx)?;
write!(self.out, "._{index}")?;
@ -3381,8 +3385,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
.or_else(|| get_inner_matrix_of_global_uniform(module, pointer, func_ctx))
{
let mut resolved = func_ctx.resolve_type(pointer, &module.types);
if let TypeInner::Pointer { base, .. } = *resolved {
resolved = &module.types[base].inner;
let ptr_tr = resolved.pointer_base_type();
if let Some(ptr_ty) =
ptr_tr.as_ref().map(|tr| tr.inner_with(&module.types))
{
resolved = ptr_ty;
}
write!(self.out, "((")?;
@ -4416,6 +4423,32 @@ pub(super) fn get_inner_matrix_data(
}
}
fn find_matrix_in_access_chain(
module: &Module,
base: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
) -> Option<Handle<crate::Expression>> {
let mut current_base = base;
loop {
let resolved_tr = func_ctx
.resolve_type(current_base, &module.types)
.pointer_base_type();
let resolved = resolved_tr.as_ref()?.inner_with(&module.types);
match *resolved {
TypeInner::Scalar(_) | TypeInner::Vector { .. } => {}
TypeInner::Matrix { .. } => return Some(current_base),
_ => return None,
}
current_base = match func_ctx.expressions[current_base] {
crate::Expression::Access { base, .. } => base,
crate::Expression::AccessIndex { base, .. } => base,
_ => return None,
}
}
}
/// Returns the matrix data if the access chain starting at `base`:
/// - starts with an expression with resolved type of [`TypeInner::Matrix`] if `direct = true`
/// - contains one or more expressions with resolved type of [`TypeInner::Array`] of [`TypeInner::Matrix`]
@ -4474,6 +4507,36 @@ pub(super) fn get_inner_matrix_of_struct_array_member(
None
}
/// Simpler version of get_inner_matrix_of_global_uniform that only looks at the
/// immediate expression, rather than traversing an access chain.
fn get_global_uniform_matrix(
module: &Module,
base: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
) -> Option<MatrixType> {
let base_tr = func_ctx
.resolve_type(base, &module.types)
.pointer_base_type();
let base_ty = base_tr.as_ref().map(|tr| tr.inner_with(&module.types));
match (&func_ctx.expressions[base], base_ty) {
(
&crate::Expression::GlobalVariable(handle),
Some(&TypeInner::Matrix {
columns,
rows,
scalar,
}),
) if module.global_variables[handle].space == crate::AddressSpace::Uniform => {
Some(MatrixType {
columns,
rows,
width: scalar.width,
})
}
_ => None,
}
}
/// Returns the matrix data if the access chain starting at `base`:
/// - starts with an expression with resolved type of [`TypeInner::Matrix`]
/// - contains zero or more expressions with resolved type of [`TypeInner::Array`] of [`TypeInner::Matrix`]

3
naga/tests/in/wgsl/access.wgsl
View File

@ -35,6 +35,9 @@ var<uniform> baz: Baz;
var<storage,read_write> qux: vec2<i32>;
fn test_matrix_within_struct_accesses() {
// Test HLSL accesses to Cx2 matrices. There are additional tests
// in `hlsl_mat_cx2.wgsl`.
var idx = 1;
idx--;

1
naga/tests/in/wgsl/hlsl_mat_cx2.toml Normal file
View File

@ -0,0 +1 @@
targets = "HLSL"

177
naga/tests/in/wgsl/hlsl_mat_cx2.wgsl Normal file
View File

@ -0,0 +1,177 @@
// Test HLSL handling of N-by-2 matrices.
// See the doc comment on `naga::back::hlsl` for details.
//
// There are additional tests in `access.wgsl`.
//
// Tests that we don't apply this handling to other sizes are in hlsl_mat_cx3.wgsl.
// Access type (3rd item in variable names)
// S = Struct
// M = Matrix
// C = Column
// E = Element
// Index type (4th item in variable names)
// C = Constant
// V = Variable
alias Mat = mat2x2<f32>;
@group(0) @binding(0)
var<storage, read_write> s_m: Mat;
@group(0) @binding(1)
var<uniform> u_m: Mat;
fn access_m() {
var idx = 1;
idx--;
// loads from storage
let l_s_m = s_m;
let l_s_c_c = s_m[0];
let l_s_c_v = s_m[idx];
let l_s_e_cc = s_m[0][0];
let l_s_e_cv = s_m[0][idx];
let l_s_e_vc = s_m[idx][0];
let l_s_e_vv = s_m[idx][idx];
// loads from uniform
let l_u_m = u_m;
let l_u_c_c = u_m[0];
let l_u_c_v = u_m[idx];
let l_u_e_cc = u_m[0][0];
let l_u_e_cv = u_m[0][idx];
let l_u_e_vc = u_m[idx][0];
let l_u_e_vv = u_m[idx][idx];
// stores to storage
s_m = l_u_m;
s_m[0] = l_u_c_c;
s_m[idx] = l_u_c_v;
s_m[0][0] = l_u_e_cc;
s_m[0][idx] = l_u_e_cv;
s_m[idx][0] = l_u_e_vc;
s_m[idx][idx] = l_u_e_vv;
}
struct StructWithMat {
m: Mat,
}
@group(1) @binding(0)
var<storage, read_write> s_sm: StructWithMat;
@group(1) @binding(1)
var<uniform> u_sm: StructWithMat;
fn access_sm() {
var idx = 1;
idx--;
// loads from storage
let l_s_s = s_sm;
let l_s_m = s_sm.m;
let l_s_c_c = s_sm.m[0];
let l_s_c_v = s_sm.m[idx];
let l_s_e_cc = s_sm.m[0][0];
let l_s_e_cv = s_sm.m[0][idx];
let l_s_e_vc = s_sm.m[idx][0];
let l_s_e_vv = s_sm.m[idx][idx];
// loads from uniform
let l_u_s = u_sm;
let l_u_m = u_sm.m;
let l_u_c_c = u_sm.m[0];
let l_u_c_v = u_sm.m[idx];
let l_u_e_cc = u_sm.m[0][0];
let l_u_e_cv = u_sm.m[0][idx];
let l_u_e_vc = u_sm.m[idx][0];
let l_u_e_vv = u_sm.m[idx][idx];
// stores to storage
s_sm = l_u_s;
s_sm.m = l_u_m;
s_sm.m[0] = l_u_c_c;
s_sm.m[idx] = l_u_c_v;
s_sm.m[0][0] = l_u_e_cc;
s_sm.m[0][idx] = l_u_e_cv;
s_sm.m[idx][0] = l_u_e_vc;
s_sm.m[idx][idx] = l_u_e_vv;
}
struct StructWithArrayOfStructOfMat {
a: array<StructWithMat, 4>,
}
@group(2) @binding(0)
var<storage, read_write> s_sasm: StructWithArrayOfStructOfMat;
@group(2) @binding(1)
var<uniform> u_sasm: StructWithArrayOfStructOfMat;
fn access_sasm() {
var idx = 1;
idx--;
// loads from storage
let l_s_s = s_sasm;
let l_s_a = s_sasm.a;
let l_s_m_c = s_sasm.a[0].m;
let l_s_m_v = s_sasm.a[idx].m;
let l_s_c_cc = s_sasm.a[0].m[0];
let l_s_c_cv = s_sasm.a[0].m[idx];
let l_s_c_vc = s_sasm.a[idx].m[0];
let l_s_c_vv = s_sasm.a[idx].m[idx];
let l_s_e_ccc = s_sasm.a[0].m[0][0];
let l_s_e_ccv = s_sasm.a[0].m[0][idx];
let l_s_e_cvc = s_sasm.a[0].m[idx][0];
let l_s_e_cvv = s_sasm.a[0].m[idx][idx];
let l_s_e_vcc = s_sasm.a[idx].m[0][0];
let l_s_e_vcv = s_sasm.a[idx].m[0][idx];
let l_s_e_vvc = s_sasm.a[idx].m[idx][0];
let l_s_e_vvv = s_sasm.a[idx].m[idx][idx];
// loads from uniform
let l_u_s = u_sasm;
let l_u_a = u_sasm.a;
let l_u_m_c = u_sasm.a[0].m;
let l_u_m_v = u_sasm.a[idx].m;
let l_u_c_cc = u_sasm.a[0].m[0];
let l_u_c_cv = u_sasm.a[0].m[idx];
let l_u_c_vc = u_sasm.a[idx].m[0];
let l_u_c_vv = u_sasm.a[idx].m[idx];
let l_u_e_ccc = u_sasm.a[0].m[0][0];
let l_u_e_ccv = u_sasm.a[0].m[0][idx];
let l_u_e_cvc = u_sasm.a[0].m[idx][0];
let l_u_e_cvv = u_sasm.a[0].m[idx][idx];
let l_u_e_vcc = u_sasm.a[idx].m[0][0];
let l_u_e_vcv = u_sasm.a[idx].m[0][idx];
let l_u_e_vvc = u_sasm.a[idx].m[idx][0];
let l_u_e_vvv = u_sasm.a[idx].m[idx][idx];
// stores to storage
s_sasm = l_u_s;
s_sasm.a = l_u_a;
s_sasm.a[0].m = l_u_m_c;
s_sasm.a[idx].m = l_u_m_v;
s_sasm.a[0].m[0] = l_u_c_cc;
s_sasm.a[0].m[idx] = l_u_c_cv;
s_sasm.a[idx].m[0] = l_u_c_vc;
s_sasm.a[idx].m[idx] = l_u_c_vv;
s_sasm.a[0].m[0][0] = l_u_e_ccc;
s_sasm.a[0].m[0][idx] = l_u_e_ccv;
s_sasm.a[0].m[idx][0] = l_u_e_cvc;
s_sasm.a[0].m[idx][idx] = l_u_e_cvv;
s_sasm.a[idx].m[0][0] = l_u_e_vcc;
s_sasm.a[idx].m[0][idx] = l_u_e_vcv;
s_sasm.a[idx].m[idx][0] = l_u_e_vvc;
s_sasm.a[idx].m[idx][idx] = l_u_e_vvv;
}
@compute @workgroup_size(1)
fn main() {
access_m();
access_sm();
access_sasm();
}

1
naga/tests/in/wgsl/hlsl_mat_cx3.toml Normal file
View File

@ -0,0 +1 @@
targets = "HLSL"

173
naga/tests/in/wgsl/hlsl_mat_cx3.wgsl Normal file
View File

@ -0,0 +1,173 @@
// Test HLSL handling of N-by-3 matrices. These should not receive the special
// treatment that N-by-2 matrices receive (which is tested in hlsl_mat_cx2).
// Access type (3rd item in variable names)
// S = Struct
// M = Matrix
// C = Column
// E = Element
// Index type (4th item in variable names)
// C = Constant
// V = Variable
alias Mat = mat3x3<f32>;
@group(0) @binding(0)
var<storage, read_write> s_m: Mat;
@group(0) @binding(1)
var<uniform> u_m: Mat;
fn access_m() {
var idx = 1;
idx--;
// loads from storage
let l_s_m = s_m;
let l_s_c_c = s_m[0];
let l_s_c_v = s_m[idx];
let l_s_e_cc = s_m[0][0];
let l_s_e_cv = s_m[0][idx];
let l_s_e_vc = s_m[idx][0];
let l_s_e_vv = s_m[idx][idx];
// loads from uniform
let l_u_m = u_m;
let l_u_c_c = u_m[0];
let l_u_c_v = u_m[idx];
let l_u_e_cc = u_m[0][0];
let l_u_e_cv = u_m[0][idx];
let l_u_e_vc = u_m[idx][0];
let l_u_e_vv = u_m[idx][idx];
// stores to storage
s_m = l_u_m;
s_m[0] = l_u_c_c;
s_m[idx] = l_u_c_v;
s_m[0][0] = l_u_e_cc;
s_m[0][idx] = l_u_e_cv;
s_m[idx][0] = l_u_e_vc;
s_m[idx][idx] = l_u_e_vv;
}
struct StructWithMat {
m: Mat,
}
@group(1) @binding(0)
var<storage, read_write> s_sm: StructWithMat;
@group(1) @binding(1)
var<uniform> u_sm: StructWithMat;
fn access_sm() {
var idx = 1;
idx--;
// loads from storage
let l_s_s = s_sm;
let l_s_m = s_sm.m;
let l_s_c_c = s_sm.m[0];
let l_s_c_v = s_sm.m[idx];
let l_s_e_cc = s_sm.m[0][0];
let l_s_e_cv = s_sm.m[0][idx];
let l_s_e_vc = s_sm.m[idx][0];
let l_s_e_vv = s_sm.m[idx][idx];
// loads from uniform
let l_u_s = u_sm;
let l_u_m = u_sm.m;
let l_u_c_c = u_sm.m[0];
let l_u_c_v = u_sm.m[idx];
let l_u_e_cc = u_sm.m[0][0];
let l_u_e_cv = u_sm.m[0][idx];
let l_u_e_vc = u_sm.m[idx][0];
let l_u_e_vv = u_sm.m[idx][idx];
// stores to storage
s_sm = l_u_s;
s_sm.m = l_u_m;
s_sm.m[0] = l_u_c_c;
s_sm.m[idx] = l_u_c_v;
s_sm.m[0][0] = l_u_e_cc;
s_sm.m[0][idx] = l_u_e_cv;
s_sm.m[idx][0] = l_u_e_vc;
s_sm.m[idx][idx] = l_u_e_vv;
}
struct StructWithArrayOfStructOfMat {
a: array<StructWithMat, 4>,
}
@group(2) @binding(0)
var<storage, read_write> s_sasm: StructWithArrayOfStructOfMat;
@group(2) @binding(1)
var<uniform> u_sasm: StructWithArrayOfStructOfMat;
fn access_sasm() {
var idx = 1;
idx--;
// loads from storage
let l_s_s = s_sasm;
let l_s_a = s_sasm.a;
let l_s_m_c = s_sasm.a[0].m;
let l_s_m_v = s_sasm.a[idx].m;
let l_s_c_cc = s_sasm.a[0].m[0];
let l_s_c_cv = s_sasm.a[0].m[idx];
let l_s_c_vc = s_sasm.a[idx].m[0];
let l_s_c_vv = s_sasm.a[idx].m[idx];
let l_s_e_ccc = s_sasm.a[0].m[0][0];
let l_s_e_ccv = s_sasm.a[0].m[0][idx];
let l_s_e_cvc = s_sasm.a[0].m[idx][0];
let l_s_e_cvv = s_sasm.a[0].m[idx][idx];
let l_s_e_vcc = s_sasm.a[idx].m[0][0];
let l_s_e_vcv = s_sasm.a[idx].m[0][idx];
let l_s_e_vvc = s_sasm.a[idx].m[idx][0];
let l_s_e_vvv = s_sasm.a[idx].m[idx][idx];
// loads from uniform
let l_u_s = u_sasm;
let l_u_a = u_sasm.a;
let l_u_m_c = u_sasm.a[0].m;
let l_u_m_v = u_sasm.a[idx].m;
let l_u_c_cc = u_sasm.a[0].m[0];
let l_u_c_cv = u_sasm.a[0].m[idx];
let l_u_c_vc = u_sasm.a[idx].m[0];
let l_u_c_vv = u_sasm.a[idx].m[idx];
let l_u_e_ccc = u_sasm.a[0].m[0][0];
let l_u_e_ccv = u_sasm.a[0].m[0][idx];
let l_u_e_cvc = u_sasm.a[0].m[idx][0];
let l_u_e_cvv = u_sasm.a[0].m[idx][idx];
let l_u_e_vcc = u_sasm.a[idx].m[0][0];
let l_u_e_vcv = u_sasm.a[idx].m[0][idx];
let l_u_e_vvc = u_sasm.a[idx].m[idx][0];
let l_u_e_vvv = u_sasm.a[idx].m[idx][idx];
// stores to storage
s_sasm = l_u_s;
s_sasm.a = l_u_a;
s_sasm.a[0].m = l_u_m_c;
s_sasm.a[idx].m = l_u_m_v;
s_sasm.a[0].m[0] = l_u_c_cc;
s_sasm.a[0].m[idx] = l_u_c_cv;
s_sasm.a[idx].m[0] = l_u_c_vc;
s_sasm.a[idx].m[idx] = l_u_c_vv;
s_sasm.a[0].m[0][0] = l_u_e_ccc;
s_sasm.a[0].m[0][idx] = l_u_e_ccv;
s_sasm.a[0].m[idx][0] = l_u_e_cvc;
s_sasm.a[0].m[idx][idx] = l_u_e_cvv;
s_sasm.a[idx].m[0][0] = l_u_e_vcc;
s_sasm.a[idx].m[0][idx] = l_u_e_vcv;
s_sasm.a[idx].m[idx][0] = l_u_e_vvc;
s_sasm.a[idx].m[idx][idx] = l_u_e_vvv;
}
@compute @workgroup_size(1)
fn main() {
access_m();
access_sm();
access_sasm();
}

372
naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.hlsl Normal file
View File

@ -0,0 +1,372 @@
typedef struct { float2 _0; float2 _1; } __mat2x2;
float2 __get_col_of_mat2x2(__mat2x2 mat, uint idx) {
switch(idx) {
case 0: { return mat._0; }
case 1: { return mat._1; }
default: { return (float2)0; }
}
}
void __set_col_of_mat2x2(__mat2x2 mat, uint idx, float2 value) {
switch(idx) {
case 0: { mat._0 = value; break; }
case 1: { mat._1 = value; break; }
}
}
void __set_el_of_mat2x2(__mat2x2 mat, uint idx, uint vec_idx, float value) {
switch(idx) {
case 0: { mat._0[vec_idx] = value; break; }
case 1: { mat._1[vec_idx] = value; break; }
}
}
struct StructWithMat {
float2 m_0; float2 m_1;
};
struct StructWithArrayOfStructOfMat {
StructWithMat a[4];
};
RWByteAddressBuffer s_m : register(u0);
cbuffer u_m : register(b1) { __mat2x2 u_m; }
RWByteAddressBuffer s_sm : register(u0, space1);
cbuffer u_sm : register(b1, space1) { StructWithMat u_sm; }
RWByteAddressBuffer s_sasm : register(u0, space2);
cbuffer u_sasm : register(b1, space2) { StructWithArrayOfStructOfMat u_sasm; }
void access_m()
{
int idx = int(1);
int _e3 = idx;
idx = asint(asuint(_e3) - asuint(int(1)));
float2x2 l_s_m = float2x2(asfloat(s_m.Load2(0)), asfloat(s_m.Load2(8)));
float2 l_s_c_c = asfloat(s_m.Load2(0));
int _e11 = idx;
float2 l_s_c_v = asfloat(s_m.Load2(_e11*8));
float l_s_e_cc = asfloat(s_m.Load(0+0));
int _e20 = idx;
float l_s_e_cv = asfloat(s_m.Load(_e20*4+0));
int _e24 = idx;
float l_s_e_vc = asfloat(s_m.Load(0+_e24*8));
int _e29 = idx;
int _e31 = idx;
float l_s_e_vv = asfloat(s_m.Load(_e31*4+_e29*8));
float2x2 l_u_m = ((float2x2)u_m);
float2 l_u_c_c = u_m._0;
int _e40 = idx;
float2 l_u_c_v = __get_col_of_mat2x2(u_m, _e40);
float l_u_e_cc = u_m._0.x;
int _e49 = idx;
float l_u_e_cv = u_m._0[_e49];
int _e53 = idx;
float l_u_e_vc = __get_col_of_mat2x2(u_m, _e53).x;
int _e58 = idx;
int _e60 = idx;
float l_u_e_vv = __get_col_of_mat2x2(u_m, _e58)[_e60];
{
float2x2 _value2 = l_u_m;
s_m.Store2(0, asuint(_value2[0]));
s_m.Store2(8, asuint(_value2[1]));
}
s_m.Store2(0, asuint(l_u_c_c));
int _e67 = idx;
s_m.Store2(_e67*8, asuint(l_u_c_v));
s_m.Store(0+0, asuint(l_u_e_cc));
int _e74 = idx;
s_m.Store(_e74*4+0, asuint(l_u_e_cv));
int _e77 = idx;
s_m.Store(0+_e77*8, asuint(l_u_e_vc));
int _e81 = idx;
int _e83 = idx;
s_m.Store(_e83*4+_e81*8, asuint(l_u_e_vv));
return;
}
StructWithMat ConstructStructWithMat(float2x2 arg0) {
StructWithMat ret = (StructWithMat)0;
ret.m_0 = arg0[0];
ret.m_1 = arg0[1];
return ret;
}
float2x2 GetMatmOnStructWithMat(StructWithMat obj) {
return float2x2(obj.m_0, obj.m_1);
}
void SetMatmOnStructWithMat(StructWithMat obj, float2x2 mat) {
obj.m_0 = mat[0];
obj.m_1 = mat[1];
}
void SetMatVecmOnStructWithMat(StructWithMat obj, float2 vec, uint mat_idx) {
switch(mat_idx) {
case 0: { obj.m_0 = vec; break; }
case 1: { obj.m_1 = vec; break; }
}
}
void SetMatScalarmOnStructWithMat(StructWithMat obj, float scalar, uint mat_idx, uint vec_idx) {
switch(mat_idx) {
case 0: { obj.m_0[vec_idx] = scalar; break; }
case 1: { obj.m_1[vec_idx] = scalar; break; }
}
}
void access_sm()
{
int idx_1 = int(1);
int _e3 = idx_1;
idx_1 = asint(asuint(_e3) - asuint(int(1)));
StructWithMat l_s_s = ConstructStructWithMat(float2x2(asfloat(s_sm.Load2(0+0)), asfloat(s_sm.Load2(0+8))));
float2x2 l_s_m_1 = float2x2(asfloat(s_sm.Load2(0+0)), asfloat(s_sm.Load2(0+8)));
float2 l_s_c_c_1 = asfloat(s_sm.Load2(0+0));
int _e16 = idx_1;
float2 l_s_c_v_1 = asfloat(s_sm.Load2(_e16*8+0));
float l_s_e_cc_1 = asfloat(s_sm.Load(0+0+0));
int _e27 = idx_1;
float l_s_e_cv_1 = asfloat(s_sm.Load(_e27*4+0+0));
int _e32 = idx_1;
float l_s_e_vc_1 = asfloat(s_sm.Load(0+_e32*8+0));
int _e38 = idx_1;
int _e40 = idx_1;
float l_s_e_vv_1 = asfloat(s_sm.Load(_e40*4+_e38*8+0));
StructWithMat l_u_s = u_sm;
float2x2 l_u_m_1 = GetMatmOnStructWithMat(u_sm);
float2 l_u_c_c_1 = GetMatmOnStructWithMat(u_sm)[0];
int _e54 = idx_1;
float2 l_u_c_v_1 = GetMatmOnStructWithMat(u_sm)[_e54];
float l_u_e_cc_1 = GetMatmOnStructWithMat(u_sm)[0].x;
int _e65 = idx_1;
float l_u_e_cv_1 = GetMatmOnStructWithMat(u_sm)[0][_e65];
int _e70 = idx_1;
float l_u_e_vc_1 = GetMatmOnStructWithMat(u_sm)[_e70].x;
int _e76 = idx_1;
int _e78 = idx_1;
float l_u_e_vv_1 = GetMatmOnStructWithMat(u_sm)[_e76][_e78];
{
StructWithMat _value2 = l_u_s;
{
s_sm.Store2(0+0, asuint(_value2.m_0));
s_sm.Store2(0+8, asuint(_value2.m_1));
}
}
{
float2x2 _value2 = l_u_m_1;
s_sm.Store2(0+0, asuint(_value2[0]));
s_sm.Store2(0+8, asuint(_value2[1]));
}
s_sm.Store2(0+0, asuint(l_u_c_c_1));
int _e89 = idx_1;
s_sm.Store2(_e89*8+0, asuint(l_u_c_v_1));
s_sm.Store(0+0+0, asuint(l_u_e_cc_1));
int _e98 = idx_1;
s_sm.Store(_e98*4+0+0, asuint(l_u_e_cv_1));
int _e102 = idx_1;
s_sm.Store(0+_e102*8+0, asuint(l_u_e_vc_1));
int _e107 = idx_1;
int _e109 = idx_1;
s_sm.Store(_e109*4+_e107*8+0, asuint(l_u_e_vv_1));
return;
}
typedef StructWithMat ret_Constructarray4_StructWithMat_[4];
ret_Constructarray4_StructWithMat_ Constructarray4_StructWithMat_(StructWithMat arg0, StructWithMat arg1, StructWithMat arg2, StructWithMat arg3) {
StructWithMat ret[4] = { arg0, arg1, arg2, arg3 };
return ret;
}
StructWithArrayOfStructOfMat ConstructStructWithArrayOfStructOfMat(StructWithMat arg0[4]) {
StructWithArrayOfStructOfMat ret = (StructWithArrayOfStructOfMat)0;
ret.a = arg0;
return ret;
}
void access_sasm()
{
int idx_2 = int(1);
int _e3 = idx_2;
idx_2 = asint(asuint(_e3) - asuint(int(1)));
StructWithArrayOfStructOfMat l_s_s_1 = ConstructStructWithArrayOfStructOfMat(Constructarray4_StructWithMat_(ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+0+0+0)), asfloat(s_sasm.Load2(0+0+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+16+0+0)), asfloat(s_sasm.Load2(0+16+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+32+0+0)), asfloat(s_sasm.Load2(0+32+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+48+0+0)), asfloat(s_sasm.Load2(0+48+0+8))))));
StructWithMat l_s_a[4] = Constructarray4_StructWithMat_(ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+0+0+0)), asfloat(s_sasm.Load2(0+0+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+16+0+0)), asfloat(s_sasm.Load2(0+16+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+32+0+0)), asfloat(s_sasm.Load2(0+32+0+8)))), ConstructStructWithMat(float2x2(asfloat(s_sasm.Load2(0+48+0+0)), asfloat(s_sasm.Load2(0+48+0+8)))));
float2x2 l_s_m_c = float2x2(asfloat(s_sasm.Load2(0+0+0+0)), asfloat(s_sasm.Load2(0+0+0+8)));
int _e17 = idx_2;
float2x2 l_s_m_v = float2x2(asfloat(s_sasm.Load2(0+_e17*16+0+0)), asfloat(s_sasm.Load2(0+_e17*16+0+8)));
float2 l_s_c_cc = asfloat(s_sasm.Load2(0+0+0+0));
int _e31 = idx_2;
float2 l_s_c_cv = asfloat(s_sasm.Load2(_e31*8+0+0+0));
int _e36 = idx_2;
float2 l_s_c_vc = asfloat(s_sasm.Load2(0+0+_e36*16+0));
int _e43 = idx_2;
int _e46 = idx_2;
float2 l_s_c_vv = asfloat(s_sasm.Load2(_e46*8+0+_e43*16+0));
float l_s_e_ccc = asfloat(s_sasm.Load(0+0+0+0+0));
int _e61 = idx_2;
float l_s_e_ccv = asfloat(s_sasm.Load(_e61*4+0+0+0+0));
int _e68 = idx_2;
float l_s_e_cvc = asfloat(s_sasm.Load(0+_e68*8+0+0+0));
int _e76 = idx_2;
int _e78 = idx_2;
float l_s_e_cvv = asfloat(s_sasm.Load(_e78*4+_e76*8+0+0+0));
int _e83 = idx_2;
float l_s_e_vcc = asfloat(s_sasm.Load(0+0+0+_e83*16+0));
int _e91 = idx_2;
int _e95 = idx_2;
float l_s_e_vcv = asfloat(s_sasm.Load(_e95*4+0+0+_e91*16+0));
int _e100 = idx_2;
int _e103 = idx_2;
float l_s_e_vvc = asfloat(s_sasm.Load(0+_e103*8+0+_e100*16+0));
int _e109 = idx_2;
int _e112 = idx_2;
int _e114 = idx_2;
float l_s_e_vvv = asfloat(s_sasm.Load(_e114*4+_e112*8+0+_e109*16+0));
StructWithArrayOfStructOfMat l_u_s_1 = u_sasm;
StructWithMat l_u_a[4] = u_sasm.a;
float2x2 l_u_m_c = GetMatmOnStructWithMat(u_sasm.a[0]);
int _e129 = idx_2;
float2x2 l_u_m_v = GetMatmOnStructWithMat(u_sasm.a[_e129]);
float2 l_u_c_cc = GetMatmOnStructWithMat(u_sasm.a[0])[0];
int _e143 = idx_2;
float2 l_u_c_cv = GetMatmOnStructWithMat(u_sasm.a[0])[_e143];
int _e148 = idx_2;
float2 l_u_c_vc = GetMatmOnStructWithMat(u_sasm.a[_e148])[0];
int _e155 = idx_2;
int _e158 = idx_2;
float2 l_u_c_vv = GetMatmOnStructWithMat(u_sasm.a[_e155])[_e158];
float l_u_e_ccc = GetMatmOnStructWithMat(u_sasm.a[0])[0].x;
int _e173 = idx_2;
float l_u_e_ccv = GetMatmOnStructWithMat(u_sasm.a[0])[0][_e173];
int _e180 = idx_2;
float l_u_e_cvc = GetMatmOnStructWithMat(u_sasm.a[0])[_e180].x;
int _e188 = idx_2;
int _e190 = idx_2;
float l_u_e_cvv = GetMatmOnStructWithMat(u_sasm.a[0])[_e188][_e190];
int _e195 = idx_2;
float l_u_e_vcc = GetMatmOnStructWithMat(u_sasm.a[_e195])[0].x;
int _e203 = idx_2;
int _e207 = idx_2;
float l_u_e_vcv = GetMatmOnStructWithMat(u_sasm.a[_e203])[0][_e207];
int _e212 = idx_2;
int _e215 = idx_2;
float l_u_e_vvc = GetMatmOnStructWithMat(u_sasm.a[_e212])[_e215].x;
int _e221 = idx_2;
int _e224 = idx_2;
int _e226 = idx_2;
float l_u_e_vvv = GetMatmOnStructWithMat(u_sasm.a[_e221])[_e224][_e226];
{
StructWithArrayOfStructOfMat _value2 = l_u_s_1;
{
StructWithMat _value3[4] = _value2.a;
{
StructWithMat _value4 = _value3[0];
{
s_sasm.Store2(0+0+0+0, asuint(_value4.m_0));
s_sasm.Store2(0+0+0+8, asuint(_value4.m_1));
}
}
{
StructWithMat _value4 = _value3[1];
{
s_sasm.Store2(0+16+0+0, asuint(_value4.m_0));
s_sasm.Store2(0+16+0+8, asuint(_value4.m_1));
}
}
{
StructWithMat _value4 = _value3[2];
{
s_sasm.Store2(0+32+0+0, asuint(_value4.m_0));
s_sasm.Store2(0+32+0+8, asuint(_value4.m_1));
}
}
{
StructWithMat _value4 = _value3[3];
{
s_sasm.Store2(0+48+0+0, asuint(_value4.m_0));
s_sasm.Store2(0+48+0+8, asuint(_value4.m_1));
}
}
}
}
{
StructWithMat _value2[4] = l_u_a;
{
StructWithMat _value3 = _value2[0];
{
s_sasm.Store2(0+0+0+0, asuint(_value3.m_0));
s_sasm.Store2(0+0+0+8, asuint(_value3.m_1));
}
}
{
StructWithMat _value3 = _value2[1];
{
s_sasm.Store2(0+16+0+0, asuint(_value3.m_0));
s_sasm.Store2(0+16+0+8, asuint(_value3.m_1));
}
}
{
StructWithMat _value3 = _value2[2];
{
s_sasm.Store2(0+32+0+0, asuint(_value3.m_0));
s_sasm.Store2(0+32+0+8, asuint(_value3.m_1));
}
}
{
StructWithMat _value3 = _value2[3];
{
s_sasm.Store2(0+48+0+0, asuint(_value3.m_0));
s_sasm.Store2(0+48+0+8, asuint(_value3.m_1));
}
}
}
{
float2x2 _value2 = l_u_m_c;
s_sasm.Store2(0+0+0+0, asuint(_value2[0]));
s_sasm.Store2(0+0+0+8, asuint(_value2[1]));
}
int _e238 = idx_2;
{
float2x2 _value2 = l_u_m_v;
s_sasm.Store2(0+_e238*16+0+0, asuint(_value2[0]));
s_sasm.Store2(0+_e238*16+0+8, asuint(_value2[1]));
}
s_sasm.Store2(0+0+0+0, asuint(l_u_c_cc));
int _e250 = idx_2;
s_sasm.Store2(_e250*8+0+0+0, asuint(l_u_c_cv));
int _e254 = idx_2;
s_sasm.Store2(0+0+_e254*16+0, asuint(l_u_c_vc));
int _e260 = idx_2;
int _e263 = idx_2;
s_sasm.Store2(_e263*8+0+_e260*16+0, asuint(l_u_c_vv));
s_sasm.Store(0+0+0+0+0, asuint(l_u_e_ccc));
int _e276 = idx_2;
s_sasm.Store(_e276*4+0+0+0+0, asuint(l_u_e_ccv));
int _e282 = idx_2;
s_sasm.Store(0+_e282*8+0+0+0, asuint(l_u_e_cvc));
int _e289 = idx_2;
int _e291 = idx_2;
s_sasm.Store(_e291*4+_e289*8+0+0+0, asuint(l_u_e_cvv));
int _e295 = idx_2;
s_sasm.Store(0+0+0+_e295*16+0, asuint(l_u_e_vcc));
int _e302 = idx_2;
int _e306 = idx_2;
s_sasm.Store(_e306*4+0+0+_e302*16+0, asuint(l_u_e_vcv));
int _e310 = idx_2;
int _e313 = idx_2;
s_sasm.Store(0+_e313*8+0+_e310*16+0, asuint(l_u_e_vvc));
int _e318 = idx_2;
int _e321 = idx_2;
int _e323 = idx_2;
s_sasm.Store(_e323*4+_e321*8+0+_e318*16+0, asuint(l_u_e_vvv));
return;
}
[numthreads(1, 1, 1)]
void main()
{
access_m();
access_sm();
access_sasm();
return;
}

12
naga/tests/out/hlsl/wgsl-hlsl_mat_cx2.ron Normal file
View File

@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"main",
target_profile:"cs_5_1",
),
],
)

350
naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.hlsl Normal file
View File

@ -0,0 +1,350 @@
struct StructWithMat {
row_major float3x3 m;
int _end_pad_0;
};
struct StructWithArrayOfStructOfMat {
StructWithMat a[4];
};
RWByteAddressBuffer s_m : register(u0);
cbuffer u_m : register(b1) { row_major float3x3 u_m; }
RWByteAddressBuffer s_sm : register(u0, space1);
cbuffer u_sm : register(b1, space1) { StructWithMat u_sm; }
RWByteAddressBuffer s_sasm : register(u0, space2);
cbuffer u_sasm : register(b1, space2) { StructWithArrayOfStructOfMat u_sasm; }
void access_m()
{
int idx = int(1);
int _e3 = idx;
idx = asint(asuint(_e3) - asuint(int(1)));
float3x3 l_s_m = float3x3(asfloat(s_m.Load3(0)), asfloat(s_m.Load3(16)), asfloat(s_m.Load3(32)));
float3 l_s_c_c = asfloat(s_m.Load3(0));
int _e11 = idx;
float3 l_s_c_v = asfloat(s_m.Load3(_e11*16));
float l_s_e_cc = asfloat(s_m.Load(0+0));
int _e20 = idx;
float l_s_e_cv = asfloat(s_m.Load(_e20*4+0));
int _e24 = idx;
float l_s_e_vc = asfloat(s_m.Load(0+_e24*16));
int _e29 = idx;
int _e31 = idx;
float l_s_e_vv = asfloat(s_m.Load(_e31*4+_e29*16));
float3x3 l_u_m = u_m;
float3 l_u_c_c = u_m[0];
int _e40 = idx;
float3 l_u_c_v = u_m[_e40];
float l_u_e_cc = u_m[0].x;
int _e49 = idx;
float l_u_e_cv = u_m[0][_e49];
int _e53 = idx;
float l_u_e_vc = u_m[_e53].x;
int _e58 = idx;
int _e60 = idx;
float l_u_e_vv = u_m[_e58][_e60];
{
float3x3 _value2 = l_u_m;
s_m.Store3(0, asuint(_value2[0]));
s_m.Store3(16, asuint(_value2[1]));
s_m.Store3(32, asuint(_value2[2]));
}
s_m.Store3(0, asuint(l_u_c_c));
int _e67 = idx;
s_m.Store3(_e67*16, asuint(l_u_c_v));
s_m.Store(0+0, asuint(l_u_e_cc));
int _e74 = idx;
s_m.Store(_e74*4+0, asuint(l_u_e_cv));
int _e77 = idx;
s_m.Store(0+_e77*16, asuint(l_u_e_vc));
int _e81 = idx;
int _e83 = idx;
s_m.Store(_e83*4+_e81*16, asuint(l_u_e_vv));
return;
}
StructWithMat ConstructStructWithMat(float3x3 arg0) {
StructWithMat ret = (StructWithMat)0;
ret.m = arg0;
return ret;
}
void access_sm()
{
int idx_1 = int(1);
int _e3 = idx_1;
idx_1 = asint(asuint(_e3) - asuint(int(1)));
StructWithMat l_s_s = ConstructStructWithMat(float3x3(asfloat(s_sm.Load3(0+0)), asfloat(s_sm.Load3(0+16)), asfloat(s_sm.Load3(0+32))));
float3x3 l_s_m_1 = float3x3(asfloat(s_sm.Load3(0+0)), asfloat(s_sm.Load3(0+16)), asfloat(s_sm.Load3(0+32)));
float3 l_s_c_c_1 = asfloat(s_sm.Load3(0+0));
int _e16 = idx_1;
float3 l_s_c_v_1 = asfloat(s_sm.Load3(_e16*16+0));
float l_s_e_cc_1 = asfloat(s_sm.Load(0+0+0));
int _e27 = idx_1;
float l_s_e_cv_1 = asfloat(s_sm.Load(_e27*4+0+0));
int _e32 = idx_1;
float l_s_e_vc_1 = asfloat(s_sm.Load(0+_e32*16+0));
int _e38 = idx_1;
int _e40 = idx_1;
float l_s_e_vv_1 = asfloat(s_sm.Load(_e40*4+_e38*16+0));
StructWithMat l_u_s = u_sm;
float3x3 l_u_m_1 = u_sm.m;
float3 l_u_c_c_1 = u_sm.m[0];
int _e54 = idx_1;
float3 l_u_c_v_1 = u_sm.m[_e54];
float l_u_e_cc_1 = u_sm.m[0].x;
int _e65 = idx_1;
float l_u_e_cv_1 = u_sm.m[0][_e65];
int _e70 = idx_1;
float l_u_e_vc_1 = u_sm.m[_e70].x;
int _e76 = idx_1;
int _e78 = idx_1;
float l_u_e_vv_1 = u_sm.m[_e76][_e78];
{
StructWithMat _value2 = l_u_s;
{
float3x3 _value3 = _value2.m;
s_sm.Store3(0+0, asuint(_value3[0]));
s_sm.Store3(0+16, asuint(_value3[1]));
s_sm.Store3(0+32, asuint(_value3[2]));
}
}
{
float3x3 _value2 = l_u_m_1;
s_sm.Store3(0+0, asuint(_value2[0]));
s_sm.Store3(0+16, asuint(_value2[1]));
s_sm.Store3(0+32, asuint(_value2[2]));
}
s_sm.Store3(0+0, asuint(l_u_c_c_1));
int _e89 = idx_1;
s_sm.Store3(_e89*16+0, asuint(l_u_c_v_1));
s_sm.Store(0+0+0, asuint(l_u_e_cc_1));
int _e98 = idx_1;
s_sm.Store(_e98*4+0+0, asuint(l_u_e_cv_1));
int _e102 = idx_1;
s_sm.Store(0+_e102*16+0, asuint(l_u_e_vc_1));
int _e107 = idx_1;
int _e109 = idx_1;
s_sm.Store(_e109*4+_e107*16+0, asuint(l_u_e_vv_1));
return;
}
typedef StructWithMat ret_Constructarray4_StructWithMat_[4];
ret_Constructarray4_StructWithMat_ Constructarray4_StructWithMat_(StructWithMat arg0, StructWithMat arg1, StructWithMat arg2, StructWithMat arg3) {
StructWithMat ret[4] = { arg0, arg1, arg2, arg3 };
return ret;
}
StructWithArrayOfStructOfMat ConstructStructWithArrayOfStructOfMat(StructWithMat arg0[4]) {
StructWithArrayOfStructOfMat ret = (StructWithArrayOfStructOfMat)0;
ret.a = arg0;
return ret;
}
void access_sasm()
{
int idx_2 = int(1);
int _e3 = idx_2;
idx_2 = asint(asuint(_e3) - asuint(int(1)));
StructWithArrayOfStructOfMat l_s_s_1 = ConstructStructWithArrayOfStructOfMat(Constructarray4_StructWithMat_(ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+0+0+0)), asfloat(s_sasm.Load3(0+0+0+16)), asfloat(s_sasm.Load3(0+0+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+48+0+0)), asfloat(s_sasm.Load3(0+48+0+16)), asfloat(s_sasm.Load3(0+48+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+96+0+0)), asfloat(s_sasm.Load3(0+96+0+16)), asfloat(s_sasm.Load3(0+96+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+144+0+0)), asfloat(s_sasm.Load3(0+144+0+16)), asfloat(s_sasm.Load3(0+144+0+32))))));
StructWithMat l_s_a[4] = Constructarray4_StructWithMat_(ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+0+0+0)), asfloat(s_sasm.Load3(0+0+0+16)), asfloat(s_sasm.Load3(0+0+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+48+0+0)), asfloat(s_sasm.Load3(0+48+0+16)), asfloat(s_sasm.Load3(0+48+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+96+0+0)), asfloat(s_sasm.Load3(0+96+0+16)), asfloat(s_sasm.Load3(0+96+0+32)))), ConstructStructWithMat(float3x3(asfloat(s_sasm.Load3(0+144+0+0)), asfloat(s_sasm.Load3(0+144+0+16)), asfloat(s_sasm.Load3(0+144+0+32)))));
float3x3 l_s_m_c = float3x3(asfloat(s_sasm.Load3(0+0+0+0)), asfloat(s_sasm.Load3(0+0+0+16)), asfloat(s_sasm.Load3(0+0+0+32)));
int _e17 = idx_2;
float3x3 l_s_m_v = float3x3(asfloat(s_sasm.Load3(0+_e17*48+0+0)), asfloat(s_sasm.Load3(0+_e17*48+0+16)), asfloat(s_sasm.Load3(0+_e17*48+0+32)));
float3 l_s_c_cc = asfloat(s_sasm.Load3(0+0+0+0));
int _e31 = idx_2;
float3 l_s_c_cv = asfloat(s_sasm.Load3(_e31*16+0+0+0));
int _e36 = idx_2;
float3 l_s_c_vc = asfloat(s_sasm.Load3(0+0+_e36*48+0));
int _e43 = idx_2;
int _e46 = idx_2;
float3 l_s_c_vv = asfloat(s_sasm.Load3(_e46*16+0+_e43*48+0));
float l_s_e_ccc = asfloat(s_sasm.Load(0+0+0+0+0));
int _e61 = idx_2;
float l_s_e_ccv = asfloat(s_sasm.Load(_e61*4+0+0+0+0));
int _e68 = idx_2;
float l_s_e_cvc = asfloat(s_sasm.Load(0+_e68*16+0+0+0));
int _e76 = idx_2;
int _e78 = idx_2;
float l_s_e_cvv = asfloat(s_sasm.Load(_e78*4+_e76*16+0+0+0));
int _e83 = idx_2;
float l_s_e_vcc = asfloat(s_sasm.Load(0+0+0+_e83*48+0));
int _e91 = idx_2;
int _e95 = idx_2;
float l_s_e_vcv = asfloat(s_sasm.Load(_e95*4+0+0+_e91*48+0));
int _e100 = idx_2;
int _e103 = idx_2;
float l_s_e_vvc = asfloat(s_sasm.Load(0+_e103*16+0+_e100*48+0));
int _e109 = idx_2;
int _e112 = idx_2;
int _e114 = idx_2;
float l_s_e_vvv = asfloat(s_sasm.Load(_e114*4+_e112*16+0+_e109*48+0));
StructWithArrayOfStructOfMat l_u_s_1 = u_sasm;
StructWithMat l_u_a[4] = u_sasm.a;
float3x3 l_u_m_c = u_sasm.a[0].m;
int _e129 = idx_2;
float3x3 l_u_m_v = u_sasm.a[_e129].m;
float3 l_u_c_cc = u_sasm.a[0].m[0];
int _e143 = idx_2;
float3 l_u_c_cv = u_sasm.a[0].m[_e143];
int _e148 = idx_2;
float3 l_u_c_vc = u_sasm.a[_e148].m[0];
int _e155 = idx_2;
int _e158 = idx_2;
float3 l_u_c_vv = u_sasm.a[_e155].m[_e158];
float l_u_e_ccc = u_sasm.a[0].m[0].x;
int _e173 = idx_2;
float l_u_e_ccv = u_sasm.a[0].m[0][_e173];
int _e180 = idx_2;
float l_u_e_cvc = u_sasm.a[0].m[_e180].x;
int _e188 = idx_2;
int _e190 = idx_2;
float l_u_e_cvv = u_sasm.a[0].m[_e188][_e190];
int _e195 = idx_2;
float l_u_e_vcc = u_sasm.a[_e195].m[0].x;
int _e203 = idx_2;
int _e207 = idx_2;
float l_u_e_vcv = u_sasm.a[_e203].m[0][_e207];
int _e212 = idx_2;
int _e215 = idx_2;
float l_u_e_vvc = u_sasm.a[_e212].m[_e215].x;
int _e221 = idx_2;
int _e224 = idx_2;
int _e226 = idx_2;
float l_u_e_vvv = u_sasm.a[_e221].m[_e224][_e226];
{
StructWithArrayOfStructOfMat _value2 = l_u_s_1;
{
StructWithMat _value3[4] = _value2.a;
{
StructWithMat _value4 = _value3[0];
{
float3x3 _value5 = _value4.m;
s_sasm.Store3(0+0+0+0, asuint(_value5[0]));
s_sasm.Store3(0+0+0+16, asuint(_value5[1]));
s_sasm.Store3(0+0+0+32, asuint(_value5[2]));
}
}
{
StructWithMat _value4 = _value3[1];
{
float3x3 _value5 = _value4.m;
s_sasm.Store3(0+48+0+0, asuint(_value5[0]));
s_sasm.Store3(0+48+0+16, asuint(_value5[1]));
s_sasm.Store3(0+48+0+32, asuint(_value5[2]));
}
}
{
StructWithMat _value4 = _value3[2];
{
float3x3 _value5 = _value4.m;
s_sasm.Store3(0+96+0+0, asuint(_value5[0]));
s_sasm.Store3(0+96+0+16, asuint(_value5[1]));
s_sasm.Store3(0+96+0+32, asuint(_value5[2]));
}
}
{
StructWithMat _value4 = _value3[3];
{
float3x3 _value5 = _value4.m;
s_sasm.Store3(0+144+0+0, asuint(_value5[0]));
s_sasm.Store3(0+144+0+16, asuint(_value5[1]));
s_sasm.Store3(0+144+0+32, asuint(_value5[2]));
}
}
}
}
{
StructWithMat _value2[4] = l_u_a;
{
StructWithMat _value3 = _value2[0];
{
float3x3 _value4 = _value3.m;
s_sasm.Store3(0+0+0+0, asuint(_value4[0]));
s_sasm.Store3(0+0+0+16, asuint(_value4[1]));
s_sasm.Store3(0+0+0+32, asuint(_value4[2]));
}
}
{
StructWithMat _value3 = _value2[1];
{
float3x3 _value4 = _value3.m;
s_sasm.Store3(0+48+0+0, asuint(_value4[0]));
s_sasm.Store3(0+48+0+16, asuint(_value4[1]));
s_sasm.Store3(0+48+0+32, asuint(_value4[2]));
}
}
{
StructWithMat _value3 = _value2[2];
{
float3x3 _value4 = _value3.m;
s_sasm.Store3(0+96+0+0, asuint(_value4[0]));
s_sasm.Store3(0+96+0+16, asuint(_value4[1]));
s_sasm.Store3(0+96+0+32, asuint(_value4[2]));
}
}
{
StructWithMat _value3 = _value2[3];
{
float3x3 _value4 = _value3.m;
s_sasm.Store3(0+144+0+0, asuint(_value4[0]));
s_sasm.Store3(0+144+0+16, asuint(_value4[1]));
s_sasm.Store3(0+144+0+32, asuint(_value4[2]));
}
}
}
{
float3x3 _value2 = l_u_m_c;
s_sasm.Store3(0+0+0+0, asuint(_value2[0]));
s_sasm.Store3(0+0+0+16, asuint(_value2[1]));
s_sasm.Store3(0+0+0+32, asuint(_value2[2]));
}
int _e238 = idx_2;
{
float3x3 _value2 = l_u_m_v;
s_sasm.Store3(0+_e238*48+0+0, asuint(_value2[0]));
s_sasm.Store3(0+_e238*48+0+16, asuint(_value2[1]));
s_sasm.Store3(0+_e238*48+0+32, asuint(_value2[2]));
}
s_sasm.Store3(0+0+0+0, asuint(l_u_c_cc));
int _e250 = idx_2;
s_sasm.Store3(_e250*16+0+0+0, asuint(l_u_c_cv));
int _e254 = idx_2;
s_sasm.Store3(0+0+_e254*48+0, asuint(l_u_c_vc));
int _e260 = idx_2;
int _e263 = idx_2;
s_sasm.Store3(_e263*16+0+_e260*48+0, asuint(l_u_c_vv));
s_sasm.Store(0+0+0+0+0, asuint(l_u_e_ccc));
int _e276 = idx_2;
s_sasm.Store(_e276*4+0+0+0+0, asuint(l_u_e_ccv));
int _e282 = idx_2;
s_sasm.Store(0+_e282*16+0+0+0, asuint(l_u_e_cvc));
int _e289 = idx_2;
int _e291 = idx_2;
s_sasm.Store(_e291*4+_e289*16+0+0+0, asuint(l_u_e_cvv));
int _e295 = idx_2;
s_sasm.Store(0+0+0+_e295*48+0, asuint(l_u_e_vcc));
int _e302 = idx_2;
int _e306 = idx_2;
s_sasm.Store(_e306*4+0+0+_e302*48+0, asuint(l_u_e_vcv));
int _e310 = idx_2;
int _e313 = idx_2;
s_sasm.Store(0+_e313*16+0+_e310*48+0, asuint(l_u_e_vvc));
int _e318 = idx_2;
int _e321 = idx_2;
int _e323 = idx_2;
s_sasm.Store(_e323*4+_e321*16+0+_e318*48+0, asuint(l_u_e_vvv));
return;
}
[numthreads(1, 1, 1)]
void main()
{
access_m();
access_sm();
access_sasm();
return;
}

12
naga/tests/out/hlsl/wgsl-hlsl_mat_cx3.ron Normal file
View File

@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"main",
target_profile:"cs_5_1",
),
],
)

391
naga/tests/out/spv/wgsl-access.spvasm
View File

@ -49,6 +49,9 @@ var<uniform> baz: Baz;
var<storage,read_write> qux: vec2<i32>;
fn test_matrix_within_struct_accesses() {
// Test HLSL accesses to Cx2 matrices. There are additional tests
// in `hlsl_mat_cx2.wgsl`.
var idx = 1;
idx--;
@ -576,91 +579,91 @@ OpDecorate %388 Location 0
%69 = OpAccessChain %68 %56 %48
OpBranch %98
%98 = OpLabel
OpLine %3 40 5
OpLine %3 43 5
%99 = OpLoad %6 %94
%100 = OpISub %6 %99 %70
OpLine %3 40 5
OpLine %3 43 5
OpStore %94 %100
OpLine %3 43 14
OpLine %3 46 14
%102 = OpAccessChain %101 %69 %48
%103 = OpLoad %22 %102
OpLine %3 44 14
OpLine %3 44 14
OpLine %3 47 14
OpLine %3 47 14
%105 = OpAccessChain %104 %69 %48 %48
%106 = OpLoad %13 %105
OpLine %3 45 14
OpLine %3 48 14
%107 = OpLoad %6 %94
%108 = OpAccessChain %104 %69 %48 %107
%109 = OpLoad %13 %108
OpLine %3 46 14
OpLine %3 46 14
OpLine %3 46 14
OpLine %3 49 14
OpLine %3 49 14
OpLine %3 49 14
%111 = OpAccessChain %110 %69 %48 %48 %44
%112 = OpLoad %9 %111
OpLine %3 47 14
OpLine %3 47 14
OpLine %3 50 14
OpLine %3 50 14
%113 = OpLoad %6 %94
%114 = OpAccessChain %110 %69 %48 %48 %113
%115 = OpLoad %9 %114
OpLine %3 48 14
OpLine %3 51 14
%116 = OpLoad %6 %94
OpLine %3 48 14
OpLine %3 51 14
%117 = OpAccessChain %110 %69 %48 %116 %44
%118 = OpLoad %9 %117
OpLine %3 49 14
OpLine %3 52 14
%119 = OpLoad %6 %94
%120 = OpLoad %6 %94
%121 = OpAccessChain %110 %69 %48 %119 %120
%122 = OpLoad %9 %121
OpLine %3 51 29
OpLine %3 51 45
OpLine %3 51 13
OpLine %3 53 5
OpLine %3 54 29
OpLine %3 54 45
OpLine %3 54 13
OpLine %3 56 5
%123 = OpLoad %6 %94
%124 = OpIAdd %6 %123 %70
OpLine %3 53 5
OpLine %3 56 5
OpStore %94 %124
OpLine %3 56 5
OpLine %3 56 23
OpLine %3 56 39
OpLine %3 56 11
OpLine %3 56 5
OpLine %3 59 5
OpLine %3 59 23
OpLine %3 59 39
OpLine %3 59 11
OpLine %3 59 5
%126 = OpAccessChain %125 %96 %48
OpStore %126 %85
OpLine %3 57 5
OpLine %3 57 5
OpLine %3 57 14
OpLine %3 57 5
OpLine %3 60 5
OpLine %3 60 5
OpLine %3 60 14
OpLine %3 60 5
%128 = OpAccessChain %127 %96 %48 %48
OpStore %128 %87
OpLine %3 58 5
OpLine %3 61 5
%129 = OpLoad %6 %94
OpLine %3 58 16
OpLine %3 58 5
OpLine %3 61 16
OpLine %3 61 5
%130 = OpAccessChain %127 %96 %48 %129
OpStore %130 %89
OpLine %3 59 5
OpLine %3 59 5
OpLine %3 59 5
OpLine %3 59 5
OpLine %3 62 5
OpLine %3 62 5
OpLine %3 62 5
OpLine %3 62 5
%131 = OpAccessChain %28 %96 %48 %48 %44
OpStore %131 %90
OpLine %3 60 5
OpLine %3 60 5
OpLine %3 63 5
OpLine %3 63 5
%132 = OpLoad %6 %94
OpLine %3 60 5
OpLine %3 63 5
%133 = OpAccessChain %28 %96 %48 %48 %132
OpStore %133 %91
OpLine %3 61 5
OpLine %3 64 5
%134 = OpLoad %6 %94
OpLine %3 61 5
OpLine %3 61 5
OpLine %3 64 5
OpLine %3 64 5
%135 = OpAccessChain %28 %96 %48 %134 %44
OpStore %135 %92
OpLine %3 62 5
OpLine %3 65 5
%136 = OpLoad %6 %94
%137 = OpLoad %6 %94
OpLine %3 62 5
OpLine %3 65 5
%138 = OpAccessChain %28 %96 %48 %136 %137
OpStore %138 %93
OpReturn
@ -672,111 +675,111 @@ OpFunctionEnd
%142 = OpAccessChain %141 %62 %48
OpBranch %153
%153 = OpLabel
OpLine %3 75 5
OpLine %3 78 5
%154 = OpLoad %6 %150
%155 = OpISub %6 %154 %70
OpLine %3 75 5
OpLine %3 78 5
OpStore %150 %155
OpLine %3 78 14
OpLine %3 81 14
%157 = OpAccessChain %156 %142 %48
%158 = OpLoad %26 %157
OpLine %3 79 14
OpLine %3 79 14
OpLine %3 82 14
OpLine %3 82 14
%160 = OpAccessChain %159 %142 %48 %48
%161 = OpLoad %25 %160
OpLine %3 80 14
OpLine %3 80 14
OpLine %3 80 14
OpLine %3 83 14
OpLine %3 83 14
OpLine %3 83 14
%162 = OpAccessChain %104 %142 %48 %48 %48
%163 = OpLoad %13 %162
OpLine %3 81 14
OpLine %3 81 14
OpLine %3 84 14
OpLine %3 84 14
%164 = OpLoad %6 %150
%165 = OpAccessChain %104 %142 %48 %48 %164
%166 = OpLoad %13 %165
OpLine %3 82 14
OpLine %3 82 14
OpLine %3 82 14
OpLine %3 82 14
OpLine %3 85 14
OpLine %3 85 14
OpLine %3 85 14
OpLine %3 85 14
%167 = OpAccessChain %110 %142 %48 %48 %48 %44
%168 = OpLoad %9 %167
OpLine %3 83 14
OpLine %3 83 14
OpLine %3 83 14
OpLine %3 86 14
OpLine %3 86 14
OpLine %3 86 14
%169 = OpLoad %6 %150
%170 = OpAccessChain %110 %142 %48 %48 %48 %169
%171 = OpLoad %9 %170
OpLine %3 84 14
OpLine %3 84 14
OpLine %3 87 14
OpLine %3 87 14
%172 = OpLoad %6 %150
OpLine %3 84 14
OpLine %3 87 14
%173 = OpAccessChain %110 %142 %48 %48 %172 %44
%174 = OpLoad %9 %173
OpLine %3 85 14
OpLine %3 85 14
OpLine %3 88 14
OpLine %3 88 14
%175 = OpLoad %6 %150
%176 = OpLoad %6 %150
%177 = OpAccessChain %110 %142 %48 %48 %175 %176
%178 = OpLoad %9 %177
OpLine %3 87 13
OpLine %3 89 5
OpLine %3 90 13
OpLine %3 92 5
%179 = OpLoad %6 %150
%180 = OpIAdd %6 %179 %70
OpLine %3 89 5
OpLine %3 92 5
OpStore %150 %180
OpLine %3 92 5
OpLine %3 92 5
OpLine %3 95 5
OpLine %3 95 5
%182 = OpAccessChain %181 %151 %48
OpStore %182 %143
OpLine %3 93 5
OpLine %3 93 5
OpLine %3 93 27
OpLine %3 93 43
OpLine %3 93 59
OpLine %3 93 15
OpLine %3 93 5
OpLine %3 96 5
OpLine %3 96 5
OpLine %3 96 27
OpLine %3 96 43
OpLine %3 96 59
OpLine %3 96 15
OpLine %3 96 5
%184 = OpAccessChain %183 %151 %48 %48
OpStore %184 %149
OpLine %3 94 5
OpLine %3 94 5
OpLine %3 94 5
OpLine %3 94 18
OpLine %3 94 5
OpLine %3 97 5
OpLine %3 97 5
OpLine %3 97 5
OpLine %3 97 18
OpLine %3 97 5
%185 = OpAccessChain %127 %151 %48 %48 %48
OpStore %185 %87
OpLine %3 95 5
OpLine %3 95 5
OpLine %3 98 5
OpLine %3 98 5
%186 = OpLoad %6 %150
OpLine %3 95 20
OpLine %3 95 5
OpLine %3 98 20
OpLine %3 98 5
%187 = OpAccessChain %127 %151 %48 %48 %186
OpStore %187 %89
OpLine %3 96 5
OpLine %3 96 5
OpLine %3 96 5
OpLine %3 96 5
OpLine %3 96 5
OpLine %3 99 5
OpLine %3 99 5
OpLine %3 99 5
OpLine %3 99 5
OpLine %3 99 5
%188 = OpAccessChain %28 %151 %48 %48 %48 %44
OpStore %188 %90
OpLine %3 97 5
OpLine %3 97 5
OpLine %3 97 5
OpLine %3 100 5
OpLine %3 100 5
OpLine %3 100 5
%189 = OpLoad %6 %150
OpLine %3 97 5
OpLine %3 100 5
%190 = OpAccessChain %28 %151 %48 %48 %48 %189
OpStore %190 %91
OpLine %3 98 5
OpLine %3 98 5
OpLine %3 101 5
OpLine %3 101 5
%191 = OpLoad %6 %150
OpLine %3 98 5
OpLine %3 98 5
OpLine %3 101 5
OpLine %3 101 5
%192 = OpAccessChain %28 %151 %48 %48 %191 %44
OpStore %192 %92
OpLine %3 99 5
OpLine %3 99 5
OpLine %3 102 5
OpLine %3 102 5
%193 = OpLoad %6 %150
%194 = OpLoad %6 %150
OpLine %3 99 5
OpLine %3 102 5
%195 = OpAccessChain %28 %151 %48 %48 %193 %194
OpStore %195 %93
OpReturn
@ -786,7 +789,7 @@ OpFunctionEnd
%196 = OpLabel
OpBranch %200
%200 = OpLabel
OpLine %3 102 22
OpLine %3 105 22
%201 = OpLoad %9 %197
OpReturnValue %201
OpFunctionEnd
@ -795,9 +798,9 @@ OpFunctionEnd
%202 = OpLabel
OpBranch %206
%206 = OpLabel
OpLine %3 107 12
OpLine %3 110 12
%207 = OpCompositeExtract %29 %203 4
OpLine %3 107 12
OpLine %3 110 12
%208 = OpCompositeExtract %9 %207 9
OpReturnValue %208
OpFunctionEnd
@ -806,7 +809,7 @@ OpFunctionEnd
%209 = OpLabel
OpBranch %214
%214 = OpLabel
OpLine %3 156 5
OpLine %3 159 5
OpStore %210 %213
OpReturn
OpFunctionEnd
@ -815,11 +818,11 @@ OpFunctionEnd
%215 = OpLabel
OpBranch %222
%222 = OpLabel
OpLine %3 160 32
OpLine %3 160 43
OpLine %3 160 32
OpLine %3 160 12
OpLine %3 160 5
OpLine %3 163 32
OpLine %3 163 43
OpLine %3 163 32
OpLine %3 163 12
OpLine %3 163 5
OpStore %216 %221
OpReturn
OpFunctionEnd
@ -829,13 +832,13 @@ OpFunctionEnd
%230 = OpVariable %36 Function %228
OpBranch %231
%231 = OpLabel
OpLine %3 165 5
%232 = OpFunctionCall %2 %211 %229
OpLine %3 167 35
OpLine %3 167 46
OpLine %3 167 35
OpLine %3 167 15
OpLine %3 168 5
%232 = OpFunctionCall %2 %211 %229
OpLine %3 170 35
OpLine %3 170 46
OpLine %3 170 35
OpLine %3 170 15
OpLine %3 171 5
%233 = OpFunctionCall %2 %217 %230
OpReturn
OpFunctionEnd
@ -844,7 +847,7 @@ OpFunctionEnd
%234 = OpLabel
OpBranch %238
%238 = OpLabel
OpLine %3 176 10
OpLine %3 179 10
%239 = OpAccessChain %34 %235 %48
%240 = OpLoad %4 %239
OpReturnValue %240
@ -854,8 +857,8 @@ OpFunctionEnd
%241 = OpLabel
OpBranch %245
%245 = OpLabel
OpLine %3 180 3
OpLine %3 180 3
OpLine %3 183 3
OpLine %3 183 3
%246 = OpAccessChain %34 %242 %48
OpStore %246 %17
OpReturn
@ -865,7 +868,7 @@ OpFunctionEnd
%247 = OpLabel
OpBranch %251
%251 = OpLabel
OpLine %3 184 10
OpLine %3 187 10
%252 = OpAccessChain %34 %248 %44
%253 = OpLoad %4 %252
OpReturnValue %253
@ -875,8 +878,8 @@ OpFunctionEnd
%254 = OpLabel
OpBranch %258
%258 = OpLabel
OpLine %3 188 3
OpLine %3 188 3
OpLine %3 191 3
OpLine %3 191 3
%259 = OpAccessChain %34 %255 %44
OpStore %259 %17
OpReturn
@ -887,13 +890,13 @@ OpFunctionEnd
%264 = OpVariable %41 Function %265
OpBranch %266
%266 = OpLabel
OpLine %3 193 4
OpLine %3 196 4
%267 = OpFunctionCall %2 %243 %262
OpLine %3 194 4
%268 = OpFunctionCall %4 %236 %262
OpLine %3 197 4
%268 = OpFunctionCall %4 %236 %262
OpLine %3 200 4
%269 = OpFunctionCall %2 %256 %264
OpLine %3 198 4
OpLine %3 201 4
%270 = OpFunctionCall %4 %249 %264
OpReturn
OpFunctionEnd
@ -903,11 +906,11 @@ OpFunctionEnd
%275 = OpVariable %276 Function %277
OpBranch %278
%278 = OpLabel
OpLine %3 202 13
OpLine %3 205 13
%279 = OpCompositeConstruct %43 %272
OpLine %3 202 5
OpLine %3 205 5
OpStore %275 %279
OpLine %3 204 12
OpLine %3 207 12
%281 = OpAccessChain %280 %275 %48
%282 = OpLoad %42 %281
OpReturnValue %282
@ -917,8 +920,8 @@ OpFunctionEnd
%288 = OpVariable %289 Function %287
OpBranch %290
%290 = OpLabel
OpLine %3 210 16
OpLine %3 212 12
OpLine %3 213 16
OpLine %3 215 12
%291 = OpAccessChain %95 %288 %48
%292 = OpLoad %6 %291
OpReturnValue %292
@ -927,19 +930,19 @@ OpFunctionEnd
%293 = OpLabel
OpBranch %296
%296 = OpLabel
OpLine %3 222 17
OpLine %3 225 17
%297 = OpCompositeExtract %46 %295 0
OpLine %3 223 20
OpLine %3 226 20
%298 = OpCompositeExtract %6 %297 0
OpLine %3 225 9
OpLine %3 228 9
%299 = OpCompositeExtract %4 %295 1
%300 = OpBitcast %4 %298
%301 = OpINotEqual %42 %299 %300
OpLine %3 225 5
OpLine %3 228 5
OpSelectionMerge %302 None
OpBranchConditional %301 %302 %302
%302 = OpLabel
OpLine %3 229 12
OpLine %3 232 12
%303 = OpCompositeExtract %46 %295 0
%304 = OpCompositeExtract %6 %303 0
OpReturnValue %304
@ -951,27 +954,27 @@ OpFunctionEnd
%312 = OpVariable %95 Function %313
OpBranch %314
%314 = OpLabel
OpLine %3 235 17
OpLine %3 238 17
%315 = OpAccessChain %310 %307 %48
%316 = OpLoad %46 %315
OpLine %3 235 5
OpLine %3 238 5
OpStore %309 %316
OpLine %3 236 20
OpLine %3 239 20
%317 = OpAccessChain %95 %309 %48
%318 = OpLoad %6 %317
OpLine %3 236 5
OpLine %3 239 5
OpStore %312 %318
OpLine %3 238 9
OpLine %3 241 9
%319 = OpAccessChain %34 %307 %44
%320 = OpLoad %4 %319
%321 = OpLoad %6 %312
%322 = OpBitcast %4 %321
%323 = OpINotEqual %42 %320 %322
OpLine %3 238 5
OpLine %3 241 5
OpSelectionMerge %324 None
OpBranchConditional %323 %324 %324
%324 = OpLabel
OpLine %3 242 12
OpLine %3 245 12
%325 = OpAccessChain %95 %307 %48 %48
%326 = OpLoad %6 %325
OpReturnValue %326
@ -988,58 +991,58 @@ OpBranch %348
%348 = OpLabel
OpLine %3 1 1
%349 = OpLoad %9 %344
OpLine %3 115 5
OpStore %344 %71
OpLine %3 117 9
%350 = OpLoad %7 %52
OpLine %3 118 5
OpStore %344 %71
OpLine %3 120 9
%350 = OpLoad %7 %52
OpLine %3 121 5
%351 = OpFunctionCall %2 %66
OpLine %3 119 5
OpLine %3 122 5
%352 = OpFunctionCall %2 %140
OpLine %3 122 19
OpLine %3 125 19
%354 = OpAccessChain %353 %54 %48
%355 = OpLoad %10 %354
OpLine %3 123 15
OpLine %3 126 15
%357 = OpAccessChain %356 %54 %40
%358 = OpLoad %19 %357
OpLine %3 125 13
OpLine %3 128 13
%361 = OpAccessChain %360 %54 %48 %339 %48
%362 = OpLoad %9 %361
OpLine %3 126 13
OpLine %3 126 22
OpLine %3 129 13
OpLine %3 129 22
%364 = OpArrayLength %4 %54 5
OpLine %3 126 13
OpLine %3 129 13
%365 = OpISub %4 %364 %15
%368 = OpAccessChain %367 %54 %31 %365 %48
%369 = OpLoad %6 %368
OpLine %3 127 13
OpLine %3 130 13
%370 = OpLoad %24 %336
OpLine %3 130 56
OpLine %3 130 56
OpLine %3 131 21
OpLine %3 133 56
OpLine %3 133 56
OpLine %3 134 21
%371 = OpFunctionCall %9 %198 %344
OpLine %3 134 31
OpLine %3 137 31
%374 = OpExtInst %9 %1 FClamp %362 %372 %373
%375 = OpConvertFToS %6 %374
OpLine %3 134 14
OpLine %3 137 14
%376 = OpCompositeConstruct %33 %369 %375 %340 %341 %342
OpLine %3 134 5
OpLine %3 137 5
OpStore %345 %376
OpLine %3 135 5
OpLine %3 138 5
%377 = OpIAdd %4 %330 %44
OpLine %3 135 5
OpLine %3 138 5
%378 = OpAccessChain %95 %345 %377
OpStore %378 %286
OpLine %3 136 17
OpLine %3 139 17
%379 = OpAccessChain %95 %345 %330
%380 = OpLoad %6 %379
OpLine %3 138 5
OpLine %3 141 5
%381 = OpFunctionCall %9 %204 %343
OpLine %3 140 22
OpLine %3 143 22
%383 = OpCompositeConstruct %382 %380 %380 %380 %380
%384 = OpConvertSToF %32 %383
%385 = OpMatrixTimesVector %11 %355 %384
OpLine %3 140 12
OpLine %3 143 12
%386 = OpCompositeConstruct %32 %385 %73
OpStore %331 %386
OpReturn
@ -1049,33 +1052,33 @@ OpFunctionEnd
%390 = OpAccessChain %335 %59 %48
OpBranch %401
%401 = OpLabel
OpLine %3 146 5
OpLine %3 146 5
OpLine %3 146 5
OpLine %3 149 5
OpLine %3 149 5
OpLine %3 149 5
%402 = OpAccessChain %360 %54 %48 %44 %15
OpStore %402 %71
OpLine %3 147 5
OpLine %3 147 31
OpLine %3 147 47
OpLine %3 147 63
OpLine %3 147 19
OpLine %3 147 5
OpLine %3 150 5
OpLine %3 150 31
OpLine %3 150 47
OpLine %3 150 63
OpLine %3 150 19
OpLine %3 150 5
%403 = OpAccessChain %353 %54 %48
OpStore %403 %395
OpLine %3 148 5
OpLine %3 148 35
OpLine %3 148 15
OpLine %3 148 5
OpLine %3 151 5
OpLine %3 151 35
OpLine %3 151 15
OpLine %3 151 5
%404 = OpAccessChain %356 %54 %40
OpStore %404 %398
OpLine %3 149 5
OpLine %3 149 5
OpLine %3 149 5
OpLine %3 152 5
OpLine %3 152 5
OpLine %3 152 5
%405 = OpAccessChain %367 %54 %31 %44 %48
OpStore %405 %70
OpLine %3 150 5
OpLine %3 153 5
OpStore %390 %399
OpLine %3 152 12
OpLine %3 155 12
OpStore %388 %400
OpReturn
OpFunctionEnd
@ -1083,17 +1086,17 @@ OpFunctionEnd
%406 = OpLabel
OpBranch %409
%409 = OpLabel
OpLine %3 247 5
%410 = OpFunctionCall %2 %224
OpLine %3 248 5
%411 = OpFunctionCall %2 %261
OpLine %3 249 5
%412 = OpFunctionCall %42 %273 %408
OpLine %3 250 5
%413 = OpFunctionCall %6 %284
%410 = OpFunctionCall %2 %224
OpLine %3 251 5
%414 = OpFunctionCall %6 %294
%411 = OpFunctionCall %2 %261
OpLine %3 252 5
%412 = OpFunctionCall %42 %273 %408
OpLine %3 253 5
%413 = OpFunctionCall %6 %284
OpLine %3 254 5
%414 = OpFunctionCall %6 %294
OpLine %3 255 5
%415 = OpFunctionCall %6 %306
OpReturn
OpFunctionEnd