[multiplanar] Add an external texture end2end test. This CL adds an external texture end2end test. Two variants are added, one that just uses regular multiplanar and one that will trigger YCBCR handling in the SPIR-V backend. Bug: 491363837 Change-Id: Iee7369b7f5ee9376b33d302be4ee8ac77fd40676 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/296935 Reviewed-by: James Price <jrprice@google.com> Commit-Queue: dan sinclair <dsinclair@chromium.org>

diff --git a/.clang-tidy b/.clang-tidy
index e258664..5354ab9 100644
--- a/.clang-tidy
+++ b/.clang-tidy

@@ -192,6 +192,7 @@
           -readability-convert-member-functions-to-static,\
           -readability-else-after-return,\
           -readability-enum-initial-value,\
+          -readability-fn_size,\
           -readability-function-cognitive-complexity,\
           -readability-function-size,\
           -readability-identifier-length,\

diff --git a/CPPLINT.cfg b/CPPLINT.cfg
index c5222f9..4ec811a 100644
--- a/CPPLINT.cfg
+++ b/CPPLINT.cfg

@@ -26,3 +26,5 @@
 # WGSL constructor-casts (in DAWN_MULTILINE) look like function-style casts.
 # Disable this globally for ease of use, as it's also not very important in C++.
 filter=-readability/casting
+
+filter=-readability/fn_size

diff --git a/src/tint/api/helpers/generate_bindings.cc b/src/tint/api/helpers/generate_bindings.cc
index 40120b9..cf94729 100644
--- a/src/tint/api/helpers/generate_bindings.cc
+++ b/src/tint/api/helpers/generate_bindings.cc

@@ -28,7 +28,6 @@
 #include "src/tint/api/helpers/generate_bindings.h"
 
 #include <algorithm>
-#include <iostream>
 
 #include "src/tint/api/common/binding_point.h"
 #include "src/tint/lang/core/ir/module.h"
@@ -44,7 +43,8 @@
 Bindings GenerateBindings(const core::ir::Module& module,
                           const std::string& ep,
                           bool set_group_to_zero,
-                          bool flatten_bindings) {
+                          bool flatten_bindings,
+                          std::unordered_set<tint::BindingPoint> ycbcr_bindings) {
     Bindings bindings{};
 
     uint32_t next_buffer_idx = 0;
@@ -152,10 +152,17 @@
             uint32_t g = set_group_to_zero ? 0 : bp.group;
 
             tint::BindingPoint plane0{.group = g, .binding = next_texture_idx++};
-            tint::BindingPoint plane1{.group = g, .binding = next_texture_idx++};
             tint::BindingPoint metadata{.group = g, .binding = next_buffer_idx++};
-            bindings.external_texture.emplace(bp,
-                                              ExternalMultiplanarTexture{metadata, plane0, plane1});
+
+            if (ycbcr_bindings.contains(bp)) {
+                tint::BindingPoint sampler{.group = g, .binding = next_sampler_idx++};
+                bindings.external_texture.emplace(bp,
+                                                  ExternalYCBCRTexture{metadata, plane0, sampler});
+            } else {
+                tint::BindingPoint plane1{.group = g, .binding = next_texture_idx++};
+                bindings.external_texture.emplace(
+                    bp, ExternalMultiplanarTexture{metadata, plane0, plane1});
+            }
         }
     } else {
         for (auto bp : ext_tex_bps) {
@@ -165,8 +172,14 @@
             tint::BindingPoint plane0{.group = g, .binding = bp.binding};
             tint::BindingPoint plane1{.group = g, .binding = next_num++};
             tint::BindingPoint metadata{.group = g, .binding = next_num++};
-            bindings.external_texture.emplace(bp,
-                                              ExternalMultiplanarTexture{metadata, plane0, plane1});
+
+            if (ycbcr_bindings.contains(bp)) {
+                bindings.external_texture.emplace(bp,
+                                                  ExternalYCBCRTexture{metadata, plane0, plane1});
+            } else {
+                bindings.external_texture.emplace(
+                    bp, ExternalMultiplanarTexture{metadata, plane0, plane1});
+            }
         }
     }
 

diff --git a/src/tint/api/helpers/generate_bindings.h b/src/tint/api/helpers/generate_bindings.h
index 9878522..dcaf001 100644
--- a/src/tint/api/helpers/generate_bindings.h
+++ b/src/tint/api/helpers/generate_bindings.h

@@ -29,6 +29,7 @@
 #define SRC_TINT_API_HELPERS_GENERATE_BINDINGS_H_
 
 #include <string>
+#include <unordered_set>
 
 #include "src/tint/api/common/bindings.h"
 
@@ -43,11 +44,13 @@
 /// @param module the module to generate from
 /// @param set_group_to_zero if true, the group used for bindings will always be zero
 /// @param flatten_bindings if true, the bindings will remap to count from 0
+/// @param ycbcr_bindings flags a binding as YCBCR data.
 /// @returns the bindings
 Bindings GenerateBindings(const core::ir::Module& module,
                           const std::string& ep,
                           bool set_group_to_zero,
-                          bool flatten_bindings);
+                          bool flatten_bindings,
+                          std::unordered_set<tint::BindingPoint> ycbcr_bindings = {});
 
 }  // namespace tint
 

diff --git a/src/tint/cmd/tint/main.cc b/src/tint/cmd/tint/main.cc
index ac3d8fa..233b55a 100644
--- a/src/tint/cmd/tint/main.cc
+++ b/src/tint/cmd/tint/main.cc

@@ -162,6 +162,8 @@
 #if TINT_BUILD_SPV_WRITER
     bool use_storage_input_output_16 = true;
     tint::spirv::writer::SpvVersion spirv_version = tint::spirv::writer::SpvVersion::kSpv13;
+
+    std::unordered_set<tint::BindingPoint> ycbcr_bindings;
 #endif  // TINT_BULD_SPV_WRITER
 
 #if TINT_BUILD_MSL_WRITER
@@ -445,6 +447,15 @@
 Valid values are 1.3 and 1.4)",
         version_enum_names, Default{tint::spirv::writer::SpvVersion::kSpv13});
     TINT_DEFER(opts->spirv_version = *spirv_version.value);
+
+    auto& ycbcr_binding_data = options.Add<StringOption>(
+        "ycbcr-bindings",
+        "Allows setting an external texture as YCBCR. "
+        "This allows specifying that the external texture at the given binding point should be "
+        "considered a YCBCR texture. Entries are provided as binding point pairs (group, binding) "
+        "(e.g. 1,2). Multiple entries should be separated with a space.",
+        Default{""});
+
 #endif  // TINT_BUILD_SPV_WRITER
 
 #if TINT_BUILD_GLSL_WRITER
@@ -616,6 +627,37 @@
     }
 #endif  // TINT_BUILD_SPV_READER
 
+#if TINT_BUILD_SPV_WRITER
+    if (!ycbcr_binding_data.value->empty()) {
+        auto str_to_bp = [](const std::string_view& str) -> std::optional<tint::BindingPoint> {
+            auto parts = tint::Split(str, ",");
+            if (parts.Length() != 2) {
+                std::cerr << "A binding point requires a 'group,binding' pair, found "
+                          << parts.Length() << " components instead of 2.\n";
+                return std::nullopt;
+            }
+
+            uint32_t group = 0;
+            std::from_chars(parts[0].data(), parts[0].data() + parts[0].size(), group);
+
+            uint32_t binding = 0;
+            std::from_chars(parts[1].data(), parts[1].data() + parts[1].size(), binding);
+
+            return {tint::BindingPoint{group, binding}};
+        };
+
+        for (auto mapping : tint::Split(*ycbcr_binding_data.value, " ")) {
+            auto opt_src = str_to_bp(mapping);
+            if (!opt_src.has_value()) {
+                return false;
+            }
+            tint::BindingPoint src_bp = opt_src.value();
+            opts->ycbcr_bindings.emplace(src_bp);
+        }
+    }
+
+#endif
+
 #if TINT_BUILD_MSL_WRITER
     if (arg_buffer.value.has_value()) {
         for (auto ab : tint::Split(*arg_buffer.value, ",")) {
@@ -931,7 +973,8 @@
         offset += 8;
     }
 
-    gen_options.bindings = tint::GenerateBindings(ir, options.ep_name, false, false);
+    gen_options.bindings =
+        tint::GenerateBindings(ir, options.ep_name, false, false, options.ycbcr_bindings);
     gen_options.resource_table = tint::core::ir::transform::GenerateResourceTableConfig(ir);
 
     // Enable the Vulkan Memory Model if needed.

diff --git a/test/tint/types/texture/external/external.wgsl b/test/tint/types/texture/external/external.wgsl
new file mode 100644
index 0000000..f543079
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var t_f : texture_external;
+
+@group(1) @binding(0) var<storage, read_write> out : vec4f;
+
+@compute @workgroup_size(1)
+fn main() {
+  out = textureLoad(t_f, vec2i(0));
+}

diff --git a/test/tint/types/texture/external/external.wgsl.expected.dxc.hlsl b/test/tint/types/texture/external/external.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..244f20e
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl.expected.dxc.hlsl

@@ -0,0 +1,127 @@
+struct tint_GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+  float3x2 sampleTransform;
+  float3x2 loadTransform;
+  float2 samplePlane0RectMin;
+  float2 samplePlane0RectMax;
+  float2 samplePlane1RectMin;
+  float2 samplePlane1RectMax;
+  uint2 apparentSize;
+  float2 plane1CoordFactor;
+};
+
+
+cbuffer cbuffer_t_f_params : register(b2) {
+  uint4 t_f_params[17];
+};
+Texture2D<float4> t_f_plane0 : register(t0);
+Texture2D<float4> t_f_plane1 : register(t1);
+RWByteAddressBuffer v_1 : register(u0, space1);
+uint2 tint_v2f32_to_v2u32(float2 value) {
+  return uint2(clamp(value, (0.0f).xx, (4294967040.0f).xx));
+}
+
+float3 tint_GammaCorrection(float3 v, tint_GammaTransferParams params) {
+  float3 v_2 = float3((params.G).xxx);
+  float3 v_3 = float3((params.D).xxx);
+  float3 v_4 = float3(sign(v));
+  return (((abs(v) < v_3)) ? ((v_4 * ((params.C * abs(v)) + params.F))) : ((v_4 * (pow(((params.A * abs(v)) + params.B), v_2) + params.E))));
+}
+
+float4 tint_TextureLoadMultiplanarExternal(Texture2D<float4> plane_0, Texture2D<float4> plane_1, tint_ExternalTextureParams params, uint2 coords) {
+  float2 v_5 = round(mul(float3(float2(min(coords, params.apparentSize)), 1.0f), params.loadTransform));
+  uint2 v_6 = tint_v2f32_to_v2u32(v_5);
+  float3 v_7 = (0.0f).xxx;
+  float v_8 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    int2 v_9 = int2(v_6);
+    float4 v_10 = plane_0.Load(int3(v_9, int(0u)));
+    v_7 = v_10.xyz;
+    v_8 = v_10.w;
+  } else {
+    int2 v_11 = int2(v_6);
+    float v_12 = plane_0.Load(int3(v_11, int(0u))).x;
+    int2 v_13 = int2(tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor)));
+    v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_12, plane_1.Load(int3(v_13, int(0u))).xy, 1.0f));
+    v_8 = 1.0f;
+  }
+  float3 v_14 = v_7;
+  float3 v_15 = (0.0f).xxx;
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    tint_GammaTransferParams v_16 = params.gammaDecodeParams;
+    tint_GammaTransferParams v_17 = params.gammaEncodeParams;
+    v_15 = tint_GammaCorrection(mul(tint_GammaCorrection(v_14, v_16), params.gamutConversionMatrix), v_17);
+  } else {
+    v_15 = v_14;
+  }
+  return float4(v_15, v_8);
+}
+
+float3x2 v_18(uint start_byte_offset) {
+  uint4 v_19 = t_f_params[(start_byte_offset / 16u)];
+  float2 v_20 = asfloat((((((start_byte_offset & 15u) >> 2u) == 2u)) ? (v_19.zw) : (v_19.xy)));
+  uint4 v_21 = t_f_params[((8u + start_byte_offset) / 16u)];
+  float2 v_22 = asfloat(((((((8u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_21.zw) : (v_21.xy)));
+  uint4 v_23 = t_f_params[((16u + start_byte_offset) / 16u)];
+  return float3x2(v_20, v_22, asfloat(((((((16u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_23.zw) : (v_23.xy))));
+}
+
+float3x3 v_24(uint start_byte_offset) {
+  return float3x3(asfloat(t_f_params[(start_byte_offset / 16u)].xyz), asfloat(t_f_params[((16u + start_byte_offset) / 16u)].xyz), asfloat(t_f_params[((32u + start_byte_offset) / 16u)].xyz));
+}
+
+tint_GammaTransferParams v_25(uint start_byte_offset) {
+  tint_GammaTransferParams v_26 = {asfloat(t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)]), asfloat(t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) & 15u) >> 2u)]), t_f_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) & 15u) >> 2u)]};
+  return v_26;
+}
+
+float3x4 v_27(uint start_byte_offset) {
+  return float3x4(asfloat(t_f_params[(start_byte_offset / 16u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)]), asfloat(t_f_params[((32u + start_byte_offset) / 16u)]));
+}
+
+tint_ExternalTextureParams v_28(uint start_byte_offset) {
+  uint v_29 = t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)];
+  uint v_30 = t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)];
+  float3x4 v_31 = v_27((16u + start_byte_offset));
+  tint_GammaTransferParams v_32 = v_25((64u + start_byte_offset));
+  tint_GammaTransferParams v_33 = v_25((96u + start_byte_offset));
+  float3x3 v_34 = v_24((128u + start_byte_offset));
+  float3x2 v_35 = v_18((176u + start_byte_offset));
+  float3x2 v_36 = v_18((200u + start_byte_offset));
+  uint4 v_37 = t_f_params[((224u + start_byte_offset) / 16u)];
+  float2 v_38 = asfloat(((((((224u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_37.zw) : (v_37.xy)));
+  uint4 v_39 = t_f_params[((232u + start_byte_offset) / 16u)];
+  float2 v_40 = asfloat(((((((232u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_39.zw) : (v_39.xy)));
+  uint4 v_41 = t_f_params[((240u + start_byte_offset) / 16u)];
+  float2 v_42 = asfloat(((((((240u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_41.zw) : (v_41.xy)));
+  uint4 v_43 = t_f_params[((248u + start_byte_offset) / 16u)];
+  float2 v_44 = asfloat(((((((248u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_43.zw) : (v_43.xy)));
+  uint4 v_45 = t_f_params[((256u + start_byte_offset) / 16u)];
+  uint2 v_46 = ((((((256u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_45.zw) : (v_45.xy));
+  uint4 v_47 = t_f_params[((264u + start_byte_offset) / 16u)];
+  tint_ExternalTextureParams v_48 = {v_29, v_30, v_31, v_32, v_33, v_34, v_35, v_36, v_38, v_40, v_42, v_44, v_46, asfloat(((((((264u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_47.zw) : (v_47.xy)))};
+  return v_48;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_ExternalTextureParams v_49 = v_28(0u);
+  v_1.Store4(0u, asuint(tint_TextureLoadMultiplanarExternal(t_f_plane0, t_f_plane1, v_49, uint2((int(0)).xx))));
+}
+

diff --git a/test/tint/types/texture/external/external.wgsl.expected.fxc.hlsl b/test/tint/types/texture/external/external.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..244f20e
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl.expected.fxc.hlsl

@@ -0,0 +1,127 @@
+struct tint_GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+  float3x2 sampleTransform;
+  float3x2 loadTransform;
+  float2 samplePlane0RectMin;
+  float2 samplePlane0RectMax;
+  float2 samplePlane1RectMin;
+  float2 samplePlane1RectMax;
+  uint2 apparentSize;
+  float2 plane1CoordFactor;
+};
+
+
+cbuffer cbuffer_t_f_params : register(b2) {
+  uint4 t_f_params[17];
+};
+Texture2D<float4> t_f_plane0 : register(t0);
+Texture2D<float4> t_f_plane1 : register(t1);
+RWByteAddressBuffer v_1 : register(u0, space1);
+uint2 tint_v2f32_to_v2u32(float2 value) {
+  return uint2(clamp(value, (0.0f).xx, (4294967040.0f).xx));
+}
+
+float3 tint_GammaCorrection(float3 v, tint_GammaTransferParams params) {
+  float3 v_2 = float3((params.G).xxx);
+  float3 v_3 = float3((params.D).xxx);
+  float3 v_4 = float3(sign(v));
+  return (((abs(v) < v_3)) ? ((v_4 * ((params.C * abs(v)) + params.F))) : ((v_4 * (pow(((params.A * abs(v)) + params.B), v_2) + params.E))));
+}
+
+float4 tint_TextureLoadMultiplanarExternal(Texture2D<float4> plane_0, Texture2D<float4> plane_1, tint_ExternalTextureParams params, uint2 coords) {
+  float2 v_5 = round(mul(float3(float2(min(coords, params.apparentSize)), 1.0f), params.loadTransform));
+  uint2 v_6 = tint_v2f32_to_v2u32(v_5);
+  float3 v_7 = (0.0f).xxx;
+  float v_8 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    int2 v_9 = int2(v_6);
+    float4 v_10 = plane_0.Load(int3(v_9, int(0u)));
+    v_7 = v_10.xyz;
+    v_8 = v_10.w;
+  } else {
+    int2 v_11 = int2(v_6);
+    float v_12 = plane_0.Load(int3(v_11, int(0u))).x;
+    int2 v_13 = int2(tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor)));
+    v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_12, plane_1.Load(int3(v_13, int(0u))).xy, 1.0f));
+    v_8 = 1.0f;
+  }
+  float3 v_14 = v_7;
+  float3 v_15 = (0.0f).xxx;
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    tint_GammaTransferParams v_16 = params.gammaDecodeParams;
+    tint_GammaTransferParams v_17 = params.gammaEncodeParams;
+    v_15 = tint_GammaCorrection(mul(tint_GammaCorrection(v_14, v_16), params.gamutConversionMatrix), v_17);
+  } else {
+    v_15 = v_14;
+  }
+  return float4(v_15, v_8);
+}
+
+float3x2 v_18(uint start_byte_offset) {
+  uint4 v_19 = t_f_params[(start_byte_offset / 16u)];
+  float2 v_20 = asfloat((((((start_byte_offset & 15u) >> 2u) == 2u)) ? (v_19.zw) : (v_19.xy)));
+  uint4 v_21 = t_f_params[((8u + start_byte_offset) / 16u)];
+  float2 v_22 = asfloat(((((((8u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_21.zw) : (v_21.xy)));
+  uint4 v_23 = t_f_params[((16u + start_byte_offset) / 16u)];
+  return float3x2(v_20, v_22, asfloat(((((((16u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_23.zw) : (v_23.xy))));
+}
+
+float3x3 v_24(uint start_byte_offset) {
+  return float3x3(asfloat(t_f_params[(start_byte_offset / 16u)].xyz), asfloat(t_f_params[((16u + start_byte_offset) / 16u)].xyz), asfloat(t_f_params[((32u + start_byte_offset) / 16u)].xyz));
+}
+
+tint_GammaTransferParams v_25(uint start_byte_offset) {
+  tint_GammaTransferParams v_26 = {asfloat(t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)]), asfloat(t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) & 15u) >> 2u)]), t_f_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) & 15u) >> 2u)]};
+  return v_26;
+}
+
+float3x4 v_27(uint start_byte_offset) {
+  return float3x4(asfloat(t_f_params[(start_byte_offset / 16u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)]), asfloat(t_f_params[((32u + start_byte_offset) / 16u)]));
+}
+
+tint_ExternalTextureParams v_28(uint start_byte_offset) {
+  uint v_29 = t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)];
+  uint v_30 = t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)];
+  float3x4 v_31 = v_27((16u + start_byte_offset));
+  tint_GammaTransferParams v_32 = v_25((64u + start_byte_offset));
+  tint_GammaTransferParams v_33 = v_25((96u + start_byte_offset));
+  float3x3 v_34 = v_24((128u + start_byte_offset));
+  float3x2 v_35 = v_18((176u + start_byte_offset));
+  float3x2 v_36 = v_18((200u + start_byte_offset));
+  uint4 v_37 = t_f_params[((224u + start_byte_offset) / 16u)];
+  float2 v_38 = asfloat(((((((224u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_37.zw) : (v_37.xy)));
+  uint4 v_39 = t_f_params[((232u + start_byte_offset) / 16u)];
+  float2 v_40 = asfloat(((((((232u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_39.zw) : (v_39.xy)));
+  uint4 v_41 = t_f_params[((240u + start_byte_offset) / 16u)];
+  float2 v_42 = asfloat(((((((240u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_41.zw) : (v_41.xy)));
+  uint4 v_43 = t_f_params[((248u + start_byte_offset) / 16u)];
+  float2 v_44 = asfloat(((((((248u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_43.zw) : (v_43.xy)));
+  uint4 v_45 = t_f_params[((256u + start_byte_offset) / 16u)];
+  uint2 v_46 = ((((((256u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_45.zw) : (v_45.xy));
+  uint4 v_47 = t_f_params[((264u + start_byte_offset) / 16u)];
+  tint_ExternalTextureParams v_48 = {v_29, v_30, v_31, v_32, v_33, v_34, v_35, v_36, v_38, v_40, v_42, v_44, v_46, asfloat(((((((264u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_47.zw) : (v_47.xy)))};
+  return v_48;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_ExternalTextureParams v_49 = v_28(0u);
+  v_1.Store4(0u, asuint(tint_TextureLoadMultiplanarExternal(t_f_plane0, t_f_plane1, v_49, uint2((int(0)).xx))));
+}
+

diff --git a/test/tint/types/texture/external/external.wgsl.expected.glsl b/test/tint/types/texture/external/external.wgsl.expected.glsl
new file mode 100644
index 0000000..09e28a8
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl.expected.glsl

@@ -0,0 +1,123 @@
+#version 310 es
+
+
+struct tint_GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+  mat3x2 sampleTransform;
+  mat3x2 loadTransform;
+  vec2 samplePlane0RectMin;
+  vec2 samplePlane0RectMax;
+  vec2 samplePlane1RectMin;
+  vec2 samplePlane1RectMax;
+  uvec2 apparentSize;
+  vec2 plane1CoordFactor;
+};
+
+layout(binding = 3, std140)
+uniform t_f_params_block_1_ubo {
+  uvec4 inner[17];
+} v_1;
+layout(binding = 0, std430)
+buffer out_block_1_ssbo {
+  vec4 inner;
+} v_2;
+uniform highp sampler2D t_f_plane0;
+uniform highp sampler2D t_f_plane1;
+vec3 tint_GammaCorrection(vec3 v, tint_GammaTransferParams params) {
+  vec3 v_3 = vec3(params.G);
+  return mix((sign(v) * (pow(((params.A * abs(v)) + params.B), v_3) + params.E)), (sign(v) * ((params.C * abs(v)) + params.F)), lessThan(abs(v), vec3(params.D)));
+}
+vec4 tint_TextureLoadMultiplanarExternal(tint_ExternalTextureParams params, uvec2 coords) {
+  vec2 v_4 = round((params.loadTransform * vec3(vec2(min(coords, params.apparentSize)), 1.0f)));
+  uvec2 v_5 = uvec2(v_4);
+  vec3 v_6 = vec3(0.0f);
+  float v_7 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    ivec2 v_8 = ivec2(v_5);
+    vec4 v_9 = texelFetch(t_f_plane0, v_8, int(0u));
+    v_6 = v_9.xyz;
+    v_7 = v_9.w;
+  } else {
+    ivec2 v_10 = ivec2(v_5);
+    float v_11 = texelFetch(t_f_plane0, v_10, int(0u)).x;
+    ivec2 v_12 = ivec2(uvec2((v_4 * params.plane1CoordFactor)));
+    v_6 = (vec4(v_11, texelFetch(t_f_plane1, v_12, int(0u)).xy, 1.0f) * params.yuvToRgbConversionMatrix);
+    v_7 = 1.0f;
+  }
+  vec3 v_13 = v_6;
+  vec3 v_14 = vec3(0.0f);
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    v_14 = tint_GammaCorrection((params.gamutConversionMatrix * tint_GammaCorrection(v_13, params.gammaDecodeParams)), params.gammaEncodeParams);
+  } else {
+    v_14 = v_13;
+  }
+  return vec4(v_14, v_7);
+}
+mat3x2 v_15(uint start_byte_offset) {
+  uvec4 v_16 = v_1.inner[(start_byte_offset / 16u)];
+  vec2 v_17 = uintBitsToFloat(mix(v_16.xy, v_16.zw, bvec2((((start_byte_offset & 15u) >> 2u) == 2u))));
+  uvec4 v_18 = v_1.inner[((8u + start_byte_offset) / 16u)];
+  vec2 v_19 = uintBitsToFloat(mix(v_18.xy, v_18.zw, bvec2(((((8u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_20 = v_1.inner[((16u + start_byte_offset) / 16u)];
+  return mat3x2(v_17, v_19, uintBitsToFloat(mix(v_20.xy, v_20.zw, bvec2(((((16u + start_byte_offset) & 15u) >> 2u) == 2u)))));
+}
+mat3 v_21(uint start_byte_offset) {
+  return mat3(uintBitsToFloat(v_1.inner[(start_byte_offset / 16u)].xyz), uintBitsToFloat(v_1.inner[((16u + start_byte_offset) / 16u)].xyz), uintBitsToFloat(v_1.inner[((32u + start_byte_offset) / 16u)].xyz));
+}
+tint_GammaTransferParams v_22(uint start_byte_offset) {
+  uvec4 v_23 = v_1.inner[(start_byte_offset / 16u)];
+  uvec4 v_24 = v_1.inner[((4u + start_byte_offset) / 16u)];
+  uvec4 v_25 = v_1.inner[((8u + start_byte_offset) / 16u)];
+  uvec4 v_26 = v_1.inner[((12u + start_byte_offset) / 16u)];
+  uvec4 v_27 = v_1.inner[((16u + start_byte_offset) / 16u)];
+  uvec4 v_28 = v_1.inner[((20u + start_byte_offset) / 16u)];
+  uvec4 v_29 = v_1.inner[((24u + start_byte_offset) / 16u)];
+  uvec4 v_30 = v_1.inner[((28u + start_byte_offset) / 16u)];
+  return tint_GammaTransferParams(uintBitsToFloat(v_23[((start_byte_offset & 15u) >> 2u)]), uintBitsToFloat(v_24[(((4u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_25[(((8u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_26[(((12u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_27[(((16u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_28[(((20u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_29[(((24u + start_byte_offset) & 15u) >> 2u)]), v_30[(((28u + start_byte_offset) & 15u) >> 2u)]);
+}
+mat3x4 v_31(uint start_byte_offset) {
+  return mat3x4(uintBitsToFloat(v_1.inner[(start_byte_offset / 16u)]), uintBitsToFloat(v_1.inner[((16u + start_byte_offset) / 16u)]), uintBitsToFloat(v_1.inner[((32u + start_byte_offset) / 16u)]));
+}
+tint_ExternalTextureParams v_32(uint start_byte_offset) {
+  uvec4 v_33 = v_1.inner[(start_byte_offset / 16u)];
+  uvec4 v_34 = v_1.inner[((4u + start_byte_offset) / 16u)];
+  mat3x4 v_35 = v_31((16u + start_byte_offset));
+  tint_GammaTransferParams v_36 = v_22((64u + start_byte_offset));
+  tint_GammaTransferParams v_37 = v_22((96u + start_byte_offset));
+  mat3 v_38 = v_21((128u + start_byte_offset));
+  mat3x2 v_39 = v_15((176u + start_byte_offset));
+  mat3x2 v_40 = v_15((200u + start_byte_offset));
+  uvec4 v_41 = v_1.inner[((224u + start_byte_offset) / 16u)];
+  vec2 v_42 = uintBitsToFloat(mix(v_41.xy, v_41.zw, bvec2(((((224u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_43 = v_1.inner[((232u + start_byte_offset) / 16u)];
+  vec2 v_44 = uintBitsToFloat(mix(v_43.xy, v_43.zw, bvec2(((((232u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_45 = v_1.inner[((240u + start_byte_offset) / 16u)];
+  vec2 v_46 = uintBitsToFloat(mix(v_45.xy, v_45.zw, bvec2(((((240u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_47 = v_1.inner[((248u + start_byte_offset) / 16u)];
+  vec2 v_48 = uintBitsToFloat(mix(v_47.xy, v_47.zw, bvec2(((((248u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_49 = v_1.inner[((256u + start_byte_offset) / 16u)];
+  uvec2 v_50 = mix(v_49.xy, v_49.zw, bvec2(((((256u + start_byte_offset) & 15u) >> 2u) == 2u)));
+  uvec4 v_51 = v_1.inner[((264u + start_byte_offset) / 16u)];
+  return tint_ExternalTextureParams(v_33[((start_byte_offset & 15u) >> 2u)], v_34[(((4u + start_byte_offset) & 15u) >> 2u)], v_35, v_36, v_37, v_38, v_39, v_40, v_42, v_44, v_46, v_48, v_50, uintBitsToFloat(mix(v_51.xy, v_51.zw, bvec2(((((264u + start_byte_offset) & 15u) >> 2u) == 2u)))));
+}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_ExternalTextureParams v_52 = v_32(0u);
+  v_2.inner = tint_TextureLoadMultiplanarExternal(v_52, min(uvec2(ivec2(0)), ((v_52.apparentSize + uvec2(1u)) - uvec2(1u))));
+}

diff --git a/test/tint/types/texture/external/external.wgsl.expected.msl b/test/tint/types/texture/external/external.wgsl.expected.msl
new file mode 100644
index 0000000..87e668a
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl.expected.msl

@@ -0,0 +1,119 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_GammaTransferParams {
+  /* 0x0000 */ float G;
+  /* 0x0004 */ float A;
+  /* 0x0008 */ float B;
+  /* 0x000c */ float C;
+  /* 0x0010 */ float D;
+  /* 0x0014 */ float E;
+  /* 0x0018 */ float F;
+  /* 0x001c */ uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+  float3x2 sampleTransform;
+  float3x2 loadTransform;
+  float2 samplePlane0RectMin;
+  float2 samplePlane0RectMax;
+  float2 samplePlane1RectMin;
+  float2 samplePlane1RectMax;
+  uint2 apparentSize;
+  float2 plane1CoordFactor;
+};
+
+template<typename T, size_t N>
+struct tint_array {
+  const constant T& operator[](size_t i) const constant { return elements[i]; }
+  device T& operator[](size_t i) device { return elements[i]; }
+  const device T& operator[](size_t i) const device { return elements[i]; }
+  thread T& operator[](size_t i) thread { return elements[i]; }
+  const thread T& operator[](size_t i) const thread { return elements[i]; }
+  threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+  const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+  T elements[N];
+};
+
+struct tint_packed_vec3_f32_array_element {
+  /* 0x0000 */ packed_float3 packed;
+  /* 0x000c */ tint_array<int8_t, 4> tint_pad_1;
+};
+
+struct tint_ExternalTextureParams_packed_vec3 {
+  /* 0x0000 */ uint numPlanes;
+  /* 0x0004 */ uint doYuvToRgbConversionOnly;
+  /* 0x0008 */ tint_array<int8_t, 8> tint_pad;
+  /* 0x0010 */ float3x4 yuvToRgbConversionMatrix;
+  /* 0x0040 */ tint_GammaTransferParams gammaDecodeParams;
+  /* 0x0060 */ tint_GammaTransferParams gammaEncodeParams;
+  /* 0x0080 */ tint_array<tint_packed_vec3_f32_array_element, 3> gamutConversionMatrix;
+  /* 0x00b0 */ float3x2 sampleTransform;
+  /* 0x00c8 */ float3x2 loadTransform;
+  /* 0x00e0 */ float2 samplePlane0RectMin;
+  /* 0x00e8 */ float2 samplePlane0RectMax;
+  /* 0x00f0 */ float2 samplePlane1RectMin;
+  /* 0x00f8 */ float2 samplePlane1RectMax;
+  /* 0x0100 */ uint2 apparentSize;
+  /* 0x0108 */ float2 plane1CoordFactor;
+};
+
+struct tint_module_vars_struct {
+  const constant tint_ExternalTextureParams_packed_vec3* t_f_params;
+  texture2d<float, access::sample> t_f_plane0;
+  texture2d<float, access::sample> t_f_plane1;
+  device float4* out;
+};
+
+float3 tint_GammaCorrection(float3 v, tint_GammaTransferParams params) {
+  float3 const v_1 = float3(params.G);
+  return select((sign(v) * (powr(((params.A * abs(v)) + params.B), v_1) + params.E)), (sign(v) * ((params.C * abs(v)) + params.F)), (abs(v) < float3(params.D)));
+}
+
+float4 tint_TextureLoadMultiplanarExternal(texture2d<float, access::sample> plane_0, texture2d<float, access::sample> plane_1, tint_ExternalTextureParams params, uint2 coords) {
+  float2 const v_2 = rint((params.loadTransform * float3(float2(min(coords, params.apparentSize)), 1.0f)));
+  uint2 const v_3 = uint2(v_2);
+  float3 v_4 = 0.0f;
+  float v_5 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    float4 const v_6 = plane_0.read(v_3, 0u);
+    v_4 = v_6.xyz;
+    v_5 = v_6.w;
+  } else {
+    float const v_7 = plane_0.read(v_3, 0u).x;
+    v_4 = (float4(v_7, plane_1.read(uint2((v_2 * params.plane1CoordFactor)), 0u).xy, 1.0f) * params.yuvToRgbConversionMatrix);
+    v_5 = 1.0f;
+  }
+  float3 const v_8 = v_4;
+  float3 v_9 = 0.0f;
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    v_9 = tint_GammaCorrection((params.gamutConversionMatrix * tint_GammaCorrection(v_8, params.gammaDecodeParams)), params.gammaEncodeParams);
+  } else {
+    v_9 = v_8;
+  }
+  return float4(v_9, v_5);
+}
+
+tint_ExternalTextureParams tint_load_struct_packed_vec3(const constant tint_ExternalTextureParams_packed_vec3* const from) {
+  uint const v_10 = (*from).numPlanes;
+  uint const v_11 = (*from).doYuvToRgbConversionOnly;
+  float3x4 const v_12 = (*from).yuvToRgbConversionMatrix;
+  tint_GammaTransferParams const v_13 = (*from).gammaDecodeParams;
+  tint_GammaTransferParams const v_14 = (*from).gammaEncodeParams;
+  tint_array<tint_packed_vec3_f32_array_element, 3> const v_15 = (*from).gamutConversionMatrix;
+  float3x3 const v_16 = float3x3(float3(v_15[0u].packed), float3(v_15[1u].packed), float3(v_15[2u].packed));
+  return tint_ExternalTextureParams{.numPlanes=v_10, .doYuvToRgbConversionOnly=v_11, .yuvToRgbConversionMatrix=v_12, .gammaDecodeParams=v_13, .gammaEncodeParams=v_14, .gamutConversionMatrix=v_16, .sampleTransform=(*from).sampleTransform, .loadTransform=(*from).loadTransform, .samplePlane0RectMin=(*from).samplePlane0RectMin, .samplePlane0RectMax=(*from).samplePlane0RectMax, .samplePlane1RectMin=(*from).samplePlane1RectMin, .samplePlane1RectMax=(*from).samplePlane1RectMax, .apparentSize=(*from).apparentSize, .plane1CoordFactor=(*from).plane1CoordFactor};
+}
+
+[[max_total_threads_per_threadgroup(1)]]
+kernel void v_17(const constant tint_ExternalTextureParams_packed_vec3* t_f_params [[buffer(1)]], texture2d<float, access::sample> t_f_plane0 [[texture(0)]], texture2d<float, access::sample> t_f_plane1 [[texture(1)]], device float4* out [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t_f_params=t_f_params, .t_f_plane0=t_f_plane0, .t_f_plane1=t_f_plane1, .out=out};
+  tint_ExternalTextureParams const v_18 = tint_load_struct_packed_vec3(tint_module_vars.t_f_params);
+  (*tint_module_vars.out) = tint_TextureLoadMultiplanarExternal(tint_module_vars.t_f_plane0, tint_module_vars.t_f_plane1, v_18, min(uint2(int2(0)), ((v_18.apparentSize + uint2(1u)) - uint2(1u))));
+}

diff --git a/test/tint/types/texture/external/external.wgsl.expected.spvasm b/test/tint/types/texture/external/external.wgsl.expected.spvasm
new file mode 100644
index 0000000..f083fbc
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl.expected.spvasm

@@ -0,0 +1,492 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 400
+; Schema: 0
+               OpCapability Shader
+         %44 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpMemberName %t_f_params_block_tint_explicit_layout 0 "inner"
+               OpName %t_f_params_block_tint_explicit_layout "t_f_params_block_tint_explicit_layout"
+               OpName %t_f_plane0 "t_f_plane0"
+               OpName %t_f_plane1 "t_f_plane1"
+               OpMemberName %out_block 0 "inner"
+               OpName %out_block "out_block"
+               OpName %main "main"
+               OpMemberName %tint_ExternalTextureParams 0 "numPlanes"
+               OpMemberName %tint_ExternalTextureParams 1 "doYuvToRgbConversionOnly"
+               OpMemberName %tint_ExternalTextureParams 2 "yuvToRgbConversionMatrix"
+               OpMemberName %tint_GammaTransferParams 0 "G"
+               OpMemberName %tint_GammaTransferParams 1 "A"
+               OpMemberName %tint_GammaTransferParams 2 "B"
+               OpMemberName %tint_GammaTransferParams 3 "C"
+               OpMemberName %tint_GammaTransferParams 4 "D"
+               OpMemberName %tint_GammaTransferParams 5 "E"
+               OpMemberName %tint_GammaTransferParams 6 "F"
+               OpMemberName %tint_GammaTransferParams 7 "padding"
+               OpName %tint_GammaTransferParams "tint_GammaTransferParams"
+               OpMemberName %tint_ExternalTextureParams 3 "gammaDecodeParams"
+               OpMemberName %tint_ExternalTextureParams 4 "gammaEncodeParams"
+               OpMemberName %tint_ExternalTextureParams 5 "gamutConversionMatrix"
+               OpMemberName %tint_ExternalTextureParams 6 "sampleTransform"
+               OpMemberName %tint_ExternalTextureParams 7 "loadTransform"
+               OpMemberName %tint_ExternalTextureParams 8 "samplePlane0RectMin"
+               OpMemberName %tint_ExternalTextureParams 9 "samplePlane0RectMax"
+               OpMemberName %tint_ExternalTextureParams 10 "samplePlane1RectMin"
+               OpMemberName %tint_ExternalTextureParams 11 "samplePlane1RectMax"
+               OpMemberName %tint_ExternalTextureParams 12 "apparentSize"
+               OpMemberName %tint_ExternalTextureParams 13 "plane1CoordFactor"
+               OpName %tint_ExternalTextureParams "tint_ExternalTextureParams"
+               OpName %tint_TextureLoadMultiplanarExternal "tint_TextureLoadMultiplanarExternal"
+               OpName %plane_0 "plane_0"
+               OpName %plane_1 "plane_1"
+               OpName %params "params"
+               OpName %coords "coords"
+               OpName %tint_GammaCorrection "tint_GammaCorrection"
+               OpName %v "v"
+               OpName %params_0 "params"
+               OpName %start_byte_offset "start_byte_offset"
+               OpName %start_byte_offset_0 "start_byte_offset"
+               OpName %start_byte_offset_1 "start_byte_offset"
+               OpName %start_byte_offset_2 "start_byte_offset"
+               OpName %start_byte_offset_3 "start_byte_offset"
+               OpDecorate %_arr_v4uint_uint_17 ArrayStride 16
+               OpMemberDecorate %t_f_params_block_tint_explicit_layout 0 Offset 0
+               OpDecorate %t_f_params_block_tint_explicit_layout Block
+               OpDecorate %1 DescriptorSet 0
+               OpDecorate %1 Binding 2
+               OpDecorate %1 NonWritable
+               OpDecorate %t_f_plane0 DescriptorSet 0
+               OpDecorate %t_f_plane0 Binding 0
+               OpDecorate %t_f_plane1 DescriptorSet 0
+               OpDecorate %t_f_plane1 Binding 1
+               OpMemberDecorate %out_block 0 Offset 0
+               OpDecorate %out_block Block
+               OpDecorate %13 DescriptorSet 1
+               OpDecorate %13 Binding 0
+               OpDecorate %13 Coherent
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+    %uint_17 = OpConstant %uint 17
+%_arr_v4uint_uint_17 = OpTypeArray %v4uint %uint_17
+%t_f_params_block_tint_explicit_layout = OpTypeStruct %_arr_v4uint_uint_17
+%_ptr_Uniform_t_f_params_block_tint_explicit_layout = OpTypePointer Uniform %t_f_params_block_tint_explicit_layout
+          %1 = OpVariable %_ptr_Uniform_t_f_params_block_tint_explicit_layout Uniform
+      %float = OpTypeFloat 32
+         %10 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
+ %t_f_plane0 = OpVariable %_ptr_UniformConstant_10 UniformConstant
+ %t_f_plane1 = OpVariable %_ptr_UniformConstant_10 UniformConstant
+    %v4float = OpTypeVector %float 4
+  %out_block = OpTypeStruct %v4float
+%_ptr_StorageBuffer_out_block = OpTypePointer StorageBuffer %out_block
+         %13 = OpVariable %_ptr_StorageBuffer_out_block StorageBuffer
+       %void = OpTypeVoid
+         %19 = OpTypeFunction %void
+%mat3v4float = OpTypeMatrix %v4float 3
+%tint_GammaTransferParams = OpTypeStruct %float %float %float %float %float %float %float %uint
+    %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+     %v2uint = OpTypeVector %uint 2
+%tint_ExternalTextureParams = OpTypeStruct %uint %uint %mat3v4float %tint_GammaTransferParams %tint_GammaTransferParams %mat3v3float %mat3v2float %mat3v2float %v2float %v2float %v2float %v2float %v2uint %v2float
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+         %36 = OpConstantComposite %v2uint %uint_1 %uint_1
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %40 = OpConstantNull %v2int
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+         %53 = OpTypeFunction %v4float %10 %10 %tint_ExternalTextureParams %v2uint
+    %float_1 = OpConstant %float 1
+       %bool = OpTypeBool
+        %101 = OpTypeFunction %v3float %v3float %tint_GammaTransferParams
+     %v3bool = OpTypeVector %bool 3
+        %129 = OpTypeFunction %tint_ExternalTextureParams %uint
+    %uint_16 = OpConstant %uint 16
+%_ptr_Uniform_v4uint = OpTypePointer Uniform %v4uint
+    %uint_15 = OpConstant %uint 15
+     %uint_2 = OpConstant %uint 2
+     %uint_4 = OpConstant %uint 4
+    %uint_64 = OpConstant %uint 64
+    %uint_96 = OpConstant %uint 96
+   %uint_128 = OpConstant %uint 128
+   %uint_176 = OpConstant %uint 176
+   %uint_200 = OpConstant %uint 200
+   %uint_224 = OpConstant %uint 224
+     %v2bool = OpTypeVector %bool 2
+   %uint_232 = OpConstant %uint 232
+   %uint_240 = OpConstant %uint 240
+   %uint_248 = OpConstant %uint 248
+   %uint_256 = OpConstant %uint 256
+   %uint_264 = OpConstant %uint 264
+        %250 = OpTypeFunction %mat3v4float %uint
+    %uint_32 = OpConstant %uint 32
+        %269 = OpTypeFunction %tint_GammaTransferParams %uint
+     %uint_8 = OpConstant %uint 8
+    %uint_12 = OpConstant %uint 12
+    %uint_20 = OpConstant %uint 20
+    %uint_24 = OpConstant %uint 24
+    %uint_28 = OpConstant %uint 28
+        %340 = OpTypeFunction %mat3v3float %uint
+     %v3uint = OpTypeVector %uint 3
+        %362 = OpTypeFunction %mat3v2float %uint
+       %main = OpFunction %void None %19
+         %20 = OpLabel
+         %21 = OpLoad %10 %t_f_plane0 None
+         %22 = OpLoad %10 %t_f_plane1 None
+         %23 = OpFunctionCall %tint_ExternalTextureParams %32 %uint_0
+         %34 = OpCompositeExtract %v2uint %23 12
+         %35 = OpIAdd %v2uint %34 %36
+         %38 = OpISub %v2uint %35 %36
+         %39 = OpBitcast %v2uint %40
+         %43 = OpExtInst %v2uint %44 UMin %39 %38
+         %45 = OpFunctionCall %v4float %tint_TextureLoadMultiplanarExternal %21 %22 %23 %43
+         %47 = OpAccessChain %_ptr_StorageBuffer_v4float %13 %uint_0
+               OpStore %47 %45 None
+               OpReturn
+               OpFunctionEnd
+%tint_TextureLoadMultiplanarExternal = OpFunction %v4float None %53
+    %plane_0 = OpFunctionParameter %10
+    %plane_1 = OpFunctionParameter %10
+     %params = OpFunctionParameter %tint_ExternalTextureParams
+     %coords = OpFunctionParameter %v2uint
+         %54 = OpLabel
+         %55 = OpCompositeExtract %uint %params 1
+         %56 = OpCompositeExtract %mat3v4float %params 2
+         %57 = OpCompositeExtract %mat3v2float %params 7
+         %58 = OpCompositeExtract %v2uint %params 12
+         %59 = OpCompositeExtract %v2float %params 13
+         %60 = OpExtInst %v2uint %44 UMin %coords %58
+         %61 = OpConvertUToF %v2float %60
+         %62 = OpCompositeConstruct %v3float %61 %float_1
+         %64 = OpMatrixTimesVector %v2float %57 %62
+         %65 = OpExtInst %v2float %44 RoundEven %64
+         %66 = OpConvertFToU %v2uint %65
+         %67 = OpCompositeExtract %uint %params 0
+         %68 = OpIEqual %bool %67 %uint_1
+               OpSelectionMerge %70 None
+               OpBranchConditional %68 %71 %72
+         %71 = OpLabel
+         %85 = OpImageFetch %v4float %plane_0 %66 Lod %uint_0
+         %74 = OpVectorShuffle %v3float %85 %85 0 1 2
+         %77 = OpCompositeExtract %float %85 3
+               OpBranch %70
+         %72 = OpLabel
+         %86 = OpImageFetch %v4float %plane_0 %66 Lod %uint_0
+         %87 = OpCompositeExtract %float %86 0
+         %88 = OpFMul %v2float %65 %59
+         %89 = OpConvertFToU %v2uint %88
+         %90 = OpImageFetch %v4float %plane_1 %89 Lod %uint_0
+         %91 = OpVectorShuffle %v2float %90 %90 0 1
+         %92 = OpCompositeConstruct %v4float %87 %91 %float_1
+         %75 = OpVectorTimesMatrix %v3float %92 %56
+               OpBranch %70
+         %70 = OpLabel
+         %73 = OpPhi %v3float %74 %71 %75 %72
+         %76 = OpPhi %float %77 %71 %float_1 %72
+         %78 = OpIEqual %bool %55 %uint_0
+               OpSelectionMerge %79 None
+               OpBranchConditional %78 %80 %81
+         %80 = OpLabel
+         %93 = OpCompositeExtract %tint_GammaTransferParams %params 3
+         %94 = OpCompositeExtract %tint_GammaTransferParams %params 4
+         %95 = OpCompositeExtract %mat3v3float %params 5
+         %96 = OpFunctionCall %v3float %tint_GammaCorrection %73 %93
+         %98 = OpMatrixTimesVector %v3float %95 %96
+         %83 = OpFunctionCall %v3float %tint_GammaCorrection %98 %94
+               OpBranch %79
+         %81 = OpLabel
+               OpBranch %79
+         %79 = OpLabel
+         %82 = OpPhi %v3float %83 %80 %73 %81
+         %84 = OpCompositeConstruct %v4float %82 %76
+               OpReturnValue %84
+               OpFunctionEnd
+%tint_GammaCorrection = OpFunction %v3float None %101
+          %v = OpFunctionParameter %v3float
+   %params_0 = OpFunctionParameter %tint_GammaTransferParams
+        %102 = OpLabel
+        %103 = OpCompositeExtract %float %params_0 0
+        %104 = OpCompositeExtract %float %params_0 1
+        %105 = OpCompositeExtract %float %params_0 2
+        %106 = OpCompositeExtract %float %params_0 3
+        %107 = OpCompositeExtract %float %params_0 4
+        %108 = OpCompositeExtract %float %params_0 5
+        %109 = OpCompositeExtract %float %params_0 6
+        %110 = OpCompositeConstruct %v3float %103 %103 %103
+        %111 = OpCompositeConstruct %v3float %107 %107 %107
+        %112 = OpExtInst %v3float %44 FAbs %v
+        %113 = OpExtInst %v3float %44 FSign %v
+        %114 = OpFOrdLessThan %v3bool %112 %111
+        %116 = OpVectorTimesScalar %v3float %112 %106
+        %117 = OpCompositeConstruct %v3float %109 %109 %109
+        %118 = OpFAdd %v3float %116 %117
+        %119 = OpFMul %v3float %113 %118
+        %120 = OpVectorTimesScalar %v3float %112 %104
+        %121 = OpCompositeConstruct %v3float %105 %105 %105
+        %122 = OpFAdd %v3float %120 %121
+        %123 = OpExtInst %v3float %44 Pow %122 %110
+        %124 = OpCompositeConstruct %v3float %108 %108 %108
+        %125 = OpFAdd %v3float %123 %124
+        %126 = OpFMul %v3float %113 %125
+        %127 = OpSelect %v3float %114 %119 %126
+               OpReturnValue %127
+               OpFunctionEnd
+         %32 = OpFunction %tint_ExternalTextureParams None %129
+%start_byte_offset = OpFunctionParameter %uint
+        %130 = OpLabel
+        %131 = OpUDiv %uint %start_byte_offset %uint_16
+        %133 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %131
+        %135 = OpBitwiseAnd %uint %start_byte_offset %uint_15
+        %137 = OpShiftRightLogical %uint %135 %uint_2
+        %139 = OpLoad %v4uint %133 None
+        %140 = OpVectorExtractDynamic %uint %139 %137
+        %141 = OpIAdd %uint %uint_4 %start_byte_offset
+        %143 = OpUDiv %uint %141 %uint_16
+        %144 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %143
+        %145 = OpBitwiseAnd %uint %141 %uint_15
+        %146 = OpShiftRightLogical %uint %145 %uint_2
+        %147 = OpLoad %v4uint %144 None
+        %148 = OpVectorExtractDynamic %uint %147 %146
+        %149 = OpIAdd %uint %uint_16 %start_byte_offset
+        %150 = OpFunctionCall %mat3v4float %151 %149
+        %152 = OpIAdd %uint %uint_64 %start_byte_offset
+        %154 = OpFunctionCall %tint_GammaTransferParams %155 %152
+        %156 = OpIAdd %uint %uint_96 %start_byte_offset
+        %158 = OpFunctionCall %tint_GammaTransferParams %155 %156
+        %159 = OpIAdd %uint %uint_128 %start_byte_offset
+        %161 = OpFunctionCall %mat3v3float %162 %159
+        %163 = OpIAdd %uint %uint_176 %start_byte_offset
+        %165 = OpFunctionCall %mat3v2float %166 %163
+        %167 = OpIAdd %uint %uint_200 %start_byte_offset
+        %169 = OpFunctionCall %mat3v2float %166 %167
+        %170 = OpIAdd %uint %uint_224 %start_byte_offset
+        %172 = OpUDiv %uint %170 %uint_16
+        %173 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %172
+        %174 = OpLoad %v4uint %173 None
+        %175 = OpBitwiseAnd %uint %170 %uint_15
+        %176 = OpShiftRightLogical %uint %175 %uint_2
+        %177 = OpVectorShuffle %v2uint %174 %174 2 3
+        %178 = OpVectorShuffle %v2uint %174 %174 0 1
+        %179 = OpIEqual %bool %176 %uint_2
+        %181 = OpCompositeConstruct %v2bool %179 %179
+        %182 = OpSelect %v2uint %181 %177 %178
+        %183 = OpBitcast %v2float %182
+        %184 = OpIAdd %uint %uint_232 %start_byte_offset
+        %186 = OpUDiv %uint %184 %uint_16
+        %187 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %186
+        %188 = OpLoad %v4uint %187 None
+        %189 = OpBitwiseAnd %uint %184 %uint_15
+        %190 = OpShiftRightLogical %uint %189 %uint_2
+        %191 = OpVectorShuffle %v2uint %188 %188 2 3
+        %192 = OpVectorShuffle %v2uint %188 %188 0 1
+        %193 = OpIEqual %bool %190 %uint_2
+        %194 = OpCompositeConstruct %v2bool %193 %193
+        %195 = OpSelect %v2uint %194 %191 %192
+        %196 = OpBitcast %v2float %195
+        %197 = OpIAdd %uint %uint_240 %start_byte_offset
+        %199 = OpUDiv %uint %197 %uint_16
+        %200 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %199
+        %201 = OpLoad %v4uint %200 None
+        %202 = OpBitwiseAnd %uint %197 %uint_15
+        %203 = OpShiftRightLogical %uint %202 %uint_2
+        %204 = OpVectorShuffle %v2uint %201 %201 2 3
+        %205 = OpVectorShuffle %v2uint %201 %201 0 1
+        %206 = OpIEqual %bool %203 %uint_2
+        %207 = OpCompositeConstruct %v2bool %206 %206
+        %208 = OpSelect %v2uint %207 %204 %205
+        %209 = OpBitcast %v2float %208
+        %210 = OpIAdd %uint %uint_248 %start_byte_offset
+        %212 = OpUDiv %uint %210 %uint_16
+        %213 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %212
+        %214 = OpLoad %v4uint %213 None
+        %215 = OpBitwiseAnd %uint %210 %uint_15
+        %216 = OpShiftRightLogical %uint %215 %uint_2
+        %217 = OpVectorShuffle %v2uint %214 %214 2 3
+        %218 = OpVectorShuffle %v2uint %214 %214 0 1
+        %219 = OpIEqual %bool %216 %uint_2
+        %220 = OpCompositeConstruct %v2bool %219 %219
+        %221 = OpSelect %v2uint %220 %217 %218
+        %222 = OpBitcast %v2float %221
+        %223 = OpIAdd %uint %uint_256 %start_byte_offset
+        %225 = OpUDiv %uint %223 %uint_16
+        %226 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %225
+        %227 = OpLoad %v4uint %226 None
+        %228 = OpBitwiseAnd %uint %223 %uint_15
+        %229 = OpShiftRightLogical %uint %228 %uint_2
+        %230 = OpVectorShuffle %v2uint %227 %227 2 3
+        %231 = OpVectorShuffle %v2uint %227 %227 0 1
+        %232 = OpIEqual %bool %229 %uint_2
+        %233 = OpCompositeConstruct %v2bool %232 %232
+        %234 = OpSelect %v2uint %233 %230 %231
+        %235 = OpIAdd %uint %uint_264 %start_byte_offset
+        %237 = OpUDiv %uint %235 %uint_16
+        %238 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %237
+        %239 = OpLoad %v4uint %238 None
+        %240 = OpBitwiseAnd %uint %235 %uint_15
+        %241 = OpShiftRightLogical %uint %240 %uint_2
+        %242 = OpVectorShuffle %v2uint %239 %239 2 3
+        %243 = OpVectorShuffle %v2uint %239 %239 0 1
+        %244 = OpIEqual %bool %241 %uint_2
+        %245 = OpCompositeConstruct %v2bool %244 %244
+        %246 = OpSelect %v2uint %245 %242 %243
+        %247 = OpBitcast %v2float %246
+        %248 = OpCompositeConstruct %tint_ExternalTextureParams %140 %148 %150 %154 %158 %161 %165 %169 %183 %196 %209 %222 %234 %247
+               OpReturnValue %248
+               OpFunctionEnd
+        %151 = OpFunction %mat3v4float None %250
+%start_byte_offset_0 = OpFunctionParameter %uint
+        %251 = OpLabel
+        %252 = OpUDiv %uint %start_byte_offset_0 %uint_16
+        %253 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %252
+        %254 = OpLoad %v4uint %253 None
+        %255 = OpBitcast %v4float %254
+        %256 = OpIAdd %uint %uint_16 %start_byte_offset_0
+        %257 = OpUDiv %uint %256 %uint_16
+        %258 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %257
+        %259 = OpLoad %v4uint %258 None
+        %260 = OpBitcast %v4float %259
+        %261 = OpIAdd %uint %uint_32 %start_byte_offset_0
+        %263 = OpUDiv %uint %261 %uint_16
+        %264 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %263
+        %265 = OpLoad %v4uint %264 None
+        %266 = OpBitcast %v4float %265
+        %267 = OpCompositeConstruct %mat3v4float %255 %260 %266
+               OpReturnValue %267
+               OpFunctionEnd
+        %155 = OpFunction %tint_GammaTransferParams None %269
+%start_byte_offset_1 = OpFunctionParameter %uint
+        %270 = OpLabel
+        %271 = OpUDiv %uint %start_byte_offset_1 %uint_16
+        %272 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %271
+        %273 = OpBitwiseAnd %uint %start_byte_offset_1 %uint_15
+        %274 = OpShiftRightLogical %uint %273 %uint_2
+        %275 = OpLoad %v4uint %272 None
+        %276 = OpVectorExtractDynamic %uint %275 %274
+        %277 = OpBitcast %float %276
+        %278 = OpIAdd %uint %uint_4 %start_byte_offset_1
+        %279 = OpUDiv %uint %278 %uint_16
+        %280 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %279
+        %281 = OpBitwiseAnd %uint %278 %uint_15
+        %282 = OpShiftRightLogical %uint %281 %uint_2
+        %283 = OpLoad %v4uint %280 None
+        %284 = OpVectorExtractDynamic %uint %283 %282
+        %285 = OpBitcast %float %284
+        %286 = OpIAdd %uint %uint_8 %start_byte_offset_1
+        %288 = OpUDiv %uint %286 %uint_16
+        %289 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %288
+        %290 = OpBitwiseAnd %uint %286 %uint_15
+        %291 = OpShiftRightLogical %uint %290 %uint_2
+        %292 = OpLoad %v4uint %289 None
+        %293 = OpVectorExtractDynamic %uint %292 %291
+        %294 = OpBitcast %float %293
+        %295 = OpIAdd %uint %uint_12 %start_byte_offset_1
+        %297 = OpUDiv %uint %295 %uint_16
+        %298 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %297
+        %299 = OpBitwiseAnd %uint %295 %uint_15
+        %300 = OpShiftRightLogical %uint %299 %uint_2
+        %301 = OpLoad %v4uint %298 None
+        %302 = OpVectorExtractDynamic %uint %301 %300
+        %303 = OpBitcast %float %302
+        %304 = OpIAdd %uint %uint_16 %start_byte_offset_1
+        %305 = OpUDiv %uint %304 %uint_16
+        %306 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %305
+        %307 = OpBitwiseAnd %uint %304 %uint_15
+        %308 = OpShiftRightLogical %uint %307 %uint_2
+        %309 = OpLoad %v4uint %306 None
+        %310 = OpVectorExtractDynamic %uint %309 %308
+        %311 = OpBitcast %float %310
+        %312 = OpIAdd %uint %uint_20 %start_byte_offset_1
+        %314 = OpUDiv %uint %312 %uint_16
+        %315 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %314
+        %316 = OpBitwiseAnd %uint %312 %uint_15
+        %317 = OpShiftRightLogical %uint %316 %uint_2
+        %318 = OpLoad %v4uint %315 None
+        %319 = OpVectorExtractDynamic %uint %318 %317
+        %320 = OpBitcast %float %319
+        %321 = OpIAdd %uint %uint_24 %start_byte_offset_1
+        %323 = OpUDiv %uint %321 %uint_16
+        %324 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %323
+        %325 = OpBitwiseAnd %uint %321 %uint_15
+        %326 = OpShiftRightLogical %uint %325 %uint_2
+        %327 = OpLoad %v4uint %324 None
+        %328 = OpVectorExtractDynamic %uint %327 %326
+        %329 = OpBitcast %float %328
+        %330 = OpIAdd %uint %uint_28 %start_byte_offset_1
+        %332 = OpUDiv %uint %330 %uint_16
+        %333 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %332
+        %334 = OpBitwiseAnd %uint %330 %uint_15
+        %335 = OpShiftRightLogical %uint %334 %uint_2
+        %336 = OpLoad %v4uint %333 None
+        %337 = OpVectorExtractDynamic %uint %336 %335
+        %338 = OpCompositeConstruct %tint_GammaTransferParams %277 %285 %294 %303 %311 %320 %329 %337
+               OpReturnValue %338
+               OpFunctionEnd
+        %162 = OpFunction %mat3v3float None %340
+%start_byte_offset_2 = OpFunctionParameter %uint
+        %341 = OpLabel
+        %342 = OpUDiv %uint %start_byte_offset_2 %uint_16
+        %343 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %342
+        %344 = OpLoad %v4uint %343 None
+        %345 = OpVectorShuffle %v3uint %344 %344 0 1 2
+        %347 = OpBitcast %v3float %345
+        %348 = OpIAdd %uint %uint_16 %start_byte_offset_2
+        %349 = OpUDiv %uint %348 %uint_16
+        %350 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %349
+        %351 = OpLoad %v4uint %350 None
+        %352 = OpVectorShuffle %v3uint %351 %351 0 1 2
+        %353 = OpBitcast %v3float %352
+        %354 = OpIAdd %uint %uint_32 %start_byte_offset_2
+        %355 = OpUDiv %uint %354 %uint_16
+        %356 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %355
+        %357 = OpLoad %v4uint %356 None
+        %358 = OpVectorShuffle %v3uint %357 %357 0 1 2
+        %359 = OpBitcast %v3float %358
+        %360 = OpCompositeConstruct %mat3v3float %347 %353 %359
+               OpReturnValue %360
+               OpFunctionEnd
+        %166 = OpFunction %mat3v2float None %362
+%start_byte_offset_3 = OpFunctionParameter %uint
+        %363 = OpLabel
+        %364 = OpUDiv %uint %start_byte_offset_3 %uint_16
+        %365 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %364
+        %366 = OpLoad %v4uint %365 None
+        %367 = OpBitwiseAnd %uint %start_byte_offset_3 %uint_15
+        %368 = OpShiftRightLogical %uint %367 %uint_2
+        %369 = OpVectorShuffle %v2uint %366 %366 2 3
+        %370 = OpVectorShuffle %v2uint %366 %366 0 1
+        %371 = OpIEqual %bool %368 %uint_2
+        %372 = OpCompositeConstruct %v2bool %371 %371
+        %373 = OpSelect %v2uint %372 %369 %370
+        %374 = OpBitcast %v2float %373
+        %375 = OpIAdd %uint %uint_8 %start_byte_offset_3
+        %376 = OpUDiv %uint %375 %uint_16
+        %377 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %376
+        %378 = OpLoad %v4uint %377 None
+        %379 = OpBitwiseAnd %uint %375 %uint_15
+        %380 = OpShiftRightLogical %uint %379 %uint_2
+        %381 = OpVectorShuffle %v2uint %378 %378 2 3
+        %382 = OpVectorShuffle %v2uint %378 %378 0 1
+        %383 = OpIEqual %bool %380 %uint_2
+        %384 = OpCompositeConstruct %v2bool %383 %383
+        %385 = OpSelect %v2uint %384 %381 %382
+        %386 = OpBitcast %v2float %385
+        %387 = OpIAdd %uint %uint_16 %start_byte_offset_3
+        %388 = OpUDiv %uint %387 %uint_16
+        %389 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %388
+        %390 = OpLoad %v4uint %389 None
+        %391 = OpBitwiseAnd %uint %387 %uint_15
+        %392 = OpShiftRightLogical %uint %391 %uint_2
+        %393 = OpVectorShuffle %v2uint %390 %390 2 3
+        %394 = OpVectorShuffle %v2uint %390 %390 0 1
+        %395 = OpIEqual %bool %392 %uint_2
+        %396 = OpCompositeConstruct %v2bool %395 %395
+        %397 = OpSelect %v2uint %396 %393 %394
+        %398 = OpBitcast %v2float %397
+        %399 = OpCompositeConstruct %mat3v2float %374 %386 %398
+               OpReturnValue %399
+               OpFunctionEnd

diff --git a/test/tint/types/texture/external/external.wgsl.expected.wgsl b/test/tint/types/texture/external/external.wgsl.expected.wgsl
new file mode 100644
index 0000000..f543079
--- /dev/null
+++ b/test/tint/types/texture/external/external.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+@group(0) @binding(0) var t_f : texture_external;
+
+@group(1) @binding(0) var<storage, read_write> out : vec4f;
+
+@compute @workgroup_size(1)
+fn main() {
+  out = textureLoad(t_f, vec2i(0));
+}

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl b/test/tint/types/texture/external/external_ycbcr.wgsl
new file mode 100644
index 0000000..5c3d4cb
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl

@@ -0,0 +1,8 @@
+// flags: --ycbcr-bindings=0,0
+
+@group(0) @binding(0) var t_f : texture_external;
+
+@compute @workgroup_size(1)
+fn main() {
+  var vals = textureLoad(t_f, vec2i(0));
+}

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl.expected.dxc.hlsl b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b93fc77
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.dxc.hlsl

@@ -0,0 +1,126 @@
+struct tint_GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+  float3x2 sampleTransform;
+  float3x2 loadTransform;
+  float2 samplePlane0RectMin;
+  float2 samplePlane0RectMax;
+  float2 samplePlane1RectMin;
+  float2 samplePlane1RectMax;
+  uint2 apparentSize;
+  float2 plane1CoordFactor;
+};
+
+
+cbuffer cbuffer_t_f_params : register(b2) {
+  uint4 t_f_params[17];
+};
+Texture2D<float4> t_f_plane0 : register(t0);
+Texture2D<float4> t_f_plane1 : register(t1);
+uint2 tint_v2f32_to_v2u32(float2 value) {
+  return uint2(clamp(value, (0.0f).xx, (4294967040.0f).xx));
+}
+
+float3 tint_GammaCorrection(float3 v, tint_GammaTransferParams params) {
+  float3 v_1 = float3((params.G).xxx);
+  float3 v_2 = float3((params.D).xxx);
+  float3 v_3 = float3(sign(v));
+  return (((abs(v) < v_2)) ? ((v_3 * ((params.C * abs(v)) + params.F))) : ((v_3 * (pow(((params.A * abs(v)) + params.B), v_1) + params.E))));
+}
+
+float4 tint_TextureLoadMultiplanarExternal(Texture2D<float4> plane_0, Texture2D<float4> plane_1, tint_ExternalTextureParams params, uint2 coords) {
+  float2 v_4 = round(mul(float3(float2(min(coords, params.apparentSize)), 1.0f), params.loadTransform));
+  uint2 v_5 = tint_v2f32_to_v2u32(v_4);
+  float3 v_6 = (0.0f).xxx;
+  float v_7 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    int2 v_8 = int2(v_5);
+    float4 v_9 = plane_0.Load(int3(v_8, int(0u)));
+    v_6 = v_9.xyz;
+    v_7 = v_9.w;
+  } else {
+    int2 v_10 = int2(v_5);
+    float v_11 = plane_0.Load(int3(v_10, int(0u))).x;
+    int2 v_12 = int2(tint_v2f32_to_v2u32((v_4 * params.plane1CoordFactor)));
+    v_6 = mul(params.yuvToRgbConversionMatrix, float4(v_11, plane_1.Load(int3(v_12, int(0u))).xy, 1.0f));
+    v_7 = 1.0f;
+  }
+  float3 v_13 = v_6;
+  float3 v_14 = (0.0f).xxx;
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    tint_GammaTransferParams v_15 = params.gammaDecodeParams;
+    tint_GammaTransferParams v_16 = params.gammaEncodeParams;
+    v_14 = tint_GammaCorrection(mul(tint_GammaCorrection(v_13, v_15), params.gamutConversionMatrix), v_16);
+  } else {
+    v_14 = v_13;
+  }
+  return float4(v_14, v_7);
+}
+
+float3x2 v_17(uint start_byte_offset) {
+  uint4 v_18 = t_f_params[(start_byte_offset / 16u)];
+  float2 v_19 = asfloat((((((start_byte_offset & 15u) >> 2u) == 2u)) ? (v_18.zw) : (v_18.xy)));
+  uint4 v_20 = t_f_params[((8u + start_byte_offset) / 16u)];
+  float2 v_21 = asfloat(((((((8u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_20.zw) : (v_20.xy)));
+  uint4 v_22 = t_f_params[((16u + start_byte_offset) / 16u)];
+  return float3x2(v_19, v_21, asfloat(((((((16u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_22.zw) : (v_22.xy))));
+}
+
+float3x3 v_23(uint start_byte_offset) {
+  return float3x3(asfloat(t_f_params[(start_byte_offset / 16u)].xyz), asfloat(t_f_params[((16u + start_byte_offset) / 16u)].xyz), asfloat(t_f_params[((32u + start_byte_offset) / 16u)].xyz));
+}
+
+tint_GammaTransferParams v_24(uint start_byte_offset) {
+  tint_GammaTransferParams v_25 = {asfloat(t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)]), asfloat(t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) & 15u) >> 2u)]), t_f_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) & 15u) >> 2u)]};
+  return v_25;
+}
+
+float3x4 v_26(uint start_byte_offset) {
+  return float3x4(asfloat(t_f_params[(start_byte_offset / 16u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)]), asfloat(t_f_params[((32u + start_byte_offset) / 16u)]));
+}
+
+tint_ExternalTextureParams v_27(uint start_byte_offset) {
+  uint v_28 = t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)];
+  uint v_29 = t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)];
+  float3x4 v_30 = v_26((16u + start_byte_offset));
+  tint_GammaTransferParams v_31 = v_24((64u + start_byte_offset));
+  tint_GammaTransferParams v_32 = v_24((96u + start_byte_offset));
+  float3x3 v_33 = v_23((128u + start_byte_offset));
+  float3x2 v_34 = v_17((176u + start_byte_offset));
+  float3x2 v_35 = v_17((200u + start_byte_offset));
+  uint4 v_36 = t_f_params[((224u + start_byte_offset) / 16u)];
+  float2 v_37 = asfloat(((((((224u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_36.zw) : (v_36.xy)));
+  uint4 v_38 = t_f_params[((232u + start_byte_offset) / 16u)];
+  float2 v_39 = asfloat(((((((232u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_38.zw) : (v_38.xy)));
+  uint4 v_40 = t_f_params[((240u + start_byte_offset) / 16u)];
+  float2 v_41 = asfloat(((((((240u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_40.zw) : (v_40.xy)));
+  uint4 v_42 = t_f_params[((248u + start_byte_offset) / 16u)];
+  float2 v_43 = asfloat(((((((248u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_42.zw) : (v_42.xy)));
+  uint4 v_44 = t_f_params[((256u + start_byte_offset) / 16u)];
+  uint2 v_45 = ((((((256u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_44.zw) : (v_44.xy));
+  uint4 v_46 = t_f_params[((264u + start_byte_offset) / 16u)];
+  tint_ExternalTextureParams v_47 = {v_28, v_29, v_30, v_31, v_32, v_33, v_34, v_35, v_37, v_39, v_41, v_43, v_45, asfloat(((((((264u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_46.zw) : (v_46.xy)))};
+  return v_47;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_ExternalTextureParams v_48 = v_27(0u);
+  float4 vals = tint_TextureLoadMultiplanarExternal(t_f_plane0, t_f_plane1, v_48, uint2((int(0)).xx));
+}
+

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl.expected.fxc.hlsl b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b93fc77
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.fxc.hlsl

@@ -0,0 +1,126 @@
+struct tint_GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+  float3x2 sampleTransform;
+  float3x2 loadTransform;
+  float2 samplePlane0RectMin;
+  float2 samplePlane0RectMax;
+  float2 samplePlane1RectMin;
+  float2 samplePlane1RectMax;
+  uint2 apparentSize;
+  float2 plane1CoordFactor;
+};
+
+
+cbuffer cbuffer_t_f_params : register(b2) {
+  uint4 t_f_params[17];
+};
+Texture2D<float4> t_f_plane0 : register(t0);
+Texture2D<float4> t_f_plane1 : register(t1);
+uint2 tint_v2f32_to_v2u32(float2 value) {
+  return uint2(clamp(value, (0.0f).xx, (4294967040.0f).xx));
+}
+
+float3 tint_GammaCorrection(float3 v, tint_GammaTransferParams params) {
+  float3 v_1 = float3((params.G).xxx);
+  float3 v_2 = float3((params.D).xxx);
+  float3 v_3 = float3(sign(v));
+  return (((abs(v) < v_2)) ? ((v_3 * ((params.C * abs(v)) + params.F))) : ((v_3 * (pow(((params.A * abs(v)) + params.B), v_1) + params.E))));
+}
+
+float4 tint_TextureLoadMultiplanarExternal(Texture2D<float4> plane_0, Texture2D<float4> plane_1, tint_ExternalTextureParams params, uint2 coords) {
+  float2 v_4 = round(mul(float3(float2(min(coords, params.apparentSize)), 1.0f), params.loadTransform));
+  uint2 v_5 = tint_v2f32_to_v2u32(v_4);
+  float3 v_6 = (0.0f).xxx;
+  float v_7 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    int2 v_8 = int2(v_5);
+    float4 v_9 = plane_0.Load(int3(v_8, int(0u)));
+    v_6 = v_9.xyz;
+    v_7 = v_9.w;
+  } else {
+    int2 v_10 = int2(v_5);
+    float v_11 = plane_0.Load(int3(v_10, int(0u))).x;
+    int2 v_12 = int2(tint_v2f32_to_v2u32((v_4 * params.plane1CoordFactor)));
+    v_6 = mul(params.yuvToRgbConversionMatrix, float4(v_11, plane_1.Load(int3(v_12, int(0u))).xy, 1.0f));
+    v_7 = 1.0f;
+  }
+  float3 v_13 = v_6;
+  float3 v_14 = (0.0f).xxx;
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    tint_GammaTransferParams v_15 = params.gammaDecodeParams;
+    tint_GammaTransferParams v_16 = params.gammaEncodeParams;
+    v_14 = tint_GammaCorrection(mul(tint_GammaCorrection(v_13, v_15), params.gamutConversionMatrix), v_16);
+  } else {
+    v_14 = v_13;
+  }
+  return float4(v_14, v_7);
+}
+
+float3x2 v_17(uint start_byte_offset) {
+  uint4 v_18 = t_f_params[(start_byte_offset / 16u)];
+  float2 v_19 = asfloat((((((start_byte_offset & 15u) >> 2u) == 2u)) ? (v_18.zw) : (v_18.xy)));
+  uint4 v_20 = t_f_params[((8u + start_byte_offset) / 16u)];
+  float2 v_21 = asfloat(((((((8u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_20.zw) : (v_20.xy)));
+  uint4 v_22 = t_f_params[((16u + start_byte_offset) / 16u)];
+  return float3x2(v_19, v_21, asfloat(((((((16u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_22.zw) : (v_22.xy))));
+}
+
+float3x3 v_23(uint start_byte_offset) {
+  return float3x3(asfloat(t_f_params[(start_byte_offset / 16u)].xyz), asfloat(t_f_params[((16u + start_byte_offset) / 16u)].xyz), asfloat(t_f_params[((32u + start_byte_offset) / 16u)].xyz));
+}
+
+tint_GammaTransferParams v_24(uint start_byte_offset) {
+  tint_GammaTransferParams v_25 = {asfloat(t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)]), asfloat(t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) & 15u) >> 2u)]), asfloat(t_f_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) & 15u) >> 2u)]), t_f_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) & 15u) >> 2u)]};
+  return v_25;
+}
+
+float3x4 v_26(uint start_byte_offset) {
+  return float3x4(asfloat(t_f_params[(start_byte_offset / 16u)]), asfloat(t_f_params[((16u + start_byte_offset) / 16u)]), asfloat(t_f_params[((32u + start_byte_offset) / 16u)]));
+}
+
+tint_ExternalTextureParams v_27(uint start_byte_offset) {
+  uint v_28 = t_f_params[(start_byte_offset / 16u)][((start_byte_offset & 15u) >> 2u)];
+  uint v_29 = t_f_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) & 15u) >> 2u)];
+  float3x4 v_30 = v_26((16u + start_byte_offset));
+  tint_GammaTransferParams v_31 = v_24((64u + start_byte_offset));
+  tint_GammaTransferParams v_32 = v_24((96u + start_byte_offset));
+  float3x3 v_33 = v_23((128u + start_byte_offset));
+  float3x2 v_34 = v_17((176u + start_byte_offset));
+  float3x2 v_35 = v_17((200u + start_byte_offset));
+  uint4 v_36 = t_f_params[((224u + start_byte_offset) / 16u)];
+  float2 v_37 = asfloat(((((((224u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_36.zw) : (v_36.xy)));
+  uint4 v_38 = t_f_params[((232u + start_byte_offset) / 16u)];
+  float2 v_39 = asfloat(((((((232u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_38.zw) : (v_38.xy)));
+  uint4 v_40 = t_f_params[((240u + start_byte_offset) / 16u)];
+  float2 v_41 = asfloat(((((((240u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_40.zw) : (v_40.xy)));
+  uint4 v_42 = t_f_params[((248u + start_byte_offset) / 16u)];
+  float2 v_43 = asfloat(((((((248u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_42.zw) : (v_42.xy)));
+  uint4 v_44 = t_f_params[((256u + start_byte_offset) / 16u)];
+  uint2 v_45 = ((((((256u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_44.zw) : (v_44.xy));
+  uint4 v_46 = t_f_params[((264u + start_byte_offset) / 16u)];
+  tint_ExternalTextureParams v_47 = {v_28, v_29, v_30, v_31, v_32, v_33, v_34, v_35, v_37, v_39, v_41, v_43, v_45, asfloat(((((((264u + start_byte_offset) & 15u) >> 2u) == 2u)) ? (v_46.zw) : (v_46.xy)))};
+  return v_47;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_ExternalTextureParams v_48 = v_27(0u);
+  float4 vals = tint_TextureLoadMultiplanarExternal(t_f_plane0, t_f_plane1, v_48, uint2((int(0)).xx));
+}
+

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl.expected.glsl b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.glsl
new file mode 100644
index 0000000..ba1f1c6
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.glsl

@@ -0,0 +1,119 @@
+#version 310 es
+
+
+struct tint_GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+  mat3x2 sampleTransform;
+  mat3x2 loadTransform;
+  vec2 samplePlane0RectMin;
+  vec2 samplePlane0RectMax;
+  vec2 samplePlane1RectMin;
+  vec2 samplePlane1RectMax;
+  uvec2 apparentSize;
+  vec2 plane1CoordFactor;
+};
+
+layout(binding = 2, std140)
+uniform t_f_params_block_1_ubo {
+  uvec4 inner[17];
+} v_1;
+uniform highp sampler2D t_f_plane0;
+uniform highp sampler2D t_f_plane1;
+vec3 tint_GammaCorrection(vec3 v, tint_GammaTransferParams params) {
+  vec3 v_2 = vec3(params.G);
+  return mix((sign(v) * (pow(((params.A * abs(v)) + params.B), v_2) + params.E)), (sign(v) * ((params.C * abs(v)) + params.F)), lessThan(abs(v), vec3(params.D)));
+}
+vec4 tint_TextureLoadMultiplanarExternal(tint_ExternalTextureParams params, uvec2 coords) {
+  vec2 v_3 = round((params.loadTransform * vec3(vec2(min(coords, params.apparentSize)), 1.0f)));
+  uvec2 v_4 = uvec2(v_3);
+  vec3 v_5 = vec3(0.0f);
+  float v_6 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    ivec2 v_7 = ivec2(v_4);
+    vec4 v_8 = texelFetch(t_f_plane0, v_7, int(0u));
+    v_5 = v_8.xyz;
+    v_6 = v_8.w;
+  } else {
+    ivec2 v_9 = ivec2(v_4);
+    float v_10 = texelFetch(t_f_plane0, v_9, int(0u)).x;
+    ivec2 v_11 = ivec2(uvec2((v_3 * params.plane1CoordFactor)));
+    v_5 = (vec4(v_10, texelFetch(t_f_plane1, v_11, int(0u)).xy, 1.0f) * params.yuvToRgbConversionMatrix);
+    v_6 = 1.0f;
+  }
+  vec3 v_12 = v_5;
+  vec3 v_13 = vec3(0.0f);
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    v_13 = tint_GammaCorrection((params.gamutConversionMatrix * tint_GammaCorrection(v_12, params.gammaDecodeParams)), params.gammaEncodeParams);
+  } else {
+    v_13 = v_12;
+  }
+  return vec4(v_13, v_6);
+}
+mat3x2 v_14(uint start_byte_offset) {
+  uvec4 v_15 = v_1.inner[(start_byte_offset / 16u)];
+  vec2 v_16 = uintBitsToFloat(mix(v_15.xy, v_15.zw, bvec2((((start_byte_offset & 15u) >> 2u) == 2u))));
+  uvec4 v_17 = v_1.inner[((8u + start_byte_offset) / 16u)];
+  vec2 v_18 = uintBitsToFloat(mix(v_17.xy, v_17.zw, bvec2(((((8u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_19 = v_1.inner[((16u + start_byte_offset) / 16u)];
+  return mat3x2(v_16, v_18, uintBitsToFloat(mix(v_19.xy, v_19.zw, bvec2(((((16u + start_byte_offset) & 15u) >> 2u) == 2u)))));
+}
+mat3 v_20(uint start_byte_offset) {
+  return mat3(uintBitsToFloat(v_1.inner[(start_byte_offset / 16u)].xyz), uintBitsToFloat(v_1.inner[((16u + start_byte_offset) / 16u)].xyz), uintBitsToFloat(v_1.inner[((32u + start_byte_offset) / 16u)].xyz));
+}
+tint_GammaTransferParams v_21(uint start_byte_offset) {
+  uvec4 v_22 = v_1.inner[(start_byte_offset / 16u)];
+  uvec4 v_23 = v_1.inner[((4u + start_byte_offset) / 16u)];
+  uvec4 v_24 = v_1.inner[((8u + start_byte_offset) / 16u)];
+  uvec4 v_25 = v_1.inner[((12u + start_byte_offset) / 16u)];
+  uvec4 v_26 = v_1.inner[((16u + start_byte_offset) / 16u)];
+  uvec4 v_27 = v_1.inner[((20u + start_byte_offset) / 16u)];
+  uvec4 v_28 = v_1.inner[((24u + start_byte_offset) / 16u)];
+  uvec4 v_29 = v_1.inner[((28u + start_byte_offset) / 16u)];
+  return tint_GammaTransferParams(uintBitsToFloat(v_22[((start_byte_offset & 15u) >> 2u)]), uintBitsToFloat(v_23[(((4u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_24[(((8u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_25[(((12u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_26[(((16u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_27[(((20u + start_byte_offset) & 15u) >> 2u)]), uintBitsToFloat(v_28[(((24u + start_byte_offset) & 15u) >> 2u)]), v_29[(((28u + start_byte_offset) & 15u) >> 2u)]);
+}
+mat3x4 v_30(uint start_byte_offset) {
+  return mat3x4(uintBitsToFloat(v_1.inner[(start_byte_offset / 16u)]), uintBitsToFloat(v_1.inner[((16u + start_byte_offset) / 16u)]), uintBitsToFloat(v_1.inner[((32u + start_byte_offset) / 16u)]));
+}
+tint_ExternalTextureParams v_31(uint start_byte_offset) {
+  uvec4 v_32 = v_1.inner[(start_byte_offset / 16u)];
+  uvec4 v_33 = v_1.inner[((4u + start_byte_offset) / 16u)];
+  mat3x4 v_34 = v_30((16u + start_byte_offset));
+  tint_GammaTransferParams v_35 = v_21((64u + start_byte_offset));
+  tint_GammaTransferParams v_36 = v_21((96u + start_byte_offset));
+  mat3 v_37 = v_20((128u + start_byte_offset));
+  mat3x2 v_38 = v_14((176u + start_byte_offset));
+  mat3x2 v_39 = v_14((200u + start_byte_offset));
+  uvec4 v_40 = v_1.inner[((224u + start_byte_offset) / 16u)];
+  vec2 v_41 = uintBitsToFloat(mix(v_40.xy, v_40.zw, bvec2(((((224u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_42 = v_1.inner[((232u + start_byte_offset) / 16u)];
+  vec2 v_43 = uintBitsToFloat(mix(v_42.xy, v_42.zw, bvec2(((((232u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_44 = v_1.inner[((240u + start_byte_offset) / 16u)];
+  vec2 v_45 = uintBitsToFloat(mix(v_44.xy, v_44.zw, bvec2(((((240u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_46 = v_1.inner[((248u + start_byte_offset) / 16u)];
+  vec2 v_47 = uintBitsToFloat(mix(v_46.xy, v_46.zw, bvec2(((((248u + start_byte_offset) & 15u) >> 2u) == 2u))));
+  uvec4 v_48 = v_1.inner[((256u + start_byte_offset) / 16u)];
+  uvec2 v_49 = mix(v_48.xy, v_48.zw, bvec2(((((256u + start_byte_offset) & 15u) >> 2u) == 2u)));
+  uvec4 v_50 = v_1.inner[((264u + start_byte_offset) / 16u)];
+  return tint_ExternalTextureParams(v_32[((start_byte_offset & 15u) >> 2u)], v_33[(((4u + start_byte_offset) & 15u) >> 2u)], v_34, v_35, v_36, v_37, v_38, v_39, v_41, v_43, v_45, v_47, v_49, uintBitsToFloat(mix(v_50.xy, v_50.zw, bvec2(((((264u + start_byte_offset) & 15u) >> 2u) == 2u)))));
+}
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_ExternalTextureParams v_51 = v_31(0u);
+  vec4 vals = tint_TextureLoadMultiplanarExternal(v_51, min(uvec2(ivec2(0)), ((v_51.apparentSize + uvec2(1u)) - uvec2(1u))));
+}

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl.expected.msl b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.msl
new file mode 100644
index 0000000..5dcc4f6
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.msl

@@ -0,0 +1,118 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_GammaTransferParams {
+  /* 0x0000 */ float G;
+  /* 0x0004 */ float A;
+  /* 0x0008 */ float B;
+  /* 0x000c */ float C;
+  /* 0x0010 */ float D;
+  /* 0x0014 */ float E;
+  /* 0x0018 */ float F;
+  /* 0x001c */ uint padding;
+};
+
+struct tint_ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  tint_GammaTransferParams gammaDecodeParams;
+  tint_GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+  float3x2 sampleTransform;
+  float3x2 loadTransform;
+  float2 samplePlane0RectMin;
+  float2 samplePlane0RectMax;
+  float2 samplePlane1RectMin;
+  float2 samplePlane1RectMax;
+  uint2 apparentSize;
+  float2 plane1CoordFactor;
+};
+
+template<typename T, size_t N>
+struct tint_array {
+  const constant T& operator[](size_t i) const constant { return elements[i]; }
+  device T& operator[](size_t i) device { return elements[i]; }
+  const device T& operator[](size_t i) const device { return elements[i]; }
+  thread T& operator[](size_t i) thread { return elements[i]; }
+  const thread T& operator[](size_t i) const thread { return elements[i]; }
+  threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+  const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+  T elements[N];
+};
+
+struct tint_packed_vec3_f32_array_element {
+  /* 0x0000 */ packed_float3 packed;
+  /* 0x000c */ tint_array<int8_t, 4> tint_pad_1;
+};
+
+struct tint_ExternalTextureParams_packed_vec3 {
+  /* 0x0000 */ uint numPlanes;
+  /* 0x0004 */ uint doYuvToRgbConversionOnly;
+  /* 0x0008 */ tint_array<int8_t, 8> tint_pad;
+  /* 0x0010 */ float3x4 yuvToRgbConversionMatrix;
+  /* 0x0040 */ tint_GammaTransferParams gammaDecodeParams;
+  /* 0x0060 */ tint_GammaTransferParams gammaEncodeParams;
+  /* 0x0080 */ tint_array<tint_packed_vec3_f32_array_element, 3> gamutConversionMatrix;
+  /* 0x00b0 */ float3x2 sampleTransform;
+  /* 0x00c8 */ float3x2 loadTransform;
+  /* 0x00e0 */ float2 samplePlane0RectMin;
+  /* 0x00e8 */ float2 samplePlane0RectMax;
+  /* 0x00f0 */ float2 samplePlane1RectMin;
+  /* 0x00f8 */ float2 samplePlane1RectMax;
+  /* 0x0100 */ uint2 apparentSize;
+  /* 0x0108 */ float2 plane1CoordFactor;
+};
+
+struct tint_module_vars_struct {
+  const constant tint_ExternalTextureParams_packed_vec3* t_f_params;
+  texture2d<float, access::sample> t_f_plane0;
+  texture2d<float, access::sample> t_f_plane1;
+};
+
+float3 tint_GammaCorrection(float3 v, tint_GammaTransferParams params) {
+  float3 const v_1 = float3(params.G);
+  return select((sign(v) * (powr(((params.A * abs(v)) + params.B), v_1) + params.E)), (sign(v) * ((params.C * abs(v)) + params.F)), (abs(v) < float3(params.D)));
+}
+
+float4 tint_TextureLoadMultiplanarExternal(texture2d<float, access::sample> plane_0, texture2d<float, access::sample> plane_1, tint_ExternalTextureParams params, uint2 coords) {
+  float2 const v_2 = rint((params.loadTransform * float3(float2(min(coords, params.apparentSize)), 1.0f)));
+  uint2 const v_3 = uint2(v_2);
+  float3 v_4 = 0.0f;
+  float v_5 = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    float4 const v_6 = plane_0.read(v_3, 0u);
+    v_4 = v_6.xyz;
+    v_5 = v_6.w;
+  } else {
+    float const v_7 = plane_0.read(v_3, 0u).x;
+    v_4 = (float4(v_7, plane_1.read(uint2((v_2 * params.plane1CoordFactor)), 0u).xy, 1.0f) * params.yuvToRgbConversionMatrix);
+    v_5 = 1.0f;
+  }
+  float3 const v_8 = v_4;
+  float3 v_9 = 0.0f;
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    v_9 = tint_GammaCorrection((params.gamutConversionMatrix * tint_GammaCorrection(v_8, params.gammaDecodeParams)), params.gammaEncodeParams);
+  } else {
+    v_9 = v_8;
+  }
+  return float4(v_9, v_5);
+}
+
+tint_ExternalTextureParams tint_load_struct_packed_vec3(const constant tint_ExternalTextureParams_packed_vec3* const from) {
+  uint const v_10 = (*from).numPlanes;
+  uint const v_11 = (*from).doYuvToRgbConversionOnly;
+  float3x4 const v_12 = (*from).yuvToRgbConversionMatrix;
+  tint_GammaTransferParams const v_13 = (*from).gammaDecodeParams;
+  tint_GammaTransferParams const v_14 = (*from).gammaEncodeParams;
+  tint_array<tint_packed_vec3_f32_array_element, 3> const v_15 = (*from).gamutConversionMatrix;
+  float3x3 const v_16 = float3x3(float3(v_15[0u].packed), float3(v_15[1u].packed), float3(v_15[2u].packed));
+  return tint_ExternalTextureParams{.numPlanes=v_10, .doYuvToRgbConversionOnly=v_11, .yuvToRgbConversionMatrix=v_12, .gammaDecodeParams=v_13, .gammaEncodeParams=v_14, .gamutConversionMatrix=v_16, .sampleTransform=(*from).sampleTransform, .loadTransform=(*from).loadTransform, .samplePlane0RectMin=(*from).samplePlane0RectMin, .samplePlane0RectMax=(*from).samplePlane0RectMax, .samplePlane1RectMin=(*from).samplePlane1RectMin, .samplePlane1RectMax=(*from).samplePlane1RectMax, .apparentSize=(*from).apparentSize, .plane1CoordFactor=(*from).plane1CoordFactor};
+}
+
+[[max_total_threads_per_threadgroup(1)]]
+kernel void v_17(const constant tint_ExternalTextureParams_packed_vec3* t_f_params [[buffer(0)]], texture2d<float, access::sample> t_f_plane0 [[texture(0)]], texture2d<float, access::sample> t_f_plane1 [[texture(1)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t_f_params=t_f_params, .t_f_plane0=t_f_plane0, .t_f_plane1=t_f_plane1};
+  tint_ExternalTextureParams const v_18 = tint_load_struct_packed_vec3(tint_module_vars.t_f_params);
+  float4 vals = tint_TextureLoadMultiplanarExternal(tint_module_vars.t_f_plane0, tint_module_vars.t_f_plane1, v_18, min(uint2(int2(0)), ((v_18.apparentSize + uint2(1u)) - uint2(1u))));
+}

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl.expected.spvasm b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.spvasm
new file mode 100644
index 0000000..a2c423e
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.spvasm

@@ -0,0 +1,464 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 383
+; Schema: 0
+               OpCapability Shader
+         %43 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpMemberName %t_f_params_block_tint_explicit_layout 0 "inner"
+               OpName %t_f_params_block_tint_explicit_layout "t_f_params_block_tint_explicit_layout"
+               OpName %t_f "t_f"
+               OpName %t_f_ycbcr_sampler "t_f_ycbcr_sampler"
+               OpName %main "main"
+               OpMemberName %tint_ExternalTextureParams 0 "numPlanes"
+               OpMemberName %tint_ExternalTextureParams 1 "doYuvToRgbConversionOnly"
+               OpMemberName %tint_ExternalTextureParams 2 "yuvToRgbConversionMatrix"
+               OpMemberName %tint_GammaTransferParams 0 "G"
+               OpMemberName %tint_GammaTransferParams 1 "A"
+               OpMemberName %tint_GammaTransferParams 2 "B"
+               OpMemberName %tint_GammaTransferParams 3 "C"
+               OpMemberName %tint_GammaTransferParams 4 "D"
+               OpMemberName %tint_GammaTransferParams 5 "E"
+               OpMemberName %tint_GammaTransferParams 6 "F"
+               OpMemberName %tint_GammaTransferParams 7 "padding"
+               OpName %tint_GammaTransferParams "tint_GammaTransferParams"
+               OpMemberName %tint_ExternalTextureParams 3 "gammaDecodeParams"
+               OpMemberName %tint_ExternalTextureParams 4 "gammaEncodeParams"
+               OpMemberName %tint_ExternalTextureParams 5 "gamutConversionMatrix"
+               OpMemberName %tint_ExternalTextureParams 6 "sampleTransform"
+               OpMemberName %tint_ExternalTextureParams 7 "loadTransform"
+               OpMemberName %tint_ExternalTextureParams 8 "samplePlane0RectMin"
+               OpMemberName %tint_ExternalTextureParams 9 "samplePlane0RectMax"
+               OpMemberName %tint_ExternalTextureParams 10 "samplePlane1RectMin"
+               OpMemberName %tint_ExternalTextureParams 11 "samplePlane1RectMax"
+               OpMemberName %tint_ExternalTextureParams 12 "apparentSize"
+               OpMemberName %tint_ExternalTextureParams 13 "plane1CoordFactor"
+               OpName %tint_ExternalTextureParams "tint_ExternalTextureParams"
+               OpName %vals "vals"
+               OpName %tint_TextureLoadYcbcrExternal "tint_TextureLoadYcbcrExternal"
+               OpName %texture "texture"
+               OpName %params "params"
+               OpName %coords "coords"
+               OpName %tint_GammaCorrection "tint_GammaCorrection"
+               OpName %v "v"
+               OpName %params_0 "params"
+               OpName %start_byte_offset "start_byte_offset"
+               OpName %start_byte_offset_0 "start_byte_offset"
+               OpName %start_byte_offset_1 "start_byte_offset"
+               OpName %start_byte_offset_2 "start_byte_offset"
+               OpName %start_byte_offset_3 "start_byte_offset"
+               OpDecorate %_arr_v4uint_uint_17 ArrayStride 16
+               OpMemberDecorate %t_f_params_block_tint_explicit_layout 0 Offset 0
+               OpDecorate %t_f_params_block_tint_explicit_layout Block
+               OpDecorate %1 DescriptorSet 0
+               OpDecorate %1 Binding 2
+               OpDecorate %1 NonWritable
+               OpDecorate %t_f DescriptorSet 0
+               OpDecorate %t_f Binding 0
+               OpDecorate %t_f_ycbcr_sampler DescriptorSet 0
+               OpDecorate %t_f_ycbcr_sampler Binding 1
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+    %uint_17 = OpConstant %uint 17
+%_arr_v4uint_uint_17 = OpTypeArray %v4uint %uint_17
+%t_f_params_block_tint_explicit_layout = OpTypeStruct %_arr_v4uint_uint_17
+%_ptr_Uniform_t_f_params_block_tint_explicit_layout = OpTypePointer Uniform %t_f_params_block_tint_explicit_layout
+          %1 = OpVariable %_ptr_Uniform_t_f_params_block_tint_explicit_layout Uniform
+      %float = OpTypeFloat 32
+         %10 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
+        %t_f = OpVariable %_ptr_UniformConstant_10 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+%t_f_ycbcr_sampler = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %void = OpTypeVoid
+         %17 = OpTypeFunction %void
+    %v4float = OpTypeVector %float 4
+%mat3v4float = OpTypeMatrix %v4float 3
+%tint_GammaTransferParams = OpTypeStruct %float %float %float %float %float %float %float %uint
+    %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+     %v2uint = OpTypeVector %uint 2
+%tint_ExternalTextureParams = OpTypeStruct %uint %uint %mat3v4float %tint_GammaTransferParams %tint_GammaTransferParams %mat3v3float %mat3v2float %mat3v2float %v2float %v2float %v2float %v2float %v2uint %v2float
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+         %35 = OpConstantComposite %v2uint %uint_1 %uint_1
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %39 = OpConstantNull %v2int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %51 = OpTypeFunction %v4float %10 %tint_ExternalTextureParams %v2uint
+    %float_1 = OpConstant %float 1
+       %bool = OpTypeBool
+         %84 = OpTypeFunction %v3float %v3float %tint_GammaTransferParams
+     %v3bool = OpTypeVector %bool 3
+        %112 = OpTypeFunction %tint_ExternalTextureParams %uint
+    %uint_16 = OpConstant %uint 16
+%_ptr_Uniform_v4uint = OpTypePointer Uniform %v4uint
+    %uint_15 = OpConstant %uint 15
+     %uint_2 = OpConstant %uint 2
+     %uint_4 = OpConstant %uint 4
+    %uint_64 = OpConstant %uint 64
+    %uint_96 = OpConstant %uint 96
+   %uint_128 = OpConstant %uint 128
+   %uint_176 = OpConstant %uint 176
+   %uint_200 = OpConstant %uint 200
+   %uint_224 = OpConstant %uint 224
+     %v2bool = OpTypeVector %bool 2
+   %uint_232 = OpConstant %uint 232
+   %uint_240 = OpConstant %uint 240
+   %uint_248 = OpConstant %uint 248
+   %uint_256 = OpConstant %uint 256
+   %uint_264 = OpConstant %uint 264
+        %233 = OpTypeFunction %mat3v4float %uint
+    %uint_32 = OpConstant %uint 32
+        %252 = OpTypeFunction %tint_GammaTransferParams %uint
+     %uint_8 = OpConstant %uint 8
+    %uint_12 = OpConstant %uint 12
+    %uint_20 = OpConstant %uint 20
+    %uint_24 = OpConstant %uint 24
+    %uint_28 = OpConstant %uint 28
+        %323 = OpTypeFunction %mat3v3float %uint
+     %v3uint = OpTypeVector %uint 3
+        %345 = OpTypeFunction %mat3v2float %uint
+       %main = OpFunction %void None %17
+         %18 = OpLabel
+       %vals = OpVariable %_ptr_Function_v4float Function
+         %19 = OpLoad %10 %t_f None
+         %20 = OpLoad %14 %t_f_ycbcr_sampler None
+         %21 = OpFunctionCall %tint_ExternalTextureParams %31 %uint_0
+         %33 = OpCompositeExtract %v2uint %21 12
+         %34 = OpIAdd %v2uint %33 %35
+         %37 = OpISub %v2uint %34 %35
+         %38 = OpBitcast %v2uint %39
+         %42 = OpExtInst %v2uint %43 UMin %38 %37
+         %44 = OpFunctionCall %v4float %tint_TextureLoadYcbcrExternal %19 %21 %42
+               OpStore %vals %44
+               OpReturn
+               OpFunctionEnd
+%tint_TextureLoadYcbcrExternal = OpFunction %v4float None %51
+    %texture = OpFunctionParameter %10
+     %params = OpFunctionParameter %tint_ExternalTextureParams
+     %coords = OpFunctionParameter %v2uint
+         %52 = OpLabel
+         %53 = OpCompositeExtract %uint %params 1
+         %54 = OpCompositeExtract %mat3v4float %params 2
+         %55 = OpCompositeExtract %mat3v2float %params 7
+         %56 = OpCompositeExtract %v2uint %params 12
+         %57 = OpExtInst %v2uint %43 UMin %coords %56
+         %58 = OpConvertUToF %v2float %57
+         %59 = OpCompositeConstruct %v3float %58 %float_1
+         %61 = OpMatrixTimesVector %v2float %55 %59
+         %62 = OpExtInst %v2float %43 RoundEven %61
+         %63 = OpConvertFToU %v2uint %62
+         %64 = OpImageFetch %v4float %texture %63 Lod %uint_0
+         %65 = OpVectorShuffle %v3float %64 %64 0 1 2
+         %66 = OpCompositeConstruct %v4float %65 %float_1
+         %67 = OpVectorTimesMatrix %v3float %66 %54
+         %68 = OpIEqual %bool %53 %uint_0
+               OpSelectionMerge %70 None
+               OpBranchConditional %68 %71 %72
+         %71 = OpLabel
+         %76 = OpCompositeExtract %tint_GammaTransferParams %params 3
+         %77 = OpCompositeExtract %tint_GammaTransferParams %params 4
+         %78 = OpCompositeExtract %mat3v3float %params 5
+         %79 = OpFunctionCall %v3float %tint_GammaCorrection %67 %76
+         %81 = OpMatrixTimesVector %v3float %78 %79
+         %74 = OpFunctionCall %v3float %tint_GammaCorrection %81 %77
+               OpBranch %70
+         %72 = OpLabel
+               OpBranch %70
+         %70 = OpLabel
+         %73 = OpPhi %v3float %74 %71 %67 %72
+         %75 = OpCompositeConstruct %v4float %73 %float_1
+               OpReturnValue %75
+               OpFunctionEnd
+%tint_GammaCorrection = OpFunction %v3float None %84
+          %v = OpFunctionParameter %v3float
+   %params_0 = OpFunctionParameter %tint_GammaTransferParams
+         %85 = OpLabel
+         %86 = OpCompositeExtract %float %params_0 0
+         %87 = OpCompositeExtract %float %params_0 1
+         %88 = OpCompositeExtract %float %params_0 2
+         %89 = OpCompositeExtract %float %params_0 3
+         %90 = OpCompositeExtract %float %params_0 4
+         %91 = OpCompositeExtract %float %params_0 5
+         %92 = OpCompositeExtract %float %params_0 6
+         %93 = OpCompositeConstruct %v3float %86 %86 %86
+         %94 = OpCompositeConstruct %v3float %90 %90 %90
+         %95 = OpExtInst %v3float %43 FAbs %v
+         %96 = OpExtInst %v3float %43 FSign %v
+         %97 = OpFOrdLessThan %v3bool %95 %94
+         %99 = OpVectorTimesScalar %v3float %95 %89
+        %100 = OpCompositeConstruct %v3float %92 %92 %92
+        %101 = OpFAdd %v3float %99 %100
+        %102 = OpFMul %v3float %96 %101
+        %103 = OpVectorTimesScalar %v3float %95 %87
+        %104 = OpCompositeConstruct %v3float %88 %88 %88
+        %105 = OpFAdd %v3float %103 %104
+        %106 = OpExtInst %v3float %43 Pow %105 %93
+        %107 = OpCompositeConstruct %v3float %91 %91 %91
+        %108 = OpFAdd %v3float %106 %107
+        %109 = OpFMul %v3float %96 %108
+        %110 = OpSelect %v3float %97 %102 %109
+               OpReturnValue %110
+               OpFunctionEnd
+         %31 = OpFunction %tint_ExternalTextureParams None %112
+%start_byte_offset = OpFunctionParameter %uint
+        %113 = OpLabel
+        %114 = OpUDiv %uint %start_byte_offset %uint_16
+        %116 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %114
+        %118 = OpBitwiseAnd %uint %start_byte_offset %uint_15
+        %120 = OpShiftRightLogical %uint %118 %uint_2
+        %122 = OpLoad %v4uint %116 None
+        %123 = OpVectorExtractDynamic %uint %122 %120
+        %124 = OpIAdd %uint %uint_4 %start_byte_offset
+        %126 = OpUDiv %uint %124 %uint_16
+        %127 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %126
+        %128 = OpBitwiseAnd %uint %124 %uint_15
+        %129 = OpShiftRightLogical %uint %128 %uint_2
+        %130 = OpLoad %v4uint %127 None
+        %131 = OpVectorExtractDynamic %uint %130 %129
+        %132 = OpIAdd %uint %uint_16 %start_byte_offset
+        %133 = OpFunctionCall %mat3v4float %134 %132
+        %135 = OpIAdd %uint %uint_64 %start_byte_offset
+        %137 = OpFunctionCall %tint_GammaTransferParams %138 %135
+        %139 = OpIAdd %uint %uint_96 %start_byte_offset
+        %141 = OpFunctionCall %tint_GammaTransferParams %138 %139
+        %142 = OpIAdd %uint %uint_128 %start_byte_offset
+        %144 = OpFunctionCall %mat3v3float %145 %142
+        %146 = OpIAdd %uint %uint_176 %start_byte_offset
+        %148 = OpFunctionCall %mat3v2float %149 %146
+        %150 = OpIAdd %uint %uint_200 %start_byte_offset
+        %152 = OpFunctionCall %mat3v2float %149 %150
+        %153 = OpIAdd %uint %uint_224 %start_byte_offset
+        %155 = OpUDiv %uint %153 %uint_16
+        %156 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %155
+        %157 = OpLoad %v4uint %156 None
+        %158 = OpBitwiseAnd %uint %153 %uint_15
+        %159 = OpShiftRightLogical %uint %158 %uint_2
+        %160 = OpVectorShuffle %v2uint %157 %157 2 3
+        %161 = OpVectorShuffle %v2uint %157 %157 0 1
+        %162 = OpIEqual %bool %159 %uint_2
+        %164 = OpCompositeConstruct %v2bool %162 %162
+        %165 = OpSelect %v2uint %164 %160 %161
+        %166 = OpBitcast %v2float %165
+        %167 = OpIAdd %uint %uint_232 %start_byte_offset
+        %169 = OpUDiv %uint %167 %uint_16
+        %170 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %169
+        %171 = OpLoad %v4uint %170 None
+        %172 = OpBitwiseAnd %uint %167 %uint_15
+        %173 = OpShiftRightLogical %uint %172 %uint_2
+        %174 = OpVectorShuffle %v2uint %171 %171 2 3
+        %175 = OpVectorShuffle %v2uint %171 %171 0 1
+        %176 = OpIEqual %bool %173 %uint_2
+        %177 = OpCompositeConstruct %v2bool %176 %176
+        %178 = OpSelect %v2uint %177 %174 %175
+        %179 = OpBitcast %v2float %178
+        %180 = OpIAdd %uint %uint_240 %start_byte_offset
+        %182 = OpUDiv %uint %180 %uint_16
+        %183 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %182
+        %184 = OpLoad %v4uint %183 None
+        %185 = OpBitwiseAnd %uint %180 %uint_15
+        %186 = OpShiftRightLogical %uint %185 %uint_2
+        %187 = OpVectorShuffle %v2uint %184 %184 2 3
+        %188 = OpVectorShuffle %v2uint %184 %184 0 1
+        %189 = OpIEqual %bool %186 %uint_2
+        %190 = OpCompositeConstruct %v2bool %189 %189
+        %191 = OpSelect %v2uint %190 %187 %188
+        %192 = OpBitcast %v2float %191
+        %193 = OpIAdd %uint %uint_248 %start_byte_offset
+        %195 = OpUDiv %uint %193 %uint_16
+        %196 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %195
+        %197 = OpLoad %v4uint %196 None
+        %198 = OpBitwiseAnd %uint %193 %uint_15
+        %199 = OpShiftRightLogical %uint %198 %uint_2
+        %200 = OpVectorShuffle %v2uint %197 %197 2 3
+        %201 = OpVectorShuffle %v2uint %197 %197 0 1
+        %202 = OpIEqual %bool %199 %uint_2
+        %203 = OpCompositeConstruct %v2bool %202 %202
+        %204 = OpSelect %v2uint %203 %200 %201
+        %205 = OpBitcast %v2float %204
+        %206 = OpIAdd %uint %uint_256 %start_byte_offset
+        %208 = OpUDiv %uint %206 %uint_16
+        %209 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %208
+        %210 = OpLoad %v4uint %209 None
+        %211 = OpBitwiseAnd %uint %206 %uint_15
+        %212 = OpShiftRightLogical %uint %211 %uint_2
+        %213 = OpVectorShuffle %v2uint %210 %210 2 3
+        %214 = OpVectorShuffle %v2uint %210 %210 0 1
+        %215 = OpIEqual %bool %212 %uint_2
+        %216 = OpCompositeConstruct %v2bool %215 %215
+        %217 = OpSelect %v2uint %216 %213 %214
+        %218 = OpIAdd %uint %uint_264 %start_byte_offset
+        %220 = OpUDiv %uint %218 %uint_16
+        %221 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %220
+        %222 = OpLoad %v4uint %221 None
+        %223 = OpBitwiseAnd %uint %218 %uint_15
+        %224 = OpShiftRightLogical %uint %223 %uint_2
+        %225 = OpVectorShuffle %v2uint %222 %222 2 3
+        %226 = OpVectorShuffle %v2uint %222 %222 0 1
+        %227 = OpIEqual %bool %224 %uint_2
+        %228 = OpCompositeConstruct %v2bool %227 %227
+        %229 = OpSelect %v2uint %228 %225 %226
+        %230 = OpBitcast %v2float %229
+        %231 = OpCompositeConstruct %tint_ExternalTextureParams %123 %131 %133 %137 %141 %144 %148 %152 %166 %179 %192 %205 %217 %230
+               OpReturnValue %231
+               OpFunctionEnd
+        %134 = OpFunction %mat3v4float None %233
+%start_byte_offset_0 = OpFunctionParameter %uint
+        %234 = OpLabel
+        %235 = OpUDiv %uint %start_byte_offset_0 %uint_16
+        %236 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %235
+        %237 = OpLoad %v4uint %236 None
+        %238 = OpBitcast %v4float %237
+        %239 = OpIAdd %uint %uint_16 %start_byte_offset_0
+        %240 = OpUDiv %uint %239 %uint_16
+        %241 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %240
+        %242 = OpLoad %v4uint %241 None
+        %243 = OpBitcast %v4float %242
+        %244 = OpIAdd %uint %uint_32 %start_byte_offset_0
+        %246 = OpUDiv %uint %244 %uint_16
+        %247 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %246
+        %248 = OpLoad %v4uint %247 None
+        %249 = OpBitcast %v4float %248
+        %250 = OpCompositeConstruct %mat3v4float %238 %243 %249
+               OpReturnValue %250
+               OpFunctionEnd
+        %138 = OpFunction %tint_GammaTransferParams None %252
+%start_byte_offset_1 = OpFunctionParameter %uint
+        %253 = OpLabel
+        %254 = OpUDiv %uint %start_byte_offset_1 %uint_16
+        %255 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %254
+        %256 = OpBitwiseAnd %uint %start_byte_offset_1 %uint_15
+        %257 = OpShiftRightLogical %uint %256 %uint_2
+        %258 = OpLoad %v4uint %255 None
+        %259 = OpVectorExtractDynamic %uint %258 %257
+        %260 = OpBitcast %float %259
+        %261 = OpIAdd %uint %uint_4 %start_byte_offset_1
+        %262 = OpUDiv %uint %261 %uint_16
+        %263 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %262
+        %264 = OpBitwiseAnd %uint %261 %uint_15
+        %265 = OpShiftRightLogical %uint %264 %uint_2
+        %266 = OpLoad %v4uint %263 None
+        %267 = OpVectorExtractDynamic %uint %266 %265
+        %268 = OpBitcast %float %267
+        %269 = OpIAdd %uint %uint_8 %start_byte_offset_1
+        %271 = OpUDiv %uint %269 %uint_16
+        %272 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %271
+        %273 = OpBitwiseAnd %uint %269 %uint_15
+        %274 = OpShiftRightLogical %uint %273 %uint_2
+        %275 = OpLoad %v4uint %272 None
+        %276 = OpVectorExtractDynamic %uint %275 %274
+        %277 = OpBitcast %float %276
+        %278 = OpIAdd %uint %uint_12 %start_byte_offset_1
+        %280 = OpUDiv %uint %278 %uint_16
+        %281 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %280
+        %282 = OpBitwiseAnd %uint %278 %uint_15
+        %283 = OpShiftRightLogical %uint %282 %uint_2
+        %284 = OpLoad %v4uint %281 None
+        %285 = OpVectorExtractDynamic %uint %284 %283
+        %286 = OpBitcast %float %285
+        %287 = OpIAdd %uint %uint_16 %start_byte_offset_1
+        %288 = OpUDiv %uint %287 %uint_16
+        %289 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %288
+        %290 = OpBitwiseAnd %uint %287 %uint_15
+        %291 = OpShiftRightLogical %uint %290 %uint_2
+        %292 = OpLoad %v4uint %289 None
+        %293 = OpVectorExtractDynamic %uint %292 %291
+        %294 = OpBitcast %float %293
+        %295 = OpIAdd %uint %uint_20 %start_byte_offset_1
+        %297 = OpUDiv %uint %295 %uint_16
+        %298 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %297
+        %299 = OpBitwiseAnd %uint %295 %uint_15
+        %300 = OpShiftRightLogical %uint %299 %uint_2
+        %301 = OpLoad %v4uint %298 None
+        %302 = OpVectorExtractDynamic %uint %301 %300
+        %303 = OpBitcast %float %302
+        %304 = OpIAdd %uint %uint_24 %start_byte_offset_1
+        %306 = OpUDiv %uint %304 %uint_16
+        %307 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %306
+        %308 = OpBitwiseAnd %uint %304 %uint_15
+        %309 = OpShiftRightLogical %uint %308 %uint_2
+        %310 = OpLoad %v4uint %307 None
+        %311 = OpVectorExtractDynamic %uint %310 %309
+        %312 = OpBitcast %float %311
+        %313 = OpIAdd %uint %uint_28 %start_byte_offset_1
+        %315 = OpUDiv %uint %313 %uint_16
+        %316 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %315
+        %317 = OpBitwiseAnd %uint %313 %uint_15
+        %318 = OpShiftRightLogical %uint %317 %uint_2
+        %319 = OpLoad %v4uint %316 None
+        %320 = OpVectorExtractDynamic %uint %319 %318
+        %321 = OpCompositeConstruct %tint_GammaTransferParams %260 %268 %277 %286 %294 %303 %312 %320
+               OpReturnValue %321
+               OpFunctionEnd
+        %145 = OpFunction %mat3v3float None %323
+%start_byte_offset_2 = OpFunctionParameter %uint
+        %324 = OpLabel
+        %325 = OpUDiv %uint %start_byte_offset_2 %uint_16
+        %326 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %325
+        %327 = OpLoad %v4uint %326 None
+        %328 = OpVectorShuffle %v3uint %327 %327 0 1 2
+        %330 = OpBitcast %v3float %328
+        %331 = OpIAdd %uint %uint_16 %start_byte_offset_2
+        %332 = OpUDiv %uint %331 %uint_16
+        %333 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %332
+        %334 = OpLoad %v4uint %333 None
+        %335 = OpVectorShuffle %v3uint %334 %334 0 1 2
+        %336 = OpBitcast %v3float %335
+        %337 = OpIAdd %uint %uint_32 %start_byte_offset_2
+        %338 = OpUDiv %uint %337 %uint_16
+        %339 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %338
+        %340 = OpLoad %v4uint %339 None
+        %341 = OpVectorShuffle %v3uint %340 %340 0 1 2
+        %342 = OpBitcast %v3float %341
+        %343 = OpCompositeConstruct %mat3v3float %330 %336 %342
+               OpReturnValue %343
+               OpFunctionEnd
+        %149 = OpFunction %mat3v2float None %345
+%start_byte_offset_3 = OpFunctionParameter %uint
+        %346 = OpLabel
+        %347 = OpUDiv %uint %start_byte_offset_3 %uint_16
+        %348 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %347
+        %349 = OpLoad %v4uint %348 None
+        %350 = OpBitwiseAnd %uint %start_byte_offset_3 %uint_15
+        %351 = OpShiftRightLogical %uint %350 %uint_2
+        %352 = OpVectorShuffle %v2uint %349 %349 2 3
+        %353 = OpVectorShuffle %v2uint %349 %349 0 1
+        %354 = OpIEqual %bool %351 %uint_2
+        %355 = OpCompositeConstruct %v2bool %354 %354
+        %356 = OpSelect %v2uint %355 %352 %353
+        %357 = OpBitcast %v2float %356
+        %358 = OpIAdd %uint %uint_8 %start_byte_offset_3
+        %359 = OpUDiv %uint %358 %uint_16
+        %360 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %359
+        %361 = OpLoad %v4uint %360 None
+        %362 = OpBitwiseAnd %uint %358 %uint_15
+        %363 = OpShiftRightLogical %uint %362 %uint_2
+        %364 = OpVectorShuffle %v2uint %361 %361 2 3
+        %365 = OpVectorShuffle %v2uint %361 %361 0 1
+        %366 = OpIEqual %bool %363 %uint_2
+        %367 = OpCompositeConstruct %v2bool %366 %366
+        %368 = OpSelect %v2uint %367 %364 %365
+        %369 = OpBitcast %v2float %368
+        %370 = OpIAdd %uint %uint_16 %start_byte_offset_3
+        %371 = OpUDiv %uint %370 %uint_16
+        %372 = OpAccessChain %_ptr_Uniform_v4uint %1 %uint_0 %371
+        %373 = OpLoad %v4uint %372 None
+        %374 = OpBitwiseAnd %uint %370 %uint_15
+        %375 = OpShiftRightLogical %uint %374 %uint_2
+        %376 = OpVectorShuffle %v2uint %373 %373 2 3
+        %377 = OpVectorShuffle %v2uint %373 %373 0 1
+        %378 = OpIEqual %bool %375 %uint_2
+        %379 = OpCompositeConstruct %v2bool %378 %378
+        %380 = OpSelect %v2uint %379 %376 %377
+        %381 = OpBitcast %v2float %380
+        %382 = OpCompositeConstruct %mat3v2float %357 %369 %381
+               OpReturnValue %382
+               OpFunctionEnd

diff --git a/test/tint/types/texture/external/external_ycbcr.wgsl.expected.wgsl b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.wgsl
new file mode 100644
index 0000000..d31becd
--- /dev/null
+++ b/test/tint/types/texture/external/external_ycbcr.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@group(0) @binding(0) var t_f : texture_external;
+
+@compute @workgroup_size(1)
+fn main() {
+  var vals = textureLoad(t_f, vec2i(0));
+}