[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

@@ -18,6 +18,7 @@ pub const RESERVED: &[&str] = &[
     "bool",
     "char",
     "int",
+    "uint",
     "long",
     "float",
     "double",

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<W: Write> Writer<W> {
     /// 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<W: Write> Writer<W> {
                     // 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<W: Write> Writer<W> {
         writeln!(self.out, "#include <metal_stdlib>")?;
         writeln!(self.out, "#include <simd/simd.h>")?;
         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<W: Write> Writer<W> {
                 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, "}};")?;

tests/out/msl/access.msl

@@ -2,8 +2,10 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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)]]
 ) {

tests/out/msl/bits.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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<uint>(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<half2>(_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;
 }

tests/out/msl/boids.msl

@@ -2,9 +2,11 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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;

tests/out/msl/bounds-check-restrict.msl

@@ -2,8 +2,10 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct _mslBufferSizes {
-    metal::uint size0;
+    uint size0;
 };
 
 struct type_1 {

tests/out/msl/bounds-check-zero-atomic.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct DefaultConstructible {
     template<typename T>
     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;
 }

tests/out/msl/bounds-check-zero.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct DefaultConstructible {
     template<typename T>
     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<int>(metal::sin(static_cast<float>(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<int>(metal::sin(static_cast<float>(i_12) / 100.0) * 100.0);
-    if (metal::uint(_e10) < 10) {
+    if (uint(_e10) < 10) {
         globals.a.inner[_e10] = v_6;
     }
     return;

tests/out/msl/collatz.msl

@@ -2,38 +2,40 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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;
 }

tests/out/msl/control-flow.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 
 void switch_default_break(
     int i

tests/out/msl/empty-global-name.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct type_1 {
     int member;
 };

tests/out/msl/empty.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 
 kernel void main_(
 ) {

tests/out/msl/extra.msl

@@ -2,14 +2,16 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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) {

tests/out/msl/functions.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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);

tests/out/msl/globals.msl

@@ -2,8 +2,10 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct _mslBufferSizes {
-    metal::uint size3;
+    uint size3;
 };
 
 constexpr constant bool Foo_2 = true;

tests/out/msl/image.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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<int>(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<int>(local_id.z), static_cast<int>(local_id.z) + 1);
-    metal::uint4 value6_ = image_1d_src.read(metal::uint(static_cast<int>(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<int>(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<metal::int2>(local_id_1.xy)) % metal::int2(10, 20);
     float val = image_depth_multisampled_src.read(metal::uint2(itc_1), static_cast<int>(local_id_1.z));
-    image_dst.write(metal::uint4(static_cast<metal::uint>(val)), metal::uint(itc_1.x));
+    image_dst.write(metal::uint4(static_cast<uint>(val)), uint(itc_1.x));
     return;
 }
 

tests/out/msl/interface.msl

@@ -2,21 +2,23 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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<float>(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

tests/out/msl/interpolate.msl

@@ -2,9 +2,11 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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]];

tests/out/msl/math-functions.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 
 vertex void main_(
 ) {

tests/out/msl/operators.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct Foo {
     metal::float4 a;
     int b;

tests/out/msl/policy-mix.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct DefaultConstructible {
     template<typename T>
     operator T() && {

tests/out/msl/quad-vert.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 struct type_5 {
     float inner[1u];
 };

tests/out/msl/quad.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 constexpr constant float c_scale = 1.2000000476837158;
 struct VertexOutput {
     metal::float2 uv;

tests/out/msl/shadow.msl

@@ -2,8 +2,10 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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<float, metal::access::sample> 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));

tests/out/msl/skybox.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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_;

tests/out/msl/standard.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 
 struct derivativesInput {
 };

tests/out/msl/texture-arg.msl

@@ -2,6 +2,8 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+using metal::uint;
+
 
 metal::float4 test(
     metal::texture2d<float, metal::access::sample> Passed_Texture,