From 688ad474f71cce856a6dbd792dccbb3b702da477 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Mon, 21 Feb 2022 13:22:54 -0800 Subject: [PATCH] [msl-out] Write 'using metal::uint' at the top, to work around bug. (#1740) This lets us remove some special cases where `uint` must be fully qualified, even though other similar types are not. --- src/back/msl/keywords.rs | 1 + src/back/msl/writer.rs | 13 +++++----- tests/out/msl/access.msl | 6 +++-- tests/out/msl/bits.msl | 22 ++++++++-------- tests/out/msl/boids.msl | 20 ++++++++------- tests/out/msl/bounds-check-restrict.msl | 4 ++- tests/out/msl/bounds-check-zero-atomic.msl | 16 +++++++----- tests/out/msl/bounds-check-zero.msl | 30 ++++++++++++---------- tests/out/msl/collatz.msl | 30 ++++++++++++---------- tests/out/msl/control-flow.msl | 2 ++ tests/out/msl/empty-global-name.msl | 2 ++ tests/out/msl/empty.msl | 2 ++ tests/out/msl/extra.msl | 8 +++--- tests/out/msl/functions.msl | 4 ++- tests/out/msl/globals.msl | 4 ++- tests/out/msl/image.msl | 8 +++--- tests/out/msl/interface.msl | 24 +++++++++-------- tests/out/msl/interpolate.msl | 8 +++--- tests/out/msl/math-functions.msl | 2 ++ tests/out/msl/operators.msl | 2 ++ tests/out/msl/policy-mix.msl | 2 ++ tests/out/msl/quad-vert.msl | 2 ++ tests/out/msl/quad.msl | 2 ++ tests/out/msl/shadow.msl | 18 +++++++------ tests/out/msl/skybox.msl | 4 ++- tests/out/msl/standard.msl | 2 ++ tests/out/msl/texture-arg.msl | 2 ++ 27 files changed, 145 insertions(+), 95 deletions(-) diff --git a/src/back/msl/keywords.rs b/src/back/msl/keywords.rs index a62113c097..7a66b94aa9 100644 --- a/src/back/msl/keywords.rs +++ b/src/back/msl/keywords.rs @@ -18,6 +18,7 @@ pub const RESERVED: &[&str] = &[ "bool", "char", "int", + "uint", "long", "float", "double", diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 32e003c656..a99c9db2a4 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -39,10 +39,6 @@ fn put_numeric_type( sizes: &[crate::VectorSize], ) -> Result<(), FmtError> { match (kind, sizes) { - (crate::ScalarKind::Uint, &[]) => { - // Work around Metal compiler bug: scalar `uint` requires namespace. - write!(out, "{}::uint", NAMESPACE) - } (kind, &[]) => { write!(out, "{}", kind.to_msl_name()) } @@ -1434,7 +1430,7 @@ impl Writer { /// The text written is of the form: /// /// ```ignore - /// {level}{prefix}metal::uint(i) < 4 && metal::uint(j) < 10 + /// {level}{prefix}uint(i) < 4 && uint(j) < 10 /// ``` /// /// where `{level}` and `{prefix}` are the arguments to this function. For [`Store`] @@ -1494,7 +1490,7 @@ impl Writer { // Check that the index falls within bounds. Do this with a single // comparison, by casting the index to `uint` first, so that negative // indices become large positive values. - write!(self.out, "{}::uint(", NAMESPACE)?; + write!(self.out, "uint(")?; self.put_index(index, context, true)?; self.out.write_str(") < ")?; match length { @@ -2351,6 +2347,9 @@ impl Writer { writeln!(self.out, "#include ")?; writeln!(self.out, "#include ")?; writeln!(self.out)?; + // Work around Metal bug where `uint` is not available by default + writeln!(self.out, "using {}::uint;", NAMESPACE)?; + writeln!(self.out)?; if options .bounds_check_policies @@ -2372,7 +2371,7 @@ impl Writer { writeln!(self.out, "struct _mslBufferSizes {{")?; for idx in indices { - writeln!(self.out, "{}{}::uint size{};", back::INDENT, NAMESPACE, idx)?; + writeln!(self.out, "{}uint size{};", back::INDENT, idx)?; } writeln!(self.out, "}};")?; diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index 39033f3fdd..d4d80a50b6 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -2,8 +2,10 @@ #include #include +using metal::uint; + struct _mslBufferSizes { - metal::uint size0; + uint size0; }; struct AlignedWrapper { @@ -41,7 +43,7 @@ struct fooOutput { metal::float4 member [[position]]; }; vertex fooOutput foo( - metal::uint vi [[vertex_id]] + uint vi [[vertex_id]] , device Bar& bar [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { diff --git a/tests/out/msl/bits.msl b/tests/out/msl/bits.msl index 05574f7170..cc3a5660ae 100644 --- a/tests/out/msl/bits.msl +++ b/tests/out/msl/bits.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + kernel void main_( ) { @@ -9,7 +11,7 @@ kernel void main_( metal::int2 i2_; metal::int3 i3_; metal::int4 i4_; - metal::uint u = 0u; + uint u = 0u; metal::uint2 u2_; metal::uint3 u3_; metal::uint4 u4_; @@ -33,15 +35,15 @@ kernel void main_( u = metal::pack_float_to_unorm2x16(_e34); metal::float2 _e36 = f2_; u = as_type(half2(_e36)); - metal::uint _e38 = u; + uint _e38 = u; f4_ = metal::unpack_snorm4x8_to_float(_e38); - metal::uint _e40 = u; + uint _e40 = u; f4_ = metal::unpack_unorm4x8_to_float(_e40); - metal::uint _e42 = u; + uint _e42 = u; f2_ = metal::unpack_snorm2x16_to_float(_e42); - metal::uint _e44 = u; + uint _e44 = u; f2_ = metal::unpack_unorm2x16_to_float(_e44); - metal::uint _e46 = u; + uint _e46 = u; f2_ = float2(as_type(_e46)); int _e48 = i; int _e49 = i; @@ -55,8 +57,8 @@ kernel void main_( metal::int4 _e63 = i4_; metal::int4 _e64 = i4_; i4_ = metal::insert_bits(_e63, _e64, 5u, 10u); - metal::uint _e68 = u; - metal::uint _e69 = u; + uint _e68 = u; + uint _e69 = u; u = metal::insert_bits(_e68, _e69, 5u, 10u); metal::uint2 _e73 = u2_; metal::uint2 _e74 = u2_; @@ -75,7 +77,7 @@ kernel void main_( i3_ = metal::extract_bits(_e96, 5u, 10u); metal::int4 _e100 = i4_; i4_ = metal::extract_bits(_e100, 5u, 10u); - metal::uint _e104 = u; + uint _e104 = u; u = metal::extract_bits(_e104, 5u, 10u); metal::uint2 _e108 = u2_; u2_ = metal::extract_bits(_e108, 5u, 10u); @@ -89,7 +91,7 @@ kernel void main_( i2_ = (((1 + int2(metal::ctz(_e122))) % 33) - 1); metal::int3 _e124 = i3_; i3_ = (((1 + int3(metal::clz(_e124))) % 33) - 1); - metal::uint _e126 = u; + uint _e126 = u; i = (((1 + int(metal::clz(_e126))) % 33) - 1); return; } diff --git a/tests/out/msl/boids.msl b/tests/out/msl/boids.msl index 9f60489978..d18d733ac7 100644 --- a/tests/out/msl/boids.msl +++ b/tests/out/msl/boids.msl @@ -2,9 +2,11 @@ #include #include +using metal::uint; + struct _mslBufferSizes { - metal::uint size1; - metal::uint size2; + uint size1; + uint size2; }; constexpr constant unsigned NUM_PARTICLES = 1500u; @@ -44,8 +46,8 @@ kernel void main_( int cVelCount = 0; metal::float2 pos; metal::float2 vel; - metal::uint i = 0u; - metal::uint index = global_invocation_id.x; + uint i = 0u; + uint index = global_invocation_id.x; if (index >= NUM_PARTICLES) { return; } @@ -59,22 +61,22 @@ kernel void main_( bool loop_init = true; while(true) { if (!loop_init) { - metal::uint _e86 = i; + uint _e86 = i; i = _e86 + 1u; } loop_init = false; - metal::uint _e37 = i; + uint _e37 = i; if (_e37 >= NUM_PARTICLES) { break; } - metal::uint _e39 = i; + uint _e39 = i; if (_e39 == index) { continue; } - metal::uint _e42 = i; + uint _e42 = i; metal::float2 _e45 = particlesSrc.particles[_e42].pos; pos = _e45; - metal::uint _e47 = i; + uint _e47 = i; metal::float2 _e50 = particlesSrc.particles[_e47].vel; vel = _e50; metal::float2 _e51 = pos; diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl index 93bbf910be..e82ebf814d 100644 --- a/tests/out/msl/bounds-check-restrict.msl +++ b/tests/out/msl/bounds-check-restrict.msl @@ -2,8 +2,10 @@ #include #include +using metal::uint; + struct _mslBufferSizes { - metal::uint size0; + uint size0; }; struct type_1 { diff --git a/tests/out/msl/bounds-check-zero-atomic.msl b/tests/out/msl/bounds-check-zero-atomic.msl index 038bef6bfc..c76ba1b549 100644 --- a/tests/out/msl/bounds-check-zero-atomic.msl +++ b/tests/out/msl/bounds-check-zero-atomic.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct DefaultConstructible { template operator T() && { @@ -9,7 +11,7 @@ struct DefaultConstructible { } }; struct _mslBufferSizes { - metal::uint size0; + uint size0; }; struct type_1 { @@ -22,28 +24,28 @@ struct Globals { type_2 c; }; -metal::uint fetch_add_atomic( +uint fetch_add_atomic( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - metal::uint _e3 = metal::atomic_fetch_add_explicit(&globals.a, 1u, metal::memory_order_relaxed); + uint _e3 = metal::atomic_fetch_add_explicit(&globals.a, 1u, metal::memory_order_relaxed); return _e3; } -metal::uint fetch_add_atomic_static_sized_array( +uint fetch_add_atomic_static_sized_array( int i, device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - metal::uint _e5 = metal::uint(i) < 10 ? metal::atomic_fetch_add_explicit(&globals.b.inner[i], 1u, metal::memory_order_relaxed) : DefaultConstructible(); + uint _e5 = uint(i) < 10 ? metal::atomic_fetch_add_explicit(&globals.b.inner[i], 1u, metal::memory_order_relaxed) : DefaultConstructible(); return _e5; } -metal::uint fetch_add_atomic_dynamic_sized_array( +uint fetch_add_atomic_dynamic_sized_array( int i_1, device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - metal::uint _e5 = metal::uint(i_1) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_fetch_add_explicit(&globals.c[i_1], 1u, metal::memory_order_relaxed) : DefaultConstructible(); + uint _e5 = uint(i_1) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_fetch_add_explicit(&globals.c[i_1], 1u, metal::memory_order_relaxed) : DefaultConstructible(); return _e5; } diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl index b61e523aab..1a6bcc1578 100644 --- a/tests/out/msl/bounds-check-zero.msl +++ b/tests/out/msl/bounds-check-zero.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct DefaultConstructible { template operator T() && { @@ -9,7 +11,7 @@ struct DefaultConstructible { } }; struct _mslBufferSizes { - metal::uint size0; + uint size0; }; struct type_1 { @@ -29,7 +31,7 @@ float index_array( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - float _e4 = metal::uint(i) < 10 ? globals.a.inner[i] : DefaultConstructible(); + float _e4 = uint(i) < 10 ? globals.a.inner[i] : DefaultConstructible(); return _e4; } @@ -38,7 +40,7 @@ float index_dynamic_array( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - float _e4 = metal::uint(i_1) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4 ? globals.d[i_1] : DefaultConstructible(); + float _e4 = uint(i_1) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4 ? globals.d[i_1] : DefaultConstructible(); return _e4; } @@ -47,7 +49,7 @@ float index_vector( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - float _e4 = metal::uint(i_2) < 4 ? globals.v[i_2] : DefaultConstructible(); + float _e4 = uint(i_2) < 4 ? globals.v[i_2] : DefaultConstructible(); return _e4; } @@ -55,7 +57,7 @@ float index_vector_by_value( metal::float4 v, int i_3 ) { - return metal::uint(i_3) < 4 ? v[i_3] : DefaultConstructible(); + return uint(i_3) < 4 ? v[i_3] : DefaultConstructible(); } metal::float4 index_matrix( @@ -63,7 +65,7 @@ metal::float4 index_matrix( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - metal::float4 _e4 = metal::uint(i_4) < 3 ? globals.m[i_4] : DefaultConstructible(); + metal::float4 _e4 = uint(i_4) < 3 ? globals.m[i_4] : DefaultConstructible(); return _e4; } @@ -73,7 +75,7 @@ float index_twice( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - float _e6 = metal::uint(j) < 4 && metal::uint(i_5) < 3 ? globals.m[i_5][j] : DefaultConstructible(); + float _e6 = uint(j) < 4 && uint(i_5) < 3 ? globals.m[i_5][j] : DefaultConstructible(); return _e6; } @@ -83,7 +85,7 @@ float index_expensive( constant _mslBufferSizes& _buffer_sizes ) { int _e9 = static_cast(metal::sin(static_cast(i_6) / 100.0) * 100.0); - float _e11 = metal::uint(_e9) < 10 ? globals.a.inner[_e9] : DefaultConstructible(); + float _e11 = uint(_e9) < 10 ? globals.a.inner[_e9] : DefaultConstructible(); return _e11; } @@ -103,7 +105,7 @@ void set_array( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - if (metal::uint(i_7) < 10) { + if (uint(i_7) < 10) { globals.a.inner[i_7] = v_1; } return; @@ -115,7 +117,7 @@ void set_dynamic_array( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - if (metal::uint(i_8) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4) { + if (uint(i_8) < 1 + (_buffer_sizes.size0 - 112 - 4) / 4) { globals.d[i_8] = v_2; } return; @@ -127,7 +129,7 @@ void set_vector( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - if (metal::uint(i_9) < 4) { + if (uint(i_9) < 4) { globals.v[i_9] = v_3; } return; @@ -139,7 +141,7 @@ void set_matrix( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - if (metal::uint(i_10) < 3) { + if (uint(i_10) < 3) { globals.m[i_10] = v_4; } return; @@ -152,7 +154,7 @@ void set_index_twice( device Globals& globals, constant _mslBufferSizes& _buffer_sizes ) { - if (metal::uint(j_1) < 4 && metal::uint(i_11) < 3) { + if (uint(j_1) < 4 && uint(i_11) < 3) { globals.m[i_11][j_1] = v_5; } return; @@ -165,7 +167,7 @@ void set_expensive( constant _mslBufferSizes& _buffer_sizes ) { int _e10 = static_cast(metal::sin(static_cast(i_12) / 100.0) * 100.0); - if (metal::uint(_e10) < 10) { + if (uint(_e10) < 10) { globals.a.inner[_e10] = v_6; } return; diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl index 3a511481a1..a1a0788ada 100644 --- a/tests/out/msl/collatz.msl +++ b/tests/out/msl/collatz.msl @@ -2,38 +2,40 @@ #include #include +using metal::uint; + struct _mslBufferSizes { - metal::uint size0; + uint size0; }; -typedef metal::uint type_1[1]; +typedef uint type_1[1]; struct PrimeIndices { type_1 data; }; -metal::uint collatz_iterations( - metal::uint n_base +uint collatz_iterations( + uint n_base ) { - metal::uint n; - metal::uint i = 0u; + uint n; + uint i = 0u; n = n_base; while(true) { - metal::uint _e5 = n; + uint _e5 = n; if (_e5 <= 1u) { break; } - metal::uint _e8 = n; + uint _e8 = n; if ((_e8 % 2u) == 0u) { - metal::uint _e13 = n; + uint _e13 = n; n = _e13 / 2u; } else { - metal::uint _e17 = n; + uint _e17 = n; n = (3u * _e17) + 1u; } - metal::uint _e21 = i; + uint _e21 = i; i = _e21 + 1u; } - metal::uint _e24 = i; + uint _e24 = i; return _e24; } @@ -44,8 +46,8 @@ kernel void main_( , device PrimeIndices& v_indices [[user(fake0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { - metal::uint _e8 = v_indices.data[global_id.x]; - metal::uint _e9 = collatz_iterations(_e8); + uint _e8 = v_indices.data[global_id.x]; + uint _e9 = collatz_iterations(_e8); v_indices.data[global_id.x] = _e9; return; } diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl index e0d2f0cebc..8fb23b5b76 100644 --- a/tests/out/msl/control-flow.msl +++ b/tests/out/msl/control-flow.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + void switch_default_break( int i diff --git a/tests/out/msl/empty-global-name.msl b/tests/out/msl/empty-global-name.msl index 3b9dfec5fc..6001dba93f 100644 --- a/tests/out/msl/empty-global-name.msl +++ b/tests/out/msl/empty-global-name.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct type_1 { int member; }; diff --git a/tests/out/msl/empty.msl b/tests/out/msl/empty.msl index c3bb1a71b9..4f8bf9f5e9 100644 --- a/tests/out/msl/empty.msl +++ b/tests/out/msl/empty.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + kernel void main_( ) { diff --git a/tests/out/msl/extra.msl b/tests/out/msl/extra.msl index e13f0c2ac6..867c7e5a00 100644 --- a/tests/out/msl/extra.msl +++ b/tests/out/msl/extra.msl @@ -2,14 +2,16 @@ #include #include +using metal::uint; + struct PushConstants { - metal::uint index; + uint index; char _pad1[12]; metal::float2 double_; }; struct FragmentIn { metal::float4 color; - metal::uint primitive_index; + uint primitive_index; }; struct main_Input { @@ -20,7 +22,7 @@ struct main_Output { }; fragment main_Output main_( main_Input varyings [[stage_in]] -, metal::uint primitive_index [[primitive_id]] +, uint primitive_index [[primitive_id]] ) { const FragmentIn in = { varyings.color, primitive_index }; if ((in.primitive_index % 2u) == 0u) { diff --git a/tests/out/msl/functions.msl b/tests/out/msl/functions.msl index 73c5a9fa64..31f57c656e 100644 --- a/tests/out/msl/functions.msl +++ b/tests/out/msl/functions.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + metal::float2 test_fma( ) { @@ -18,7 +20,7 @@ int test_integer_dot_product( int c_2_ = ( + a_2_.x * b_2_.x + a_2_.y * b_2_.y); metal::uint3 a_3_ = metal::uint3(1u); metal::uint3 b_3_ = metal::uint3(1u); - metal::uint c_3_ = ( + a_3_.x * b_3_.x + a_3_.y * b_3_.y + a_3_.z * b_3_.z); + uint c_3_ = ( + a_3_.x * b_3_.x + a_3_.y * b_3_.y + a_3_.z * b_3_.z); metal::int4 _e11 = metal::int4(4); metal::int4 _e13 = metal::int4(2); int c_4_ = ( + _e11.x * _e13.x + _e11.y * _e13.y + _e11.z * _e13.z + _e11.w * _e13.w); diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index b56bc647dc..8f4ce07ade 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -2,8 +2,10 @@ #include #include +using metal::uint; + struct _mslBufferSizes { - metal::uint size3; + uint size3; }; constexpr constant bool Foo_2 = true; diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index 025f2a6a90..74d75ff7d6 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + constant metal::int2 const_type_9_ = {3, 1}; struct main_Input { @@ -21,8 +23,8 @@ kernel void main_( metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast(local_id.z)); metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc)); metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), static_cast(local_id.z), static_cast(local_id.z) + 1); - metal::uint4 value6_ = image_1d_src.read(metal::uint(static_cast(local_id.x))); - image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, metal::uint(itc.x)); + metal::uint4 value6_ = image_1d_src.read(uint(static_cast(local_id.x))); + image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, uint(itc.x)); return; } @@ -38,7 +40,7 @@ kernel void depth_load( metal::int2 dim_1 = int2(image_storage_src.get_width(), image_storage_src.get_height()); metal::int2 itc_1 = (dim_1 * static_cast(local_id_1.xy)) % metal::int2(10, 20); float val = image_depth_multisampled_src.read(metal::uint2(itc_1), static_cast(local_id_1.z)); - image_dst.write(metal::uint4(static_cast(val)), metal::uint(itc_1.x)); + image_dst.write(metal::uint4(static_cast(val)), uint(itc_1.x)); return; } diff --git a/tests/out/msl/interface.msl b/tests/out/msl/interface.msl index 6461aee75a..dbc6871984 100644 --- a/tests/out/msl/interface.msl +++ b/tests/out/msl/interface.msl @@ -2,21 +2,23 @@ #include #include +using metal::uint; + struct VertexOutput { metal::float4 position; float varying; }; struct FragmentOutput { float depth; - metal::uint sample_mask; + uint sample_mask; float color; }; struct type_4 { - metal::uint inner[1]; + uint inner[1]; }; struct vertex_Input { - metal::uint color [[attribute(10)]]; + uint color [[attribute(10)]]; }; struct vertex_Output { metal::float4 position [[position]]; @@ -25,11 +27,11 @@ struct vertex_Output { }; vertex vertex_Output vertex_( vertex_Input varyings [[stage_in]] -, metal::uint vertex_index [[vertex_id]] -, metal::uint instance_index [[instance_id]] +, uint vertex_index [[vertex_id]] +, uint instance_index [[instance_id]] ) { const auto color = varyings.color; - metal::uint tmp = (vertex_index + instance_index) + color; + uint tmp = (vertex_index + instance_index) + color; const auto _tmp = VertexOutput {metal::float4(1.0), static_cast(tmp)}; return vertex_Output { _tmp.position, _tmp.varying, 1.0 }; } @@ -40,18 +42,18 @@ struct fragment_Input { }; struct fragment_Output { float depth [[depth(any)]]; - metal::uint sample_mask [[sample_mask]]; + uint sample_mask [[sample_mask]]; float color [[color(0)]]; }; fragment fragment_Output fragment_( fragment_Input varyings_1 [[stage_in]] , metal::float4 position [[position]] , bool front_facing [[front_facing]] -, metal::uint sample_index [[sample_id]] -, metal::uint sample_mask [[sample_mask]] +, uint sample_index [[sample_id]] +, uint sample_mask [[sample_mask]] ) { const VertexOutput in = { position, varyings_1.varying }; - metal::uint mask = sample_mask & (1u << sample_index); + uint mask = sample_mask & (1u << sample_index); float color_1 = front_facing ? 1.0 : 0.0; const auto _tmp = FragmentOutput {in.varying, mask, color_1}; return fragment_Output { _tmp.depth, _tmp.sample_mask, _tmp.color }; @@ -63,7 +65,7 @@ struct compute_Input { kernel void compute_( metal::uint3 global_id [[thread_position_in_grid]] , metal::uint3 local_id [[thread_position_in_threadgroup]] -, metal::uint local_index [[thread_index_in_threadgroup]] +, uint local_index [[thread_index_in_threadgroup]] , metal::uint3 wg_id [[threadgroup_position_in_grid]] , metal::uint3 num_wgs [[threadgroups_per_grid]] , threadgroup type_4& output diff --git a/tests/out/msl/interpolate.msl b/tests/out/msl/interpolate.msl index 0f89a547ba..659b9d3d23 100644 --- a/tests/out/msl/interpolate.msl +++ b/tests/out/msl/interpolate.msl @@ -2,9 +2,11 @@ #include #include +using metal::uint; + struct FragmentInput { metal::float4 position; - metal::uint flat; + uint flat; float linear; metal::float2 linear_centroid; metal::float3 linear_sample; @@ -15,7 +17,7 @@ struct FragmentInput { struct vert_mainOutput { metal::float4 position [[position]]; - metal::uint flat [[user(loc0), flat]]; + uint flat [[user(loc0), flat]]; float linear [[user(loc1), center_no_perspective]]; metal::float2 linear_centroid [[user(loc2), centroid_no_perspective]]; metal::float3 linear_sample [[user(loc3), sample_no_perspective]]; @@ -41,7 +43,7 @@ vertex vert_mainOutput vert_main( struct frag_mainInput { - metal::uint flat [[user(loc0), flat]]; + uint flat [[user(loc0), flat]]; float linear [[user(loc1), center_no_perspective]]; metal::float2 linear_centroid [[user(loc2), centroid_no_perspective]]; metal::float3 linear_sample [[user(loc3), sample_no_perspective]]; diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index ee0acbca84..fe1fd5573f 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + vertex void main_( ) { diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index eebd1e73b8..1d76848a2d 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct Foo { metal::float4 a; int b; diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl index 893c700ca0..9020c64fe0 100644 --- a/tests/out/msl/policy-mix.msl +++ b/tests/out/msl/policy-mix.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct DefaultConstructible { template operator T() && { diff --git a/tests/out/msl/quad-vert.msl b/tests/out/msl/quad-vert.msl index 87114a8d74..ade1791db1 100644 --- a/tests/out/msl/quad-vert.msl +++ b/tests/out/msl/quad-vert.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct type_5 { float inner[1u]; }; diff --git a/tests/out/msl/quad.msl b/tests/out/msl/quad.msl index a03d61bf57..a6beb30b8c 100644 --- a/tests/out/msl/quad.msl +++ b/tests/out/msl/quad.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + constexpr constant float c_scale = 1.2000000476837158; struct VertexOutput { metal::float2 uv; diff --git a/tests/out/msl/shadow.msl b/tests/out/msl/shadow.msl index f7978728d2..58d03e957d 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -2,8 +2,10 @@ #include #include +using metal::uint; + struct _mslBufferSizes { - metal::uint size1; + uint size1; }; constexpr constant unsigned c_max_lights = 10u; @@ -22,7 +24,7 @@ struct Lights { constant metal::float3 c_ambient = {0.05000000074505806, 0.05000000074505806, 0.05000000074505806}; float fetch_shadow( - metal::uint light_id, + uint light_id, metal::float4 homogeneous_coords, metal::depth2d_array t_shadow, metal::sampler sampler_shadow @@ -54,23 +56,23 @@ fragment fs_mainOutput fs_main( const auto raw_normal = varyings.raw_normal; const auto position = varyings.position; metal::float3 color = c_ambient; - metal::uint i = 0u; + uint i = 0u; metal::float3 normal = metal::normalize(raw_normal); bool loop_init = true; while(true) { if (!loop_init) { - metal::uint _e40 = i; + uint _e40 = i; i = _e40 + 1u; } loop_init = false; - metal::uint _e12 = i; - metal::uint _e15 = u_globals.num_lights.x; + uint _e12 = i; + uint _e15 = u_globals.num_lights.x; if (_e12 >= metal::min(_e15, c_max_lights)) { break; } - metal::uint _e19 = i; + uint _e19 = i; Light light = s_lights.data[_e19]; - metal::uint _e22 = i; + uint _e22 = i; float _e25 = fetch_shadow(_e22, light.proj * position, t_shadow, sampler_shadow); metal::float3 light_dir = metal::normalize(light.pos.xyz - position.xyz); float diffuse = metal::max(0.0, metal::dot(normal, light_dir)); diff --git a/tests/out/msl/skybox.msl b/tests/out/msl/skybox.msl index 26b6812a42..b6e2ce5b1a 100644 --- a/tests/out/msl/skybox.msl +++ b/tests/out/msl/skybox.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct VertexOutput { metal::float4 position; metal::float3 uv; @@ -18,7 +20,7 @@ struct vs_mainOutput { metal::float3 uv [[user(loc0), center_perspective]]; }; vertex vs_mainOutput vs_main( - metal::uint vertex_index [[vertex_id]] + uint vertex_index [[vertex_id]] , constant Data& r_data [[buffer(0)]] ) { int tmp1_; diff --git a/tests/out/msl/standard.msl b/tests/out/msl/standard.msl index 0164252374..d21930c061 100644 --- a/tests/out/msl/standard.msl +++ b/tests/out/msl/standard.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + struct derivativesInput { }; diff --git a/tests/out/msl/texture-arg.msl b/tests/out/msl/texture-arg.msl index 20e76207d3..ba0be05e18 100644 --- a/tests/out/msl/texture-arg.msl +++ b/tests/out/msl/texture-arg.msl @@ -2,6 +2,8 @@ #include #include +using metal::uint; + metal::float4 test( metal::texture2d Passed_Texture,