Move tint::transform::Robustness to a santizier transform

There's no good reason for this to be public.
Move it into the writers, and expose a 'disable_robustness' option to
turn it off. This can be expanded to hold more fine-grain control in the
future.

Change-Id: I6ea6e54a27b2ae0fbcba5fdf45539063045cc15a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/122203
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>

include/tint/tint.h

@@ -32,7 +32,6 @@
 #include "src/tint/transform/manager.h"
 #include "src/tint/transform/multiplanar_external_texture.h"
 #include "src/tint/transform/renamer.h"
-#include "src/tint/transform/robustness.h"
 #include "src/tint/transform/single_entry_point.h"
 #include "src/tint/transform/substitute_override.h"
 #include "src/tint/transform/vertex_pulling.h"

src/dawn/native/d3d12/ShaderModuleD3D12.cpp

@@ -331,10 +331,6 @@ ResultOrError<std::string> TranslateToHLSL(
             std::move(r.substituteOverrideConfig).value());
     }
 
-    if (r.isRobustnessEnabled) {
-        transformManager.Add<tint::transform::Robustness>();
-    }
-
     transformManager.Add<tint::transform::BindingRemapper>();
 
     // D3D12 registers like `t3` and `c3` have the same bindingOffset number in
@@ -383,6 +379,7 @@ ResultOrError<std::string> TranslateToHLSL(
     }
 
     tint::writer::hlsl::Options options;
+    options.disable_robustness = !r.isRobustnessEnabled;
     options.disable_workgroup_init = r.disableWorkgroupInit;
     if (r.usesNumWorkgroups) {
         options.root_constant_binding_point =

src/dawn/native/metal/ShaderModuleMTL.mm

@@ -222,9 +222,6 @@ ResultOrError<CacheResult<MslCompilation>> TranslateToMSL(
                     std::move(r.substituteOverrideConfig).value());
             }
 
-            if (r.isRobustnessEnabled) {
-                transformManager.Add<tint::transform::Robustness>();
-            }
             transformManager.Add<BindingRemapper>();
             transformInputs.Add<BindingRemapper::Remappings>(std::move(r.bindingPoints),
                                                              BindingRemapper::AccessControls{},
@@ -262,6 +259,7 @@ ResultOrError<CacheResult<MslCompilation>> TranslateToMSL(
             }
 
             tint::writer::msl::Options options;
+            options.disable_robustness = !r.isRobustnessEnabled;
             options.buffer_size_ubo_index = kBufferLengthBufferSlot;
             options.fixed_sample_mask = r.sampleMask;
             options.disable_workgroup_init = r.disableWorkgroupInit;

src/dawn/native/null/DeviceNull.cpp

@@ -391,8 +391,6 @@ MaybeError ComputePipeline::Initialize() {
     tint::transform::Manager transformManager;
     tint::transform::DataMap transformInputs;
 
-    transformManager.Add<tint::transform::Robustness>();
-
     if (!computeStage.metadata->overrides.empty()) {
         transformManager.Add<tint::transform::SingleEntryPoint>();
         transformInputs.Add<tint::transform::SingleEntryPoint::Config>(

src/dawn/native/opengl/ShaderModuleGL.cpp

@@ -228,6 +228,9 @@ ResultOrError<GLuint> ShaderModule::CompileShader(const OpenGLFunctions& gl,
             tintOptions.version = tint::writer::glsl::Version(ToTintGLStandard(r.glVersionStandard),
                                                               r.glVersionMajor, r.glVersionMinor);
 
+            // TODO(crbug.com/dawn/1686): Robustness causes shader compilation failures.
+            tintOptions.disable_robustness = true;
+
             // When textures are accessed without a sampler (e.g., textureLoad()),
             // GetSamplerTextureUses() will return this sentinel value.
             BindingPoint placeholderBindingPoint{static_cast<uint32_t>(kMaxBindGroupsTyped), 0};

src/dawn/native/vulkan/ShaderModuleVk.cpp

@@ -293,10 +293,6 @@ ResultOrError<ShaderModule::ModuleAndSpirv> ShaderModule::GetHandleAndSpirv(
                     std::move(r.substituteOverrideConfig).value());
             }
 
-            if (r.isRobustnessEnabled) {
-                transformManager.append(std::make_unique<tint::transform::Robustness>());
-            }
-
             // Run the binding remapper after SingleEntryPoint to avoid collisions with
             // unused entryPoints.
             transformManager.append(std::make_unique<tint::transform::BindingRemapper>());
@@ -344,6 +340,7 @@ ResultOrError<ShaderModule::ModuleAndSpirv> ShaderModule::GetHandleAndSpirv(
             }
 
             tint::writer::spirv::Options options;
+            options.disable_robustness = !r.isRobustnessEnabled;
             options.emit_vertex_point_size = true;
             options.disable_workgroup_init = r.disableWorkgroupInit;
             options.use_zero_initialize_workgroup_memory_extension =

src/tint/cmd/main.cc

@@ -79,6 +79,7 @@ struct Options {
     bool print_hash = false;
     bool demangle = false;
     bool dump_inspector_bindings = false;
+    bool enable_robustness = false;
 
     std::unordered_set<uint32_t> skip_hash;
 
@@ -533,6 +534,7 @@ bool GenerateSpirv(const tint::Program* program, const Options& options) {
 #if TINT_BUILD_SPV_WRITER
     // TODO(jrprice): Provide a way for the user to set non-default options.
     tint::writer::spirv::Options gen_options;
+    gen_options.disable_robustness = !options.enable_robustness;
     gen_options.disable_workgroup_init = options.disable_workgroup_init;
     gen_options.generate_external_texture_bindings = true;
     auto result = tint::writer::spirv::Generate(program, gen_options);
@@ -639,6 +641,7 @@ bool GenerateMsl(const tint::Program* program, const Options& options) {
 
     // TODO(jrprice): Provide a way for the user to set non-default options.
     tint::writer::msl::Options gen_options;
+    gen_options.disable_robustness = !options.enable_robustness;
     gen_options.disable_workgroup_init = options.disable_workgroup_init;
     gen_options.generate_external_texture_bindings = true;
     auto result = tint::writer::msl::Generate(input_program, gen_options);
@@ -699,6 +702,7 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) {
 #if TINT_BUILD_HLSL_WRITER
     // TODO(jrprice): Provide a way for the user to set non-default options.
     tint::writer::hlsl::Options gen_options;
+    gen_options.disable_robustness = !options.enable_robustness;
     gen_options.disable_workgroup_init = options.disable_workgroup_init;
     gen_options.generate_external_texture_bindings = true;
     gen_options.root_constant_binding_point = options.hlsl_root_constant_binding_point;
@@ -838,6 +842,7 @@ bool GenerateGlsl(const tint::Program* program, const Options& options) {
 
     auto generate = [&](const tint::Program* prg, const std::string entry_point_name) -> bool {
         tint::writer::glsl::Options gen_options;
+        gen_options.disable_robustness = !options.enable_robustness;
         gen_options.generate_external_texture_bindings = true;
         auto result = tint::writer::glsl::Generate(prg, gen_options, entry_point_name);
         if (!result.success) {
@@ -946,8 +951,9 @@ int main(int argc, const char** argv) {
              return true;
          }},
         {"robustness",
-         [](tint::inspector::Inspector&, tint::transform::Manager& m, tint::transform::DataMap&) {
-             m.Add<tint::transform::Robustness>();
+         [&](tint::inspector::Inspector&, tint::transform::Manager&,
+             tint::transform::DataMap&) {  // enabled via writer option
+             options.enable_robustness = true;
              return true;
          }},
         {"substitute_override",

src/tint/fuzzers/tint_reader_writer_fuzzer.h

@@ -43,16 +43,14 @@ class ReaderWriterFuzzer : public CommonFuzzer {
         CommonFuzzer::SetTransformManager(tm, inputs);
     }
 
-    /// Pass through to the CommonFuzzer implementation, but will setup a
-    /// robustness transform, if no other transforms have been set.
-    /// @param data buffer of data that will interpreted as a byte array or string
-    ///             depending on the shader input format.
+    /// Pass through to the CommonFuzzer implementation.
+    /// @param data buffer of data that will interpreted as a byte array or string depending on the
+    /// shader input format.
     /// @param size number of elements in buffer
     /// @returns 0, this is what libFuzzer expects
     int Run(const uint8_t* data, size_t size) {
         if (!tm_set_) {
             tb_ = std::make_unique<TransformBuilder>(data, size);
-            tb_->AddTransform<tint::transform::Robustness>();
             SetTransformManager(tb_->manager(), tb_->data_map());
         }
 

src/tint/fuzzers/tint_robustness_fuzzer.cc

@@ -15,6 +15,7 @@
 #include "src/tint/fuzzers/fuzzer_init.h"
 #include "src/tint/fuzzers/tint_common_fuzzer.h"
 #include "src/tint/fuzzers/transform_builder.h"
+#include "src/tint/transform/robustness.h"
 
 namespace tint::fuzzers {
 

src/tint/fuzzers/transform_builder.h

@@ -22,6 +22,7 @@
 
 #include "src/tint/fuzzers/data_builder.h"
 #include "src/tint/fuzzers/shuffle_transform.h"
+#include "src/tint/transform/robustness.h"
 
 namespace tint::fuzzers {
 
@@ -61,7 +62,6 @@ class TransformBuilder {
     /// Helper that invokes Add*Transform for all of the platform independent
     /// passes.
     void AddPlatformIndependentPasses() {
-        AddTransform<transform::Robustness>();
         AddTransform<transform::FirstIndexOffset>();
         AddTransform<transform::BindingRemapper>();
         AddTransform<transform::Renamer>();

src/tint/transform/pad_structs.h

@@ -22,6 +22,8 @@ namespace tint::transform {
 /// This transform turns all explicit alignment and sizing into padding
 /// members of structs. This is required for GLSL ES, since it not support
 /// the offset= decoration.
+///
+/// @note This transform requires the CanonicalizeEntryPointIO transform to have been run first.
 class PadStructs final : public Castable<PadStructs, Transform> {
   public:
     /// Constructor

src/tint/transform/robustness.h

@@ -31,6 +31,7 @@ namespace tint::transform {
 /// the bounds of the array. Any access before the start of the array will clamp
 /// to zero and any access past the end of the array will clamp to
 /// (array length - 1).
+/// @note This transform must come before the BuiltinPolyfill transform
 class Robustness final : public Castable<Robustness, Transform> {
   public:
     /// Address space to be skipped in the transform

src/tint/writer/glsl/generator.h

@@ -63,6 +63,9 @@ struct Options {
     /// transform
     std::unordered_map<sem::BindingPoint, builtin::Access> access_controls;
 
+    /// Set to `true` to disable software robustness that prevents out-of-bounds accesses.
+    bool disable_robustness = false;
+
     /// If true, then validation will be disabled for binding point collisions
     /// generated by the BindingRemapper transform
     bool allow_collisions = false;
@@ -75,6 +78,13 @@ struct Options {
 
     /// The GLSL version to emit
     Version version;
+
+    /// Reflect the fields of this class so that it can be used by tint::ForeachField()
+    TINT_REFLECT(disable_robustness,
+                 allow_collisions,
+                 disable_workgroup_init,
+                 generate_external_texture_bindings,
+                 version);
 };
 
 /// The result produced when generating GLSL.

src/tint/writer/glsl/generator_impl.cc

@@ -57,6 +57,7 @@
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
 #include "src/tint/transform/renamer.h"
+#include "src/tint/transform/robustness.h"
 #include "src/tint/transform/simplify_pointers.h"
 #include "src/tint/transform/single_entry_point.h"
 #include "src/tint/transform/std140.h"
@@ -149,6 +150,33 @@ SanitizedResult Sanitize(const Program* in,
     // ExpandCompoundAssignment must come before BuiltinPolyfill
     manager.Add<transform::ExpandCompoundAssignment>();
 
+    if (!entry_point.empty()) {
+        manager.Add<transform::SingleEntryPoint>();
+        data.Add<transform::SingleEntryPoint::Config>(entry_point);
+    }
+    manager.Add<transform::Renamer>();
+    data.Add<transform::Renamer::Config>(transform::Renamer::Target::kGlslKeywords,
+                                         /* preserve_unicode */ false);
+
+    manager.Add<transform::PreservePadding>();  // Must come before DirectVariableAccess
+
+    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
+    manager.Add<transform::DirectVariableAccess>();
+
+    manager.Add<transform::PromoteSideEffectsToDecl>();
+
+    if (!options.disable_robustness) {
+        // Robustness must come before BuiltinPolyfill
+        manager.Add<transform::Robustness>();
+    }
+
+    if (options.generate_external_texture_bindings) {
+        // Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
+        auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
+        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+        manager.Add<transform::MultiplanarExternalTexture>();
+    }
+
     {  // Builtin polyfills
         transform::BuiltinPolyfill::Builtins polyfills;
         polyfills.acosh = transform::BuiltinPolyfill::Level::kRangeCheck;
@@ -169,26 +197,16 @@ SanitizedResult Sanitize(const Program* in,
         manager.Add<transform::BuiltinPolyfill>();
     }
 
-    if (!entry_point.empty()) {
-        manager.Add<transform::SingleEntryPoint>();
-        data.Add<transform::SingleEntryPoint::Config>(entry_point);
-    }
-    manager.Add<transform::Renamer>();
-    data.Add<transform::Renamer::Config>(transform::Renamer::Target::kGlslKeywords,
-                                         /* preserve_unicode */ false);
-
-    manager.Add<transform::PreservePadding>();  // Must come before DirectVariableAccess
-
-    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
-    manager.Add<transform::DirectVariableAccess>();
-
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
         // ZeroInitWorkgroupMemory may inject new builtin parameters.
         manager.Add<transform::ZeroInitWorkgroupMemory>();
     }
+
+    // CanonicalizeEntryPointIO must come after Robustness
     manager.Add<transform::CanonicalizeEntryPointIO>();
-    manager.Add<transform::PromoteSideEffectsToDecl>();
+
+    // PadStructs must come after CanonicalizeEntryPointIO
     manager.Add<transform::PadStructs>();
 
     // DemoteToHelper must come after PromoteSideEffectsToDecl and ExpandCompoundAssignment.
@@ -196,12 +214,6 @@ SanitizedResult Sanitize(const Program* in,
 
     manager.Add<transform::RemovePhonies>();
 
-    if (options.generate_external_texture_bindings) {
-        auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
-        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
-    }
-    manager.Add<transform::MultiplanarExternalTexture>();
-
     data.Add<transform::CombineSamplers::BindingInfo>(options.binding_map,
                                                       options.placeholder_binding_point);
     manager.Add<transform::CombineSamplers>();

src/tint/writer/glsl/test_helper.h

@@ -32,6 +32,14 @@ class TestHelperBase : public BODY, public ProgramBuilder {
     TestHelperBase() = default;
     ~TestHelperBase() override = default;
 
+    /// @returns the default generator options for SanitizeAndBuild(), if no explicit options are
+    /// provided.
+    static Options DefaultOptions() {
+        Options opts;
+        opts.disable_robustness = true;
+        return opts;
+    }
+
     /// Builds the program and returns a GeneratorImpl from the program.
     /// @note The generator is only built once. Multiple calls to Build() will
     /// return the same GeneratorImpl without rebuilding.
@@ -60,7 +68,8 @@ class TestHelperBase : public BODY, public ProgramBuilder {
     /// @param version the GLSL version
     /// @param options the GLSL backend options
     /// @return the built generator
-    GeneratorImpl& SanitizeAndBuild(Version version = Version(), const Options& options = {}) {
+    GeneratorImpl& SanitizeAndBuild(Version version = Version(),
+                                    const Options& options = DefaultOptions()) {
         if (gen_) {
             return *gen_;
         }

src/tint/writer/hlsl/generator.h

@@ -49,23 +49,32 @@ struct Options {
     /// @returns this Options
     Options& operator=(const Options&);
 
+    /// Set to `true` to disable software robustness that prevents out-of-bounds accesses.
+    bool disable_robustness = false;
+
     /// The binding point to use for information passed via root constants.
     std::optional<sem::BindingPoint> root_constant_binding_point;
+
     /// Set to `true` to disable workgroup memory zero initialization
     bool disable_workgroup_init = false;
+
     /// Set to 'true' to generates binding mappings for external textures
     bool generate_external_texture_bindings = false;
+
     /// Options used to specify a mapping of binding points to indices into a UBO
     /// from which to load buffer sizes.
     ArrayLengthFromUniformOptions array_length_from_uniform = {};
+
     /// Interstage locations actually used as inputs in the next stage of the pipeline.
     /// This is potentially used for truncating unused interstage outputs at current shader stage.
     std::bitset<16> interstage_locations;
+
     /// Set to `true` to generate polyfill for `reflect` builtin for vec2<f32>
     bool polyfill_reflect_vec2_f32 = false;
 
     /// Reflect the fields of this class so that it can be used by tint::ForeachField()
-    TINT_REFLECT(root_constant_binding_point,
+    TINT_REFLECT(disable_robustness,
+                 root_constant_binding_point,
                  disable_workgroup_init,
                  generate_external_texture_bindings,
                  array_length_from_uniform);

src/tint/writer/hlsl/generator_impl.cc

@@ -57,6 +57,7 @@
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_continue_in_switch.h"
 #include "src/tint/transform/remove_phonies.h"
+#include "src/tint/transform/robustness.h"
 #include "src/tint/transform/simplify_pointers.h"
 #include "src/tint/transform/truncate_interstage_variables.h"
 #include "src/tint/transform/unshadow.h"
@@ -165,6 +166,33 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
     // ExpandCompoundAssignment must come before BuiltinPolyfill
     manager.Add<transform::ExpandCompoundAssignment>();
 
+    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
+
+    manager.Add<transform::DirectVariableAccess>();
+
+    // LocalizeStructArrayAssignment must come after:
+    // * SimplifyPointers, because it assumes assignment to arrays in structs are
+    // done directly, not indirectly.
+    // TODO(crbug.com/tint/1340): See if we can get rid of the duplicate
+    // SimplifyPointers transform. Can't do it right now because
+    // LocalizeStructArrayAssignment introduces pointers.
+    manager.Add<transform::SimplifyPointers>();
+    manager.Add<transform::LocalizeStructArrayAssignment>();
+
+    manager.Add<transform::PromoteSideEffectsToDecl>();
+
+    if (!options.disable_robustness) {
+        // Robustness must come before BuiltinPolyfill
+        manager.Add<transform::Robustness>();
+    }
+
+    if (options.generate_external_texture_bindings) {
+        // Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
+        auto new_bindings_map = GenerateExternalTextureBindings(in);
+        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+        manager.Add<transform::MultiplanarExternalTexture>();
+    }
+
     {  // Builtin polyfills
         transform::BuiltinPolyfill::Builtins polyfills;
         polyfills.acosh = transform::BuiltinPolyfill::Level::kFull;
@@ -189,37 +217,13 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         manager.Add<transform::BuiltinPolyfill>();
     }
 
-    // Build the config for the internal ArrayLengthFromUniform transform.
-    auto& array_length_from_uniform = options.array_length_from_uniform;
-    transform::ArrayLengthFromUniform::Config array_length_from_uniform_cfg(
-        array_length_from_uniform.ubo_binding);
-    array_length_from_uniform_cfg.bindpoint_to_size_index =
-        array_length_from_uniform.bindpoint_to_size_index;
-
-    if (options.generate_external_texture_bindings) {
-        auto new_bindings_map = GenerateExternalTextureBindings(in);
-        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
-    }
-    manager.Add<transform::MultiplanarExternalTexture>();
-
-    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
-
-    manager.Add<transform::DirectVariableAccess>();
-
-    // LocalizeStructArrayAssignment must come after:
-    // * SimplifyPointers, because it assumes assignment to arrays in structs are
-    // done directly, not indirectly.
-    // TODO(crbug.com/tint/1340): See if we can get rid of the duplicate
-    // SimplifyPointers transform. Can't do it right now because
-    // LocalizeStructArrayAssignment introduces pointers.
-    manager.Add<transform::SimplifyPointers>();
-    manager.Add<transform::LocalizeStructArrayAssignment>();
-
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
         // ZeroInitWorkgroupMemory may inject new builtin parameters.
         manager.Add<transform::ZeroInitWorkgroupMemory>();
     }
+
+    // CanonicalizeEntryPointIO must come after Robustness
     manager.Add<transform::CanonicalizeEntryPointIO>();
 
     if (options.interstage_locations.any()) {
@@ -246,11 +250,17 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
     // assumes that num_workgroups builtins only appear as struct members and are
     // only accessed directly via member accessors.
     manager.Add<transform::NumWorkgroupsFromUniform>();
-    manager.Add<transform::PromoteSideEffectsToDecl>();
     manager.Add<transform::VectorizeScalarMatrixInitializers>();
     manager.Add<transform::SimplifyPointers>();
     manager.Add<transform::RemovePhonies>();
 
+    // Build the config for the internal ArrayLengthFromUniform transform.
+    auto& array_length_from_uniform = options.array_length_from_uniform;
+    transform::ArrayLengthFromUniform::Config array_length_from_uniform_cfg(
+        array_length_from_uniform.ubo_binding);
+    array_length_from_uniform_cfg.bindpoint_to_size_index =
+        array_length_from_uniform.bindpoint_to_size_index;
+
     // DemoteToHelper must come after CanonicalizeEntryPointIO, PromoteSideEffectsToDecl, and
     // ExpandCompoundAssignment.
     // TODO(crbug.com/tint/1752): This is only necessary when FXC is being used.

src/tint/writer/hlsl/test_helper.h

@@ -34,6 +34,14 @@ class TestHelperBase : public BODY, public ProgramBuilder {
     TestHelperBase() = default;
     ~TestHelperBase() override = default;
 
+    /// @returns the default generator options for SanitizeAndBuild(), if no explicit options are
+    /// provided.
+    static Options DefaultOptions() {
+        Options opts;
+        opts.disable_robustness = true;
+        return opts;
+    }
+
     /// Builds the program and returns a GeneratorImpl from the program.
     /// @note The generator is only built once. Multiple calls to Build() will
     /// return the same GeneratorImpl without rebuilding.
@@ -60,7 +68,7 @@ class TestHelperBase : public BODY, public ProgramBuilder {
     /// @note The generator is only built once. Multiple calls to Build() will
     /// return the same GeneratorImpl without rebuilding.
     /// @return the built generator
-    GeneratorImpl& SanitizeAndBuild(const Options& options = {}) {
+    GeneratorImpl& SanitizeAndBuild(const Options& options = DefaultOptions()) {
         if (gen_) {
             return *gen_;
         }

src/tint/writer/msl/generator.h

@@ -44,6 +44,9 @@ struct Options {
     /// @returns this Options
     Options& operator=(const Options&);
 
+    /// Set to `true` to disable software robustness that prevents out-of-bounds accesses.
+    bool disable_robustness = false;
+
     /// The index to use when generating a UBO to receive storage buffer sizes.
     /// Defaults to 30, which is the last valid buffer slot.
     uint32_t buffer_size_ubo_index = 30;
@@ -67,7 +70,8 @@ struct Options {
     ArrayLengthFromUniformOptions array_length_from_uniform = {};
 
     /// Reflect the fields of this class so that it can be used by tint::ForeachField()
-    TINT_REFLECT(buffer_size_ubo_index,
+    TINT_REFLECT(disable_robustness,
+                 buffer_size_ubo_index,
                  fixed_sample_mask,
                  emit_vertex_point_size,
                  disable_workgroup_init,

src/tint/writer/msl/generator_impl.cc

@@ -53,6 +53,7 @@
 #include "src/tint/transform/promote_initializers_to_let.h"
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
+#include "src/tint/transform/robustness.h"
 #include "src/tint/transform/simplify_pointers.h"
 #include "src/tint/transform/unshadow.h"
 #include "src/tint/transform/vectorize_scalar_matrix_initializers.h"
@@ -171,24 +172,6 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
     // ExpandCompoundAssignment must come before BuiltinPolyfill
     manager.Add<transform::ExpandCompoundAssignment>();
 
-    {  // Builtin polyfills
-        transform::BuiltinPolyfill::Builtins polyfills;
-        polyfills.acosh = transform::BuiltinPolyfill::Level::kRangeCheck;
-        polyfills.atanh = transform::BuiltinPolyfill::Level::kRangeCheck;
-        polyfills.bitshift_modulo = true;  // crbug.com/tint/1543
-        polyfills.clamp_int = true;
-        polyfills.extract_bits = transform::BuiltinPolyfill::Level::kClampParameters;
-        polyfills.first_leading_bit = true;
-        polyfills.first_trailing_bit = true;
-        polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
-        polyfills.int_div_mod = true;
-        polyfills.sign_int = true;
-        polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
-        polyfills.workgroup_uniform_load = true;
-        data.Add<transform::BuiltinPolyfill::Config>(polyfills);
-        manager.Add<transform::BuiltinPolyfill>();
-    }
-
     // Build the config for the internal ArrayLengthFromUniform transform.
     auto& array_length_from_uniform = options.array_length_from_uniform;
     transform::ArrayLengthFromUniform::Config array_length_from_uniform_cfg(
@@ -217,23 +200,51 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         transform::CanonicalizeEntryPointIO::ShaderStyle::kMsl, options.fixed_sample_mask,
         options.emit_vertex_point_size);
 
-    if (options.generate_external_texture_bindings) {
-        auto new_bindings_map = GenerateExternalTextureBindings(in);
-        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
-    }
-    manager.Add<transform::MultiplanarExternalTexture>();
-
     manager.Add<transform::PreservePadding>();
 
     manager.Add<transform::Unshadow>();
 
+    manager.Add<transform::PromoteSideEffectsToDecl>();
+
+    if (!options.disable_robustness) {
+        // Robustness must come before BuiltinPolyfill
+        manager.Add<transform::Robustness>();
+    }
+
+    {  // Builtin polyfills
+        transform::BuiltinPolyfill::Builtins polyfills;
+        polyfills.acosh = transform::BuiltinPolyfill::Level::kRangeCheck;
+        polyfills.atanh = transform::BuiltinPolyfill::Level::kRangeCheck;
+        polyfills.bitshift_modulo = true;  // crbug.com/tint/1543
+        polyfills.clamp_int = true;
+        polyfills.extract_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+        polyfills.first_leading_bit = true;
+        polyfills.first_trailing_bit = true;
+        polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+        polyfills.int_div_mod = true;
+        polyfills.sign_int = true;
+        polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
+        polyfills.workgroup_uniform_load = true;
+        data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+        manager.Add<transform::BuiltinPolyfill>();
+    }
+
+    if (options.generate_external_texture_bindings) {
+        // Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
+        auto new_bindings_map = GenerateExternalTextureBindings(in);
+        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+        manager.Add<transform::MultiplanarExternalTexture>();
+    }
+
     if (!options.disable_workgroup_init) {
         // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
         // ZeroInitWorkgroupMemory may inject new builtin parameters.
         manager.Add<transform::ZeroInitWorkgroupMemory>();
     }
+
+    // CanonicalizeEntryPointIO must come after Robustness
     manager.Add<transform::CanonicalizeEntryPointIO>();
-    manager.Add<transform::PromoteSideEffectsToDecl>();
+
     manager.Add<transform::PromoteInitializersToLet>();
 
     // DemoteToHelper must come after PromoteSideEffectsToDecl and ExpandCompoundAssignment.

src/tint/writer/msl/test_helper.h

@@ -33,6 +33,14 @@ class TestHelperBase : public BASE, public ProgramBuilder {
     TestHelperBase() = default;
     ~TestHelperBase() override = default;
 
+    /// @returns the default generator options for SanitizeAndBuild(), if no explicit options are
+    /// provided.
+    static Options DefaultOptions() {
+        Options opts;
+        opts.disable_robustness = true;
+        return opts;
+    }
+
     /// Builds and returns a GeneratorImpl from the program.
     /// @note The generator is only built once. Multiple calls to Build() will
     /// return the same GeneratorImpl without rebuilding.
@@ -59,7 +67,7 @@ class TestHelperBase : public BASE, public ProgramBuilder {
     /// @note The generator is only built once. Multiple calls to Build() will
     /// return the same GeneratorImpl without rebuilding.
     /// @return the built generator
-    GeneratorImpl& SanitizeAndBuild(const Options& options = {}) {
+    GeneratorImpl& SanitizeAndBuild(const Options& options = DefaultOptions()) {
         if (gen_) {
             return *gen_;
         }

src/tint/writer/spirv/builder_builtin_test.cc

@@ -2107,7 +2107,7 @@ OpEntryPoint Fragment %24 "a_func"
 OpExecutionMode %24 OriginUpperLeft
 OpName %5 "v"
 OpName %8 "tint_quantizeToF16"
-OpName %9 "v_1"
+OpName %9 "v"
 OpName %24 "a_func"
 %2 = OpTypeFloat 32
 %1 = OpTypeVector %2 3

src/tint/writer/spirv/generator.h

@@ -35,6 +35,9 @@ namespace tint::writer::spirv {
 
 /// Configuration options used for generating SPIR-V.
 struct Options {
+    /// Set to `true` to disable software robustness that prevents out-of-bounds accesses.
+    bool disable_robustness = false;
+
     /// Set to `true` to generate a PointSize builtin and have it set to 1.0
     /// from all vertex shaders in the module.
     bool emit_vertex_point_size = true;
@@ -50,7 +53,8 @@ struct Options {
     bool use_zero_initialize_workgroup_memory_extension = false;
 
     /// Reflect the fields of this class so that it can be used by tint::ForeachField()
-    TINT_REFLECT(emit_vertex_point_size,
+    TINT_REFLECT(disable_robustness,
+                 emit_vertex_point_size,
                  disable_workgroup_init,
                  generate_external_texture_bindings,
                  use_zero_initialize_workgroup_memory_extension);

src/tint/writer/spirv/generator_impl.cc

@@ -32,6 +32,7 @@
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
 #include "src/tint/transform/remove_unreachable_statements.h"
+#include "src/tint/transform/robustness.h"
 #include "src/tint/transform/simplify_pointers.h"
 #include "src/tint/transform/std140.h"
 #include "src/tint/transform/unshadow.h"
@@ -53,7 +54,34 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
     // ExpandCompoundAssignment must come before BuiltinPolyfill
     manager.Add<transform::ExpandCompoundAssignment>();
 
+    manager.Add<transform::PreservePadding>();  // Must come before DirectVariableAccess
+
+    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
+
+    manager.Add<transform::RemoveUnreachableStatements>();
+    manager.Add<transform::PromoteSideEffectsToDecl>();
+    manager.Add<transform::SimplifyPointers>();  // Required for arrayLength()
+    manager.Add<transform::RemovePhonies>();
+    manager.Add<transform::VectorizeScalarMatrixInitializers>();
+    manager.Add<transform::VectorizeMatrixConversions>();
+    manager.Add<transform::WhileToLoop>();  // ZeroInitWorkgroupMemory
+    manager.Add<transform::MergeReturn>();
+
+    if (!options.disable_robustness) {
+        // Robustness must come before BuiltinPolyfill
+        manager.Add<transform::Robustness>();
+    }
+
+    if (options.generate_external_texture_bindings) {
+        // Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
+        auto new_bindings_map = GenerateExternalTextureBindings(in);
+        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+        manager.Add<transform::MultiplanarExternalTexture>();
+    }
+
     {  // Builtin polyfills
+        // BuiltinPolyfill must come before DirectVariableAccess, due to the use of pointer
+        // parameter for workgroupUniformLoad()
         transform::BuiltinPolyfill::Builtins polyfills;
         polyfills.acosh = transform::BuiltinPolyfill::Level::kRangeCheck;
         polyfills.atanh = transform::BuiltinPolyfill::Level::kRangeCheck;
@@ -75,18 +103,11 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         manager.Add<transform::BuiltinPolyfill>();
     }
 
-    if (options.generate_external_texture_bindings) {
-        auto new_bindings_map = GenerateExternalTextureBindings(in);
-        data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
-    }
-    manager.Add<transform::MultiplanarExternalTexture>();
-
-    manager.Add<transform::PreservePadding>();  // Must come before DirectVariableAccess
-
-    manager.Add<transform::Unshadow>();  // Must come before DirectVariableAccess
     bool disable_workgroup_init_in_sanitizer =
         options.disable_workgroup_init || options.use_zero_initialize_workgroup_memory_extension;
     if (!disable_workgroup_init_in_sanitizer) {
+        // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
+        // ZeroInitWorkgroupMemory may inject new builtin parameters.
         manager.Add<transform::ZeroInitWorkgroupMemory>();
     }
 
@@ -98,16 +119,11 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         manager.Add<transform::DirectVariableAccess>();
     }
 
-    manager.Add<transform::RemoveUnreachableStatements>();
-    manager.Add<transform::PromoteSideEffectsToDecl>();
-    manager.Add<transform::SimplifyPointers>();  // Required for arrayLength()
-    manager.Add<transform::RemovePhonies>();
-    manager.Add<transform::VectorizeScalarMatrixInitializers>();
-    manager.Add<transform::VectorizeMatrixConversions>();
-    manager.Add<transform::WhileToLoop>();  // ZeroInitWorkgroupMemory
-    manager.Add<transform::MergeReturn>();
+    // CanonicalizeEntryPointIO must come after Robustness
     manager.Add<transform::CanonicalizeEntryPointIO>();
     manager.Add<transform::AddEmptyEntryPoint>();
+
+    // AddBlockAttribute must come after MultiplanarExternalTexture
     manager.Add<transform::AddBlockAttribute>();
 
     // DemoteToHelper must come after CanonicalizeEntryPointIO, PromoteSideEffectsToDecl, and

src/tint/writer/spirv/test_helper.h

@@ -33,6 +33,14 @@ class TestHelperBase : public ProgramBuilder, public BASE {
     TestHelperBase() = default;
     ~TestHelperBase() override = default;
 
+    /// @returns the default generator options for SanitizeAndBuild(), if no explicit options are
+    /// provided.
+    static Options DefaultOptions() {
+        Options opts;
+        opts.disable_robustness = true;
+        return opts;
+    }
+
     /// Builds and returns a spirv::Builder from the program.
     /// @note The spirv::Builder is only built once. Multiple calls to Build()
     /// will return the same spirv::Builder without rebuilding.
@@ -59,7 +67,7 @@ class TestHelperBase : public ProgramBuilder, public BASE {
     /// @note The spirv::Builder is only built once. Multiple calls to Build()
     /// will return the same spirv::Builder without rebuilding.
     /// @return the built spirv::Builder
-    spirv::Builder& SanitizeAndBuild(const Options& options = {}) {
+    spirv::Builder& SanitizeAndBuild(const Options& options = DefaultOptions()) {
         if (spirv_builder) {
             return *spirv_builder;
         }

test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl

@@ -215,12 +215,12 @@ kernel void tint_symbol(device S_tint_packed_vec3* tint_symbol_5 [[buffer(0)]])
   (*(tint_symbol_5)).mat4x4_f16 = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
   tint_array<float3, 2> const tint_symbol_1 = tint_array<float3, 2>{};
   assign_and_preserve_padding_6(&((*(tint_symbol_5)).arr2_vec3_f32), tint_symbol_1);
-  tint_array<half4x2, 2> const tint_symbol_4 = tint_array<half4x2, 2>{};
-  (*(tint_symbol_5)).arr2_mat4x2_f16 = tint_symbol_4;
-  Inner const tint_symbol_2 = Inner{};
-  assign_and_preserve_padding_7(&((*(tint_symbol_5)).struct_inner), tint_symbol_2);
-  tint_array<Inner, 4> const tint_symbol_3 = tint_array<Inner, 4>{};
-  assign_and_preserve_padding_8(&((*(tint_symbol_5)).array_struct_inner), tint_symbol_3);
+  tint_array<half4x2, 2> const tint_symbol_2 = tint_array<half4x2, 2>{};
+  (*(tint_symbol_5)).arr2_mat4x2_f16 = tint_symbol_2;
+  Inner const tint_symbol_3 = Inner{};
+  assign_and_preserve_padding_7(&((*(tint_symbol_5)).struct_inner), tint_symbol_3);
+  tint_array<Inner, 4> const tint_symbol_4 = tint_array<Inner, 4>{};
+  assign_and_preserve_padding_8(&((*(tint_symbol_5)).array_struct_inner), tint_symbol_4);
   return;
 }
 

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

@@ -7,9 +7,7 @@ struct tint_symbol_1 {
 };
 
 void f_inner(uint3 v) {
-  const uint tint_symbol_2 = v.x;
-  const uint tint_symbol_3 = tint_mod(v.y, 1u);
-  const uint l = (tint_symbol_2 << (tint_symbol_3 & 31u));
+  const uint l = (v.x << (tint_mod(v.y, 1u) & 31u));
 }
 
 [numthreads(1, 1, 1)]

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

@@ -7,9 +7,7 @@ struct tint_symbol_1 {
 };
 
 void f_inner(uint3 v) {
-  const uint tint_symbol_2 = v.x;
-  const uint tint_symbol_3 = tint_mod(v.y, 1u);
-  const uint l = (tint_symbol_2 << (tint_symbol_3 & 31u));
+  const uint l = (v.x << (tint_mod(v.y, 1u) & 31u));
 }
 
 [numthreads(1, 1, 1)]

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

@@ -5,9 +5,7 @@ uint tint_mod(uint lhs, uint rhs) {
 }
 
 void f(uvec3 v) {
-  uint tint_symbol = v.x;
-  uint tint_symbol_1 = tint_mod(v.y, 1u);
-  uint l = (tint_symbol << (tint_symbol_1 & 31u));
+  uint l = (v.x << (tint_mod(v.y, 1u) & 31u));
 }
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

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

@@ -6,9 +6,7 @@ uint tint_mod(uint lhs, uint rhs) {
 }
 
 void f_inner(uint3 v) {
-  uint const tint_symbol = v[0];
-  uint const tint_symbol_1 = tint_mod(v[1], 1u);
-  uint const l = (tint_symbol << (tint_symbol_1 & 31u));
+  uint const l = (v[0] << (tint_mod(v[1], 1u) & 31u));
 }
 
 kernel void f(uint3 v [[thread_position_in_grid]]) {

test/tint/bug/tint/1076.wgsl.expected.dxc.hlsl

@@ -16,13 +16,13 @@ FragIn main_inner(FragIn tint_symbol, float b) {
   if ((tint_symbol.mask == 0u)) {
     return tint_symbol;
   }
-  const FragIn tint_symbol_4 = {b, 1u};
-  return tint_symbol_4;
+  const FragIn tint_symbol_5 = {b, 1u};
+  return tint_symbol_5;
 }
 
 tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
-  const FragIn tint_symbol_5 = {tint_symbol_1.a, tint_symbol_1.mask};
-  const FragIn inner_result = main_inner(tint_symbol_5, tint_symbol_1.b);
+  const FragIn tint_symbol_4 = {tint_symbol_1.a, tint_symbol_1.mask};
+  const FragIn inner_result = main_inner(tint_symbol_4, tint_symbol_1.b);
   tint_symbol_3 wrapper_result = (tint_symbol_3)0;
   wrapper_result.a = inner_result.a;
   wrapper_result.mask = inner_result.mask;

test/tint/bug/tint/1076.wgsl.expected.fxc.hlsl

@@ -16,13 +16,13 @@ FragIn main_inner(FragIn tint_symbol, float b) {
   if ((tint_symbol.mask == 0u)) {
     return tint_symbol;
   }
-  const FragIn tint_symbol_4 = {b, 1u};
-  return tint_symbol_4;
+  const FragIn tint_symbol_5 = {b, 1u};
+  return tint_symbol_5;
 }
 
 tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
-  const FragIn tint_symbol_5 = {tint_symbol_1.a, tint_symbol_1.mask};
-  const FragIn inner_result = main_inner(tint_symbol_5, tint_symbol_1.b);
+  const FragIn tint_symbol_4 = {tint_symbol_1.a, tint_symbol_1.mask};
+  const FragIn inner_result = main_inner(tint_symbol_4, tint_symbol_1.b);
   tint_symbol_3 wrapper_result = (tint_symbol_3)0;
   wrapper_result.a = inner_result.a;
   wrapper_result.mask = inner_result.mask;

test/tint/bug/tint/1076.wgsl.expected.msl

@@ -20,13 +20,13 @@ FragIn tint_symbol_inner(FragIn in, float b) {
   if ((in.mask == 0u)) {
     return in;
   }
-  FragIn const tint_symbol_4 = {.a=b, .mask=1u};
-  return tint_symbol_4;
+  FragIn const tint_symbol_5 = {.a=b, .mask=1u};
+  return tint_symbol_5;
 }
 
 fragment tint_symbol_3 tint_symbol(uint mask [[sample_mask]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
-  FragIn const tint_symbol_5 = {.a=tint_symbol_1.a, .mask=mask};
-  FragIn const inner_result = tint_symbol_inner(tint_symbol_5, tint_symbol_1.b);
+  FragIn const tint_symbol_4 = {.a=tint_symbol_1.a, .mask=mask};
+  FragIn const inner_result = tint_symbol_inner(tint_symbol_4, tint_symbol_1.b);
   tint_symbol_3 wrapper_result = {};
   wrapper_result.a = inner_result.a;
   wrapper_result.mask = inner_result.mask;

test/tint/bug/tint/1113.wgsl.expected.dxc.hlsl

@@ -67,7 +67,7 @@ void doIgnore() {
   int g55 = LUTatomicLoad(0u);
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -103,12 +103,12 @@ void main_count_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(128, 1, 1)]
-void main_count(tint_symbol_1 tint_symbol) {
-  main_count_inner(tint_symbol.GlobalInvocationID);
+void main_count(tint_symbol_2 tint_symbol_1) {
+  main_count_inner(tint_symbol_1.GlobalInvocationID);
   return;
 }
 
-struct tint_symbol_3 {
+struct tint_symbol_4 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -135,19 +135,19 @@ void main_create_lut_inner(uint3 GlobalInvocationID) {
   uint numTriangles = countersatomicLoad((4u * voxelIndex));
   int offset = -1;
   if ((numTriangles > 0u)) {
-    const uint tint_symbol_6 = dbgatomicAdd(0u, numTriangles);
-    offset = int(tint_symbol_6);
+    const uint tint_symbol = dbgatomicAdd(0u, numTriangles);
+    offset = int(tint_symbol);
   }
   LUTatomicStore((4u * voxelIndex), offset);
 }
 
 [numthreads(128, 1, 1)]
-void main_create_lut(tint_symbol_3 tint_symbol_2) {
-  main_create_lut_inner(tint_symbol_2.GlobalInvocationID);
+void main_create_lut(tint_symbol_4 tint_symbol_3) {
+  main_create_lut_inner(tint_symbol_3.GlobalInvocationID);
   return;
 }
 
-struct tint_symbol_5 {
+struct tint_symbol_6 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -177,7 +177,7 @@ void main_sort_triangles_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(128, 1, 1)]
-void main_sort_triangles(tint_symbol_5 tint_symbol_4) {
-  main_sort_triangles_inner(tint_symbol_4.GlobalInvocationID);
+void main_sort_triangles(tint_symbol_6 tint_symbol_5) {
+  main_sort_triangles_inner(tint_symbol_5.GlobalInvocationID);
   return;
 }

test/tint/bug/tint/1113.wgsl.expected.fxc.hlsl

@@ -67,7 +67,7 @@ void doIgnore() {
   int g55 = LUTatomicLoad(0u);
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -103,12 +103,12 @@ void main_count_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(128, 1, 1)]
-void main_count(tint_symbol_1 tint_symbol) {
-  main_count_inner(tint_symbol.GlobalInvocationID);
+void main_count(tint_symbol_2 tint_symbol_1) {
+  main_count_inner(tint_symbol_1.GlobalInvocationID);
   return;
 }
 
-struct tint_symbol_3 {
+struct tint_symbol_4 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -135,19 +135,19 @@ void main_create_lut_inner(uint3 GlobalInvocationID) {
   uint numTriangles = countersatomicLoad((4u * voxelIndex));
   int offset = -1;
   if ((numTriangles > 0u)) {
-    const uint tint_symbol_6 = dbgatomicAdd(0u, numTriangles);
-    offset = int(tint_symbol_6);
+    const uint tint_symbol = dbgatomicAdd(0u, numTriangles);
+    offset = int(tint_symbol);
   }
   LUTatomicStore((4u * voxelIndex), offset);
 }
 
 [numthreads(128, 1, 1)]
-void main_create_lut(tint_symbol_3 tint_symbol_2) {
-  main_create_lut_inner(tint_symbol_2.GlobalInvocationID);
+void main_create_lut(tint_symbol_4 tint_symbol_3) {
+  main_create_lut_inner(tint_symbol_3.GlobalInvocationID);
   return;
 }
 
-struct tint_symbol_5 {
+struct tint_symbol_6 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -177,7 +177,7 @@ void main_sort_triangles_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(128, 1, 1)]
-void main_sort_triangles(tint_symbol_5 tint_symbol_4) {
-  main_sort_triangles_inner(tint_symbol_4.GlobalInvocationID);
+void main_sort_triangles(tint_symbol_6 tint_symbol_5) {
+  main_sort_triangles_inner(tint_symbol_5.GlobalInvocationID);
   return;
 }

test/tint/bug/tint/1520.spvasm.expected.dxc.hlsl

@@ -25,8 +25,7 @@ bool test_int_S1_c0_b() {
   ok = true;
   x_41 = false;
   if (true) {
-    const int4 tint_symbol_3 = tint_div((0).xxxx, int4(x_27, x_27, x_27, x_27));
-    x_40 = all((tint_symbol_3 == (0).xxxx));
+    x_40 = all((tint_div((0).xxxx, int4(x_27, x_27, x_27, x_27)) == (0).xxxx));
     x_41 = x_40;
   }
   ok = x_41;
@@ -152,8 +151,8 @@ main_out main_inner(bool sk_Clockwise_param, float4 vcolor_S0_param) {
   sk_Clockwise = sk_Clockwise_param;
   vcolor_S0 = vcolor_S0_param;
   main_1();
-  const main_out tint_symbol_4 = {sk_FragColor};
-  return tint_symbol_4;
+  const main_out tint_symbol_3 = {sk_FragColor};
+  return tint_symbol_3;
 }
 
 tint_symbol_2 main(tint_symbol_1 tint_symbol) {

test/tint/bug/tint/1520.spvasm.expected.fxc.hlsl

@@ -25,8 +25,7 @@ bool test_int_S1_c0_b() {
   ok = true;
   x_41 = false;
   if (true) {
-    const int4 tint_symbol_3 = tint_div((0).xxxx, int4(x_27, x_27, x_27, x_27));
-    x_40 = all((tint_symbol_3 == (0).xxxx));
+    x_40 = all((tint_div((0).xxxx, int4(x_27, x_27, x_27, x_27)) == (0).xxxx));
     x_41 = x_40;
   }
   ok = x_41;
@@ -152,8 +151,8 @@ main_out main_inner(bool sk_Clockwise_param, float4 vcolor_S0_param) {
   sk_Clockwise = sk_Clockwise_param;
   vcolor_S0 = vcolor_S0_param;
   main_1();
-  const main_out tint_symbol_4 = {sk_FragColor};
-  return tint_symbol_4;
+  const main_out tint_symbol_3 = {sk_FragColor};
+  return tint_symbol_3;
 }
 
 tint_symbol_2 main(tint_symbol_1 tint_symbol) {

test/tint/bug/tint/1520.spvasm.expected.glsl

@@ -44,8 +44,7 @@ bool test_int_S1_c0_b() {
   ok = true;
   x_41 = false;
   if (true) {
-    ivec4 tint_symbol_1 = tint_div(ivec4(0), ivec4(x_27, x_27, x_27, x_27));
-    x_40 = all(equal(tint_symbol_1, ivec4(0)));
+    x_40 = all(equal(tint_div(ivec4(0), ivec4(x_27, x_27, x_27, x_27)), ivec4(0)));
     x_41 = x_40;
   }
   ok = x_41;
@@ -164,8 +163,8 @@ main_out tint_symbol(bool sk_Clockwise_param, vec4 vcolor_S0_param) {
   sk_Clockwise = sk_Clockwise_param;
   vcolor_S0 = vcolor_S0_param;
   main_1();
-  main_out tint_symbol_2 = main_out(sk_FragColor);
-  return tint_symbol_2;
+  main_out tint_symbol_1 = main_out(sk_FragColor);
+  return tint_symbol_1;
 }
 
 void main() {

test/tint/bug/tint/1520.spvasm.expected.msl

@@ -39,7 +39,7 @@ int4 tint_div(int4 lhs, int4 rhs) {
   return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
 }
 
-bool test_int_S1_c0_b(const constant UniformBuffer_tint_packed_vec3* const tint_symbol_6) {
+bool test_int_S1_c0_b(const constant UniformBuffer_tint_packed_vec3* const tint_symbol_5) {
   int unknown = 0;
   bool ok = false;
   int4 val = 0;
@@ -49,14 +49,13 @@ bool test_int_S1_c0_b(const constant UniformBuffer_tint_packed_vec3* const tint_
   bool x_55 = false;
   bool x_65 = false;
   bool x_66 = false;
-  float const x_26 = (*(tint_symbol_6)).unknownInput_S1_c0;
+  float const x_26 = (*(tint_symbol_5)).unknownInput_S1_c0;
   int const x_27 = int(x_26);
   unknown = x_27;
   ok = true;
   x_41 = false;
   if (true) {
-    int4 const tint_symbol_4 = tint_div(int4(0), int4(x_27, x_27, x_27, x_27));
-    x_40 = all((tint_symbol_4 == int4(0)));
+    x_40 = all((tint_div(int4(0), int4(x_27, x_27, x_27, x_27)) == int4(0)));
     x_41 = x_40;
   }
   ok = x_41;
@@ -93,7 +92,7 @@ bool test_int_S1_c0_b(const constant UniformBuffer_tint_packed_vec3* const tint_
   return x_66;
 }
 
-void main_1(thread float4* const tint_symbol_7, const constant UniformBuffer_tint_packed_vec3* const tint_symbol_8, thread float4* const tint_symbol_9) {
+void main_1(thread float4* const tint_symbol_6, const constant UniformBuffer_tint_packed_vec3* const tint_symbol_7, thread float4* const tint_symbol_8) {
   float4 outputColor_S0 = 0.0f;
   float4 output_S1 = 0.0f;
   float x_8_unknown = 0.0f;
@@ -108,9 +107,9 @@ void main_1(thread float4* const tint_symbol_7, const constant UniformBuffer_tin
   bool x_111 = false;
   bool x_114 = false;
   bool x_115 = false;
-  float4 const x_72 = *(tint_symbol_7);
+  float4 const x_72 = *(tint_symbol_6);
   outputColor_S0 = x_72;
-  float const x_77 = (*(tint_symbol_8)).unknownInput_S1_c0;
+  float const x_77 = (*(tint_symbol_7)).unknownInput_S1_c0;
   x_8_unknown = x_77;
   x_9_ok = true;
   x_87 = false;
@@ -151,19 +150,19 @@ void main_1(thread float4* const tint_symbol_7, const constant UniformBuffer_tin
   x_9_ok = x_111;
   x_115 = false;
   if (x_111) {
-    x_114 = test_int_S1_c0_b(tint_symbol_8);
+    x_114 = test_int_S1_c0_b(tint_symbol_7);
     x_115 = x_114;
   }
   if (x_115) {
-    float4 const x_122 = (*(tint_symbol_8)).ucolorGreen_S1_c0;
+    float4 const x_122 = (*(tint_symbol_7)).ucolorGreen_S1_c0;
     x_116 = x_122;
   } else {
-    float4 const x_124 = (*(tint_symbol_8)).ucolorRed_S1_c0;
+    float4 const x_124 = (*(tint_symbol_7)).ucolorRed_S1_c0;
     x_116 = x_124;
   }
   float4 const x_125 = x_116;
   output_S1 = x_125;
-  *(tint_symbol_9) = x_125;
+  *(tint_symbol_8) = x_125;
   return;
 }
 
@@ -179,19 +178,19 @@ struct tint_symbol_3 {
   float4 sk_FragColor_1 [[color(0)]];
 };
 
-main_out tint_symbol_inner(bool sk_Clockwise_param, float4 vcolor_S0_param, thread float4* const tint_symbol_11, const constant UniformBuffer_tint_packed_vec3* const tint_symbol_12, thread float4* const tint_symbol_13) {
-  thread bool tint_symbol_10 = false;
-  tint_symbol_10 = sk_Clockwise_param;
-  *(tint_symbol_11) = vcolor_S0_param;
-  main_1(tint_symbol_11, tint_symbol_12, tint_symbol_13);
-  main_out const tint_symbol_5 = {.sk_FragColor_1=*(tint_symbol_13)};
-  return tint_symbol_5;
+main_out tint_symbol_inner(bool sk_Clockwise_param, float4 vcolor_S0_param, thread float4* const tint_symbol_10, const constant UniformBuffer_tint_packed_vec3* const tint_symbol_11, thread float4* const tint_symbol_12) {
+  thread bool tint_symbol_9 = false;
+  tint_symbol_9 = sk_Clockwise_param;
+  *(tint_symbol_10) = vcolor_S0_param;
+  main_1(tint_symbol_10, tint_symbol_11, tint_symbol_12);
+  main_out const tint_symbol_4 = {.sk_FragColor_1=*(tint_symbol_12)};
+  return tint_symbol_4;
 }
 
-fragment tint_symbol_3 tint_symbol(const constant UniformBuffer_tint_packed_vec3* tint_symbol_15 [[buffer(0)]], bool sk_Clockwise_param [[front_facing]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
-  thread float4 tint_symbol_14 = 0.0f;
-  thread float4 tint_symbol_16 = 0.0f;
-  main_out const inner_result = tint_symbol_inner(sk_Clockwise_param, tint_symbol_1.vcolor_S0_param, &(tint_symbol_14), tint_symbol_15, &(tint_symbol_16));
+fragment tint_symbol_3 tint_symbol(const constant UniformBuffer_tint_packed_vec3* tint_symbol_14 [[buffer(0)]], bool sk_Clockwise_param [[front_facing]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  thread float4 tint_symbol_13 = 0.0f;
+  thread float4 tint_symbol_15 = 0.0f;
+  main_out const inner_result = tint_symbol_inner(sk_Clockwise_param, tint_symbol_1.vcolor_S0_param, &(tint_symbol_13), tint_symbol_14, &(tint_symbol_15));
   tint_symbol_3 wrapper_result = {};
   wrapper_result.sk_FragColor_1 = inner_result.sk_FragColor_1;
   return wrapper_result;

test/tint/bug/tint/1520.spvasm.expected.spvasm

@@ -161,11 +161,11 @@
                OpSelectionMerge %68 None
                OpBranchConditional %true %69 %68
          %69 = OpLabel
-         %71 = OpCompositeConstruct %v4int %66 %66 %66 %66
-         %70 = OpFunctionCall %v4int %tint_div %31 %71
-         %73 = OpIEqual %v4bool %70 %31
-         %72 = OpAll %bool %73
-               OpStore %x_40 %72
+         %72 = OpCompositeConstruct %v4int %66 %66 %66 %66
+         %71 = OpFunctionCall %v4int %tint_div %31 %72
+         %73 = OpIEqual %v4bool %71 %31
+         %70 = OpAll %bool %73
+               OpStore %x_40 %70
          %74 = OpLoad %bool %x_40
                OpStore %x_41 %74
                OpBranch %68

test/tint/bug/tint/1563.wgsl.expected.wgsl

@@ -1,7 +1,7 @@
 fn foo() -> f32 {
   let oob = 99;
-  let b = vec4<f32>()[min(u32(oob), 3u)];
+  let b = vec4<f32>()[oob];
   var v : vec4<f32>;
-  v[min(u32(oob), 3u)] = b;
+  v[oob] = b;
   return b;
 }

test/tint/bug/tint/1725.wgsl.expected.wgsl

@@ -5,5 +5,5 @@ fn tint_symbol_1(@builtin(local_invocation_index) tint_symbol_2 : u32) {
   let tint_symbol_3 = 0;
   let tint_symbol_4 = 0;
   let tint_symbol_5 = 0;
-  let tint_symbol_6 = tint_symbol[min(tint_symbol_2, (arrayLength(&(tint_symbol)) - 1u))];
+  let tint_symbol_6 = tint_symbol[tint_symbol_2];
 }

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

@@ -1,4 +1,8 @@
-int2 tint_clamp(int2 e, int2 low, int2 high) {
+int tint_clamp(int e, int low, int high) {
+  return min(max(e, low), high);
+}
+
+int2 tint_clamp_1(int2 e, int2 low, int2 high) {
   return min(max(e, low), high);
 }
 
@@ -40,9 +44,27 @@ float4 textureLoadExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, i
   const int2 coord1 = (coord >> (1u).xx);
   float3 color = float3(0.0f, 0.0f, 0.0f);
   if ((params.numPlanes == 1u)) {
-    color = plane0.Load(int3(coord, 0)).rgb;
+    int3 tint_tmp;
+    int3 tint_tmp_1;
+    plane0.GetDimensions(0, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
+    plane0.GetDimensions(tint_clamp(0, 0, int((uint(tint_tmp_1.z) - 1u))), tint_tmp.x, tint_tmp.y, tint_tmp.z);
+    int3 tint_tmp_2;
+    plane0.GetDimensions(0, tint_tmp_2.x, tint_tmp_2.y, tint_tmp_2.z);
+    color = plane0.Load(int3(tint_clamp_1(coord, (0).xx, int2((uint2(tint_tmp.xy) - (1u).xx))), tint_clamp(0, 0, int((uint(tint_tmp_2.z) - 1u))))).rgb;
   } else {
-    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(coord, 0)).r, plane1.Load(int3(coord1, 0)).rg, 1.0f));
+    int3 tint_tmp_3;
+    int3 tint_tmp_4;
+    plane0.GetDimensions(0, tint_tmp_4.x, tint_tmp_4.y, tint_tmp_4.z);
+    plane0.GetDimensions(tint_clamp(0, 0, int((uint(tint_tmp_4.z) - 1u))), tint_tmp_3.x, tint_tmp_3.y, tint_tmp_3.z);
+    int3 tint_tmp_5;
+    plane0.GetDimensions(0, tint_tmp_5.x, tint_tmp_5.y, tint_tmp_5.z);
+    int3 tint_tmp_6;
+    int3 tint_tmp_7;
+    plane1.GetDimensions(0, tint_tmp_7.x, tint_tmp_7.y, tint_tmp_7.z);
+    plane1.GetDimensions(tint_clamp(0, 0, int((uint(tint_tmp_7.z) - 1u))), tint_tmp_6.x, tint_tmp_6.y, tint_tmp_6.z);
+    int3 tint_tmp_8;
+    plane1.GetDimensions(0, tint_tmp_8.x, tint_tmp_8.y, tint_tmp_8.z);
+    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(tint_clamp_1(coord, (0).xx, int2((uint2(tint_tmp_3.xy) - (1u).xx))), tint_clamp(0, 0, int((uint(tint_tmp_5.z) - 1u))))).r, plane1.Load(int3(tint_clamp_1(coord1, (0).xx, int2((uint2(tint_tmp_6.xy) - (1u).xx))), tint_clamp(0, 0, int((uint(tint_tmp_8.z) - 1u))))).rg, 1.0f));
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -68,8 +90,8 @@ GammaTransferParams ext_tex_params_load_4(uint offset) {
   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_4 = {asfloat(ext_tex_params[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(ext_tex_params[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(ext_tex_params[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(ext_tex_params[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(ext_tex_params[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(ext_tex_params[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(ext_tex_params[scalar_offset_9 / 4][scalar_offset_9 % 4]), ext_tex_params[scalar_offset_10 / 4][scalar_offset_10 % 4]};
-  return tint_symbol_4;
+  const GammaTransferParams tint_symbol = {asfloat(ext_tex_params[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(ext_tex_params[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(ext_tex_params[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(ext_tex_params[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(ext_tex_params[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(ext_tex_params[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(ext_tex_params[scalar_offset_9 / 4][scalar_offset_9 % 4]), ext_tex_params[scalar_offset_10 / 4][scalar_offset_10 % 4]};
+  return tint_symbol;
 }
 
 float3x3 ext_tex_params_load_6(uint offset) {
@@ -92,27 +114,19 @@ float3x2 ext_tex_params_load_8(uint offset) {
 ExternalTextureParams ext_tex_params_load(uint offset) {
   const uint scalar_offset_17 = ((offset + 0u)) / 4;
   const uint scalar_offset_18 = ((offset + 4u)) / 4;
-  const ExternalTextureParams tint_symbol_5 = {ext_tex_params[scalar_offset_17 / 4][scalar_offset_17 % 4], ext_tex_params[scalar_offset_18 / 4][scalar_offset_18 % 4], ext_tex_params_load_2((offset + 16u)), ext_tex_params_load_4((offset + 64u)), ext_tex_params_load_4((offset + 96u)), ext_tex_params_load_6((offset + 128u)), ext_tex_params_load_8((offset + 176u))};
-  return tint_symbol_5;
+  const ExternalTextureParams tint_symbol_1 = {ext_tex_params[scalar_offset_17 / 4][scalar_offset_17 % 4], ext_tex_params[scalar_offset_18 / 4][scalar_offset_18 % 4], ext_tex_params_load_2((offset + 16u)), ext_tex_params_load_4((offset + 64u)), ext_tex_params_load_4((offset + 96u)), ext_tex_params_load_6((offset + 128u)), ext_tex_params_load_8((offset + 176u))};
+  return tint_symbol_1;
 }
 
 [numthreads(1, 1, 1)]
 void main() {
-  int2 tint_tmp;
-  t.GetDimensions(tint_tmp.x, tint_tmp.y);
-  const int2 tint_symbol = tint_clamp((10).xx, (0).xx, int2((uint2(tint_tmp) - (1u).xx)));
-  float4 red = textureLoadExternal(t, ext_tex_plane_1, tint_symbol, ext_tex_params_load(0u));
-  int2 tint_tmp_1;
-  outImage.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
-  const int2 tint_symbol_1 = tint_clamp((0).xx, (0).xx, int2((uint2(tint_tmp_1) - (1u).xx)));
-  outImage[tint_symbol_1] = red;
-  int2 tint_tmp_2;
-  t.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
-  const int2 tint_symbol_2 = tint_clamp(int2(70, 118), (0).xx, int2((uint2(tint_tmp_2) - (1u).xx)));
-  float4 green = textureLoadExternal(t, ext_tex_plane_1, tint_symbol_2, ext_tex_params_load(0u));
-  int2 tint_tmp_3;
-  outImage.GetDimensions(tint_tmp_3.x, tint_tmp_3.y);
-  const int2 tint_symbol_3 = tint_clamp(int2(1, 0), (0).xx, int2((uint2(tint_tmp_3) - (1u).xx)));
-  outImage[tint_symbol_3] = green;
+  float4 red = textureLoadExternal(t, ext_tex_plane_1, (10).xx, ext_tex_params_load(0u));
+  int2 tint_tmp_9;
+  outImage.GetDimensions(tint_tmp_9.x, tint_tmp_9.y);
+  outImage[tint_clamp_1((0).xx, (0).xx, int2((uint2(tint_tmp_9) - (1u).xx)))] = red;
+  float4 green = textureLoadExternal(t, ext_tex_plane_1, int2(70, 118), ext_tex_params_load(0u));
+  int2 tint_tmp_10;
+  outImage.GetDimensions(tint_tmp_10.x, tint_tmp_10.y);
+  outImage[tint_clamp_1(int2(1, 0), (0).xx, int2((uint2(tint_tmp_10) - (1u).xx)))] = green;
   return;
 }

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

@@ -1,4 +1,8 @@
-int2 tint_clamp(int2 e, int2 low, int2 high) {
+int tint_clamp(int e, int low, int high) {
+  return min(max(e, low), high);
+}
+
+int2 tint_clamp_1(int2 e, int2 low, int2 high) {
   return min(max(e, low), high);
 }
 
@@ -40,9 +44,27 @@ float4 textureLoadExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, i
   const int2 coord1 = (coord >> (1u).xx);
   float3 color = float3(0.0f, 0.0f, 0.0f);
   if ((params.numPlanes == 1u)) {
-    color = plane0.Load(int3(coord, 0)).rgb;
+    int3 tint_tmp;
+    int3 tint_tmp_1;
+    plane0.GetDimensions(0, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
+    plane0.GetDimensions(tint_clamp(0, 0, int((uint(tint_tmp_1.z) - 1u))), tint_tmp.x, tint_tmp.y, tint_tmp.z);
+    int3 tint_tmp_2;
+    plane0.GetDimensions(0, tint_tmp_2.x, tint_tmp_2.y, tint_tmp_2.z);
+    color = plane0.Load(int3(tint_clamp_1(coord, (0).xx, int2((uint2(tint_tmp.xy) - (1u).xx))), tint_clamp(0, 0, int((uint(tint_tmp_2.z) - 1u))))).rgb;
   } else {
-    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(coord, 0)).r, plane1.Load(int3(coord1, 0)).rg, 1.0f));
+    int3 tint_tmp_3;
+    int3 tint_tmp_4;
+    plane0.GetDimensions(0, tint_tmp_4.x, tint_tmp_4.y, tint_tmp_4.z);
+    plane0.GetDimensions(tint_clamp(0, 0, int((uint(tint_tmp_4.z) - 1u))), tint_tmp_3.x, tint_tmp_3.y, tint_tmp_3.z);
+    int3 tint_tmp_5;
+    plane0.GetDimensions(0, tint_tmp_5.x, tint_tmp_5.y, tint_tmp_5.z);
+    int3 tint_tmp_6;
+    int3 tint_tmp_7;
+    plane1.GetDimensions(0, tint_tmp_7.x, tint_tmp_7.y, tint_tmp_7.z);
+    plane1.GetDimensions(tint_clamp(0, 0, int((uint(tint_tmp_7.z) - 1u))), tint_tmp_6.x, tint_tmp_6.y, tint_tmp_6.z);
+    int3 tint_tmp_8;
+    plane1.GetDimensions(0, tint_tmp_8.x, tint_tmp_8.y, tint_tmp_8.z);
+    color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(tint_clamp_1(coord, (0).xx, int2((uint2(tint_tmp_3.xy) - (1u).xx))), tint_clamp(0, 0, int((uint(tint_tmp_5.z) - 1u))))).r, plane1.Load(int3(tint_clamp_1(coord1, (0).xx, int2((uint2(tint_tmp_6.xy) - (1u).xx))), tint_clamp(0, 0, int((uint(tint_tmp_8.z) - 1u))))).rg, 1.0f));
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -68,8 +90,8 @@ GammaTransferParams ext_tex_params_load_4(uint offset) {
   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_4 = {asfloat(ext_tex_params[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(ext_tex_params[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(ext_tex_params[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(ext_tex_params[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(ext_tex_params[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(ext_tex_params[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(ext_tex_params[scalar_offset_9 / 4][scalar_offset_9 % 4]), ext_tex_params[scalar_offset_10 / 4][scalar_offset_10 % 4]};
-  return tint_symbol_4;
+  const GammaTransferParams tint_symbol = {asfloat(ext_tex_params[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(ext_tex_params[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(ext_tex_params[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(ext_tex_params[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(ext_tex_params[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(ext_tex_params[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(ext_tex_params[scalar_offset_9 / 4][scalar_offset_9 % 4]), ext_tex_params[scalar_offset_10 / 4][scalar_offset_10 % 4]};
+  return tint_symbol;
 }
 
 float3x3 ext_tex_params_load_6(uint offset) {
@@ -92,27 +114,19 @@ float3x2 ext_tex_params_load_8(uint offset) {
 ExternalTextureParams ext_tex_params_load(uint offset) {
   const uint scalar_offset_17 = ((offset + 0u)) / 4;
   const uint scalar_offset_18 = ((offset + 4u)) / 4;
-  const ExternalTextureParams tint_symbol_5 = {ext_tex_params[scalar_offset_17 / 4][scalar_offset_17 % 4], ext_tex_params[scalar_offset_18 / 4][scalar_offset_18 % 4], ext_tex_params_load_2((offset + 16u)), ext_tex_params_load_4((offset + 64u)), ext_tex_params_load_4((offset + 96u)), ext_tex_params_load_6((offset + 128u)), ext_tex_params_load_8((offset + 176u))};
-  return tint_symbol_5;
+  const ExternalTextureParams tint_symbol_1 = {ext_tex_params[scalar_offset_17 / 4][scalar_offset_17 % 4], ext_tex_params[scalar_offset_18 / 4][scalar_offset_18 % 4], ext_tex_params_load_2((offset + 16u)), ext_tex_params_load_4((offset + 64u)), ext_tex_params_load_4((offset + 96u)), ext_tex_params_load_6((offset + 128u)), ext_tex_params_load_8((offset + 176u))};
+  return tint_symbol_1;
 }
 
 [numthreads(1, 1, 1)]
 void main() {
-  int2 tint_tmp;
-  t.GetDimensions(tint_tmp.x, tint_tmp.y);
-  const int2 tint_symbol = tint_clamp((10).xx, (0).xx, int2((uint2(tint_tmp) - (1u).xx)));
-  float4 red = textureLoadExternal(t, ext_tex_plane_1, tint_symbol, ext_tex_params_load(0u));
-  int2 tint_tmp_1;
-  outImage.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
-  const int2 tint_symbol_1 = tint_clamp((0).xx, (0).xx, int2((uint2(tint_tmp_1) - (1u).xx)));
-  outImage[tint_symbol_1] = red;
-  int2 tint_tmp_2;
-  t.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
-  const int2 tint_symbol_2 = tint_clamp(int2(70, 118), (0).xx, int2((uint2(tint_tmp_2) - (1u).xx)));
-  float4 green = textureLoadExternal(t, ext_tex_plane_1, tint_symbol_2, ext_tex_params_load(0u));
-  int2 tint_tmp_3;
-  outImage.GetDimensions(tint_tmp_3.x, tint_tmp_3.y);
-  const int2 tint_symbol_3 = tint_clamp(int2(1, 0), (0).xx, int2((uint2(tint_tmp_3) - (1u).xx)));
-  outImage[tint_symbol_3] = green;
+  float4 red = textureLoadExternal(t, ext_tex_plane_1, (10).xx, ext_tex_params_load(0u));
+  int2 tint_tmp_9;
+  outImage.GetDimensions(tint_tmp_9.x, tint_tmp_9.y);
+  outImage[tint_clamp_1((0).xx, (0).xx, int2((uint2(tint_tmp_9) - (1u).xx)))] = red;
+  float4 green = textureLoadExternal(t, ext_tex_plane_1, int2(70, 118), ext_tex_params_load(0u));
+  int2 tint_tmp_10;
+  outImage.GetDimensions(tint_tmp_10.x, tint_tmp_10.y);
+  outImage[tint_clamp_1(int2(1, 0), (0).xx, int2((uint2(tint_tmp_10) - (1u).xx)))] = green;
   return;
 }

test/tint/bug/tint/1739.wgsl.expected.glsl

@@ -1,3 +1,5 @@
+SKIP: FAILED
+
 #version 310 es
 
 struct GammaTransferParams {
@@ -57,9 +59,9 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
   ivec2 coord1 = (coord >> uvec2(1u));
   vec3 color = vec3(0.0f, 0.0f, 0.0f);
   if ((params.numPlanes == 1u)) {
-    color = texelFetch(plane0_1, coord, 0).rgb;
+    color = texelFetch(plane0_1, clamp(coord, ivec2(0), ivec2((uvec2(uvec2(textureSize(plane0_1, clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))))) - uvec2(1u)))), clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))).rgb;
   } else {
-    color = (vec4(texelFetch(plane0_1, coord, 0).r, texelFetch(plane1_1, coord1, 0).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (vec4(texelFetch(plane0_1, clamp(coord, ivec2(0), ivec2((uvec2(uvec2(textureSize(plane0_1, clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))))) - uvec2(1u)))), clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))).r, texelFetch(plane1_1, clamp(coord1, ivec2(0), ivec2((uvec2(uvec2(textureSize(plane1_1, clamp(0, 0, int((uint(uint(textureQueryLevels(plane1_1))) - 1u)))))) - uvec2(1u)))), clamp(0, 0, int((uint(uint(textureQueryLevels(plane1_1))) - 1u)))).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -76,9 +78,9 @@ ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 va
 }
 
 void tint_symbol() {
-  vec4 red = textureLoadExternal(t_2, ext_tex_plane_1_1, clamp(ivec2(10), ivec2(0), ivec2((uvec2(uvec2(textureSize(t_2, 0))) - uvec2(1u)))), conv_ExternalTextureParams(ext_tex_params.inner));
+  vec4 red = textureLoadExternal(t_2, ext_tex_plane_1_1, ivec2(10), conv_ExternalTextureParams(ext_tex_params.inner));
   imageStore(outImage, clamp(ivec2(0), ivec2(0), ivec2((uvec2(uvec2(imageSize(outImage))) - uvec2(1u)))), red);
-  vec4 green = textureLoadExternal(t_2, ext_tex_plane_1_1, clamp(ivec2(70, 118), ivec2(0), ivec2((uvec2(uvec2(textureSize(t_2, 0))) - uvec2(1u)))), conv_ExternalTextureParams(ext_tex_params.inner));
+  vec4 green = textureLoadExternal(t_2, ext_tex_plane_1_1, ivec2(70, 118), conv_ExternalTextureParams(ext_tex_params.inner));
   imageStore(outImage, clamp(ivec2(1, 0), ivec2(0), ivec2((uvec2(uvec2(imageSize(outImage))) - uvec2(1u)))), green);
   return;
 }
@@ -88,3 +90,10 @@ void main() {
   tint_symbol();
   return;
 }
+Error parsing GLSL shader:
+ERROR: 0:60: 'textureQueryLevels' : no matching overloaded function found 
+ERROR: 0:60: '' : compilation terminated 
+ERROR: 2 compilation errors.  No code generated.
+
+
+

test/tint/bug/tint/1739.wgsl.expected.msl

@@ -72,7 +72,11 @@ ExternalTextureParams tint_unpack_vec3_in_composite_1(ExternalTextureParams_tint
   return result;
 }
 
-int2 tint_clamp(int2 e, int2 low, int2 high) {
+int tint_clamp(int e, int low, int high) {
+  return min(max(e, low), high);
+}
+
+int2 tint_clamp_1(int2 e, int2 low, int2 high) {
   return min(max(e, low), high);
 }
 
@@ -87,9 +91,9 @@ float4 textureLoadExternal(texture2d<float, access::sample> plane0, texture2d<fl
   int2 const coord1 = (coord >> uint2(1u));
   float3 color = 0.0f;
   if ((params.numPlanes == 1u)) {
-    color = plane0.read(uint2(coord), 0).rgb;
+    color = plane0.read(uint2(tint_clamp_1(coord, int2(0), int2((uint2(uint2(plane0.get_width(tint_clamp(0, 0, int((uint(plane0.get_num_mip_levels()) - 1u)))), plane0.get_height(tint_clamp(0, 0, int((uint(plane0.get_num_mip_levels()) - 1u)))))) - uint2(1u))))), tint_clamp(0, 0, int((uint(plane0.get_num_mip_levels()) - 1u)))).rgb;
   } else {
-    color = (float4(plane0.read(uint2(coord), 0)[0], plane1.read(uint2(coord1), 0).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+    color = (float4(plane0.read(uint2(tint_clamp_1(coord, int2(0), int2((uint2(uint2(plane0.get_width(tint_clamp(0, 0, int((uint(plane0.get_num_mip_levels()) - 1u)))), plane0.get_height(tint_clamp(0, 0, int((uint(plane0.get_num_mip_levels()) - 1u)))))) - uint2(1u))))), tint_clamp(0, 0, int((uint(plane0.get_num_mip_levels()) - 1u))))[0], plane1.read(uint2(tint_clamp_1(coord1, int2(0), int2((uint2(uint2(plane1.get_width(tint_clamp(0, 0, int((uint(plane1.get_num_mip_levels()) - 1u)))), plane1.get_height(tint_clamp(0, 0, int((uint(plane1.get_num_mip_levels()) - 1u)))))) - uint2(1u))))), tint_clamp(0, 0, int((uint(plane1.get_num_mip_levels()) - 1u)))).rg, 1.0f) * params.yuvToRgbConversionMatrix);
   }
   if ((params.doYuvToRgbConversionOnly == 0u)) {
     color = gammaCorrection(color, params.gammaDecodeParams);
@@ -99,15 +103,11 @@ float4 textureLoadExternal(texture2d<float, access::sample> plane0, texture2d<fl
   return float4(color, 1.0f);
 }
 
-kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]], texture2d<float, access::sample> tint_symbol_6 [[texture(1)]], const constant ExternalTextureParams_tint_packed_vec3* tint_symbol_7 [[buffer(0)]], texture2d<float, access::write> tint_symbol_8 [[texture(2)]]) {
-  int2 const tint_symbol_1 = tint_clamp(int2(10), int2(0), int2((uint2(uint2(tint_symbol_5.get_width(), tint_symbol_5.get_height())) - uint2(1u))));
-  float4 red = textureLoadExternal(tint_symbol_5, tint_symbol_6, tint_symbol_1, tint_unpack_vec3_in_composite_1(*(tint_symbol_7)));
-  int2 const tint_symbol_2 = tint_clamp(int2(0), int2(0), int2((uint2(uint2(tint_symbol_8.get_width(), tint_symbol_8.get_height())) - uint2(1u))));
-  tint_symbol_8.write(red, uint2(tint_symbol_2));
-  int2 const tint_symbol_3 = tint_clamp(int2(70, 118), int2(0), int2((uint2(uint2(tint_symbol_5.get_width(), tint_symbol_5.get_height())) - uint2(1u))));
-  float4 green = textureLoadExternal(tint_symbol_5, tint_symbol_6, tint_symbol_3, tint_unpack_vec3_in_composite_1(*(tint_symbol_7)));
-  int2 const tint_symbol_4 = tint_clamp(int2(1, 0), int2(0), int2((uint2(uint2(tint_symbol_8.get_width(), tint_symbol_8.get_height())) - uint2(1u))));
-  tint_symbol_8.write(green, uint2(tint_symbol_4));
+kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_1 [[texture(0)]], texture2d<float, access::sample> tint_symbol_2 [[texture(1)]], const constant ExternalTextureParams_tint_packed_vec3* tint_symbol_3 [[buffer(0)]], texture2d<float, access::write> tint_symbol_4 [[texture(2)]]) {
+  float4 red = textureLoadExternal(tint_symbol_1, tint_symbol_2, int2(10), tint_unpack_vec3_in_composite_1(*(tint_symbol_3)));
+  tint_symbol_4.write(red, uint2(tint_clamp_1(int2(0), int2(0), int2((uint2(uint2(tint_symbol_4.get_width(), tint_symbol_4.get_height())) - uint2(1u))))));
+  float4 green = textureLoadExternal(tint_symbol_1, tint_symbol_2, int2(70, 118), tint_unpack_vec3_in_composite_1(*(tint_symbol_3)));
+  tint_symbol_4.write(green, uint2(tint_clamp_1(int2(1, 0), int2(0), int2((uint2(uint2(tint_symbol_4.get_width(), tint_symbol_4.get_height())) - uint2(1u))))));
   return;
 }
 

test/tint/bug/tint/1739.wgsl.expected.spvasm

@@ -1,11 +1,11 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 197
+; Bound: 237
 ; Schema: 0
                OpCapability Shader
                OpCapability ImageQuery
-         %29 = OpExtInstImport "GLSL.std.450"
+         %28 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
@@ -38,6 +38,10 @@
                OpName %e "e"
                OpName %low "low"
                OpName %high "high"
+               OpName %tint_clamp_1 "tint_clamp_1"
+               OpName %e_0 "e"
+               OpName %low_0 "low"
+               OpName %high_0 "high"
                OpName %gammaCorrection "gammaCorrection"
                OpName %v "v"
                OpName %params "params"
@@ -126,198 +130,233 @@
 %_ptr_UniformConstant_19 = OpTypePointer UniformConstant %19
    %outImage = OpVariable %_ptr_UniformConstant_19 UniformConstant
         %int = OpTypeInt 32 1
+         %20 = OpTypeFunction %int %int %int %int
       %v2int = OpTypeVector %int 2
-         %20 = OpTypeFunction %v2int %v2int %v2int %v2int
-         %31 = OpTypeFunction %v3float %v3float %GammaTransferParams
+         %30 = OpTypeFunction %v2int %v2int %v2int %v2int
+         %39 = OpTypeFunction %v3float %v3float %GammaTransferParams
        %bool = OpTypeBool
      %v3bool = OpTypeVector %bool 3
 %_ptr_Function_v3float = OpTypePointer Function %v3float
-         %50 = OpConstantNull %v3float
+         %58 = OpConstantNull %v3float
 %mat3v2float = OpTypeMatrix %v2float 3
 %ExternalTextureParams = OpTypeStruct %uint %uint %mat3v4float %GammaTransferParams %GammaTransferParams %mat3v3float %mat3v2float
-         %70 = OpTypeFunction %v4float %3 %3 %v2int %ExternalTextureParams
+         %78 = OpTypeFunction %v4float %3 %3 %v2int %ExternalTextureParams
      %v2uint = OpTypeVector %uint 2
      %uint_1 = OpConstant %uint 1
-         %81 = OpConstantComposite %v2uint %uint_1 %uint_1
-         %90 = OpConstantNull %int
+         %89 = OpConstantComposite %v2uint %uint_1 %uint_1
+         %99 = OpConstantNull %v2int
+        %104 = OpConstantNull %int
     %float_1 = OpConstant %float 1
-        %103 = OpConstantNull %uint
-        %121 = OpTypeFunction %ExternalTextureParams %ExternalTextureParams_std140
+        %157 = OpConstantNull %uint
+        %175 = OpTypeFunction %ExternalTextureParams %ExternalTextureParams_std140
        %void = OpTypeVoid
-        %136 = OpTypeFunction %void
+        %190 = OpTypeFunction %void
      %int_10 = OpConstant %int 10
-        %142 = OpConstantComposite %v2int %int_10 %int_10
-        %143 = OpConstantNull %v2int
-      %int_0 = OpConstant %int 0
+        %198 = OpConstantComposite %v2int %int_10 %int_10
      %uint_0 = OpConstant %uint 0
 %_ptr_Uniform_ExternalTextureParams_std140 = OpTypePointer Uniform %ExternalTextureParams_std140
 %_ptr_Function_v4float = OpTypePointer Function %v4float
-        %160 = OpConstantNull %v4float
+        %206 = OpConstantNull %v4float
      %int_70 = OpConstant %int 70
     %int_118 = OpConstant %int 118
-        %173 = OpConstantComposite %v2int %int_70 %int_118
+        %221 = OpConstantComposite %v2int %int_70 %int_118
       %int_1 = OpConstant %int 1
-        %188 = OpConstantComposite %v2int %int_1 %90
- %tint_clamp = OpFunction %v2int None %20
-          %e = OpFunctionParameter %v2int
-        %low = OpFunctionParameter %v2int
-       %high = OpFunctionParameter %v2int
-         %27 = OpLabel
-         %30 = OpExtInst %v2int %29 SMax %e %low
-         %28 = OpExtInst %v2int %29 SMin %30 %high
-               OpReturnValue %28
+        %230 = OpConstantComposite %v2int %int_1 %104
+ %tint_clamp = OpFunction %int None %20
+          %e = OpFunctionParameter %int
+        %low = OpFunctionParameter %int
+       %high = OpFunctionParameter %int
+         %26 = OpLabel
+         %29 = OpExtInst %int %28 SMax %e %low
+         %27 = OpExtInst %int %28 SMin %29 %high
+               OpReturnValue %27
                OpFunctionEnd
-%gammaCorrection = OpFunction %v3float None %31
+%tint_clamp_1 = OpFunction %v2int None %30
+        %e_0 = OpFunctionParameter %v2int
+      %low_0 = OpFunctionParameter %v2int
+     %high_0 = OpFunctionParameter %v2int
+         %36 = OpLabel
+         %38 = OpExtInst %v2int %28 SMax %e_0 %low_0
+         %37 = OpExtInst %v2int %28 SMin %38 %high_0
+               OpReturnValue %37
+               OpFunctionEnd
+%gammaCorrection = OpFunction %v3float None %39
           %v = OpFunctionParameter %v3float
      %params = OpFunctionParameter %GammaTransferParams
-         %35 = OpLabel
-         %48 = OpVariable %_ptr_Function_v3float Function %50
-         %60 = OpVariable %_ptr_Function_v3float Function %50
-         %66 = OpVariable %_ptr_Function_v3float Function %50
-         %36 = OpExtInst %v3float %29 FAbs %v
-         %37 = OpCompositeExtract %float %params 4
-         %38 = OpCompositeConstruct %v3float %37 %37 %37
-         %39 = OpFOrdLessThan %v3bool %36 %38
-         %42 = OpExtInst %v3float %29 FSign %v
-         %43 = OpCompositeExtract %float %params 3
-         %44 = OpExtInst %v3float %29 FAbs %v
-         %45 = OpVectorTimesScalar %v3float %44 %43
-         %46 = OpCompositeExtract %float %params 6
-         %51 = OpCompositeConstruct %v3float %46 %46 %46
-         %47 = OpFAdd %v3float %45 %51
-         %52 = OpFMul %v3float %42 %47
-         %53 = OpExtInst %v3float %29 FSign %v
-         %55 = OpCompositeExtract %float %params 1
-         %56 = OpExtInst %v3float %29 FAbs %v
-         %57 = OpVectorTimesScalar %v3float %56 %55
-         %58 = OpCompositeExtract %float %params 2
-         %61 = OpCompositeConstruct %v3float %58 %58 %58
-         %59 = OpFAdd %v3float %57 %61
-         %62 = OpCompositeExtract %float %params 0
-         %63 = OpCompositeConstruct %v3float %62 %62 %62
-         %54 = OpExtInst %v3float %29 Pow %59 %63
-         %64 = OpCompositeExtract %float %params 5
-         %67 = OpCompositeConstruct %v3float %64 %64 %64
-         %65 = OpFAdd %v3float %54 %67
-         %68 = OpFMul %v3float %53 %65
-         %69 = OpSelect %v3float %39 %52 %68
-               OpReturnValue %69
+         %43 = OpLabel
+         %56 = OpVariable %_ptr_Function_v3float Function %58
+         %68 = OpVariable %_ptr_Function_v3float Function %58
+         %74 = OpVariable %_ptr_Function_v3float Function %58
+         %44 = OpExtInst %v3float %28 FAbs %v
+         %45 = OpCompositeExtract %float %params 4
+         %46 = OpCompositeConstruct %v3float %45 %45 %45
+         %47 = OpFOrdLessThan %v3bool %44 %46
+         %50 = OpExtInst %v3float %28 FSign %v
+         %51 = OpCompositeExtract %float %params 3
+         %52 = OpExtInst %v3float %28 FAbs %v
+         %53 = OpVectorTimesScalar %v3float %52 %51
+         %54 = OpCompositeExtract %float %params 6
+         %59 = OpCompositeConstruct %v3float %54 %54 %54
+         %55 = OpFAdd %v3float %53 %59
+         %60 = OpFMul %v3float %50 %55
+         %61 = OpExtInst %v3float %28 FSign %v
+         %63 = OpCompositeExtract %float %params 1
+         %64 = OpExtInst %v3float %28 FAbs %v
+         %65 = OpVectorTimesScalar %v3float %64 %63
+         %66 = OpCompositeExtract %float %params 2
+         %69 = OpCompositeConstruct %v3float %66 %66 %66
+         %67 = OpFAdd %v3float %65 %69
+         %70 = OpCompositeExtract %float %params 0
+         %71 = OpCompositeConstruct %v3float %70 %70 %70
+         %62 = OpExtInst %v3float %28 Pow %67 %71
+         %72 = OpCompositeExtract %float %params 5
+         %75 = OpCompositeConstruct %v3float %72 %72 %72
+         %73 = OpFAdd %v3float %62 %75
+         %76 = OpFMul %v3float %61 %73
+         %77 = OpSelect %v3float %47 %60 %76
+               OpReturnValue %77
                OpFunctionEnd
-%textureLoadExternal = OpFunction %v4float None %70
+%textureLoadExternal = OpFunction %v4float None %78
      %plane0 = OpFunctionParameter %3
      %plane1 = OpFunctionParameter %3
       %coord = OpFunctionParameter %v2int
    %params_0 = OpFunctionParameter %ExternalTextureParams
-         %78 = OpLabel
-      %color = OpVariable %_ptr_Function_v3float Function %50
-         %82 = OpShiftRightArithmetic %v2int %coord %81
-         %84 = OpCompositeExtract %uint %params_0 0
-         %85 = OpIEqual %bool %84 %uint_1
-               OpSelectionMerge %86 None
-               OpBranchConditional %85 %87 %88
-         %87 = OpLabel
-         %89 = OpImageFetch %v4float %plane0 %coord Lod %90
-         %91 = OpVectorShuffle %v3float %89 %89 0 1 2
-               OpStore %color %91
-               OpBranch %86
-         %88 = OpLabel
-         %92 = OpImageFetch %v4float %plane0 %coord Lod %90
-         %93 = OpCompositeExtract %float %92 0
-         %94 = OpImageFetch %v4float %plane1 %82 Lod %90
-         %95 = OpVectorShuffle %v2float %94 %94 0 1
-         %96 = OpCompositeExtract %float %95 0
-         %97 = OpCompositeExtract %float %95 1
-         %99 = OpCompositeConstruct %v4float %93 %96 %97 %float_1
-        %100 = OpCompositeExtract %mat3v4float %params_0 2
-        %101 = OpVectorTimesMatrix %v3float %99 %100
-               OpStore %color %101
-               OpBranch %86
          %86 = OpLabel
-        %102 = OpCompositeExtract %uint %params_0 1
-        %104 = OpIEqual %bool %102 %103
-               OpSelectionMerge %105 None
-               OpBranchConditional %104 %106 %105
-        %106 = OpLabel
-        %108 = OpLoad %v3float %color
-        %109 = OpCompositeExtract %GammaTransferParams %params_0 3
-        %107 = OpFunctionCall %v3float %gammaCorrection %108 %109
-               OpStore %color %107
-        %110 = OpCompositeExtract %mat3v3float %params_0 5
-        %111 = OpLoad %v3float %color
-        %112 = OpMatrixTimesVector %v3float %110 %111
-               OpStore %color %112
-        %114 = OpLoad %v3float %color
-        %115 = OpCompositeExtract %GammaTransferParams %params_0 4
-        %113 = OpFunctionCall %v3float %gammaCorrection %114 %115
-               OpStore %color %113
-               OpBranch %105
-        %105 = OpLabel
-        %116 = OpLoad %v3float %color
-        %117 = OpCompositeExtract %float %116 0
-        %118 = OpCompositeExtract %float %116 1
-        %119 = OpCompositeExtract %float %116 2
-        %120 = OpCompositeConstruct %v4float %117 %118 %119 %float_1
-               OpReturnValue %120
+      %color = OpVariable %_ptr_Function_v3float Function %58
+         %90 = OpShiftRightArithmetic %v2int %coord %89
+         %92 = OpCompositeExtract %uint %params_0 0
+         %93 = OpIEqual %bool %92 %uint_1
+               OpSelectionMerge %94 None
+               OpBranchConditional %93 %95 %96
+         %95 = OpLabel
+        %107 = OpImageQueryLevels %uint %plane0
+        %108 = OpISub %uint %107 %uint_1
+        %105 = OpBitcast %int %108
+        %103 = OpFunctionCall %int %tint_clamp %104 %104 %105
+        %102 = OpImageQuerySizeLod %v2uint %plane0 %103
+        %109 = OpISub %v2uint %102 %89
+        %100 = OpBitcast %v2int %109
+         %98 = OpFunctionCall %v2int %tint_clamp_1 %coord %99 %100
+        %113 = OpImageQueryLevels %uint %plane0
+        %114 = OpISub %uint %113 %uint_1
+        %111 = OpBitcast %int %114
+        %110 = OpFunctionCall %int %tint_clamp %104 %104 %111
+         %97 = OpImageFetch %v4float %plane0 %98 Lod %110
+        %115 = OpVectorShuffle %v3float %97 %97 0 1 2
+               OpStore %color %115
+               OpBranch %94
+         %96 = OpLabel
+        %124 = OpImageQueryLevels %uint %plane0
+        %125 = OpISub %uint %124 %uint_1
+        %122 = OpBitcast %int %125
+        %121 = OpFunctionCall %int %tint_clamp %104 %104 %122
+        %120 = OpImageQuerySizeLod %v2uint %plane0 %121
+        %126 = OpISub %v2uint %120 %89
+        %118 = OpBitcast %v2int %126
+        %117 = OpFunctionCall %v2int %tint_clamp_1 %coord %99 %118
+        %130 = OpImageQueryLevels %uint %plane0
+        %131 = OpISub %uint %130 %uint_1
+        %128 = OpBitcast %int %131
+        %127 = OpFunctionCall %int %tint_clamp %104 %104 %128
+        %116 = OpImageFetch %v4float %plane0 %117 Lod %127
+        %132 = OpCompositeExtract %float %116 0
+        %141 = OpImageQueryLevels %uint %plane1
+        %142 = OpISub %uint %141 %uint_1
+        %139 = OpBitcast %int %142
+        %138 = OpFunctionCall %int %tint_clamp %104 %104 %139
+        %137 = OpImageQuerySizeLod %v2uint %plane1 %138
+        %143 = OpISub %v2uint %137 %89
+        %135 = OpBitcast %v2int %143
+        %134 = OpFunctionCall %v2int %tint_clamp_1 %90 %99 %135
+        %147 = OpImageQueryLevels %uint %plane1
+        %148 = OpISub %uint %147 %uint_1
+        %145 = OpBitcast %int %148
+        %144 = OpFunctionCall %int %tint_clamp %104 %104 %145
+        %133 = OpImageFetch %v4float %plane1 %134 Lod %144
+        %149 = OpVectorShuffle %v2float %133 %133 0 1
+        %150 = OpCompositeExtract %float %149 0
+        %151 = OpCompositeExtract %float %149 1
+        %153 = OpCompositeConstruct %v4float %132 %150 %151 %float_1
+        %154 = OpCompositeExtract %mat3v4float %params_0 2
+        %155 = OpVectorTimesMatrix %v3float %153 %154
+               OpStore %color %155
+               OpBranch %94
+         %94 = OpLabel
+        %156 = OpCompositeExtract %uint %params_0 1
+        %158 = OpIEqual %bool %156 %157
+               OpSelectionMerge %159 None
+               OpBranchConditional %158 %160 %159
+        %160 = OpLabel
+        %162 = OpLoad %v3float %color
+        %163 = OpCompositeExtract %GammaTransferParams %params_0 3
+        %161 = OpFunctionCall %v3float %gammaCorrection %162 %163
+               OpStore %color %161
+        %164 = OpCompositeExtract %mat3v3float %params_0 5
+        %165 = OpLoad %v3float %color
+        %166 = OpMatrixTimesVector %v3float %164 %165
+               OpStore %color %166
+        %168 = OpLoad %v3float %color
+        %169 = OpCompositeExtract %GammaTransferParams %params_0 4
+        %167 = OpFunctionCall %v3float %gammaCorrection %168 %169
+               OpStore %color %167
+               OpBranch %159
+        %159 = OpLabel
+        %170 = OpLoad %v3float %color
+        %171 = OpCompositeExtract %float %170 0
+        %172 = OpCompositeExtract %float %170 1
+        %173 = OpCompositeExtract %float %170 2
+        %174 = OpCompositeConstruct %v4float %171 %172 %173 %float_1
+               OpReturnValue %174
                OpFunctionEnd
-%conv_ExternalTextureParams = OpFunction %ExternalTextureParams None %121
+%conv_ExternalTextureParams = OpFunction %ExternalTextureParams None %175
         %val = OpFunctionParameter %ExternalTextureParams_std140
-        %124 = OpLabel
-        %125 = OpCompositeExtract %uint %val 0
-        %126 = OpCompositeExtract %uint %val 1
-        %127 = OpCompositeExtract %mat3v4float %val 2
-        %128 = OpCompositeExtract %GammaTransferParams %val 3
-        %129 = OpCompositeExtract %GammaTransferParams %val 4
-        %130 = OpCompositeExtract %mat3v3float %val 5
-        %131 = OpCompositeExtract %v2float %val 6
-        %132 = OpCompositeExtract %v2float %val 7
-        %133 = OpCompositeExtract %v2float %val 8
-        %134 = OpCompositeConstruct %mat3v2float %131 %132 %133
-        %135 = OpCompositeConstruct %ExternalTextureParams %125 %126 %127 %128 %129 %130 %134
-               OpReturnValue %135
+        %178 = OpLabel
+        %179 = OpCompositeExtract %uint %val 0
+        %180 = OpCompositeExtract %uint %val 1
+        %181 = OpCompositeExtract %mat3v4float %val 2
+        %182 = OpCompositeExtract %GammaTransferParams %val 3
+        %183 = OpCompositeExtract %GammaTransferParams %val 4
+        %184 = OpCompositeExtract %mat3v3float %val 5
+        %185 = OpCompositeExtract %v2float %val 6
+        %186 = OpCompositeExtract %v2float %val 7
+        %187 = OpCompositeExtract %v2float %val 8
+        %188 = OpCompositeConstruct %mat3v2float %185 %186 %187
+        %189 = OpCompositeConstruct %ExternalTextureParams %179 %180 %181 %182 %183 %184 %188
+               OpReturnValue %189
                OpFunctionEnd
-       %main = OpFunction %void None %136
-        %139 = OpLabel
-        %red = OpVariable %_ptr_Function_v4float Function %160
-      %green = OpVariable %_ptr_Function_v4float Function %160
-        %147 = OpLoad %3 %t
-        %146 = OpImageQuerySizeLod %v2uint %147 %int_0
-        %149 = OpISub %v2uint %146 %81
-        %144 = OpBitcast %v2int %149
-        %140 = OpFunctionCall %v2int %tint_clamp %142 %143 %144
-        %151 = OpLoad %3 %t
-        %152 = OpLoad %3 %ext_tex_plane_1
-        %156 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
-        %157 = OpLoad %ExternalTextureParams_std140 %156
-        %153 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %157
-        %150 = OpFunctionCall %v4float %textureLoadExternal %151 %152 %140 %153
-               OpStore %red %150
-        %165 = OpLoad %19 %outImage
-        %164 = OpImageQuerySize %v2uint %165
-        %166 = OpISub %v2uint %164 %81
-        %162 = OpBitcast %v2int %166
-        %161 = OpFunctionCall %v2int %tint_clamp %143 %143 %162
-        %168 = OpLoad %19 %outImage
-        %169 = OpLoad %v4float %red
-               OpImageWrite %168 %161 %169
-        %177 = OpLoad %3 %t
-        %176 = OpImageQuerySizeLod %v2uint %177 %int_0
-        %178 = OpISub %v2uint %176 %81
-        %174 = OpBitcast %v2int %178
-        %170 = OpFunctionCall %v2int %tint_clamp %173 %143 %174
-        %180 = OpLoad %3 %t
-        %181 = OpLoad %3 %ext_tex_plane_1
-        %183 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
-        %184 = OpLoad %ExternalTextureParams_std140 %183
-        %182 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %184
-        %179 = OpFunctionCall %v4float %textureLoadExternal %180 %181 %170 %182
-               OpStore %green %179
-        %192 = OpLoad %19 %outImage
-        %191 = OpImageQuerySize %v2uint %192
-        %193 = OpISub %v2uint %191 %81
-        %189 = OpBitcast %v2int %193
-        %186 = OpFunctionCall %v2int %tint_clamp %188 %143 %189
-        %195 = OpLoad %19 %outImage
-        %196 = OpLoad %v4float %green
-               OpImageWrite %195 %186 %196
+       %main = OpFunction %void None %190
+        %193 = OpLabel
+        %red = OpVariable %_ptr_Function_v4float Function %206
+      %green = OpVariable %_ptr_Function_v4float Function %206
+        %195 = OpLoad %3 %t
+        %196 = OpLoad %3 %ext_tex_plane_1
+        %202 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
+        %203 = OpLoad %ExternalTextureParams_std140 %202
+        %199 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %203
+        %194 = OpFunctionCall %v4float %textureLoadExternal %195 %196 %198 %199
+               OpStore %red %194
+        %208 = OpLoad %19 %outImage
+        %213 = OpLoad %19 %outImage
+        %212 = OpImageQuerySize %v2uint %213
+        %214 = OpISub %v2uint %212 %89
+        %210 = OpBitcast %v2int %214
+        %209 = OpFunctionCall %v2int %tint_clamp_1 %99 %99 %210
+        %215 = OpLoad %v4float %red
+               OpImageWrite %208 %209 %215
+        %217 = OpLoad %3 %t
+        %218 = OpLoad %3 %ext_tex_plane_1
+        %223 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
+        %224 = OpLoad %ExternalTextureParams_std140 %223
+        %222 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %224
+        %216 = OpFunctionCall %v4float %textureLoadExternal %217 %218 %221 %222
+               OpStore %green %216
+        %227 = OpLoad %19 %outImage
+        %234 = OpLoad %19 %outImage
+        %233 = OpImageQuerySize %v2uint %234
+        %235 = OpISub %v2uint %233 %89
+        %231 = OpBitcast %v2int %235
+        %228 = OpFunctionCall %v2int %tint_clamp_1 %230 %99 %231
+        %236 = OpLoad %v4float %green
+               OpImageWrite %227 %228 %236
                OpReturn
                OpFunctionEnd

test/tint/bug/tint/1739.wgsl.expected.wgsl

@@ -52,9 +52,9 @@ fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord
 
 @compute @workgroup_size(1)
 fn main() {
-  var red : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, clamp(vec2<i32>(10, 10), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(t)) - vec2(1)))), ext_tex_params);
-  textureStore(outImage, clamp(vec2<i32>(0, 0), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(outImage)) - vec2(1)))), red);
-  var green : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, clamp(vec2<i32>(70, 118), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(t)) - vec2(1)))), ext_tex_params);
-  textureStore(outImage, clamp(vec2<i32>(1, 0), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(outImage)) - vec2(1)))), green);
+  var red : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, vec2<i32>(10, 10), ext_tex_params);
+  textureStore(outImage, vec2<i32>(0, 0), red);
+  var green : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, vec2<i32>(70, 118), ext_tex_params);
+  textureStore(outImage, vec2<i32>(1, 0), green);
   return;
 }

test/tint/bug/tint/221.wgsl.expected.dxc.hlsl

@@ -12,8 +12,7 @@ void main() {
       break;
     }
     const uint p_save = i;
-    const uint tint_symbol = tint_mod(i, 2u);
-    if ((tint_symbol == 0u)) {
+    if ((tint_mod(i, 2u) == 0u)) {
       {
         b.Store((4u + (4u * p_save)), asuint((b.Load((4u + (4u * p_save))) * 2u)));
         i = (i + 1u);

test/tint/bug/tint/221.wgsl.expected.fxc.hlsl

@@ -12,8 +12,7 @@ void main() {
       break;
     }
     const uint p_save = i;
-    const uint tint_symbol = tint_mod(i, 2u);
-    if ((tint_symbol == 0u)) {
+    if ((tint_mod(i, 2u) == 0u)) {
       {
         b.Store((4u + (4u * p_save)), asuint((b.Load((4u + (4u * p_save))) * 2u)));
         i = (i + 1u);

test/tint/bug/tint/221.wgsl.expected.glsl

@@ -20,8 +20,7 @@ void tint_symbol() {
       break;
     }
     uint p_save = i;
-    uint tint_symbol_1 = tint_mod(i, 2u);
-    if ((tint_symbol_1 == 0u)) {
+    if ((tint_mod(i, 2u) == 0u)) {
       {
         b.inner.data[p_save] = (b.inner.data[p_save] * 2u);
         i = (i + 1u);

test/tint/bug/tint/221.wgsl.expected.msl

@@ -23,24 +23,23 @@ uint tint_mod(uint lhs, uint rhs) {
   return (lhs % select(rhs, 1u, (rhs == 0u)));
 }
 
-kernel void tint_symbol(device Buf* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device Buf* tint_symbol_1 [[buffer(0)]]) {
   uint i = 0u;
   while (true) {
-    if ((i >= (*(tint_symbol_2)).count)) {
+    if ((i >= (*(tint_symbol_1)).count)) {
       break;
     }
     uint const p_save = i;
-    uint const tint_symbol_1 = tint_mod(i, 2u);
-    if ((tint_symbol_1 == 0u)) {
+    if ((tint_mod(i, 2u) == 0u)) {
       {
-        (*(tint_symbol_2)).data[p_save] = ((*(tint_symbol_2)).data[p_save] * 2u);
+        (*(tint_symbol_1)).data[p_save] = ((*(tint_symbol_1)).data[p_save] * 2u);
         i = (i + 1u);
       }
       continue;
     }
-    (*(tint_symbol_2)).data[p_save] = 0u;
+    (*(tint_symbol_1)).data[p_save] = 0u;
     {
-      (*(tint_symbol_2)).data[p_save] = ((*(tint_symbol_2)).data[p_save] * 2u);
+      (*(tint_symbol_1)).data[p_save] = ((*(tint_symbol_1)).data[p_save] * 2u);
       i = (i + 1u);
     }
   }

test/tint/bug/tint/534.wgsl.expected.dxc.hlsl

@@ -13,7 +13,7 @@ uint ConvertToFp16FloatValue(float fp32) {
   return 1u;
 }
 
-struct tint_symbol_2 {
+struct tint_symbol_3 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -33,8 +33,8 @@ void main_inner(uint3 GlobalInvocationID) {
   uint4 dstColorBits = uint4(dstColor);
   {
     for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
-      const uint tint_symbol_3 = ConvertToFp16FloatValue(srcColor[i]);
-      set_uint4(srcColorBits, i, tint_symbol_3);
+      const uint tint_symbol_1 = ConvertToFp16FloatValue(srcColor[i]);
+      set_uint4(srcColorBits, i, tint_symbol_1);
       bool tint_tmp_1 = success;
       if (tint_tmp_1) {
         tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
@@ -51,7 +51,7 @@ void main_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
-  main_inner(tint_symbol_1.GlobalInvocationID);
+void main(tint_symbol_3 tint_symbol_2) {
+  main_inner(tint_symbol_2.GlobalInvocationID);
   return;
 }

test/tint/bug/tint/534.wgsl.expected.fxc.hlsl

@@ -13,7 +13,7 @@ uint ConvertToFp16FloatValue(float fp32) {
   return 1u;
 }
 
-struct tint_symbol_2 {
+struct tint_symbol_3 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -33,8 +33,8 @@ void main_inner(uint3 GlobalInvocationID) {
   uint4 dstColorBits = uint4(dstColor);
   {
     for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
-      const uint tint_symbol_3 = ConvertToFp16FloatValue(srcColor[i]);
-      set_uint4(srcColorBits, i, tint_symbol_3);
+      const uint tint_symbol_1 = ConvertToFp16FloatValue(srcColor[i]);
+      set_uint4(srcColorBits, i, tint_symbol_1);
       bool tint_tmp_1 = success;
       if (tint_tmp_1) {
         tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
@@ -51,7 +51,7 @@ void main_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
-  main_inner(tint_symbol_1.GlobalInvocationID);
+void main(tint_symbol_3 tint_symbol_2) {
+  main_inner(tint_symbol_2.GlobalInvocationID);
   return;
 }

test/tint/bug/tint/913.wgsl.expected.dxc.hlsl

@@ -9,7 +9,7 @@ bool aboutEqual(float value, float expect) {
   return (abs((value - expect)) < 0.001f);
 }
 
-struct tint_symbol_2 {
+struct tint_symbol_8 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -49,33 +49,33 @@ void main_inner(uint3 GlobalInvocationID) {
     const float4 srcColor = src.Load(int3(int2(srcTexCoord), 0));
     const float4 dstColor = tint_symbol.Load(int3(int2(dstTexCoord), 0));
     if ((uniforms[0].y == 2u)) {
-      bool tint_symbol_4 = success;
-      if (tint_symbol_4) {
-        tint_symbol_4 = aboutEqual(dstColor.r, srcColor.r);
+      bool tint_symbol_2 = success;
+      if (tint_symbol_2) {
+        tint_symbol_2 = aboutEqual(dstColor.r, srcColor.r);
       }
-      bool tint_symbol_3 = tint_symbol_4;
-      if (tint_symbol_3) {
-        tint_symbol_3 = aboutEqual(dstColor.g, srcColor.g);
+      bool tint_symbol_1 = tint_symbol_2;
+      if (tint_symbol_1) {
+        tint_symbol_1 = aboutEqual(dstColor.g, srcColor.g);
       }
-      success = tint_symbol_3;
+      success = tint_symbol_1;
     } else {
-      bool tint_symbol_8 = success;
-      if (tint_symbol_8) {
-        tint_symbol_8 = aboutEqual(dstColor.r, srcColor.r);
-      }
-      bool tint_symbol_7 = tint_symbol_8;
-      if (tint_symbol_7) {
-        tint_symbol_7 = aboutEqual(dstColor.g, srcColor.g);
-      }
-      bool tint_symbol_6 = tint_symbol_7;
+      bool tint_symbol_6 = success;
       if (tint_symbol_6) {
-        tint_symbol_6 = aboutEqual(dstColor.b, srcColor.b);
+        tint_symbol_6 = aboutEqual(dstColor.r, srcColor.r);
       }
       bool tint_symbol_5 = tint_symbol_6;
       if (tint_symbol_5) {
-        tint_symbol_5 = aboutEqual(dstColor.a, srcColor.a);
+        tint_symbol_5 = aboutEqual(dstColor.g, srcColor.g);
       }
-      success = tint_symbol_5;
+      bool tint_symbol_4 = tint_symbol_5;
+      if (tint_symbol_4) {
+        tint_symbol_4 = aboutEqual(dstColor.b, srcColor.b);
+      }
+      bool tint_symbol_3 = tint_symbol_4;
+      if (tint_symbol_3) {
+        tint_symbol_3 = aboutEqual(dstColor.a, srcColor.a);
+      }
+      success = tint_symbol_3;
     }
   }
   const uint outputIndex = ((GlobalInvocationID.y * dstSize.x) + GlobalInvocationID.x);
@@ -87,7 +87,7 @@ void main_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
-  main_inner(tint_symbol_1.GlobalInvocationID);
+void main(tint_symbol_8 tint_symbol_7) {
+  main_inner(tint_symbol_7.GlobalInvocationID);
   return;
 }

test/tint/bug/tint/913.wgsl.expected.fxc.hlsl

@@ -9,7 +9,7 @@ bool aboutEqual(float value, float expect) {
   return (abs((value - expect)) < 0.001f);
 }
 
-struct tint_symbol_2 {
+struct tint_symbol_8 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
@@ -49,33 +49,33 @@ void main_inner(uint3 GlobalInvocationID) {
     const float4 srcColor = src.Load(int3(int2(srcTexCoord), 0));
     const float4 dstColor = tint_symbol.Load(int3(int2(dstTexCoord), 0));
     if ((uniforms[0].y == 2u)) {
-      bool tint_symbol_4 = success;
-      if (tint_symbol_4) {
-        tint_symbol_4 = aboutEqual(dstColor.r, srcColor.r);
+      bool tint_symbol_2 = success;
+      if (tint_symbol_2) {
+        tint_symbol_2 = aboutEqual(dstColor.r, srcColor.r);
       }
-      bool tint_symbol_3 = tint_symbol_4;
-      if (tint_symbol_3) {
-        tint_symbol_3 = aboutEqual(dstColor.g, srcColor.g);
+      bool tint_symbol_1 = tint_symbol_2;
+      if (tint_symbol_1) {
+        tint_symbol_1 = aboutEqual(dstColor.g, srcColor.g);
       }
-      success = tint_symbol_3;
+      success = tint_symbol_1;
     } else {
-      bool tint_symbol_8 = success;
-      if (tint_symbol_8) {
-        tint_symbol_8 = aboutEqual(dstColor.r, srcColor.r);
-      }
-      bool tint_symbol_7 = tint_symbol_8;
-      if (tint_symbol_7) {
-        tint_symbol_7 = aboutEqual(dstColor.g, srcColor.g);
-      }
-      bool tint_symbol_6 = tint_symbol_7;
+      bool tint_symbol_6 = success;
       if (tint_symbol_6) {
-        tint_symbol_6 = aboutEqual(dstColor.b, srcColor.b);
+        tint_symbol_6 = aboutEqual(dstColor.r, srcColor.r);
       }
       bool tint_symbol_5 = tint_symbol_6;
       if (tint_symbol_5) {
-        tint_symbol_5 = aboutEqual(dstColor.a, srcColor.a);
+        tint_symbol_5 = aboutEqual(dstColor.g, srcColor.g);
       }
-      success = tint_symbol_5;
+      bool tint_symbol_4 = tint_symbol_5;
+      if (tint_symbol_4) {
+        tint_symbol_4 = aboutEqual(dstColor.b, srcColor.b);
+      }
+      bool tint_symbol_3 = tint_symbol_4;
+      if (tint_symbol_3) {
+        tint_symbol_3 = aboutEqual(dstColor.a, srcColor.a);
+      }
+      success = tint_symbol_3;
     }
   }
   const uint outputIndex = ((GlobalInvocationID.y * dstSize.x) + GlobalInvocationID.x);
@@ -87,7 +87,7 @@ void main_inner(uint3 GlobalInvocationID) {
 }
 
 [numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
-  main_inner(tint_symbol_1.GlobalInvocationID);
+void main(tint_symbol_8 tint_symbol_7) {
+  main_inner(tint_symbol_7.GlobalInvocationID);
   return;
 }

test/tint/bug/tint/914.wgsl.expected.dxc.hlsl

@@ -47,7 +47,7 @@ uint tint_div(uint lhs, uint rhs) {
   return (lhs / ((rhs == 0u) ? 1u : rhs));
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_3 {
   uint3 local_id : SV_GroupThreadID;
   uint local_invocation_index : SV_GroupIndex;
   uint3 global_id : SV_DispatchThreadID;
@@ -67,8 +67,7 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
   const uint tileCol = (local_id.x * 4u);
   const uint globalRow = (global_id.y * 4u);
   const uint globalCol = (global_id.x * 4u);
-  const uint tint_symbol_2 = tint_div((uniforms[0].y - 1u), 64u);
-  const uint numTiles = (tint_symbol_2 + 1u);
+  const uint numTiles = (tint_div((uniforms[0].y - 1u), 64u) + 1u);
   float acc[16] = (float[16])0;
   float ACached = 0.0f;
   float BCached[4] = (float[4])0;
@@ -89,8 +88,8 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
             for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
               const uint inputRow = (tileRow + innerRow);
               const uint inputCol = (tileColA + innerCol);
-              const float tint_symbol_3 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
-              mm_Asub[inputRow][inputCol] = tint_symbol_3;
+              const float tint_symbol = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
+              mm_Asub[inputRow][inputCol] = tint_symbol;
             }
           }
         }
@@ -101,8 +100,8 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
             for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
               const uint inputRow = (tileRowB + innerRow);
               const uint inputCol = (tileCol + innerCol);
-              const float tint_symbol_4 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
-              mm_Bsub[innerCol][inputCol] = tint_symbol_4;
+              const float tint_symbol_1 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
+              mm_Bsub[innerCol][inputCol] = tint_symbol_1;
             }
           }
         }
@@ -144,7 +143,7 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
 }
 
 [numthreads(16, 16, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  main_inner(tint_symbol.local_id, tint_symbol.global_id, tint_symbol.local_invocation_index);
+void main(tint_symbol_3 tint_symbol_2) {
+  main_inner(tint_symbol_2.local_id, tint_symbol_2.global_id, tint_symbol_2.local_invocation_index);
   return;
 }

test/tint/bug/tint/914.wgsl.expected.fxc.hlsl

@@ -47,7 +47,7 @@ uint tint_div(uint lhs, uint rhs) {
   return (lhs / ((rhs == 0u) ? 1u : rhs));
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_3 {
   uint3 local_id : SV_GroupThreadID;
   uint local_invocation_index : SV_GroupIndex;
   uint3 global_id : SV_DispatchThreadID;
@@ -67,8 +67,7 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
   const uint tileCol = (local_id.x * 4u);
   const uint globalRow = (global_id.y * 4u);
   const uint globalCol = (global_id.x * 4u);
-  const uint tint_symbol_2 = tint_div((uniforms[0].y - 1u), 64u);
-  const uint numTiles = (tint_symbol_2 + 1u);
+  const uint numTiles = (tint_div((uniforms[0].y - 1u), 64u) + 1u);
   float acc[16] = (float[16])0;
   float ACached = 0.0f;
   float BCached[4] = (float[4])0;
@@ -89,8 +88,8 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
             for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
               const uint inputRow = (tileRow + innerRow);
               const uint inputCol = (tileColA + innerCol);
-              const float tint_symbol_3 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
-              mm_Asub[inputRow][inputCol] = tint_symbol_3;
+              const float tint_symbol = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
+              mm_Asub[inputRow][inputCol] = tint_symbol;
             }
           }
         }
@@ -101,8 +100,8 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
             for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
               const uint inputRow = (tileRowB + innerRow);
               const uint inputCol = (tileCol + innerCol);
-              const float tint_symbol_4 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
-              mm_Bsub[innerCol][inputCol] = tint_symbol_4;
+              const float tint_symbol_1 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
+              mm_Bsub[innerCol][inputCol] = tint_symbol_1;
             }
           }
         }
@@ -144,7 +143,7 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
 }
 
 [numthreads(16, 16, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  main_inner(tint_symbol.local_id, tint_symbol.global_id, tint_symbol.local_invocation_index);
+void main(tint_symbol_3 tint_symbol_2) {
+  main_inner(tint_symbol_2.local_id, tint_symbol_2.global_id, tint_symbol_2.local_invocation_index);
   return;
 }

test/tint/bug/tint/914.wgsl.expected.glsl

@@ -78,8 +78,7 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
   uint tileCol = (local_id.x * 4u);
   uint globalRow = (global_id.y * 4u);
   uint globalCol = (global_id.x * 4u);
-  uint tint_symbol_1 = tint_div((uniforms.inner.dimInner - 1u), 64u);
-  uint numTiles = (tint_symbol_1 + 1u);
+  uint numTiles = (tint_div((uniforms.inner.dimInner - 1u), 64u) + 1u);
   float acc[16] = float[16](0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
   float ACached = 0.0f;
   float BCached[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
@@ -100,8 +99,8 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
             for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
               uint inputRow = (tileRow + innerRow);
               uint inputCol = (tileColA + innerCol);
-              float tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
-              mm_Asub[inputRow][inputCol] = tint_symbol_2;
+              float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol));
+              mm_Asub[inputRow][inputCol] = tint_symbol_1;
             }
           }
         }
@@ -112,8 +111,8 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
             for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
               uint inputRow = (tileRowB + innerRow);
               uint inputCol = (tileCol + innerCol);
-              float tint_symbol_3 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
-              mm_Bsub[innerCol][inputCol] = tint_symbol_3;
+              float tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol));
+              mm_Bsub[innerCol][inputCol] = tint_symbol_2;
             }
           }
         }

test/tint/bug/tint/914.wgsl.expected.msl

@@ -24,26 +24,26 @@ struct Matrix {
   /* 0x0000 */ tint_array<float, 1> numbers;
 };
 
-float mm_readA(uint row, uint col, const constant Uniforms* const tint_symbol_4, const device Matrix* const tint_symbol_5) {
-  if (((row < (*(tint_symbol_4)).dimAOuter) && (col < (*(tint_symbol_4)).dimInner))) {
-    float const result = (*(tint_symbol_5)).numbers[((row * (*(tint_symbol_4)).dimInner) + col)];
+float mm_readA(uint row, uint col, const constant Uniforms* const tint_symbol_3, const device Matrix* const tint_symbol_4) {
+  if (((row < (*(tint_symbol_3)).dimAOuter) && (col < (*(tint_symbol_3)).dimInner))) {
+    float const result = (*(tint_symbol_4)).numbers[((row * (*(tint_symbol_3)).dimInner) + col)];
     return result;
   }
   return 0.0f;
 }
 
-float mm_readB(uint row, uint col, const constant Uniforms* const tint_symbol_6, const device Matrix* const tint_symbol_7) {
-  if (((row < (*(tint_symbol_6)).dimInner) && (col < (*(tint_symbol_6)).dimBOuter))) {
-    float const result = (*(tint_symbol_7)).numbers[((row * (*(tint_symbol_6)).dimBOuter) + col)];
+float mm_readB(uint row, uint col, const constant Uniforms* const tint_symbol_5, const device Matrix* const tint_symbol_6) {
+  if (((row < (*(tint_symbol_5)).dimInner) && (col < (*(tint_symbol_5)).dimBOuter))) {
+    float const result = (*(tint_symbol_6)).numbers[((row * (*(tint_symbol_5)).dimBOuter) + col)];
     return result;
   }
   return 0.0f;
 }
 
-void mm_write(uint row, uint col, float value, const constant Uniforms* const tint_symbol_8, device Matrix* const tint_symbol_9) {
-  if (((row < (*(tint_symbol_8)).dimAOuter) && (col < (*(tint_symbol_8)).dimBOuter))) {
-    uint const index = (col + (row * (*(tint_symbol_8)).dimBOuter));
-    (*(tint_symbol_9)).numbers[index] = value;
+void mm_write(uint row, uint col, float value, const constant Uniforms* const tint_symbol_7, device Matrix* const tint_symbol_8) {
+  if (((row < (*(tint_symbol_7)).dimAOuter) && (col < (*(tint_symbol_7)).dimBOuter))) {
+    uint const index = (col + (row * (*(tint_symbol_7)).dimBOuter));
+    (*(tint_symbol_8)).numbers[index] = value;
   }
 }
 
@@ -51,20 +51,19 @@ uint tint_div(uint lhs, uint rhs) {
   return (lhs / select(rhs, 1u, (rhs == 0u)));
 }
 
-void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_10, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_11, const constant Uniforms* const tint_symbol_12, const device Matrix* const tint_symbol_13, const device Matrix* const tint_symbol_14, device Matrix* const tint_symbol_15) {
+void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_9, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_10, const constant Uniforms* const tint_symbol_11, const device Matrix* const tint_symbol_12, const device Matrix* const tint_symbol_13, device Matrix* const tint_symbol_14) {
   for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
     uint const i = (idx / 64u);
     uint const i_1 = (idx % 64u);
+    (*(tint_symbol_9))[i][i_1] = 0.0f;
     (*(tint_symbol_10))[i][i_1] = 0.0f;
-    (*(tint_symbol_11))[i][i_1] = 0.0f;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   uint const tileRow = (local_id[1] * 4u);
   uint const tileCol = (local_id[0] * 4u);
   uint const globalRow = (global_id[1] * 4u);
   uint const globalCol = (global_id[0] * 4u);
-  uint const tint_symbol_1 = tint_div(((*(tint_symbol_12)).dimInner - 1u), 64u);
-  uint const numTiles = (tint_symbol_1 + 1u);
+  uint const numTiles = (tint_div(((*(tint_symbol_11)).dimInner - 1u), 64u) + 1u);
   tint_array<float, 16> acc = {};
   float ACached = 0.0f;
   tint_array<float, 4> BCached = {};
@@ -80,25 +79,25 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in
       for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
         uint const inputRow = (tileRow + innerRow);
         uint const inputCol = (tileColA + innerCol);
-        float const tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_12, tint_symbol_13);
-        (*(tint_symbol_10))[inputRow][inputCol] = tint_symbol_2;
+        float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_11, tint_symbol_12);
+        (*(tint_symbol_9))[inputRow][inputCol] = tint_symbol_1;
       }
     }
     for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
       for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
         uint const inputRow = (tileRowB + innerRow);
         uint const inputCol = (tileCol + innerCol);
-        float const tint_symbol_3 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_symbol_12, tint_symbol_14);
-        (*(tint_symbol_11))[innerCol][inputCol] = tint_symbol_3;
+        float const tint_symbol_2 = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
+        (*(tint_symbol_10))[innerCol][inputCol] = tint_symbol_2;
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
     for(uint k = 0u; (k < 64u); k = (k + 1u)) {
       for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
-        BCached[inner] = (*(tint_symbol_11))[k][(tileCol + inner)];
+        BCached[inner] = (*(tint_symbol_10))[k][(tileCol + inner)];
       }
       for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
-        ACached = (*(tint_symbol_10))[(tileRow + innerRow)][k];
+        ACached = (*(tint_symbol_9))[(tileRow + innerRow)][k];
         for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
           uint const index = ((innerRow * 4u) + innerCol);
           acc[index] = (acc[index] + (ACached * BCached[innerCol]));
@@ -110,15 +109,15 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in
   for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
     for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
       uint const index = ((innerRow * 4u) + innerCol);
-      mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_12, tint_symbol_15);
+      mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_11, tint_symbol_14);
     }
   }
 }
 
-kernel void tint_symbol(const constant Uniforms* tint_symbol_18 [[buffer(0)]], const device Matrix* tint_symbol_19 [[buffer(2)]], const device Matrix* tint_symbol_20 [[buffer(3)]], device Matrix* tint_symbol_21 [[buffer(1)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(const constant Uniforms* tint_symbol_17 [[buffer(0)]], const device Matrix* tint_symbol_18 [[buffer(2)]], const device Matrix* tint_symbol_19 [[buffer(3)]], device Matrix* tint_symbol_20 [[buffer(1)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_15;
   threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_16;
-  threadgroup tint_array<tint_array<float, 64>, 64> tint_symbol_17;
-  tint_symbol_inner(local_id, global_id, local_invocation_index, &(tint_symbol_16), &(tint_symbol_17), tint_symbol_18, tint_symbol_19, tint_symbol_20, tint_symbol_21);
+  tint_symbol_inner(local_id, global_id, local_invocation_index, &(tint_symbol_15), &(tint_symbol_16), tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
   return;
 }
 

test/tint/bug/tint/980.wgsl.expected.dxc.hlsl

@@ -10,17 +10,17 @@ float3 Bad(uint index, float3 rd) {
 
 RWByteAddressBuffer io : register(u0, space0);
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint idx : SV_GroupIndex;
 };
 
 void main_inner(uint idx) {
-  const float3 tint_symbol_2 = Bad(io.Load(12u), asfloat(io.Load3(0u)));
-  io.Store3(0u, asuint(tint_symbol_2));
+  const float3 tint_symbol = Bad(io.Load(12u), asfloat(io.Load3(0u)));
+  io.Store3(0u, asuint(tint_symbol));
 }
 
 [numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  main_inner(tint_symbol.idx);
+void main(tint_symbol_2 tint_symbol_1) {
+  main_inner(tint_symbol_1.idx);
   return;
 }

test/tint/bug/tint/980.wgsl.expected.fxc.hlsl

@@ -10,17 +10,17 @@ float3 Bad(uint index, float3 rd) {
 
 RWByteAddressBuffer io : register(u0, space0);
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint idx : SV_GroupIndex;
 };
 
 void main_inner(uint idx) {
-  const float3 tint_symbol_2 = Bad(io.Load(12u), asfloat(io.Load3(0u)));
-  io.Store3(0u, asuint(tint_symbol_2));
+  const float3 tint_symbol = Bad(io.Load(12u), asfloat(io.Load3(0u)));
+  io.Store3(0u, asuint(tint_symbol));
 }
 
 [numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  main_inner(tint_symbol.idx);
+void main(tint_symbol_2 tint_symbol_1) {
+  main_inner(tint_symbol_1.idx);
   return;
 }

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl

@@ -20,11 +20,8 @@ void compute_main_inner(uint local_invocation_index_2) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
-    const uint tint_symbol_2 = tint_div(x_31, 2u);
-    const uint tint_symbol_3 = tint_mod(x_33, 2u);
-    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
+    InterlockedExchange(wg[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl

@@ -20,11 +20,8 @@ void compute_main_inner(uint local_invocation_index_2) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
-    const uint tint_symbol_2 = tint_div(x_31, 2u);
-    const uint tint_symbol_3 = tint_mod(x_33, 2u);
-    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
+    InterlockedExchange(wg[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl

@@ -21,10 +21,7 @@ void compute_main_inner(uint local_invocation_index_2) {
     uint x_31 = idx;
     uint x_33 = idx;
     uint x_35 = idx;
-    uint tint_symbol = tint_div(x_31, 2u);
-    uint tint_symbol_1 = tint_mod(x_33, 2u);
-    uint tint_symbol_2 = tint_mod(x_35, 1u);
-    atomicExchange(wg[tint_symbol][tint_symbol_1][tint_symbol_2], 0u);
+    atomicExchange(wg[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)], 0u);
     {
       uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl

@@ -22,7 +22,7 @@ uint tint_mod(uint lhs, uint rhs) {
   return (lhs % select(rhs, 1u, (rhs == 0u)));
 }
 
-void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index_2;
   while (true) {
@@ -33,42 +33,39 @@ void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<ti
     uint const x_31 = idx;
     uint const x_33 = idx;
     uint const x_35 = idx;
-    uint const tint_symbol = tint_div(x_31, 2u);
-    uint const tint_symbol_1 = tint_mod(x_33, 2u);
-    uint const tint_symbol_2 = tint_mod(x_35, 1u);
-    atomic_store_explicit(&((*(tint_symbol_3))[tint_symbol][tint_symbol_1][tint_symbol_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol))[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)]), 0u, memory_order_relaxed);
     {
       uint const x_42 = idx;
       idx = (x_42 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol_3))[2][1][0]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_4, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_5) {
-  uint const x_57 = *(tint_symbol_4);
-  compute_main_inner(x_57, tint_symbol_5);
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
+  uint const x_57 = *(tint_symbol_1);
+  compute_main_inner(x_57, tint_symbol_2);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_6, thread uint* const tint_symbol_7) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
-    atomic_store_explicit(&((*(tint_symbol_6))[i][i_1][i_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  *(tint_symbol_7) = local_invocation_index_1_param;
-  compute_main_1(tint_symbol_7, tint_symbol_6);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_8;
-  thread uint tint_symbol_9 = 0u;
-  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_8), &(tint_symbol_9));
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
   return;
 }
 

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.spvasm

@@ -96,10 +96,10 @@
          %48 = OpLoad %uint %idx
          %49 = OpLoad %uint %idx
          %50 = OpLoad %uint %idx
-         %51 = OpFunctionCall %uint %tint_div %48 %uint_2
-         %52 = OpFunctionCall %uint %tint_mod %49 %uint_2
-         %53 = OpFunctionCall %uint %tint_mod %50 %uint_1
-         %58 = OpAccessChain %_ptr_Workgroup_uint %wg %51 %52 %53
+         %54 = OpFunctionCall %uint %tint_div %48 %uint_2
+         %55 = OpFunctionCall %uint %tint_mod %49 %uint_2
+         %56 = OpFunctionCall %uint %tint_mod %50 %uint_1
+         %58 = OpAccessChain %_ptr_Workgroup_uint %wg %54 %55 %56
                OpAtomicStore %58 %uint_2 %uint_0 %6
                OpBranch %40
          %40 = OpLabel

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl

@@ -20,11 +20,8 @@ void compute_main_inner(uint local_invocation_index_2) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
-    const uint tint_symbol_2 = tint_div(x_31, 2u);
-    const uint tint_symbol_3 = tint_mod(x_33, 2u);
-    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
+    InterlockedExchange(wg[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl

@@ -20,11 +20,8 @@ void compute_main_inner(uint local_invocation_index_2) {
     const uint x_31 = idx;
     const uint x_33 = idx;
     const uint x_35 = idx;
-    const uint tint_symbol_2 = tint_div(x_31, 2u);
-    const uint tint_symbol_3 = tint_mod(x_33, 2u);
-    const uint tint_symbol_4 = tint_mod(x_35, 1u);
     uint atomic_result = 0u;
-    InterlockedExchange(wg[tint_symbol_2][tint_symbol_3][tint_symbol_4], 0u, atomic_result);
+    InterlockedExchange(wg[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)], 0u, atomic_result);
     {
       const uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl

@@ -21,10 +21,7 @@ void compute_main_inner(uint local_invocation_index_2) {
     uint x_31 = idx;
     uint x_33 = idx;
     uint x_35 = idx;
-    uint tint_symbol = tint_div(x_31, 2u);
-    uint tint_symbol_1 = tint_mod(x_33, 2u);
-    uint tint_symbol_2 = tint_mod(x_35, 1u);
-    atomicExchange(wg[tint_symbol][tint_symbol_1][tint_symbol_2], 0u);
+    atomicExchange(wg[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)], 0u);
     {
       uint x_42 = idx;
       idx = (x_42 + 1u);

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl

@@ -22,7 +22,7 @@ uint tint_mod(uint lhs, uint rhs) {
   return (lhs % select(rhs, 1u, (rhs == 0u)));
 }
 
-void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3) {
+void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index_2;
   while (true) {
@@ -33,42 +33,39 @@ void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<ti
     uint const x_31 = idx;
     uint const x_33 = idx;
     uint const x_35 = idx;
-    uint const tint_symbol = tint_div(x_31, 2u);
-    uint const tint_symbol_1 = tint_mod(x_33, 2u);
-    uint const tint_symbol_2 = tint_mod(x_35, 1u);
-    atomic_store_explicit(&((*(tint_symbol_3))[tint_symbol][tint_symbol_1][tint_symbol_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol))[tint_div(x_31, 2u)][tint_mod(x_33, 2u)][tint_mod(x_35, 1u)]), 0u, memory_order_relaxed);
     {
       uint const x_42 = idx;
       idx = (x_42 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol_3))[2][1][0]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_4, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_5) {
-  uint const x_57 = *(tint_symbol_4);
-  compute_main_inner(x_57, tint_symbol_5);
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
+  uint const x_57 = *(tint_symbol_1);
+  compute_main_inner(x_57, tint_symbol_2);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_6, thread uint* const tint_symbol_7) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
-    atomic_store_explicit(&((*(tint_symbol_6))[i][i_1][i_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  *(tint_symbol_7) = local_invocation_index_1_param;
-  compute_main_1(tint_symbol_7, tint_symbol_6);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_8;
-  thread uint tint_symbol_9 = 0u;
-  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_8), &(tint_symbol_9));
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
   return;
 }
 

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.spvasm

@@ -96,10 +96,10 @@
          %48 = OpLoad %uint %idx
          %49 = OpLoad %uint %idx
          %50 = OpLoad %uint %idx
-         %51 = OpFunctionCall %uint %tint_div %48 %uint_2
-         %52 = OpFunctionCall %uint %tint_mod %49 %uint_2
-         %53 = OpFunctionCall %uint %tint_mod %50 %uint_1
-         %58 = OpAccessChain %_ptr_Workgroup_uint %wg %51 %52 %53
+         %54 = OpFunctionCall %uint %tint_div %48 %uint_2
+         %55 = OpFunctionCall %uint %tint_mod %49 %uint_2
+         %56 = OpFunctionCall %uint %tint_mod %50 %uint_1
+         %58 = OpAccessChain %_ptr_Workgroup_uint %wg %54 %55 %56
                OpAtomicStore %58 %uint_2 %uint_0 %6
                OpBranch %40
          %40 = OpLabel

test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl

@@ -16,8 +16,8 @@ void atomicCompareExchangeWeak_e88938() {
   int atomic_compare_value = 1;
   InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resulti32 tint_symbol_2 = atomic_result;
-  const int old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resulti32 tint_symbol = atomic_result;
+  const int old_value_1 = tint_symbol.old_value;
   const int x_18 = old_value_1;
   const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)};
   res = tint_symbol_3;
@@ -38,7 +38,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -53,7 +53,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl

@@ -16,8 +16,8 @@ void atomicCompareExchangeWeak_e88938() {
   int atomic_compare_value = 1;
   InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resulti32 tint_symbol_2 = atomic_result;
-  const int old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resulti32 tint_symbol = atomic_result;
+  const int old_value_1 = tint_symbol.old_value;
   const int x_18 = old_value_1;
   const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)};
   res = tint_symbol_3;
@@ -38,7 +38,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -53,7 +53,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl

@@ -16,8 +16,8 @@ void atomicCompareExchangeWeak_83580d() {
   uint atomic_compare_value = 1u;
   InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resultu32 tint_symbol_2 = atomic_result;
-  const uint old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resultu32 tint_symbol = atomic_result;
+  const uint old_value_1 = tint_symbol.old_value;
   const uint x_17 = old_value_1;
   const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)};
   res = tint_symbol_3;
@@ -38,7 +38,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -53,7 +53,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl

@@ -16,8 +16,8 @@ void atomicCompareExchangeWeak_83580d() {
   uint atomic_compare_value = 1u;
   InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resultu32 tint_symbol_2 = atomic_result;
-  const uint old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resultu32 tint_symbol = atomic_result;
+  const uint old_value_1 = tint_symbol.old_value;
   const uint x_17 = old_value_1;
   const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)};
   res = tint_symbol_3;
@@ -38,7 +38,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -53,7 +53,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl

@@ -22,8 +22,8 @@ void atomicCompareExchangeWeak_e88938() {
   int atomic_compare_value = x_23;
   InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resulti32 tint_symbol_2 = atomic_result;
-  const int old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resulti32 tint_symbol = atomic_result;
+  const int old_value_1 = tint_symbol.old_value;
   const int x_24 = old_value_1;
   const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)};
   res = tint_symbol_3;
@@ -44,7 +44,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -59,7 +59,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl

@@ -22,8 +22,8 @@ void atomicCompareExchangeWeak_e88938() {
   int atomic_compare_value = x_23;
   InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resulti32 tint_symbol_2 = atomic_result;
-  const int old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resulti32 tint_symbol = atomic_result;
+  const int old_value_1 = tint_symbol.old_value;
   const int x_24 = old_value_1;
   const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)};
   res = tint_symbol_3;
@@ -44,7 +44,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -59,7 +59,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl

@@ -22,8 +22,8 @@ void atomicCompareExchangeWeak_83580d() {
   uint atomic_compare_value = x_22;
   InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resultu32 tint_symbol_2 = atomic_result;
-  const uint old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resultu32 tint_symbol = atomic_result;
+  const uint old_value_1 = tint_symbol.old_value;
   const uint x_23 = old_value_1;
   const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)};
   res = tint_symbol_3;
@@ -44,7 +44,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -59,7 +59,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl

@@ -22,8 +22,8 @@ void atomicCompareExchangeWeak_83580d() {
   uint atomic_compare_value = x_22;
   InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_value);
   atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  const atomic_compare_exchange_resultu32 tint_symbol_2 = atomic_result;
-  const uint old_value_1 = tint_symbol_2.old_value;
+  const atomic_compare_exchange_resultu32 tint_symbol = atomic_result;
+  const uint old_value_1 = tint_symbol.old_value;
   const uint x_23 = old_value_1;
   const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)};
   res = tint_symbol_3;
@@ -44,7 +44,7 @@ void compute_main_1() {
   return;
 }
 
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   uint local_invocation_index_1_param : SV_GroupIndex;
 };
 
@@ -59,7 +59,7 @@ void compute_main_inner_1(uint local_invocation_index_1_param) {
 }
 
 [numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+void compute_main(tint_symbol_2 tint_symbol_1) {
+  compute_main_inner_1(tint_symbol_1.local_invocation_index_1_param);
   return;
 }

test/tint/builtins/gen/literal/textureDimensions/cdc6c9.wgsl.expected.glsl

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -72,16 +80,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -89,6 +103,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -124,16 +140,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -141,6 +163,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {

test/tint/builtins/gen/literal/textureLoad/1bfdfb.wgsl.expected.glsl

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -63,7 +71,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, uve
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_1bfdfb() {
@@ -100,16 +108,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -117,6 +131,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -149,7 +165,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, uve
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_1bfdfb() {
@@ -180,16 +196,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -197,6 +219,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -229,7 +253,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, uve
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_1bfdfb() {

test/tint/builtins/gen/literal/textureLoad/8acf41.wgsl.expected.glsl

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -63,7 +71,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_8acf41() {
@@ -100,16 +108,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -117,6 +131,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -149,7 +165,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_8acf41() {
@@ -180,16 +196,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -197,6 +219,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -229,7 +253,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_8acf41() {

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

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 3, std140) uniform ext_tex_params_block_std140_ubo {
@@ -72,7 +80,7 @@ uniform highp sampler2D ext_tex_plane_1_1;
 uniform highp sampler2D arg_0_arg_1;
 uniform highp sampler2D ext_tex_plane_1_arg_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureSampleBaseClampToEdge_7c04e6() {
@@ -109,16 +117,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -126,6 +140,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 3, std140) uniform ext_tex_params_block_std140_ubo {
@@ -167,7 +183,7 @@ uniform highp sampler2D ext_tex_plane_1_1;
 uniform highp sampler2D arg_0_arg_1;
 uniform highp sampler2D ext_tex_plane_1_arg_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureSampleBaseClampToEdge_7c04e6() {
@@ -198,16 +214,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -215,6 +237,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 3, std140) uniform ext_tex_params_block_std140_ubo {
@@ -256,7 +280,7 @@ uniform highp sampler2D ext_tex_plane_1_1;
 uniform highp sampler2D arg_0_arg_1;
 uniform highp sampler2D ext_tex_plane_1_arg_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureSampleBaseClampToEdge_7c04e6() {

test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl

@@ -1,14 +1,14 @@
-groupshared uint arg_0;
-
-uint tint_workgroupUniformLoad_arg_0() {
+uint tint_workgroupUniformLoad(inout uint p) {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = arg_0;
+  const uint result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared uint arg_0;
+
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad_arg_0();
+  uint res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.fxc.hlsl

@@ -1,14 +1,14 @@
-groupshared uint arg_0;
-
-uint tint_workgroupUniformLoad_arg_0() {
+uint tint_workgroupUniformLoad(inout uint p) {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = arg_0;
+  const uint result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared uint arg_0;
+
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad_arg_0();
+  uint res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/37307c.wgsl.expected.glsl

@@ -1,15 +1,15 @@
 #version 310 es
 
-shared uint arg_0;
-uint tint_workgroupUniformLoad_arg_0() {
+uint tint_workgroupUniformLoad(inout uint p) {
   barrier();
-  uint result = arg_0;
+  uint result = p;
   barrier();
   return result;
 }
 
+shared uint arg_0;
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad_arg_0();
+  uint res = tint_workgroupUniformLoad(arg_0);
 }
 
 void compute_main(uint local_invocation_index) {

test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.dxc.hlsl

@@ -1,14 +1,14 @@
-groupshared float arg_0;
-
-float tint_workgroupUniformLoad_arg_0() {
+float tint_workgroupUniformLoad(inout float p) {
   GroupMemoryBarrierWithGroupSync();
-  const float result = arg_0;
+  const float result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared float arg_0;
+
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad_arg_0();
+  float res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.fxc.hlsl

@@ -1,14 +1,14 @@
-groupshared float arg_0;
-
-float tint_workgroupUniformLoad_arg_0() {
+float tint_workgroupUniformLoad(inout float p) {
   GroupMemoryBarrierWithGroupSync();
-  const float result = arg_0;
+  const float result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared float arg_0;
+
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad_arg_0();
+  float res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/7a857c.wgsl.expected.glsl

@@ -1,15 +1,15 @@
 #version 310 es
 
-shared float arg_0;
-float tint_workgroupUniformLoad_arg_0() {
+float tint_workgroupUniformLoad(inout float p) {
   barrier();
-  float result = arg_0;
+  float result = p;
   barrier();
   return result;
 }
 
+shared float arg_0;
 void workgroupUniformLoad_7a857c() {
-  float res = tint_workgroupUniformLoad_arg_0();
+  float res = tint_workgroupUniformLoad(arg_0);
 }
 
 void compute_main(uint local_invocation_index) {

test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.dxc.hlsl

@@ -1,14 +1,14 @@
-groupshared int arg_0;
-
-int tint_workgroupUniformLoad_arg_0() {
+int tint_workgroupUniformLoad(inout int p) {
   GroupMemoryBarrierWithGroupSync();
-  const int result = arg_0;
+  const int result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared int arg_0;
+
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad_arg_0();
+  int res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.fxc.hlsl

@@ -1,14 +1,14 @@
-groupshared int arg_0;
-
-int tint_workgroupUniformLoad_arg_0() {
+int tint_workgroupUniformLoad(inout int p) {
   GroupMemoryBarrierWithGroupSync();
-  const int result = arg_0;
+  const int result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared int arg_0;
+
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad_arg_0();
+  int res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/9d33de.wgsl.expected.glsl

@@ -1,15 +1,15 @@
 #version 310 es
 
-shared int arg_0;
-int tint_workgroupUniformLoad_arg_0() {
+int tint_workgroupUniformLoad(inout int p) {
   barrier();
-  int result = arg_0;
+  int result = p;
   barrier();
   return result;
 }
 
+shared int arg_0;
 void workgroupUniformLoad_9d33de() {
-  int res = tint_workgroupUniformLoad_arg_0();
+  int res = tint_workgroupUniformLoad(arg_0);
 }
 
 void compute_main(uint local_invocation_index) {

test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.dxc.hlsl

@@ -1,14 +1,14 @@
-groupshared float16_t arg_0;
-
-float16_t tint_workgroupUniformLoad_arg_0() {
+float16_t tint_workgroupUniformLoad(inout float16_t p) {
   GroupMemoryBarrierWithGroupSync();
-  const float16_t result = arg_0;
+  const float16_t result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared float16_t arg_0;
+
 void workgroupUniformLoad_e07d08() {
-  float16_t res = tint_workgroupUniformLoad_arg_0();
+  float16_t res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {

test/tint/builtins/gen/literal/workgroupUniformLoad/e07d08.wgsl.expected.glsl

@@ -1,16 +1,16 @@
 #version 310 es
 #extension GL_AMD_gpu_shader_half_float : require
 
-shared float16_t arg_0;
-float16_t tint_workgroupUniformLoad_arg_0() {
+float16_t tint_workgroupUniformLoad(inout float16_t p) {
   barrier();
-  float16_t result = arg_0;
+  float16_t result = p;
   barrier();
   return result;
 }
 
+shared float16_t arg_0;
 void workgroupUniformLoad_e07d08() {
-  float16_t res = tint_workgroupUniformLoad_arg_0();
+  float16_t res = tint_workgroupUniformLoad(arg_0);
 }
 
 void compute_main(uint local_invocation_index) {

test/tint/builtins/gen/var/textureDimensions/cdc6c9.wgsl.expected.glsl

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -72,16 +80,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -89,6 +103,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -124,16 +140,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -141,6 +163,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {

test/tint/builtins/gen/var/textureLoad/1bfdfb.wgsl.expected.glsl

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -63,7 +71,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, uve
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_1bfdfb() {
@@ -101,16 +109,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -118,6 +132,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -150,7 +166,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, uve
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_1bfdfb() {
@@ -182,16 +198,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -199,6 +221,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -231,7 +255,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, uve
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_1bfdfb() {

test/tint/builtins/gen/var/textureLoad/8acf41.wgsl.expected.glsl

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -63,7 +71,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_8acf41() {
@@ -101,16 +109,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -118,6 +132,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -150,7 +166,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_8acf41() {
@@ -182,16 +198,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -199,6 +221,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 2, std140) uniform ext_tex_params_block_std140_ubo {
@@ -231,7 +255,7 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive
 uniform highp sampler2D arg_0_1;
 uniform highp sampler2D ext_tex_plane_1_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureLoad_8acf41() {

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

@@ -14,16 +14,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -31,6 +37,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 3, std140) uniform ext_tex_params_block_std140_ubo {
@@ -72,7 +80,7 @@ uniform highp sampler2D ext_tex_plane_1_1;
 uniform highp sampler2D arg_0_arg_1;
 uniform highp sampler2D ext_tex_plane_1_arg_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureSampleBaseClampToEdge_7c04e6() {
@@ -110,16 +118,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -127,6 +141,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 3, std140) uniform ext_tex_params_block_std140_ubo {
@@ -168,7 +184,7 @@ uniform highp sampler2D ext_tex_plane_1_1;
 uniform highp sampler2D arg_0_arg_1;
 uniform highp sampler2D ext_tex_plane_1_arg_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureSampleBaseClampToEdge_7c04e6() {
@@ -200,16 +216,22 @@ struct GammaTransferParams {
 struct ExternalTextureParams {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
   mat3 gamutConversionMatrix;
   mat3x2 coordTransformationMatrix;
+  uint pad_2;
+  uint pad_3;
 };
 
 struct ExternalTextureParams_std140 {
   uint numPlanes;
   uint doYuvToRgbConversionOnly;
+  uint pad;
+  uint pad_1;
   mat3x4 yuvToRgbConversionMatrix;
   GammaTransferParams gammaDecodeParams;
   GammaTransferParams gammaEncodeParams;
@@ -217,6 +239,8 @@ struct ExternalTextureParams_std140 {
   vec2 coordTransformationMatrix_0;
   vec2 coordTransformationMatrix_1;
   vec2 coordTransformationMatrix_2;
+  uint pad_2;
+  uint pad_3;
 };
 
 layout(binding = 3, std140) uniform ext_tex_params_block_std140_ubo {
@@ -258,7 +282,7 @@ uniform highp sampler2D ext_tex_plane_1_1;
 uniform highp sampler2D arg_0_arg_1;
 uniform highp sampler2D ext_tex_plane_1_arg_1;
 ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
-  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2));
+  return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
 }
 
 void textureSampleBaseClampToEdge_7c04e6() {

test/tint/builtins/gen/var/workgroupUniformLoad/37307c.wgsl.expected.dxc.hlsl

@@ -1,14 +1,14 @@
-groupshared uint arg_0;
-
-uint tint_workgroupUniformLoad_arg_0() {
+uint tint_workgroupUniformLoad(inout uint p) {
   GroupMemoryBarrierWithGroupSync();
-  const uint result = arg_0;
+  const uint result = p;
   GroupMemoryBarrierWithGroupSync();
   return result;
 }
 
+groupshared uint arg_0;
+
 void workgroupUniformLoad_37307c() {
-  uint res = tint_workgroupUniformLoad_arg_0();
+  uint res = tint_workgroupUniformLoad(arg_0);
 }
 
 struct tint_symbol_1 {