diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index d731ec13e4..0f78cac21a 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -2310,6 +2310,26 @@ bool GeneratorImpl::EmitPackedType(std::ostream& out,
     if (!EmitType(out, vec, "")) {
       return false;
     }
+
+    if (vec->is_float_vector() && !matrix_packed_vector_overloads_) {
+      // Overload operators for matrix-vector arithmetic where the vector
+      // operand is packed, as these overloads to not exist in the metal
+      // namespace.
+      TextBuffer b;
+      TINT_DEFER(helpers_.Append(b));
+      line(&b) << R"(template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+)";
+      matrix_packed_vector_overloads_ = true;
+    }
+
     return true;
   }
 
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index 6a4efb7f0b..9c65c158e3 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -355,6 +355,9 @@ class GeneratorImpl : public TextGenerator {
   /// True if an invariant attribute has been generated.
   bool has_invariant_ = false;
 
+  /// True if matrix-packed_vector operator overloads have been generated.
+  bool matrix_packed_vector_overloads_ = false;
+
   std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
   std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
 };
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
index 6b40e48bd5..12f9ac5ad0 100644
--- a/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct tint_array_wrapper {
   /* 0x0000 */ int4 arr[4];
 };
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
index fcb9558edc..4c36e4818e 100644
--- a/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct tint_array_wrapper {
   /* 0x0000 */ int4 arr[4];
 };
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.msl b/test/buffer/storage/static_index/read.wgsl.expected.msl
index 316ed31324..6dbc898fb7 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.msl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct Inner {
   /* 0x0000 */ int x;
 };
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.msl b/test/buffer/storage/static_index/write.wgsl.expected.msl
index eab1f4f5a9..b2b8a49593 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct Inner {
   /* 0x0000 */ int x;
 };
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
index eed156e0a6..7d55efd45c 100644
--- a/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct tint_array_wrapper {
   /* 0x0000 */ int4 arr[4];
 };
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.msl b/test/buffer/uniform/static_index/read.wgsl.expected.msl
index 1cc175ac54..d00308107d 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.msl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct Inner {
   /* 0x0000 */ int x;
 };
diff --git a/test/bug/tint/1113.wgsl.expected.msl b/test/bug/tint/1113.wgsl.expected.msl
index 1eaadd3dee..663d47fc9b 100644
--- a/test/bug/tint/1113.wgsl.expected.msl
+++ b/test/bug/tint/1113.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct Uniforms {
   /* 0x0000 */ uint numTriangles;
   /* 0x0004 */ uint gridSize;
diff --git a/test/bug/tint/1121.wgsl b/test/bug/tint/1121.wgsl
new file mode 100644
index 0000000000..75425d1154
--- /dev/null
+++ b/test/bug/tint/1121.wgsl
@@ -0,0 +1,127 @@
+// Take from here:
+// https://github.com/shrekshao/webgpu-deferred-renderer/blob/4f8bf0910793100aa8d60dbd1319bddb5357b1fa/renderer/LightCulling.js
+// With these token replacements:
+//   $NUM_TILE_LIGHT_SLOT = 64
+//   $NUM_TILES = 4
+//   $TILE_COUNT_Y = 2
+//   $TILE_COUNT_X = 2
+//   $TILE_SIZE = 16
+
+struct LightData {
+    position : vec4<f32>;
+    color : vec3<f32>;
+    radius : f32;
+};
+[[block]] struct LightsBuffer {
+    lights: array<LightData>;
+};
+[[group(0), binding(0)]] var<storage, read_write> lightsBuffer: LightsBuffer;
+struct TileLightIdData {
+    count: atomic<u32>;
+    lightId: array<u32, 64>;
+};
+[[block]] struct Tiles {
+    data: array<TileLightIdData, 4>;
+};
+[[group(1), binding(0)]] var<storage, read_write> tileLightId: Tiles;
+  
+[[block]] struct Config {
+    numLights : u32;
+    numTiles : u32;
+    tileCountX : u32;
+    tileCountY : u32;
+    numTileLightSlot : u32;
+    tileSize : u32;
+};
+[[group(2), binding(0)]] var<uniform> config: Config;
+[[block]] struct Uniforms {
+    min : vec4<f32>;
+    max : vec4<f32>;
+    // camera
+    viewMatrix : mat4x4<f32>;
+    projectionMatrix : mat4x4<f32>;
+    // Tile info
+    fullScreenSize : vec4<f32>;    // width, height
+};
+[[group(3), binding(0)]] var<uniform> uniforms: Uniforms;
+[[stage(compute), workgroup_size(64, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+    var index = GlobalInvocationID.x;
+    if (index >= config.numLights) {
+        return;
+    }
+    // Light position updating
+    lightsBuffer.lights[index].position.y = lightsBuffer.lights[index].position.y - 0.1 + 0.001 * (f32(index) - 64.0 * floor(f32(index) / 64.0));
+  
+    if (lightsBuffer.lights[index].position.y < uniforms.min.y) {
+        lightsBuffer.lights[index].position.y = uniforms.max.y;
+    }
+    // Light culling
+    // Implementation here is Tiled without per tile min-max depth
+    // You could also implement cluster culling
+    // Feel free to add more compute passes if necessary
+    // some math reference: http://www.txutxi.com/?p=444
+    var M: mat4x4<f32> = uniforms.projectionMatrix;
+    var viewNear: f32 = - M[3][2] / ( -1.0 + M[2][2]);
+    var viewFar: f32 = - M[3][2] / (1.0 + M[2][2]);
+    var lightPos = lightsBuffer.lights[index].position;
+    lightPos = uniforms.viewMatrix * lightPos;
+    lightPos = lightPos / lightPos.w;
+    var lightRadius: f32 = lightsBuffer.lights[index].radius;
+    var boxMin: vec4<f32> = lightPos - vec4<f32>(vec3<f32>(lightRadius), 0.0);
+    var boxMax: vec4<f32> = lightPos + vec4<f32>(vec3<f32>(lightRadius), 0.0);
+    var frustumPlanes: array<vec4<f32>, 6>;
+    frustumPlanes[4] = vec4<f32>(0.0, 0.0, -1.0, viewNear);    // near
+    frustumPlanes[5] = vec4<f32>(0.0, 0.0, 1.0, -viewFar);    // far
+    let TILE_SIZE: i32 = 16;
+    let TILE_COUNT_X: i32 = 2;
+    let TILE_COUNT_Y: i32 = 2;
+    for (var y : i32 = 0; y < TILE_COUNT_Y; y = y + 1) {
+        for (var x : i32 = 0; x < TILE_COUNT_X; x = x + 1) {
+            var tilePixel0Idx : vec2<i32> = vec2<i32>(x * TILE_SIZE, y * TILE_SIZE);
+            // tile position in NDC space
+            var floorCoord: vec2<f32> = 2.0 * vec2<f32>(tilePixel0Idx) / uniforms.fullScreenSize.xy - vec2<f32>(1.0);  // -1, 1
+            var ceilCoord: vec2<f32> = 2.0 * vec2<f32>(tilePixel0Idx + vec2<i32>(TILE_SIZE)) / uniforms.fullScreenSize.xy - vec2<f32>(1.0);  // -1, 1
+            var viewFloorCoord: vec2<f32> = vec2<f32>( (- viewNear * floorCoord.x - M[2][0] * viewNear) / M[0][0] , (- viewNear * floorCoord.y - M[2][1] * viewNear) / M[1][1] );
+            var viewCeilCoord: vec2<f32> = vec2<f32>( (- viewNear * ceilCoord.x - M[2][0] * viewNear) / M[0][0] , (- viewNear * ceilCoord.y - M[2][1] * viewNear) / M[1][1] );
+            frustumPlanes[0] = vec4<f32>(1.0, 0.0, - viewFloorCoord.x / viewNear, 0.0);       // left
+            frustumPlanes[1] = vec4<f32>(-1.0, 0.0, viewCeilCoord.x / viewNear, 0.0);   // right
+            frustumPlanes[2] = vec4<f32>(0.0, 1.0, - viewFloorCoord.y / viewNear, 0.0);       // bottom
+            frustumPlanes[3] = vec4<f32>(0.0, -1.0, viewCeilCoord.y / viewNear, 0.0);   // top
+            var dp: f32 = 0.0;  // dot product
+            for (var i: u32 = 0u; i < 6u; i = i + 1u)
+            {
+                var p: vec4<f32>;
+                if (frustumPlanes[i].x > 0.0) {
+                    p.x = boxMax.x;
+                } else {
+                    p.x = boxMin.x;
+                }
+                if (frustumPlanes[i].y > 0.0) {
+                    p.y = boxMax.y;
+                } else {
+                    p.y = boxMin.y;
+                }
+                if (frustumPlanes[i].z > 0.0) {
+                    p.z = boxMax.z;
+                } else {
+                    p.z = boxMin.z;
+                }
+                p.w = 1.0;
+                dp = dp + min(0.0, dot(p, frustumPlanes[i]));
+            }
+            if (dp >= 0.0) {
+                // light is overlapping with the tile
+                var tileId: u32 = u32(x + y * TILE_COUNT_X);
+                if (tileId < 0u || tileId >= config.numTiles) {
+                    continue;
+                }
+                var offset: u32 = atomicAdd(&(tileLightId.data[tileId].count), 1u);
+                if (offset >= config.numTileLightSlot) {
+                    continue;
+                }
+                tileLightId.data[tileId].lightId[offset] = GlobalInvocationID.x;
+            }
+        }
+    }
+}
diff --git a/test/bug/tint/1121.wgsl.expected.hlsl b/test/bug/tint/1121.wgsl.expected.hlsl
new file mode 100644
index 0000000000..12d8865772
--- /dev/null
+++ b/test/bug/tint/1121.wgsl.expected.hlsl
@@ -0,0 +1,115 @@
+uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+RWByteAddressBuffer lightsBuffer : register(u0, space0);
+
+RWByteAddressBuffer tileLightId : register(u0, space1);
+
+cbuffer cbuffer_config : register(b0, space2) {
+  uint4 config[2];
+};
+
+cbuffer cbuffer_uniforms : register(b0, space3) {
+  uint4 uniforms[11];
+};
+
+struct tint_symbol_1 {
+  uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+float4x4 tint_symbol_6(uint4 buffer[11], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  const uint scalar_offset_3 = ((offset + 48u)) / 4;
+  return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+void main_inner(uint3 GlobalInvocationID) {
+  uint index = GlobalInvocationID.x;
+  if ((index >= config[0].x)) {
+    return;
+  }
+  lightsBuffer.Store(((32u * index) + 4u), asuint(((asfloat(lightsBuffer.Load(((32u * index) + 4u))) - 0.100000001f) + (0.001f * (float(index) - (64.0f * floor((float(index) / 64.0f))))))));
+  if ((asfloat(lightsBuffer.Load(((32u * index) + 4u))) < asfloat(uniforms[0].y))) {
+    lightsBuffer.Store(((32u * index) + 4u), asuint(asfloat(uniforms[1].y)));
+  }
+  float4x4 M = tint_symbol_6(uniforms, 96u);
+  float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
+  float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
+  float4 lightPos = asfloat(lightsBuffer.Load4((32u * index)));
+  lightPos = mul(lightPos, tint_symbol_6(uniforms, 32u));
+  lightPos = (lightPos / lightPos.w);
+  float lightRadius = asfloat(lightsBuffer.Load(((32u * index) + 28u)));
+  float4 boxMin = (lightPos - float4(float3((lightRadius).xxx), 0.0f));
+  float4 boxMax = (lightPos + float4(float3((lightRadius).xxx), 0.0f));
+  float4 frustumPlanes[6] = (float4[6])0;
+  frustumPlanes[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
+  frustumPlanes[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
+  const int TILE_SIZE = 16;
+  const int TILE_COUNT_X = 2;
+  {
+    for(int y_1 = 0; (y_1 < 2); y_1 = (y_1 + 1)) {
+      {
+        for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = (x_1 + 1)) {
+          int2 tilePixel0Idx = int2((x_1 * TILE_SIZE), (y_1 * TILE_SIZE));
+          float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / asfloat(uniforms[10]).xy) - float2((1.0f).xx));
+          float2 ceilCoord = (((2.0f * float2((tilePixel0Idx + int2((TILE_SIZE).xx)))) / asfloat(uniforms[10]).xy) - float2((1.0f).xx));
+          float2 viewFloorCoord = float2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
+          float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
+          frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord.x) / viewNear), 0.0f);
+          frustumPlanes[1] = float4(-1.0f, 0.0f, (viewCeilCoord.x / viewNear), 0.0f);
+          frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord.y) / viewNear), 0.0f);
+          frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord.y / viewNear), 0.0f);
+          float dp = 0.0f;
+          {
+            for(uint i = 0u; (i < 6u); i = (i + 1u)) {
+              float4 p = float4(0.0f, 0.0f, 0.0f, 0.0f);
+              if ((frustumPlanes[i].x > 0.0f)) {
+                p.x = boxMax.x;
+              } else {
+                p.x = boxMin.x;
+              }
+              if ((frustumPlanes[i].y > 0.0f)) {
+                p.y = boxMax.y;
+              } else {
+                p.y = boxMin.y;
+              }
+              if ((frustumPlanes[i].z > 0.0f)) {
+                p.z = boxMax.z;
+              } else {
+                p.z = boxMin.z;
+              }
+              p.w = 1.0f;
+              dp = (dp + min(0.0f, dot(p, frustumPlanes[i])));
+            }
+          }
+          if ((dp >= 0.0f)) {
+            uint tileId = uint((x_1 + (y_1 * TILE_COUNT_X)));
+            bool tint_tmp = (tileId < 0u);
+            if (!tint_tmp) {
+              tint_tmp = (tileId >= config[0].y);
+            }
+            if ((tint_tmp)) {
+              continue;
+            }
+            uint offset = atomicAdd_1(tileLightId, (260u * tileId), 1u);
+            if ((offset >= config[1].x)) {
+              continue;
+            }
+            tileLightId.Store((((260u * tileId) + 4u) + (4u * offset)), asuint(GlobalInvocationID.x));
+          }
+        }
+      }
+    }
+  }
+}
+
+[numthreads(64, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.GlobalInvocationID);
+  return;
+}
diff --git a/test/bug/tint/1121.wgsl.expected.msl b/test/bug/tint/1121.wgsl.expected.msl
new file mode 100644
index 0000000000..c7eb13cd87
--- /dev/null
+++ b/test/bug/tint/1121.wgsl.expected.msl
@@ -0,0 +1,130 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct LightData {
+  /* 0x0000 */ float4 position;
+  /* 0x0010 */ packed_float3 color;
+  /* 0x001c */ float radius;
+};
+struct LightsBuffer {
+  /* 0x0000 */ LightData lights[1];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ uint arr[64];
+};
+struct TileLightIdData {
+  /* 0x0000 */ atomic_uint count;
+  /* 0x0004 */ tint_array_wrapper lightId;
+};
+struct tint_array_wrapper_1 {
+  /* 0x0000 */ TileLightIdData arr[4];
+};
+struct Tiles {
+  /* 0x0000 */ tint_array_wrapper_1 data;
+};
+struct Config {
+  /* 0x0000 */ uint numLights;
+  /* 0x0004 */ uint numTiles;
+  /* 0x0008 */ uint tileCountX;
+  /* 0x000c */ uint tileCountY;
+  /* 0x0010 */ uint numTileLightSlot;
+  /* 0x0014 */ uint tileSize;
+};
+struct Uniforms {
+  /* 0x0000 */ float4 min;
+  /* 0x0010 */ float4 max;
+  /* 0x0020 */ float4x4 viewMatrix;
+  /* 0x0060 */ float4x4 projectionMatrix;
+  /* 0x00a0 */ float4 fullScreenSize;
+};
+struct tint_array_wrapper_2 {
+  float4 arr[6];
+};
+
+void tint_symbol_inner(constant Config& config, constant Uniforms& uniforms, device LightsBuffer& lightsBuffer, device Tiles& tileLightId, uint3 GlobalInvocationID) {
+  uint index = GlobalInvocationID.x;
+  if ((index >= config.numLights)) {
+    return;
+  }
+  lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.100000001f) + (0.001f * (float(index) - (64.0f * floor((float(index) / 64.0f))))));
+  if ((lightsBuffer.lights[index].position.y < uniforms.min.y)) {
+    lightsBuffer.lights[index].position.y = uniforms.max.y;
+  }
+  float4x4 M = uniforms.projectionMatrix;
+  float viewNear = (-(M[3][2]) / (-1.0f + M[2][2]));
+  float viewFar = (-(M[3][2]) / (1.0f + M[2][2]));
+  float4 lightPos = lightsBuffer.lights[index].position;
+  lightPos = (uniforms.viewMatrix * lightPos);
+  lightPos = (lightPos / lightPos.w);
+  float lightRadius = lightsBuffer.lights[index].radius;
+  float4 boxMin = (lightPos - float4(float3(lightRadius), 0.0f));
+  float4 boxMax = (lightPos + float4(float3(lightRadius), 0.0f));
+  tint_array_wrapper_2 frustumPlanes = {};
+  frustumPlanes.arr[4] = float4(0.0f, 0.0f, -1.0f, viewNear);
+  frustumPlanes.arr[5] = float4(0.0f, 0.0f, 1.0f, -(viewFar));
+  int const TILE_SIZE = 16;
+  int const TILE_COUNT_X = 2;
+  int const TILE_COUNT_Y = 2;
+  for(int y_1 = 0; (y_1 < TILE_COUNT_Y); y_1 = as_type<int>((as_type<uint>(y_1) + as_type<uint>(1)))) {
+    for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
+      int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
+      float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - float2(1.0f));
+      float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / uniforms.fullScreenSize.xy) - float2(1.0f));
+      float2 viewFloorCoord = float2((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
+      float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
+      frustumPlanes.arr[0] = float4(1.0f, 0.0f, (-(viewFloorCoord.x) / viewNear), 0.0f);
+      frustumPlanes.arr[1] = float4(-1.0f, 0.0f, (viewCeilCoord.x / viewNear), 0.0f);
+      frustumPlanes.arr[2] = float4(0.0f, 1.0f, (-(viewFloorCoord.y) / viewNear), 0.0f);
+      frustumPlanes.arr[3] = float4(0.0f, -1.0f, (viewCeilCoord.y / viewNear), 0.0f);
+      float dp = 0.0f;
+      for(uint i = 0u; (i < 6u); i = (i + 1u)) {
+        float4 p = 0.0f;
+        if ((frustumPlanes.arr[i].x > 0.0f)) {
+          p.x = boxMax.x;
+        } else {
+          p.x = boxMin.x;
+        }
+        if ((frustumPlanes.arr[i].y > 0.0f)) {
+          p.y = boxMax.y;
+        } else {
+          p.y = boxMin.y;
+        }
+        if ((frustumPlanes.arr[i].z > 0.0f)) {
+          p.z = boxMax.z;
+        } else {
+          p.z = boxMin.z;
+        }
+        p.w = 1.0f;
+        dp = (dp + fmin(0.0f, dot(p, frustumPlanes.arr[i])));
+      }
+      if ((dp >= 0.0f)) {
+        uint tileId = uint(as_type<int>((as_type<uint>(x_1) + as_type<uint>(as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_COUNT_X)))))));
+        if (((tileId < 0u) || (tileId >= config.numTiles))) {
+          continue;
+        }
+        uint offset = atomic_fetch_add_explicit(&(tileLightId.data.arr[tileId].count), 1u, memory_order_relaxed);
+        if ((offset >= config.numTileLightSlot)) {
+          continue;
+        }
+        tileLightId.data.arr[tileId].lightId.arr[offset] = GlobalInvocationID.x;
+      }
+    }
+  }
+}
+
+kernel void tint_symbol(uint3 GlobalInvocationID [[thread_position_in_grid]], constant Config& config [[buffer(0)]], constant Uniforms& uniforms [[buffer(1)]], device LightsBuffer& lightsBuffer [[buffer(2)]], device Tiles& tileLightId [[buffer(3)]]) {
+  tint_symbol_inner(config, uniforms, lightsBuffer, tileLightId, GlobalInvocationID);
+  return;
+}
+
diff --git a/test/bug/tint/1121.wgsl.expected.spvasm b/test/bug/tint/1121.wgsl.expected.spvasm
new file mode 100644
index 0000000000..6d3354b7da
--- /dev/null
+++ b/test/bug/tint/1121.wgsl.expected.spvasm
@@ -0,0 +1,617 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 417
+; Schema: 0
+               OpCapability Shader
+         %60 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %GlobalInvocationID_1
+               OpExecutionMode %main LocalSize 64 1 1
+               OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
+               OpName %LightsBuffer "LightsBuffer"
+               OpMemberName %LightsBuffer 0 "lights"
+               OpName %LightData "LightData"
+               OpMemberName %LightData 0 "position"
+               OpMemberName %LightData 1 "color"
+               OpMemberName %LightData 2 "radius"
+               OpName %lightsBuffer "lightsBuffer"
+               OpName %Tiles "Tiles"
+               OpMemberName %Tiles 0 "data"
+               OpName %TileLightIdData "TileLightIdData"
+               OpMemberName %TileLightIdData 0 "count"
+               OpMemberName %TileLightIdData 1 "lightId"
+               OpName %tileLightId "tileLightId"
+               OpName %Config "Config"
+               OpMemberName %Config 0 "numLights"
+               OpMemberName %Config 1 "numTiles"
+               OpMemberName %Config 2 "tileCountX"
+               OpMemberName %Config 3 "tileCountY"
+               OpMemberName %Config 4 "numTileLightSlot"
+               OpMemberName %Config 5 "tileSize"
+               OpName %config "config"
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "min"
+               OpMemberName %Uniforms 1 "max"
+               OpMemberName %Uniforms 2 "viewMatrix"
+               OpMemberName %Uniforms 3 "projectionMatrix"
+               OpMemberName %Uniforms 4 "fullScreenSize"
+               OpName %uniforms "uniforms"
+               OpName %main_inner "main_inner"
+               OpName %GlobalInvocationID "GlobalInvocationID"
+               OpName %index "index"
+               OpName %M "M"
+               OpName %viewNear "viewNear"
+               OpName %viewFar "viewFar"
+               OpName %lightPos "lightPos"
+               OpName %lightRadius "lightRadius"
+               OpName %boxMin "boxMin"
+               OpName %boxMax "boxMax"
+               OpName %frustumPlanes "frustumPlanes"
+               OpName %y "y"
+               OpName %x "x"
+               OpName %tilePixel0Idx "tilePixel0Idx"
+               OpName %floorCoord "floorCoord"
+               OpName %ceilCoord "ceilCoord"
+               OpName %viewFloorCoord "viewFloorCoord"
+               OpName %viewCeilCoord "viewCeilCoord"
+               OpName %dp "dp"
+               OpName %i "i"
+               OpName %p "p"
+               OpName %tileId "tileId"
+               OpName %offset "offset"
+               OpName %main "main"
+               OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
+               OpDecorate %LightsBuffer Block
+               OpMemberDecorate %LightsBuffer 0 Offset 0
+               OpMemberDecorate %LightData 0 Offset 0
+               OpMemberDecorate %LightData 1 Offset 16
+               OpMemberDecorate %LightData 2 Offset 28
+               OpDecorate %_runtimearr_LightData ArrayStride 32
+               OpDecorate %lightsBuffer DescriptorSet 0
+               OpDecorate %lightsBuffer Binding 0
+               OpDecorate %Tiles Block
+               OpMemberDecorate %Tiles 0 Offset 0
+               OpMemberDecorate %TileLightIdData 0 Offset 0
+               OpMemberDecorate %TileLightIdData 1 Offset 4
+               OpDecorate %_arr_uint_uint_64 ArrayStride 4
+               OpDecorate %_arr_TileLightIdData_uint_4 ArrayStride 260
+               OpDecorate %tileLightId DescriptorSet 1
+               OpDecorate %tileLightId Binding 0
+               OpDecorate %Config Block
+               OpMemberDecorate %Config 0 Offset 0
+               OpMemberDecorate %Config 1 Offset 4
+               OpMemberDecorate %Config 2 Offset 8
+               OpMemberDecorate %Config 3 Offset 12
+               OpMemberDecorate %Config 4 Offset 16
+               OpMemberDecorate %Config 5 Offset 20
+               OpDecorate %config NonWritable
+               OpDecorate %config DescriptorSet 2
+               OpDecorate %config Binding 0
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 16
+               OpMemberDecorate %Uniforms 2 Offset 32
+               OpMemberDecorate %Uniforms 2 ColMajor
+               OpMemberDecorate %Uniforms 2 MatrixStride 16
+               OpMemberDecorate %Uniforms 3 Offset 96
+               OpMemberDecorate %Uniforms 3 ColMajor
+               OpMemberDecorate %Uniforms 3 MatrixStride 16
+               OpMemberDecorate %Uniforms 4 Offset 160
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 3
+               OpDecorate %uniforms Binding 0
+               OpDecorate %_arr_v4float_uint_6 ArrayStride 16
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+    %v3float = OpTypeVector %float 3
+  %LightData = OpTypeStruct %v4float %v3float %float
+%_runtimearr_LightData = OpTypeRuntimeArray %LightData
+%LightsBuffer = OpTypeStruct %_runtimearr_LightData
+%_ptr_StorageBuffer_LightsBuffer = OpTypePointer StorageBuffer %LightsBuffer
+%lightsBuffer = OpVariable %_ptr_StorageBuffer_LightsBuffer StorageBuffer
+    %uint_64 = OpConstant %uint 64
+%_arr_uint_uint_64 = OpTypeArray %uint %uint_64
+%TileLightIdData = OpTypeStruct %uint %_arr_uint_uint_64
+     %uint_4 = OpConstant %uint 4
+%_arr_TileLightIdData_uint_4 = OpTypeArray %TileLightIdData %uint_4
+      %Tiles = OpTypeStruct %_arr_TileLightIdData_uint_4
+%_ptr_StorageBuffer_Tiles = OpTypePointer StorageBuffer %Tiles
+%tileLightId = OpVariable %_ptr_StorageBuffer_Tiles StorageBuffer
+     %Config = OpTypeStruct %uint %uint %uint %uint %uint %uint
+%_ptr_Uniform_Config = OpTypePointer Uniform %Config
+     %config = OpVariable %_ptr_Uniform_Config Uniform
+%mat4v4float = OpTypeMatrix %v4float 4
+   %Uniforms = OpTypeStruct %v4float %v4float %mat4v4float %mat4v4float %v4float
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+         %28 = OpTypeFunction %void %v3uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %36 = OpConstantNull %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+%float_0_100000001 = OpConstant %float 0.100000001
+%float_0_00100000005 = OpConstant %float 0.00100000005
+   %float_64 = OpConstant %float 64
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+     %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+%_ptr_Function_mat4v4float = OpTypePointer Function %mat4v4float
+         %87 = OpConstantNull %mat4v4float
+        %int = OpTypeInt 32 1
+      %int_3 = OpConstant %int 3
+      %int_2 = OpConstant %int 2
+%_ptr_Function_float = OpTypePointer Function %float
+   %float_n1 = OpConstant %float -1
+        %101 = OpConstantNull %float
+    %float_1 = OpConstant %float 1
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+        %117 = OpConstantNull %v4float
+     %uint_2 = OpConstant %uint 2
+    %float_0 = OpConstant %float 0
+     %uint_6 = OpConstant %uint 6
+%_arr_v4float_uint_6 = OpTypeArray %v4float %uint_6
+%_ptr_Function__arr_v4float_uint_6 = OpTypePointer Function %_arr_v4float_uint_6
+        %156 = OpConstantNull %_arr_v4float_uint_6
+      %int_4 = OpConstant %int 4
+      %int_5 = OpConstant %int 5
+     %int_16 = OpConstant %int 16
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+        %170 = OpConstantNull %int
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+        %198 = OpConstantNull %v2int
+    %float_2 = OpConstant %float 2
+    %v2float = OpTypeVector %float 2
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+        %209 = OpConstantComposite %v2float %float_1 %float_1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+        %213 = OpConstantNull %v2float
+      %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_uint_0 = OpTypePointer StorageBuffer %uint
+        %412 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %28
+%GlobalInvocationID = OpFunctionParameter %v3uint
+         %32 = OpLabel
+      %index = OpVariable %_ptr_Function_uint Function %36
+          %M = OpVariable %_ptr_Function_mat4v4float Function %87
+   %viewNear = OpVariable %_ptr_Function_float Function %101
+    %viewFar = OpVariable %_ptr_Function_float Function %101
+   %lightPos = OpVariable %_ptr_Function_v4float Function %117
+        %127 = OpVariable %_ptr_Function_v4float Function %117
+%lightRadius = OpVariable %_ptr_Function_float Function %101
+     %boxMin = OpVariable %_ptr_Function_v4float Function %117
+     %boxMax = OpVariable %_ptr_Function_v4float Function %117
+%frustumPlanes = OpVariable %_ptr_Function__arr_v4float_uint_6 Function %156
+          %y = OpVariable %_ptr_Function_int Function %170
+          %x = OpVariable %_ptr_Function_int Function %170
+%tilePixel0Idx = OpVariable %_ptr_Function_v2int Function %198
+ %floorCoord = OpVariable %_ptr_Function_v2float Function %213
+  %ceilCoord = OpVariable %_ptr_Function_v2float Function %213
+%viewFloorCoord = OpVariable %_ptr_Function_v2float Function %213
+%viewCeilCoord = OpVariable %_ptr_Function_v2float Function %213
+         %dp = OpVariable %_ptr_Function_float Function %101
+          %i = OpVariable %_ptr_Function_uint Function %36
+          %p = OpVariable %_ptr_Function_v4float Function %117
+     %tileId = OpVariable %_ptr_Function_uint Function %36
+     %offset = OpVariable %_ptr_Function_uint Function %36
+         %33 = OpCompositeExtract %uint %GlobalInvocationID 0
+               OpStore %index %33
+         %37 = OpLoad %uint %index
+         %40 = OpAccessChain %_ptr_Uniform_uint %config %uint_0
+         %41 = OpLoad %uint %40
+         %42 = OpUGreaterThanEqual %bool %37 %41
+               OpSelectionMerge %44 None
+               OpBranchConditional %42 %45 %44
+         %45 = OpLabel
+               OpReturn
+         %44 = OpLabel
+         %46 = OpLoad %uint %index
+         %49 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %46 %uint_0 %uint_1
+         %50 = OpLoad %uint %index
+         %51 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %50 %uint_0 %uint_1
+         %52 = OpLoad %float %51
+         %54 = OpFSub %float %52 %float_0_100000001
+         %57 = OpLoad %uint %index
+         %56 = OpConvertUToF %float %57
+         %62 = OpLoad %uint %index
+         %61 = OpConvertUToF %float %62
+         %63 = OpFDiv %float %61 %float_64
+         %59 = OpExtInst %float %60 Floor %63
+         %64 = OpFMul %float %float_64 %59
+         %65 = OpFSub %float %56 %64
+         %66 = OpFMul %float %float_0_00100000005 %65
+         %67 = OpFAdd %float %54 %66
+               OpStore %49 %67
+         %68 = OpLoad %uint %index
+         %69 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %68 %uint_0 %uint_1
+         %70 = OpLoad %float %69
+         %72 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_1
+         %73 = OpLoad %float %72
+         %74 = OpFOrdLessThan %bool %70 %73
+               OpSelectionMerge %75 None
+               OpBranchConditional %74 %76 %75
+         %76 = OpLabel
+         %77 = OpLoad %uint %index
+         %78 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %77 %uint_0 %uint_1
+         %79 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_1 %uint_1
+         %80 = OpLoad %float %79
+               OpStore %78 %80
+               OpBranch %75
+         %75 = OpLabel
+         %83 = OpAccessChain %_ptr_Uniform_mat4v4float %uniforms %uint_3
+         %84 = OpLoad %mat4v4float %83
+               OpStore %M %84
+         %93 = OpAccessChain %_ptr_Function_float %M %int_3 %int_2
+         %94 = OpLoad %float %93
+         %88 = OpFNegate %float %94
+         %96 = OpAccessChain %_ptr_Function_float %M %int_2 %int_2
+         %97 = OpLoad %float %96
+         %98 = OpFAdd %float %float_n1 %97
+         %99 = OpFDiv %float %88 %98
+               OpStore %viewNear %99
+        %103 = OpAccessChain %_ptr_Function_float %M %int_3 %int_2
+        %104 = OpLoad %float %103
+        %102 = OpFNegate %float %104
+        %106 = OpAccessChain %_ptr_Function_float %M %int_2 %int_2
+        %107 = OpLoad %float %106
+        %108 = OpFAdd %float %float_1 %107
+        %109 = OpFDiv %float %102 %108
+               OpStore %viewFar %109
+        %111 = OpLoad %uint %index
+        %113 = OpAccessChain %_ptr_StorageBuffer_v4float %lightsBuffer %uint_0 %111 %uint_0
+        %114 = OpLoad %v4float %113
+               OpStore %lightPos %114
+        %119 = OpAccessChain %_ptr_Uniform_mat4v4float %uniforms %uint_2
+        %120 = OpLoad %mat4v4float %119
+        %121 = OpLoad %v4float %lightPos
+        %122 = OpMatrixTimesVector %v4float %120 %121
+               OpStore %lightPos %122
+        %123 = OpLoad %v4float %lightPos
+        %124 = OpAccessChain %_ptr_Function_float %lightPos %uint_3
+        %125 = OpLoad %float %124
+        %128 = OpCompositeConstruct %v4float %125 %125 %125 %125
+        %126 = OpFDiv %v4float %123 %128
+               OpStore %lightPos %126
+        %129 = OpLoad %uint %index
+        %130 = OpAccessChain %_ptr_StorageBuffer_float %lightsBuffer %uint_0 %129 %uint_2
+        %131 = OpLoad %float %130
+               OpStore %lightRadius %131
+        %133 = OpLoad %v4float %lightPos
+        %134 = OpLoad %float %lightRadius
+        %135 = OpCompositeConstruct %v3float %134 %134 %134
+        %136 = OpCompositeExtract %float %135 0
+        %137 = OpCompositeExtract %float %135 1
+        %138 = OpCompositeExtract %float %135 2
+        %140 = OpCompositeConstruct %v4float %136 %137 %138 %float_0
+        %141 = OpFSub %v4float %133 %140
+               OpStore %boxMin %141
+        %143 = OpLoad %v4float %lightPos
+        %144 = OpLoad %float %lightRadius
+        %145 = OpCompositeConstruct %v3float %144 %144 %144
+        %146 = OpCompositeExtract %float %145 0
+        %147 = OpCompositeExtract %float %145 1
+        %148 = OpCompositeExtract %float %145 2
+        %149 = OpCompositeConstruct %v4float %146 %147 %148 %float_0
+        %150 = OpFAdd %v4float %143 %149
+               OpStore %boxMax %150
+        %158 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_4
+        %159 = OpLoad %float %viewNear
+        %160 = OpCompositeConstruct %v4float %float_0 %float_0 %float_n1 %159
+               OpStore %158 %160
+        %162 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_5
+        %164 = OpLoad %float %viewFar
+        %163 = OpFNegate %float %164
+        %165 = OpCompositeConstruct %v4float %float_0 %float_0 %float_1 %163
+               OpStore %162 %165
+               OpStore %y %int_0
+               OpBranch %171
+        %171 = OpLabel
+               OpLoopMerge %172 %173 None
+               OpBranch %174
+        %174 = OpLabel
+        %176 = OpLoad %int %y
+        %177 = OpSLessThan %bool %176 %int_2
+        %175 = OpLogicalNot %bool %177
+               OpSelectionMerge %178 None
+               OpBranchConditional %175 %179 %178
+        %179 = OpLabel
+               OpBranch %172
+        %178 = OpLabel
+               OpStore %x %int_0
+               OpBranch %181
+        %181 = OpLabel
+               OpLoopMerge %182 %183 None
+               OpBranch %184
+        %184 = OpLabel
+        %186 = OpLoad %int %x
+        %187 = OpSLessThan %bool %186 %int_2
+        %185 = OpLogicalNot %bool %187
+               OpSelectionMerge %188 None
+               OpBranchConditional %185 %189 %188
+        %189 = OpLabel
+               OpBranch %182
+        %188 = OpLabel
+        %191 = OpLoad %int %x
+        %192 = OpIMul %int %191 %int_16
+        %193 = OpLoad %int %y
+        %194 = OpIMul %int %193 %int_16
+        %195 = OpCompositeConstruct %v2int %192 %194
+               OpStore %tilePixel0Idx %195
+        %202 = OpLoad %v2int %tilePixel0Idx
+        %200 = OpConvertSToF %v2float %202
+        %203 = OpVectorTimesScalar %v2float %200 %float_2
+        %205 = OpAccessChain %_ptr_Uniform_v4float %uniforms %uint_4
+        %206 = OpLoad %v4float %205
+        %207 = OpVectorShuffle %v2float %206 %206 0 1
+        %208 = OpFDiv %v2float %203 %207
+        %210 = OpFSub %v2float %208 %209
+               OpStore %floorCoord %210
+        %215 = OpLoad %v2int %tilePixel0Idx
+        %216 = OpCompositeConstruct %v2int %int_16 %int_16
+        %217 = OpIAdd %v2int %215 %216
+        %214 = OpConvertSToF %v2float %217
+        %218 = OpVectorTimesScalar %v2float %214 %float_2
+        %219 = OpAccessChain %_ptr_Uniform_v4float %uniforms %uint_4
+        %220 = OpLoad %v4float %219
+        %221 = OpVectorShuffle %v2float %220 %220 0 1
+        %222 = OpFDiv %v2float %218 %221
+        %223 = OpFSub %v2float %222 %209
+               OpStore %ceilCoord %223
+        %226 = OpLoad %float %viewNear
+        %225 = OpFNegate %float %226
+        %227 = OpAccessChain %_ptr_Function_float %floorCoord %uint_0
+        %228 = OpLoad %float %227
+        %229 = OpFMul %float %225 %228
+        %230 = OpAccessChain %_ptr_Function_float %M %int_2 %int_0
+        %231 = OpLoad %float %230
+        %232 = OpLoad %float %viewNear
+        %233 = OpFMul %float %231 %232
+        %234 = OpFSub %float %229 %233
+        %235 = OpAccessChain %_ptr_Function_float %M %int_0 %int_0
+        %236 = OpLoad %float %235
+        %237 = OpFDiv %float %234 %236
+        %239 = OpLoad %float %viewNear
+        %238 = OpFNegate %float %239
+        %240 = OpAccessChain %_ptr_Function_float %floorCoord %uint_1
+        %241 = OpLoad %float %240
+        %242 = OpFMul %float %238 %241
+        %244 = OpAccessChain %_ptr_Function_float %M %int_2 %int_1
+        %245 = OpLoad %float %244
+        %246 = OpLoad %float %viewNear
+        %247 = OpFMul %float %245 %246
+        %248 = OpFSub %float %242 %247
+        %249 = OpAccessChain %_ptr_Function_float %M %int_1 %int_1
+        %250 = OpLoad %float %249
+        %251 = OpFDiv %float %248 %250
+        %252 = OpCompositeConstruct %v2float %237 %251
+               OpStore %viewFloorCoord %252
+        %255 = OpLoad %float %viewNear
+        %254 = OpFNegate %float %255
+        %256 = OpAccessChain %_ptr_Function_float %ceilCoord %uint_0
+        %257 = OpLoad %float %256
+        %258 = OpFMul %float %254 %257
+        %259 = OpAccessChain %_ptr_Function_float %M %int_2 %int_0
+        %260 = OpLoad %float %259
+        %261 = OpLoad %float %viewNear
+        %262 = OpFMul %float %260 %261
+        %263 = OpFSub %float %258 %262
+        %264 = OpAccessChain %_ptr_Function_float %M %int_0 %int_0
+        %265 = OpLoad %float %264
+        %266 = OpFDiv %float %263 %265
+        %268 = OpLoad %float %viewNear
+        %267 = OpFNegate %float %268
+        %269 = OpAccessChain %_ptr_Function_float %ceilCoord %uint_1
+        %270 = OpLoad %float %269
+        %271 = OpFMul %float %267 %270
+        %272 = OpAccessChain %_ptr_Function_float %M %int_2 %int_1
+        %273 = OpLoad %float %272
+        %274 = OpLoad %float %viewNear
+        %275 = OpFMul %float %273 %274
+        %276 = OpFSub %float %271 %275
+        %277 = OpAccessChain %_ptr_Function_float %M %int_1 %int_1
+        %278 = OpLoad %float %277
+        %279 = OpFDiv %float %276 %278
+        %280 = OpCompositeConstruct %v2float %266 %279
+               OpStore %viewCeilCoord %280
+        %282 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_0
+        %284 = OpAccessChain %_ptr_Function_float %viewFloorCoord %uint_0
+        %285 = OpLoad %float %284
+        %283 = OpFNegate %float %285
+        %286 = OpLoad %float %viewNear
+        %287 = OpFDiv %float %283 %286
+        %288 = OpCompositeConstruct %v4float %float_1 %float_0 %287 %float_0
+               OpStore %282 %288
+        %289 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_1
+        %290 = OpAccessChain %_ptr_Function_float %viewCeilCoord %uint_0
+        %291 = OpLoad %float %290
+        %292 = OpLoad %float %viewNear
+        %293 = OpFDiv %float %291 %292
+        %294 = OpCompositeConstruct %v4float %float_n1 %float_0 %293 %float_0
+               OpStore %289 %294
+        %295 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_2
+        %297 = OpAccessChain %_ptr_Function_float %viewFloorCoord %uint_1
+        %298 = OpLoad %float %297
+        %296 = OpFNegate %float %298
+        %299 = OpLoad %float %viewNear
+        %300 = OpFDiv %float %296 %299
+        %301 = OpCompositeConstruct %v4float %float_0 %float_1 %300 %float_0
+               OpStore %295 %301
+        %302 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %int_3
+        %303 = OpAccessChain %_ptr_Function_float %viewCeilCoord %uint_1
+        %304 = OpLoad %float %303
+        %305 = OpLoad %float %viewNear
+        %306 = OpFDiv %float %304 %305
+        %307 = OpCompositeConstruct %v4float %float_0 %float_n1 %306 %float_0
+               OpStore %302 %307
+               OpStore %dp %float_0
+               OpStore %i %uint_0
+               OpBranch %310
+        %310 = OpLabel
+               OpLoopMerge %311 %312 None
+               OpBranch %313
+        %313 = OpLabel
+        %315 = OpLoad %uint %i
+        %316 = OpULessThan %bool %315 %uint_6
+        %314 = OpLogicalNot %bool %316
+               OpSelectionMerge %317 None
+               OpBranchConditional %314 %318 %317
+        %318 = OpLabel
+               OpBranch %311
+        %317 = OpLabel
+        %320 = OpLoad %uint %i
+        %321 = OpAccessChain %_ptr_Function_float %frustumPlanes %320 %uint_0
+        %322 = OpLoad %float %321
+        %323 = OpFOrdGreaterThan %bool %322 %float_0
+               OpSelectionMerge %324 None
+               OpBranchConditional %323 %325 %326
+        %325 = OpLabel
+        %327 = OpAccessChain %_ptr_Function_float %p %uint_0
+        %328 = OpAccessChain %_ptr_Function_float %boxMax %uint_0
+        %329 = OpLoad %float %328
+               OpStore %327 %329
+               OpBranch %324
+        %326 = OpLabel
+        %330 = OpAccessChain %_ptr_Function_float %p %uint_0
+        %331 = OpAccessChain %_ptr_Function_float %boxMin %uint_0
+        %332 = OpLoad %float %331
+               OpStore %330 %332
+               OpBranch %324
+        %324 = OpLabel
+        %333 = OpLoad %uint %i
+        %334 = OpAccessChain %_ptr_Function_float %frustumPlanes %333 %uint_1
+        %335 = OpLoad %float %334
+        %336 = OpFOrdGreaterThan %bool %335 %float_0
+               OpSelectionMerge %337 None
+               OpBranchConditional %336 %338 %339
+        %338 = OpLabel
+        %340 = OpAccessChain %_ptr_Function_float %p %uint_1
+        %341 = OpAccessChain %_ptr_Function_float %boxMax %uint_1
+        %342 = OpLoad %float %341
+               OpStore %340 %342
+               OpBranch %337
+        %339 = OpLabel
+        %343 = OpAccessChain %_ptr_Function_float %p %uint_1
+        %344 = OpAccessChain %_ptr_Function_float %boxMin %uint_1
+        %345 = OpLoad %float %344
+               OpStore %343 %345
+               OpBranch %337
+        %337 = OpLabel
+        %346 = OpLoad %uint %i
+        %347 = OpAccessChain %_ptr_Function_float %frustumPlanes %346 %uint_2
+        %348 = OpLoad %float %347
+        %349 = OpFOrdGreaterThan %bool %348 %float_0
+               OpSelectionMerge %350 None
+               OpBranchConditional %349 %351 %352
+        %351 = OpLabel
+        %353 = OpAccessChain %_ptr_Function_float %p %uint_2
+        %354 = OpAccessChain %_ptr_Function_float %boxMax %uint_2
+        %355 = OpLoad %float %354
+               OpStore %353 %355
+               OpBranch %350
+        %352 = OpLabel
+        %356 = OpAccessChain %_ptr_Function_float %p %uint_2
+        %357 = OpAccessChain %_ptr_Function_float %boxMin %uint_2
+        %358 = OpLoad %float %357
+               OpStore %356 %358
+               OpBranch %350
+        %350 = OpLabel
+        %359 = OpAccessChain %_ptr_Function_float %p %uint_3
+               OpStore %359 %float_1
+        %360 = OpLoad %float %dp
+        %363 = OpLoad %v4float %p
+        %364 = OpLoad %uint %i
+        %365 = OpAccessChain %_ptr_Function_v4float %frustumPlanes %364
+        %366 = OpLoad %v4float %365
+        %362 = OpDot %float %363 %366
+        %361 = OpExtInst %float %60 NMin %float_0 %362
+        %367 = OpFAdd %float %360 %361
+               OpStore %dp %367
+               OpBranch %312
+        %312 = OpLabel
+        %368 = OpLoad %uint %i
+        %369 = OpIAdd %uint %368 %uint_1
+               OpStore %i %369
+               OpBranch %310
+        %311 = OpLabel
+        %370 = OpLoad %float %dp
+        %371 = OpFOrdGreaterThanEqual %bool %370 %float_0
+               OpSelectionMerge %372 None
+               OpBranchConditional %371 %373 %372
+        %373 = OpLabel
+        %375 = OpLoad %int %x
+        %376 = OpLoad %int %y
+        %377 = OpIMul %int %376 %int_2
+        %378 = OpIAdd %int %375 %377
+        %374 = OpBitcast %uint %378
+               OpStore %tileId %374
+        %380 = OpLoad %uint %tileId
+        %381 = OpULessThan %bool %380 %uint_0
+               OpSelectionMerge %382 None
+               OpBranchConditional %381 %382 %383
+        %383 = OpLabel
+        %384 = OpLoad %uint %tileId
+        %385 = OpAccessChain %_ptr_Uniform_uint %config %uint_1
+        %386 = OpLoad %uint %385
+        %387 = OpUGreaterThanEqual %bool %384 %386
+               OpBranch %382
+        %382 = OpLabel
+        %388 = OpPhi %bool %381 %373 %387 %383
+               OpSelectionMerge %389 None
+               OpBranchConditional %388 %390 %389
+        %390 = OpLabel
+               OpBranch %183
+        %389 = OpLabel
+        %393 = OpLoad %uint %tileId
+        %395 = OpAccessChain %_ptr_StorageBuffer_uint %tileLightId %uint_0 %393 %uint_0
+        %391 = OpAtomicIAdd %uint %395 %uint_1 %uint_0 %uint_1
+               OpStore %offset %391
+        %397 = OpLoad %uint %offset
+        %398 = OpAccessChain %_ptr_Uniform_uint %config %uint_4
+        %399 = OpLoad %uint %398
+        %400 = OpUGreaterThanEqual %bool %397 %399
+               OpSelectionMerge %401 None
+               OpBranchConditional %400 %402 %401
+        %402 = OpLabel
+               OpBranch %183
+        %401 = OpLabel
+        %403 = OpLoad %uint %tileId
+        %404 = OpLoad %uint %offset
+        %406 = OpAccessChain %_ptr_StorageBuffer_uint_0 %tileLightId %uint_0 %403 %uint_1 %404
+        %407 = OpCompositeExtract %uint %GlobalInvocationID 0
+               OpStore %406 %407
+               OpBranch %372
+        %372 = OpLabel
+               OpBranch %183
+        %183 = OpLabel
+        %408 = OpLoad %int %x
+        %409 = OpIAdd %int %408 %int_1
+               OpStore %x %409
+               OpBranch %181
+        %182 = OpLabel
+               OpBranch %173
+        %173 = OpLabel
+        %410 = OpLoad %int %y
+        %411 = OpIAdd %int %410 %int_1
+               OpStore %y %411
+               OpBranch %171
+        %172 = OpLabel
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %412
+        %414 = OpLabel
+        %416 = OpLoad %v3uint %GlobalInvocationID_1
+        %415 = OpFunctionCall %void %main_inner %416
+               OpReturn
+               OpFunctionEnd
diff --git a/test/bug/tint/1121.wgsl.expected.wgsl b/test/bug/tint/1121.wgsl.expected.wgsl
new file mode 100644
index 0000000000..65efc9a4d9
--- /dev/null
+++ b/test/bug/tint/1121.wgsl.expected.wgsl
@@ -0,0 +1,119 @@
+struct LightData {
+  position : vec4<f32>;
+  color : vec3<f32>;
+  radius : f32;
+};
+
+[[block]]
+struct LightsBuffer {
+  lights : array<LightData>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> lightsBuffer : LightsBuffer;
+
+struct TileLightIdData {
+  count : atomic<u32>;
+  lightId : array<u32, 64>;
+};
+
+[[block]]
+struct Tiles {
+  data : array<TileLightIdData, 4>;
+};
+
+[[group(1), binding(0)]] var<storage, read_write> tileLightId : Tiles;
+
+[[block]]
+struct Config {
+  numLights : u32;
+  numTiles : u32;
+  tileCountX : u32;
+  tileCountY : u32;
+  numTileLightSlot : u32;
+  tileSize : u32;
+};
+
+[[group(2), binding(0)]] var<uniform> config : Config;
+
+[[block]]
+struct Uniforms {
+  min : vec4<f32>;
+  max : vec4<f32>;
+  viewMatrix : mat4x4<f32>;
+  projectionMatrix : mat4x4<f32>;
+  fullScreenSize : vec4<f32>;
+};
+
+[[group(3), binding(0)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(64, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+  var index = GlobalInvocationID.x;
+  if ((index >= config.numLights)) {
+    return;
+  }
+  lightsBuffer.lights[index].position.y = ((lightsBuffer.lights[index].position.y - 0.100000001) + (0.001 * (f32(index) - (64.0 * floor((f32(index) / 64.0))))));
+  if ((lightsBuffer.lights[index].position.y < uniforms.min.y)) {
+    lightsBuffer.lights[index].position.y = uniforms.max.y;
+  }
+  var M : mat4x4<f32> = uniforms.projectionMatrix;
+  var viewNear : f32 = (-(M[3][2]) / (-1.0 + M[2][2]));
+  var viewFar : f32 = (-(M[3][2]) / (1.0 + M[2][2]));
+  var lightPos = lightsBuffer.lights[index].position;
+  lightPos = (uniforms.viewMatrix * lightPos);
+  lightPos = (lightPos / lightPos.w);
+  var lightRadius : f32 = lightsBuffer.lights[index].radius;
+  var boxMin : vec4<f32> = (lightPos - vec4<f32>(vec3<f32>(lightRadius), 0.0));
+  var boxMax : vec4<f32> = (lightPos + vec4<f32>(vec3<f32>(lightRadius), 0.0));
+  var frustumPlanes : array<vec4<f32>, 6>;
+  frustumPlanes[4] = vec4<f32>(0.0, 0.0, -1.0, viewNear);
+  frustumPlanes[5] = vec4<f32>(0.0, 0.0, 1.0, -(viewFar));
+  let TILE_SIZE : i32 = 16;
+  let TILE_COUNT_X : i32 = 2;
+  let TILE_COUNT_Y : i32 = 2;
+  for(var y : i32 = 0; (y < TILE_COUNT_Y); y = (y + 1)) {
+    for(var x : i32 = 0; (x < TILE_COUNT_X); x = (x + 1)) {
+      var tilePixel0Idx : vec2<i32> = vec2<i32>((x * TILE_SIZE), (y * TILE_SIZE));
+      var floorCoord : vec2<f32> = (((2.0 * vec2<f32>(tilePixel0Idx)) / uniforms.fullScreenSize.xy) - vec2<f32>(1.0));
+      var ceilCoord : vec2<f32> = (((2.0 * vec2<f32>((tilePixel0Idx + vec2<i32>(TILE_SIZE)))) / uniforms.fullScreenSize.xy) - vec2<f32>(1.0));
+      var viewFloorCoord : vec2<f32> = vec2<f32>((((-(viewNear) * floorCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
+      var viewCeilCoord : vec2<f32> = vec2<f32>((((-(viewNear) * ceilCoord.x) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord.y) - (M[2][1] * viewNear)) / M[1][1]));
+      frustumPlanes[0] = vec4<f32>(1.0, 0.0, (-(viewFloorCoord.x) / viewNear), 0.0);
+      frustumPlanes[1] = vec4<f32>(-1.0, 0.0, (viewCeilCoord.x / viewNear), 0.0);
+      frustumPlanes[2] = vec4<f32>(0.0, 1.0, (-(viewFloorCoord.y) / viewNear), 0.0);
+      frustumPlanes[3] = vec4<f32>(0.0, -1.0, (viewCeilCoord.y / viewNear), 0.0);
+      var dp : f32 = 0.0;
+      for(var i : u32 = 0u; (i < 6u); i = (i + 1u)) {
+        var p : vec4<f32>;
+        if ((frustumPlanes[i].x > 0.0)) {
+          p.x = boxMax.x;
+        } else {
+          p.x = boxMin.x;
+        }
+        if ((frustumPlanes[i].y > 0.0)) {
+          p.y = boxMax.y;
+        } else {
+          p.y = boxMin.y;
+        }
+        if ((frustumPlanes[i].z > 0.0)) {
+          p.z = boxMax.z;
+        } else {
+          p.z = boxMin.z;
+        }
+        p.w = 1.0;
+        dp = (dp + min(0.0, dot(p, frustumPlanes[i])));
+      }
+      if ((dp >= 0.0)) {
+        var tileId : u32 = u32((x + (y * TILE_COUNT_X)));
+        if (((tileId < 0u) || (tileId >= config.numTiles))) {
+          continue;
+        }
+        var offset : u32 = atomicAdd(&(tileLightId.data[tileId].count), 1u);
+        if ((offset >= config.numTileLightSlot)) {
+          continue;
+        }
+        tileLightId.data[tileId].lightId[offset] = GlobalInvocationID.x;
+      }
+    }
+  }
+}
diff --git a/test/bug/tint/294.wgsl.expected.msl b/test/bug/tint/294.wgsl.expected.msl
index fb5e0c6f64..065aee7fda 100644
--- a/test/bug/tint/294.wgsl.expected.msl
+++ b/test/bug/tint/294.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct Light {
   /* 0x0000 */ packed_float3 position;
   /* 0x000c */ int8_t tint_pad[4];
diff --git a/test/bug/tint/948.wgsl.expected.msl b/test/bug/tint/948.wgsl.expected.msl
index eefb6f401e..a0e0430bd8 100644
--- a/test/bug/tint/948.wgsl.expected.msl
+++ b/test/bug/tint/948.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct LeftOver {
   /* 0x0000 */ float time;
   /* 0x0004 */ uint padding;
diff --git a/test/bug/tint/949.wgsl.expected.msl b/test/bug/tint/949.wgsl.expected.msl
index 63e609783e..776865dfa5 100644
--- a/test/bug/tint/949.wgsl.expected.msl
+++ b/test/bug/tint/949.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct lightingInfo {
   float3 diffuse;
   float3 specular;
diff --git a/test/bug/tint/980.wgsl.expected.msl b/test/bug/tint/980.wgsl.expected.msl
index 27ccfef96a..4c9c0714e2 100644
--- a/test/bug/tint/980.wgsl.expected.msl
+++ b/test/bug/tint/980.wgsl.expected.msl
@@ -1,6 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
 struct S {
   /* 0x0000 */ packed_float3 v;
   /* 0x000c */ uint i;
diff --git a/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl
new file mode 100644
index 0000000000..ccfdaa1439
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+    matrix : mat3x2<f32>;
+    vector : vec3<f32>;
+};
+[[group(0), binding(0)]] var<uniform> data: S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = data.matrix * data.vector;
+}
diff --git a/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..6e12e5cfc5
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.hlsl
@@ -0,0 +1,18 @@
+cbuffer cbuffer_data : register(b0, space0) {
+  uint4 data[3];
+};
+
+float3x2 tint_symbol_2(uint4 buffer[3], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  uint4 ubo_load = buffer[scalar_offset / 4];
+  const uint scalar_offset_1 = ((offset + 8u)) / 4;
+  uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+  const uint scalar_offset_2 = ((offset + 16u)) / 4;
+  uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
+  return float3x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+void main() {
+  const float2 x = mul(asfloat(data[2].xyz), tint_symbol_2(data, 0u));
+  return;
+}
diff --git a/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000000..409b87b668
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl
@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct S {
+  /* 0x0000 */ float3x2 tint_symbol;
+  /* 0x0018 */ int8_t tint_pad[8];
+  /* 0x0020 */ packed_float3 vector;
+  /* 0x002c */ int8_t tint_pad_1[4];
+};
+
+fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
+  float2 const x = (data.tint_symbol * data.vector);
+  return;
+}
+
diff --git a/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..ea2ff2624b
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.spvasm
@@ -0,0 +1,45 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpName %S "S"
+               OpMemberName %S 0 "matrix"
+               OpMemberName %S 1 "vector"
+               OpName %data "data"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 0 ColMajor
+               OpMemberDecorate %S 0 MatrixStride 8
+               OpMemberDecorate %S 1 Offset 32
+               OpDecorate %data NonWritable
+               OpDecorate %data DescriptorSet 0
+               OpDecorate %data Binding 0
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+    %v3float = OpTypeVector %float 3
+          %S = OpTypeStruct %mat3v2float %v3float
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+       %data = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_Uniform_mat3v2float %data %uint_0
+         %16 = OpLoad %mat3v2float %15
+         %19 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
+         %20 = OpLoad %v3float %19
+         %21 = OpMatrixTimesVector %v2float %16 %20
+               OpReturn
+               OpFunctionEnd
diff --git a/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..e2551fa950
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+  matrix : mat3x2<f32>;
+  vector : vec3<f32>;
+};
+
+[[group(0), binding(0)]] var<uniform> data : S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = (data.matrix * data.vector);
+}
diff --git a/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl
new file mode 100644
index 0000000000..2193faf357
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+    matrix : mat3x3<f32>;
+    vector : vec3<f32>;
+};
+[[group(0), binding(0)]] var<uniform> data: S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = data.matrix * data.vector;
+}
diff --git a/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..c680c08b29
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.hlsl
@@ -0,0 +1,15 @@
+cbuffer cbuffer_data : register(b0, space0) {
+  uint4 data[4];
+};
+
+float3x3 tint_symbol_2(uint4 buffer[4], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  return float3x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz), asfloat(buffer[scalar_offset_2 / 4].xyz));
+}
+
+void main() {
+  const float3 x = mul(asfloat(data[3].xyz), tint_symbol_2(data, 0u));
+  return;
+}
diff --git a/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000000..edbe4639b7
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct S {
+  /* 0x0000 */ float3x3 tint_symbol;
+  /* 0x0030 */ packed_float3 vector;
+  /* 0x003c */ int8_t tint_pad[4];
+};
+
+fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
+  float3 const x = (data.tint_symbol * data.vector);
+  return;
+}
+
diff --git a/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..468d2e9c11
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 21
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpName %S "S"
+               OpMemberName %S 0 "matrix"
+               OpMemberName %S 1 "vector"
+               OpName %data "data"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 0 ColMajor
+               OpMemberDecorate %S 0 MatrixStride 16
+               OpMemberDecorate %S 1 Offset 48
+               OpDecorate %data NonWritable
+               OpDecorate %data DescriptorSet 0
+               OpDecorate %data Binding 0
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+          %S = OpTypeStruct %mat3v3float %v3float
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+       %data = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat3v3float = OpTypePointer Uniform %mat3v3float
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat3v3float %data %uint_0
+         %15 = OpLoad %mat3v3float %14
+         %18 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
+         %19 = OpLoad %v3float %18
+         %20 = OpMatrixTimesVector %v3float %15 %19
+               OpReturn
+               OpFunctionEnd
diff --git a/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..e60691cadf
--- /dev/null
+++ b/test/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+  matrix : mat3x3<f32>;
+  vector : vec3<f32>;
+};
+
+[[group(0), binding(0)]] var<uniform> data : S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = (data.matrix * data.vector);
+}
diff --git a/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl
new file mode 100644
index 0000000000..498b3ddb8d
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+    matrix : mat3x3<f32>;
+    vector : vec3<f32>;
+};
+[[group(0), binding(0)]] var<uniform> data: S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = data.vector * data.matrix;
+}
diff --git a/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.hlsl b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..4bd768f2a7
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.hlsl
@@ -0,0 +1,15 @@
+cbuffer cbuffer_data : register(b0, space0) {
+  uint4 data[4];
+};
+
+float3x3 tint_symbol_3(uint4 buffer[4], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  return float3x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz), asfloat(buffer[scalar_offset_2 / 4].xyz));
+}
+
+void main() {
+  const float3 x = mul(tint_symbol_3(data, 0u), asfloat(data[3].xyz));
+  return;
+}
diff --git a/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl
new file mode 100644
index 0000000000..d4a9b71d5c
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct S {
+  /* 0x0000 */ float3x3 tint_symbol;
+  /* 0x0030 */ packed_float3 vector;
+  /* 0x003c */ int8_t tint_pad[4];
+};
+
+fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
+  float3 const x = (data.vector * data.tint_symbol);
+  return;
+}
+
diff --git a/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.spvasm b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..1820ed869a
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 21
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpName %S "S"
+               OpMemberName %S 0 "matrix"
+               OpMemberName %S 1 "vector"
+               OpName %data "data"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 0 ColMajor
+               OpMemberDecorate %S 0 MatrixStride 16
+               OpMemberDecorate %S 1 Offset 48
+               OpDecorate %data NonWritable
+               OpDecorate %data DescriptorSet 0
+               OpDecorate %data Binding 0
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+          %S = OpTypeStruct %mat3v3float %v3float
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+       %data = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat3v3float = OpTypePointer Uniform %mat3v3float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
+         %15 = OpLoad %v3float %14
+         %18 = OpAccessChain %_ptr_Uniform_mat3v3float %data %uint_0
+         %19 = OpLoad %mat3v3float %18
+         %20 = OpVectorTimesMatrix %v3float %15 %19
+               OpReturn
+               OpFunctionEnd
diff --git a/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.wgsl b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..3a8b95764c
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+  matrix : mat3x3<f32>;
+  vector : vec3<f32>;
+};
+
+[[group(0), binding(0)]] var<uniform> data : S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = (data.vector * data.matrix);
+}
diff --git a/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl
new file mode 100644
index 0000000000..97844ada53
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl
@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+    matrix : mat4x3<f32>;
+    vector : vec3<f32>;
+};
+[[group(0), binding(0)]] var<uniform> data: S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = data.vector * data.matrix;
+}
diff --git a/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.hlsl b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..574af42142
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+cbuffer cbuffer_data : register(b0, space0) {
+  uint4 data[5];
+};
+
+float4x3 tint_symbol_3(uint4 buffer[5], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  const uint scalar_offset_3 = ((offset + 48u)) / 4;
+  return float4x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz), asfloat(buffer[scalar_offset_2 / 4].xyz), asfloat(buffer[scalar_offset_3 / 4].xyz));
+}
+
+void main() {
+  const float4 x = mul(tint_symbol_3(data, 0u), asfloat(data[4].xyz));
+  return;
+}
diff --git a/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl
new file mode 100644
index 0000000000..bb2e75f7e3
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline auto operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline auto operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct S {
+  /* 0x0000 */ float4x3 tint_symbol;
+  /* 0x0040 */ packed_float3 vector;
+  /* 0x004c */ int8_t tint_pad[4];
+};
+
+fragment void tint_symbol_1(constant S& data [[buffer(0)]]) {
+  float4 const x = (data.vector * data.tint_symbol);
+  return;
+}
+
diff --git a/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.spvasm b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..2efe5701c8
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.spvasm
@@ -0,0 +1,45 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpName %S "S"
+               OpMemberName %S 0 "matrix"
+               OpMemberName %S 1 "vector"
+               OpName %data "data"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 0 ColMajor
+               OpMemberDecorate %S 0 MatrixStride 16
+               OpMemberDecorate %S 1 Offset 64
+               OpDecorate %data NonWritable
+               OpDecorate %data DescriptorSet 0
+               OpDecorate %data Binding 0
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat4v3float = OpTypeMatrix %v3float 4
+          %S = OpTypeStruct %mat4v3float %v3float
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+       %data = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat4v3float = OpTypePointer Uniform %mat4v3float
+    %v4float = OpTypeVector %float 4
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_v3float %data %uint_1
+         %15 = OpLoad %v3float %14
+         %18 = OpAccessChain %_ptr_Uniform_mat4v3float %data %uint_0
+         %19 = OpLoad %mat4v3float %18
+         %20 = OpVectorTimesMatrix %v4float %15 %19
+               OpReturn
+               OpFunctionEnd
diff --git a/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.wgsl b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..2f567cc670
--- /dev/null
+++ b/test/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+  matrix : mat4x3<f32>;
+  vector : vec3<f32>;
+};
+
+[[group(0), binding(0)]] var<uniform> data : S;
+
+[[stage(fragment)]]
+fn main() {
+  let x = (data.vector * data.matrix);
+}