IR: Fix robustness transform on textureLoad of sampled and depth textures

For sampled and depth textures, which contain a 'level' argument, the
robustness transform is supposed to clamp 'coords' using the dimensions
at the clamped level, but it was looking up dimensions at level 0
instead.

Bug: 42250751
Bug: 42251045
Bug: 378541479
Change-Id: I0e7fd6148417b248a9b584ae19818e9027306b63
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/214514
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>

src/tint/lang/core/ir/transform/robustness.cc

diff --git a/src/tint/lang/core/ir/transform/robustness.cc b/src/tint/lang/core/ir/transform/robustness.cc
index 079a6cd..b9dd071 100644
--- a/src/tint/lang/core/ir/transform/robustness.cc
+++ b/src/tint/lang/core/ir/transform/robustness.cc
@@ -309,7 +309,6 @@
                 break;
             }
             case core::BuiltinFn::kTextureLoad: {
-                clamp_coords(1u);
                 uint32_t next_arg = 2u;
                 if (type::IsTextureArray(texture->Dim())) {
                     clamp_array_index(next_arg++);
@@ -317,6 +316,7 @@
                 if (texture->IsAnyOf<type::SampledTexture, type::DepthTexture>()) {
                     clamp_level(next_arg++);
                 }
+                clamp_coords(1u);  // Must run after clamp_level
                 break;
             }
             case core::BuiltinFn::kTextureStore: {

src/tint/lang/core/ir/transform/robustness_test.cc

diff --git a/src/tint/lang/core/ir/transform/robustness_test.cc b/src/tint/lang/core/ir/transform/robustness_test.cc
index a9ecbf8..791101e 100644
--- a/src/tint/lang/core/ir/transform/robustness_test.cc
+++ b/src/tint/lang/core/ir/transform/robustness_test.cc
@@ -2146,28 +2146,28 @@
 %load_signed = func(%coords:i32, %level:i32):vec4<f32> {
   $B2: {
     %5:texture_1d<f32> = load %texture
-    %6:u32 = textureDimensions %5
+    %6:u32 = textureNumLevels %5
     %7:u32 = sub %6, 1u
-    %8:u32 = convert %coords
+    %8:u32 = convert %level
     %9:u32 = min %8, %7
-    %10:u32 = textureNumLevels %5
+    %10:u32 = textureDimensions %5, %9
     %11:u32 = sub %10, 1u
-    %12:u32 = convert %level
+    %12:u32 = convert %coords
     %13:u32 = min %12, %11
-    %14:vec4<f32> = textureLoad %5, %9, %13
+    %14:vec4<f32> = textureLoad %5, %13, %9
     ret %14
   }
 }
 %load_unsigned = func(%coords_1:u32, %level_1:u32):vec4<f32> {  # %coords_1: 'coords', %level_1: 'level'
   $B3: {
     %18:texture_1d<f32> = load %texture
-    %19:u32 = textureDimensions %18
+    %19:u32 = textureNumLevels %18
     %20:u32 = sub %19, 1u
-    %21:u32 = min %coords_1, %20
-    %22:u32 = textureNumLevels %18
+    %21:u32 = min %level_1, %20
+    %22:u32 = textureDimensions %18, %21
     %23:u32 = sub %22, 1u
-    %24:u32 = min %level_1, %23
-    %25:vec4<f32> = textureLoad %18, %21, %24
+    %24:u32 = min %coords_1, %23
+    %25:vec4<f32> = textureLoad %18, %24, %21
     ret %25
   }
 }
@@ -2243,28 +2243,28 @@
 %load_signed = func(%coords:vec2<i32>, %level:i32):vec4<f32> {
   $B2: {
     %5:texture_2d<f32> = load %texture
-    %6:vec2<u32> = textureDimensions %5
-    %7:vec2<u32> = sub %6, vec2<u32>(1u)
-    %8:vec2<u32> = convert %coords
-    %9:vec2<u32> = min %8, %7
-    %10:u32 = textureNumLevels %5
-    %11:u32 = sub %10, 1u
-    %12:u32 = convert %level
-    %13:u32 = min %12, %11
-    %14:vec4<f32> = textureLoad %5, %9, %13
+    %6:u32 = textureNumLevels %5
+    %7:u32 = sub %6, 1u
+    %8:u32 = convert %level
+    %9:u32 = min %8, %7
+    %10:vec2<u32> = textureDimensions %5, %9
+    %11:vec2<u32> = sub %10, vec2<u32>(1u)
+    %12:vec2<u32> = convert %coords
+    %13:vec2<u32> = min %12, %11
+    %14:vec4<f32> = textureLoad %5, %13, %9
     ret %14
   }
 }
 %load_unsigned = func(%coords_1:vec2<u32>, %level_1:u32):vec4<f32> {  # %coords_1: 'coords', %level_1: 'level'
   $B3: {
     %18:texture_2d<f32> = load %texture
-    %19:vec2<u32> = textureDimensions %18
-    %20:vec2<u32> = sub %19, vec2<u32>(1u)
-    %21:vec2<u32> = min %coords_1, %20
-    %22:u32 = textureNumLevels %18
-    %23:u32 = sub %22, 1u
-    %24:u32 = min %level_1, %23
-    %25:vec4<f32> = textureLoad %18, %21, %24
+    %19:u32 = textureNumLevels %18
+    %20:u32 = sub %19, 1u
+    %21:u32 = min %level_1, %20
+    %22:vec2<u32> = textureDimensions %18, %21
+    %23:vec2<u32> = sub %22, vec2<u32>(1u)
+    %24:vec2<u32> = min %coords_1, %23
+    %25:vec4<f32> = textureLoad %18, %24, %21
     ret %25
   }
 }
@@ -2343,35 +2343,35 @@
 %load_signed = func(%coords:vec2<i32>, %layer:i32, %level:i32):vec4<f32> {
   $B2: {
     %6:texture_2d_array<f32> = load %texture
-    %7:vec2<u32> = textureDimensions %6
-    %8:vec2<u32> = sub %7, vec2<u32>(1u)
-    %9:vec2<u32> = convert %coords
-    %10:vec2<u32> = min %9, %8
-    %11:u32 = textureNumLayers %6
+    %7:u32 = textureNumLayers %6
+    %8:u32 = sub %7, 1u
+    %9:u32 = convert %layer
+    %10:u32 = min %9, %8
+    %11:u32 = textureNumLevels %6
     %12:u32 = sub %11, 1u
-    %13:u32 = convert %layer
+    %13:u32 = convert %level
     %14:u32 = min %13, %12
-    %15:u32 = textureNumLevels %6
-    %16:u32 = sub %15, 1u
-    %17:u32 = convert %level
-    %18:u32 = min %17, %16
-    %19:vec4<f32> = textureLoad %6, %10, %14, %18
+    %15:vec2<u32> = textureDimensions %6, %14
+    %16:vec2<u32> = sub %15, vec2<u32>(1u)
+    %17:vec2<u32> = convert %coords
+    %18:vec2<u32> = min %17, %16
+    %19:vec4<f32> = textureLoad %6, %18, %10, %14
     ret %19
   }
 }
 %load_unsigned = func(%coords_1:vec2<u32>, %layer_1:u32, %level_1:u32):vec4<f32> {  # %coords_1: 'coords', %layer_1: 'layer', %level_1: 'level'
   $B3: {
     %24:texture_2d_array<f32> = load %texture
-    %25:vec2<u32> = textureDimensions %24
-    %26:vec2<u32> = sub %25, vec2<u32>(1u)
-    %27:vec2<u32> = min %coords_1, %26
-    %28:u32 = textureNumLayers %24
+    %25:u32 = textureNumLayers %24
+    %26:u32 = sub %25, 1u
+    %27:u32 = min %layer_1, %26
+    %28:u32 = textureNumLevels %24
     %29:u32 = sub %28, 1u
-    %30:u32 = min %layer_1, %29
-    %31:u32 = textureNumLevels %24
-    %32:u32 = sub %31, 1u
-    %33:u32 = min %level_1, %32
-    %34:vec4<f32> = textureLoad %24, %27, %30, %33
+    %30:u32 = min %level_1, %29
+    %31:vec2<u32> = textureDimensions %24, %30
+    %32:vec2<u32> = sub %31, vec2<u32>(1u)
+    %33:vec2<u32> = min %coords_1, %32
+    %34:vec4<f32> = textureLoad %24, %33, %27, %30
     ret %34
   }
 }
@@ -2447,28 +2447,28 @@
 %load_signed = func(%coords:vec3<i32>, %level:i32):vec4<f32> {
   $B2: {
     %5:texture_3d<f32> = load %texture
-    %6:vec3<u32> = textureDimensions %5
-    %7:vec3<u32> = sub %6, vec3<u32>(1u)
-    %8:vec3<u32> = convert %coords
-    %9:vec3<u32> = min %8, %7
-    %10:u32 = textureNumLevels %5
-    %11:u32 = sub %10, 1u
-    %12:u32 = convert %level
-    %13:u32 = min %12, %11
-    %14:vec4<f32> = textureLoad %5, %9, %13
+    %6:u32 = textureNumLevels %5
+    %7:u32 = sub %6, 1u
+    %8:u32 = convert %level
+    %9:u32 = min %8, %7
+    %10:vec3<u32> = textureDimensions %5, %9
+    %11:vec3<u32> = sub %10, vec3<u32>(1u)
+    %12:vec3<u32> = convert %coords
+    %13:vec3<u32> = min %12, %11
+    %14:vec4<f32> = textureLoad %5, %13, %9
     ret %14
   }
 }
 %load_unsigned = func(%coords_1:vec3<u32>, %level_1:u32):vec4<f32> {  # %coords_1: 'coords', %level_1: 'level'
   $B3: {
     %18:texture_3d<f32> = load %texture
-    %19:vec3<u32> = textureDimensions %18
-    %20:vec3<u32> = sub %19, vec3<u32>(1u)
-    %21:vec3<u32> = min %coords_1, %20
-    %22:u32 = textureNumLevels %18
-    %23:u32 = sub %22, 1u
-    %24:u32 = min %level_1, %23
-    %25:vec4<f32> = textureLoad %18, %21, %24
+    %19:u32 = textureNumLevels %18
+    %20:u32 = sub %19, 1u
+    %21:u32 = min %level_1, %20
+    %22:vec3<u32> = textureDimensions %18, %21
+    %23:vec3<u32> = sub %22, vec3<u32>(1u)
+    %24:vec3<u32> = min %coords_1, %23
+    %25:vec4<f32> = textureLoad %18, %24, %21
     ret %25
   }
 }
@@ -2632,28 +2632,28 @@
 %load_signed = func(%coords:vec2<i32>, %level:i32):f32 {
   $B2: {
     %5:texture_depth_2d = load %texture
-    %6:vec2<u32> = textureDimensions %5
-    %7:vec2<u32> = sub %6, vec2<u32>(1u)
-    %8:vec2<u32> = convert %coords
-    %9:vec2<u32> = min %8, %7
-    %10:u32 = textureNumLevels %5
-    %11:u32 = sub %10, 1u
-    %12:u32 = convert %level
-    %13:u32 = min %12, %11
-    %14:f32 = textureLoad %5, %9, %13
+    %6:u32 = textureNumLevels %5
+    %7:u32 = sub %6, 1u
+    %8:u32 = convert %level
+    %9:u32 = min %8, %7
+    %10:vec2<u32> = textureDimensions %5, %9
+    %11:vec2<u32> = sub %10, vec2<u32>(1u)
+    %12:vec2<u32> = convert %coords
+    %13:vec2<u32> = min %12, %11
+    %14:f32 = textureLoad %5, %13, %9
     ret %14
   }
 }
 %load_unsigned = func(%coords_1:vec2<u32>, %level_1:u32):f32 {  # %coords_1: 'coords', %level_1: 'level'
   $B3: {
     %18:texture_depth_2d = load %texture
-    %19:vec2<u32> = textureDimensions %18
-    %20:vec2<u32> = sub %19, vec2<u32>(1u)
-    %21:vec2<u32> = min %coords_1, %20
-    %22:u32 = textureNumLevels %18
-    %23:u32 = sub %22, 1u
-    %24:u32 = min %level_1, %23
-    %25:f32 = textureLoad %18, %21, %24
+    %19:u32 = textureNumLevels %18
+    %20:u32 = sub %19, 1u
+    %21:u32 = min %level_1, %20
+    %22:vec2<u32> = textureDimensions %18, %21
+    %23:vec2<u32> = sub %22, vec2<u32>(1u)
+    %24:vec2<u32> = min %coords_1, %23
+    %25:f32 = textureLoad %18, %24, %21
     ret %25
   }
 }
@@ -2731,35 +2731,35 @@
 %load_signed = func(%coords:vec2<i32>, %layer:i32, %level:i32):f32 {
   $B2: {
     %6:texture_depth_2d_array = load %texture
-    %7:vec2<u32> = textureDimensions %6
-    %8:vec2<u32> = sub %7, vec2<u32>(1u)
-    %9:vec2<u32> = convert %coords
-    %10:vec2<u32> = min %9, %8
-    %11:u32 = textureNumLayers %6
+    %7:u32 = textureNumLayers %6
+    %8:u32 = sub %7, 1u
+    %9:u32 = convert %layer
+    %10:u32 = min %9, %8
+    %11:u32 = textureNumLevels %6
     %12:u32 = sub %11, 1u
-    %13:u32 = convert %layer
+    %13:u32 = convert %level
     %14:u32 = min %13, %12
-    %15:u32 = textureNumLevels %6
-    %16:u32 = sub %15, 1u
-    %17:u32 = convert %level
-    %18:u32 = min %17, %16
-    %19:f32 = textureLoad %6, %10, %14, %18
+    %15:vec2<u32> = textureDimensions %6, %14
+    %16:vec2<u32> = sub %15, vec2<u32>(1u)
+    %17:vec2<u32> = convert %coords
+    %18:vec2<u32> = min %17, %16
+    %19:f32 = textureLoad %6, %18, %10, %14
     ret %19
   }
 }
 %load_unsigned = func(%coords_1:vec2<u32>, %layer_1:u32, %level_1:u32):f32 {  # %coords_1: 'coords', %layer_1: 'layer', %level_1: 'level'
   $B3: {
     %24:texture_depth_2d_array = load %texture
-    %25:vec2<u32> = textureDimensions %24
-    %26:vec2<u32> = sub %25, vec2<u32>(1u)
-    %27:vec2<u32> = min %coords_1, %26
-    %28:u32 = textureNumLayers %24
+    %25:u32 = textureNumLayers %24
+    %26:u32 = sub %25, 1u
+    %27:u32 = min %layer_1, %26
+    %28:u32 = textureNumLevels %24
     %29:u32 = sub %28, 1u
-    %30:u32 = min %layer_1, %29
-    %31:u32 = textureNumLevels %24
-    %32:u32 = sub %31, 1u
-    %33:u32 = min %level_1, %32
-    %34:f32 = textureLoad %24, %27, %30, %33
+    %30:u32 = min %level_1, %29
+    %31:vec2<u32> = textureDimensions %24, %30
+    %32:vec2<u32> = sub %31, vec2<u32>(1u)
+    %33:vec2<u32> = min %coords_1, %32
+    %34:f32 = textureLoad %24, %33, %27, %30
     ret %34
   }
 }
@@ -3191,28 +3191,28 @@
 %load_signed = func(%coords:vec2<i32>, %layer:i32):vec4<f32> {
   $B2: {
     %5:texture_storage_2d_array<rgba8unorm, read_write> = load %texture
-    %6:vec2<u32> = textureDimensions %5
-    %7:vec2<u32> = sub %6, vec2<u32>(1u)
-    %8:vec2<u32> = convert %coords
-    %9:vec2<u32> = min %8, %7
-    %10:u32 = textureNumLayers %5
-    %11:u32 = sub %10, 1u
-    %12:u32 = convert %layer
-    %13:u32 = min %12, %11
-    %14:vec4<f32> = textureLoad %5, %9, %13
+    %6:u32 = textureNumLayers %5
+    %7:u32 = sub %6, 1u
+    %8:u32 = convert %layer
+    %9:u32 = min %8, %7
+    %10:vec2<u32> = textureDimensions %5
+    %11:vec2<u32> = sub %10, vec2<u32>(1u)
+    %12:vec2<u32> = convert %coords
+    %13:vec2<u32> = min %12, %11
+    %14:vec4<f32> = textureLoad %5, %13, %9
     ret %14
   }
 }
 %load_unsigned = func(%coords_1:vec2<u32>, %layer_1:u32):vec4<f32> {  # %coords_1: 'coords', %layer_1: 'layer'
   $B3: {
     %18:texture_storage_2d_array<rgba8unorm, read_write> = load %texture
-    %19:vec2<u32> = textureDimensions %18
-    %20:vec2<u32> = sub %19, vec2<u32>(1u)
-    %21:vec2<u32> = min %coords_1, %20
-    %22:u32 = textureNumLayers %18
-    %23:u32 = sub %22, 1u
-    %24:u32 = min %layer_1, %23
-    %25:vec4<f32> = textureLoad %18, %21, %24
+    %19:u32 = textureNumLayers %18
+    %20:u32 = sub %19, 1u
+    %21:u32 = min %layer_1, %20
+    %22:vec2<u32> = textureDimensions %18
+    %23:vec2<u32> = sub %22, vec2<u32>(1u)
+    %24:vec2<u32> = min %coords_1, %23
+    %25:vec4<f32> = textureLoad %18, %24, %21
     ret %25
   }
 }

src/tint/lang/spirv/writer/texture_builtin_test.cc

diff --git a/src/tint/lang/spirv/writer/texture_builtin_test.cc b/src/tint/lang/spirv/writer/texture_builtin_test.cc
index 432dc0d..0a4da78 100644
--- a/src/tint/lang/spirv/writer/texture_builtin_test.cc
+++ b/src/tint/lang/spirv/writer/texture_builtin_test.cc
@@ -2006,14 +2006,14 @@
 
     ASSERT_TRUE(Generate()) << Error() << output_;
     EXPECT_INST(R"(
-         %13 = OpImageQuerySizeLod %v2uint %texture %uint_0
-         %15 = OpISub %v2uint %13 %16
-         %18 = OpExtInst %v2uint %19 UMin %coords %15
-         %20 = OpImageQueryLevels %uint %texture
-         %21 = OpISub %uint %20 %uint_1
-         %22 = OpBitcast %uint %level
-         %23 = OpExtInst %uint %19 UMin %22 %21
-     %result = OpImageFetch %v4float %texture %18 Lod %23
+         %13 = OpImageQueryLevels %uint %texture
+         %14 = OpISub %uint %13 %uint_1
+         %16 = OpBitcast %uint %level
+         %17 = OpExtInst %uint %18 UMin %16 %14
+         %19 = OpImageQuerySizeLod %v2uint %texture %17
+         %20 = OpISub %v2uint %19 %21
+         %22 = OpExtInst %v2uint %18 UMin %coords %20
+     %result = OpImageFetch %v4float %texture %22 Lod %17
 )");
 }
 

test/tint/bug/chromium/378541479.wgsl

diff --git a/test/tint/bug/chromium/378541479.wgsl b/test/tint/bug/chromium/378541479.wgsl
new file mode 100644
index 0000000..8badf52
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl
@@ -0,0 +1,10 @@
+// flags: --transform robustness
+
+@group(0) @binding(0) var<uniform> level : u32;
+@group(0) @binding(1) var<uniform> coords : vec2<u32>;
+@group(0) @binding(2) var tex: texture_depth_2d;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  var res: f32 = textureLoad(tex, coords, level);
+}

test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6ceba9a
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+cbuffer cbuffer_level : register(b0) {
+  uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+  uint4 coords[1];
+};
+Texture2D tex : register(t2);
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  float res = tex.Load(uint3(coords[0].xy, level[0].x)).x;
+  return;
+}

test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6ceba9a
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl
@@ -0,0 +1,13 @@
+cbuffer cbuffer_level : register(b0) {
+  uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+  uint4 coords[1];
+};
+Texture2D tex : register(t2);
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  float res = tex.Load(uint3(coords[0].xy, level[0].x)).x;
+  return;
+}

test/tint/bug/chromium/378541479.wgsl.expected.glsl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.glsl b/test/tint/bug/chromium/378541479.wgsl.expected.glsl
new file mode 100644
index 0000000..bb0ee27
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+
+struct TintTextureUniformData {
+  uint tint_builtin_value_0;
+};
+
+layout(binding = 0, std140)
+uniform level_block_1_ubo {
+  uint inner;
+} v;
+layout(binding = 1, std140)
+uniform coords_block_1_ubo {
+  uvec2 inner;
+} v_1;
+layout(binding = 0, std140)
+uniform tint_symbol_1_ubo {
+  TintTextureUniformData inner;
+} v_2;
+uniform highp sampler2D tex;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  uvec2 v_3 = v_1.inner;
+  uint v_4 = min(v.inner, (v_2.inner.tint_builtin_value_0 - 1u));
+  ivec2 v_5 = ivec2(min(v_3, (uvec2(textureSize(tex, int(v_4))) - uvec2(1u))));
+  float res = texelFetch(tex, v_5, int(v_4)).x;
+}

test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..72b8b84
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,21 @@
+
+cbuffer cbuffer_level : register(b0) {
+  uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+  uint4 coords[1];
+};
+Texture2D tex : register(t2);
+[numthreads(1, 1, 1)]
+void compute_main() {
+  uint2 v = coords[0u].xy;
+  uint v_1 = level[0u].x;
+  uint3 v_2 = (0u).xxx;
+  tex.GetDimensions(0u, v_2[0u], v_2[1u], v_2[2u]);
+  uint v_3 = min(v_1, (v_2.z - 1u));
+  uint3 v_4 = (0u).xxx;
+  tex.GetDimensions(uint(v_3), v_4[0u], v_4[1u], v_4[2u]);
+  int2 v_5 = int2(min(v, (v_4.xy - (1u).xx)));
+  float res = tex.Load(int3(v_5, int(v_3))).x;
+}
+

test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..72b8b84
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,21 @@
+
+cbuffer cbuffer_level : register(b0) {
+  uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+  uint4 coords[1];
+};
+Texture2D tex : register(t2);
+[numthreads(1, 1, 1)]
+void compute_main() {
+  uint2 v = coords[0u].xy;
+  uint v_1 = level[0u].x;
+  uint3 v_2 = (0u).xxx;
+  tex.GetDimensions(0u, v_2[0u], v_2[1u], v_2[2u]);
+  uint v_3 = min(v_1, (v_2.z - 1u));
+  uint3 v_4 = (0u).xxx;
+  tex.GetDimensions(uint(v_3), v_4[0u], v_4[1u], v_4[2u]);
+  int2 v_5 = int2(min(v, (v_4.xy - (1u).xx)));
+  float res = tex.Load(int3(v_5, int(v_3))).x;
+}
+

test/tint/bug/chromium/378541479.wgsl.expected.ir.msl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.ir.msl b/test/tint/bug/chromium/378541479.wgsl.expected.ir.msl
new file mode 100644
index 0000000..b1edc57
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.ir.msl
@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+  const constant uint* level;
+  const constant uint2* coords;
+  depth2d<float, access::sample> tex;
+};
+
+kernel void compute_main(const constant uint* level [[buffer(1)]], const constant uint2* coords [[buffer(0)]], depth2d<float, access::sample> tex [[texture(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.level=level, .coords=coords, .tex=tex};
+  uint2 const v = (*tint_module_vars.coords);
+  uint const v_1 = (*tint_module_vars.level);
+  uint const v_2 = min(v_1, (tint_module_vars.tex.get_num_mip_levels() - 1u));
+  uint const v_3 = tint_module_vars.tex.get_width(v_2);
+  float res = tint_module_vars.tex.read(min(v, (uint2(v_3, tint_module_vars.tex.get_height(v_2)) - uint2(1u))), v_2);
+}

test/tint/bug/chromium/378541479.wgsl.expected.msl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.msl b/test/tint/bug/chromium/378541479.wgsl.expected.msl
new file mode 100644
index 0000000..0ac2d2a
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void compute_main(const constant uint* tint_symbol [[buffer(1)]], depth2d<float, access::sample> tint_symbol_1 [[texture(0)]], const constant uint2* tint_symbol_2 [[buffer(0)]]) {
+  uint const level_idx = min(uint(*(tint_symbol)), (tint_symbol_1.get_num_mip_levels() - 1u));
+  float res = tint_symbol_1.read(uint2(min(*(tint_symbol_2), (uint2(tint_symbol_1.get_width(level_idx), tint_symbol_1.get_height(level_idx)) - uint2(1u)))), level_idx);
+  return;
+}
+

test/tint/bug/chromium/378541479.wgsl.expected.spvasm

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.spvasm b/test/tint/bug/chromium/378541479.wgsl.expected.spvasm
new file mode 100644
index 0000000..b8a3b11
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.spvasm
@@ -0,0 +1,70 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 39
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %29 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpMemberName %level_block 0 "inner"
+               OpName %level_block "level_block"
+               OpMemberName %coords_block 0 "inner"
+               OpName %coords_block "coords_block"
+               OpName %tex "tex"
+               OpName %compute_main "compute_main"
+               OpName %res "res"
+               OpMemberDecorate %level_block 0 Offset 0
+               OpDecorate %level_block Block
+               OpDecorate %1 DescriptorSet 0
+               OpDecorate %1 Binding 0
+               OpDecorate %1 NonWritable
+               OpMemberDecorate %coords_block 0 Offset 0
+               OpDecorate %coords_block Block
+               OpDecorate %5 DescriptorSet 0
+               OpDecorate %5 Binding 1
+               OpDecorate %5 NonWritable
+               OpDecorate %tex DescriptorSet 0
+               OpDecorate %tex Binding 2
+       %uint = OpTypeInt 32 0
+%level_block = OpTypeStruct %uint
+%_ptr_Uniform_level_block = OpTypePointer Uniform %level_block
+          %1 = OpVariable %_ptr_Uniform_level_block Uniform
+     %v2uint = OpTypeVector %uint 2
+%coords_block = OpTypeStruct %v2uint
+%_ptr_Uniform_coords_block = OpTypePointer Uniform %coords_block
+          %5 = OpVariable %_ptr_Uniform_coords_block Uniform
+      %float = OpTypeFloat 32
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+        %tex = OpVariable %_ptr_UniformConstant_11 UniformConstant
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+%_ptr_Uniform_v2uint = OpTypePointer Uniform %v2uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+         %32 = OpConstantComposite %v2uint %uint_1 %uint_1
+    %v4float = OpTypeVector %float 4
+%_ptr_Function_float = OpTypePointer Function %float
+%compute_main = OpFunction %void None %15
+         %16 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function
+         %17 = OpLoad %11 %tex None
+         %18 = OpAccessChain %_ptr_Uniform_v2uint %5 %uint_0
+         %21 = OpLoad %v2uint %18 None
+         %22 = OpAccessChain %_ptr_Uniform_uint %1 %uint_0
+         %24 = OpLoad %uint %22 None
+         %25 = OpImageQueryLevels %uint %17
+         %26 = OpISub %uint %25 %uint_1
+         %28 = OpExtInst %uint %29 UMin %24 %26
+         %30 = OpImageQuerySizeLod %v2uint %17 %28
+         %31 = OpISub %v2uint %30 %32
+         %33 = OpExtInst %v2uint %29 UMin %21 %31
+         %34 = OpImageFetch %v4float %17 %33 Lod %28
+         %36 = OpCompositeExtract %float %34 0
+               OpStore %res %36
+               OpReturn
+               OpFunctionEnd

test/tint/bug/chromium/378541479.wgsl.expected.wgsl

diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.wgsl b/test/tint/bug/chromium/378541479.wgsl.expected.wgsl
new file mode 100644
index 0000000..705ce78
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<uniform> level : u32;
+
+@group(0) @binding(1) var<uniform> coords : vec2<u32>;
+
+@group(0) @binding(2) var tex : texture_depth_2d;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  var res : f32 = textureLoad(tex, coords, level);
+}

test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl

diff --git a/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl
index e6e322dc..32a4cf0 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl
@@ -51,112 +51,109 @@
   float3 v_7 = (0.0f).xxx;
   float v_8 = 0.0f;
   if ((params.numPlanes == 1u)) {
-    uint2 v_9 = (0u).xx;
-    plane_0.GetDimensions(v_9[0u], v_9[1u]);
-    uint2 v_10 = min(v_6, (v_9 - (1u).xx));
+    uint3 v_9 = (0u).xxx;
+    plane_0.GetDimensions(0u, v_9[0u], v_9[1u], v_9[2u]);
+    uint v_10 = min(0u, (v_9.z - 1u));
     uint3 v_11 = (0u).xxx;
-    plane_0.GetDimensions(0u, v_11[0u], v_11[1u], v_11[2u]);
-    uint v_12 = min(0u, (v_11.z - 1u));
-    int2 v_13 = int2(v_10);
-    float4 v_14 = float4(plane_0.Load(int3(v_13, int(v_12))));
-    v_7 = v_14.xyz;
-    v_8 = v_14[3u];
+    plane_0.GetDimensions(uint(v_10), v_11[0u], v_11[1u], v_11[2u]);
+    int2 v_12 = int2(min(v_6, (v_11.xy - (1u).xx)));
+    float4 v_13 = float4(plane_0.Load(int3(v_12, int(v_10))));
+    v_7 = v_13.xyz;
+    v_8 = v_13[3u];
   } else {
-    uint2 v_15 = (0u).xx;
-    plane_0.GetDimensions(v_15[0u], v_15[1u]);
-    uint2 v_16 = min(v_6, (v_15 - (1u).xx));
-    uint3 v_17 = (0u).xxx;
-    plane_0.GetDimensions(0u, v_17[0u], v_17[1u], v_17[2u]);
-    uint v_18 = min(0u, (v_17.z - 1u));
-    int2 v_19 = int2(v_16);
-    float v_20 = float4(plane_0.Load(int3(v_19, int(v_18))))[0u];
-    uint2 v_21 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
-    uint2 v_22 = (0u).xx;
-    plane_1.GetDimensions(v_22[0u], v_22[1u]);
-    uint2 v_23 = min(v_21, (v_22 - (1u).xx));
-    uint3 v_24 = (0u).xxx;
-    plane_1.GetDimensions(0u, v_24[0u], v_24[1u], v_24[2u]);
-    uint v_25 = min(0u, (v_24.z - 1u));
-    int2 v_26 = int2(v_23);
-    v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_20, float4(plane_1.Load(int3(v_26, int(v_25)))).xy, 1.0f));
+    uint3 v_14 = (0u).xxx;
+    plane_0.GetDimensions(0u, v_14[0u], v_14[1u], v_14[2u]);
+    uint v_15 = min(0u, (v_14.z - 1u));
+    uint3 v_16 = (0u).xxx;
+    plane_0.GetDimensions(uint(v_15), v_16[0u], v_16[1u], v_16[2u]);
+    int2 v_17 = int2(min(v_6, (v_16.xy - (1u).xx)));
+    float v_18 = float4(plane_0.Load(int3(v_17, int(v_15))))[0u];
+    uint2 v_19 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
+    uint3 v_20 = (0u).xxx;
+    plane_1.GetDimensions(0u, v_20[0u], v_20[1u], v_20[2u]);
+    uint v_21 = min(0u, (v_20.z - 1u));
+    uint3 v_22 = (0u).xxx;
+    plane_1.GetDimensions(uint(v_21), v_22[0u], v_22[1u], v_22[2u]);
+    int2 v_23 = int2(min(v_19, (v_22.xy - (1u).xx)));
+    v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_18, float4(plane_1.Load(int3(v_23, int(v_21)))).xy, 1.0f));
     v_8 = 1.0f;
   }
-  float3 v_27 = v_7;
-  float3 v_28 = (0.0f).xxx;
+  float3 v_24 = v_7;
+  float3 v_25 = (0.0f).xxx;
   if ((params.doYuvToRgbConversionOnly == 0u)) {
-    tint_GammaTransferParams v_29 = params.gammaDecodeParams;
-    tint_GammaTransferParams v_30 = params.gammaEncodeParams;
-    v_28 = tint_GammaCorrection(mul(tint_GammaCorrection(v_27, v_29), params.gamutConversionMatrix), v_30);
+    tint_GammaTransferParams v_26 = params.gammaDecodeParams;
+    tint_GammaTransferParams v_27 = params.gammaEncodeParams;
+    v_25 = tint_GammaCorrection(mul(tint_GammaCorrection(v_24, v_26), params.gamutConversionMatrix), v_27);
   } else {
-    v_28 = v_27;
+    v_25 = v_24;
   }
-  return float4(v_28, v_8);
+  return float4(v_25, v_8);
 }
 
-float3x2 v_31(uint start_byte_offset) {
-  uint4 v_32 = t_params[(start_byte_offset / 16u)];
-  float2 v_33 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_32.zw) : (v_32.xy)));
-  uint4 v_34 = t_params[((8u + start_byte_offset) / 16u)];
-  float2 v_35 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_34.zw) : (v_34.xy)));
-  uint4 v_36 = t_params[((16u + start_byte_offset) / 16u)];
-  return float3x2(v_33, v_35, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_36.zw) : (v_36.xy))));
+float3x2 v_28(uint start_byte_offset) {
+  uint4 v_29 = t_params[(start_byte_offset / 16u)];
+  float2 v_30 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_29.zw) : (v_29.xy)));
+  uint4 v_31 = t_params[((8u + start_byte_offset) / 16u)];
+  float2 v_32 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_31.zw) : (v_31.xy)));
+  uint4 v_33 = t_params[((16u + start_byte_offset) / 16u)];
+  return float3x2(v_30, v_32, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_33.zw) : (v_33.xy))));
 }
 
-float3x3 v_37(uint start_byte_offset) {
-  float3 v_38 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
-  float3 v_39 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
-  return float3x3(v_38, v_39, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
+float3x3 v_34(uint start_byte_offset) {
+  float3 v_35 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
+  float3 v_36 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
+  return float3x3(v_35, v_36, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
 }
 
-tint_GammaTransferParams v_40(uint start_byte_offset) {
-  float v_41 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
-  float v_42 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
-  float v_43 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
-  float v_44 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
-  float v_45 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
-  float v_46 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
-  float v_47 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
-  tint_GammaTransferParams v_48 = {v_41, v_42, v_43, v_44, v_45, v_46, v_47, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
-  return v_48;
+tint_GammaTransferParams v_37(uint start_byte_offset) {
+  float v_38 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
+  float v_39 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
+  float v_40 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
+  float v_41 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
+  float v_42 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
+  float v_43 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
+  float v_44 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
+  tint_GammaTransferParams v_45 = {v_38, v_39, v_40, v_41, v_42, v_43, v_44, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
+  return v_45;
 }
 
-float3x4 v_49(uint start_byte_offset) {
-  float4 v_50 = asfloat(t_params[(start_byte_offset / 16u)]);
-  float4 v_51 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
-  return float3x4(v_50, v_51, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
+float3x4 v_46(uint start_byte_offset) {
+  float4 v_47 = asfloat(t_params[(start_byte_offset / 16u)]);
+  float4 v_48 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
+  return float3x4(v_47, v_48, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
 }
 
-tint_ExternalTextureParams v_52(uint start_byte_offset) {
-  uint v_53 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
-  uint v_54 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
-  float3x4 v_55 = v_49((16u + start_byte_offset));
-  tint_GammaTransferParams v_56 = v_40((64u + start_byte_offset));
-  tint_GammaTransferParams v_57 = v_40((96u + start_byte_offset));
-  float3x3 v_58 = v_37((128u + start_byte_offset));
-  float3x2 v_59 = v_31((176u + start_byte_offset));
-  float3x2 v_60 = v_31((200u + start_byte_offset));
-  uint4 v_61 = t_params[((224u + start_byte_offset) / 16u)];
-  float2 v_62 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_61.zw) : (v_61.xy)));
-  uint4 v_63 = t_params[((232u + start_byte_offset) / 16u)];
-  float2 v_64 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_63.zw) : (v_63.xy)));
-  uint4 v_65 = t_params[((240u + start_byte_offset) / 16u)];
-  float2 v_66 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_65.zw) : (v_65.xy)));
-  uint4 v_67 = t_params[((248u + start_byte_offset) / 16u)];
-  float2 v_68 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_67.zw) : (v_67.xy)));
-  uint4 v_69 = t_params[((256u + start_byte_offset) / 16u)];
-  uint2 v_70 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_69.zw) : (v_69.xy));
-  uint4 v_71 = t_params[((264u + start_byte_offset) / 16u)];
-  tint_ExternalTextureParams v_72 = {v_53, v_54, v_55, v_56, v_57, v_58, v_59, v_60, v_62, v_64, v_66, v_68, v_70, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_71.zw) : (v_71.xy)))};
-  return v_72;
+tint_ExternalTextureParams v_49(uint start_byte_offset) {
+  uint v_50 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
+  uint v_51 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
+  float3x4 v_52 = v_46((16u + start_byte_offset));
+  tint_GammaTransferParams v_53 = v_37((64u + start_byte_offset));
+  tint_GammaTransferParams v_54 = v_37((96u + start_byte_offset));
+  float3x3 v_55 = v_34((128u + start_byte_offset));
+  float3x2 v_56 = v_28((176u + start_byte_offset));
+  float3x2 v_57 = v_28((200u + start_byte_offset));
+  uint4 v_58 = t_params[((224u + start_byte_offset) / 16u)];
+  float2 v_59 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_58.zw) : (v_58.xy)));
+  uint4 v_60 = t_params[((232u + start_byte_offset) / 16u)];
+  float2 v_61 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_60.zw) : (v_60.xy)));
+  uint4 v_62 = t_params[((240u + start_byte_offset) / 16u)];
+  float2 v_63 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_62.zw) : (v_62.xy)));
+  uint4 v_64 = t_params[((248u + start_byte_offset) / 16u)];
+  float2 v_65 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_64.zw) : (v_64.xy)));
+  uint4 v_66 = t_params[((256u + start_byte_offset) / 16u)];
+  uint2 v_67 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_66.zw) : (v_66.xy));
+  uint4 v_68 = t_params[((264u + start_byte_offset) / 16u)];
+  tint_ExternalTextureParams v_69 = {v_50, v_51, v_52, v_53, v_54, v_55, v_56, v_57, v_59, v_61, v_63, v_65, v_67, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_68.zw) : (v_68.xy)))};
+  return v_69;
 }
 
 [numthreads(1, 1, 1)]
 void main() {
-  tint_ExternalTextureParams v_73 = v_52(0u);
-  float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_73, uint2((int(10)).xx));
+  tint_ExternalTextureParams v_70 = v_49(0u);
+  float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_70, uint2((int(10)).xx));
   outImage[(int(0)).xx] = red;
-  tint_ExternalTextureParams v_74 = v_52(0u);
-  float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_74, uint2(int2(int(70), int(118))));
+  tint_ExternalTextureParams v_71 = v_49(0u);
+  float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_71, uint2(int2(int(70), int(118))));
   outImage[int2(int(1), int(0))] = green;
 }
 

test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl

diff --git a/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl
index e6e322dc..32a4cf0 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl
@@ -51,112 +51,109 @@
   float3 v_7 = (0.0f).xxx;
   float v_8 = 0.0f;
   if ((params.numPlanes == 1u)) {
-    uint2 v_9 = (0u).xx;
-    plane_0.GetDimensions(v_9[0u], v_9[1u]);
-    uint2 v_10 = min(v_6, (v_9 - (1u).xx));
+    uint3 v_9 = (0u).xxx;
+    plane_0.GetDimensions(0u, v_9[0u], v_9[1u], v_9[2u]);
+    uint v_10 = min(0u, (v_9.z - 1u));
     uint3 v_11 = (0u).xxx;
-    plane_0.GetDimensions(0u, v_11[0u], v_11[1u], v_11[2u]);
-    uint v_12 = min(0u, (v_11.z - 1u));
-    int2 v_13 = int2(v_10);
-    float4 v_14 = float4(plane_0.Load(int3(v_13, int(v_12))));
-    v_7 = v_14.xyz;
-    v_8 = v_14[3u];
+    plane_0.GetDimensions(uint(v_10), v_11[0u], v_11[1u], v_11[2u]);
+    int2 v_12 = int2(min(v_6, (v_11.xy - (1u).xx)));
+    float4 v_13 = float4(plane_0.Load(int3(v_12, int(v_10))));
+    v_7 = v_13.xyz;
+    v_8 = v_13[3u];
   } else {
-    uint2 v_15 = (0u).xx;
-    plane_0.GetDimensions(v_15[0u], v_15[1u]);
-    uint2 v_16 = min(v_6, (v_15 - (1u).xx));
-    uint3 v_17 = (0u).xxx;
-    plane_0.GetDimensions(0u, v_17[0u], v_17[1u], v_17[2u]);
-    uint v_18 = min(0u, (v_17.z - 1u));
-    int2 v_19 = int2(v_16);
-    float v_20 = float4(plane_0.Load(int3(v_19, int(v_18))))[0u];
-    uint2 v_21 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
-    uint2 v_22 = (0u).xx;
-    plane_1.GetDimensions(v_22[0u], v_22[1u]);
-    uint2 v_23 = min(v_21, (v_22 - (1u).xx));
-    uint3 v_24 = (0u).xxx;
-    plane_1.GetDimensions(0u, v_24[0u], v_24[1u], v_24[2u]);
-    uint v_25 = min(0u, (v_24.z - 1u));
-    int2 v_26 = int2(v_23);
-    v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_20, float4(plane_1.Load(int3(v_26, int(v_25)))).xy, 1.0f));
+    uint3 v_14 = (0u).xxx;
+    plane_0.GetDimensions(0u, v_14[0u], v_14[1u], v_14[2u]);
+    uint v_15 = min(0u, (v_14.z - 1u));
+    uint3 v_16 = (0u).xxx;
+    plane_0.GetDimensions(uint(v_15), v_16[0u], v_16[1u], v_16[2u]);
+    int2 v_17 = int2(min(v_6, (v_16.xy - (1u).xx)));
+    float v_18 = float4(plane_0.Load(int3(v_17, int(v_15))))[0u];
+    uint2 v_19 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
+    uint3 v_20 = (0u).xxx;
+    plane_1.GetDimensions(0u, v_20[0u], v_20[1u], v_20[2u]);
+    uint v_21 = min(0u, (v_20.z - 1u));
+    uint3 v_22 = (0u).xxx;
+    plane_1.GetDimensions(uint(v_21), v_22[0u], v_22[1u], v_22[2u]);
+    int2 v_23 = int2(min(v_19, (v_22.xy - (1u).xx)));
+    v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_18, float4(plane_1.Load(int3(v_23, int(v_21)))).xy, 1.0f));
     v_8 = 1.0f;
   }
-  float3 v_27 = v_7;
-  float3 v_28 = (0.0f).xxx;
+  float3 v_24 = v_7;
+  float3 v_25 = (0.0f).xxx;
   if ((params.doYuvToRgbConversionOnly == 0u)) {
-    tint_GammaTransferParams v_29 = params.gammaDecodeParams;
-    tint_GammaTransferParams v_30 = params.gammaEncodeParams;
-    v_28 = tint_GammaCorrection(mul(tint_GammaCorrection(v_27, v_29), params.gamutConversionMatrix), v_30);
+    tint_GammaTransferParams v_26 = params.gammaDecodeParams;
+    tint_GammaTransferParams v_27 = params.gammaEncodeParams;
+    v_25 = tint_GammaCorrection(mul(tint_GammaCorrection(v_24, v_26), params.gamutConversionMatrix), v_27);
   } else {
-    v_28 = v_27;
+    v_25 = v_24;
   }
-  return float4(v_28, v_8);
+  return float4(v_25, v_8);
 }
 
-float3x2 v_31(uint start_byte_offset) {
-  uint4 v_32 = t_params[(start_byte_offset / 16u)];
-  float2 v_33 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_32.zw) : (v_32.xy)));
-  uint4 v_34 = t_params[((8u + start_byte_offset) / 16u)];
-  float2 v_35 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_34.zw) : (v_34.xy)));
-  uint4 v_36 = t_params[((16u + start_byte_offset) / 16u)];
-  return float3x2(v_33, v_35, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_36.zw) : (v_36.xy))));
+float3x2 v_28(uint start_byte_offset) {
+  uint4 v_29 = t_params[(start_byte_offset / 16u)];
+  float2 v_30 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_29.zw) : (v_29.xy)));
+  uint4 v_31 = t_params[((8u + start_byte_offset) / 16u)];
+  float2 v_32 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_31.zw) : (v_31.xy)));
+  uint4 v_33 = t_params[((16u + start_byte_offset) / 16u)];
+  return float3x2(v_30, v_32, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_33.zw) : (v_33.xy))));
 }
 
-float3x3 v_37(uint start_byte_offset) {
-  float3 v_38 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
-  float3 v_39 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
-  return float3x3(v_38, v_39, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
+float3x3 v_34(uint start_byte_offset) {
+  float3 v_35 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
+  float3 v_36 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
+  return float3x3(v_35, v_36, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
 }
 
-tint_GammaTransferParams v_40(uint start_byte_offset) {
-  float v_41 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
-  float v_42 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
-  float v_43 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
-  float v_44 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
-  float v_45 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
-  float v_46 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
-  float v_47 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
-  tint_GammaTransferParams v_48 = {v_41, v_42, v_43, v_44, v_45, v_46, v_47, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
-  return v_48;
+tint_GammaTransferParams v_37(uint start_byte_offset) {
+  float v_38 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
+  float v_39 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
+  float v_40 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
+  float v_41 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
+  float v_42 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
+  float v_43 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
+  float v_44 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
+  tint_GammaTransferParams v_45 = {v_38, v_39, v_40, v_41, v_42, v_43, v_44, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
+  return v_45;
 }
 
-float3x4 v_49(uint start_byte_offset) {
-  float4 v_50 = asfloat(t_params[(start_byte_offset / 16u)]);
-  float4 v_51 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
-  return float3x4(v_50, v_51, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
+float3x4 v_46(uint start_byte_offset) {
+  float4 v_47 = asfloat(t_params[(start_byte_offset / 16u)]);
+  float4 v_48 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
+  return float3x4(v_47, v_48, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
 }
 
-tint_ExternalTextureParams v_52(uint start_byte_offset) {
-  uint v_53 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
-  uint v_54 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
-  float3x4 v_55 = v_49((16u + start_byte_offset));
-  tint_GammaTransferParams v_56 = v_40((64u + start_byte_offset));
-  tint_GammaTransferParams v_57 = v_40((96u + start_byte_offset));
-  float3x3 v_58 = v_37((128u + start_byte_offset));
-  float3x2 v_59 = v_31((176u + start_byte_offset));
-  float3x2 v_60 = v_31((200u + start_byte_offset));
-  uint4 v_61 = t_params[((224u + start_byte_offset) / 16u)];
-  float2 v_62 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_61.zw) : (v_61.xy)));
-  uint4 v_63 = t_params[((232u + start_byte_offset) / 16u)];
-  float2 v_64 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_63.zw) : (v_63.xy)));
-  uint4 v_65 = t_params[((240u + start_byte_offset) / 16u)];
-  float2 v_66 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_65.zw) : (v_65.xy)));
-  uint4 v_67 = t_params[((248u + start_byte_offset) / 16u)];
-  float2 v_68 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_67.zw) : (v_67.xy)));
-  uint4 v_69 = t_params[((256u + start_byte_offset) / 16u)];
-  uint2 v_70 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_69.zw) : (v_69.xy));
-  uint4 v_71 = t_params[((264u + start_byte_offset) / 16u)];
-  tint_ExternalTextureParams v_72 = {v_53, v_54, v_55, v_56, v_57, v_58, v_59, v_60, v_62, v_64, v_66, v_68, v_70, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_71.zw) : (v_71.xy)))};
-  return v_72;
+tint_ExternalTextureParams v_49(uint start_byte_offset) {
+  uint v_50 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
+  uint v_51 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
+  float3x4 v_52 = v_46((16u + start_byte_offset));
+  tint_GammaTransferParams v_53 = v_37((64u + start_byte_offset));
+  tint_GammaTransferParams v_54 = v_37((96u + start_byte_offset));
+  float3x3 v_55 = v_34((128u + start_byte_offset));
+  float3x2 v_56 = v_28((176u + start_byte_offset));
+  float3x2 v_57 = v_28((200u + start_byte_offset));
+  uint4 v_58 = t_params[((224u + start_byte_offset) / 16u)];
+  float2 v_59 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_58.zw) : (v_58.xy)));
+  uint4 v_60 = t_params[((232u + start_byte_offset) / 16u)];
+  float2 v_61 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_60.zw) : (v_60.xy)));
+  uint4 v_62 = t_params[((240u + start_byte_offset) / 16u)];
+  float2 v_63 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_62.zw) : (v_62.xy)));
+  uint4 v_64 = t_params[((248u + start_byte_offset) / 16u)];
+  float2 v_65 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_64.zw) : (v_64.xy)));
+  uint4 v_66 = t_params[((256u + start_byte_offset) / 16u)];
+  uint2 v_67 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_66.zw) : (v_66.xy));
+  uint4 v_68 = t_params[((264u + start_byte_offset) / 16u)];
+  tint_ExternalTextureParams v_69 = {v_50, v_51, v_52, v_53, v_54, v_55, v_56, v_57, v_59, v_61, v_63, v_65, v_67, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_68.zw) : (v_68.xy)))};
+  return v_69;
 }
 
 [numthreads(1, 1, 1)]
 void main() {
-  tint_ExternalTextureParams v_73 = v_52(0u);
-  float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_73, uint2((int(10)).xx));
+  tint_ExternalTextureParams v_70 = v_49(0u);
+  float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_70, uint2((int(10)).xx));
   outImage[(int(0)).xx] = red;
-  tint_ExternalTextureParams v_74 = v_52(0u);
-  float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_74, uint2(int2(int(70), int(118))));
+  tint_ExternalTextureParams v_71 = v_49(0u);
+  float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_71, uint2(int2(int(70), int(118))));
   outImage[int2(int(1), int(0))] = green;
 }