tint: Implement textureSampleBaseClampToEdge

Fixed: tint:1671
Change-Id: Iaae5b5d571a4401c0255de727245bf8dbbe06740
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/102642
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

docs/tint/origin-trial-changes.md

@@ -1,5 +1,15 @@
 # Tint changes during Origin Trial
 
+## Changes for M108
+
+### New features
+
+* `textureSampleBaseClampToEdge()` has been implemented. [tint:1671](crbug.com/tint/1671)
+
+### Deprecated Features
+
+* The `external_texture` overload of `textureSampleLevel()` has been deprecated. Use `textureSampleBaseClampToEdge()` instead. [tint:1671](crbug.com/tint/1671)
+
 ## Changes for M107
 
 ### New features

src/tint/intrinsics.def

@@ -659,7 +659,9 @@ fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords:
 fn textureSampleLevel(texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: i32, level: i32, @const offset: vec2<i32>) -> f32
 fn textureSampleLevel(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>, level: i32) -> f32
 fn textureSampleLevel(texture: texture_depth_cube_array,sampler: sampler, coords: vec3<f32>, array_index: i32, level: i32) -> f32
-fn textureSampleLevel(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+@deprecated fn textureSampleLevel(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureSampleBaseClampToEdge(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureSampleBaseClampToEdge(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
 fn textureStore(texture: texture_storage_1d<f32_texel_format, write>, coords: i32, value: vec4<f32>)
 fn textureStore(texture: texture_storage_2d<f32_texel_format, write>, coords: vec2<i32>, value: vec4<f32>)
 fn textureStore(texture: texture_storage_2d_array<f32_texel_format, write>, coords: vec2<i32>, array_index: i32, value: vec4<f32>)

src/tint/sem/builtin.cc

@@ -57,11 +57,17 @@ bool IsDerivativeBuiltin(BuiltinType i) {
 }
 
 bool IsTextureBuiltin(BuiltinType i) {
-    return IsImageQueryBuiltin(i) || i == BuiltinType::kTextureLoad ||
+    return IsImageQueryBuiltin(i) ||                           //
-           i == BuiltinType::kTextureGather || i == BuiltinType::kTextureGatherCompare ||
+           i == BuiltinType::kTextureGather ||                 //
-           i == BuiltinType::kTextureSample || i == BuiltinType::kTextureSampleLevel ||
+           i == BuiltinType::kTextureGatherCompare ||          //
-           i == BuiltinType::kTextureSampleBias || i == BuiltinType::kTextureSampleCompare ||
+           i == BuiltinType::kTextureLoad ||                   //
-           i == BuiltinType::kTextureSampleCompareLevel || i == BuiltinType::kTextureSampleGrad ||
+           i == BuiltinType::kTextureSample ||                 //
+           i == BuiltinType::kTextureSampleBaseClampToEdge ||  //
+           i == BuiltinType::kTextureSampleBias ||             //
+           i == BuiltinType::kTextureSampleCompare ||          //
+           i == BuiltinType::kTextureSampleCompareLevel ||     //
+           i == BuiltinType::kTextureSampleGrad ||             //
+           i == BuiltinType::kTextureSampleLevel ||            //
            i == BuiltinType::kTextureStore;
 }
 

src/tint/sem/builtin_type.cc

@@ -318,6 +318,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
     if (name == "textureSampleLevel") {
         return BuiltinType::kTextureSampleLevel;
     }
+    if (name == "textureSampleBaseClampToEdge") {
+        return BuiltinType::kTextureSampleBaseClampToEdge;
+    }
     if (name == "textureStore") {
         return BuiltinType::kTextureStore;
     }
@@ -558,6 +561,8 @@ const char* str(BuiltinType i) {
             return "textureSampleGrad";
         case BuiltinType::kTextureSampleLevel:
             return "textureSampleLevel";
+        case BuiltinType::kTextureSampleBaseClampToEdge:
+            return "textureSampleBaseClampToEdge";
         case BuiltinType::kTextureStore:
             return "textureStore";
         case BuiltinType::kTextureLoad:

src/tint/sem/builtin_type.h

@@ -128,6 +128,7 @@ enum class BuiltinType {
     kTextureSampleCompareLevel,
     kTextureSampleGrad,
     kTextureSampleLevel,
+    kTextureSampleBaseClampToEdge,
     kTextureStore,
     kTextureLoad,
     kAtomicLoad,

src/tint/sem/call_target.cc

@@ -34,7 +34,7 @@ CallTarget::CallTarget(const CallTarget&) = default;
 CallTarget::~CallTarget() = default;
 
 CallTargetSignature::CallTargetSignature(const sem::Type* ret_ty,
-                                         utils::VectorRef<const Parameter*> params)
+                                         utils::VectorRef<const sem::Parameter*> params)
     : return_type(ret_ty), parameters(std::move(params)) {}
 CallTargetSignature::CallTargetSignature(const CallTargetSignature&) = default;
 CallTargetSignature::~CallTargetSignature() = default;

src/tint/sem/call_target.h

@@ -57,6 +57,14 @@ struct CallTargetSignature {
     /// @returns the index of the parameter with the given usage, or -1 if no
     /// parameter with the given usage exists.
     int IndexOf(ParameterUsage usage) const;
+
+    /// @param usage  the parameter usage to find
+    /// @returns the the parameter with the given usage, or nullptr if no parameter with the given
+    /// usage exists.
+    inline const sem::Parameter* Parameter(ParameterUsage usage) const {
+        auto idx = IndexOf(usage);
+        return (idx >= 0) ? parameters[static_cast<size_t>(idx)] : nullptr;
+    }
 };
 
 /// CallTarget is the base for callable functions, builtins, type constructors

src/tint/transform/builtin_polyfill.cc

@@ -512,6 +512,29 @@ struct BuiltinPolyfill::State {
         return name;
     }
 
+    /// Builds the polyfill function for the `textureSampleBaseClampToEdge` builtin, when the
+    /// texture type is texture_2d<f32>.
+    /// @return the polyfill function name
+    Symbol textureSampleBaseClampToEdge_2d_f32() {
+        auto name = b.Symbols().New("tint_textureSampleBaseClampToEdge");
+        auto body = utils::Vector{
+            b.Decl(b.Let("dims",
+                         b.Construct(b.ty.vec2<f32>(), b.Call("textureDimensions", "t", 0_a)))),
+            b.Decl(b.Let("half_texel", b.Div(b.vec2<f32>(0.5_a), "dims"))),
+            b.Decl(
+                b.Let("clamped", b.Call("clamp", "coord", "half_texel", b.Sub(1_a, "half_texel")))),
+            b.Return(b.Call("textureSampleLevel", "t", "s", "clamped", 0_a)),
+        };
+        b.Func(name,
+               utils::Vector{
+                   b.Param("t", b.ty.sampled_texture(ast::TextureDimension::k2d, b.ty.f32())),
+                   b.Param("s", b.ty.sampler(ast::SamplerKind::kSampler)),
+                   b.Param("coord", b.ty.vec2<f32>()),
+               },
+               b.ty.vec4<f32>(), body);
+        return name;
+    }
+
   private:
     /// @returns the AST type for the given sem type
     const ast::Type* T(const sem::Type* ty) const { return CreateASTTypeFor(ctx, ty); }
@@ -597,6 +620,15 @@ bool BuiltinPolyfill::ShouldRun(const Program* program, const DataMap& data) con
                                 return true;
                             }
                             break;
+                        case sem::BuiltinType::kTextureSampleBaseClampToEdge:
+                            if (builtins.texture_sample_base_clamp_to_edge_2d_f32) {
+                                auto& sig = builtin->Signature();
+                                auto* tex = sig.Parameter(sem::ParameterUsage::kTexture);
+                                if (auto* stex = tex->Type()->As<sem::SampledTexture>()) {
+                                    return stex->type()->Is<sem::F32>();
+                                }
+                            }
+                            break;
                         default:
                             break;
                     }
@@ -690,6 +722,19 @@ void BuiltinPolyfill::Run(CloneContext& ctx, const DataMap& data, DataMap&) cons
                             });
                         }
                         break;
+                    case sem::BuiltinType::kTextureSampleBaseClampToEdge:
+                        if (builtins.texture_sample_base_clamp_to_edge_2d_f32) {
+                            auto& sig = builtin->Signature();
+                            auto* tex = sig.Parameter(sem::ParameterUsage::kTexture);
+                            if (auto* stex = tex->Type()->As<sem::SampledTexture>()) {
+                                if (stex->type()->Is<sem::F32>()) {
+                                    polyfill = utils::GetOrCreate(polyfills, builtin, [&] {
+                                        return s.textureSampleBaseClampToEdge_2d_f32();
+                                    });
+                                }
+                            }
+                        }
+                        break;
                     default:
                         break;
                 }

src/tint/transform/builtin_polyfill.h

@@ -61,6 +61,8 @@ class BuiltinPolyfill final : public Castable<BuiltinPolyfill, Transform> {
         Level insert_bits = Level::kNone;
         /// Should `saturate()` be polyfilled?
         bool saturate = false;
+        /// Should `textureSampleBaseClampToEdge()` be polyfilled for texture_2d<f32> textures?
+        bool texture_sample_base_clamp_to_edge_2d_f32 = false;
     };
 
     /// Config is consumed by the BuiltinPolyfill transform.

src/tint/transform/builtin_polyfill_test.cc

@@ -1549,5 +1549,76 @@ fn f() {
     EXPECT_EQ(expect, str(got));
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// textureSampleBaseClampToEdge
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillTextureSampleBaseClampToEdge_2d_f32() {
+    BuiltinPolyfill::Builtins builtins;
+    builtins.texture_sample_base_clamp_to_edge_2d_f32 = true;
+    DataMap data;
+    data.Add<BuiltinPolyfill::Config>(builtins);
+    return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunTextureSampleBaseClampToEdge_2d_f32) {
+    auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+@group(0) @binding(1) var s : sampler;
+
+fn f() {
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5));
+}
+)";
+
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+    EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillTextureSampleBaseClampToEdge_2d_f32()));
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunTextureSampleBaseClampToEdge_external) {
+    auto* src = R"(
+@group(0) @binding(0) var t : texture_external;
+@group(0) @binding(1) var s : sampler;
+
+fn f() {
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5));
+}
+)";
+
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src, polyfillTextureSampleBaseClampToEdge_2d_f32()));
+}
+
+TEST_F(BuiltinPolyfillTest, TextureSampleBaseClampToEdge_2d_f32_f32) {
+    auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+@group(0) @binding(1) var s : sampler;
+
+fn f() {
+  let r = textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5));
+}
+)";
+
+    auto* expect = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+fn tint_textureSampleBaseClampToEdge(t : texture_2d<f32>, s : sampler, coord : vec2<f32>) -> vec4<f32> {
+  let dims = vec2<f32>(textureDimensions(t, 0));
+  let half_texel = (vec2<f32>(0.5) / dims);
+  let clamped = clamp(coord, half_texel, (1 - half_texel));
+  return textureSampleLevel(t, s, clamped, 0);
+}
+
+fn f() {
+  let r = tint_textureSampleBaseClampToEdge(t, s, vec2<f32>(0.5));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillTextureSampleBaseClampToEdge_2d_f32());
+
+    EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace tint::transform

src/tint/transform/multiplanar_external_texture.cc

@@ -31,8 +31,8 @@ using namespace tint::number_suffixes;  // NOLINT
 namespace tint::transform {
 namespace {
 
-/// This struct stores symbols for new bindings created as a result of
+/// This struct stores symbols for new bindings created as a result of transforming a
-/// transforming a texture_external instance.
+/// texture_external instance.
 struct NewBindingSymbols {
     Symbol params;
     Symbol plane_0;
@@ -64,11 +64,14 @@ struct MultiplanarExternalTexture::State {
     /// Symbol for the textureSampleExternal function
     Symbol texture_sample_external_sym;
 
+    /// Symbol for the textureSampleExternalDEPRECATED function
+    Symbol texture_sample_external_deprecated_sym;
+
     /// Symbol for the gammaCorrection function
     Symbol gamma_correction_sym;
 
-    /// Storage for new bindings that have been created corresponding to an
+    /// Storage for new bindings that have been created corresponding to an original
-    /// original texture_external binding.
+    /// texture_external binding.
     std::unordered_map<const sem::Variable*, NewBindingSymbols> new_binding_symbols;
 
     /// Constructor
@@ -82,33 +85,30 @@ struct MultiplanarExternalTexture::State {
     void Process() {
         auto& sem = ctx.src->Sem();
 
-        // For each texture_external binding, we replace it with a texture_2d<f32>
+        // For each texture_external binding, we replace it with a texture_2d<f32> binding and
-        // binding and create two additional bindings (one texture_2d<f32> to
+        // create two additional bindings (one texture_2d<f32> to represent the secondary plane and
-        // represent the secondary plane and one uniform buffer for the
+        // one uniform buffer for the ExternalTextureParams struct).
-        // ExternalTextureParams struct).
         for (auto* global : ctx.src->AST().GlobalVariables()) {
             auto* sem_var = sem.Get<sem::GlobalVariable>(global);
             if (!sem_var->Type()->UnwrapRef()->Is<sem::ExternalTexture>()) {
                 continue;
             }
 
-            // If the attributes are empty, then this must be a texture_external
+            // If the attributes are empty, then this must be a texture_external passed as a
-            // passed as a function parameter. These variables are transformed
+            // function parameter. These variables are transformed elsewhere.
-            // elsewhere.
             if (global->attributes.IsEmpty()) {
                 continue;
             }
 
-            // If we find a texture_external binding, we know we must emit the
+            // If we find a texture_external binding, we know we must emit the ExternalTextureParams
-            // ExternalTextureParams struct.
+            // struct.
             if (!params_struct_sym.IsValid()) {
                 createExtTexParamsStructs();
             }
 
-            // The binding points for the newly introduced bindings must have been
+            // The binding points for the newly introduced bindings must have been provided to this
-            // provided to this transform. We fetch the new binding points by
+            // transform. We fetch the new binding points by providing the original texture_external
-            // providing the original texture_external binding points into the
+            // binding points into the passed map.
-            // passed map.
             BindingPoint bp = sem_var->BindingPoint();
 
             BindingsMap::const_iterator it = new_binding_points->bindings_map.find(bp);
@@ -122,10 +122,9 @@ struct MultiplanarExternalTexture::State {
 
             BindingPoints bps = it->second;
 
-            // Symbols for the newly created bindings must be saved so they can be
+            // Symbols for the newly created bindings must be saved so they can be passed as
-            // passed as parameters later. These are placed in a map and keyed by
+            // parameters later. These are placed in a map and keyed by the source symbol associated
-            // the source symbol associated with the texture_external binding that
+            // with the texture_external binding that corresponds with the new destination bindings.
-            // corresponds with the new destination bindings.
             // NewBindingSymbols new_binding_syms;
             auto& syms = new_binding_symbols[sem_var];
             syms.plane_0 = ctx.Clone(global->symbol);
@@ -137,8 +136,7 @@ struct MultiplanarExternalTexture::State {
                         ast::StorageClass::kUniform, b.Group(AInt(bps.params.group)),
                         b.Binding(AInt(bps.params.binding)));
 
-            // Replace the original texture_external binding with a texture_2d<f32>
+            // Replace the original texture_external binding with a texture_2d<f32> binding.
-            // binding.
             auto cloned_attributes = ctx.Clone(global->attributes);
             const ast::Expression* cloned_constructor = ctx.Clone(global->constructor);
 
@@ -148,23 +146,22 @@ struct MultiplanarExternalTexture::State {
             ctx.Replace(global, replacement);
         }
 
-        // We must update all the texture_external parameters for user declared
+        // We must update all the texture_external parameters for user declared functions.
-        // functions.
         for (auto* fn : ctx.src->AST().Functions()) {
             for (const ast::Variable* param : fn->params) {
                 if (auto* sem_var = sem.Get(param)) {
                     if (!sem_var->Type()->UnwrapRef()->Is<sem::ExternalTexture>()) {
                         continue;
                     }
-                    // If we find a texture_external, we must ensure the
+                    // If we find a texture_external, we must ensure the ExternalTextureParams
-                    // ExternalTextureParams struct exists.
+                    // struct exists.
                     if (!params_struct_sym.IsValid()) {
                         createExtTexParamsStructs();
                     }
-                    // When a texture_external is found, we insert all components
+                    // When a texture_external is found, we insert all components the
-                    // the texture_external into the parameter list. We must also place
+                    // texture_external into the parameter list. We must also place the new symbols
-                    // the new symbols into the transform state so they can be used when
+                    // into the transform state so they can be used when transforming function
-                    // transforming function calls.
+                    // calls.
                     auto& syms = new_binding_symbols[sem_var];
                     syms.plane_0 = ctx.Clone(param->symbol);
                     syms.plane_1 = b.Symbols().New("ext_tex_plane_1");
@@ -180,8 +177,8 @@ struct MultiplanarExternalTexture::State {
             }
         }
 
-        // Transform the original textureLoad and textureSampleLevel calls into
+        // Transform the external texture builtin calls into calls to the external texture
-        // textureLoadExternal and textureSampleExternal calls.
+        // functions.
         ctx.ReplaceAll([&](const ast::CallExpression* expr) -> const ast::CallExpression* {
             auto* call = sem.Get(expr)->UnwrapMaterialize()->As<sem::Call>();
             auto* builtin = call->Target()->As<sem::Builtin>();
@@ -192,28 +189,29 @@ struct MultiplanarExternalTexture::State {
                 if (auto* var_user = sem.Get<sem::VariableUser>(expr->args[0])) {
                     auto it = new_binding_symbols.find(var_user->Variable());
                     if (it == new_binding_symbols.end()) {
-                        // If valid new binding locations were not provided earlier, we
+                        // If valid new binding locations were not provided earlier, we would have
-                        // would have been unable to create these symbols. An error
+                        // been unable to create these symbols. An error message was emitted
-                        // message was emitted earlier, so just return early to avoid
+                        // earlier, so just return early to avoid internal compiler errors and
-                        // internal compiler errors and retain a clean error message.
+                        // retain a clean error message.
                         return nullptr;
                     }
                     auto& syms = it->second;
 
-                    if (builtin->Type() == sem::BuiltinType::kTextureLoad) {
+                    switch (builtin->Type()) {
-                        return createTexLdExt(expr, syms);
+                        case sem::BuiltinType::kTextureLoad:
-                    }
+                            return createTextureLoad(expr, syms);
-
+                        case sem::BuiltinType::kTextureSampleLevel:
-                    if (builtin->Type() == sem::BuiltinType::kTextureSampleLevel) {
+                            return createTextureSampleLevel(expr, syms);
-                        return createTexSmpExt(expr, syms);
+                        case sem::BuiltinType::kTextureSampleBaseClampToEdge:
+                            return createTextureSampleBaseClampToEdge(expr, syms);
+                        default:
+                            break;
                     }
                 }
-
             } else if (call->Target()->Is<sem::Function>()) {
-                // The call expression may be to a user-defined function that
+                // The call expression may be to a user-defined function that contains a
-                // contains a texture_external parameter. These need to be expanded
+                // texture_external parameter. These need to be expanded out to multiple plane
-                // out to multiple plane textures and the texture parameters
+                // textures and the texture parameters structure.
-                // structure.
                 for (auto* arg : expr->args) {
                     if (auto* var_user = sem.Get<sem::VariableUser>(arg)) {
                         // Check if a parameter is a texture_external by trying to find
@@ -296,38 +294,69 @@ struct MultiplanarExternalTexture::State {
             });
     }
 
-    /// Constructs a StatementList containing all the statements making up the
+    /// Constructs a StatementList containing all the statements making up the body of the texture
-    /// bodies of the textureSampleExternal and textureLoadExternal functions.
+    /// builtin function.
     /// @param call_type determines which function body to generate
     /// @returns a statement list that makes of the body of the chosen function
-    auto createTexFnExtStatementList(sem::BuiltinType call_type) {
+    auto buildTextureBuiltinBody(sem::BuiltinType call_type) {
+        utils::Vector<const ast::Statement*, 16> stmts;
         const ast::CallExpression* single_plane_call = nullptr;
         const ast::CallExpression* plane_0_call = nullptr;
         const ast::CallExpression* plane_1_call = nullptr;
-        if (call_type == sem::BuiltinType::kTextureSampleLevel) {
+        switch (call_type) {
-            // textureSampleLevel(plane0, smp, coord.xy, 0.0);
+            case sem::BuiltinType::kTextureSampleLevel:
-            single_plane_call = b.Call("textureSampleLevel", "plane0", "smp", "coord", 0_f);
+                // TODO(crbug.com/tint/1671): DEPRECATED
-            // textureSampleLevel(plane0, smp, coord.xy, 0.0);
+                // textureSampleLevel(plane0, smp, coord, 0.0);
-            plane_0_call = b.Call("textureSampleLevel", "plane0", "smp", "coord", 0_f);
+                single_plane_call = b.Call("textureSampleLevel", "plane0", "smp", "coord", 0_f);
-            // textureSampleLevel(plane1, smp, coord.xy, 0.0);
+                // textureSampleLevel(plane0, smp, coord, 0.0);
-            plane_1_call = b.Call("textureSampleLevel", "plane1", "smp", "coord", 0_f);
+                plane_0_call = b.Call("textureSampleLevel", "plane0", "smp", "coord", 0_f);
-        } else if (call_type == sem::BuiltinType::kTextureLoad) {
+                // textureSampleLevel(plane1, smp, coord, 0.0);
-            // textureLoad(plane0, coords.xy, 0);
+                plane_1_call = b.Call("textureSampleLevel", "plane1", "smp", "coord", 0_f);
-            single_plane_call = b.Call("textureLoad", "plane0", "coord", 0_i);
+                break;
-            // textureLoad(plane0, coords.xy, 0);
+            case sem::BuiltinType::kTextureSampleBaseClampToEdge:
-            plane_0_call = b.Call("textureLoad", "plane0", "coord", 0_i);
+                stmts.Push(b.Decl(b.Let(
-            // textureLoad(plane1, coords.xy, 0);
+                    "plane0_dims",
-            plane_1_call = b.Call("textureLoad", "plane1", "coord", 0_i);
+                    b.Construct(b.ty.vec2<f32>(), b.Call("textureDimensions", "plane0", 0_a)))));
-        } else {
+                stmts.Push(
-            TINT_ICE(Transform, b.Diagnostics()) << "unhandled builtin: " << call_type;
+                    b.Decl(b.Let("plane0_half_texel", b.Div(b.vec2<f32>(0.5_a), "plane0_dims"))));
+                stmts.Push(
+                    b.Decl(b.Let("plane0_clamped", b.Call("clamp", "coord", "plane0_half_texel",
+                                                          b.Sub(1_a, "plane0_half_texel")))));
+                stmts.Push(b.Decl(b.Let(
+                    "plane1_dims",
+                    b.Construct(b.ty.vec2<f32>(), b.Call("textureDimensions", "plane1", 0_a)))));
+                stmts.Push(
+                    b.Decl(b.Let("plane1_half_texel", b.Div(b.vec2<f32>(0.5_a), "plane1_dims"))));
+                stmts.Push(
+                    b.Decl(b.Let("plane1_clamped", b.Call("clamp", "coord", "plane1_half_texel",
+                                                          b.Sub(1_a, "plane1_half_texel")))));
+
+                // textureSampleLevel(plane0, smp, plane0_clamped, 0.0);
+                single_plane_call =
+                    b.Call("textureSampleLevel", "plane0", "smp", "plane0_clamped", 0_f);
+                // textureSampleLevel(plane0, smp, plane0_clamped, 0.0);
+                plane_0_call = b.Call("textureSampleLevel", "plane0", "smp", "plane0_clamped", 0_f);
+                // textureSampleLevel(plane1, smp, plane1_clamped, 0.0);
+                plane_1_call = b.Call("textureSampleLevel", "plane1", "smp", "plane1_clamped", 0_f);
+                break;
+            case sem::BuiltinType::kTextureLoad:
+                // textureLoad(plane0, coord, 0);
+                single_plane_call = b.Call("textureLoad", "plane0", "coord", 0_i);
+                // textureLoad(plane0, coord, 0);
+                plane_0_call = b.Call("textureLoad", "plane0", "coord", 0_i);
+                // textureLoad(plane1, coord, 0);
+                plane_1_call = b.Call("textureLoad", "plane1", "coord", 0_i);
+                break;
+            default:
+                TINT_ICE(Transform, b.Diagnostics()) << "unhandled builtin: " << call_type;
         }
 
-        return utils::Vector{
+        // var color: vec3<f32>;
-            // var color: vec3<f32>;
+        stmts.Push(b.Decl(b.Var("color", b.ty.vec3(b.ty.f32()))));
-            b.Decl(b.Var("color", b.ty.vec3(b.ty.f32()))),
+
-            // if ((params.numPlanes == 1u))
+        // if ((params.numPlanes == 1u))
-            b.If(b.create<ast::BinaryExpression>(
+        stmts.Push(
-                     ast::BinaryOp::kEqual, b.MemberAccessor("params", "numPlanes"), b.Expr(1_u)),
+            b.If(b.Equal(b.MemberAccessor("params", "numPlanes"), b.Expr(1_u)),
                  b.Block(
                      // color = textureLoad(plane0, coord, 0).rgb;
                      b.Assign("color", b.MemberAccessor(single_plane_call, "rgb"))),
@@ -337,11 +366,11 @@ struct MultiplanarExternalTexture::State {
                      b.Assign("color",
                               b.Mul(b.vec4<f32>(b.MemberAccessor(plane_0_call, "r"),
                                                 b.MemberAccessor(plane_1_call, "rg"), 1_f),
-                                    b.MemberAccessor("params", "yuvToRgbConversionMatrix")))))),
+                                    b.MemberAccessor("params", "yuvToRgbConversionMatrix")))))));
-            // if (params.doYuvToRgbConversionOnly == 0u)
+
-            b.If(b.create<ast::BinaryExpression>(
+        // if (params.doYuvToRgbConversionOnly == 0u)
-                     ast::BinaryOp::kEqual, b.MemberAccessor("params", "doYuvToRgbConversionOnly"),
+        stmts.Push(
-                     b.Expr(0_u)),
+            b.If(b.Equal(b.MemberAccessor("params", "doYuvToRgbConversionOnly"), b.Expr(0_u)),
                  b.Block(
                      // color = gammaConversion(color, gammaDecodeParams);
                      b.Assign("color", b.Call("gammaCorrection", "color",
@@ -351,18 +380,21 @@ struct MultiplanarExternalTexture::State {
                               b.Mul(b.MemberAccessor("params", "gamutConversionMatrix"), "color")),
                      // color = gammaConversion(color, gammaEncodeParams);
                      b.Assign("color", b.Call("gammaCorrection", "color",
-                                              b.MemberAccessor("params", "gammaEncodeParams"))))),
+                                              b.MemberAccessor("params", "gammaEncodeParams"))))));
-            // return vec4<f32>(color, 1.f);
+
-            b.Return(b.vec4<f32>("color", 1_f))};
+        // return vec4<f32>(color, 1.f);
+        stmts.Push(b.Return(b.vec4<f32>("color", 1_f)));
+
+        return stmts;
     }
 
-    /// Creates the textureSampleExternal function if needed and returns a call
+    /// Creates the textureSampleExternal function if needed and returns a call expression to it.
-    /// expression to it.
+    /// TODO(crbug.com/tint/1671): DEPRECATED: Replaced with createTextureSampleBaseClampToEdge().
     /// @param expr the call expression being transformed
     /// @param syms the expanded symbols to be used in the new call
     /// @returns a call expression to textureSampleExternal
-    const ast::CallExpression* createTexSmpExt(const ast::CallExpression* expr,
+    const ast::CallExpression* createTextureSampleLevel(const ast::CallExpression* expr,
-                                               NewBindingSymbols syms) {
+                                                        NewBindingSymbols syms) {
         const ast::Expression* plane_0_binding_param = ctx.Clone(expr->args[0]);
 
         if (expr->args.Length() != 3) {
@@ -391,9 +423,7 @@ struct MultiplanarExternalTexture::State {
                     b.Param("params", b.ty.type_name(params_struct_sym)),
                 },
                 b.ty.vec4(b.ty.f32()),
-                utils::Vector{
+                buildTextureBuiltinBody(sem::BuiltinType::kTextureSampleLevel));
-                    createTexFnExtStatementList(sem::BuiltinType::kTextureSampleLevel),
-                });
         }
 
         const ast::IdentifierExpression* exp = b.Expr(texture_sample_external_sym);
@@ -406,13 +436,60 @@ struct MultiplanarExternalTexture::State {
                            });
     }
 
-    /// Creates the textureLoadExternal function if needed and returns a call
+    /// Creates the textureSampleExternal function if needed and returns a call expression to it.
-    /// expression to it.
+    /// @param expr the call expression being transformed
+    /// @param syms the expanded symbols to be used in the new call
+    /// @returns a call expression to textureSampleExternal
+    const ast::CallExpression* createTextureSampleBaseClampToEdge(const ast::CallExpression* expr,
+                                                                  NewBindingSymbols syms) {
+        const ast::Expression* plane_0_binding_param = ctx.Clone(expr->args[0]);
+
+        if (expr->args.Length() != 3) {
+            TINT_ICE(Transform, b.Diagnostics())
+                << "expected textureSampleBaseClampToEdge call with a "
+                   "texture_external to have 3 parameters, found "
+                << expr->args.Length() << " parameters";
+        }
+
+        // TextureSampleExternal calls the gammaCorrection function, so ensure it
+        // exists.
+        if (!gamma_correction_sym.IsValid()) {
+            createGammaCorrectionFn();
+        }
+
+        if (!texture_sample_external_sym.IsValid()) {
+            texture_sample_external_sym = b.Symbols().New("textureSampleExternal");
+
+            // Emit the textureSampleExternal function.
+            b.Func(
+                texture_sample_external_sym,
+                utils::Vector{
+                    b.Param("plane0", b.ty.sampled_texture(ast::TextureDimension::k2d, b.ty.f32())),
+                    b.Param("plane1", b.ty.sampled_texture(ast::TextureDimension::k2d, b.ty.f32())),
+                    b.Param("smp", b.ty.sampler(ast::SamplerKind::kSampler)),
+                    b.Param("coord", b.ty.vec2(b.ty.f32())),
+                    b.Param("params", b.ty.type_name(params_struct_sym)),
+                },
+                b.ty.vec4(b.ty.f32()),
+                buildTextureBuiltinBody(sem::BuiltinType::kTextureSampleBaseClampToEdge));
+        }
+
+        const ast::IdentifierExpression* exp = b.Expr(texture_sample_external_sym);
+        return b.Call(exp, utils::Vector{
+                               plane_0_binding_param,
+                               b.Expr(syms.plane_1),
+                               ctx.Clone(expr->args[1]),
+                               ctx.Clone(expr->args[2]),
+                               b.Expr(syms.params),
+                           });
+    }
+
+    /// Creates the textureLoadExternal function if needed and returns a call expression to it.
     /// @param expr the call expression being transformed
     /// @param syms the expanded symbols to be used in the new call
     /// @returns a call expression to textureLoadExternal
-    const ast::CallExpression* createTexLdExt(const ast::CallExpression* expr,
+    const ast::CallExpression* createTextureLoad(const ast::CallExpression* expr,
-                                              NewBindingSymbols syms) {
+                                                 NewBindingSymbols syms) {
         const ast::Expression* plane_0_binding_param = ctx.Clone(expr->args[0]);
 
         if (expr->args.Length() != 2) {
@@ -440,10 +517,8 @@ struct MultiplanarExternalTexture::State {
                     b.Param("coord", b.ty.vec2(b.ty.i32())),
                     b.Param("params", b.ty.type_name(params_struct_sym)),
                 },
-                b.ty.vec4(b.ty.f32()),
+                b.ty.vec4(b.ty.f32()),  //
-                utils::Vector{
+                buildTextureBuiltinBody(sem::BuiltinType::kTextureLoad));
-                    createTexFnExtStatementList(sem::BuiltinType::kTextureLoad),
-                });
         }
 
         return b.Call(texture_load_external_sym, plane_0_binding_param, syms.plane_1,
@@ -469,13 +544,11 @@ bool MultiplanarExternalTexture::ShouldRun(const Program* program, const DataMap
     return false;
 }
 
-// Within this transform, an instance of a texture_external binding is unpacked
+// Within this transform, an instance of a texture_external binding is unpacked into two
-// into two texture_2d<f32> bindings representing two possible planes of a
+// texture_2d<f32> bindings representing two possible planes of a single texture and a uniform
-// single texture and a uniform buffer binding representing a struct of
+// buffer binding representing a struct of parameters. Calls to texture builtins that contain a
-// parameters. Calls to textureLoad or textureSampleLevel that contain a
+// texture_external parameter will be transformed into a newly generated version of the function,
-// texture_external parameter will be transformed into a newly generated version
+// which can perform the desired operation on a single RGBA plane or on separate Y and UV planes.
-// of the function, which can perform the desired operation on a single RGBA
-// plane or on separate Y and UV planes.
 void MultiplanarExternalTexture::Run(CloneContext& ctx, const DataMap& inputs, DataMap&) const {
     auto* new_binding_points = inputs.Get<NewBindingPoints>();
 

src/tint/transform/multiplanar_external_texture_test.cc

@@ -49,9 +49,8 @@ fn f(ext_tex : texture_external) {}
     EXPECT_TRUE(ShouldRun<MultiplanarExternalTexture>(src));
 }
 
-// Running the transform without passing in data for the new bindings should
+// Running the transform without passing in data for the new bindings should result in an error.
-// result in an error.
+TEST_F(MultiplanarExternalTextureTest, ErrorNoPassedData_SampleLevel) {
-TEST_F(MultiplanarExternalTextureTest, ErrorNoPassedData) {
     auto* src = R"(
 @group(0) @binding(0) var s : sampler;
 @group(0) @binding(1) var ext_tex : texture_external;
@@ -68,8 +67,26 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
+// Running the transform without passing in data for the new bindings should result in an error.
+TEST_F(MultiplanarExternalTextureTest, ErrorNoPassedData_SampleBaseClampToEdge) {
+    auto* src = R"(
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy);
+}
+)";
+    auto* expect =
+        R"(error: missing new binding point data for tint::transform::MultiplanarExternalTexture)";
+
+    auto got = Run<MultiplanarExternalTexture>(src);
+    EXPECT_EQ(expect, str(got));
+}
+
 // Running the transform with incorrect binding data should result in an error.
-TEST_F(MultiplanarExternalTextureTest, ErrorIncorrectBindingPont) {
+TEST_F(MultiplanarExternalTextureTest, ErrorIncorrectBindingPont_SampleLevel) {
     auto* src = R"(
 @group(0) @binding(0) var s : sampler;
 @group(0) @binding(1) var ext_tex : texture_external;
@@ -91,6 +108,29 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
+// Running the transform with incorrect binding data should result in an error.
+TEST_F(MultiplanarExternalTextureTest, ErrorIncorrectBindingPont_SampleBaseClampToEdge) {
+    auto* src = R"(
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy);
+}
+)";
+
+    auto* expect = R"(error: missing new binding points for texture_external at binding {0,1})";
+
+    DataMap data;
+    // This bindings map specifies 0,0 as the location of the texture_external,
+    // which is incorrect.
+    data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+        MultiplanarExternalTexture::BindingsMap{{{0, 0}, {{0, 1}, {0, 2}}}});
+    auto got = Run<MultiplanarExternalTexture>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with a textureDimensions call.
 TEST_F(MultiplanarExternalTextureTest, Dimensions) {
     auto* src = R"(
@@ -277,6 +317,88 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
+// Test that the transform works with a textureSampleBaseClampToEdge call.
+TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleBaseClampToEdge) {
+    auto* src = R"(
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy);
+}
+)";
+
+    auto* expect = R"(
+struct GammaTransferParams {
+  G : f32,
+  A : f32,
+  B : f32,
+  C : f32,
+  D : f32,
+  E : f32,
+  F : f32,
+  padding : u32,
+}
+
+struct ExternalTextureParams {
+  numPlanes : u32,
+  doYuvToRgbConversionOnly : u32,
+  yuvToRgbConversionMatrix : mat3x4<f32>,
+  gammaDecodeParams : GammaTransferParams,
+  gammaEncodeParams : GammaTransferParams,
+  gamutConversionMatrix : mat3x3<f32>,
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@group(0) @binding(0) var s : sampler;
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+
+fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
+  let cond = (abs(v) < vec3<f32>(params.D));
+  let t = (sign(v) * ((params.C * abs(v)) + params.F));
+  let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params);
+}
+)";
+
+    DataMap data;
+    data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+        MultiplanarExternalTexture::BindingsMap{{{0, 1}, {{0, 2}, {0, 3}}}});
+    auto got = Run<MultiplanarExternalTexture>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
 // Test that the transform works with a textureSampleLevel call.
 TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleLevel_OutOfOrder) {
     auto* src = R"(
@@ -353,6 +475,88 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
+// Test that the transform works with a textureSampleBaseClampToEdge call.
+TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleBaseClampToEdge_OutOfOrder) {
+    auto* src = R"(
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy);
+}
+
+@group(0) @binding(1) var ext_tex : texture_external;
+@group(0) @binding(0) var s : sampler;
+)";
+
+    auto* expect = R"(
+struct GammaTransferParams {
+  G : f32,
+  A : f32,
+  B : f32,
+  C : f32,
+  D : f32,
+  E : f32,
+  F : f32,
+  padding : u32,
+}
+
+struct ExternalTextureParams {
+  numPlanes : u32,
+  doYuvToRgbConversionOnly : u32,
+  yuvToRgbConversionMatrix : mat3x4<f32>,
+  gammaDecodeParams : GammaTransferParams,
+  gammaEncodeParams : GammaTransferParams,
+  gamutConversionMatrix : mat3x3<f32>,
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
+  let cond = (abs(v) < vec3<f32>(params.D));
+  let t = (sign(v) * ((params.C * abs(v)) + params.F));
+  let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params);
+}
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(0) var s : sampler;
+)";
+
+    DataMap data;
+    data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+        MultiplanarExternalTexture::BindingsMap{{{0, 1}, {{0, 2}, {0, 3}}}});
+    auto got = Run<MultiplanarExternalTexture>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with a textureLoad call.
 TEST_F(MultiplanarExternalTextureTest, BasicTextureLoad) {
     auto* src = R"(
@@ -499,8 +703,7 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
-// Tests that the transform works with both a textureSampleLevel and textureLoad
+// Tests that the transform works with both a textureSampleLevel and textureLoad call.
-// call.
 TEST_F(MultiplanarExternalTextureTest, TextureSampleAndTextureLoad) {
     auto* src = R"(
 @group(0) @binding(0) var s : sampler;
@@ -591,8 +794,104 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
-// Tests that the transform works with both a textureSampleLevel and textureLoad
+// Tests that the transform works with both a textureSampleBaseClampToEdge and textureLoad call.
-// call.
+TEST_F(MultiplanarExternalTextureTest, TextureSampleBaseClampToEdgeAndTextureLoad) {
+    auto* src = R"(
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy) + textureLoad(ext_tex, vec2<i32>(1, 1));
+}
+)";
+
+    auto* expect = R"(
+struct GammaTransferParams {
+  G : f32,
+  A : f32,
+  B : f32,
+  C : f32,
+  D : f32,
+  E : f32,
+  F : f32,
+  padding : u32,
+}
+
+struct ExternalTextureParams {
+  numPlanes : u32,
+  doYuvToRgbConversionOnly : u32,
+  yuvToRgbConversionMatrix : mat3x4<f32>,
+  gammaDecodeParams : GammaTransferParams,
+  gammaEncodeParams : GammaTransferParams,
+  gamutConversionMatrix : mat3x3<f32>,
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@group(0) @binding(0) var s : sampler;
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+
+fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
+  let cond = (abs(v) < vec3<f32>(params.D));
+  let t = (sign(v) * ((params.C * abs(v)) + params.F));
+  let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord : vec2<i32>, params : ExternalTextureParams) -> vec4<f32> {
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureLoad(plane0, coord, 0i).rgb;
+  } else {
+    color = (vec4<f32>(textureLoad(plane0, coord, 0i).r, textureLoad(plane1, coord, 0i).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return (textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params));
+}
+)";
+
+    DataMap data;
+    data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+        MultiplanarExternalTexture::BindingsMap{{{0, 1}, {{0, 2}, {0, 3}}}});
+    auto got = Run<MultiplanarExternalTexture>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
+// Tests that the transform works with both a textureSampleLevel and textureLoad call.
 TEST_F(MultiplanarExternalTextureTest, TextureSampleAndTextureLoad_OutOfOrder) {
     auto* src = R"(
 @fragment
@@ -683,6 +982,103 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the transform works with both a textureSampleBaseClampToEdge and textureLoad call.
+TEST_F(MultiplanarExternalTextureTest, TextureSampleBaseClampToEdgeAndTextureLoad_OutOfOrder) {
+    auto* src = R"(
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy) + textureLoad(ext_tex, vec2<i32>(1, 1));
+}
+
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+)";
+
+    auto* expect = R"(
+struct GammaTransferParams {
+  G : f32,
+  A : f32,
+  B : f32,
+  C : f32,
+  D : f32,
+  E : f32,
+  F : f32,
+  padding : u32,
+}
+
+struct ExternalTextureParams {
+  numPlanes : u32,
+  doYuvToRgbConversionOnly : u32,
+  yuvToRgbConversionMatrix : mat3x4<f32>,
+  gammaDecodeParams : GammaTransferParams,
+  gammaEncodeParams : GammaTransferParams,
+  gamutConversionMatrix : mat3x3<f32>,
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
+  let cond = (abs(v) < vec3<f32>(params.D));
+  let t = (sign(v) * ((params.C * abs(v)) + params.F));
+  let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord : vec2<i32>, params : ExternalTextureParams) -> vec4<f32> {
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureLoad(plane0, coord, 0i).rgb;
+  } else {
+    color = (vec4<f32>(textureLoad(plane0, coord, 0i).r, textureLoad(plane1, coord, 0i).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return (textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params));
+}
+
+@group(0) @binding(0) var s : sampler;
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+)";
+
+    DataMap data;
+    data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+        MultiplanarExternalTexture::BindingsMap{{{0, 1}, {{0, 2}, {0, 3}}}});
+    auto got = Run<MultiplanarExternalTexture>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with many instances of texture_external.
 TEST_F(MultiplanarExternalTextureTest, ManyTextureSampleLevel) {
     auto* src = R"(
@@ -784,12 +1180,122 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
     EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the transform works with many instances of texture_external.
+TEST_F(MultiplanarExternalTextureTest, ManyTextureSampleBaseClampToEdge) {
+    auto* src = R"(
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+@group(0) @binding(2) var ext_tex_1 : texture_external;
+@group(0) @binding(3) var ext_tex_2 : texture_external;
+@group(1) @binding(0) var ext_tex_3 : texture_external;
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleBaseClampToEdge(ext_tex, s, coord.xy) +
+         textureSampleBaseClampToEdge(ext_tex_1, s, coord.xy) +
+         textureSampleBaseClampToEdge(ext_tex_2, s, coord.xy) +
+         textureSampleBaseClampToEdge(ext_tex_3, s, coord.xy);
+}
+)";
+
+    auto* expect = R"(
+struct GammaTransferParams {
+  G : f32,
+  A : f32,
+  B : f32,
+  C : f32,
+  D : f32,
+  E : f32,
+  F : f32,
+  padding : u32,
+}
+
+struct ExternalTextureParams {
+  numPlanes : u32,
+  doYuvToRgbConversionOnly : u32,
+  yuvToRgbConversionMatrix : mat3x4<f32>,
+  gammaDecodeParams : GammaTransferParams,
+  gammaEncodeParams : GammaTransferParams,
+  gamutConversionMatrix : mat3x3<f32>,
+}
+
+@group(0) @binding(4) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(5) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@group(0) @binding(6) var ext_tex_plane_1_1 : texture_2d<f32>;
+
+@group(0) @binding(7) var<uniform> ext_tex_params_1 : ExternalTextureParams;
+
+@group(0) @binding(8) var ext_tex_plane_1_2 : texture_2d<f32>;
+
+@group(0) @binding(9) var<uniform> ext_tex_params_2 : ExternalTextureParams;
+
+@group(1) @binding(1) var ext_tex_plane_1_3 : texture_2d<f32>;
+
+@group(1) @binding(2) var<uniform> ext_tex_params_3 : ExternalTextureParams;
+
+@group(0) @binding(0) var s : sampler;
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(2) var ext_tex_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var ext_tex_2 : texture_2d<f32>;
+
+@group(1) @binding(0) var ext_tex_3 : texture_2d<f32>;
+
+fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
+  let cond = (abs(v) < vec3<f32>(params.D));
+  let t = (sign(v) * ((params.C * abs(v)) + params.F));
+  let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1u)) {
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4<f32>(color, 1.0f);
+}
+
+@fragment
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return (((textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureSampleExternal(ext_tex_1, ext_tex_plane_1_1, s, coord.xy, ext_tex_params_1)) + textureSampleExternal(ext_tex_2, ext_tex_plane_1_2, s, coord.xy, ext_tex_params_2)) + textureSampleExternal(ext_tex_3, ext_tex_plane_1_3, s, coord.xy, ext_tex_params_3));
+}
+)";
+
+    DataMap data;
+    data.Add<MultiplanarExternalTexture::NewBindingPoints>(MultiplanarExternalTexture::BindingsMap{
+        {{0, 1}, {{0, 4}, {0, 5}}},
+        {{0, 2}, {{0, 6}, {0, 7}}},
+        {{0, 3}, {{0, 8}, {0, 9}}},
+        {{1, 0}, {{1, 1}, {1, 2}}},
+    });
+    auto got = Run<MultiplanarExternalTexture>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the texture_external passed as a function parameter produces the
 // correct output.
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParam) {
     auto* src = R"(
 fn f(t : texture_external, s : sampler) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : texture_external;
@@ -834,11 +1340,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -879,7 +1391,7 @@ fn main() {
 }
 
 fn f(t : texture_external, s : sampler) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : texture_external;
@@ -924,11 +1436,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -959,7 +1477,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsSecondParam) {
     auto* src = R"(
 fn f(s : sampler, t : texture_external) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : texture_external;
@@ -1004,11 +1522,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -1044,8 +1568,8 @@ fn main() {
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamMultiple) {
     auto* src = R"(
 fn f(t : texture_external, s : sampler, t2 : texture_external) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
-  textureSampleLevel(t2, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t2, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : texture_external;
@@ -1095,11 +1619,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -1144,8 +1674,8 @@ fn main() {
 }
 
 fn f(t : texture_external, s : sampler, t2 : texture_external) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
-  textureSampleLevel(t2, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t2, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : texture_external;
@@ -1196,11 +1726,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -1235,7 +1771,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested) {
     auto* src = R"(
 fn nested(t : texture_external, s : sampler) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 fn f(t : texture_external, s : sampler) {
@@ -1284,11 +1820,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -1323,12 +1865,12 @@ fn main() {
     EXPECT_EQ(expect, str(got));
 }
 
-// Tests that the texture_external passed to as a parameter to multiple
+// Tests that the texture_external passed to as a parameter to multiple functions produces the
-// functions produces the correct output.
+// correct output.
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested_OutOfOrder) {
     auto* src = R"(
 fn nested(t : texture_external, s : sampler) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 fn f(t : texture_external, s : sampler) {
@@ -1377,11 +1919,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -1464,7 +2012,7 @@ TEST_F(MultiplanarExternalTextureTest, ExternalTextureAlias) {
 type ET = texture_external;
 
 fn f(t : ET, s : sampler) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : ET;
@@ -1511,11 +2059,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -1555,7 +2109,7 @@ fn main() {
 }
 
 fn f(t : ET, s : sampler) {
-  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleBaseClampToEdge(t, s, vec2<f32>(1.0, 2.0));
 }
 
 @group(0) @binding(0) var ext_tex : ET;
@@ -1602,11 +2156,17 @@ fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
 }
 
 fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  let plane0_dims = vec2<f32>(textureDimensions(plane0, 0));
+  let plane0_half_texel = (vec2<f32>(0.5) / plane0_dims);
+  let plane0_clamped = clamp(coord, plane0_half_texel, (1 - plane0_half_texel));
+  let plane1_dims = vec2<f32>(textureDimensions(plane1, 0));
+  let plane1_half_texel = (vec2<f32>(0.5) / plane1_dims);
+  let plane1_clamped = clamp(coord, plane1_half_texel, (1 - plane1_half_texel));
   var color : vec3<f32>;
   if ((params.numPlanes == 1u)) {
-    color = textureSampleLevel(plane0, smp, coord, 0.0f).rgb;
+    color = textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).rgb;
   } else {
-    color = (vec4<f32>(textureSampleLevel(plane0, smp, coord, 0.0f).r, textureSampleLevel(plane1, smp, coord, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4<f32>(textureSampleLevel(plane0, smp, plane0_clamped, 0.0f).r, textureSampleLevel(plane1, smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);

src/tint/writer/glsl/generator_impl.cc

@@ -192,6 +192,7 @@ SanitizedResult Sanitize(const Program* in,
         polyfills.first_trailing_bit = true;
         polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
         polyfills.saturate = true;
+        polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

src/tint/writer/hlsl/generator_impl.cc

@@ -175,6 +175,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.first_leading_bit = true;
         polyfills.first_trailing_bit = true;
         polyfills.insert_bits = transform::BuiltinPolyfill::Level::kFull;
+        polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

src/tint/writer/msl/generator_impl.cc

@@ -172,6 +172,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.first_leading_bit = true;
         polyfills.first_trailing_bit = true;
         polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+        polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

src/tint/writer/spirv/generator_impl.cc

@@ -58,6 +58,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.first_trailing_bit = true;
         polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
         polyfills.saturate = true;
+        polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+@group(1) @binding(0) var arg_0: texture_external;
+@group(1) @binding(1) var arg_1: sampler;
+
+// fn textureSampleBaseClampToEdge(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureSampleBaseClampToEdge_7c04e6() {
+  var res: vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, vec2<f32>());
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_7c04e6();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.dxc.hlsl

@@ -0,0 +1,122 @@
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+};
+
+Texture2D<float4> ext_tex_plane_1 : register(t2, space1);
+cbuffer cbuffer_ext_tex_params : register(b3, space1) {
+  uint4 ext_tex_params[11];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float3 gammaCorrection(float3 v, GammaTransferParams params) {
+  const bool3 cond = (abs(v) < float3((params.D).xxx));
+  const float3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  const float3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E));
+  return (cond ? t : f);
+}
+
+float4 textureSampleExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, SamplerState smp, float2 coord, ExternalTextureParams params) {
+  int3 tint_tmp;
+  plane0.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 plane0_dims = float2(tint_tmp.xy);
+  const float2 plane0_half_texel = ((0.5f).xx / plane0_dims);
+  const float2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  int3 tint_tmp_1;
+  plane1.GetDimensions(0, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
+  const float2 plane1_dims = float2(tint_tmp_1.xy);
+  const float2 plane1_half_texel = ((0.5f).xx / plane1_dims);
+  const float2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  float3 color = float3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = plane0.SampleLevel(smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.SampleLevel(smp, plane0_clamped, 0.0f).r, plane1.SampleLevel(smp, plane1_clamped, 0.0f).rg, 1.0f));
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = mul(color, params.gamutConversionMatrix);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return float4(color, 1.0f);
+}
+
+float3x4 tint_symbol_3(uint4 buffer[11], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  return float3x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]));
+}
+
+GammaTransferParams tint_symbol_5(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_3 = ((offset + 0u)) / 4;
+  const uint scalar_offset_4 = ((offset + 4u)) / 4;
+  const uint scalar_offset_5 = ((offset + 8u)) / 4;
+  const uint scalar_offset_6 = ((offset + 12u)) / 4;
+  const uint scalar_offset_7 = ((offset + 16u)) / 4;
+  const uint scalar_offset_8 = ((offset + 20u)) / 4;
+  const uint scalar_offset_9 = ((offset + 24u)) / 4;
+  const uint scalar_offset_10 = ((offset + 28u)) / 4;
+  const GammaTransferParams tint_symbol_9 = {asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(buffer[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(buffer[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(buffer[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(buffer[scalar_offset_9 / 4][scalar_offset_9 % 4]), buffer[scalar_offset_10 / 4][scalar_offset_10 % 4]};
+  return tint_symbol_9;
+}
+
+float3x3 tint_symbol_7(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_11 = ((offset + 0u)) / 4;
+  const uint scalar_offset_12 = ((offset + 16u)) / 4;
+  const uint scalar_offset_13 = ((offset + 32u)) / 4;
+  return float3x3(asfloat(buffer[scalar_offset_11 / 4].xyz), asfloat(buffer[scalar_offset_12 / 4].xyz), asfloat(buffer[scalar_offset_13 / 4].xyz));
+}
+
+ExternalTextureParams tint_symbol_1(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_14 = ((offset + 0u)) / 4;
+  const uint scalar_offset_15 = ((offset + 4u)) / 4;
+  const ExternalTextureParams tint_symbol_10 = {buffer[scalar_offset_14 / 4][scalar_offset_14 % 4], buffer[scalar_offset_15 / 4][scalar_offset_15 % 4], tint_symbol_3(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 64u)), tint_symbol_5(buffer, (offset + 96u)), tint_symbol_7(buffer, (offset + 128u))};
+  return tint_symbol_10;
+}
+
+void textureSampleBaseClampToEdge_7c04e6() {
+  float4 res = textureSampleExternal(arg_0, ext_tex_plane_1, arg_1, (0.0f).xx, tint_symbol_1(ext_tex_params, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.fxc.hlsl

@@ -0,0 +1,122 @@
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+};
+
+Texture2D<float4> ext_tex_plane_1 : register(t2, space1);
+cbuffer cbuffer_ext_tex_params : register(b3, space1) {
+  uint4 ext_tex_params[11];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float3 gammaCorrection(float3 v, GammaTransferParams params) {
+  const bool3 cond = (abs(v) < float3((params.D).xxx));
+  const float3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  const float3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E));
+  return (cond ? t : f);
+}
+
+float4 textureSampleExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, SamplerState smp, float2 coord, ExternalTextureParams params) {
+  int3 tint_tmp;
+  plane0.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 plane0_dims = float2(tint_tmp.xy);
+  const float2 plane0_half_texel = ((0.5f).xx / plane0_dims);
+  const float2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  int3 tint_tmp_1;
+  plane1.GetDimensions(0, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
+  const float2 plane1_dims = float2(tint_tmp_1.xy);
+  const float2 plane1_half_texel = ((0.5f).xx / plane1_dims);
+  const float2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  float3 color = float3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = plane0.SampleLevel(smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.SampleLevel(smp, plane0_clamped, 0.0f).r, plane1.SampleLevel(smp, plane1_clamped, 0.0f).rg, 1.0f));
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = mul(color, params.gamutConversionMatrix);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return float4(color, 1.0f);
+}
+
+float3x4 tint_symbol_3(uint4 buffer[11], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  return float3x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]));
+}
+
+GammaTransferParams tint_symbol_5(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_3 = ((offset + 0u)) / 4;
+  const uint scalar_offset_4 = ((offset + 4u)) / 4;
+  const uint scalar_offset_5 = ((offset + 8u)) / 4;
+  const uint scalar_offset_6 = ((offset + 12u)) / 4;
+  const uint scalar_offset_7 = ((offset + 16u)) / 4;
+  const uint scalar_offset_8 = ((offset + 20u)) / 4;
+  const uint scalar_offset_9 = ((offset + 24u)) / 4;
+  const uint scalar_offset_10 = ((offset + 28u)) / 4;
+  const GammaTransferParams tint_symbol_9 = {asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(buffer[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(buffer[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(buffer[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(buffer[scalar_offset_9 / 4][scalar_offset_9 % 4]), buffer[scalar_offset_10 / 4][scalar_offset_10 % 4]};
+  return tint_symbol_9;
+}
+
+float3x3 tint_symbol_7(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_11 = ((offset + 0u)) / 4;
+  const uint scalar_offset_12 = ((offset + 16u)) / 4;
+  const uint scalar_offset_13 = ((offset + 32u)) / 4;
+  return float3x3(asfloat(buffer[scalar_offset_11 / 4].xyz), asfloat(buffer[scalar_offset_12 / 4].xyz), asfloat(buffer[scalar_offset_13 / 4].xyz));
+}
+
+ExternalTextureParams tint_symbol_1(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_14 = ((offset + 0u)) / 4;
+  const uint scalar_offset_15 = ((offset + 4u)) / 4;
+  const ExternalTextureParams tint_symbol_10 = {buffer[scalar_offset_14 / 4][scalar_offset_14 % 4], buffer[scalar_offset_15 / 4][scalar_offset_15 % 4], tint_symbol_3(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 64u)), tint_symbol_5(buffer, (offset + 96u)), tint_symbol_7(buffer, (offset + 128u))};
+  return tint_symbol_10;
+}
+
+void textureSampleBaseClampToEdge_7c04e6() {
+  float4 res = textureSampleExternal(arg_0, ext_tex_plane_1, arg_1, (0.0f).xx, tint_symbol_1(ext_tex_params, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.glsl

@@ -0,0 +1,220 @@
+#version 310 es
+
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+};
+
+layout(binding = 3, std140) uniform ext_tex_params_block_ubo {
+  ExternalTextureParams inner;
+} ext_tex_params;
+
+vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
+  bvec3 cond = lessThan(abs(v), vec3(params.D));
+  vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
+  return mix(f, t, cond);
+}
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+  vec2 plane0_dims = vec2(textureSize(plane0_1, 0));
+  vec2 plane0_half_texel = (vec2(0.5f) / plane0_dims);
+  vec2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  vec2 plane1_dims = vec2(textureSize(plane1_1, 0));
+  vec2 plane1_half_texel = (vec2(0.5f) / plane1_dims);
+  vec2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  vec3 color = vec3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = textureLod(plane0_smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4(textureLod(plane0_smp, plane0_clamped, 0.0f).r, textureLod(plane1_smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4(color, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleBaseClampToEdge_7c04e6() {
+  vec4 res = textureSampleExternal(arg_0_1, ext_tex_plane_1_1, arg_0_arg_1, ext_tex_plane_1_arg_1, vec2(0.0f), ext_tex_params.inner);
+}
+
+vec4 vertex_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+};
+
+layout(binding = 3, std140) uniform ext_tex_params_block_ubo {
+  ExternalTextureParams inner;
+} ext_tex_params;
+
+vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
+  bvec3 cond = lessThan(abs(v), vec3(params.D));
+  vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
+  return mix(f, t, cond);
+}
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+  vec2 plane0_dims = vec2(textureSize(plane0_1, 0));
+  vec2 plane0_half_texel = (vec2(0.5f) / plane0_dims);
+  vec2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  vec2 plane1_dims = vec2(textureSize(plane1_1, 0));
+  vec2 plane1_half_texel = (vec2(0.5f) / plane1_dims);
+  vec2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  vec3 color = vec3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = textureLod(plane0_smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4(textureLod(plane0_smp, plane0_clamped, 0.0f).r, textureLod(plane1_smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4(color, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleBaseClampToEdge_7c04e6() {
+  vec4 res = textureSampleExternal(arg_0_1, ext_tex_plane_1_1, arg_0_arg_1, ext_tex_plane_1_arg_1, vec2(0.0f), ext_tex_params.inner);
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+};
+
+layout(binding = 3, std140) uniform ext_tex_params_block_ubo {
+  ExternalTextureParams inner;
+} ext_tex_params;
+
+vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
+  bvec3 cond = lessThan(abs(v), vec3(params.D));
+  vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
+  return mix(f, t, cond);
+}
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+  vec2 plane0_dims = vec2(textureSize(plane0_1, 0));
+  vec2 plane0_half_texel = (vec2(0.5f) / plane0_dims);
+  vec2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  vec2 plane1_dims = vec2(textureSize(plane1_1, 0));
+  vec2 plane1_half_texel = (vec2(0.5f) / plane1_dims);
+  vec2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  vec3 color = vec3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = textureLod(plane0_smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4(textureLod(plane0_smp, plane0_clamped, 0.0f).r, textureLod(plane1_smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4(color, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleBaseClampToEdge_7c04e6() {
+  vec4 res = textureSampleExternal(arg_0_1, ext_tex_plane_1_1, arg_0_arg_1, ext_tex_plane_1_arg_1, vec2(0.0f), ext_tex_params.inner);
+}
+
+void compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.msl

@@ -0,0 +1,95 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+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 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 ExternalTextureParams {
+  /* 0x0000 */ uint numPlanes;
+  /* 0x0004 */ uint doYuvToRgbConversionOnly;
+  /* 0x0008 */ tint_array<int8_t, 8> tint_pad;
+  /* 0x0010 */ float3x4 yuvToRgbConversionMatrix;
+  /* 0x0040 */ GammaTransferParams gammaDecodeParams;
+  /* 0x0060 */ GammaTransferParams gammaEncodeParams;
+  /* 0x0080 */ float3x3 gamutConversionMatrix;
+};
+
+float3 gammaCorrection(float3 v, GammaTransferParams params) {
+  bool3 const cond = (fabs(v) < float3(params.D));
+  float3 const t = (sign(v) * ((params.C * fabs(v)) + params.F));
+  float3 const f = (sign(v) * (pow(((params.A * fabs(v)) + params.B), float3(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+float4 textureSampleExternal(texture2d<float, access::sample> plane0, texture2d<float, access::sample> plane1, sampler smp, float2 coord, ExternalTextureParams params) {
+  float2 const plane0_dims = float2(int2(plane0.get_width(0), plane0.get_height(0)));
+  float2 const plane0_half_texel = (float2(0.5f) / plane0_dims);
+  float2 const plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  float2 const plane1_dims = float2(int2(plane1.get_width(0), plane1.get_height(0)));
+  float2 const plane1_half_texel = (float2(0.5f) / plane1_dims);
+  float2 const plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  float3 color = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    color = float4(plane0.sample(smp, plane0_clamped, level(0.0f))).rgb;
+  } else {
+    color = (float4(plane0.sample(smp, plane0_clamped, level(0.0f))[0], float4(plane1.sample(smp, plane1_clamped, level(0.0f))).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return float4(color, 1.0f);
+}
+
+void textureSampleBaseClampToEdge_7c04e6(texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2, sampler tint_symbol_3, const constant ExternalTextureParams* const tint_symbol_4) {
+  float4 res = textureSampleExternal(tint_symbol_1, tint_symbol_2, tint_symbol_3, float2(0.0f), *(tint_symbol_4));
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_5, texture2d<float, access::sample> tint_symbol_6, sampler tint_symbol_7, const constant ExternalTextureParams* const tint_symbol_8) {
+  textureSampleBaseClampToEdge_7c04e6(tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_9 [[texture(0)]], texture2d<float, access::sample> tint_symbol_10 [[texture(1)]], sampler tint_symbol_11 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_12 [[buffer(2)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_13 [[texture(0)]], texture2d<float, access::sample> tint_symbol_14 [[texture(1)]], sampler tint_symbol_15 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_16 [[buffer(2)]]) {
+  textureSampleBaseClampToEdge_7c04e6(tint_symbol_13, tint_symbol_14, tint_symbol_15, tint_symbol_16);
+  return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_17 [[texture(0)]], texture2d<float, access::sample> tint_symbol_18 [[texture(1)]], sampler tint_symbol_19 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_20 [[buffer(2)]]) {
+  textureSampleBaseClampToEdge_7c04e6(tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
+  return;
+}
+

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.spvasm

@@ -0,0 +1,276 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 164
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %31 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %ext_tex_plane_1 "ext_tex_plane_1"
+               OpName %ext_tex_params_block "ext_tex_params_block"
+               OpMemberName %ext_tex_params_block 0 "inner"
+               OpName %ExternalTextureParams "ExternalTextureParams"
+               OpMemberName %ExternalTextureParams 0 "numPlanes"
+               OpMemberName %ExternalTextureParams 1 "doYuvToRgbConversionOnly"
+               OpMemberName %ExternalTextureParams 2 "yuvToRgbConversionMatrix"
+               OpMemberName %ExternalTextureParams 3 "gammaDecodeParams"
+               OpName %GammaTransferParams "GammaTransferParams"
+               OpMemberName %GammaTransferParams 0 "G"
+               OpMemberName %GammaTransferParams 1 "A"
+               OpMemberName %GammaTransferParams 2 "B"
+               OpMemberName %GammaTransferParams 3 "C"
+               OpMemberName %GammaTransferParams 4 "D"
+               OpMemberName %GammaTransferParams 5 "E"
+               OpMemberName %GammaTransferParams 6 "F"
+               OpMemberName %GammaTransferParams 7 "padding"
+               OpMemberName %ExternalTextureParams 4 "gammaEncodeParams"
+               OpMemberName %ExternalTextureParams 5 "gamutConversionMatrix"
+               OpName %ext_tex_params "ext_tex_params"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %gammaCorrection "gammaCorrection"
+               OpName %v "v"
+               OpName %params "params"
+               OpName %textureSampleExternal "textureSampleExternal"
+               OpName %plane0 "plane0"
+               OpName %plane1 "plane1"
+               OpName %smp "smp"
+               OpName %coord "coord"
+               OpName %params_0 "params"
+               OpName %color "color"
+               OpName %textureSampleBaseClampToEdge_7c04e6 "textureSampleBaseClampToEdge_7c04e6"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %ext_tex_plane_1 DescriptorSet 1
+               OpDecorate %ext_tex_plane_1 Binding 2
+               OpDecorate %ext_tex_params_block Block
+               OpMemberDecorate %ext_tex_params_block 0 Offset 0
+               OpMemberDecorate %ExternalTextureParams 0 Offset 0
+               OpMemberDecorate %ExternalTextureParams 1 Offset 4
+               OpMemberDecorate %ExternalTextureParams 2 Offset 16
+               OpMemberDecorate %ExternalTextureParams 2 ColMajor
+               OpMemberDecorate %ExternalTextureParams 2 MatrixStride 16
+               OpMemberDecorate %ExternalTextureParams 3 Offset 64
+               OpMemberDecorate %GammaTransferParams 0 Offset 0
+               OpMemberDecorate %GammaTransferParams 1 Offset 4
+               OpMemberDecorate %GammaTransferParams 2 Offset 8
+               OpMemberDecorate %GammaTransferParams 3 Offset 12
+               OpMemberDecorate %GammaTransferParams 4 Offset 16
+               OpMemberDecorate %GammaTransferParams 5 Offset 20
+               OpMemberDecorate %GammaTransferParams 6 Offset 24
+               OpMemberDecorate %GammaTransferParams 7 Offset 28
+               OpMemberDecorate %ExternalTextureParams 4 Offset 96
+               OpMemberDecorate %ExternalTextureParams 5 Offset 128
+               OpMemberDecorate %ExternalTextureParams 5 ColMajor
+               OpMemberDecorate %ExternalTextureParams 5 MatrixStride 16
+               OpDecorate %ext_tex_params NonWritable
+               OpDecorate %ext_tex_params DescriptorSet 1
+               OpDecorate %ext_tex_params Binding 3
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+       %uint = OpTypeInt 32 0
+%mat3v4float = OpTypeMatrix %v4float 3
+%GammaTransferParams = OpTypeStruct %float %float %float %float %float %float %float %uint
+    %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+%ExternalTextureParams = OpTypeStruct %uint %uint %mat3v4float %GammaTransferParams %GammaTransferParams %mat3v3float
+%ext_tex_params_block = OpTypeStruct %ExternalTextureParams
+%_ptr_Uniform_ext_tex_params_block = OpTypePointer Uniform %ext_tex_params_block
+%ext_tex_params = OpVariable %_ptr_Uniform_ext_tex_params_block Uniform
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %24 = OpTypeSampler
+%_ptr_UniformConstant_24 = OpTypePointer UniformConstant %24
+      %arg_1 = OpVariable %_ptr_UniformConstant_24 UniformConstant
+         %25 = OpTypeFunction %v3float %v3float %GammaTransferParams
+       %bool = OpTypeBool
+     %v3bool = OpTypeVector %bool 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %45 = OpConstantNull %v3float
+    %v2float = OpTypeVector %float 2
+         %65 = OpTypeFunction %v4float %11 %11 %24 %v2float %ExternalTextureParams
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %78 = OpConstantNull %int
+  %float_0_5 = OpConstant %float 0.5
+         %80 = OpConstantComposite %v2float %float_0_5 %float_0_5
+    %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %87 = OpConstantNull %v2float
+     %uint_1 = OpConstant %uint 1
+        %104 = OpTypeSampledImage %11
+        %119 = OpConstantNull %uint
+       %void = OpTypeVoid
+        %137 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+        %151 = OpTypeFunction %v4float
+%gammaCorrection = OpFunction %v3float None %25
+          %v = OpFunctionParameter %v3float
+     %params = OpFunctionParameter %GammaTransferParams
+         %29 = OpLabel
+         %43 = OpVariable %_ptr_Function_v3float Function %45
+         %55 = OpVariable %_ptr_Function_v3float Function %45
+         %61 = OpVariable %_ptr_Function_v3float Function %45
+         %30 = OpExtInst %v3float %31 FAbs %v
+         %32 = OpCompositeExtract %float %params 4
+         %33 = OpCompositeConstruct %v3float %32 %32 %32
+         %34 = OpFOrdLessThan %v3bool %30 %33
+         %37 = OpExtInst %v3float %31 FSign %v
+         %38 = OpCompositeExtract %float %params 3
+         %39 = OpExtInst %v3float %31 FAbs %v
+         %40 = OpVectorTimesScalar %v3float %39 %38
+         %41 = OpCompositeExtract %float %params 6
+         %46 = OpCompositeConstruct %v3float %41 %41 %41
+         %42 = OpFAdd %v3float %40 %46
+         %47 = OpFMul %v3float %37 %42
+         %48 = OpExtInst %v3float %31 FSign %v
+         %50 = OpCompositeExtract %float %params 1
+         %51 = OpExtInst %v3float %31 FAbs %v
+         %52 = OpVectorTimesScalar %v3float %51 %50
+         %53 = OpCompositeExtract %float %params 2
+         %56 = OpCompositeConstruct %v3float %53 %53 %53
+         %54 = OpFAdd %v3float %52 %56
+         %57 = OpCompositeExtract %float %params 0
+         %58 = OpCompositeConstruct %v3float %57 %57 %57
+         %49 = OpExtInst %v3float %31 Pow %54 %58
+         %59 = OpCompositeExtract %float %params 5
+         %62 = OpCompositeConstruct %v3float %59 %59 %59
+         %60 = OpFAdd %v3float %49 %62
+         %63 = OpFMul %v3float %48 %60
+         %64 = OpSelect %v3float %34 %47 %63
+               OpReturnValue %64
+               OpFunctionEnd
+%textureSampleExternal = OpFunction %v4float None %65
+     %plane0 = OpFunctionParameter %11
+     %plane1 = OpFunctionParameter %11
+        %smp = OpFunctionParameter %24
+      %coord = OpFunctionParameter %v2float
+   %params_0 = OpFunctionParameter %ExternalTextureParams
+         %73 = OpLabel
+         %85 = OpVariable %_ptr_Function_v2float Function %87
+         %94 = OpVariable %_ptr_Function_v2float Function %87
+      %color = OpVariable %_ptr_Function_v3float Function %45
+         %75 = OpImageQuerySizeLod %v2int %plane0 %78
+         %74 = OpConvertSToF %v2float %75
+         %81 = OpFDiv %v2float %80 %74
+         %88 = OpCompositeConstruct %v2float %float_1 %float_1
+         %84 = OpFSub %v2float %88 %81
+         %82 = OpExtInst %v2float %31 NClamp %coord %81 %84
+         %90 = OpImageQuerySizeLod %v2int %plane1 %78
+         %89 = OpConvertSToF %v2float %90
+         %91 = OpFDiv %v2float %80 %89
+         %95 = OpCompositeConstruct %v2float %float_1 %float_1
+         %93 = OpFSub %v2float %95 %91
+         %92 = OpExtInst %v2float %31 NClamp %coord %91 %93
+         %97 = OpCompositeExtract %uint %params_0 0
+         %99 = OpIEqual %bool %97 %uint_1
+               OpSelectionMerge %100 None
+               OpBranchConditional %99 %101 %102
+        %101 = OpLabel
+        %105 = OpSampledImage %104 %plane0 %smp
+        %103 = OpImageSampleExplicitLod %v4float %105 %82 Lod %8
+        %106 = OpVectorShuffle %v3float %103 %103 0 1 2
+               OpStore %color %106
+               OpBranch %100
+        %102 = OpLabel
+        %108 = OpSampledImage %104 %plane0 %smp
+        %107 = OpImageSampleExplicitLod %v4float %108 %82 Lod %8
+        %109 = OpCompositeExtract %float %107 0
+        %111 = OpSampledImage %104 %plane1 %smp
+        %110 = OpImageSampleExplicitLod %v4float %111 %92 Lod %8
+        %112 = OpVectorShuffle %v2float %110 %110 0 1
+        %113 = OpCompositeExtract %float %112 0
+        %114 = OpCompositeExtract %float %112 1
+        %115 = OpCompositeConstruct %v4float %109 %113 %114 %float_1
+        %116 = OpCompositeExtract %mat3v4float %params_0 2
+        %117 = OpVectorTimesMatrix %v3float %115 %116
+               OpStore %color %117
+               OpBranch %100
+        %100 = OpLabel
+        %118 = OpCompositeExtract %uint %params_0 1
+        %120 = OpIEqual %bool %118 %119
+               OpSelectionMerge %121 None
+               OpBranchConditional %120 %122 %121
+        %122 = OpLabel
+        %124 = OpLoad %v3float %color
+        %125 = OpCompositeExtract %GammaTransferParams %params_0 3
+        %123 = OpFunctionCall %v3float %gammaCorrection %124 %125
+               OpStore %color %123
+        %126 = OpCompositeExtract %mat3v3float %params_0 5
+        %127 = OpLoad %v3float %color
+        %128 = OpMatrixTimesVector %v3float %126 %127
+               OpStore %color %128
+        %130 = OpLoad %v3float %color
+        %131 = OpCompositeExtract %GammaTransferParams %params_0 4
+        %129 = OpFunctionCall %v3float %gammaCorrection %130 %131
+               OpStore %color %129
+               OpBranch %121
+        %121 = OpLabel
+        %132 = OpLoad %v3float %color
+        %133 = OpCompositeExtract %float %132 0
+        %134 = OpCompositeExtract %float %132 1
+        %135 = OpCompositeExtract %float %132 2
+        %136 = OpCompositeConstruct %v4float %133 %134 %135 %float_1
+               OpReturnValue %136
+               OpFunctionEnd
+%textureSampleBaseClampToEdge_7c04e6 = OpFunction %void None %137
+        %140 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+        %142 = OpLoad %11 %arg_0
+        %143 = OpLoad %11 %ext_tex_plane_1
+        %144 = OpLoad %24 %arg_1
+        %147 = OpAccessChain %_ptr_Uniform_ExternalTextureParams %ext_tex_params %uint_0
+        %148 = OpLoad %ExternalTextureParams %147
+        %141 = OpFunctionCall %v4float %textureSampleExternal %142 %143 %144 %87 %148
+               OpStore %res %141
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %151
+        %153 = OpLabel
+        %154 = OpFunctionCall %void %textureSampleBaseClampToEdge_7c04e6
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %137
+        %156 = OpLabel
+        %157 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %157
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %137
+        %159 = OpLabel
+        %160 = OpFunctionCall %void %textureSampleBaseClampToEdge_7c04e6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %137
+        %162 = OpLabel
+        %163 = OpFunctionCall %void %textureSampleBaseClampToEdge_7c04e6
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+@group(1) @binding(0) var arg_0 : texture_external;
+
+@group(1) @binding(1) var arg_1 : sampler;
+
+fn textureSampleBaseClampToEdge_7c04e6() {
+  var res : vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, vec2<f32>());
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_7c04e6();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+@group(1) @binding(0) var arg_0: texture_2d<f32>;
+@group(1) @binding(1) var arg_1: sampler;
+
+// fn textureSampleBaseClampToEdge(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureSampleBaseClampToEdge_9ca02c() {
+  var res: vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, vec2<f32>());
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_9ca02c();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.dxc.hlsl

@@ -0,0 +1,42 @@
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float4 tint_textureSampleBaseClampToEdge(Texture2D<float4> t, SamplerState s, float2 coord) {
+  int3 tint_tmp;
+  t.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 dims = float2(tint_tmp.xy);
+  const float2 half_texel = ((0.5f).xx / dims);
+  const float2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return t.SampleLevel(s, clamped, 0.0f);
+}
+
+void textureSampleBaseClampToEdge_9ca02c() {
+  float4 res = tint_textureSampleBaseClampToEdge(arg_0, arg_1, (0.0f).xx);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.fxc.hlsl

@@ -0,0 +1,42 @@
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float4 tint_textureSampleBaseClampToEdge(Texture2D<float4> t, SamplerState s, float2 coord) {
+  int3 tint_tmp;
+  t.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 dims = float2(tint_tmp.xy);
+  const float2 half_texel = ((0.5f).xx / dims);
+  const float2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return t.SampleLevel(s, clamped, 0.0f);
+}
+
+void textureSampleBaseClampToEdge_9ca02c() {
+  float4 res = tint_textureSampleBaseClampToEdge(arg_0, arg_1, (0.0f).xx);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.glsl

@@ -0,0 +1,79 @@
+#version 310 es
+
+
+vec4 tint_textureSampleBaseClampToEdge(highp sampler2D t_1, highp sampler2D t_s, vec2 coord) {
+  vec2 dims = vec2(textureSize(t_1, 0));
+  vec2 half_texel = (vec2(0.5f) / dims);
+  vec2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return textureLod(t_s, clamped, 0.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D arg_0_arg_1;
+void textureSampleBaseClampToEdge_9ca02c() {
+  vec4 res = tint_textureSampleBaseClampToEdge(arg_0_1, arg_0_arg_1, vec2(0.0f));
+}
+
+vec4 vertex_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+
+vec4 tint_textureSampleBaseClampToEdge(highp sampler2D t_1, highp sampler2D t_s, vec2 coord) {
+  vec2 dims = vec2(textureSize(t_1, 0));
+  vec2 half_texel = (vec2(0.5f) / dims);
+  vec2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return textureLod(t_s, clamped, 0.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D arg_0_arg_1;
+void textureSampleBaseClampToEdge_9ca02c() {
+  vec4 res = tint_textureSampleBaseClampToEdge(arg_0_1, arg_0_arg_1, vec2(0.0f));
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+
+vec4 tint_textureSampleBaseClampToEdge(highp sampler2D t_1, highp sampler2D t_s, vec2 coord) {
+  vec2 dims = vec2(textureSize(t_1, 0));
+  vec2 half_texel = (vec2(0.5f) / dims);
+  vec2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return textureLod(t_s, clamped, 0.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D arg_0_arg_1;
+void textureSampleBaseClampToEdge_9ca02c() {
+  vec4 res = tint_textureSampleBaseClampToEdge(arg_0_1, arg_0_arg_1, vec2(0.0f));
+}
+
+void compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.msl

@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float4 tint_textureSampleBaseClampToEdge(texture2d<float, access::sample> t, sampler s, float2 coord) {
+  float2 const dims = float2(int2(t.get_width(0), t.get_height(0)));
+  float2 const half_texel = (float2(0.5f) / dims);
+  float2 const clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return t.sample(s, clamped, level(0.0f));
+}
+
+void textureSampleBaseClampToEdge_9ca02c(texture2d<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float4 res = tint_textureSampleBaseClampToEdge(tint_symbol_1, tint_symbol_2, float2(0.0f));
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureSampleBaseClampToEdge_9ca02c(tint_symbol_3, tint_symbol_4);
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureSampleBaseClampToEdge_9ca02c(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureSampleBaseClampToEdge_9ca02c(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.spvasm

@@ -0,0 +1,110 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 63
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %31 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %tint_textureSampleBaseClampToEdge "tint_textureSampleBaseClampToEdge"
+               OpName %t "t"
+               OpName %s "s"
+               OpName %coord "coord"
+               OpName %textureSampleBaseClampToEdge_9ca02c "textureSampleBaseClampToEdge_9ca02c"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_1 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+    %v2float = OpTypeVector %float 2
+         %15 = OpTypeFunction %v4float %11 %14 %v2float
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %26 = OpConstantNull %int
+  %float_0_5 = OpConstant %float 0.5
+         %28 = OpConstantComposite %v2float %float_0_5 %float_0_5
+    %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %36 = OpConstantNull %v2float
+         %39 = OpTypeSampledImage %11
+       %void = OpTypeVoid
+         %41 = OpTypeFunction %void
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %50 = OpTypeFunction %v4float
+%tint_textureSampleBaseClampToEdge = OpFunction %v4float None %15
+          %t = OpFunctionParameter %11
+          %s = OpFunctionParameter %14
+      %coord = OpFunctionParameter %v2float
+         %21 = OpLabel
+         %34 = OpVariable %_ptr_Function_v2float Function %36
+         %23 = OpImageQuerySizeLod %v2int %t %26
+         %22 = OpConvertSToF %v2float %23
+         %29 = OpFDiv %v2float %28 %22
+         %37 = OpCompositeConstruct %v2float %float_1 %float_1
+         %33 = OpFSub %v2float %37 %29
+         %30 = OpExtInst %v2float %31 NClamp %coord %29 %33
+         %40 = OpSampledImage %39 %t %s
+         %38 = OpImageSampleExplicitLod %v4float %40 %30 Lod %8
+               OpReturnValue %38
+               OpFunctionEnd
+%textureSampleBaseClampToEdge_9ca02c = OpFunction %void None %41
+         %44 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %46 = OpLoad %11 %arg_0
+         %47 = OpLoad %14 %arg_1
+         %45 = OpFunctionCall %v4float %tint_textureSampleBaseClampToEdge %46 %47 %36
+               OpStore %res %45
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %50
+         %52 = OpLabel
+         %53 = OpFunctionCall %void %textureSampleBaseClampToEdge_9ca02c
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %41
+         %55 = OpLabel
+         %56 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %56
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %41
+         %58 = OpLabel
+         %59 = OpFunctionCall %void %textureSampleBaseClampToEdge_9ca02c
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %41
+         %61 = OpLabel
+         %62 = OpFunctionCall %void %textureSampleBaseClampToEdge_9ca02c
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+@group(1) @binding(0) var arg_0 : texture_2d<f32>;
+
+@group(1) @binding(1) var arg_1 : sampler;
+
+fn textureSampleBaseClampToEdge_9ca02c() {
+  var res : vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, vec2<f32>());
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_9ca02c();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}

test/tint/builtins/gen/literal/textureSampleLevel/979816.wgsl.expected.dxc.hlsl

@@ -1,3 +1,7 @@
+builtins/gen/literal/textureSampleLevel/979816.wgsl:28:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, vec2<f32>());
+                       ^^^^^^^^^^^^^^^^^^
+
 struct GammaTransferParams {
   float G;
   float A;

test/tint/builtins/gen/literal/textureSampleLevel/979816.wgsl.expected.fxc.hlsl

@@ -1,3 +1,7 @@
+builtins/gen/literal/textureSampleLevel/979816.wgsl:28:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, vec2<f32>());
+                       ^^^^^^^^^^^^^^^^^^
+
 struct GammaTransferParams {
   float G;
   float A;

test/tint/builtins/gen/literal/textureSampleLevel/979816.wgsl.expected.glsl

@@ -1,3 +1,7 @@
+builtins/gen/literal/textureSampleLevel/979816.wgsl:28:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, vec2<f32>());
+                       ^^^^^^^^^^^^^^^^^^
+
 #version 310 es
 
 struct GammaTransferParams {

test/tint/builtins/gen/literal/textureSampleLevel/979816.wgsl.expected.msl

@@ -1,3 +1,7 @@
+builtins/gen/literal/textureSampleLevel/979816.wgsl:28:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, vec2<f32>());
+                       ^^^^^^^^^^^^^^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;

test/tint/builtins/gen/literal/textureSampleLevel/979816.wgsl.expected.spvasm

@@ -1,3 +1,7 @@
+builtins/gen/literal/textureSampleLevel/979816.wgsl:28:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, vec2<f32>());
+                       ^^^^^^^^^^^^^^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

test/tint/builtins/gen/literal/textureSampleLevel/979816.wgsl.expected.wgsl

@@ -1,3 +1,7 @@
+builtins/gen/literal/textureSampleLevel/979816.wgsl:28:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, vec2<f32>());
+                       ^^^^^^^^^^^^^^^^^^
+
 @group(1) @binding(0) var arg_0 : texture_external;
 
 @group(1) @binding(1) var arg_1 : sampler;

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl

@@ -0,0 +1,46 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+@group(1) @binding(0) var arg_0: texture_external;
+@group(1) @binding(1) var arg_1: sampler;
+
+// fn textureSampleBaseClampToEdge(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureSampleBaseClampToEdge_7c04e6() {
+  var arg_2 = vec2<f32>();
+  var res: vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_7c04e6();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.dxc.hlsl

@@ -0,0 +1,123 @@
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+};
+
+Texture2D<float4> ext_tex_plane_1 : register(t2, space1);
+cbuffer cbuffer_ext_tex_params : register(b3, space1) {
+  uint4 ext_tex_params[11];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float3 gammaCorrection(float3 v, GammaTransferParams params) {
+  const bool3 cond = (abs(v) < float3((params.D).xxx));
+  const float3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  const float3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E));
+  return (cond ? t : f);
+}
+
+float4 textureSampleExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, SamplerState smp, float2 coord, ExternalTextureParams params) {
+  int3 tint_tmp;
+  plane0.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 plane0_dims = float2(tint_tmp.xy);
+  const float2 plane0_half_texel = ((0.5f).xx / plane0_dims);
+  const float2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  int3 tint_tmp_1;
+  plane1.GetDimensions(0, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
+  const float2 plane1_dims = float2(tint_tmp_1.xy);
+  const float2 plane1_half_texel = ((0.5f).xx / plane1_dims);
+  const float2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  float3 color = float3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = plane0.SampleLevel(smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.SampleLevel(smp, plane0_clamped, 0.0f).r, plane1.SampleLevel(smp, plane1_clamped, 0.0f).rg, 1.0f));
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = mul(color, params.gamutConversionMatrix);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return float4(color, 1.0f);
+}
+
+float3x4 tint_symbol_3(uint4 buffer[11], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  return float3x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]));
+}
+
+GammaTransferParams tint_symbol_5(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_3 = ((offset + 0u)) / 4;
+  const uint scalar_offset_4 = ((offset + 4u)) / 4;
+  const uint scalar_offset_5 = ((offset + 8u)) / 4;
+  const uint scalar_offset_6 = ((offset + 12u)) / 4;
+  const uint scalar_offset_7 = ((offset + 16u)) / 4;
+  const uint scalar_offset_8 = ((offset + 20u)) / 4;
+  const uint scalar_offset_9 = ((offset + 24u)) / 4;
+  const uint scalar_offset_10 = ((offset + 28u)) / 4;
+  const GammaTransferParams tint_symbol_9 = {asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(buffer[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(buffer[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(buffer[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(buffer[scalar_offset_9 / 4][scalar_offset_9 % 4]), buffer[scalar_offset_10 / 4][scalar_offset_10 % 4]};
+  return tint_symbol_9;
+}
+
+float3x3 tint_symbol_7(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_11 = ((offset + 0u)) / 4;
+  const uint scalar_offset_12 = ((offset + 16u)) / 4;
+  const uint scalar_offset_13 = ((offset + 32u)) / 4;
+  return float3x3(asfloat(buffer[scalar_offset_11 / 4].xyz), asfloat(buffer[scalar_offset_12 / 4].xyz), asfloat(buffer[scalar_offset_13 / 4].xyz));
+}
+
+ExternalTextureParams tint_symbol_1(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_14 = ((offset + 0u)) / 4;
+  const uint scalar_offset_15 = ((offset + 4u)) / 4;
+  const ExternalTextureParams tint_symbol_10 = {buffer[scalar_offset_14 / 4][scalar_offset_14 % 4], buffer[scalar_offset_15 / 4][scalar_offset_15 % 4], tint_symbol_3(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 64u)), tint_symbol_5(buffer, (offset + 96u)), tint_symbol_7(buffer, (offset + 128u))};
+  return tint_symbol_10;
+}
+
+void textureSampleBaseClampToEdge_7c04e6() {
+  float2 arg_2 = (0.0f).xx;
+  float4 res = textureSampleExternal(arg_0, ext_tex_plane_1, arg_1, arg_2, tint_symbol_1(ext_tex_params, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.fxc.hlsl

@@ -0,0 +1,123 @@
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  float3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  float3x3 gamutConversionMatrix;
+};
+
+Texture2D<float4> ext_tex_plane_1 : register(t2, space1);
+cbuffer cbuffer_ext_tex_params : register(b3, space1) {
+  uint4 ext_tex_params[11];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float3 gammaCorrection(float3 v, GammaTransferParams params) {
+  const bool3 cond = (abs(v) < float3((params.D).xxx));
+  const float3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  const float3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E));
+  return (cond ? t : f);
+}
+
+float4 textureSampleExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, SamplerState smp, float2 coord, ExternalTextureParams params) {
+  int3 tint_tmp;
+  plane0.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 plane0_dims = float2(tint_tmp.xy);
+  const float2 plane0_half_texel = ((0.5f).xx / plane0_dims);
+  const float2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  int3 tint_tmp_1;
+  plane1.GetDimensions(0, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
+  const float2 plane1_dims = float2(tint_tmp_1.xy);
+  const float2 plane1_half_texel = ((0.5f).xx / plane1_dims);
+  const float2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  float3 color = float3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = plane0.SampleLevel(smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.SampleLevel(smp, plane0_clamped, 0.0f).r, plane1.SampleLevel(smp, plane1_clamped, 0.0f).rg, 1.0f));
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = mul(color, params.gamutConversionMatrix);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return float4(color, 1.0f);
+}
+
+float3x4 tint_symbol_3(uint4 buffer[11], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  return float3x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]));
+}
+
+GammaTransferParams tint_symbol_5(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_3 = ((offset + 0u)) / 4;
+  const uint scalar_offset_4 = ((offset + 4u)) / 4;
+  const uint scalar_offset_5 = ((offset + 8u)) / 4;
+  const uint scalar_offset_6 = ((offset + 12u)) / 4;
+  const uint scalar_offset_7 = ((offset + 16u)) / 4;
+  const uint scalar_offset_8 = ((offset + 20u)) / 4;
+  const uint scalar_offset_9 = ((offset + 24u)) / 4;
+  const uint scalar_offset_10 = ((offset + 28u)) / 4;
+  const GammaTransferParams tint_symbol_9 = {asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(buffer[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(buffer[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(buffer[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(buffer[scalar_offset_9 / 4][scalar_offset_9 % 4]), buffer[scalar_offset_10 / 4][scalar_offset_10 % 4]};
+  return tint_symbol_9;
+}
+
+float3x3 tint_symbol_7(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_11 = ((offset + 0u)) / 4;
+  const uint scalar_offset_12 = ((offset + 16u)) / 4;
+  const uint scalar_offset_13 = ((offset + 32u)) / 4;
+  return float3x3(asfloat(buffer[scalar_offset_11 / 4].xyz), asfloat(buffer[scalar_offset_12 / 4].xyz), asfloat(buffer[scalar_offset_13 / 4].xyz));
+}
+
+ExternalTextureParams tint_symbol_1(uint4 buffer[11], uint offset) {
+  const uint scalar_offset_14 = ((offset + 0u)) / 4;
+  const uint scalar_offset_15 = ((offset + 4u)) / 4;
+  const ExternalTextureParams tint_symbol_10 = {buffer[scalar_offset_14 / 4][scalar_offset_14 % 4], buffer[scalar_offset_15 / 4][scalar_offset_15 % 4], tint_symbol_3(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 64u)), tint_symbol_5(buffer, (offset + 96u)), tint_symbol_7(buffer, (offset + 128u))};
+  return tint_symbol_10;
+}
+
+void textureSampleBaseClampToEdge_7c04e6() {
+  float2 arg_2 = (0.0f).xx;
+  float4 res = textureSampleExternal(arg_0, ext_tex_plane_1, arg_1, arg_2, tint_symbol_1(ext_tex_params, 0u));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return;
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.glsl

@@ -0,0 +1,223 @@
+#version 310 es
+
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+};
+
+layout(binding = 3, std140) uniform ext_tex_params_block_ubo {
+  ExternalTextureParams inner;
+} ext_tex_params;
+
+vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
+  bvec3 cond = lessThan(abs(v), vec3(params.D));
+  vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
+  return mix(f, t, cond);
+}
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+  vec2 plane0_dims = vec2(textureSize(plane0_1, 0));
+  vec2 plane0_half_texel = (vec2(0.5f) / plane0_dims);
+  vec2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  vec2 plane1_dims = vec2(textureSize(plane1_1, 0));
+  vec2 plane1_half_texel = (vec2(0.5f) / plane1_dims);
+  vec2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  vec3 color = vec3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = textureLod(plane0_smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4(textureLod(plane0_smp, plane0_clamped, 0.0f).r, textureLod(plane1_smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4(color, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleBaseClampToEdge_7c04e6() {
+  vec2 arg_2 = vec2(0.0f);
+  vec4 res = textureSampleExternal(arg_0_1, ext_tex_plane_1_1, arg_0_arg_1, ext_tex_plane_1_arg_1, arg_2, ext_tex_params.inner);
+}
+
+vec4 vertex_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+};
+
+layout(binding = 3, std140) uniform ext_tex_params_block_ubo {
+  ExternalTextureParams inner;
+} ext_tex_params;
+
+vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
+  bvec3 cond = lessThan(abs(v), vec3(params.D));
+  vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
+  return mix(f, t, cond);
+}
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+  vec2 plane0_dims = vec2(textureSize(plane0_1, 0));
+  vec2 plane0_half_texel = (vec2(0.5f) / plane0_dims);
+  vec2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  vec2 plane1_dims = vec2(textureSize(plane1_1, 0));
+  vec2 plane1_half_texel = (vec2(0.5f) / plane1_dims);
+  vec2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  vec3 color = vec3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = textureLod(plane0_smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4(textureLod(plane0_smp, plane0_clamped, 0.0f).r, textureLod(plane1_smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4(color, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleBaseClampToEdge_7c04e6() {
+  vec2 arg_2 = vec2(0.0f);
+  vec4 res = textureSampleExternal(arg_0_1, ext_tex_plane_1_1, arg_0_arg_1, ext_tex_plane_1_arg_1, arg_2, ext_tex_params.inner);
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct GammaTransferParams {
+  float G;
+  float A;
+  float B;
+  float C;
+  float D;
+  float E;
+  float F;
+  uint padding;
+};
+
+struct ExternalTextureParams {
+  uint numPlanes;
+  uint doYuvToRgbConversionOnly;
+  mat3x4 yuvToRgbConversionMatrix;
+  GammaTransferParams gammaDecodeParams;
+  GammaTransferParams gammaEncodeParams;
+  mat3 gamutConversionMatrix;
+};
+
+layout(binding = 3, std140) uniform ext_tex_params_block_ubo {
+  ExternalTextureParams inner;
+} ext_tex_params;
+
+vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
+  bvec3 cond = lessThan(abs(v), vec3(params.D));
+  vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
+  vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
+  return mix(f, t, cond);
+}
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+  vec2 plane0_dims = vec2(textureSize(plane0_1, 0));
+  vec2 plane0_half_texel = (vec2(0.5f) / plane0_dims);
+  vec2 plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  vec2 plane1_dims = vec2(textureSize(plane1_1, 0));
+  vec2 plane1_half_texel = (vec2(0.5f) / plane1_dims);
+  vec2 plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  vec3 color = vec3(0.0f, 0.0f, 0.0f);
+  if ((params.numPlanes == 1u)) {
+    color = textureLod(plane0_smp, plane0_clamped, 0.0f).rgb;
+  } else {
+    color = (vec4(textureLod(plane0_smp, plane0_clamped, 0.0f).r, textureLod(plane1_smp, plane1_clamped, 0.0f).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return vec4(color, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleBaseClampToEdge_7c04e6() {
+  vec2 arg_2 = vec2(0.0f);
+  vec4 res = textureSampleExternal(arg_0_1, ext_tex_plane_1_1, arg_0_arg_1, ext_tex_plane_1_arg_1, arg_2, ext_tex_params.inner);
+}
+
+void compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.msl

@@ -0,0 +1,96 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+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 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 ExternalTextureParams {
+  /* 0x0000 */ uint numPlanes;
+  /* 0x0004 */ uint doYuvToRgbConversionOnly;
+  /* 0x0008 */ tint_array<int8_t, 8> tint_pad;
+  /* 0x0010 */ float3x4 yuvToRgbConversionMatrix;
+  /* 0x0040 */ GammaTransferParams gammaDecodeParams;
+  /* 0x0060 */ GammaTransferParams gammaEncodeParams;
+  /* 0x0080 */ float3x3 gamutConversionMatrix;
+};
+
+float3 gammaCorrection(float3 v, GammaTransferParams params) {
+  bool3 const cond = (fabs(v) < float3(params.D));
+  float3 const t = (sign(v) * ((params.C * fabs(v)) + params.F));
+  float3 const f = (sign(v) * (pow(((params.A * fabs(v)) + params.B), float3(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+float4 textureSampleExternal(texture2d<float, access::sample> plane0, texture2d<float, access::sample> plane1, sampler smp, float2 coord, ExternalTextureParams params) {
+  float2 const plane0_dims = float2(int2(plane0.get_width(0), plane0.get_height(0)));
+  float2 const plane0_half_texel = (float2(0.5f) / plane0_dims);
+  float2 const plane0_clamped = clamp(coord, plane0_half_texel, (1.0f - plane0_half_texel));
+  float2 const plane1_dims = float2(int2(plane1.get_width(0), plane1.get_height(0)));
+  float2 const plane1_half_texel = (float2(0.5f) / plane1_dims);
+  float2 const plane1_clamped = clamp(coord, plane1_half_texel, (1.0f - plane1_half_texel));
+  float3 color = 0.0f;
+  if ((params.numPlanes == 1u)) {
+    color = float4(plane0.sample(smp, plane0_clamped, level(0.0f))).rgb;
+  } else {
+    color = (float4(plane0.sample(smp, plane0_clamped, level(0.0f))[0], float4(plane1.sample(smp, plane1_clamped, level(0.0f))).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+  }
+  if ((params.doYuvToRgbConversionOnly == 0u)) {
+    color = gammaCorrection(color, params.gammaDecodeParams);
+    color = (params.gamutConversionMatrix * color);
+    color = gammaCorrection(color, params.gammaEncodeParams);
+  }
+  return float4(color, 1.0f);
+}
+
+void textureSampleBaseClampToEdge_7c04e6(texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2, sampler tint_symbol_3, const constant ExternalTextureParams* const tint_symbol_4) {
+  float2 arg_2 = float2(0.0f);
+  float4 res = textureSampleExternal(tint_symbol_1, tint_symbol_2, tint_symbol_3, arg_2, *(tint_symbol_4));
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_5, texture2d<float, access::sample> tint_symbol_6, sampler tint_symbol_7, const constant ExternalTextureParams* const tint_symbol_8) {
+  textureSampleBaseClampToEdge_7c04e6(tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_9 [[texture(0)]], texture2d<float, access::sample> tint_symbol_10 [[texture(1)]], sampler tint_symbol_11 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_12 [[buffer(2)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_13 [[texture(0)]], texture2d<float, access::sample> tint_symbol_14 [[texture(1)]], sampler tint_symbol_15 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_16 [[buffer(2)]]) {
+  textureSampleBaseClampToEdge_7c04e6(tint_symbol_13, tint_symbol_14, tint_symbol_15, tint_symbol_16);
+  return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_17 [[texture(0)]], texture2d<float, access::sample> tint_symbol_18 [[texture(1)]], sampler tint_symbol_19 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_20 [[buffer(2)]]) {
+  textureSampleBaseClampToEdge_7c04e6(tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
+  return;
+}
+

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.spvasm

@@ -0,0 +1,280 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 166
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %31 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %ext_tex_plane_1 "ext_tex_plane_1"
+               OpName %ext_tex_params_block "ext_tex_params_block"
+               OpMemberName %ext_tex_params_block 0 "inner"
+               OpName %ExternalTextureParams "ExternalTextureParams"
+               OpMemberName %ExternalTextureParams 0 "numPlanes"
+               OpMemberName %ExternalTextureParams 1 "doYuvToRgbConversionOnly"
+               OpMemberName %ExternalTextureParams 2 "yuvToRgbConversionMatrix"
+               OpMemberName %ExternalTextureParams 3 "gammaDecodeParams"
+               OpName %GammaTransferParams "GammaTransferParams"
+               OpMemberName %GammaTransferParams 0 "G"
+               OpMemberName %GammaTransferParams 1 "A"
+               OpMemberName %GammaTransferParams 2 "B"
+               OpMemberName %GammaTransferParams 3 "C"
+               OpMemberName %GammaTransferParams 4 "D"
+               OpMemberName %GammaTransferParams 5 "E"
+               OpMemberName %GammaTransferParams 6 "F"
+               OpMemberName %GammaTransferParams 7 "padding"
+               OpMemberName %ExternalTextureParams 4 "gammaEncodeParams"
+               OpMemberName %ExternalTextureParams 5 "gamutConversionMatrix"
+               OpName %ext_tex_params "ext_tex_params"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %gammaCorrection "gammaCorrection"
+               OpName %v "v"
+               OpName %params "params"
+               OpName %textureSampleExternal "textureSampleExternal"
+               OpName %plane0 "plane0"
+               OpName %plane1 "plane1"
+               OpName %smp "smp"
+               OpName %coord "coord"
+               OpName %params_0 "params"
+               OpName %color "color"
+               OpName %textureSampleBaseClampToEdge_7c04e6 "textureSampleBaseClampToEdge_7c04e6"
+               OpName %arg_2 "arg_2"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %ext_tex_plane_1 DescriptorSet 1
+               OpDecorate %ext_tex_plane_1 Binding 2
+               OpDecorate %ext_tex_params_block Block
+               OpMemberDecorate %ext_tex_params_block 0 Offset 0
+               OpMemberDecorate %ExternalTextureParams 0 Offset 0
+               OpMemberDecorate %ExternalTextureParams 1 Offset 4
+               OpMemberDecorate %ExternalTextureParams 2 Offset 16
+               OpMemberDecorate %ExternalTextureParams 2 ColMajor
+               OpMemberDecorate %ExternalTextureParams 2 MatrixStride 16
+               OpMemberDecorate %ExternalTextureParams 3 Offset 64
+               OpMemberDecorate %GammaTransferParams 0 Offset 0
+               OpMemberDecorate %GammaTransferParams 1 Offset 4
+               OpMemberDecorate %GammaTransferParams 2 Offset 8
+               OpMemberDecorate %GammaTransferParams 3 Offset 12
+               OpMemberDecorate %GammaTransferParams 4 Offset 16
+               OpMemberDecorate %GammaTransferParams 5 Offset 20
+               OpMemberDecorate %GammaTransferParams 6 Offset 24
+               OpMemberDecorate %GammaTransferParams 7 Offset 28
+               OpMemberDecorate %ExternalTextureParams 4 Offset 96
+               OpMemberDecorate %ExternalTextureParams 5 Offset 128
+               OpMemberDecorate %ExternalTextureParams 5 ColMajor
+               OpMemberDecorate %ExternalTextureParams 5 MatrixStride 16
+               OpDecorate %ext_tex_params NonWritable
+               OpDecorate %ext_tex_params DescriptorSet 1
+               OpDecorate %ext_tex_params Binding 3
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+       %uint = OpTypeInt 32 0
+%mat3v4float = OpTypeMatrix %v4float 3
+%GammaTransferParams = OpTypeStruct %float %float %float %float %float %float %float %uint
+    %v3float = OpTypeVector %float 3
+%mat3v3float = OpTypeMatrix %v3float 3
+%ExternalTextureParams = OpTypeStruct %uint %uint %mat3v4float %GammaTransferParams %GammaTransferParams %mat3v3float
+%ext_tex_params_block = OpTypeStruct %ExternalTextureParams
+%_ptr_Uniform_ext_tex_params_block = OpTypePointer Uniform %ext_tex_params_block
+%ext_tex_params = OpVariable %_ptr_Uniform_ext_tex_params_block Uniform
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %24 = OpTypeSampler
+%_ptr_UniformConstant_24 = OpTypePointer UniformConstant %24
+      %arg_1 = OpVariable %_ptr_UniformConstant_24 UniformConstant
+         %25 = OpTypeFunction %v3float %v3float %GammaTransferParams
+       %bool = OpTypeBool
+     %v3bool = OpTypeVector %bool 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %45 = OpConstantNull %v3float
+    %v2float = OpTypeVector %float 2
+         %65 = OpTypeFunction %v4float %11 %11 %24 %v2float %ExternalTextureParams
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %78 = OpConstantNull %int
+  %float_0_5 = OpConstant %float 0.5
+         %80 = OpConstantComposite %v2float %float_0_5 %float_0_5
+    %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %87 = OpConstantNull %v2float
+     %uint_1 = OpConstant %uint 1
+        %104 = OpTypeSampledImage %11
+        %119 = OpConstantNull %uint
+       %void = OpTypeVoid
+        %137 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+        %153 = OpTypeFunction %v4float
+%gammaCorrection = OpFunction %v3float None %25
+          %v = OpFunctionParameter %v3float
+     %params = OpFunctionParameter %GammaTransferParams
+         %29 = OpLabel
+         %43 = OpVariable %_ptr_Function_v3float Function %45
+         %55 = OpVariable %_ptr_Function_v3float Function %45
+         %61 = OpVariable %_ptr_Function_v3float Function %45
+         %30 = OpExtInst %v3float %31 FAbs %v
+         %32 = OpCompositeExtract %float %params 4
+         %33 = OpCompositeConstruct %v3float %32 %32 %32
+         %34 = OpFOrdLessThan %v3bool %30 %33
+         %37 = OpExtInst %v3float %31 FSign %v
+         %38 = OpCompositeExtract %float %params 3
+         %39 = OpExtInst %v3float %31 FAbs %v
+         %40 = OpVectorTimesScalar %v3float %39 %38
+         %41 = OpCompositeExtract %float %params 6
+         %46 = OpCompositeConstruct %v3float %41 %41 %41
+         %42 = OpFAdd %v3float %40 %46
+         %47 = OpFMul %v3float %37 %42
+         %48 = OpExtInst %v3float %31 FSign %v
+         %50 = OpCompositeExtract %float %params 1
+         %51 = OpExtInst %v3float %31 FAbs %v
+         %52 = OpVectorTimesScalar %v3float %51 %50
+         %53 = OpCompositeExtract %float %params 2
+         %56 = OpCompositeConstruct %v3float %53 %53 %53
+         %54 = OpFAdd %v3float %52 %56
+         %57 = OpCompositeExtract %float %params 0
+         %58 = OpCompositeConstruct %v3float %57 %57 %57
+         %49 = OpExtInst %v3float %31 Pow %54 %58
+         %59 = OpCompositeExtract %float %params 5
+         %62 = OpCompositeConstruct %v3float %59 %59 %59
+         %60 = OpFAdd %v3float %49 %62
+         %63 = OpFMul %v3float %48 %60
+         %64 = OpSelect %v3float %34 %47 %63
+               OpReturnValue %64
+               OpFunctionEnd
+%textureSampleExternal = OpFunction %v4float None %65
+     %plane0 = OpFunctionParameter %11
+     %plane1 = OpFunctionParameter %11
+        %smp = OpFunctionParameter %24
+      %coord = OpFunctionParameter %v2float
+   %params_0 = OpFunctionParameter %ExternalTextureParams
+         %73 = OpLabel
+         %85 = OpVariable %_ptr_Function_v2float Function %87
+         %94 = OpVariable %_ptr_Function_v2float Function %87
+      %color = OpVariable %_ptr_Function_v3float Function %45
+         %75 = OpImageQuerySizeLod %v2int %plane0 %78
+         %74 = OpConvertSToF %v2float %75
+         %81 = OpFDiv %v2float %80 %74
+         %88 = OpCompositeConstruct %v2float %float_1 %float_1
+         %84 = OpFSub %v2float %88 %81
+         %82 = OpExtInst %v2float %31 NClamp %coord %81 %84
+         %90 = OpImageQuerySizeLod %v2int %plane1 %78
+         %89 = OpConvertSToF %v2float %90
+         %91 = OpFDiv %v2float %80 %89
+         %95 = OpCompositeConstruct %v2float %float_1 %float_1
+         %93 = OpFSub %v2float %95 %91
+         %92 = OpExtInst %v2float %31 NClamp %coord %91 %93
+         %97 = OpCompositeExtract %uint %params_0 0
+         %99 = OpIEqual %bool %97 %uint_1
+               OpSelectionMerge %100 None
+               OpBranchConditional %99 %101 %102
+        %101 = OpLabel
+        %105 = OpSampledImage %104 %plane0 %smp
+        %103 = OpImageSampleExplicitLod %v4float %105 %82 Lod %8
+        %106 = OpVectorShuffle %v3float %103 %103 0 1 2
+               OpStore %color %106
+               OpBranch %100
+        %102 = OpLabel
+        %108 = OpSampledImage %104 %plane0 %smp
+        %107 = OpImageSampleExplicitLod %v4float %108 %82 Lod %8
+        %109 = OpCompositeExtract %float %107 0
+        %111 = OpSampledImage %104 %plane1 %smp
+        %110 = OpImageSampleExplicitLod %v4float %111 %92 Lod %8
+        %112 = OpVectorShuffle %v2float %110 %110 0 1
+        %113 = OpCompositeExtract %float %112 0
+        %114 = OpCompositeExtract %float %112 1
+        %115 = OpCompositeConstruct %v4float %109 %113 %114 %float_1
+        %116 = OpCompositeExtract %mat3v4float %params_0 2
+        %117 = OpVectorTimesMatrix %v3float %115 %116
+               OpStore %color %117
+               OpBranch %100
+        %100 = OpLabel
+        %118 = OpCompositeExtract %uint %params_0 1
+        %120 = OpIEqual %bool %118 %119
+               OpSelectionMerge %121 None
+               OpBranchConditional %120 %122 %121
+        %122 = OpLabel
+        %124 = OpLoad %v3float %color
+        %125 = OpCompositeExtract %GammaTransferParams %params_0 3
+        %123 = OpFunctionCall %v3float %gammaCorrection %124 %125
+               OpStore %color %123
+        %126 = OpCompositeExtract %mat3v3float %params_0 5
+        %127 = OpLoad %v3float %color
+        %128 = OpMatrixTimesVector %v3float %126 %127
+               OpStore %color %128
+        %130 = OpLoad %v3float %color
+        %131 = OpCompositeExtract %GammaTransferParams %params_0 4
+        %129 = OpFunctionCall %v3float %gammaCorrection %130 %131
+               OpStore %color %129
+               OpBranch %121
+        %121 = OpLabel
+        %132 = OpLoad %v3float %color
+        %133 = OpCompositeExtract %float %132 0
+        %134 = OpCompositeExtract %float %132 1
+        %135 = OpCompositeExtract %float %132 2
+        %136 = OpCompositeConstruct %v4float %133 %134 %135 %float_1
+               OpReturnValue %136
+               OpFunctionEnd
+%textureSampleBaseClampToEdge_7c04e6 = OpFunction %void None %137
+        %140 = OpLabel
+      %arg_2 = OpVariable %_ptr_Function_v2float Function %87
+        %res = OpVariable %_ptr_Function_v4float Function %5
+               OpStore %arg_2 %87
+        %143 = OpLoad %11 %arg_0
+        %144 = OpLoad %11 %ext_tex_plane_1
+        %145 = OpLoad %24 %arg_1
+        %146 = OpLoad %v2float %arg_2
+        %149 = OpAccessChain %_ptr_Uniform_ExternalTextureParams %ext_tex_params %uint_0
+        %150 = OpLoad %ExternalTextureParams %149
+        %142 = OpFunctionCall %v4float %textureSampleExternal %143 %144 %145 %146 %150
+               OpStore %res %142
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %153
+        %155 = OpLabel
+        %156 = OpFunctionCall %void %textureSampleBaseClampToEdge_7c04e6
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %137
+        %158 = OpLabel
+        %159 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %159
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %137
+        %161 = OpLabel
+        %162 = OpFunctionCall %void %textureSampleBaseClampToEdge_7c04e6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %137
+        %164 = OpLabel
+        %165 = OpFunctionCall %void %textureSampleBaseClampToEdge_7c04e6
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/7c04e6.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+@group(1) @binding(0) var arg_0 : texture_external;
+
+@group(1) @binding(1) var arg_1 : sampler;
+
+fn textureSampleBaseClampToEdge_7c04e6() {
+  var arg_2 = vec2<f32>();
+  var res : vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_7c04e6();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_7c04e6();
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl

@@ -0,0 +1,46 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+@group(1) @binding(0) var arg_0: texture_2d<f32>;
+@group(1) @binding(1) var arg_1: sampler;
+
+// fn textureSampleBaseClampToEdge(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+fn textureSampleBaseClampToEdge_9ca02c() {
+  var arg_2 = vec2<f32>();
+  var res: vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_9ca02c();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.dxc.hlsl

@@ -0,0 +1,43 @@
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float4 tint_textureSampleBaseClampToEdge(Texture2D<float4> t, SamplerState s, float2 coord) {
+  int3 tint_tmp;
+  t.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 dims = float2(tint_tmp.xy);
+  const float2 half_texel = ((0.5f).xx / dims);
+  const float2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return t.SampleLevel(s, clamped, 0.0f);
+}
+
+void textureSampleBaseClampToEdge_9ca02c() {
+  float2 arg_2 = (0.0f).xx;
+  float4 res = tint_textureSampleBaseClampToEdge(arg_0, arg_1, arg_2);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.fxc.hlsl

@@ -0,0 +1,43 @@
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+
+float4 tint_textureSampleBaseClampToEdge(Texture2D<float4> t, SamplerState s, float2 coord) {
+  int3 tint_tmp;
+  t.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const float2 dims = float2(tint_tmp.xy);
+  const float2 half_texel = ((0.5f).xx / dims);
+  const float2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return t.SampleLevel(s, clamped, 0.0f);
+}
+
+void textureSampleBaseClampToEdge_9ca02c() {
+  float2 arg_2 = (0.0f).xx;
+  float4 res = tint_textureSampleBaseClampToEdge(arg_0, arg_1, arg_2);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return;
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.glsl

@@ -0,0 +1,82 @@
+#version 310 es
+
+
+vec4 tint_textureSampleBaseClampToEdge(highp sampler2D t_1, highp sampler2D t_s, vec2 coord) {
+  vec2 dims = vec2(textureSize(t_1, 0));
+  vec2 half_texel = (vec2(0.5f) / dims);
+  vec2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return textureLod(t_s, clamped, 0.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D arg_0_arg_1;
+void textureSampleBaseClampToEdge_9ca02c() {
+  vec2 arg_2 = vec2(0.0f);
+  vec4 res = tint_textureSampleBaseClampToEdge(arg_0_1, arg_0_arg_1, arg_2);
+}
+
+vec4 vertex_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+
+vec4 tint_textureSampleBaseClampToEdge(highp sampler2D t_1, highp sampler2D t_s, vec2 coord) {
+  vec2 dims = vec2(textureSize(t_1, 0));
+  vec2 half_texel = (vec2(0.5f) / dims);
+  vec2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return textureLod(t_s, clamped, 0.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D arg_0_arg_1;
+void textureSampleBaseClampToEdge_9ca02c() {
+  vec2 arg_2 = vec2(0.0f);
+  vec4 res = tint_textureSampleBaseClampToEdge(arg_0_1, arg_0_arg_1, arg_2);
+}
+
+void fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+
+vec4 tint_textureSampleBaseClampToEdge(highp sampler2D t_1, highp sampler2D t_s, vec2 coord) {
+  vec2 dims = vec2(textureSize(t_1, 0));
+  vec2 half_texel = (vec2(0.5f) / dims);
+  vec2 clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return textureLod(t_s, clamped, 0.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D arg_0_arg_1;
+void textureSampleBaseClampToEdge_9ca02c() {
+  vec2 arg_2 = vec2(0.0f);
+  vec4 res = tint_textureSampleBaseClampToEdge(arg_0_1, arg_0_arg_1, arg_2);
+}
+
+void compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.msl

@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+float4 tint_textureSampleBaseClampToEdge(texture2d<float, access::sample> t, sampler s, float2 coord) {
+  float2 const dims = float2(int2(t.get_width(0), t.get_height(0)));
+  float2 const half_texel = (float2(0.5f) / dims);
+  float2 const clamped = clamp(coord, half_texel, (1.0f - half_texel));
+  return t.sample(s, clamped, level(0.0f));
+}
+
+void textureSampleBaseClampToEdge_9ca02c(texture2d<float, access::sample> tint_symbol_1, sampler tint_symbol_2) {
+  float2 arg_2 = float2(0.0f);
+  float4 res = tint_textureSampleBaseClampToEdge(tint_symbol_1, tint_symbol_2, arg_2);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
+  textureSampleBaseClampToEdge_9ca02c(tint_symbol_3, tint_symbol_4);
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]], sampler tint_symbol_6 [[sampler(0)]]) {
+  float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_7 [[texture(0)]], sampler tint_symbol_8 [[sampler(0)]]) {
+  textureSampleBaseClampToEdge_9ca02c(tint_symbol_7, tint_symbol_8);
+  return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_9 [[texture(0)]], sampler tint_symbol_10 [[sampler(0)]]) {
+  textureSampleBaseClampToEdge_9ca02c(tint_symbol_9, tint_symbol_10);
+  return;
+}
+

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.spvasm

@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 65
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %31 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %arg_0 "arg_0"
+               OpName %arg_1 "arg_1"
+               OpName %tint_textureSampleBaseClampToEdge "tint_textureSampleBaseClampToEdge"
+               OpName %t "t"
+               OpName %s "s"
+               OpName %coord "coord"
+               OpName %textureSampleBaseClampToEdge_9ca02c "textureSampleBaseClampToEdge_9ca02c"
+               OpName %arg_2 "arg_2"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %arg_1 DescriptorSet 1
+               OpDecorate %arg_1 Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+         %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+      %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+         %14 = OpTypeSampler
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+      %arg_1 = OpVariable %_ptr_UniformConstant_14 UniformConstant
+    %v2float = OpTypeVector %float 2
+         %15 = OpTypeFunction %v4float %11 %14 %v2float
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %26 = OpConstantNull %int
+  %float_0_5 = OpConstant %float 0.5
+         %28 = OpConstantComposite %v2float %float_0_5 %float_0_5
+    %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %36 = OpConstantNull %v2float
+         %39 = OpTypeSampledImage %11
+       %void = OpTypeVoid
+         %41 = OpTypeFunction %void
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %52 = OpTypeFunction %v4float
+%tint_textureSampleBaseClampToEdge = OpFunction %v4float None %15
+          %t = OpFunctionParameter %11
+          %s = OpFunctionParameter %14
+      %coord = OpFunctionParameter %v2float
+         %21 = OpLabel
+         %34 = OpVariable %_ptr_Function_v2float Function %36
+         %23 = OpImageQuerySizeLod %v2int %t %26
+         %22 = OpConvertSToF %v2float %23
+         %29 = OpFDiv %v2float %28 %22
+         %37 = OpCompositeConstruct %v2float %float_1 %float_1
+         %33 = OpFSub %v2float %37 %29
+         %30 = OpExtInst %v2float %31 NClamp %coord %29 %33
+         %40 = OpSampledImage %39 %t %s
+         %38 = OpImageSampleExplicitLod %v4float %40 %30 Lod %8
+               OpReturnValue %38
+               OpFunctionEnd
+%textureSampleBaseClampToEdge_9ca02c = OpFunction %void None %41
+         %44 = OpLabel
+      %arg_2 = OpVariable %_ptr_Function_v2float Function %36
+        %res = OpVariable %_ptr_Function_v4float Function %5
+               OpStore %arg_2 %36
+         %47 = OpLoad %11 %arg_0
+         %48 = OpLoad %14 %arg_1
+         %49 = OpLoad %v2float %arg_2
+         %46 = OpFunctionCall %v4float %tint_textureSampleBaseClampToEdge %47 %48 %49
+               OpStore %res %46
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %52
+         %54 = OpLabel
+         %55 = OpFunctionCall %void %textureSampleBaseClampToEdge_9ca02c
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %41
+         %57 = OpLabel
+         %58 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %58
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %41
+         %60 = OpLabel
+         %61 = OpFunctionCall %void %textureSampleBaseClampToEdge_9ca02c
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %41
+         %63 = OpLabel
+         %64 = OpFunctionCall %void %textureSampleBaseClampToEdge_9ca02c
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/var/textureSampleBaseClampToEdge/9ca02c.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+@group(1) @binding(0) var arg_0 : texture_2d<f32>;
+
+@group(1) @binding(1) var arg_1 : sampler;
+
+fn textureSampleBaseClampToEdge_9ca02c() {
+  var arg_2 = vec2<f32>();
+  var res : vec4<f32> = textureSampleBaseClampToEdge(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  textureSampleBaseClampToEdge_9ca02c();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureSampleBaseClampToEdge_9ca02c();
+}

test/tint/builtins/gen/var/textureSampleLevel/979816.wgsl.expected.dxc.hlsl

@@ -1,3 +1,7 @@
+builtins/gen/var/textureSampleLevel/979816.wgsl:29:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, arg_2);
+                       ^^^^^^^^^^^^^^^^^^
+
 struct GammaTransferParams {
   float G;
   float A;

test/tint/builtins/gen/var/textureSampleLevel/979816.wgsl.expected.fxc.hlsl

@@ -1,3 +1,7 @@
+builtins/gen/var/textureSampleLevel/979816.wgsl:29:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, arg_2);
+                       ^^^^^^^^^^^^^^^^^^
+
 struct GammaTransferParams {
   float G;
   float A;

test/tint/builtins/gen/var/textureSampleLevel/979816.wgsl.expected.glsl

@@ -1,3 +1,7 @@
+builtins/gen/var/textureSampleLevel/979816.wgsl:29:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, arg_2);
+                       ^^^^^^^^^^^^^^^^^^
+
 #version 310 es
 
 struct GammaTransferParams {

test/tint/builtins/gen/var/textureSampleLevel/979816.wgsl.expected.msl

@@ -1,3 +1,7 @@
+builtins/gen/var/textureSampleLevel/979816.wgsl:29:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, arg_2);
+                       ^^^^^^^^^^^^^^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;

test/tint/builtins/gen/var/textureSampleLevel/979816.wgsl.expected.spvasm

@@ -1,3 +1,7 @@
+builtins/gen/var/textureSampleLevel/979816.wgsl:29:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, arg_2);
+                       ^^^^^^^^^^^^^^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

test/tint/builtins/gen/var/textureSampleLevel/979816.wgsl.expected.wgsl

@@ -1,3 +1,7 @@
+builtins/gen/var/textureSampleLevel/979816.wgsl:29:24 warning: use of deprecated builtin
+  var res: vec4<f32> = textureSampleLevel(arg_0, arg_1, arg_2);
+                       ^^^^^^^^^^^^^^^^^^
+
 @group(1) @binding(0) var arg_0 : texture_external;
 
 @group(1) @binding(1) var arg_1 : sampler;