test: Add FXC test cases for indexing arrays in structs

Also move:
  test/fxc_bugs/vector_assignment_in_loop
to
  test/bug/fxc/vector_assignment_in_loop

Change-Id: I7bbfc476fdb7a3296025609625e322fed8d16285
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59444
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>

test/bug/fxc/dyn_array_idx/read/function.wgsl

@@ -0,0 +1,19 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var s : S;
+  result.out = s.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.hlsl

@@ -0,0 +1,16 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void f() {
+  S s = (S)0;
+  result.Store(0u, asuint(s.data[asint(ubo[0].x)]));
+  return;
+}

test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  S s = {};
+  result.out = s.data.arr[ubo.dynamic_idx];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.spvasm

@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %f "f"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Function_S = OpTypePointer Function %S
+         %18 = OpConstantNull %S
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Function_int = OpTypePointer Function %int
+          %f = OpFunction %void None %8
+         %11 = OpLabel
+          %s = OpVariable %_ptr_Function_S Function %18
+         %21 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %24 = OpLoad %int %23
+         %26 = OpAccessChain %_ptr_Function_int %s %uint_0 %24
+         %27 = OpLoad %int %26
+               OpStore %21 %27
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.wgsl

@@ -0,0 +1,23 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var s : S;
+  result.out = s.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/private.wgsl

@@ -0,0 +1,20 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+var<private> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = s.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.hlsl

@@ -0,0 +1,16 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+static S s = (S)0;
+
+[numthreads(1, 1, 1)]
+void f() {
+  result.Store(0u, asuint(s.data[asint(ubo[0].x)]));
+  return;
+}

test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  thread S tint_symbol = {};
+  result.out = tint_symbol.data.arr[ubo.dynamic_idx];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.spvasm

@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Private_S = OpTypePointer Private %S
+         %14 = OpConstantNull %S
+          %s = OpVariable %_ptr_Private_S Private %14
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Private_int = OpTypePointer Private %int
+          %f = OpFunction %void None %15
+         %18 = OpLabel
+         %21 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %24 = OpLoad %int %23
+         %26 = OpAccessChain %_ptr_Private_int %s %uint_0 %24
+         %27 = OpLoad %int %26
+               OpStore %21 %27
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+var<private> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = s.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/storage.wgsl

@@ -0,0 +1,21 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(2)]] var<storage, read_write> result: Result;
+
+[[block]]
+struct SSBO {
+  data: array<i32, 4>;
+};
+[[group(0), binding(1)]] var<storage, read_write> ssbo: SSBO;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = ssbo.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.hlsl

@@ -0,0 +1,13 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+RWByteAddressBuffer result : register(u2, space0);
+
+RWByteAddressBuffer ssbo : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void f() {
+  result.Store(0u, asuint(asint(ssbo.Load((4u * uint(asint(ubo[0].x)))))));
+  return;
+}

test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ int arr[4];
+};
+struct SSBO {
+  /* 0x0000 */ tint_array_wrapper data;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(2)]], device SSBO& ssbo [[buffer(1)]]) {
+  result.out = ssbo.data.arr[ubo.dynamic_idx];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.spvasm

@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "data"
+               OpName %ssbo "ssbo"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 2
+               OpDecorate %SSBO Block
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpDecorate %_arr_int_uint_4 ArrayStride 4
+               OpDecorate %ssbo DescriptorSet 0
+               OpDecorate %ssbo Binding 1
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+       %SSBO = OpTypeStruct %_arr_int_uint_4
+%_ptr_StorageBuffer_SSBO = OpTypePointer StorageBuffer %SSBO
+       %ssbo = OpVariable %_ptr_StorageBuffer_SSBO StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+          %f = OpFunction %void None %14
+         %17 = OpLabel
+         %20 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %22 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %23 = OpLoad %int %22
+         %24 = OpAccessChain %_ptr_StorageBuffer_int %ssbo %uint_0 %23
+         %25 = OpLoad %int %24
+               OpStore %20 %25
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.wgsl

@@ -0,0 +1,25 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(2)]] var<storage, read_write> result : Result;
+
+[[block]]
+struct SSBO {
+  data : array<i32, 4>;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> ssbo : SSBO;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = ssbo.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/uniform.wgsl

@@ -0,0 +1,16 @@
+[[block]]
+struct UBO {
+  data: [[stride(16)]] array<i32, 4>;
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(2)]] var<storage, read_write> result: Result;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = ubo.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.hlsl

@@ -0,0 +1,12 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[5];
+};
+
+RWByteAddressBuffer result : register(u2, space0);
+
+[numthreads(1, 1, 1)]
+void f() {
+  const uint scalar_offset = ((16u * uint(asint(ubo[4].x)))) / 4;
+  result.Store(0u, asuint(asint(ubo[scalar_offset / 4][scalar_offset % 4])));
+  return;
+}

test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_padded_array_element {
+  /* 0x0000 */ int el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct UBO {
+  /* 0x0000 */ tint_array_wrapper data;
+  /* 0x0040 */ int dynamic_idx;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(2)]]) {
+  result.out = ubo.data.arr[ubo.dynamic_idx].el;
+  return;
+}
+

test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm

@@ -0,0 +1,54 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "data"
+               OpMemberName %UBO 1 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %_arr_int_uint_4 ArrayStride 16
+               OpMemberDecorate %UBO 1 Offset 64
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 2
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+        %UBO = OpTypeStruct %_arr_int_uint_4 %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+          %f = OpFunction %void None %11
+         %14 = OpLabel
+         %17 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %20 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1
+         %21 = OpLoad %int %20
+         %22 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %21
+         %23 = OpLoad %int %22
+               OpStore %17 %23
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+[[block]]
+struct UBO {
+  data : [[stride(16)]] array<i32, 4>;
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(2)]] var<storage, read_write> result : Result;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = ubo.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl

@@ -0,0 +1,20 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+var<workgroup> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = s.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.hlsl

@@ -0,0 +1,29 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+groupshared S s;
+
+struct tint_symbol_2 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void f(tint_symbol_2 tint_symbol_1) {
+  const uint local_invocation_index = tint_symbol_1.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    {
+      for(int i = 0; (i < 64); i = (i + 1)) {
+        s.data[i] = 0;
+      }
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  result.Store(0u, asuint(s.data[asint(ubo[0].x)]));
+  return;
+}

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  threadgroup S tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    S const tint_symbol_1 = {};
+    tint_symbol_2 = tint_symbol_1;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  result.out = tint_symbol_2.data.arr[ubo.dynamic_idx];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.spvasm

@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f" %tint_symbol
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpName %tint_symbol "tint_symbol"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+          %s = OpVariable %_ptr_Workgroup_S Workgroup
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %26 = OpConstantNull %S
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+          %f = OpFunction %void None %16
+         %19 = OpLabel
+         %20 = OpLoad %uint %tint_symbol
+         %22 = OpIEqual %bool %20 %uint_0
+               OpSelectionMerge %24 None
+               OpBranchConditional %22 %25 %24
+         %25 = OpLabel
+               OpStore %s %26
+               OpBranch %24
+         %24 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %31 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %33 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %34 = OpLoad %int %33
+         %36 = OpAccessChain %_ptr_Workgroup_int %s %uint_0 %34
+         %37 = OpLoad %int %36
+               OpStore %31 %37
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+var<workgroup> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  result.out = s.data[ubo.dynamic_idx];
+}

test/bug/fxc/dyn_array_idx/write/function.wgsl

@@ -0,0 +1,20 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var s : S;
+  s.data[ubo.dynamic_idx] = 1;
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void f() {
+  S s = (S)0;
+  s.data[asint(ubo[0].x)] = 1;
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}
+C:\src\tint\test\Shader@0x000002C587F87250(14,3-25): error X3500: array reference cannot be used as an l-value; not natively addressable
+

test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.msl

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  S s = {};
+  s.data.arr[ubo.dynamic_idx] = 1;
+  result.out = s.data.arr[3];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.spvasm

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %f "f"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Function_S = OpTypePointer Function %S
+         %18 = OpConstantNull %S
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Function_int = OpTypePointer Function %int
+      %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_3 = OpConstant %int 3
+          %f = OpFunction %void None %8
+         %11 = OpLabel
+          %s = OpVariable %_ptr_Function_S Function %18
+         %21 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %22 = OpLoad %int %21
+         %24 = OpAccessChain %_ptr_Function_int %s %uint_0 %22
+               OpStore %24 %int_1
+         %27 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %29 = OpAccessChain %_ptr_Function_int %s %uint_0 %int_3
+         %30 = OpLoad %int %29
+               OpStore %27 %30
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var s : S;
+  s.data[ubo.dynamic_idx] = 1;
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl

@@ -0,0 +1,24 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+fn x(p : ptr<function, S>) {
+  (*p).data[ubo.dynamic_idx] = 1;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var s : S;
+  x(&s);
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+SKIP: FAILED
+
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+
+void x(inout S p) {
+  p.data[asint(ubo[0].x)] = 1;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  S s = (S)0;
+  x(s);
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}
+C:\src\tint\test\Shader@0x000002338E919910(12,3-25): error X3500: array reference cannot be used as an l-value; not natively addressable
+

test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+void x(constant UBO& ubo, thread S* const p) {
+  (*(p)).data.arr[ubo.dynamic_idx] = 1;
+}
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  S s = {};
+  x(ubo, &(s));
+  result.out = s.data.arr[3];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm

@@ -0,0 +1,73 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %x "x"
+               OpName %p "p"
+               OpName %f "f"
+               OpName %s "s"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Function_S = OpTypePointer Function %S
+          %8 = OpTypeFunction %void %_ptr_Function_S
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Function_int = OpTypePointer Function %int
+      %int_1 = OpConstant %int 1
+         %26 = OpTypeFunction %void
+         %30 = OpConstantNull %S
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_3 = OpConstant %int 3
+          %x = OpFunction %void None %8
+          %p = OpFunctionParameter %_ptr_Function_S
+         %17 = OpLabel
+         %21 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %22 = OpLoad %int %21
+         %24 = OpAccessChain %_ptr_Function_int %p %uint_0 %22
+               OpStore %24 %int_1
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %26
+         %28 = OpLabel
+          %s = OpVariable %_ptr_Function_S Function %30
+         %31 = OpFunctionCall %void %x %s
+         %34 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %36 = OpAccessChain %_ptr_Function_int %s %uint_0 %int_3
+         %37 = OpLoad %int %36
+               OpStore %34 %37
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.wgsl

@@ -0,0 +1,28 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+fn x(p : ptr<function, S>) {
+  (*(p)).data[ubo.dynamic_idx] = 1;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var s : S;
+  x(&(s));
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/private.wgsl

@@ -0,0 +1,21 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+var<private> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  s.data[ubo.dynamic_idx] = 1;
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+SKIP: FAILED
+
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+static S s = (S)0;
+
+[numthreads(1, 1, 1)]
+void f() {
+  s.data[asint(ubo[0].x)] = 1;
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}
+C:\src\tint\test\Shader@0x000001D94063C630(14,3-25): error X3500: array reference cannot be used as an l-value; not natively addressable
+

test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.msl

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  thread S tint_symbol = {};
+  tint_symbol.data.arr[ubo.dynamic_idx] = 1;
+  result.out = tint_symbol.data.arr[3];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.spvasm

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Private_S = OpTypePointer Private %S
+         %14 = OpConstantNull %S
+          %s = OpVariable %_ptr_Private_S Private %14
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Private_int = OpTypePointer Private %int
+      %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_3 = OpConstant %int 3
+          %f = OpFunction %void None %15
+         %18 = OpLabel
+         %21 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %22 = OpLoad %int %21
+         %24 = OpAccessChain %_ptr_Private_int %s %uint_0 %22
+               OpStore %24 %int_1
+         %27 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %29 = OpAccessChain %_ptr_Private_int %s %uint_0 %int_3
+         %30 = OpLoad %int %29
+               OpStore %27 %30
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.wgsl

@@ -0,0 +1,25 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+var<private> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  s.data[ubo.dynamic_idx] = 1;
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl

@@ -0,0 +1,25 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+var<private> s : S;
+
+fn x(p : ptr<private, S>) {
+  (*p).data[ubo.dynamic_idx] = 1;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  x(&s);
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+SKIP: FAILED
+
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+static S s = (S)0;
+
+void x(inout S p) {
+  p.data[asint(ubo[0].x)] = 1;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  x(s);
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}
+C:\src\tint\test\Shader@0x0000018B80081AA0(13,3-25): error X3500: array reference cannot be used as an l-value; not natively addressable
+

test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+void x(constant UBO& ubo, thread S* const p) {
+  (*(p)).data.arr[ubo.dynamic_idx] = 1;
+}
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  thread S tint_symbol = {};
+  x(ubo, &(tint_symbol));
+  result.out = tint_symbol.data.arr[3];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm

@@ -0,0 +1,73 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpName %x "x"
+               OpName %p "p"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Private_S = OpTypePointer Private %S
+         %14 = OpConstantNull %S
+          %s = OpVariable %_ptr_Private_S Private %14
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void %_ptr_Private_S
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Private_int = OpTypePointer Private %int
+      %int_1 = OpConstant %int 1
+         %28 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_3 = OpConstant %int 3
+          %x = OpFunction %void None %15
+          %p = OpFunctionParameter %_ptr_Private_S
+         %19 = OpLabel
+         %23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %24 = OpLoad %int %23
+         %26 = OpAccessChain %_ptr_Private_int %p %uint_0 %24
+               OpStore %26 %int_1
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %28
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %x %s
+         %34 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %36 = OpAccessChain %_ptr_Private_int %s %uint_0 %int_3
+         %37 = OpLoad %int %36
+               OpStore %34 %37
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.wgsl

@@ -0,0 +1,29 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+var<private> s : S;
+
+fn x(p : ptr<private, S>) {
+  (*(p)).data[ubo.dynamic_idx] = 1;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  x(&(s));
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/storage.wgsl

@@ -0,0 +1,22 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(2)]] var<storage, read_write> result: Result;
+
+[[block]]
+struct SSBO {
+  data: array<i32, 4>;
+};
+[[group(0), binding(1)]] var<storage, read_write> ssbo: SSBO;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  ssbo.data[ubo.dynamic_idx] = 1;
+  result.out = ssbo.data[3];
+}

test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.hlsl

@@ -0,0 +1,14 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+RWByteAddressBuffer result : register(u2, space0);
+
+RWByteAddressBuffer ssbo : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void f() {
+  ssbo.Store((4u * uint(asint(ubo[0].x))), asuint(1));
+  result.Store(0u, asuint(asint(ssbo.Load(12u))));
+  return;
+}

test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl

@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ int arr[4];
+};
+struct SSBO {
+  /* 0x0000 */ tint_array_wrapper data;
+};
+
+kernel void f(constant UBO& ubo [[buffer(0)]], device SSBO& ssbo [[buffer(1)]], device Result& result [[buffer(2)]]) {
+  ssbo.data.arr[ubo.dynamic_idx] = 1;
+  result.out = ssbo.data.arr[3];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.spvasm

@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %SSBO "SSBO"
+               OpMemberName %SSBO 0 "data"
+               OpName %ssbo "ssbo"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 2
+               OpDecorate %SSBO Block
+               OpMemberDecorate %SSBO 0 Offset 0
+               OpDecorate %_arr_int_uint_4 ArrayStride 4
+               OpDecorate %ssbo DescriptorSet 0
+               OpDecorate %ssbo Binding 1
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+       %SSBO = OpTypeStruct %_arr_int_uint_4
+%_ptr_StorageBuffer_SSBO = OpTypePointer StorageBuffer %SSBO
+       %ssbo = OpVariable %_ptr_StorageBuffer_SSBO StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+      %int_3 = OpConstant %int 3
+          %f = OpFunction %void None %14
+         %17 = OpLabel
+         %20 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %21 = OpLoad %int %20
+         %23 = OpAccessChain %_ptr_StorageBuffer_int %ssbo %uint_0 %21
+               OpStore %23 %int_1
+         %25 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %27 = OpAccessChain %_ptr_StorageBuffer_int %ssbo %uint_0 %int_3
+         %28 = OpLoad %int %27
+               OpStore %25 %28
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.wgsl

@@ -0,0 +1,26 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(2)]] var<storage, read_write> result : Result;
+
+[[block]]
+struct SSBO {
+  data : array<i32, 4>;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> ssbo : SSBO;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  ssbo.data[ubo.dynamic_idx] = 1;
+  result.out = ssbo.data[3];
+}

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl

@@ -0,0 +1,21 @@
+[[block]]
+struct UBO {
+  dynamic_idx: i32;
+};
+[[group(0), binding(0)]] var<uniform> ubo: UBO;
+struct S {
+  data: array<i32, 64>;
+};
+[[block]]
+struct Result {
+  out: i32;
+};
+[[group(0), binding(1)]] var<storage, read_write> result: Result;
+
+var<workgroup> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  s.data[ubo.dynamic_idx] = 1;
+  result.out = s.data[3];
+}

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+groupshared S s;
+
+struct tint_symbol_2 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void f(tint_symbol_2 tint_symbol_1) {
+  const uint local_invocation_index = tint_symbol_1.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    {
+      for(int i = 0; (i < 64); i = (i + 1)) {
+        s.data[i] = 0;
+      }
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  s.data[asint(ubo[0].x)] = 1;
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl

@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct UBO {
+  /* 0x0000 */ int dynamic_idx;
+};
+struct tint_array_wrapper {
+  int arr[64];
+};
+struct S {
+  tint_array_wrapper data;
+};
+struct Result {
+  /* 0x0000 */ int out;
+};
+
+kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  threadgroup S tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    S const tint_symbol_1 = {};
+    tint_symbol_2 = tint_symbol_1;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  tint_symbol_2.data.arr[ubo.dynamic_idx] = 1;
+  result.out = tint_symbol_2.data.arr[3];
+  return;
+}
+

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.spvasm

@@ -0,0 +1,80 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 41
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f" %tint_symbol
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "dynamic_idx"
+               OpName %ubo "ubo"
+               OpName %Result "Result"
+               OpMemberName %Result 0 "out"
+               OpName %result "result"
+               OpName %S "S"
+               OpMemberName %S 0 "data"
+               OpName %s "s"
+               OpName %tint_symbol "tint_symbol"
+               OpName %f "f"
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %ubo Binding 0
+               OpDecorate %Result Block
+               OpMemberDecorate %Result 0 Offset 0
+               OpDecorate %result DescriptorSet 0
+               OpDecorate %result Binding 1
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_int_uint_64 ArrayStride 4
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+        %UBO = OpTypeStruct %int
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+     %Result = OpTypeStruct %int
+%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
+     %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
+       %uint = OpTypeInt 32 0
+    %uint_64 = OpConstant %uint 64
+%_arr_int_uint_64 = OpTypeArray %int %uint_64
+          %S = OpTypeStruct %_arr_int_uint_64
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+          %s = OpVariable %_ptr_Workgroup_S Workgroup
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %26 = OpConstantNull %S
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+      %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_3 = OpConstant %int 3
+          %f = OpFunction %void None %16
+         %19 = OpLabel
+         %20 = OpLoad %uint %tint_symbol
+         %22 = OpIEqual %bool %20 %uint_0
+               OpSelectionMerge %24 None
+               OpBranchConditional %22 %25 %24
+         %25 = OpLabel
+               OpStore %s %26
+               OpBranch %24
+         %24 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %31 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0
+         %32 = OpLoad %int %31
+         %34 = OpAccessChain %_ptr_Workgroup_int %s %uint_0 %32
+               OpStore %34 %int_1
+         %37 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %39 = OpAccessChain %_ptr_Workgroup_int %s %uint_0 %int_3
+         %40 = OpLoad %int %39
+               OpStore %37 %40
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.wgsl

@@ -0,0 +1,25 @@
+[[block]]
+struct UBO {
+  dynamic_idx : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> ubo : UBO;
+
+struct S {
+  data : array<i32, 64>;
+};
+
+[[block]]
+struct Result {
+  out : i32;
+};
+
+[[group(0), binding(1)]] var<storage, read_write> result : Result;
+
+var<workgroup> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  s.data[ubo.dynamic_idx] = 1;
+  result.out = s.data[3];
+}

test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl

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

test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl

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

test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl

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

test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl

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

test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl

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

test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl

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

test/bug/tint/998.wgsl.expected.hlsl

@@ -1,3 +1,5 @@
+SKIP: FAILED
+
 cbuffer cbuffer_constants : register(b0, space1) {
   uint4 constants[1];
 };
@@ -15,3 +17,5 @@ void main() {
   s.data[constants[0].x] = 0u;
   return;
 }
+C:\src\tint\test\Shader@0x000002B2827B9810(15,3-24): error X3500: array reference cannot be used as an l-value; not natively addressable
+