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

Fix hlsl output for writes to scalar/vector storage buffer

Browse Source
This commit is contained in:
Hasan Ali
2022-05-10 21:22:20 +01:00
committed by Teodor Tanasoaia
parent b3d5e6d807
commit 205ea6cc5c
10 changed files with 464 additions and 421 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -44,6 +44,9 @@ impl<W: fmt::Write> super::Writer<'_, W> {
chain: &[SubAccess],
func_ctx: &FunctionCtx,
) -> BackendResult {
if chain.is_empty() {
write!(self.out, "0")?;
}
for (i, access) in chain.iter().enumerate() {
if i != 0 {
write!(self.out, "+")?;

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

@@ -11,12 +11,14 @@
resources: {
(group: 0, binding: 0): (buffer: Some(0), mutable: false),
(group: 0, binding: 1): (buffer: Some(1), mutable: false),
(group: 0, binding: 2): (buffer: Some(2), mutable: false),
},
sizes_buffer: Some(24),
),
fs: (
resources: {
(group: 0, binding: 0): (buffer: Some(0), mutable: true),
(group: 0, binding: 2): (buffer: Some(2), mutable: true),
},
sizes_buffer: Some(24),
),

5
tests/in/access.wgsl
View File

@@ -22,6 +22,9 @@ struct Baz {
@group(0) @binding(1)
var<uniform> baz: Baz;
@group(0) @binding(2)
var<storage,read_write> qux: vec2<i32>;
fn test_matrix_within_struct_accesses() {
var idx = 9;
@@ -73,6 +76,7 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let index = 3u;
let b = bar._matrix[index].x;
let a = bar.data[arrayLength(&bar.data) - 2u].value;
let c = qux;
// test pointer types
let data_pointer: ptr<storage, i32, read_write> = &bar.data[0].value;
@@ -95,6 +99,7 @@ fn foo_frag() -> @location(0) vec4<f32> {
bar._matrix = mat4x3<f32>(vec3<f32>(0.0), vec3<f32>(1.0), vec3<f32>(2.0), vec3<f32>(3.0));
bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
bar.data[1].value = 1;
qux = vec2<i32>();
return vec4<f32>(0.0);
}

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

@@ -21,8 +21,8 @@ layout(std430) buffer Bar_block_0Compute {
float read_from_private(inout float foo_1) {
float _e3 = foo_1;
return _e3;
float _e4 = foo_1;
return _e4;
}
float test_arr_as_arg(float a[5][10]) {
@@ -32,22 +32,22 @@ float test_arr_as_arg(float a[5][10]) {
void main() {
int tmp = 0;
int value = _group_0_binding_0_cs.atom;
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;
int _e8 = atomicAdd(_group_0_binding_0_cs.atom, 5);
tmp = _e8;
int _e11 = atomicAdd(_group_0_binding_0_cs.atom, -5);
tmp = _e11;
int _e14 = atomicAnd(_group_0_binding_0_cs.atom, 5);
tmp = _e14;
int _e17 = atomicOr(_group_0_binding_0_cs.atom, 5);
tmp = _e17;
int _e20 = atomicXor(_group_0_binding_0_cs.atom, 5);
tmp = _e20;
int _e23 = atomicMin(_group_0_binding_0_cs.atom, 5);
tmp = _e23;
int _e26 = atomicMax(_group_0_binding_0_cs.atom, 5);
tmp = _e26;
int _e29 = atomicExchange(_group_0_binding_0_cs.atom, 5);
tmp = _e29;
_group_0_binding_0_cs.atom = value;
return;
}

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

@@ -17,11 +17,13 @@ layout(std430) buffer Bar_block_0Fragment {
AlignedWrapper data[];
} _group_0_binding_0_fs;
layout(std430) buffer type_9_block_1Fragment { ivec2 _group_0_binding_2_fs; };
layout(location = 0) out vec4 _fs2p_location0;
float read_from_private(inout float foo_1) {
float _e3 = foo_1;
return _e3;
float _e4 = foo_1;
return _e4;
}
float test_arr_as_arg(float a[5][10]) {
@@ -33,6 +35,7 @@ void main() {
_group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));
_group_0_binding_0_fs.arr = uvec2[2](uvec2(0u), uvec2(1u));
_group_0_binding_0_fs.data[1].value = 1;
_group_0_binding_2_fs = ivec2(0, 0);
_fs2p_location0 = vec4(0.0);
return;
}

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

@@ -19,45 +19,47 @@ layout(std430) buffer Bar_block_0Vertex {
uniform Baz_block_1Vertex { Baz _group_0_binding_1_vs; };
layout(std430) buffer type_9_block_2Vertex { ivec2 _group_0_binding_2_vs; };
void test_matrix_within_struct_accesses() {
int idx = 9;
Baz t = Baz(mat3x2(0.0));
int _e4 = idx;
idx = (_e4 - 1);
int _e5 = idx;
idx = (_e5 - 1);
mat3x2 unnamed = _group_0_binding_1_vs.m;
vec2 unnamed_1 = _group_0_binding_1_vs.m[0];
int _e14 = idx;
vec2 unnamed_2 = _group_0_binding_1_vs.m[_e14];
int _e15 = idx;
vec2 unnamed_2 = _group_0_binding_1_vs.m[_e15];
float unnamed_3 = _group_0_binding_1_vs.m[0][1];
int _e26 = idx;
float unnamed_4 = _group_0_binding_1_vs.m[0][_e26];
int _e30 = idx;
float unnamed_5 = _group_0_binding_1_vs.m[_e30][1];
int _e36 = idx;
int _e38 = idx;
float unnamed_6 = _group_0_binding_1_vs.m[_e36][_e38];
int _e27 = idx;
float unnamed_4 = _group_0_binding_1_vs.m[0][_e27];
int _e31 = idx;
float unnamed_5 = _group_0_binding_1_vs.m[_e31][1];
int _e37 = idx;
int _e39 = idx;
float unnamed_6 = _group_0_binding_1_vs.m[_e37][_e39];
t = Baz(mat3x2(vec2(1.0), vec2(2.0), vec2(3.0)));
int _e50 = idx;
idx = (_e50 + 1);
int _e51 = idx;
idx = (_e51 + 1);
t.m = mat3x2(vec2(6.0), vec2(5.0), vec2(4.0));
t.m[0] = vec2(9.0);
int _e67 = idx;
t.m[_e67] = vec2(90.0);
int _e68 = idx;
t.m[_e68] = vec2(90.0);
t.m[0][1] = 10.0;
int _e80 = idx;
t.m[0][_e80] = 20.0;
int _e84 = idx;
t.m[_e84][1] = 30.0;
int _e90 = idx;
int _e92 = idx;
t.m[_e90][_e92] = 40.0;
int _e81 = idx;
t.m[0][_e81] = 20.0;
int _e85 = idx;
t.m[_e85][1] = 30.0;
int _e91 = idx;
int _e93 = idx;
t.m[_e91][_e93] = 40.0;
return;
}
float read_from_private(inout float foo_1) {
float _e3 = foo_1;
return _e3;
float _e4 = foo_1;
return _e4;
}
float test_arr_as_arg(float a[5][10]) {
@@ -75,11 +77,12 @@ void main() {
uvec2 arr[2] = _group_0_binding_0_vs.arr;
float b = _group_0_binding_0_vs._matrix[3][0];
int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
float _e28 = read_from_private(foo);
ivec2 c_1 = _group_0_binding_2_vs;
float _e30 = read_from_private(foo);
c = int[5](a_1, int(b), 3, 4, 5);
c[(vi + 1u)] = 42;
int value = c[vi];
float _e42 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
float _e44 = test_arr_as_arg(float[5][10](float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), float[10](0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
gl_Position = vec4((_matrix * vec4(ivec4(value))), 2.0);
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
return;

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

@@ -20,6 +20,7 @@ float Constructarray5_array10_float__(float arg0[10], float arg1[10], float arg2
RWByteAddressBuffer bar : register(u0);
cbuffer baz : register(b1) { Baz baz; }
RWByteAddressBuffer qux : register(u2);
float3x2 GetMatmOnBaz(Baz obj) {
return float3x2(obj.m_0, obj.m_1, obj.m_2);
@@ -60,42 +61,42 @@ void test_matrix_within_struct_accesses()
int idx = 9;
Baz t = (Baz)0;
int _expr4 = idx;
idx = (_expr4 - 1);
int _expr5 = idx;
idx = (_expr5 - 1);
float3x2 unnamed = GetMatmOnBaz(baz);
float2 unnamed_1 = GetMatmOnBaz(baz)[0];
int _expr14 = idx;
float2 unnamed_2 = GetMatmOnBaz(baz)[_expr14];
int _expr15 = idx;
float2 unnamed_2 = GetMatmOnBaz(baz)[_expr15];
float unnamed_3 = GetMatmOnBaz(baz)[0][1];
int _expr26 = idx;
float unnamed_4 = GetMatmOnBaz(baz)[0][_expr26];
int _expr30 = idx;
float unnamed_5 = GetMatmOnBaz(baz)[_expr30][1];
int _expr36 = idx;
int _expr38 = idx;
float unnamed_6 = GetMatmOnBaz(baz)[_expr36][_expr38];
int _expr27 = idx;
float unnamed_4 = GetMatmOnBaz(baz)[0][_expr27];
int _expr31 = idx;
float unnamed_5 = GetMatmOnBaz(baz)[_expr31][1];
int _expr37 = idx;
int _expr39 = idx;
float unnamed_6 = GetMatmOnBaz(baz)[_expr37][_expr39];
t = ConstructBaz(float3x2((1.0).xx, (2.0).xx, (3.0).xx));
int _expr50 = idx;
idx = (_expr50 + 1);
int _expr51 = idx;
idx = (_expr51 + 1);
SetMatmOnBaz(t, float3x2((6.0).xx, (5.0).xx, (4.0).xx));
t.m_0 = (9.0).xx;
int _expr67 = idx;
SetMatVecmOnBaz(t, (90.0).xx, _expr67);
int _expr68 = idx;
SetMatVecmOnBaz(t, (90.0).xx, _expr68);
t.m_0[1] = 10.0;
int _expr80 = idx;
t.m_0[_expr80] = 20.0;
int _expr84 = idx;
SetMatScalarmOnBaz(t, 30.0, _expr84, 1);
int _expr90 = idx;
int _expr92 = idx;
SetMatScalarmOnBaz(t, 40.0, _expr90, _expr92);
int _expr81 = idx;
t.m_0[_expr81] = 20.0;
int _expr85 = idx;
SetMatScalarmOnBaz(t, 30.0, _expr85, 1);
int _expr91 = idx;
int _expr93 = idx;
SetMatScalarmOnBaz(t, 40.0, _expr91, _expr93);
return;
}
float read_from_private(inout float foo_1)
{
float _expr3 = foo_1;
return _expr3;
float _expr4 = foo_1;
return _expr4;
}
float test_arr_as_arg(float a[5][10])
@@ -127,14 +128,15 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))};
float b = asfloat(bar.Load(0+48+0));
int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120));
const float _e28 = read_from_private(foo);
int2 c_1 = asint(qux.Load2(0));
const float _e30 = read_from_private(foo);
{
int _result[5]=Constructarray5_int_(a_1, int(b), 3, 4, 5);
for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
}
c[(vi + 1u)] = 42;
int value = c[vi];
const float _e42 = test_arr_as_arg(Constructarray5_array10_float__(Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
const float _e44 = test_arr_as_arg(Constructarray5_array10_float__(Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), Constructarray10_float_(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
return float4(mul(float4((value).xxxx), _matrix), 2.0);
}
@@ -159,6 +161,7 @@ float4 foo_frag() : SV_Target0
bar.Store2(104+8, asuint(_value2[1]));
}
bar.Store(0+8+120, asuint(1));
qux.Store2(0, asuint(int2(0, 0)));
return (0.0).xxxx;
}
@@ -168,22 +171,22 @@ void atomics()
int tmp = (int)0;
int value_1 = asint(bar.Load(96));
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;
int _e8; bar.InterlockedAdd(96, 5, _e8);
tmp = _e8;
int _e11; bar.InterlockedAdd(96, -5, _e11);
tmp = _e11;
int _e14; bar.InterlockedAnd(96, 5, _e14);
tmp = _e14;
int _e17; bar.InterlockedOr(96, 5, _e17);
tmp = _e17;
int _e20; bar.InterlockedXor(96, 5, _e20);
tmp = _e20;
int _e23; bar.InterlockedMin(96, 5, _e23);
tmp = _e23;
int _e26; bar.InterlockedMax(96, 5, _e26);
tmp = _e26;
int _e29; bar.InterlockedExchange(96, 5, _e29);
tmp = _e29;
bar.Store(96, asuint(value_1));
return;
}

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

@@ -29,64 +29,65 @@ struct Bar {
struct Baz {
metal::float3x2 m;
};
struct type_11 {
struct type_12 {
float inner[10];
};
struct type_12 {
type_11 inner[5];
struct type_13 {
type_12 inner[5];
};
struct type_16 {
struct type_17 {
int inner[5];
};
constant type_11 const_type_11_ = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
constant type_12 const_type_12_ = {const_type_11_, const_type_11_, const_type_11_, const_type_11_, const_type_11_};
constant type_12 const_type_12_ = {0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0};
constant type_13 const_type_13_ = {const_type_12_, const_type_12_, const_type_12_, const_type_12_, const_type_12_};
constant metal::int2 const_type_9_ = {0, 0};
void test_matrix_within_struct_accesses(
constant Baz& baz
) {
int idx = 9;
Baz t = {};
int _e4 = idx;
idx = _e4 - 1;
int _e5 = idx;
idx = _e5 - 1;
metal::float3x2 unnamed = baz.m;
metal::float2 unnamed_1 = baz.m[0];
int _e14 = idx;
metal::float2 unnamed_2 = baz.m[_e14];
int _e15 = idx;
metal::float2 unnamed_2 = baz.m[_e15];
float unnamed_3 = baz.m[0].y;
int _e26 = idx;
float unnamed_4 = baz.m[0][_e26];
int _e30 = idx;
float unnamed_5 = baz.m[_e30].y;
int _e36 = idx;
int _e38 = idx;
float unnamed_6 = baz.m[_e36][_e38];
int _e27 = idx;
float unnamed_4 = baz.m[0][_e27];
int _e31 = idx;
float unnamed_5 = baz.m[_e31].y;
int _e37 = idx;
int _e39 = idx;
float unnamed_6 = baz.m[_e37][_e39];
t = Baz {metal::float3x2(metal::float2(1.0), metal::float2(2.0), metal::float2(3.0))};
int _e50 = idx;
idx = _e50 + 1;
int _e51 = idx;
idx = _e51 + 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 _e67 = idx;
t.m[_e67] = metal::float2(90.0);
int _e68 = idx;
t.m[_e68] = metal::float2(90.0);
t.m[0].y = 10.0;
int _e80 = idx;
t.m[0][_e80] = 20.0;
int _e84 = idx;
t.m[_e84].y = 30.0;
int _e90 = idx;
int _e92 = idx;
t.m[_e90][_e92] = 40.0;
int _e81 = idx;
t.m[0][_e81] = 20.0;
int _e85 = idx;
t.m[_e85].y = 30.0;
int _e91 = idx;
int _e93 = idx;
t.m[_e91][_e93] = 40.0;
return;
}
float read_from_private(
thread float& foo_1
) {
float _e3 = foo_1;
return _e3;
float _e4 = foo_1;
return _e4;
}
float test_arr_as_arg(
type_12 a
type_13 a
) {
return a.inner[4].inner[9];
}
@@ -100,10 +101,11 @@ vertex foo_vertOutput foo_vert(
uint vi [[vertex_id]]
, device Bar const& bar [[buffer(0)]]
, constant Baz& baz [[buffer(1)]]
, device metal::int2 const& qux [[buffer(2)]]
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
float foo = 0.0;
type_16 c = {};
type_17 c = {};
float baz_1 = foo;
foo = 1.0;
test_matrix_within_struct_accesses(baz);
@@ -111,11 +113,12 @@ vertex foo_vertOutput foo_vert(
type_6 arr = bar.arr;
float b = bar._matrix[3].x;
int a_1 = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value;
float _e28 = read_from_private(foo);
for(int _i=0; _i<5; ++_i) c.inner[_i] = type_16 {a_1, static_cast<int>(b), 3, 4, 5}.inner[_i];
metal::int2 c_1 = qux;
float _e30 = read_from_private(foo);
for(int _i=0; _i<5; ++_i) c.inner[_i] = type_17 {a_1, static_cast<int>(b), 3, 4, 5}.inner[_i];
c.inner[vi + 1u] = 42;
int value = c.inner[vi];
float _e42 = test_arr_as_arg(const_type_12_);
float _e44 = test_arr_as_arg(const_type_13_);
return foo_vertOutput { metal::float4(_matrix * static_cast<metal::float4>(metal::int4(value)), 2.0) };
}
@@ -125,12 +128,14 @@ struct foo_fragOutput {
};
fragment foo_fragOutput foo_frag(
device Bar& bar [[buffer(0)]]
, device metal::int2& qux [[buffer(2)]]
, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
) {
bar._matrix[1].z = 1.0;
bar._matrix = metal::float4x3(metal::float3(0.0), metal::float3(1.0), metal::float3(2.0), metal::float3(3.0));
for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_6 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
bar.data[1].value = 1;
qux = const_type_9_;
return foo_fragOutput { metal::float4(0.0) };
}
@@ -141,22 +146,22 @@ kernel void atomics(
) {
int tmp = {};
int value_1 = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed);
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;
int _e8 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e8;
int _e11 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e11;
int _e14 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e14;
int _e17 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e17;
int _e20 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e20;
int _e23 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e23;
int _e26 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e26;
int _e29 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e29;
metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed);
return;
}

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

@@ -1,16 +1,16 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 221
; Bound: 230
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %143 "foo_vert" %138 %141
OpEntryPoint Fragment %180 "foo_frag" %179
OpEntryPoint GLCompute %198 "atomics"
OpExecutionMode %180 OriginUpperLeft
OpExecutionMode %198 LocalSize 1 1 1
OpEntryPoint Vertex %149 "foo_vert" %144 %147
OpEntryPoint Fragment %188 "foo_frag" %187
OpEntryPoint GLCompute %207 "atomics"
OpExecutionMode %188 OriginUpperLeft
OpExecutionMode %207 LocalSize 1 1 1
OpSource GLSL 450
OpMemberName %32 0 "value"
OpName %32 "AlignedWrapper"
@@ -22,22 +22,23 @@ OpMemberName %41 4 "data"
OpName %41 "Bar"
OpMemberName %43 0 "m"
OpName %43 "Baz"
OpName %52 "bar"
OpName %54 "baz"
OpName %57 "idx"
OpName %59 "t"
OpName %63 "test_matrix_within_struct_accesses"
OpName %121 "foo"
OpName %122 "read_from_private"
OpName %127 "a"
OpName %128 "test_arr_as_arg"
OpName %133 "foo"
OpName %134 "c"
OpName %138 "vi"
OpName %143 "foo_vert"
OpName %180 "foo_frag"
OpName %195 "tmp"
OpName %198 "atomics"
OpName %54 "bar"
OpName %56 "baz"
OpName %59 "qux"
OpName %62 "idx"
OpName %64 "t"
OpName %68 "test_matrix_within_struct_accesses"
OpName %127 "foo"
OpName %128 "read_from_private"
OpName %133 "a"
OpName %134 "test_arr_as_arg"
OpName %139 "foo"
OpName %140 "c"
OpName %144 "vi"
OpName %149 "foo_vert"
OpName %188 "foo_frag"
OpName %204 "tmp"
OpName %207 "atomics"
OpMemberDecorate %32 0 Offset 0
OpDecorate %37 ArrayStride 16
OpDecorate %39 ArrayStride 8
@@ -54,19 +55,23 @@ OpMemberDecorate %41 4 Offset 120
OpMemberDecorate %43 0 Offset 0
OpMemberDecorate %43 0 ColMajor
OpMemberDecorate %43 0 MatrixStride 8
OpDecorate %45 ArrayStride 4
OpDecorate %46 ArrayStride 40
OpDecorate %49 ArrayStride 4
OpDecorate %52 DescriptorSet 0
OpDecorate %52 Binding 0
OpDecorate %41 Block
OpDecorate %46 ArrayStride 4
OpDecorate %47 ArrayStride 40
OpDecorate %50 ArrayStride 4
OpDecorate %54 DescriptorSet 0
OpDecorate %54 Binding 1
OpDecorate %55 Block
OpMemberDecorate %55 0 Offset 0
OpDecorate %138 BuiltIn VertexIndex
OpDecorate %141 BuiltIn Position
OpDecorate %179 Location 0
OpDecorate %54 Binding 0
OpDecorate %41 Block
OpDecorate %56 DescriptorSet 0
OpDecorate %56 Binding 1
OpDecorate %57 Block
OpMemberDecorate %57 0 Offset 0
OpDecorate %59 DescriptorSet 0
OpDecorate %59 Binding 2
OpDecorate %60 Block
OpMemberDecorate %60 0 Offset 0
OpDecorate %144 BuiltIn VertexIndex
OpDecorate %147 BuiltIn Position
OpDecorate %187 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 2
@@ -109,226 +114,236 @@ OpDecorate %179 Location 0
%41 = OpTypeStruct %33 %37 %4 %39 %40
%42 = OpTypeMatrix %36 3
%43 = OpTypeStruct %42
%44 = OpTypePointer Function %9
%45 = OpTypeArray %9 %21
%46 = OpTypeArray %45 %22
%47 = OpTypeVector %9 4
%48 = OpTypePointer StorageBuffer %4
%49 = OpTypeArray %4 %22
%50 = OpConstantComposite %45 %24 %24 %24 %24 %24 %24 %24 %24 %24 %24
%51 = OpConstantComposite %46 %50 %50 %50 %50 %50
%53 = OpTypePointer StorageBuffer %41
%52 = OpVariable %53 StorageBuffer
%55 = OpTypeStruct %43
%56 = OpTypePointer Uniform %55
%54 = OpVariable %56 Uniform
%58 = OpTypePointer Function %4
%60 = OpTypePointer Function %43
%61 = OpConstantNull %43
%64 = OpTypeFunction %2
%65 = OpTypePointer Uniform %43
%70 = OpTypePointer Uniform %42
%73 = OpTypePointer Uniform %36
%79 = OpTypePointer Uniform %9
%99 = OpTypePointer Function %42
%105 = OpTypePointer Function %36
%111 = OpTypePointer Function %9
%123 = OpTypeFunction %9 %44
%129 = OpTypeFunction %9 %46
%135 = OpTypePointer Function %49
%136 = OpConstantNull %49
%139 = OpTypePointer Input %26
%138 = OpVariable %139 Input
%142 = OpTypePointer Output %47
%141 = OpVariable %142 Output
%148 = OpTypePointer StorageBuffer %33
%151 = OpTypePointer StorageBuffer %39
%154 = OpTypePointer StorageBuffer %34
%155 = OpTypePointer StorageBuffer %9
%158 = OpTypePointer StorageBuffer %40
%161 = OpTypePointer StorageBuffer %32
%162 = OpConstant %26 4
%173 = OpTypeVector %4 4
%179 = OpVariable %142 Output
%196 = OpConstantNull %4
%200 = OpTypePointer StorageBuffer %4
%203 = OpConstant %26 64
%63 = OpFunction %2 None %64
%62 = OpLabel
%57 = OpVariable %58 Function %5
%59 = OpVariable %60 Function %61
%66 = OpAccessChain %65 %54 %31
OpBranch %67
%44 = OpTypeVector %4 2
%45 = OpTypePointer Function %9
%46 = OpTypeArray %9 %21
%47 = OpTypeArray %46 %22
%48 = OpTypeVector %9 4
%49 = OpTypePointer StorageBuffer %4
%50 = OpTypeArray %4 %22
%51 = OpConstantComposite %46 %24 %24 %24 %24 %24 %24 %24 %24 %24 %24
%52 = OpConstantComposite %47 %51 %51 %51 %51 %51
%53 = OpConstantComposite %44 %7 %7
%55 = OpTypePointer StorageBuffer %41
%54 = OpVariable %55 StorageBuffer
%57 = OpTypeStruct %43
%58 = OpTypePointer Uniform %57
%56 = OpVariable %58 Uniform
%60 = OpTypeStruct %44
%61 = OpTypePointer StorageBuffer %60
%59 = OpVariable %61 StorageBuffer
%63 = OpTypePointer Function %4
%65 = OpTypePointer Function %43
%66 = OpConstantNull %43
%69 = OpTypeFunction %2
%70 = OpTypePointer Uniform %43
%72 = OpTypePointer StorageBuffer %44
%76 = OpTypePointer Uniform %42
%79 = OpTypePointer Uniform %36
%85 = OpTypePointer Uniform %9
%105 = OpTypePointer Function %42
%111 = OpTypePointer Function %36
%117 = OpTypePointer Function %9
%129 = OpTypeFunction %9 %45
%135 = OpTypeFunction %9 %47
%141 = OpTypePointer Function %50
%142 = OpConstantNull %50
%145 = OpTypePointer Input %26
%144 = OpVariable %145 Input
%148 = OpTypePointer Output %48
%147 = OpVariable %148 Output
%155 = OpTypePointer StorageBuffer %33
%158 = OpTypePointer StorageBuffer %39
%161 = OpTypePointer StorageBuffer %34
%162 = OpTypePointer StorageBuffer %9
%165 = OpTypePointer StorageBuffer %40
%168 = OpTypePointer StorageBuffer %32
%169 = OpConstant %26 4
%181 = OpTypeVector %4 4
%187 = OpVariable %148 Output
%205 = OpConstantNull %4
%209 = OpTypePointer StorageBuffer %4
%212 = OpConstant %26 64
%68 = OpFunction %2 None %69
%67 = OpLabel
%68 = OpLoad %4 %57
%69 = OpISub %4 %68 %6
OpStore %57 %69
%71 = OpAccessChain %70 %66 %31
%72 = OpLoad %42 %71
%74 = OpAccessChain %73 %66 %31 %31
%75 = OpLoad %36 %74
%76 = OpLoad %4 %57
%77 = OpAccessChain %73 %66 %31 %76
%78 = OpLoad %36 %77
%80 = OpAccessChain %79 %66 %31 %31 %29
%81 = OpLoad %9 %80
%82 = OpLoad %4 %57
%83 = OpAccessChain %79 %66 %31 %31 %82
%84 = OpLoad %9 %83
%85 = OpLoad %4 %57
%86 = OpAccessChain %79 %66 %31 %85 %29
%62 = OpVariable %63 Function %5
%64 = OpVariable %65 Function %66
%71 = OpAccessChain %70 %56 %31
OpBranch %73
%73 = OpLabel
%74 = OpLoad %4 %62
%75 = OpISub %4 %74 %6
OpStore %62 %75
%77 = OpAccessChain %76 %71 %31
%78 = OpLoad %42 %77
%80 = OpAccessChain %79 %71 %31 %31
%81 = OpLoad %36 %80
%82 = OpLoad %4 %62
%83 = OpAccessChain %79 %71 %31 %82
%84 = OpLoad %36 %83
%86 = OpAccessChain %85 %71 %31 %31 %29
%87 = OpLoad %9 %86
%88 = OpLoad %4 %57
%89 = OpLoad %4 %57
%90 = OpAccessChain %79 %66 %31 %88 %89
%91 = OpLoad %9 %90
%92 = OpCompositeConstruct %36 %8 %8
%93 = OpCompositeConstruct %36 %10 %10
%94 = OpCompositeConstruct %36 %11 %11
%95 = OpCompositeConstruct %42 %92 %93 %94
%96 = OpCompositeConstruct %43 %95
OpStore %59 %96
%97 = OpLoad %4 %57
%98 = OpIAdd %4 %97 %6
OpStore %57 %98
%100 = OpCompositeConstruct %36 %12 %12
%101 = OpCompositeConstruct %36 %13 %13
%102 = OpCompositeConstruct %36 %14 %14
%103 = OpCompositeConstruct %42 %100 %101 %102
%104 = OpAccessChain %99 %59 %31
OpStore %104 %103
%106 = OpCompositeConstruct %36 %15 %15
%107 = OpAccessChain %105 %59 %31 %31
OpStore %107 %106
%108 = OpLoad %4 %57
%109 = OpCompositeConstruct %36 %16 %16
%110 = OpAccessChain %105 %59 %31 %108
%88 = OpLoad %4 %62
%89 = OpAccessChain %85 %71 %31 %31 %88
%90 = OpLoad %9 %89
%91 = OpLoad %4 %62
%92 = OpAccessChain %85 %71 %31 %91 %29
%93 = OpLoad %9 %92
%94 = OpLoad %4 %62
%95 = OpLoad %4 %62
%96 = OpAccessChain %85 %71 %31 %94 %95
%97 = OpLoad %9 %96
%98 = OpCompositeConstruct %36 %8 %8
%99 = OpCompositeConstruct %36 %10 %10
%100 = OpCompositeConstruct %36 %11 %11
%101 = OpCompositeConstruct %42 %98 %99 %100
%102 = OpCompositeConstruct %43 %101
OpStore %64 %102
%103 = OpLoad %4 %62
%104 = OpIAdd %4 %103 %6
OpStore %62 %104
%106 = OpCompositeConstruct %36 %12 %12
%107 = OpCompositeConstruct %36 %13 %13
%108 = OpCompositeConstruct %36 %14 %14
%109 = OpCompositeConstruct %42 %106 %107 %108
%110 = OpAccessChain %105 %64 %31
OpStore %110 %109
%112 = OpAccessChain %111 %59 %31 %31 %29
OpStore %112 %17
%113 = OpLoad %4 %57
%114 = OpAccessChain %111 %59 %31 %31 %113
OpStore %114 %18
%115 = OpLoad %4 %57
%116 = OpAccessChain %111 %59 %31 %115 %29
OpStore %116 %19
%117 = OpLoad %4 %57
%118 = OpLoad %4 %57
%119 = OpAccessChain %111 %59 %31 %117 %118
OpStore %119 %20
%112 = OpCompositeConstruct %36 %15 %15
%113 = OpAccessChain %111 %64 %31 %31
OpStore %113 %112
%114 = OpLoad %4 %62
%115 = OpCompositeConstruct %36 %16 %16
%116 = OpAccessChain %111 %64 %31 %114
OpStore %116 %115
%118 = OpAccessChain %117 %64 %31 %31 %29
OpStore %118 %17
%119 = OpLoad %4 %62
%120 = OpAccessChain %117 %64 %31 %31 %119
OpStore %120 %18
%121 = OpLoad %4 %62
%122 = OpAccessChain %117 %64 %31 %121 %29
OpStore %122 %19
%123 = OpLoad %4 %62
%124 = OpLoad %4 %62
%125 = OpAccessChain %117 %64 %31 %123 %124
OpStore %125 %20
OpReturn
OpFunctionEnd
%122 = OpFunction %9 None %123
%121 = OpFunctionParameter %44
%120 = OpLabel
OpBranch %124
%124 = OpLabel
%125 = OpLoad %9 %121
OpReturnValue %125
OpFunctionEnd
%128 = OpFunction %9 None %129
%127 = OpFunctionParameter %46
%127 = OpFunctionParameter %45
%126 = OpLabel
OpBranch %130
%130 = OpLabel
%131 = OpCompositeExtract %45 %127 4
%132 = OpCompositeExtract %9 %131 9
OpReturnValue %132
%131 = OpLoad %9 %127
OpReturnValue %131
OpFunctionEnd
%143 = OpFunction %2 None %64
%137 = OpLabel
%133 = OpVariable %44 Function %24
%134 = OpVariable %135 Function %136
%140 = OpLoad %26 %138
%144 = OpAccessChain %65 %54 %31
OpBranch %145
%145 = OpLabel
%146 = OpLoad %9 %133
OpStore %133 %8
%147 = OpFunctionCall %2 %63
%149 = OpAccessChain %148 %52 %31
%150 = OpLoad %33 %149
%152 = OpAccessChain %151 %52 %25
%153 = OpLoad %39 %152
%156 = OpAccessChain %155 %52 %31 %25 %31
%157 = OpLoad %9 %156
%159 = OpArrayLength %26 %52 4
%160 = OpISub %26 %159 %27
%163 = OpAccessChain %48 %52 %162 %160 %31
%164 = OpLoad %4 %163
%165 = OpFunctionCall %9 %122 %133
%166 = OpConvertFToS %4 %157
%167 = OpCompositeConstruct %49 %164 %166 %28 %23 %22
OpStore %134 %167
%168 = OpIAdd %26 %140 %29
%169 = OpAccessChain %58 %134 %168
OpStore %169 %30
%170 = OpAccessChain %58 %134 %140
%134 = OpFunction %9 None %135
%133 = OpFunctionParameter %47
%132 = OpLabel
OpBranch %136
%136 = OpLabel
%137 = OpCompositeExtract %46 %133 4
%138 = OpCompositeExtract %9 %137 9
OpReturnValue %138
OpFunctionEnd
%149 = OpFunction %2 None %69
%143 = OpLabel
%139 = OpVariable %45 Function %24
%140 = OpVariable %141 Function %142
%146 = OpLoad %26 %144
%150 = OpAccessChain %70 %56 %31
%151 = OpAccessChain %72 %59 %31
OpBranch %152
%152 = OpLabel
%153 = OpLoad %9 %139
OpStore %139 %8
%154 = OpFunctionCall %2 %68
%156 = OpAccessChain %155 %54 %31
%157 = OpLoad %33 %156
%159 = OpAccessChain %158 %54 %25
%160 = OpLoad %39 %159
%163 = OpAccessChain %162 %54 %31 %25 %31
%164 = OpLoad %9 %163
%166 = OpArrayLength %26 %54 4
%167 = OpISub %26 %166 %27
%170 = OpAccessChain %49 %54 %169 %167 %31
%171 = OpLoad %4 %170
%172 = OpFunctionCall %9 %128 %51
%174 = OpCompositeConstruct %173 %171 %171 %171 %171
%175 = OpConvertSToF %47 %174
%176 = OpMatrixTimesVector %34 %150 %175
%177 = OpCompositeConstruct %47 %176 %10
OpStore %141 %177
%172 = OpLoad %44 %151
%173 = OpFunctionCall %9 %128 %139
%174 = OpConvertFToS %4 %164
%175 = OpCompositeConstruct %50 %171 %174 %28 %23 %22
OpStore %140 %175
%176 = OpIAdd %26 %146 %29
%177 = OpAccessChain %63 %140 %176
OpStore %177 %30
%178 = OpAccessChain %63 %140 %146
%179 = OpLoad %4 %178
%180 = OpFunctionCall %9 %134 %52
%182 = OpCompositeConstruct %181 %179 %179 %179 %179
%183 = OpConvertSToF %48 %182
%184 = OpMatrixTimesVector %34 %157 %183
%185 = OpCompositeConstruct %48 %184 %10
OpStore %147 %185
OpReturn
OpFunctionEnd
%180 = OpFunction %2 None %64
%178 = OpLabel
OpBranch %181
%181 = OpLabel
%182 = OpAccessChain %155 %52 %31 %29 %27
OpStore %182 %8
%183 = OpCompositeConstruct %34 %24 %24 %24
%184 = OpCompositeConstruct %34 %8 %8 %8
%185 = OpCompositeConstruct %34 %10 %10 %10
%186 = OpCompositeConstruct %34 %11 %11 %11
%187 = OpCompositeConstruct %33 %183 %184 %185 %186
%188 = OpAccessChain %148 %52 %31
OpStore %188 %187
%189 = OpCompositeConstruct %38 %31 %31
%190 = OpCompositeConstruct %38 %29 %29
%191 = OpCompositeConstruct %39 %189 %190
%192 = OpAccessChain %151 %52 %25
OpStore %192 %191
%193 = OpAccessChain %48 %52 %162 %29 %31
OpStore %193 %6
%194 = OpCompositeConstruct %47 %24 %24 %24 %24
OpStore %179 %194
%188 = OpFunction %2 None %69
%186 = OpLabel
%189 = OpAccessChain %72 %59 %31
OpBranch %190
%190 = OpLabel
%191 = OpAccessChain %162 %54 %31 %29 %27
OpStore %191 %8
%192 = OpCompositeConstruct %34 %24 %24 %24
%193 = OpCompositeConstruct %34 %8 %8 %8
%194 = OpCompositeConstruct %34 %10 %10 %10
%195 = OpCompositeConstruct %34 %11 %11 %11
%196 = OpCompositeConstruct %33 %192 %193 %194 %195
%197 = OpAccessChain %155 %54 %31
OpStore %197 %196
%198 = OpCompositeConstruct %38 %31 %31
%199 = OpCompositeConstruct %38 %29 %29
%200 = OpCompositeConstruct %39 %198 %199
%201 = OpAccessChain %158 %54 %25
OpStore %201 %200
%202 = OpAccessChain %49 %54 %169 %29 %31
OpStore %202 %6
OpStore %189 %53
%203 = OpCompositeConstruct %48 %24 %24 %24 %24
OpStore %187 %203
OpReturn
OpFunctionEnd
%198 = OpFunction %2 None %64
%197 = OpLabel
%195 = OpVariable %58 Function %196
OpBranch %199
%199 = OpLabel
%201 = OpAccessChain %200 %52 %27
%202 = OpAtomicLoad %4 %201 %6 %203
%205 = OpAccessChain %200 %52 %27
%204 = OpAtomicIAdd %4 %205 %6 %203 %22
OpStore %195 %204
%207 = OpAccessChain %200 %52 %27
%206 = OpAtomicISub %4 %207 %6 %203 %22
OpStore %195 %206
%209 = OpAccessChain %200 %52 %27
%208 = OpAtomicAnd %4 %209 %6 %203 %22
OpStore %195 %208
%211 = OpAccessChain %200 %52 %27
%210 = OpAtomicOr %4 %211 %6 %203 %22
OpStore %195 %210
%213 = OpAccessChain %200 %52 %27
%212 = OpAtomicXor %4 %213 %6 %203 %22
OpStore %195 %212
%215 = OpAccessChain %200 %52 %27
%214 = OpAtomicSMin %4 %215 %6 %203 %22
OpStore %195 %214
%217 = OpAccessChain %200 %52 %27
%216 = OpAtomicSMax %4 %217 %6 %203 %22
OpStore %195 %216
%219 = OpAccessChain %200 %52 %27
%218 = OpAtomicExchange %4 %219 %6 %203 %22
OpStore %195 %218
%220 = OpAccessChain %200 %52 %27
OpAtomicStore %220 %6 %203 %202
%207 = OpFunction %2 None %69
%206 = OpLabel
%204 = OpVariable %63 Function %205
OpBranch %208
%208 = OpLabel
%210 = OpAccessChain %209 %54 %27
%211 = OpAtomicLoad %4 %210 %6 %212
%214 = OpAccessChain %209 %54 %27
%213 = OpAtomicIAdd %4 %214 %6 %212 %22
OpStore %204 %213
%216 = OpAccessChain %209 %54 %27
%215 = OpAtomicISub %4 %216 %6 %212 %22
OpStore %204 %215
%218 = OpAccessChain %209 %54 %27
%217 = OpAtomicAnd %4 %218 %6 %212 %22
OpStore %204 %217
%220 = OpAccessChain %209 %54 %27
%219 = OpAtomicOr %4 %220 %6 %212 %22
OpStore %204 %219
%222 = OpAccessChain %209 %54 %27
%221 = OpAtomicXor %4 %222 %6 %212 %22
OpStore %204 %221
%224 = OpAccessChain %209 %54 %27
%223 = OpAtomicSMin %4 %224 %6 %212 %22
OpStore %204 %223
%226 = OpAccessChain %209 %54 %27
%225 = OpAtomicSMax %4 %226 %6 %212 %22
OpStore %204 %225
%228 = OpAccessChain %209 %54 %27
%227 = OpAtomicExchange %4 %228 %6 %212 %22
OpStore %204 %227
%229 = OpAccessChain %209 %54 %27
OpAtomicStore %229 %6 %212 %211
OpReturn
OpFunctionEnd

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

@@ -18,46 +18,48 @@ struct Baz {
var<storage, read_write> bar: Bar;
@group(0) @binding(1)
var<uniform> baz: Baz;
@group(0) @binding(2)
var<storage, read_write> qux: vec2<i32>;
fn test_matrix_within_struct_accesses() {
var idx: i32 = 9;
var t: Baz;
let _e4 = idx;
idx = (_e4 - 1);
let _e5 = idx;
idx = (_e5 - 1);
_ = baz.m;
_ = baz.m[0];
let _e14 = idx;
_ = baz.m[_e14];
let _e15 = idx;
_ = baz.m[_e15];
_ = baz.m[0][1];
let _e26 = idx;
_ = baz.m[0][_e26];
let _e30 = idx;
_ = baz.m[_e30][1];
let _e36 = idx;
let _e38 = idx;
_ = baz.m[_e36][_e38];
let _e27 = idx;
_ = baz.m[0][_e27];
let _e31 = idx;
_ = baz.m[_e31][1];
let _e37 = idx;
let _e39 = idx;
_ = baz.m[_e37][_e39];
t = Baz(mat3x2<f32>(vec2<f32>(1.0), vec2<f32>(2.0), vec2<f32>(3.0)));
let _e50 = idx;
idx = (_e50 + 1);
let _e51 = idx;
idx = (_e51 + 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 _e67 = idx;
t.m[_e67] = vec2<f32>(90.0);
let _e68 = idx;
t.m[_e68] = vec2<f32>(90.0);
t.m[0][1] = 10.0;
let _e80 = idx;
t.m[0][_e80] = 20.0;
let _e84 = idx;
t.m[_e84][1] = 30.0;
let _e90 = idx;
let _e92 = idx;
t.m[_e90][_e92] = 40.0;
let _e81 = idx;
t.m[0][_e81] = 20.0;
let _e85 = idx;
t.m[_e85][1] = 30.0;
let _e91 = idx;
let _e93 = idx;
t.m[_e91][_e93] = 40.0;
return;
}
fn read_from_private(foo_1: ptr<function, f32>) -> f32 {
let _e3 = (*foo_1);
return _e3;
let _e4 = (*foo_1);
return _e4;
}
fn test_arr_as_arg(a: array<array<f32,10>,5>) -> f32 {
@@ -76,12 +78,13 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let arr = bar.arr;
let b = bar._matrix[3][0];
let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
let c_1 = qux;
let data_pointer = (&bar.data[0].value);
let _e28 = read_from_private((&foo));
let _e30 = read_from_private((&foo));
c = array<i32,5>(a_1, i32(b), 3, 4, 5);
c[(vi + 1u)] = 42;
let value = c[vi];
let _e42 = test_arr_as_arg(array<array<f32,10>,5>(array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
let _e44 = test_arr_as_arg(array<array<f32,10>,5>(array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array<f32,10>(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0)));
return vec4<f32>((_matrix * vec4<f32>(vec4<i32>(value))), 2.0);
}
@@ -91,6 +94,7 @@ fn foo_frag() -> @location(0) vec4<f32> {
bar._matrix = mat4x3<f32>(vec3<f32>(0.0), vec3<f32>(1.0), vec3<f32>(2.0), vec3<f32>(3.0));
bar.arr = array<vec2<u32>,2>(vec2<u32>(0u), vec2<u32>(1u));
bar.data[1].value = 1;
qux = vec2<i32>(0, 0);
return vec4<f32>(0.0);
}
@@ -99,22 +103,22 @@ fn atomics() {
var tmp: i32;
let value_1 = atomicLoad((&bar.atom));
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;
let _e8 = atomicAdd((&bar.atom), 5);
tmp = _e8;
let _e11 = atomicSub((&bar.atom), 5);
tmp = _e11;
let _e14 = atomicAnd((&bar.atom), 5);
tmp = _e14;
let _e17 = atomicOr((&bar.atom), 5);
tmp = _e17;
let _e20 = atomicXor((&bar.atom), 5);
tmp = _e20;
let _e23 = atomicMin((&bar.atom), 5);
tmp = _e23;
let _e26 = atomicMax((&bar.atom), 5);
tmp = _e26;
let _e29 = atomicExchange((&bar.atom), 5);
tmp = _e29;
atomicStore((&bar.atom), value_1);
return;
}