From 79845371d3916b08689da374533c280246f6efc0 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Sun, 6 Mar 2022 09:08:40 -0800 Subject: [PATCH] msl: qualify read-only storage with const --- src/back/msl/writer.rs | 39 ++-- tests/in/access.param.ron | 6 + tests/in/access.wgsl | 19 +- tests/in/globals.wgsl | 3 +- tests/out/glsl/access.atomics.Compute.glsl | 4 +- tests/out/glsl/access.foo_frag.Fragment.glsl | 32 +++ ...ertex.glsl => access.foo_vert.Vertex.glsl} | 16 +- tests/out/glsl/globals.main.Compute.glsl | 3 +- tests/out/hlsl/access.hlsl | 33 ++-- tests/out/hlsl/access.hlsl.config | 4 +- tests/out/hlsl/globals.hlsl | 3 +- tests/out/msl/access.msl | 39 ++-- tests/out/msl/boids.msl | 2 +- tests/out/msl/bounds-check-restrict.msl | 14 +- tests/out/msl/bounds-check-zero.msl | 14 +- tests/out/msl/globals.msl | 3 +- tests/out/msl/policy-mix.msl | 2 +- tests/out/msl/shadow.msl | 2 +- tests/out/spv/access.spvasm | 183 ++++++++++-------- tests/out/spv/globals.spvasm | 166 ++++++++-------- tests/out/wgsl/access.wgsl | 27 +-- tests/out/wgsl/globals.wgsl | 3 +- 22 files changed, 354 insertions(+), 263 deletions(-) create mode 100644 tests/out/glsl/access.foo_frag.Fragment.glsl rename tests/out/glsl/{access.foo.Vertex.glsl => access.foo_vert.Vertex.glsl} (67%) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index cd86d98834..c04abe01a9 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -238,13 +238,12 @@ impl<'a> TypedGlobalVariable<'a> { let (space, access, reference) = match var.space.to_msl_name() { Some(space) if self.reference => { - let access = match var.space { - crate::AddressSpace::Private | crate::AddressSpace::WorkGroup - if !self.usage.contains(valid::GlobalUse::WRITE) => - { - "const" - } - _ => "", + let access = if var.space.needs_access_qualifier() + && !self.usage.contains(valid::GlobalUse::WRITE) + { + "const" + } else { + "" }; (space, access, "&") } @@ -401,13 +400,25 @@ impl crate::AddressSpace { /// passed through any functions called from the entry point. fn needs_pass_through(&self) -> bool { match *self { - crate::AddressSpace::Uniform - | crate::AddressSpace::Storage { .. } - | crate::AddressSpace::Private - | crate::AddressSpace::WorkGroup - | crate::AddressSpace::PushConstant - | crate::AddressSpace::Handle => true, - crate::AddressSpace::Function => false, + Self::Uniform + | Self::Storage { .. } + | Self::Private + | Self::WorkGroup + | Self::PushConstant + | Self::Handle => true, + Self::Function => false, + } + } + + /// Returns true if the address space may need a "const" qualifier. + fn needs_access_qualifier(&self) -> bool { + match *self { + //Note: we are ignoring the storage access here, and instead + // rely on the actual use of a global by functions. This means we + // may end up with "const" even if the binding is read-write, + // and that should be OK. + Self::Storage { .. } | Self::Private | Self::WorkGroup => true, + Self::Uniform | Self::PushConstant | Self::Handle | Self::Function => false, } } diff --git a/tests/in/access.param.ron b/tests/in/access.param.ron index 3e9214f474..84595754e4 100644 --- a/tests/in/access.param.ron +++ b/tests/in/access.param.ron @@ -8,6 +8,12 @@ lang_version: (2, 0), per_stage_map: ( vs: ( + resources: { + (group: 0, binding: 0): (buffer: Some(0), mutable: false), + }, + sizes_buffer: Some(24), + ), + fs: ( resources: { (group: 0, binding: 0): (buffer: Some(0), mutable: true), }, diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index 90c707359c..d68a74ca0e 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -20,7 +20,7 @@ fn read_from_private(foo: ptr) -> f32 { } @stage(vertex) -fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { +fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { var foo: f32 = 0.0; // We should check that backed doesn't skip this expression let baz: f32 = foo; @@ -37,12 +37,6 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { let data_pointer: ptr = &bar.data[0].value; let foo_value = read_from_private(&foo); - // test storage stores - bar.matrix[1].z = 1.0; - bar.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); - bar.arr = array, 2>(vec2(0u), vec2(1u)); - bar.data[1].value = 1; - // test array indexing var c = array(a, i32(b), 3, 4, 5); c[vi + 1u] = 42; @@ -51,6 +45,17 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { return matrix * vec4(vec4(value)); } +@stage(fragment) +fn foo_frag() -> @location(0) vec4 { + // test storage stores + bar.matrix[1].z = 1.0; + bar.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); + bar.arr = array, 2>(vec2(0u), vec2(1u)); + bar.data[1].value = 1; + + return vec4(0.0); +} + @stage(compute) @workgroup_size(1) fn atomics() { var tmp: i32; diff --git a/tests/in/globals.wgsl b/tests/in/globals.wgsl index e9f3fa77eb..7e5b725016 100644 --- a/tests/in/globals.wgsl +++ b/tests/in/globals.wgsl @@ -11,7 +11,7 @@ struct Foo { v1: f32; }; @group(0) @binding(1) -var alignment: Foo; +var alignment: Foo; @group(0) @binding(2) var dummy: array>; @@ -23,6 +23,7 @@ var float_vecs: array, 20>; fn main() { wg[3] = alignment.v1; wg[2] = alignment.v3.x; + alignment.v1 = 4.0; wg[1] = f32(arrayLength(&dummy)); atomicStore(&at, 2u); diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index 632fa26666..d1fddc7e11 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -17,8 +17,8 @@ layout(std430) buffer Bar_block_0Compute { } _group_0_binding_0_cs; -float read_from_private(inout float foo_2) { - float _e2 = foo_2; +float read_from_private(inout float foo_1) { + float _e2 = foo_1; return _e2; } diff --git a/tests/out/glsl/access.foo_frag.Fragment.glsl b/tests/out/glsl/access.foo_frag.Fragment.glsl new file mode 100644 index 0000000000..ab14704b0b --- /dev/null +++ b/tests/out/glsl/access.foo_frag.Fragment.glsl @@ -0,0 +1,32 @@ +#version 310 es + +precision highp float; +precision highp int; + +struct AlignedWrapper { + int value; +}; +layout(std430) buffer Bar_block_0Fragment { + mat4x4 matrix; + mat2x2 matrix_array[2]; + int atom; + uvec2 arr[2]; + AlignedWrapper data[]; +} _group_0_binding_0_fs; + +layout(location = 0) out vec4 _fs2p_location0; + +float read_from_private(inout float foo_1) { + float _e2 = foo_1; + return _e2; +} + +void main() { + _group_0_binding_0_fs.matrix[1][2] = 1.0; + _group_0_binding_0_fs.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); + _group_0_binding_0_fs.arr = uvec2[2](uvec2(0u), uvec2(1u)); + _group_0_binding_0_fs.data[1].value = 1; + _fs2p_location0 = vec4(0.0); + return; +} + diff --git a/tests/out/glsl/access.foo.Vertex.glsl b/tests/out/glsl/access.foo_vert.Vertex.glsl similarity index 67% rename from tests/out/glsl/access.foo.Vertex.glsl rename to tests/out/glsl/access.foo_vert.Vertex.glsl index 3508015cf1..05fd493214 100644 --- a/tests/out/glsl/access.foo.Vertex.glsl +++ b/tests/out/glsl/access.foo_vert.Vertex.glsl @@ -15,26 +15,22 @@ layout(std430) buffer Bar_block_0Vertex { } _group_0_binding_0_vs; -float read_from_private(inout float foo_2) { - float _e2 = foo_2; +float read_from_private(inout float foo_1) { + float _e2 = foo_1; return _e2; } void main() { uint vi = uint(gl_VertexID); - float foo_1 = 0.0; + float foo = 0.0; int c[5] = int[5](0, 0, 0, 0, 0); - float baz = foo_1; - foo_1 = 1.0; + float baz = foo; + foo = 1.0; mat4x4 matrix = _group_0_binding_0_vs.matrix; uvec2 arr[2] = _group_0_binding_0_vs.arr; float b = _group_0_binding_0_vs.matrix[3][0]; int a = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value; - float _e27 = read_from_private(foo_1); - _group_0_binding_0_vs.matrix[1][2] = 1.0; - _group_0_binding_0_vs.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); - _group_0_binding_0_vs.arr = uvec2[2](uvec2(0u), uvec2(1u)); - _group_0_binding_0_vs.data[1].value = 1; + float _e27 = read_from_private(foo); c = int[5](a, int(b), 3, 4, 5); c[(vi + 1u)] = 42; int value = c[vi]; diff --git a/tests/out/glsl/globals.main.Compute.glsl b/tests/out/glsl/globals.main.Compute.glsl index 46150907be..61513e8af6 100644 --- a/tests/out/glsl/globals.main.Compute.glsl +++ b/tests/out/glsl/globals.main.Compute.glsl @@ -13,7 +13,7 @@ shared float wg[10]; shared uint at_1; -layout(std430) readonly buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; }; +layout(std430) buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; }; layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; }; @@ -25,6 +25,7 @@ void main() { wg[3] = _e9; float _e14 = _group_0_binding_1_cs.v3_.x; wg[2] = _e14; + _group_0_binding_1_cs.v1_ = 4.0; wg[1] = float(uint(_group_0_binding_2_cs.length())); at_1 = 2u; return; diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index f79a6798bc..91ffd5293a 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -5,9 +5,9 @@ struct AlignedWrapper { RWByteAddressBuffer bar : register(u0); -float read_from_private(inout float foo_2) +float read_from_private(inout float foo_1) { - float _expr2 = foo_2; + float _expr2 = foo_1; return _expr2; } @@ -18,18 +18,29 @@ uint NagaBufferLengthRW(RWByteAddressBuffer buffer) return ret; } -float4 foo(uint vi : SV_VertexID) : SV_Position +float4 foo_vert(uint vi : SV_VertexID) : SV_Position { - float foo_1 = 0.0; + float foo = 0.0; int c[5] = {(int)0,(int)0,(int)0,(int)0,(int)0}; - float baz = foo_1; - foo_1 = 1.0; + float baz = foo; + foo = 1.0; float4x4 matrix_ = float4x4(asfloat(bar.Load4(0+0)), asfloat(bar.Load4(0+16)), asfloat(bar.Load4(0+32)), asfloat(bar.Load4(0+48))); uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))}; float b = asfloat(bar.Load(0+48+0)); int a = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120)); - const float _e27 = read_from_private(foo_1); + const float _e27 = read_from_private(foo); + { + int _result[5]={ a, int(b), 3, 4, 5 }; + for(int _i=0; _i<5; ++_i) c[_i] = _result[_i]; + } + c[(vi + 1u)] = 42; + int value = c[vi]; + return mul(float4(int4(value.xxxx)), matrix_); +} + +float4 foo_frag() : SV_Target0 +{ bar.Store(8+16+0, asuint(1.0)); { float4x4 _value2 = float4x4(float4(0.0.xxxx), float4(1.0.xxxx), float4(2.0.xxxx), float4(3.0.xxxx)); @@ -44,13 +55,7 @@ float4 foo(uint vi : SV_VertexID) : SV_Position bar.Store2(104+8, asuint(_value2[1])); } bar.Store(0+8+120, asuint(1)); - { - int _result[5]={ a, int(b), 3, 4, 5 }; - for(int _i=0; _i<5; ++_i) c[_i] = _result[_i]; - } - c[(vi + 1u)] = 42; - int value = c[vi]; - return mul(float4(int4(value.xxxx)), matrix_); + return float4(0.0.xxxx); } [numthreads(1, 1, 1)] diff --git a/tests/out/hlsl/access.hlsl.config b/tests/out/hlsl/access.hlsl.config index 49dcb6821d..f6a8845ad5 100644 --- a/tests/out/hlsl/access.hlsl.config +++ b/tests/out/hlsl/access.hlsl.config @@ -1,3 +1,3 @@ -vertex=(foo:vs_5_1 ) -fragment=() +vertex=(foo_vert:vs_5_1 ) +fragment=(foo_frag:ps_5_1 ) compute=(atomics:cs_5_1 ) diff --git a/tests/out/hlsl/globals.hlsl b/tests/out/hlsl/globals.hlsl index 515ff1de6b..d540235bab 100644 --- a/tests/out/hlsl/globals.hlsl +++ b/tests/out/hlsl/globals.hlsl @@ -7,7 +7,7 @@ struct Foo { groupshared float wg[10]; groupshared uint at_1; -ByteAddressBuffer alignment : register(t1); +RWByteAddressBuffer alignment : register(u1); ByteAddressBuffer dummy : register(t2); cbuffer float_vecs : register(b3) { float4 float_vecs[20]; } @@ -28,6 +28,7 @@ void main() wg[3] = _expr9; float _expr14 = asfloat(alignment.Load(0+0)); wg[2] = _expr14; + alignment.Store(12, asuint(4.0)); wg[1] = float(((NagaBufferLength(dummy) - 0) / 8)); at_1 = 2u; return; diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index d4d80a50b6..0a1c0cfb2a 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -31,39 +31,50 @@ struct type_13 { }; float read_from_private( - thread float& foo_2 + thread float& foo_1 ) { - float _e2 = foo_2; + float _e2 = foo_1; return _e2; } -struct fooInput { +struct foo_vertInput { }; -struct fooOutput { +struct foo_vertOutput { metal::float4 member [[position]]; }; -vertex fooOutput foo( +vertex foo_vertOutput foo_vert( uint vi [[vertex_id]] -, device Bar& bar [[buffer(0)]] +, device Bar const& bar [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { - float foo_1 = 0.0; + float foo = 0.0; type_13 c; - float baz = foo_1; - foo_1 = 1.0; + float baz = foo; + foo = 1.0; metal::float4x4 matrix = bar.matrix; type_6 arr = bar.arr; float b = bar.matrix[3].x; int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value; - float _e27 = read_from_private(foo_1); + float _e27 = read_from_private(foo); + for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast(b), 3, 4, 5}.inner[_i]; + c.inner[vi + 1u] = 42; + int value = c.inner[vi]; + return foo_vertOutput { matrix * static_cast(metal::int4(value)) }; +} + + +struct foo_fragOutput { + metal::float4 member_1 [[color(0)]]; +}; +fragment foo_fragOutput foo_frag( + device Bar& bar [[buffer(0)]] +, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] +) { bar.matrix[1].z = 1.0; bar.matrix = metal::float4x4(metal::float4(0.0), metal::float4(1.0), metal::float4(2.0), metal::float4(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; - for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast(b), 3, 4, 5}.inner[_i]; - c.inner[vi + 1u] = 42; - int value = c.inner[vi]; - return fooOutput { matrix * static_cast(metal::int4(value)) }; + return foo_fragOutput { metal::float4(0.0) }; } diff --git a/tests/out/msl/boids.msl b/tests/out/msl/boids.msl index d18d733ac7..bcd80ca056 100644 --- a/tests/out/msl/boids.msl +++ b/tests/out/msl/boids.msl @@ -33,7 +33,7 @@ struct main_Input { kernel void main_( metal::uint3 global_invocation_id [[thread_position_in_grid]] , constant SimParams& params [[buffer(0)]] -, device Particles& particlesSrc [[buffer(1)]] +, device Particles const& particlesSrc [[buffer(1)]] , device Particles& particlesDst [[buffer(2)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(3)]] ) { diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl index e82ebf814d..a614e59240 100644 --- a/tests/out/msl/bounds-check-restrict.msl +++ b/tests/out/msl/bounds-check-restrict.msl @@ -22,7 +22,7 @@ struct Globals { float index_array( int i, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = globals.a.inner[metal::min(unsigned(i), 9u)]; @@ -31,7 +31,7 @@ float index_array( float index_dynamic_array( int i_1, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = globals.d[metal::min(unsigned(i_1), (_buffer_sizes.size0 - 112 - 4) / 4)]; @@ -40,7 +40,7 @@ float index_dynamic_array( float index_vector( int i_2, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = globals.v[metal::min(unsigned(i_2), 3u)]; @@ -56,7 +56,7 @@ float index_vector_by_value( metal::float4 index_matrix( int i_4, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { metal::float4 _e4 = globals.m[metal::min(unsigned(i_4), 2u)]; @@ -66,7 +66,7 @@ metal::float4 index_matrix( float index_twice( int i_5, int j, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e6 = globals.m[metal::min(unsigned(i_5), 2u)][metal::min(unsigned(j), 3u)]; @@ -75,7 +75,7 @@ float index_twice( float index_expensive( int i_6, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e11 = globals.a.inner[metal::min(unsigned(static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0)), 9u)]; @@ -83,7 +83,7 @@ float index_expensive( } float index_in_bounds( - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = globals.a.inner[9]; diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl index 1a6bcc1578..b45550d5ae 100644 --- a/tests/out/msl/bounds-check-zero.msl +++ b/tests/out/msl/bounds-check-zero.msl @@ -28,7 +28,7 @@ struct Globals { float index_array( int i, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = uint(i) < 10 ? globals.a.inner[i] : DefaultConstructible(); @@ -37,7 +37,7 @@ float index_array( float index_dynamic_array( int i_1, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = uint(i_1) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4 ? globals.d[i_1] : DefaultConstructible(); @@ -46,7 +46,7 @@ float index_dynamic_array( float index_vector( int i_2, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = uint(i_2) < 4 ? globals.v[i_2] : DefaultConstructible(); @@ -62,7 +62,7 @@ float index_vector_by_value( metal::float4 index_matrix( int i_4, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { metal::float4 _e4 = uint(i_4) < 3 ? globals.m[i_4] : DefaultConstructible(); @@ -72,7 +72,7 @@ metal::float4 index_matrix( float index_twice( int i_5, int j, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e6 = uint(j) < 4 && uint(i_5) < 3 ? globals.m[i_5][j] : DefaultConstructible(); @@ -81,7 +81,7 @@ float index_twice( float index_expensive( int i_6, - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { int _e9 = static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0); @@ -90,7 +90,7 @@ float index_expensive( } float index_in_bounds( - device Globals& globals, + device Globals const& globals, constant _mslBufferSizes& _buffer_sizes ) { float _e4 = globals.a.inner[9]; diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index 8f4ce07ade..14fbc152a4 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -25,7 +25,7 @@ kernel void main_( threadgroup type_2& wg , threadgroup metal::atomic_uint& at_1 , device Foo& alignment [[user(fake0)]] -, device type_6& dummy [[user(fake0)]] +, device type_6 const& dummy [[user(fake0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { float Foo_1 = 1.0; @@ -34,6 +34,7 @@ kernel void main_( wg.inner[3] = _e9; float _e14 = metal::float3(alignment.v3_).x; wg.inner[2] = _e14; + alignment.v1_ = 4.0; wg.inner[1] = static_cast(1 + (_buffer_sizes.size3 - 0 - 8) / 8); metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed); return; diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl index 8b4ba9911d..b6b9b942bd 100644 --- a/tests/out/msl/policy-mix.msl +++ b/tests/out/msl/policy-mix.msl @@ -36,7 +36,7 @@ metal::float4 mock_function( metal::int2 c, int i, int l, - device InStorage& in_storage, + device InStorage const& in_storage, constant InUniform& in_uniform, metal::texture2d_array image_2d_array, threadgroup type_5 const& in_workgroup, diff --git a/tests/out/msl/shadow.msl b/tests/out/msl/shadow.msl index 58d03e957d..5bdd3ba977 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -48,7 +48,7 @@ struct fs_mainOutput { fragment fs_mainOutput fs_main( fs_mainInput varyings [[stage_in]] , constant Globals& u_globals [[user(fake0)]] -, device Lights& s_lights [[user(fake0)]] +, device Lights const& s_lights [[user(fake0)]] , metal::depth2d_array t_shadow [[user(fake0)]] , metal::sampler sampler_shadow [[user(fake0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index bdd6160406..21e4f513e6 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,14 +1,16 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 121 +; Bound: 126 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %51 "foo" %46 %49 -OpEntryPoint GLCompute %98 "atomics" -OpExecutionMode %98 LocalSize 1 1 1 +OpEntryPoint Vertex %51 "foo_vert" %46 %49 +OpEntryPoint Fragment %86 "foo_frag" %85 +OpEntryPoint GLCompute %103 "atomics" +OpExecutionMode %86 OriginUpperLeft +OpExecutionMode %103 LocalSize 1 1 1 OpSource GLSL 450 OpMemberName %21 0 "value" OpName %21 "AlignedWrapper" @@ -24,9 +26,10 @@ OpName %38 "read_from_private" OpName %42 "foo" OpName %43 "c" OpName %46 "vi" -OpName %51 "foo" -OpName %96 "tmp" -OpName %98 "atomics" +OpName %51 "foo_vert" +OpName %86 "foo_frag" +OpName %101 "tmp" +OpName %103 "atomics" OpMemberDecorate %21 0 Offset 0 OpDecorate %26 ArrayStride 16 OpDecorate %28 ArrayStride 8 @@ -46,6 +49,7 @@ OpDecorate %34 Binding 0 OpDecorate %30 Block OpDecorate %46 BuiltIn VertexIndex OpDecorate %49 BuiltIn Position +OpDecorate %85 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -56,15 +60,15 @@ OpDecorate %49 BuiltIn Position %8 = OpConstant %9 3 %10 = OpConstant %9 2 %11 = OpConstant %4 0 -%12 = OpConstant %4 1 -%13 = OpConstant %6 2.0 -%14 = OpConstant %6 3.0 -%15 = OpConstant %9 0 -%16 = OpConstant %9 1 -%17 = OpConstant %4 5 -%18 = OpConstant %4 3 -%19 = OpConstant %4 4 -%20 = OpConstant %4 42 +%12 = OpConstant %4 5 +%13 = OpConstant %4 3 +%14 = OpConstant %4 4 +%15 = OpConstant %9 1 +%16 = OpConstant %4 42 +%17 = OpConstant %4 1 +%18 = OpConstant %6 2.0 +%19 = OpConstant %6 3.0 +%20 = OpConstant %9 0 %21 = OpTypeStruct %4 %23 = OpTypeVector %6 4 %22 = OpTypeMatrix %23 4 @@ -77,7 +81,7 @@ OpDecorate %49 BuiltIn Position %30 = OpTypeStruct %22 %26 %4 %28 %29 %31 = OpTypePointer Function %6 %32 = OpTypePointer StorageBuffer %4 -%33 = OpTypeArray %4 %17 +%33 = OpTypeArray %4 %12 %35 = OpTypePointer StorageBuffer %30 %34 = OpVariable %35 StorageBuffer %39 = OpTypeFunction %6 %31 @@ -94,10 +98,11 @@ OpDecorate %49 BuiltIn Position %65 = OpTypePointer StorageBuffer %29 %68 = OpTypePointer StorageBuffer %21 %69 = OpConstant %9 4 -%88 = OpTypePointer Function %4 -%92 = OpTypeVector %4 4 -%100 = OpTypePointer StorageBuffer %4 -%103 = OpConstant %9 64 +%76 = OpTypePointer Function %4 +%80 = OpTypeVector %4 4 +%85 = OpVariable %50 Output +%105 = OpTypePointer StorageBuffer %4 +%108 = OpConstant %9 64 %38 = OpFunction %6 None %39 %37 = OpFunctionParameter %31 %36 = OpLabel @@ -115,79 +120,87 @@ OpBranch %53 %53 = OpLabel %54 = OpLoad %6 %42 OpStore %42 %7 -%56 = OpAccessChain %55 %34 %15 +%56 = OpAccessChain %55 %34 %20 %57 = OpLoad %22 %56 %59 = OpAccessChain %58 %34 %8 %60 = OpLoad %28 %59 -%63 = OpAccessChain %62 %34 %15 %8 %15 +%63 = OpAccessChain %62 %34 %20 %8 %20 %64 = OpLoad %6 %63 %66 = OpArrayLength %9 %34 4 %67 = OpISub %9 %66 %10 -%70 = OpAccessChain %32 %34 %69 %67 %15 +%70 = OpAccessChain %32 %34 %69 %67 %20 %71 = OpLoad %4 %70 %72 = OpFunctionCall %6 %38 %42 -%73 = OpAccessChain %62 %34 %15 %16 %10 -OpStore %73 %7 -%74 = OpCompositeConstruct %23 %5 %5 %5 %5 -%75 = OpCompositeConstruct %23 %7 %7 %7 %7 -%76 = OpCompositeConstruct %23 %13 %13 %13 %13 -%77 = OpCompositeConstruct %23 %14 %14 %14 %14 -%78 = OpCompositeConstruct %22 %74 %75 %76 %77 -%79 = OpAccessChain %55 %34 %15 -OpStore %79 %78 -%80 = OpCompositeConstruct %27 %15 %15 -%81 = OpCompositeConstruct %27 %16 %16 -%82 = OpCompositeConstruct %28 %80 %81 -%83 = OpAccessChain %58 %34 %8 -OpStore %83 %82 -%84 = OpAccessChain %32 %34 %69 %16 %15 -OpStore %84 %12 -%85 = OpConvertFToS %4 %64 -%86 = OpCompositeConstruct %33 %71 %85 %18 %19 %17 -OpStore %43 %86 -%87 = OpIAdd %9 %48 %16 -%89 = OpAccessChain %88 %43 %87 -OpStore %89 %20 -%90 = OpAccessChain %88 %43 %48 -%91 = OpLoad %4 %90 -%93 = OpCompositeConstruct %92 %91 %91 %91 %91 -%94 = OpConvertSToF %23 %93 -%95 = OpMatrixTimesVector %23 %57 %94 -OpStore %49 %95 +%73 = OpConvertFToS %4 %64 +%74 = OpCompositeConstruct %33 %71 %73 %13 %14 %12 +OpStore %43 %74 +%75 = OpIAdd %9 %48 %15 +%77 = OpAccessChain %76 %43 %75 +OpStore %77 %16 +%78 = OpAccessChain %76 %43 %48 +%79 = OpLoad %4 %78 +%81 = OpCompositeConstruct %80 %79 %79 %79 %79 +%82 = OpConvertSToF %23 %81 +%83 = OpMatrixTimesVector %23 %57 %82 +OpStore %49 %83 OpReturn OpFunctionEnd -%98 = OpFunction %2 None %52 -%97 = OpLabel -%96 = OpVariable %88 Function -OpBranch %99 -%99 = OpLabel -%101 = OpAccessChain %100 %34 %10 -%102 = OpAtomicLoad %4 %101 %12 %103 -%105 = OpAccessChain %100 %34 %10 -%104 = OpAtomicIAdd %4 %105 %12 %103 %17 -OpStore %96 %104 -%107 = OpAccessChain %100 %34 %10 -%106 = OpAtomicISub %4 %107 %12 %103 %17 -OpStore %96 %106 -%109 = OpAccessChain %100 %34 %10 -%108 = OpAtomicAnd %4 %109 %12 %103 %17 -OpStore %96 %108 -%111 = OpAccessChain %100 %34 %10 -%110 = OpAtomicOr %4 %111 %12 %103 %17 -OpStore %96 %110 -%113 = OpAccessChain %100 %34 %10 -%112 = OpAtomicXor %4 %113 %12 %103 %17 -OpStore %96 %112 -%115 = OpAccessChain %100 %34 %10 -%114 = OpAtomicSMin %4 %115 %12 %103 %17 -OpStore %96 %114 -%117 = OpAccessChain %100 %34 %10 -%116 = OpAtomicSMax %4 %117 %12 %103 %17 -OpStore %96 %116 -%119 = OpAccessChain %100 %34 %10 -%118 = OpAtomicExchange %4 %119 %12 %103 %17 -OpStore %96 %118 -%120 = OpAccessChain %100 %34 %10 -OpAtomicStore %120 %12 %103 %102 +%86 = OpFunction %2 None %52 +%84 = OpLabel +OpBranch %87 +%87 = OpLabel +%88 = OpAccessChain %62 %34 %20 %15 %10 +OpStore %88 %7 +%89 = OpCompositeConstruct %23 %5 %5 %5 %5 +%90 = OpCompositeConstruct %23 %7 %7 %7 %7 +%91 = OpCompositeConstruct %23 %18 %18 %18 %18 +%92 = OpCompositeConstruct %23 %19 %19 %19 %19 +%93 = OpCompositeConstruct %22 %89 %90 %91 %92 +%94 = OpAccessChain %55 %34 %20 +OpStore %94 %93 +%95 = OpCompositeConstruct %27 %20 %20 +%96 = OpCompositeConstruct %27 %15 %15 +%97 = OpCompositeConstruct %28 %95 %96 +%98 = OpAccessChain %58 %34 %8 +OpStore %98 %97 +%99 = OpAccessChain %32 %34 %69 %15 %20 +OpStore %99 %17 +%100 = OpCompositeConstruct %23 %5 %5 %5 %5 +OpStore %85 %100 +OpReturn +OpFunctionEnd +%103 = OpFunction %2 None %52 +%102 = OpLabel +%101 = OpVariable %76 Function +OpBranch %104 +%104 = OpLabel +%106 = OpAccessChain %105 %34 %10 +%107 = OpAtomicLoad %4 %106 %17 %108 +%110 = OpAccessChain %105 %34 %10 +%109 = OpAtomicIAdd %4 %110 %17 %108 %12 +OpStore %101 %109 +%112 = OpAccessChain %105 %34 %10 +%111 = OpAtomicISub %4 %112 %17 %108 %12 +OpStore %101 %111 +%114 = OpAccessChain %105 %34 %10 +%113 = OpAtomicAnd %4 %114 %17 %108 %12 +OpStore %101 %113 +%116 = OpAccessChain %105 %34 %10 +%115 = OpAtomicOr %4 %116 %17 %108 %12 +OpStore %101 %115 +%118 = OpAccessChain %105 %34 %10 +%117 = OpAtomicXor %4 %118 %17 %108 %12 +OpStore %101 %117 +%120 = OpAccessChain %105 %34 %10 +%119 = OpAtomicSMin %4 %120 %17 %108 %12 +OpStore %101 %119 +%122 = OpAccessChain %105 %34 %10 +%121 = OpAtomicSMax %4 %122 %17 %108 %12 +OpStore %101 %121 +%124 = OpAccessChain %105 %34 %10 +%123 = OpAtomicExchange %4 %124 %17 %108 %12 +OpStore %101 %123 +%125 = OpAccessChain %105 %34 %10 +OpAtomicStore %125 %17 %108 %107 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/globals.spvasm b/tests/out/spv/globals.spvasm index bfba6116fa..e70ef7ca3a 100644 --- a/tests/out/spv/globals.spvasm +++ b/tests/out/spv/globals.spvasm @@ -1,32 +1,31 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 66 +; Bound: 68 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %41 "main" -OpExecutionMode %41 LocalSize 1 1 1 -OpDecorate %16 ArrayStride 4 -OpMemberDecorate %18 0 Offset 0 -OpMemberDecorate %18 1 Offset 12 -OpDecorate %20 ArrayStride 8 -OpDecorate %22 ArrayStride 16 -OpDecorate %27 NonWritable -OpDecorate %27 DescriptorSet 0 -OpDecorate %27 Binding 1 -OpDecorate %28 Block -OpMemberDecorate %28 0 Offset 0 -OpDecorate %30 NonWritable -OpDecorate %30 DescriptorSet 0 -OpDecorate %30 Binding 2 -OpDecorate %31 Block -OpMemberDecorate %31 0 Offset 0 -OpDecorate %33 DescriptorSet 0 -OpDecorate %33 Binding 3 -OpDecorate %34 Block -OpMemberDecorate %34 0 Offset 0 +OpEntryPoint GLCompute %42 "main" +OpExecutionMode %42 LocalSize 1 1 1 +OpDecorate %17 ArrayStride 4 +OpMemberDecorate %19 0 Offset 0 +OpMemberDecorate %19 1 Offset 12 +OpDecorate %21 ArrayStride 8 +OpDecorate %23 ArrayStride 16 +OpDecorate %28 DescriptorSet 0 +OpDecorate %28 Binding 1 +OpDecorate %29 Block +OpMemberDecorate %29 0 Offset 0 +OpDecorate %31 NonWritable +OpDecorate %31 DescriptorSet 0 +OpDecorate %31 Binding 2 +OpDecorate %32 Block +OpMemberDecorate %32 0 Offset 0 +OpDecorate %34 DescriptorSet 0 +OpDecorate %34 Binding 3 +OpDecorate %35 Block +OpMemberDecorate %35 0 Offset 0 %2 = OpTypeVoid %4 = OpTypeBool %3 = OpConstantTrue %4 @@ -36,65 +35,68 @@ OpMemberDecorate %34 0 Offset 0 %7 = OpConstant %8 20 %9 = OpConstant %8 3 %10 = OpConstant %8 2 -%11 = OpConstant %8 1 -%12 = OpConstant %6 2 -%14 = OpTypeFloat 32 -%13 = OpConstant %14 1.0 -%15 = OpConstantTrue %4 -%16 = OpTypeArray %14 %5 -%17 = OpTypeVector %14 3 -%18 = OpTypeStruct %17 %14 -%19 = OpTypeVector %14 2 -%20 = OpTypeRuntimeArray %19 -%21 = OpTypeVector %14 4 -%22 = OpTypeArray %21 %7 -%24 = OpTypePointer Workgroup %16 -%23 = OpVariable %24 Workgroup -%26 = OpTypePointer Workgroup %6 -%25 = OpVariable %26 Workgroup -%28 = OpTypeStruct %18 -%29 = OpTypePointer StorageBuffer %28 -%27 = OpVariable %29 StorageBuffer -%31 = OpTypeStruct %20 -%32 = OpTypePointer StorageBuffer %31 -%30 = OpVariable %32 StorageBuffer -%34 = OpTypeStruct %22 -%35 = OpTypePointer Uniform %34 -%33 = OpVariable %35 Uniform -%37 = OpTypePointer Function %14 -%39 = OpTypePointer Function %4 -%42 = OpTypeFunction %2 -%43 = OpTypePointer StorageBuffer %18 -%44 = OpConstant %6 0 -%46 = OpTypePointer StorageBuffer %20 -%48 = OpTypePointer Uniform %22 -%50 = OpTypePointer Workgroup %14 -%51 = OpTypePointer StorageBuffer %14 -%52 = OpConstant %6 1 -%55 = OpConstant %6 3 -%57 = OpTypePointer StorageBuffer %17 -%58 = OpTypePointer StorageBuffer %14 -%65 = OpConstant %6 256 -%41 = OpFunction %2 None %42 -%40 = OpLabel -%36 = OpVariable %37 Function %13 -%38 = OpVariable %39 Function %15 -%45 = OpAccessChain %43 %27 %44 -%47 = OpAccessChain %46 %30 %44 -OpBranch %49 -%49 = OpLabel -%53 = OpAccessChain %51 %45 %52 -%54 = OpLoad %14 %53 -%56 = OpAccessChain %50 %23 %55 -OpStore %56 %54 -%59 = OpAccessChain %58 %45 %44 %44 -%60 = OpLoad %14 %59 -%61 = OpAccessChain %50 %23 %12 -OpStore %61 %60 -%62 = OpArrayLength %6 %30 0 -%63 = OpConvertUToF %14 %62 -%64 = OpAccessChain %50 %23 %52 -OpStore %64 %63 -OpAtomicStore %25 %10 %65 %12 +%12 = OpTypeFloat 32 +%11 = OpConstant %12 4.0 +%13 = OpConstant %8 1 +%14 = OpConstant %6 2 +%15 = OpConstant %12 1.0 +%16 = OpConstantTrue %4 +%17 = OpTypeArray %12 %5 +%18 = OpTypeVector %12 3 +%19 = OpTypeStruct %18 %12 +%20 = OpTypeVector %12 2 +%21 = OpTypeRuntimeArray %20 +%22 = OpTypeVector %12 4 +%23 = OpTypeArray %22 %7 +%25 = OpTypePointer Workgroup %17 +%24 = OpVariable %25 Workgroup +%27 = OpTypePointer Workgroup %6 +%26 = OpVariable %27 Workgroup +%29 = OpTypeStruct %19 +%30 = OpTypePointer StorageBuffer %29 +%28 = OpVariable %30 StorageBuffer +%32 = OpTypeStruct %21 +%33 = OpTypePointer StorageBuffer %32 +%31 = OpVariable %33 StorageBuffer +%35 = OpTypeStruct %23 +%36 = OpTypePointer Uniform %35 +%34 = OpVariable %36 Uniform +%38 = OpTypePointer Function %12 +%40 = OpTypePointer Function %4 +%43 = OpTypeFunction %2 +%44 = OpTypePointer StorageBuffer %19 +%45 = OpConstant %6 0 +%47 = OpTypePointer StorageBuffer %21 +%49 = OpTypePointer Uniform %23 +%51 = OpTypePointer Workgroup %12 +%52 = OpTypePointer StorageBuffer %12 +%53 = OpConstant %6 1 +%56 = OpConstant %6 3 +%58 = OpTypePointer StorageBuffer %18 +%59 = OpTypePointer StorageBuffer %12 +%67 = OpConstant %6 256 +%42 = OpFunction %2 None %43 +%41 = OpLabel +%37 = OpVariable %38 Function %15 +%39 = OpVariable %40 Function %16 +%46 = OpAccessChain %44 %28 %45 +%48 = OpAccessChain %47 %31 %45 +OpBranch %50 +%50 = OpLabel +%54 = OpAccessChain %52 %46 %53 +%55 = OpLoad %12 %54 +%57 = OpAccessChain %51 %24 %56 +OpStore %57 %55 +%60 = OpAccessChain %59 %46 %45 %45 +%61 = OpLoad %12 %60 +%62 = OpAccessChain %51 %24 %14 +OpStore %62 %61 +%63 = OpAccessChain %52 %46 %53 +OpStore %63 %11 +%64 = OpArrayLength %6 %31 0 +%65 = OpConvertUToF %12 %64 +%66 = OpAccessChain %51 %24 %53 +OpStore %66 %65 +OpAtomicStore %26 %10 %67 %14 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index 28d29440f9..1080002efc 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -13,34 +13,39 @@ struct Bar { @group(0) @binding(0) var bar: Bar; -fn read_from_private(foo_2: ptr) -> f32 { - let _e2 = (*foo_2); +fn read_from_private(foo_1: ptr) -> f32 { + let _e2 = (*foo_1); return _e2; } @stage(vertex) -fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { - var foo_1: f32 = 0.0; +fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4 { + var foo: f32 = 0.0; var c: array; - let baz = foo_1; - foo_1 = 1.0; + let baz = foo; + foo = 1.0; let matrix = bar.matrix; let arr = bar.arr; let b = bar.matrix[3][0]; let a = bar.data[(arrayLength((&bar.data)) - 2u)].value; let data_pointer = (&bar.data[0].value); - let _e27 = read_from_private((&foo_1)); - bar.matrix[1][2] = 1.0; - bar.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); - bar.arr = array,2>(vec2(0u), vec2(1u)); - bar.data[1].value = 1; + let _e27 = read_from_private((&foo)); c = array(a, i32(b), 3, 4, 5); c[(vi + 1u)] = 42; let value = c[vi]; return (matrix * vec4(vec4(value))); } +@stage(fragment) +fn foo_frag() -> @location(0) vec4 { + bar.matrix[1][2] = 1.0; + bar.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0)); + bar.arr = array,2>(vec2(0u), vec2(1u)); + bar.data[1].value = 1; + return vec4(0.0); +} + @stage(compute) @workgroup_size(1, 1, 1) fn atomics() { var tmp: i32; diff --git a/tests/out/wgsl/globals.wgsl b/tests/out/wgsl/globals.wgsl index f73c6c04d3..d1ab5cbf11 100644 --- a/tests/out/wgsl/globals.wgsl +++ b/tests/out/wgsl/globals.wgsl @@ -8,7 +8,7 @@ let Foo_2: bool = true; var wg: array; var at_1: atomic; @group(0) @binding(1) -var alignment: Foo; +var alignment: Foo; @group(0) @binding(2) var dummy: array>; @group(0) @binding(3) @@ -23,6 +23,7 @@ fn main() { wg[3] = _e9; let _e14 = alignment.v3_.x; wg[2] = _e14; + alignment.v1_ = 4.0; wg[1] = f32(arrayLength((&dummy))); atomicStore((&at_1), 2u); return;