HLSL: work around FXC failures when dynamically indexing matrices

This fixes errors like "error X3500: array reference cannot be used as
an l-value; not natively addressable". Note that FXC treats matrices
like arrays. We still get this error for dynamically indexed arrays in
structs.

Also improved HLSL assign tests, and add missing ones for vector
indexing.

Manually removed 20 e2e skip hlsl SKIP files that are now passing with
this change.

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

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl

@@ -0,0 +1,12 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var m1 : mat2x4<f32>;
+  m1[uniforms.i][0] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) {
+  switch (col) {
+    case 0:
+      mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0];
+      break;
+    case 1:
+      mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1];
+      break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  set_scalar_float2x4(m1, 0, uniforms[0].x, 1.0f);
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  float2x4 m1 = float2x4(0.0f);
+  m1[(*(tint_symbol_1)).i][0] = 1.0f;
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.spvasm

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %m1 "m1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float
+         %14 = OpConstantNull %mat2v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %m1 = OpVariable %_ptr_Function_mat2v4float Function %14
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %22 = OpAccessChain %_ptr_Function_float %m1 %18 %int_0
+               OpStore %22 %float_1
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_x.wgsl.expected.wgsl

@@ -0,0 +1,13 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var m1 : mat2x4<f32>;
+  m1[uniforms.i][0] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl

@@ -0,0 +1,12 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var m1 : mat2x4<f32>;
+  m1[uniforms.i][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) {
+  switch (col) {
+    case 0:
+      mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0];
+      break;
+    case 1:
+      mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1];
+      break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  set_scalar_float2x4(m1, uniforms[0].y, uniforms[0].x, 1.0f);
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  float2x4 m1 = float2x4(0.0f);
+  m1[(*(tint_symbol_1)).i][(*(tint_symbol_1)).j] = 1.0f;
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.spvasm

@@ -0,0 +1,48 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %m1 "m1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float
+         %14 = OpConstantNull %mat2v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %m1 = OpVariable %_ptr_Function_mat2v4float Function %14
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %21 = OpLoad %uint %20
+         %23 = OpAccessChain %_ptr_Function_float %m1 %18 %21
+               OpStore %23 %float_1
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_xy.wgsl.expected.wgsl

@@ -0,0 +1,13 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var m1 : mat2x4<f32>;
+  m1[uniforms.i][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl

@@ -0,0 +1,13 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[0][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) {
+  switch (col) {
+    case 0:
+      mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0];
+      break;
+    case 1:
+      mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1];
+      break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  set_scalar_float2x4(m1, uniforms[0].y, 0, 1.0f);
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  thread float2x4 tint_symbol_1 = float2x4(0.0f);
+  tint_symbol_1[0][(*(tint_symbol_2)).j] = 1.0f;
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.spvasm

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %m1 "m1"
+               OpName %main "main"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float
+         %10 = OpConstantNull %mat2v4float
+         %m1 = OpVariable %_ptr_Private_mat2v4float Private %10
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Private_float = OpTypePointer Private %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %19 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %20 = OpLoad %uint %19
+         %22 = OpAccessChain %_ptr_Private_float %m1 %int_0 %20
+               OpStore %22 %float_1
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_scalar_y.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[0][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl

@@ -0,0 +1,12 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var m1 : mat2x4<f32>;
+  m1[uniforms.i] = vec4<f32>(1.0);
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.hlsl

@@ -0,0 +1,17 @@
+void set_vector_float2x4(inout float2x4 mat, int col, float4 val) {
+  switch (col) {
+    case 0: mat[0] = val; break;
+    case 1: mat[1] = val; break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  set_vector_float2x4(m1, uniforms[0].x, float4((1.0f).xxxx));
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  float2x4 m1 = float2x4(0.0f);
+  m1[(*(tint_symbol_1)).i] = float4(1.0f);
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.spvasm

@@ -0,0 +1,46 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %m1 "m1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Function_mat2v4float = OpTypePointer Function %mat2v4float
+         %14 = OpConstantNull %mat2v4float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+    %float_1 = OpConstant %float 1
+         %22 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %m1 = OpVariable %_ptr_Function_mat2v4float Function %14
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %20 = OpAccessChain %_ptr_Function_v4float %m1 %18
+               OpStore %20 %22
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/local_assign_vector.wgsl.expected.wgsl

@@ -0,0 +1,13 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var m1 : mat2x4<f32>;
+  m1[uniforms.i] = vec4<f32>(1.0);
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl

@@ -0,0 +1,13 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[uniforms.i][0] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) {
+  switch (col) {
+    case 0:
+      mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0];
+      break;
+    case 1:
+      mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1];
+      break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  set_scalar_float2x4(m1, 0, uniforms[0].x, 1.0f);
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  thread float2x4 tint_symbol_1 = float2x4(0.0f);
+  tint_symbol_1[(*(tint_symbol_2)).i][0] = 1.0f;
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.spvasm

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %m1 "m1"
+               OpName %main "main"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float
+         %10 = OpConstantNull %mat2v4float
+         %m1 = OpVariable %_ptr_Private_mat2v4float Private %10
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Private_float = OpTypePointer Private %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %22 = OpAccessChain %_ptr_Private_float %m1 %18 %int_0
+               OpStore %22 %float_1
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_x.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[uniforms.i][0] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl

@@ -0,0 +1,13 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[uniforms.i][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) {
+  switch (col) {
+    case 0:
+      mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0];
+      break;
+    case 1:
+      mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1];
+      break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  set_scalar_float2x4(m1, uniforms[0].y, uniforms[0].x, 1.0f);
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  thread float2x4 tint_symbol_1 = float2x4(0.0f);
+  tint_symbol_1[(*(tint_symbol_2)).i][(*(tint_symbol_2)).j] = 1.0f;
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.spvasm

@@ -0,0 +1,48 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %m1 "m1"
+               OpName %main "main"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float
+         %10 = OpConstantNull %mat2v4float
+         %m1 = OpVariable %_ptr_Private_mat2v4float Private %10
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_Private_float = OpTypePointer Private %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %21 = OpLoad %uint %20
+         %23 = OpAccessChain %_ptr_Private_float %m1 %18 %21
+               OpStore %23 %float_1
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_xy.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[uniforms.i][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl

@@ -0,0 +1,13 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[0][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+void set_scalar_float2x4(inout float2x4 mat, int col, int row, float val) {
+  switch (col) {
+    case 0:
+      mat[0] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[0];
+      break;
+    case 1:
+      mat[1] = (row.xxxx == int4(0, 1, 2, 3)) ? val.xxxx : mat[1];
+      break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  set_scalar_float2x4(m1, uniforms[0].y, 0, 1.0f);
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  thread float2x4 tint_symbol_1 = float2x4(0.0f);
+  tint_symbol_1[0][(*(tint_symbol_2)).j] = 1.0f;
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.spvasm

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %m1 "m1"
+               OpName %main "main"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float
+         %10 = OpConstantNull %mat2v4float
+         %m1 = OpVariable %_ptr_Private_mat2v4float Private %10
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Private_float = OpTypePointer Private %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %19 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %20 = OpLoad %uint %19
+         %22 = OpAccessChain %_ptr_Private_float %m1 %int_0 %20
+               OpStore %22 %float_1
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_scalar_y.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[0][uniforms.j] = 1.0;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl

@@ -0,0 +1,13 @@
+[[block]] struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[uniforms.i] = vec4<f32>(1.0);
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.hlsl

@@ -0,0 +1,17 @@
+void set_vector_float2x4(inout float2x4 mat, int col, float4 val) {
+  switch (col) {
+    case 0: mat[0] = val; break;
+    case 1: mat[1] = val; break;
+  }
+}
+
+cbuffer cbuffer_uniforms : register(b4, space1) {
+  uint4 uniforms[1];
+};
+static float2x4 m1 = float2x4(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  set_vector_float2x4(m1, uniforms[0].x, float4((1.0f).xxxx));
+  return;
+}

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.msl

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  thread float2x4 tint_symbol_1 = float2x4(0.0f);
+  tint_symbol_1[(*(tint_symbol_2)).i] = float4(1.0f);
+  return;
+}
+

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.spvasm

@@ -0,0 +1,46 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %m1 "m1"
+               OpName %main "main"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+%_ptr_Private_mat2v4float = OpTypePointer Private %mat2v4float
+         %10 = OpConstantNull %mat2v4float
+         %m1 = OpVariable %_ptr_Private_mat2v4float Private %10
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+    %float_1 = OpConstant %float 1
+         %22 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %20 = OpAccessChain %_ptr_Private_v4float %m1 %18
+               OpStore %20 %22
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/matrix_assignment_dynamic_index/module_assign_vector.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+[[block]]
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+var<private> m1 : mat2x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  m1[uniforms.i] = vec4<f32>(1.0);
+}