From b3d1562515a9f382af61a7520a9c76da78154397 Mon Sep 17 00:00:00 2001 From: Hasan Ali Date: Tue, 10 May 2022 21:22:20 +0100 Subject: [PATCH] Fix hlsl output for writes to scalar/vector storage buffer --- src/back/hlsl/storage.rs | 3 + tests/in/access.param.ron | 2 + tests/in/access.wgsl | 5 + tests/out/glsl/access.atomics.Compute.glsl | 36 +- tests/out/glsl/access.foo_frag.Fragment.glsl | 7 +- tests/out/glsl/access.foo_vert.Vertex.glsl | 55 ++- tests/out/hlsl/access.hlsl | 87 ++-- tests/out/msl/access.msl | 107 ++-- tests/out/spv/access.spvasm | 495 ++++++++++--------- tests/out/wgsl/access.wgsl | 88 ++-- 10 files changed, 464 insertions(+), 421 deletions(-) diff --git a/src/back/hlsl/storage.rs b/src/back/hlsl/storage.rs index 658090256e..637b05d4bb 100644 --- a/src/back/hlsl/storage.rs +++ b/src/back/hlsl/storage.rs @@ -44,6 +44,9 @@ impl 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, "+")?; diff --git a/tests/in/access.param.ron b/tests/in/access.param.ron index 57cb7a632d..3e79bea26c 100644 --- a/tests/in/access.param.ron +++ b/tests/in/access.param.ron @@ -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), ), diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index c80d0349d3..ee426876d2 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -22,6 +22,9 @@ struct Baz { @group(0) @binding(1) var baz: Baz; +@group(0) @binding(2) +var qux: vec2; + fn test_matrix_within_struct_accesses() { var idx = 9; @@ -73,6 +76,7 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { 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 = &bar.data[0].value; @@ -95,6 +99,7 @@ fn foo_frag() -> @location(0) vec4 { bar._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0)); bar.arr = array, 2>(vec2(0u), vec2(1u)); bar.data[1].value = 1; + qux = vec2(); return vec4(0.0); } diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index 1a8668a820..2ab1b4fa8e 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -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; } diff --git a/tests/out/glsl/access.foo_frag.Fragment.glsl b/tests/out/glsl/access.foo_frag.Fragment.glsl index df6ac6770a..fcd3e43007 100644 --- a/tests/out/glsl/access.foo_frag.Fragment.glsl +++ b/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -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; } diff --git a/tests/out/glsl/access.foo_vert.Vertex.glsl b/tests/out/glsl/access.foo_vert.Vertex.glsl index b5b4804d76..5eead818e5 100644 --- a/tests/out/glsl/access.foo_vert.Vertex.glsl +++ b/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -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; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 89dd2ea61c..4bb113716f 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -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; } diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index 7ee2d6c132..f3f11559c2 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -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(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(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::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; } diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index 68f296d425..21e92ab2b9 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -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 +%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 -%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 +%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 \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 2b88d7aac1..752e3853d5 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -18,46 +18,48 @@ struct Baz { var bar: Bar; @group(0) @binding(1) var baz: Baz; +@group(0) @binding(2) +var qux: vec2; 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(vec2(1.0), vec2(2.0), vec2(3.0))); - let _e50 = idx; - idx = (_e50 + 1); + let _e51 = idx; + idx = (_e51 + 1); t.m = mat3x2(vec2(6.0), vec2(5.0), vec2(4.0)); t.m[0] = vec2(9.0); - let _e67 = idx; - t.m[_e67] = vec2(90.0); + let _e68 = idx; + t.m[_e68] = vec2(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) -> f32 { - let _e3 = (*foo_1); - return _e3; + let _e4 = (*foo_1); + return _e4; } fn test_arr_as_arg(a: array,5>) -> f32 { @@ -76,12 +78,13 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { 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(a_1, i32(b), 3, 4, 5); c[(vi + 1u)] = 42; let value = c[vi]; - let _e42 = test_arr_as_arg(array,5>(array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(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,5>(array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0), array(0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0))); return vec4((_matrix * vec4(vec4(value))), 2.0); } @@ -91,6 +94,7 @@ fn foo_frag() -> @location(0) vec4 { bar._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0)); bar.arr = array,2>(vec2(0u), vec2(1u)); bar.data[1].value = 1; + qux = vec2(0, 0); return vec4(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; }