Implement FXC workaround for vector access in loops resulting in failed loop unrolling

When indexing into vectors in a loop, FXC sometimes fails to determine
the max number of iterations when attempting to unroll the loop,
resulting in "error X3511: forced to unroll loop, but unrolling
failed.". We work around this by calling a function that sets the input
value at the input index into an inout vector. This seems to nudge FXC
enough for it to determine the number of loop iterations to unroll.

Bug: tint:534
Change-Id: I52cb209be29fcad8fbb91283c7be8c6e22e00656
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56140
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/val/hlsl.cc

@@ -20,9 +20,9 @@
 #include "src/utils/io/tmpfile.h"
 
 #ifdef _WIN32
+#include <Windows.h>
 #include <d3dcommon.h>
 #include <d3dcompiler.h>
-#include <windows.h>
 
 #include <wrl.h>
 using Microsoft::WRL::ComPtr;

src/writer/hlsl/generator_impl.cc

@@ -17,6 +17,7 @@
 #include <algorithm>
 #include <iomanip>
 #include <iosfwd>
+#include <set>
 #include <utility>
 #include <vector>
 
@@ -26,19 +27,23 @@
 #include "src/ast/interpolate_decoration.h"
 #include "src/ast/override_decoration.h"
 #include "src/ast/variable_decl_statement.h"
+#include "src/debug.h"
 #include "src/sem/array.h"
 #include "src/sem/atomic_type.h"
+#include "src/sem/block_statement.h"
 #include "src/sem/call.h"
 #include "src/sem/depth_texture_type.h"
 #include "src/sem/function.h"
 #include "src/sem/member_accessor_expression.h"
 #include "src/sem/multisampled_texture_type.h"
 #include "src/sem/sampled_texture_type.h"
+#include "src/sem/statement.h"
 #include "src/sem/storage_texture_type.h"
 #include "src/sem/struct.h"
 #include "src/sem/variable.h"
 #include "src/transform/calculate_array_length.h"
 #include "src/transform/hlsl.h"
+#include "src/utils/get_or_create.h"
 #include "src/utils/scoped_assignment.h"
 #include "src/writer/append_vector.h"
 #include "src/writer/float_to_string.h"
@@ -120,6 +125,10 @@ bool GeneratorImpl::Generate() {
   const TypeInfo* last_kind = nullptr;
   std::streampos last_padding_pos;
 
+  if (!FindAndEmitVectorAssignmentInLoopFunctions()) {
+    return false;
+  }
+
   for (auto* decl : builder_.AST().GlobalDeclarations()) {
     if (decl->Is<ast::Alias>()) {
       continue;  // Ignore aliases.
@@ -164,6 +173,112 @@ bool GeneratorImpl::Generate() {
   return true;
 }
 
+bool GeneratorImpl::FindAndEmitVectorAssignmentInLoopFunctions() {
+  auto is_in_loop = [](const sem::Expression* expr) {
+    auto* block = expr->Stmt()->Block();
+    if (!block) {
+      return false;
+    }
+    return block->FindFirstParent<sem::LoopBlockStatement>() != nullptr;
+  };
+
+  auto emit_function_once = [&](const sem::Vector* vec) {
+    utils::GetOrCreate(vector_assignment_in_loop_funcs_, vec, [&] {
+      std::ostringstream ss;
+      EmitType(ss, vec, tint::ast::StorageClass::kInvalid,
+               ast::Access::kUndefined, "");
+      auto func_name = UniqueIdentifier("Set_" + ss.str());
+      {
+        auto out = line();
+        out << "void " << func_name << "(inout ";
+        EmitType(out, vec, ast::StorageClass::kInvalid, ast::Access::kUndefined,
+                 "");
+        out << " vec, int idx, ";
+        EmitType(out, vec->type(), ast::StorageClass::kInvalid,
+                 ast::Access::kUndefined, "");
+        out << " val) {";
+      }
+      {
+        ScopedIndent si(this);
+        line() << "switch(idx) {";
+        {
+          ScopedIndent si2(this);
+          for (size_t i = 0; i < vec->size(); ++i) {
+            auto sidx = std::to_string(i);
+            line() << "case " + sidx + ": vec[" + sidx + "] = val; break;";
+          }
+        }
+        line() << "}";
+      }
+      line() << "}";
+      return func_name;
+    });
+  };
+
+  // Find vector assignments via an accessor expression (index) within loops so
+  // that we can replace them later with calls to setter functions. Also emit
+  // the setter functions per vector type as we find them. We do this to avoid
+  // having FCX fail to unroll loops with "error X3511: forced to unroll loop,
+  // but unrolling failed." See crbug.com/tint/534.
+
+  for (auto* ast_node : program_->ASTNodes().Objects()) {
+    auto* ast_assign = ast_node->As<ast::AssignmentStatement>();
+    if (!ast_assign) {
+      continue;
+    }
+
+    auto* ast_access_expr =
+        ast_assign->lhs()->As<ast::ArrayAccessorExpression>();
+    if (!ast_access_expr) {
+      continue;
+    }
+
+    auto* array_expr = builder_.Sem().Get(ast_access_expr->array());
+    auto* vec = array_expr->Type()->UnwrapRef()->As<sem::Vector>();
+
+    // Skip non-vectors
+    if (!vec) {
+      continue;
+    }
+
+    // Skip if not part of a loop
+    if (!is_in_loop(array_expr)) {
+      continue;
+    }
+
+    // Save this assignment along with the vector type
+    vector_assignments_in_loops_.emplace(ast_assign, vec);
+
+    // Emit the function if it hasn't already
+    emit_function_once(vec);
+  }
+
+  return true;
+}
+
+bool GeneratorImpl::EmitVectorAssignmentInLoopCall(
+    const ast::AssignmentStatement* stmt,
+    const sem::Vector* vec) {
+  auto* ast_access_expr = stmt->lhs()->As<ast::ArrayAccessorExpression>();
+
+  auto out = line();
+  out << vector_assignment_in_loop_funcs_.at(vec) << "(";
+  if (!EmitExpression(out, ast_access_expr->array())) {
+    return false;
+  }
+  out << ", ";
+  if (!EmitExpression(out, ast_access_expr->idx_expr())) {
+    return false;
+  }
+  out << ", ";
+  if (!EmitExpression(out, stmt->rhs())) {
+    return false;
+  }
+  out << ");";
+
+  return true;
+}
+
 bool GeneratorImpl::EmitArrayAccessor(std::ostream& out,
                                       ast::ArrayAccessorExpression* expr) {
   if (!EmitExpression(out, expr->array())) {
@@ -202,6 +317,11 @@ bool GeneratorImpl::EmitBitcast(std::ostream& out,
 }
 
 bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) {
+  auto iter = vector_assignments_in_loops_.find(stmt);
+  if (iter != vector_assignments_in_loops_.end()) {
+    return EmitVectorAssignmentInLoopCall(iter->first, iter->second);
+  }
+
   auto out = line();
   if (!EmitExpression(out, stmt->lhs())) {
     return false;

src/writer/hlsl/generator_impl.h

@@ -344,6 +344,21 @@ class GeneratorImpl : public TextGenerator {
   /// @returns true if the variable was emitted
   bool EmitProgramConstVariable(const ast::Variable* var);
 
+  /// Finds vector assignments via an accessor expression within loops, storing
+  /// the assignment/vector node pair in `vector_assignments_in_loops`, and
+  /// emits function definitions per vector type found. Required to work around
+  /// an FXC bug, see crbug.com/tint/534.
+  /// @returns true on success
+  bool FindAndEmitVectorAssignmentInLoopFunctions();
+  /// Emits call to vector assignment function for the input assignment
+  /// statement and vector type.
+  /// @param stmt assignment statement that corresponds to a vector assingment
+  /// via an accessor expression
+  /// @param vec the vector type being assigned to
+  /// @returns true on success
+  bool EmitVectorAssignmentInLoopCall(const ast::AssignmentStatement* stmt,
+                                      const sem::Vector* vec);
+
   /// Handles generating a builtin method name
   /// @param intrinsic the semantic info for the intrinsic
   /// @returns the name or "" if not valid
@@ -373,6 +388,10 @@ class GeneratorImpl : public TextGenerator {
 
   std::function<bool()> emit_continuing_;
   std::unordered_map<const sem::Struct*, std::string> structure_builders_;
+  std::unordered_map<const ast::AssignmentStatement*, const sem::Vector*>
+      vector_assignments_in_loops_;
+  std::unordered_map<const sem::Vector*, std::string>
+      vector_assignment_in_loop_funcs_;
 };
 
 }  // namespace hlsl

test/bug/tint/534.wgsl

@@ -0,0 +1,48 @@
+[[block]] struct Uniforms {
+    dstTextureFlipY : u32;
+    isFloat16 : u32;
+    isRGB10A2Unorm : u32;
+    channelCount : u32;
+};
+[[block]] struct OutputBuf {
+    result : [[stride(4)]] array<u32>;
+};
+[[group(0), binding(0)]] var src : texture_2d<f32>;
+[[group(0), binding(1)]] var dst : texture_2d<f32>;
+[[group(0), binding(2)]] var<storage_buffer, read_write> output : OutputBuf;
+[[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
+//[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
+// Fp16 logic
+// Infinity and NaN won't happen in this test case.
+fn ConvertToFp16FloatValue(fp32 : f32) -> u32 {
+    return 1u;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+    var size : vec2<i32> = textureDimensions(src);
+    var dstTexCoord : vec2<i32> = vec2<i32>(GlobalInvocationID.xy);
+    var srcTexCoord : vec2<i32> = dstTexCoord;
+    if (uniforms.dstTextureFlipY == 1u) {
+        srcTexCoord.y = size.y - dstTexCoord.y - 1;
+    }
+
+    var srcColor : vec4<f32> = textureLoad(src, srcTexCoord, 0);
+    var dstColor : vec4<f32> = textureLoad(dst, dstTexCoord, 0);
+    var success : bool = true;
+
+    var srcColorBits : vec4<u32>;
+    var dstColorBits : vec4<u32> = vec4<u32>(dstColor);
+
+    for (var i : u32 = 0u; i < uniforms.channelCount; i = i + 1u) {
+            srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+            success = success && (srcColorBits[i] == dstColorBits[i]);
+    }
+
+    var outputIndex : u32 = GlobalInvocationID.y * u32(size.x) + GlobalInvocationID.x;
+    if (success) {
+        output.result[outputIndex] = u32(1);
+    } else {
+        output.result[outputIndex] = u32(0);
+    }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl

@@ -0,0 +1,21 @@
+var<private> v2f : vec2<f32>;
+var<private> v3i : vec3<i32>;
+var<private> v4u : vec4<u32>;
+var<private> v2b : vec2<bool>;
+
+fn foo() {
+  for (var i : i32 = 0; i < 2; i = i + 1) {
+    v2f[i] = 1.0;
+    v3i[i] = 1;
+    v4u[i] = 1u;
+    v2b[i] = true;
+  }
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+
+  for (var i : i32 = 0; i < 2; i = i + 1) {
+      foo();
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl

@@ -0,0 +1,66 @@
+void Set_float2(inout float2 vec, int idx, float val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_int3(inout int3 vec, int idx, int val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+  }
+}
+void Set_uint4(inout uint4 vec, int idx, uint val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+    case 3: vec[3] = val; break;
+  }
+}
+void Set_bool2(inout bool2 vec, int idx, bool val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+static float2 v2f = float2(0.0f, 0.0f);
+static int3 v3i = int3(0, 0, 0);
+static uint4 v4u = uint4(0u, 0u, 0u, 0u);
+static bool2 v2b = bool2(false, false);
+
+void foo() {
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      Set_float2(v2f, i, 1.0f);
+      Set_int3(v3i, i, 1);
+      Set_uint4(v4u, i, 1u);
+      Set_bool2(v2b, i, true);
+      {
+        i = (i + 1);
+      }
+    }
+  }
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      foo();
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl

@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) {
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      (*(tint_symbol_1))[i] = 1.0f;
+      (*(tint_symbol_2))[i] = 1;
+      (*(tint_symbol_3))[i] = 1u;
+      (*(tint_symbol_4))[i] = true;
+      {
+        i = (i + 1);
+      }
+    }
+  }
+}
+
+kernel void tint_symbol() {
+  thread float2 tint_symbol_5 = 0.0f;
+  thread int3 tint_symbol_6 = 0;
+  thread uint4 tint_symbol_7 = 0u;
+  thread bool2 tint_symbol_8 = false;
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}
+

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm

@@ -0,0 +1,116 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 72
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %v2f "v2f"
+               OpName %v3i "v3i"
+               OpName %v4u "v4u"
+               OpName %v2b "v2b"
+               OpName %foo "foo"
+               OpName %i "i"
+               OpName %main "main"
+               OpName %i_0 "i"
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%_ptr_Private_v2float = OpTypePointer Private %v2float
+          %5 = OpConstantNull %v2float
+        %v2f = OpVariable %_ptr_Private_v2float Private %5
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Private_v3int = OpTypePointer Private %v3int
+         %10 = OpConstantNull %v3int
+        %v3i = OpVariable %_ptr_Private_v3int Private %10
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Private_v4uint = OpTypePointer Private %v4uint
+         %15 = OpConstantNull %v4uint
+        %v4u = OpVariable %_ptr_Private_v4uint Private %15
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+%_ptr_Private_v2bool = OpTypePointer Private %v2bool
+         %20 = OpConstantNull %v2bool
+        %v2b = OpVariable %_ptr_Private_v2bool Private %20
+       %void = OpTypeVoid
+         %21 = OpTypeFunction %void
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %28 = OpConstantNull %int
+      %int_2 = OpConstant %int 2
+%_ptr_Private_float = OpTypePointer Private %float
+    %float_1 = OpConstant %float 1
+%_ptr_Private_int = OpTypePointer Private %int
+      %int_1 = OpConstant %int 1
+%_ptr_Private_uint = OpTypePointer Private %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Private_bool = OpTypePointer Private %bool
+       %true = OpConstantTrue %bool
+        %foo = OpFunction %void None %21
+         %24 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %28
+               OpStore %i %int_0
+               OpBranch %29
+         %29 = OpLabel
+               OpLoopMerge %30 %31 None
+               OpBranch %32
+         %32 = OpLabel
+         %34 = OpLoad %int %i
+         %36 = OpSLessThan %bool %34 %int_2
+         %33 = OpLogicalNot %bool %36
+               OpSelectionMerge %37 None
+               OpBranchConditional %33 %38 %37
+         %38 = OpLabel
+               OpBranch %30
+         %37 = OpLabel
+         %39 = OpLoad %int %i
+         %41 = OpAccessChain %_ptr_Private_float %v2f %39
+               OpStore %41 %float_1
+         %43 = OpLoad %int %i
+         %45 = OpAccessChain %_ptr_Private_int %v3i %43
+               OpStore %45 %int_1
+         %47 = OpLoad %int %i
+         %49 = OpAccessChain %_ptr_Private_uint %v4u %47
+               OpStore %49 %uint_1
+         %51 = OpLoad %int %i
+         %53 = OpAccessChain %_ptr_Private_bool %v2b %51
+               OpStore %53 %true
+               OpBranch %31
+         %31 = OpLabel
+         %55 = OpLoad %int %i
+         %56 = OpIAdd %int %55 %int_1
+               OpStore %i %56
+               OpBranch %29
+         %30 = OpLabel
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %21
+         %58 = OpLabel
+        %i_0 = OpVariable %_ptr_Function_int Function %28
+               OpStore %i_0 %int_0
+               OpBranch %60
+         %60 = OpLabel
+               OpLoopMerge %61 %62 None
+               OpBranch %63
+         %63 = OpLabel
+         %65 = OpLoad %int %i_0
+         %66 = OpSLessThan %bool %65 %int_2
+         %64 = OpLogicalNot %bool %66
+               OpSelectionMerge %67 None
+               OpBranchConditional %64 %68 %67
+         %68 = OpLabel
+               OpBranch %61
+         %67 = OpLabel
+         %69 = OpFunctionCall %void %foo
+               OpBranch %62
+         %62 = OpLabel
+         %70 = OpLoad %int %i_0
+         %71 = OpIAdd %int %70 %int_1
+               OpStore %i_0 %71
+               OpBranch %60
+         %61 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl

@@ -0,0 +1,43 @@
+var<private> v2f : vec2<f32>;
+
+var<private> v3i : vec3<i32>;
+
+var<private> v4u : vec4<u32>;
+
+var<private> v2b : vec2<bool>;
+
+fn foo() {
+  {
+    var i : i32 = 0;
+    loop {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0;
+      v3i[i] = 1;
+      v4u[i] = 1u;
+      v2b[i] = true;
+
+      continuing {
+        i = (i + 1);
+      }
+    }
+  }
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  {
+    var i : i32 = 0;
+    loop {
+      if (!((i < 2))) {
+        break;
+      }
+      foo();
+
+      continuing {
+        i = (i + 1);
+      }
+    }
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl

@@ -0,0 +1,20 @@
+var<private> v2f : vec2<f32>;
+var<private> v3i : vec3<i32>;
+var<private> v4u : vec4<u32>;
+var<private> v2b : vec2<bool>;
+
+fn foo() {
+  var i : i32 = 0;
+  v2f[i] = 1.0;
+  v3i[i] = 1;
+  v4u[i] = 1u;
+  v2b[i] = true;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+
+  for (var i : i32 = 0; i < 2; i = i + 1) {
+      foo();
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl

@@ -0,0 +1,29 @@
+static float2 v2f = float2(0.0f, 0.0f);
+static int3 v3i = int3(0, 0, 0);
+static uint4 v4u = uint4(0u, 0u, 0u, 0u);
+static bool2 v2b = bool2(false, false);
+
+void foo() {
+  int i = 0;
+  v2f[i] = 1.0f;
+  v3i[i] = 1;
+  v4u[i] = 1u;
+  v2b[i] = true;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      foo();
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl

@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) {
+  int i = 0;
+  (*(tint_symbol_1))[i] = 1.0f;
+  (*(tint_symbol_2))[i] = 1;
+  (*(tint_symbol_3))[i] = 1u;
+  (*(tint_symbol_4))[i] = true;
+}
+
+kernel void tint_symbol() {
+  thread float2 tint_symbol_5 = 0.0f;
+  thread int3 tint_symbol_6 = 0;
+  thread uint4 tint_symbol_7 = 0u;
+  thread bool2 tint_symbol_8 = false;
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}
+

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm

@@ -0,0 +1,96 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 61
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %v2f "v2f"
+               OpName %v3i "v3i"
+               OpName %v4u "v4u"
+               OpName %v2b "v2b"
+               OpName %foo "foo"
+               OpName %i "i"
+               OpName %main "main"
+               OpName %i_0 "i"
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%_ptr_Private_v2float = OpTypePointer Private %v2float
+          %5 = OpConstantNull %v2float
+        %v2f = OpVariable %_ptr_Private_v2float Private %5
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Private_v3int = OpTypePointer Private %v3int
+         %10 = OpConstantNull %v3int
+        %v3i = OpVariable %_ptr_Private_v3int Private %10
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Private_v4uint = OpTypePointer Private %v4uint
+         %15 = OpConstantNull %v4uint
+        %v4u = OpVariable %_ptr_Private_v4uint Private %15
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+%_ptr_Private_v2bool = OpTypePointer Private %v2bool
+         %20 = OpConstantNull %v2bool
+        %v2b = OpVariable %_ptr_Private_v2bool Private %20
+       %void = OpTypeVoid
+         %21 = OpTypeFunction %void
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %28 = OpConstantNull %int
+%_ptr_Private_float = OpTypePointer Private %float
+    %float_1 = OpConstant %float 1
+%_ptr_Private_int = OpTypePointer Private %int
+      %int_1 = OpConstant %int 1
+%_ptr_Private_uint = OpTypePointer Private %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Private_bool = OpTypePointer Private %bool
+       %true = OpConstantTrue %bool
+      %int_2 = OpConstant %int 2
+        %foo = OpFunction %void None %21
+         %24 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %28
+               OpStore %i %int_0
+         %29 = OpLoad %int %i
+         %31 = OpAccessChain %_ptr_Private_float %v2f %29
+               OpStore %31 %float_1
+         %33 = OpLoad %int %i
+         %35 = OpAccessChain %_ptr_Private_int %v3i %33
+               OpStore %35 %int_1
+         %37 = OpLoad %int %i
+         %39 = OpAccessChain %_ptr_Private_uint %v4u %37
+               OpStore %39 %uint_1
+         %41 = OpLoad %int %i
+         %43 = OpAccessChain %_ptr_Private_bool %v2b %41
+               OpStore %43 %true
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %21
+         %46 = OpLabel
+        %i_0 = OpVariable %_ptr_Function_int Function %28
+               OpStore %i_0 %int_0
+               OpBranch %48
+         %48 = OpLabel
+               OpLoopMerge %49 %50 None
+               OpBranch %51
+         %51 = OpLabel
+         %53 = OpLoad %int %i_0
+         %55 = OpSLessThan %bool %53 %int_2
+         %52 = OpLogicalNot %bool %55
+               OpSelectionMerge %56 None
+               OpBranchConditional %52 %57 %56
+         %57 = OpLabel
+               OpBranch %49
+         %56 = OpLabel
+         %58 = OpFunctionCall %void %foo
+               OpBranch %50
+         %50 = OpLabel
+         %59 = OpLoad %int %i_0
+         %60 = OpIAdd %int %59 %int_1
+               OpStore %i_0 %60
+               OpBranch %48
+         %49 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl

@@ -0,0 +1,32 @@
+var<private> v2f : vec2<f32>;
+
+var<private> v3i : vec3<i32>;
+
+var<private> v4u : vec4<u32>;
+
+var<private> v2b : vec2<bool>;
+
+fn foo() {
+  var i : i32 = 0;
+  v2f[i] = 1.0;
+  v3i[i] = 1;
+  v4u[i] = 1u;
+  v2b[i] = true;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  {
+    var i : i32 = 0;
+    loop {
+      if (!((i < 2))) {
+        break;
+      }
+      foo();
+
+      continuing {
+        i = (i + 1);
+      }
+    }
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl

@@ -0,0 +1,30 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v3f : vec3<f32>;
+  var v4f : vec4<f32>;
+  var v2i : vec2<i32>;
+  var v3i : vec3<i32>;
+  var v4i : vec4<i32>;
+  var v2u : vec2<u32>;
+  var v3u : vec3<u32>;
+  var v4u : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v3b : vec3<bool>;
+  var v4b : vec4<bool>;
+
+  for (var i : i32 = 0; i < 2; i = i + 1) {
+    v2f[i] = 1.0;
+    v3f[i] = 1.0;
+    v4f[i] = 1.0;
+    v2i[i] = 1;
+    v3i[i] = 1;
+    v4i[i] = 1;
+    v2u[i] = 1u;
+    v3u[i] = 1u;
+    v4u[i] = 1u;
+    v2b[i] = true;
+    v3b[i] = true;
+    v4b[i] = true;
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl

@@ -0,0 +1,123 @@
+void Set_float2(inout float2 vec, int idx, float val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_float3(inout float3 vec, int idx, float val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+  }
+}
+void Set_float4(inout float4 vec, int idx, float val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+    case 3: vec[3] = val; break;
+  }
+}
+void Set_int2(inout int2 vec, int idx, int val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_int3(inout int3 vec, int idx, int val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+  }
+}
+void Set_int4(inout int4 vec, int idx, int val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+    case 3: vec[3] = val; break;
+  }
+}
+void Set_uint2(inout uint2 vec, int idx, uint val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_uint3(inout uint3 vec, int idx, uint val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+  }
+}
+void Set_uint4(inout uint4 vec, int idx, uint val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+    case 3: vec[3] = val; break;
+  }
+}
+void Set_bool2(inout bool2 vec, int idx, bool val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_bool3(inout bool3 vec, int idx, bool val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+  }
+}
+void Set_bool4(inout bool4 vec, int idx, bool val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+    case 3: vec[3] = val; break;
+  }
+}
+[numthreads(1, 1, 1)]
+void main() {
+  float2 v2f = float2(0.0f, 0.0f);
+  float3 v3f = float3(0.0f, 0.0f, 0.0f);
+  float4 v4f = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  int2 v2i = int2(0, 0);
+  int3 v3i = int3(0, 0, 0);
+  int4 v4i = int4(0, 0, 0, 0);
+  uint2 v2u = uint2(0u, 0u);
+  uint3 v3u = uint3(0u, 0u, 0u);
+  uint4 v4u = uint4(0u, 0u, 0u, 0u);
+  bool2 v2b = bool2(false, false);
+  bool3 v3b = bool3(false, false, false);
+  bool4 v4b = bool4(false, false, false, false);
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      Set_float2(v2f, i, 1.0f);
+      Set_float3(v3f, i, 1.0f);
+      Set_float4(v4f, i, 1.0f);
+      Set_int2(v2i, i, 1);
+      Set_int3(v3i, i, 1);
+      Set_int4(v4i, i, 1);
+      Set_uint2(v2u, i, 1u);
+      Set_uint3(v3u, i, 1u);
+      Set_uint4(v4u, i, 1u);
+      Set_bool2(v2b, i, true);
+      Set_bool3(v3b, i, true);
+      Set_bool4(v4b, i, true);
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl

@@ -0,0 +1,42 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  float2 v2f = 0.0f;
+  float3 v3f = 0.0f;
+  float4 v4f = 0.0f;
+  int2 v2i = 0;
+  int3 v3i = 0;
+  int4 v4i = 0;
+  uint2 v2u = 0u;
+  uint3 v3u = 0u;
+  uint4 v4u = 0u;
+  bool2 v2b = false;
+  bool3 v3b = false;
+  bool4 v4b = false;
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0f;
+      v3f[i] = 1.0f;
+      v4f[i] = 1.0f;
+      v2i[i] = 1;
+      v3i[i] = 1;
+      v4i[i] = 1;
+      v2u[i] = 1u;
+      v3u[i] = 1u;
+      v4u[i] = 1u;
+      v2b[i] = true;
+      v3b[i] = true;
+      v4b[i] = true;
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}
+

test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm

@@ -0,0 +1,150 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 104
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %v2f "v2f"
+               OpName %v3f "v3f"
+               OpName %v4f "v4f"
+               OpName %v2i "v2i"
+               OpName %v3i "v3i"
+               OpName %v4i "v4i"
+               OpName %v2u "v2u"
+               OpName %v3u "v3u"
+               OpName %v4u "v4u"
+               OpName %v2b "v2b"
+               OpName %v3b "v3b"
+               OpName %v4b "v4b"
+               OpName %i "i"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+          %9 = OpConstantNull %v2float
+    %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %13 = OpConstantNull %v3float
+    %v4float = OpTypeVector %float 4
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %17 = OpConstantNull %v4float
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %22 = OpConstantNull %v2int
+      %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %26 = OpConstantNull %v3int
+      %v4int = OpTypeVector %int 4
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %30 = OpConstantNull %v4int
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %35 = OpConstantNull %v2uint
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %39 = OpConstantNull %v3uint
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %43 = OpConstantNull %v4uint
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+%_ptr_Function_v2bool = OpTypePointer Function %v2bool
+         %48 = OpConstantNull %v2bool
+     %v3bool = OpTypeVector %bool 3
+%_ptr_Function_v3bool = OpTypePointer Function %v3bool
+         %52 = OpConstantNull %v3bool
+     %v4bool = OpTypeVector %bool 4
+%_ptr_Function_v4bool = OpTypePointer Function %v4bool
+         %56 = OpConstantNull %v4bool
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %60 = OpConstantNull %int
+      %int_2 = OpConstant %int 2
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_bool = OpTypePointer Function %bool
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+        %v2f = OpVariable %_ptr_Function_v2float Function %9
+        %v3f = OpVariable %_ptr_Function_v3float Function %13
+        %v4f = OpVariable %_ptr_Function_v4float Function %17
+        %v2i = OpVariable %_ptr_Function_v2int Function %22
+        %v3i = OpVariable %_ptr_Function_v3int Function %26
+        %v4i = OpVariable %_ptr_Function_v4int Function %30
+        %v2u = OpVariable %_ptr_Function_v2uint Function %35
+        %v3u = OpVariable %_ptr_Function_v3uint Function %39
+        %v4u = OpVariable %_ptr_Function_v4uint Function %43
+        %v2b = OpVariable %_ptr_Function_v2bool Function %48
+        %v3b = OpVariable %_ptr_Function_v3bool Function %52
+        %v4b = OpVariable %_ptr_Function_v4bool Function %56
+          %i = OpVariable %_ptr_Function_int Function %60
+               OpStore %i %int_0
+               OpBranch %61
+         %61 = OpLabel
+               OpLoopMerge %62 %63 None
+               OpBranch %64
+         %64 = OpLabel
+         %66 = OpLoad %int %i
+         %68 = OpSLessThan %bool %66 %int_2
+         %65 = OpLogicalNot %bool %68
+               OpSelectionMerge %69 None
+               OpBranchConditional %65 %70 %69
+         %70 = OpLabel
+               OpBranch %62
+         %69 = OpLabel
+         %71 = OpLoad %int %i
+         %73 = OpAccessChain %_ptr_Function_float %v2f %71
+               OpStore %73 %float_1
+         %75 = OpLoad %int %i
+         %76 = OpAccessChain %_ptr_Function_float %v3f %75
+               OpStore %76 %float_1
+         %77 = OpLoad %int %i
+         %78 = OpAccessChain %_ptr_Function_float %v4f %77
+               OpStore %78 %float_1
+         %79 = OpLoad %int %i
+         %80 = OpAccessChain %_ptr_Function_int %v2i %79
+               OpStore %80 %int_1
+         %82 = OpLoad %int %i
+         %83 = OpAccessChain %_ptr_Function_int %v3i %82
+               OpStore %83 %int_1
+         %84 = OpLoad %int %i
+         %85 = OpAccessChain %_ptr_Function_int %v4i %84
+               OpStore %85 %int_1
+         %86 = OpLoad %int %i
+         %88 = OpAccessChain %_ptr_Function_uint %v2u %86
+               OpStore %88 %uint_1
+         %90 = OpLoad %int %i
+         %91 = OpAccessChain %_ptr_Function_uint %v3u %90
+               OpStore %91 %uint_1
+         %92 = OpLoad %int %i
+         %93 = OpAccessChain %_ptr_Function_uint %v4u %92
+               OpStore %93 %uint_1
+         %94 = OpLoad %int %i
+         %96 = OpAccessChain %_ptr_Function_bool %v2b %94
+               OpStore %96 %true
+         %98 = OpLoad %int %i
+         %99 = OpAccessChain %_ptr_Function_bool %v3b %98
+               OpStore %99 %true
+        %100 = OpLoad %int %i
+        %101 = OpAccessChain %_ptr_Function_bool %v4b %100
+               OpStore %101 %true
+               OpBranch %63
+         %63 = OpLabel
+        %102 = OpLoad %int %i
+        %103 = OpIAdd %int %102 %int_1
+               OpStore %i %103
+               OpBranch %61
+         %62 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl

@@ -0,0 +1,39 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v3f : vec3<f32>;
+  var v4f : vec4<f32>;
+  var v2i : vec2<i32>;
+  var v3i : vec3<i32>;
+  var v4i : vec4<i32>;
+  var v2u : vec2<u32>;
+  var v3u : vec3<u32>;
+  var v4u : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v3b : vec3<bool>;
+  var v4b : vec4<bool>;
+  {
+    var i : i32 = 0;
+    loop {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0;
+      v3f[i] = 1.0;
+      v4f[i] = 1.0;
+      v2i[i] = 1;
+      v3i[i] = 1;
+      v4i[i] = 1;
+      v2u[i] = 1u;
+      v3u[i] = 1u;
+      v4u[i] = 1u;
+      v2b[i] = true;
+      v3b[i] = true;
+      v4b[i] = true;
+
+      continuing {
+        i = (i + 1);
+      }
+    }
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl

@@ -0,0 +1,26 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v2f_2 : vec2<f32>;
+
+  var v3i : vec3<i32>;
+  var v3i_2 : vec3<i32>;
+
+  var v4u : vec4<u32>;
+  var v4u_2 : vec4<u32>;
+
+  var v2b : vec2<bool>;
+  var v2b_2 : vec2<bool>;
+
+  for (var i : i32 = 0; i < 2; i = i + 1) {
+    v2f[i] = 1.0;
+    v3i[i] = 1;
+    v4u[i] = 1u;
+    v2b[i] = true;
+
+    v2f_2[i] = 1.0;
+    v3i_2[i] = 1;
+    v4u_2[i] = 1u;
+    v2b_2[i] = true;
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl

@@ -0,0 +1,58 @@
+void Set_float2(inout float2 vec, int idx, float val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_int3(inout int3 vec, int idx, int val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+  }
+}
+void Set_uint4(inout uint4 vec, int idx, uint val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+    case 2: vec[2] = val; break;
+    case 3: vec[3] = val; break;
+  }
+}
+void Set_bool2(inout bool2 vec, int idx, bool val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+[numthreads(1, 1, 1)]
+void main() {
+  float2 v2f = float2(0.0f, 0.0f);
+  float2 v2f_2 = float2(0.0f, 0.0f);
+  int3 v3i = int3(0, 0, 0);
+  int3 v3i_2 = int3(0, 0, 0);
+  uint4 v4u = uint4(0u, 0u, 0u, 0u);
+  uint4 v4u_2 = uint4(0u, 0u, 0u, 0u);
+  bool2 v2b = bool2(false, false);
+  bool2 v2b_2 = bool2(false, false);
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      Set_float2(v2f, i, 1.0f);
+      Set_int3(v3i, i, 1);
+      Set_uint4(v4u, i, 1u);
+      Set_bool2(v2b, i, true);
+      Set_float2(v2f_2, i, 1.0f);
+      Set_int3(v3i_2, i, 1);
+      Set_uint4(v4u_2, i, 1u);
+      Set_bool2(v2b_2, i, true);
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl

@@ -0,0 +1,34 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  float2 v2f = 0.0f;
+  float2 v2f_2 = 0.0f;
+  int3 v3i = 0;
+  int3 v3i_2 = 0;
+  uint4 v4u = 0u;
+  uint4 v4u_2 = 0u;
+  bool2 v2b = false;
+  bool2 v2b_2 = false;
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0f;
+      v3i[i] = 1;
+      v4u[i] = 1u;
+      v2b[i] = true;
+      v2f_2[i] = 1.0f;
+      v3i_2[i] = 1;
+      v4u_2[i] = 1u;
+      v2b_2[i] = true;
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  return;
+}
+

test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm

@@ -0,0 +1,106 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %v2f "v2f"
+               OpName %v2f_2 "v2f_2"
+               OpName %v3i "v3i"
+               OpName %v3i_2 "v3i_2"
+               OpName %v4u "v4u"
+               OpName %v4u_2 "v4u_2"
+               OpName %v2b "v2b"
+               OpName %v2b_2 "v2b_2"
+               OpName %i "i"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+          %9 = OpConstantNull %v2float
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %15 = OpConstantNull %v3int
+       %uint = OpTypeInt 32 0
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %21 = OpConstantNull %v4uint
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+%_ptr_Function_v2bool = OpTypePointer Function %v2bool
+         %27 = OpConstantNull %v2bool
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %32 = OpConstantNull %int
+      %int_2 = OpConstant %int 2
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_bool = OpTypePointer Function %bool
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+        %v2f = OpVariable %_ptr_Function_v2float Function %9
+      %v2f_2 = OpVariable %_ptr_Function_v2float Function %9
+        %v3i = OpVariable %_ptr_Function_v3int Function %15
+      %v3i_2 = OpVariable %_ptr_Function_v3int Function %15
+        %v4u = OpVariable %_ptr_Function_v4uint Function %21
+      %v4u_2 = OpVariable %_ptr_Function_v4uint Function %21
+        %v2b = OpVariable %_ptr_Function_v2bool Function %27
+      %v2b_2 = OpVariable %_ptr_Function_v2bool Function %27
+          %i = OpVariable %_ptr_Function_int Function %32
+               OpStore %i %int_0
+               OpBranch %33
+         %33 = OpLabel
+               OpLoopMerge %34 %35 None
+               OpBranch %36
+         %36 = OpLabel
+         %38 = OpLoad %int %i
+         %40 = OpSLessThan %bool %38 %int_2
+         %37 = OpLogicalNot %bool %40
+               OpSelectionMerge %41 None
+               OpBranchConditional %37 %42 %41
+         %42 = OpLabel
+               OpBranch %34
+         %41 = OpLabel
+         %43 = OpLoad %int %i
+         %45 = OpAccessChain %_ptr_Function_float %v2f %43
+               OpStore %45 %float_1
+         %47 = OpLoad %int %i
+         %48 = OpAccessChain %_ptr_Function_int %v3i %47
+               OpStore %48 %int_1
+         %50 = OpLoad %int %i
+         %52 = OpAccessChain %_ptr_Function_uint %v4u %50
+               OpStore %52 %uint_1
+         %54 = OpLoad %int %i
+         %56 = OpAccessChain %_ptr_Function_bool %v2b %54
+               OpStore %56 %true
+         %58 = OpLoad %int %i
+         %59 = OpAccessChain %_ptr_Function_float %v2f_2 %58
+               OpStore %59 %float_1
+         %60 = OpLoad %int %i
+         %61 = OpAccessChain %_ptr_Function_int %v3i_2 %60
+               OpStore %61 %int_1
+         %62 = OpLoad %int %i
+         %63 = OpAccessChain %_ptr_Function_uint %v4u_2 %62
+               OpStore %63 %uint_1
+         %64 = OpLoad %int %i
+         %65 = OpAccessChain %_ptr_Function_bool %v2b_2 %64
+               OpStore %65 %true
+               OpBranch %35
+         %35 = OpLabel
+         %66 = OpLoad %int %i
+         %67 = OpIAdd %int %66 %int_1
+               OpStore %i %67
+               OpBranch %33
+         %34 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl

@@ -0,0 +1,31 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v2f_2 : vec2<f32>;
+  var v3i : vec3<i32>;
+  var v3i_2 : vec3<i32>;
+  var v4u : vec4<u32>;
+  var v4u_2 : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v2b_2 : vec2<bool>;
+  {
+    var i : i32 = 0;
+    loop {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0;
+      v3i[i] = 1;
+      v4u[i] = 1u;
+      v2b[i] = true;
+      v2f_2[i] = 1.0;
+      v3i_2[i] = 1;
+      v4u_2[i] = 1u;
+      v2b_2[i] = true;
+
+      continuing {
+        i = (i + 1);
+      }
+    }
+  }
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl

@@ -0,0 +1,32 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v3f : vec3<f32>;
+  var v4f : vec4<f32>;
+  var v2i : vec2<i32>;
+  var v3i : vec3<i32>;
+  var v4i : vec4<i32>;
+  var v2u : vec2<u32>;
+  var v3u : vec3<u32>;
+  var v4u : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v3b : vec3<bool>;
+  var v4b : vec4<bool>;
+
+  for (var i : i32 = 0; i < 2; i = i + 1) {
+    v2f[i] = 1.0;
+    v2i[i] = 1;
+    v2u[i] = 1u;
+    v2b[i] = true;
+  }
+
+  var i : i32 = 0;
+  v3f[i] = 1.0;
+  v4f[i] = 1.0;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v3b[i] = true;
+  v4b[i] = true;
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl

@@ -0,0 +1,64 @@
+void Set_float2(inout float2 vec, int idx, float val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_int2(inout int2 vec, int idx, int val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_uint2(inout uint2 vec, int idx, uint val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+void Set_bool2(inout bool2 vec, int idx, bool val) {
+  switch(idx) {
+    case 0: vec[0] = val; break;
+    case 1: vec[1] = val; break;
+  }
+}
+[numthreads(1, 1, 1)]
+void main() {
+  float2 v2f = float2(0.0f, 0.0f);
+  float3 v3f = float3(0.0f, 0.0f, 0.0f);
+  float4 v4f = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  int2 v2i = int2(0, 0);
+  int3 v3i = int3(0, 0, 0);
+  int4 v4i = int4(0, 0, 0, 0);
+  uint2 v2u = uint2(0u, 0u);
+  uint3 v3u = uint3(0u, 0u, 0u);
+  uint4 v4u = uint4(0u, 0u, 0u, 0u);
+  bool2 v2b = bool2(false, false);
+  bool3 v3b = bool3(false, false, false);
+  bool4 v4b = bool4(false, false, false, false);
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      Set_float2(v2f, i, 1.0f);
+      Set_int2(v2i, i, 1);
+      Set_uint2(v2u, i, 1u);
+      Set_bool2(v2b, i, true);
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  int i = 0;
+  v3f[i] = 1.0f;
+  v4f[i] = 1.0f;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v3b[i] = true;
+  v4b[i] = true;
+  return;
+}

test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl

@@ -0,0 +1,43 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  float2 v2f = 0.0f;
+  float3 v3f = 0.0f;
+  float4 v4f = 0.0f;
+  int2 v2i = 0;
+  int3 v3i = 0;
+  int4 v4i = 0;
+  uint2 v2u = 0u;
+  uint3 v3u = 0u;
+  uint4 v4u = 0u;
+  bool2 v2b = false;
+  bool3 v3b = false;
+  bool4 v4b = false;
+  {
+    int i = 0;
+    while (true) {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0f;
+      v2i[i] = 1;
+      v2u[i] = 1u;
+      v2b[i] = true;
+      {
+        i = (i + 1);
+      }
+    }
+  }
+  int i = 0;
+  v3f[i] = 1.0f;
+  v4f[i] = 1.0f;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v3b[i] = true;
+  v4b[i] = true;
+  return;
+}
+

test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm

@@ -0,0 +1,153 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 105
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %v2f "v2f"
+               OpName %v3f "v3f"
+               OpName %v4f "v4f"
+               OpName %v2i "v2i"
+               OpName %v3i "v3i"
+               OpName %v4i "v4i"
+               OpName %v2u "v2u"
+               OpName %v3u "v3u"
+               OpName %v4u "v4u"
+               OpName %v2b "v2b"
+               OpName %v3b "v3b"
+               OpName %v4b "v4b"
+               OpName %i "i"
+               OpName %i_0 "i"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+          %9 = OpConstantNull %v2float
+    %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %13 = OpConstantNull %v3float
+    %v4float = OpTypeVector %float 4
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %17 = OpConstantNull %v4float
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %22 = OpConstantNull %v2int
+      %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %26 = OpConstantNull %v3int
+      %v4int = OpTypeVector %int 4
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %30 = OpConstantNull %v4int
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %35 = OpConstantNull %v2uint
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %39 = OpConstantNull %v3uint
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %43 = OpConstantNull %v4uint
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+%_ptr_Function_v2bool = OpTypePointer Function %v2bool
+         %48 = OpConstantNull %v2bool
+     %v3bool = OpTypeVector %bool 3
+%_ptr_Function_v3bool = OpTypePointer Function %v3bool
+         %52 = OpConstantNull %v3bool
+     %v4bool = OpTypeVector %bool 4
+%_ptr_Function_v4bool = OpTypePointer Function %v4bool
+         %56 = OpConstantNull %v4bool
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %60 = OpConstantNull %int
+      %int_2 = OpConstant %int 2
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_bool = OpTypePointer Function %bool
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+        %v2f = OpVariable %_ptr_Function_v2float Function %9
+        %v3f = OpVariable %_ptr_Function_v3float Function %13
+        %v4f = OpVariable %_ptr_Function_v4float Function %17
+        %v2i = OpVariable %_ptr_Function_v2int Function %22
+        %v3i = OpVariable %_ptr_Function_v3int Function %26
+        %v4i = OpVariable %_ptr_Function_v4int Function %30
+        %v2u = OpVariable %_ptr_Function_v2uint Function %35
+        %v3u = OpVariable %_ptr_Function_v3uint Function %39
+        %v4u = OpVariable %_ptr_Function_v4uint Function %43
+        %v2b = OpVariable %_ptr_Function_v2bool Function %48
+        %v3b = OpVariable %_ptr_Function_v3bool Function %52
+        %v4b = OpVariable %_ptr_Function_v4bool Function %56
+          %i = OpVariable %_ptr_Function_int Function %60
+        %i_0 = OpVariable %_ptr_Function_int Function %60
+               OpStore %i %int_0
+               OpBranch %61
+         %61 = OpLabel
+               OpLoopMerge %62 %63 None
+               OpBranch %64
+         %64 = OpLabel
+         %66 = OpLoad %int %i
+         %68 = OpSLessThan %bool %66 %int_2
+         %65 = OpLogicalNot %bool %68
+               OpSelectionMerge %69 None
+               OpBranchConditional %65 %70 %69
+         %70 = OpLabel
+               OpBranch %62
+         %69 = OpLabel
+         %71 = OpLoad %int %i
+         %73 = OpAccessChain %_ptr_Function_float %v2f %71
+               OpStore %73 %float_1
+         %75 = OpLoad %int %i
+         %76 = OpAccessChain %_ptr_Function_int %v2i %75
+               OpStore %76 %int_1
+         %78 = OpLoad %int %i
+         %80 = OpAccessChain %_ptr_Function_uint %v2u %78
+               OpStore %80 %uint_1
+         %82 = OpLoad %int %i
+         %84 = OpAccessChain %_ptr_Function_bool %v2b %82
+               OpStore %84 %true
+               OpBranch %63
+         %63 = OpLabel
+         %86 = OpLoad %int %i
+         %87 = OpIAdd %int %86 %int_1
+               OpStore %i %87
+               OpBranch %61
+         %62 = OpLabel
+               OpStore %i_0 %int_0
+         %89 = OpLoad %int %i_0
+         %90 = OpAccessChain %_ptr_Function_float %v3f %89
+               OpStore %90 %float_1
+         %91 = OpLoad %int %i_0
+         %92 = OpAccessChain %_ptr_Function_float %v4f %91
+               OpStore %92 %float_1
+         %93 = OpLoad %int %i_0
+         %94 = OpAccessChain %_ptr_Function_int %v3i %93
+               OpStore %94 %int_1
+         %95 = OpLoad %int %i_0
+         %96 = OpAccessChain %_ptr_Function_int %v4i %95
+               OpStore %96 %int_1
+         %97 = OpLoad %int %i_0
+         %98 = OpAccessChain %_ptr_Function_uint %v3u %97
+               OpStore %98 %uint_1
+         %99 = OpLoad %int %i_0
+        %100 = OpAccessChain %_ptr_Function_uint %v4u %99
+               OpStore %100 %uint_1
+        %101 = OpLoad %int %i_0
+        %102 = OpAccessChain %_ptr_Function_bool %v3b %101
+               OpStore %102 %true
+        %103 = OpLoad %int %i_0
+        %104 = OpAccessChain %_ptr_Function_bool %v4b %103
+               OpStore %104 %true
+               OpReturn
+               OpFunctionEnd

test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl

@@ -0,0 +1,40 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v3f : vec3<f32>;
+  var v4f : vec4<f32>;
+  var v2i : vec2<i32>;
+  var v3i : vec3<i32>;
+  var v4i : vec4<i32>;
+  var v2u : vec2<u32>;
+  var v3u : vec3<u32>;
+  var v4u : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v3b : vec3<bool>;
+  var v4b : vec4<bool>;
+  {
+    var i : i32 = 0;
+    loop {
+      if (!((i < 2))) {
+        break;
+      }
+      v2f[i] = 1.0;
+      v2i[i] = 1;
+      v2u[i] = 1u;
+      v2b[i] = true;
+
+      continuing {
+        i = (i + 1);
+      }
+    }
+  }
+  var i : i32 = 0;
+  v3f[i] = 1.0;
+  v4f[i] = 1.0;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v3b[i] = true;
+  v4b[i] = true;
+}

test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl

@@ -0,0 +1,29 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v3f : vec3<f32>;
+  var v4f : vec4<f32>;
+  var v2i : vec2<i32>;
+  var v3i : vec3<i32>;
+  var v4i : vec4<i32>;
+  var v2u : vec2<u32>;
+  var v3u : vec3<u32>;
+  var v4u : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v3b : vec3<bool>;
+  var v4b : vec4<bool>;
+
+  var i : i32 = 0;
+  v2f[i] = 1.0;
+  v3f[i] = 1.0;
+  v4f[i] = 1.0;
+  v2i[i] = 1;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v2u[i] = 1u;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v2b[i] = true;
+  v3b[i] = true;
+  v4b[i] = true;
+}

test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl

@@ -0,0 +1,29 @@
+[numthreads(1, 1, 1)]
+void main() {
+  float2 v2f = float2(0.0f, 0.0f);
+  float3 v3f = float3(0.0f, 0.0f, 0.0f);
+  float4 v4f = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  int2 v2i = int2(0, 0);
+  int3 v3i = int3(0, 0, 0);
+  int4 v4i = int4(0, 0, 0, 0);
+  uint2 v2u = uint2(0u, 0u);
+  uint3 v3u = uint3(0u, 0u, 0u);
+  uint4 v4u = uint4(0u, 0u, 0u, 0u);
+  bool2 v2b = bool2(false, false);
+  bool3 v3b = bool3(false, false, false);
+  bool4 v4b = bool4(false, false, false, false);
+  int i = 0;
+  v2f[i] = 1.0f;
+  v3f[i] = 1.0f;
+  v4f[i] = 1.0f;
+  v2i[i] = 1;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v2u[i] = 1u;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v2b[i] = true;
+  v3b[i] = true;
+  v4b[i] = true;
+  return;
+}

test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.msl

@@ -0,0 +1,32 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  float2 v2f = 0.0f;
+  float3 v3f = 0.0f;
+  float4 v4f = 0.0f;
+  int2 v2i = 0;
+  int3 v3i = 0;
+  int4 v4i = 0;
+  uint2 v2u = 0u;
+  uint3 v3u = 0u;
+  uint4 v4u = 0u;
+  bool2 v2b = false;
+  bool3 v3b = false;
+  bool4 v4b = false;
+  int i = 0;
+  v2f[i] = 1.0f;
+  v3f[i] = 1.0f;
+  v4f[i] = 1.0f;
+  v2i[i] = 1;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v2u[i] = 1u;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v2b[i] = true;
+  v3b[i] = true;
+  v4b[i] = true;
+  return;
+}
+

test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm

@@ -0,0 +1,129 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 92
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %v2f "v2f"
+               OpName %v3f "v3f"
+               OpName %v4f "v4f"
+               OpName %v2i "v2i"
+               OpName %v3i "v3i"
+               OpName %v4i "v4i"
+               OpName %v2u "v2u"
+               OpName %v3u "v3u"
+               OpName %v4u "v4u"
+               OpName %v2b "v2b"
+               OpName %v3b "v3b"
+               OpName %v4b "v4b"
+               OpName %i "i"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+          %9 = OpConstantNull %v2float
+    %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %13 = OpConstantNull %v3float
+    %v4float = OpTypeVector %float 4
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %17 = OpConstantNull %v4float
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %22 = OpConstantNull %v2int
+      %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %26 = OpConstantNull %v3int
+      %v4int = OpTypeVector %int 4
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %30 = OpConstantNull %v4int
+       %uint = OpTypeInt 32 0
+     %v2uint = OpTypeVector %uint 2
+%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %35 = OpConstantNull %v2uint
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %39 = OpConstantNull %v3uint
+     %v4uint = OpTypeVector %uint 4
+%_ptr_Function_v4uint = OpTypePointer Function %v4uint
+         %43 = OpConstantNull %v4uint
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+%_ptr_Function_v2bool = OpTypePointer Function %v2bool
+         %48 = OpConstantNull %v2bool
+     %v3bool = OpTypeVector %bool 3
+%_ptr_Function_v3bool = OpTypePointer Function %v3bool
+         %52 = OpConstantNull %v3bool
+     %v4bool = OpTypeVector %bool 4
+%_ptr_Function_v4bool = OpTypePointer Function %v4bool
+         %56 = OpConstantNull %v4bool
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %60 = OpConstantNull %int
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_bool = OpTypePointer Function %bool
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+        %v2f = OpVariable %_ptr_Function_v2float Function %9
+        %v3f = OpVariable %_ptr_Function_v3float Function %13
+        %v4f = OpVariable %_ptr_Function_v4float Function %17
+        %v2i = OpVariable %_ptr_Function_v2int Function %22
+        %v3i = OpVariable %_ptr_Function_v3int Function %26
+        %v4i = OpVariable %_ptr_Function_v4int Function %30
+        %v2u = OpVariable %_ptr_Function_v2uint Function %35
+        %v3u = OpVariable %_ptr_Function_v3uint Function %39
+        %v4u = OpVariable %_ptr_Function_v4uint Function %43
+        %v2b = OpVariable %_ptr_Function_v2bool Function %48
+        %v3b = OpVariable %_ptr_Function_v3bool Function %52
+        %v4b = OpVariable %_ptr_Function_v4bool Function %56
+          %i = OpVariable %_ptr_Function_int Function %60
+               OpStore %i %int_0
+         %61 = OpLoad %int %i
+         %63 = OpAccessChain %_ptr_Function_float %v2f %61
+               OpStore %63 %float_1
+         %65 = OpLoad %int %i
+         %66 = OpAccessChain %_ptr_Function_float %v3f %65
+               OpStore %66 %float_1
+         %67 = OpLoad %int %i
+         %68 = OpAccessChain %_ptr_Function_float %v4f %67
+               OpStore %68 %float_1
+         %69 = OpLoad %int %i
+         %70 = OpAccessChain %_ptr_Function_int %v2i %69
+               OpStore %70 %int_1
+         %72 = OpLoad %int %i
+         %73 = OpAccessChain %_ptr_Function_int %v3i %72
+               OpStore %73 %int_1
+         %74 = OpLoad %int %i
+         %75 = OpAccessChain %_ptr_Function_int %v4i %74
+               OpStore %75 %int_1
+         %76 = OpLoad %int %i
+         %78 = OpAccessChain %_ptr_Function_uint %v2u %76
+               OpStore %78 %uint_1
+         %80 = OpLoad %int %i
+         %81 = OpAccessChain %_ptr_Function_uint %v3u %80
+               OpStore %81 %uint_1
+         %82 = OpLoad %int %i
+         %83 = OpAccessChain %_ptr_Function_uint %v4u %82
+               OpStore %83 %uint_1
+         %84 = OpLoad %int %i
+         %86 = OpAccessChain %_ptr_Function_bool %v2b %84
+               OpStore %86 %true
+         %88 = OpLoad %int %i
+         %89 = OpAccessChain %_ptr_Function_bool %v3b %88
+               OpStore %89 %true
+         %90 = OpLoad %int %i
+         %91 = OpAccessChain %_ptr_Function_bool %v4b %90
+               OpStore %91 %true
+               OpReturn
+               OpFunctionEnd

test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl

@@ -0,0 +1,28 @@
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+  var v2f : vec2<f32>;
+  var v3f : vec3<f32>;
+  var v4f : vec4<f32>;
+  var v2i : vec2<i32>;
+  var v3i : vec3<i32>;
+  var v4i : vec4<i32>;
+  var v2u : vec2<u32>;
+  var v3u : vec3<u32>;
+  var v4u : vec4<u32>;
+  var v2b : vec2<bool>;
+  var v3b : vec3<bool>;
+  var v4b : vec4<bool>;
+  var i : i32 = 0;
+  v2f[i] = 1.0;
+  v3f[i] = 1.0;
+  v4f[i] = 1.0;
+  v2i[i] = 1;
+  v3i[i] = 1;
+  v4i[i] = 1;
+  v2u[i] = 1u;
+  v3u[i] = 1u;
+  v4u[i] = 1u;
+  v2b[i] = true;
+  v3b[i] = true;
+  v4b[i] = true;
+}