Add tint e2e tests for vector assignment via dynamic index
We have similar tests for matrices. Note the ir.fxc.hlsl failures. A
workaround similar to what we have on the AST path will be implemented
to address these.
Bug: 368355095
Change-Id: I489b8d089b1e497521f95f2d2de0a5632eb2a010
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/209754
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl
new file mode 100644
index 0000000..a15201f
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl
@@ -0,0 +1,7 @@
+@group(0) @binding(0) var<uniform> i : u32;
+
+@compute @workgroup_size(1)
+fn main() {
+ var v1 : vec3<f32>;
+ v1[i] = 1.0;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.dxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..bd0c260
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.dxc.hlsl
@@ -0,0 +1,14 @@
+void set_vector_element(inout float3 vec, int idx, float val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ float3 v1 = float3(0.0f, 0.0f, 0.0f);
+ set_vector_element(v1, i[0].x, 1.0f);
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.fxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..bd0c260
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.fxc.hlsl
@@ -0,0 +1,14 @@
+void set_vector_element(inout float3 vec, int idx, float val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+ float3 v1 = float3(0.0f, 0.0f, 0.0f);
+ set_vector_element(v1, i[0].x, 1.0f);
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.glsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.glsl
new file mode 100644
index 0000000..58d239d
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.glsl
@@ -0,0 +1,16 @@
+#version 310 es
+
+layout(binding = 0, std140) uniform i_block_ubo {
+ uint inner;
+} i;
+
+void tint_symbol() {
+ vec3 v1 = vec3(0.0f, 0.0f, 0.0f);
+ v1[i.inner] = 1.0f;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..24851e9
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,10 @@
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+[numthreads(1, 1, 1)]
+void main() {
+ float3 v1 = (0.0f).xxx;
+ v1[i[0u].x] = 1.0f;
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..a0e620d
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,17 @@
+SKIP: FAILED
+
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+[numthreads(1, 1, 1)]
+void main() {
+ float3 v1 = (0.0f).xxx;
+ v1[i[0u].x] = 1.0f;
+}
+
+FXC validation failure:
+<scrubbed_path>(8,3-13): error X3500: array reference cannot be used as an l-value; not natively addressable
+
+
+tint executable returned error: exit status 1
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.glsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.glsl
new file mode 100644
index 0000000..62befd3
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.glsl
@@ -0,0 +1,11 @@
+#version 310 es
+
+layout(binding = 0, std140)
+uniform tint_symbol_2_1_ubo {
+ uint tint_symbol_1;
+} v;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ vec3 v1 = vec3(0.0f);
+ v1[v.tint_symbol_1] = 1.0f;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.msl
new file mode 100644
index 0000000..164ff26
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.ir.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ const constant uint* i;
+};
+
+kernel void tint_symbol(const constant uint* i [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=i};
+ float3 v1 = 0.0f;
+ v1[(*tint_module_vars.i)] = 1.0f;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.msl
new file mode 100644
index 0000000..c4924dc
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant uint* tint_symbol_1 [[buffer(0)]]) {
+ float3 v1 = 0.0f;
+ v1[*(tint_symbol_1)] = 1.0f;
+ return;
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.spvasm b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.spvasm
new file mode 100644
index 0000000..30afeb6
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.spvasm
@@ -0,0 +1,41 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 21
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %main "main"
+ OpName %v1 "v1"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 NonWritable
+ %uint = OpTypeInt 32 0
+%tint_symbol_1 = OpTypeStruct %uint
+%_ptr_Uniform_tint_symbol_1 = OpTypePointer Uniform %tint_symbol_1
+ %1 = OpVariable %_ptr_Uniform_tint_symbol_1 Uniform
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %13 = OpConstantNull %v3float
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_float = OpTypePointer Function %float
+ %float_1 = OpConstant %float 1
+ %main = OpFunction %void None %7
+ %8 = OpLabel
+ %v1 = OpVariable %_ptr_Function_v3float Function %13
+ %14 = OpAccessChain %_ptr_Uniform_uint %1 %uint_0
+ %17 = OpLoad %uint %14 None
+ %18 = OpAccessChain %_ptr_Function_float %v1 %17
+ OpStore %18 %float_1 None
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.wgsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.wgsl
new file mode 100644
index 0000000..a15201f
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/function_var.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+@group(0) @binding(0) var<uniform> i : u32;
+
+@compute @workgroup_size(1)
+fn main() {
+ var v1 : vec3<f32>;
+ v1[i] = 1.0;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl
new file mode 100644
index 0000000..6b61eeb
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl
@@ -0,0 +1,7 @@
+@group(0) @binding(0) var<uniform> i : u32;
+var<private> v1 : vec3<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ v1[i] = 1.0;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.dxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ff5cf5b
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.dxc.hlsl
@@ -0,0 +1,14 @@
+void set_vector_element(inout float3 vec, int idx, float val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+static float3 v1 = float3(0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+ set_vector_element(v1, i[0].x, 1.0f);
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.fxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ff5cf5b
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.fxc.hlsl
@@ -0,0 +1,14 @@
+void set_vector_element(inout float3 vec, int idx, float val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+static float3 v1 = float3(0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+ set_vector_element(v1, i[0].x, 1.0f);
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.glsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.glsl
new file mode 100644
index 0000000..c63e627
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.glsl
@@ -0,0 +1,16 @@
+#version 310 es
+
+layout(binding = 0, std140) uniform i_block_ubo {
+ uint inner;
+} i;
+
+vec3 v1 = vec3(0.0f, 0.0f, 0.0f);
+void tint_symbol() {
+ v1[i.inner] = 1.0f;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..bb0d174
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,10 @@
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+static float3 v1 = (0.0f).xxx;
+[numthreads(1, 1, 1)]
+void main() {
+ v1[i[0u].x] = 1.0f;
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4545f7c
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,17 @@
+SKIP: FAILED
+
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+static float3 v1 = (0.0f).xxx;
+[numthreads(1, 1, 1)]
+void main() {
+ v1[i[0u].x] = 1.0f;
+}
+
+FXC validation failure:
+<scrubbed_path>(8,3-13): error X3500: array reference cannot be used as an l-value; not natively addressable
+
+
+tint executable returned error: exit status 1
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.glsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.glsl
new file mode 100644
index 0000000..a43e8a3
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.glsl
@@ -0,0 +1,11 @@
+#version 310 es
+
+layout(binding = 0, std140)
+uniform tint_symbol_2_1_ubo {
+ uint tint_symbol_1;
+} v;
+vec3 v1 = vec3(0.0f);
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v1[v.tint_symbol_1] = 1.0f;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.msl
new file mode 100644
index 0000000..16c0883
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.ir.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ const constant uint* i;
+ thread float3* v1;
+};
+
+kernel void tint_symbol(const constant uint* i [[buffer(0)]]) {
+ thread float3 v1 = 0.0f;
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=i, .v1=(&v1)};
+ (*tint_module_vars.v1)[(*tint_module_vars.i)] = 1.0f;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.msl
new file mode 100644
index 0000000..2165cfa
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_private_vars_struct {
+ float3 v1;
+};
+
+kernel void tint_symbol(const constant uint* tint_symbol_1 [[buffer(0)]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ tint_private_vars.v1[*(tint_symbol_1)] = 1.0f;
+ return;
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.spvasm b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.spvasm
new file mode 100644
index 0000000..933b33f
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.spvasm
@@ -0,0 +1,41 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 21
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %v1 "v1"
+ OpName %main "main"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 NonWritable
+ %uint = OpTypeInt 32 0
+%tint_symbol_1 = OpTypeStruct %uint
+%_ptr_Uniform_tint_symbol_1 = OpTypePointer Uniform %tint_symbol_1
+ %1 = OpVariable %_ptr_Uniform_tint_symbol_1 Uniform
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%_ptr_Private_v3float = OpTypePointer Private %v3float
+ %9 = OpConstantNull %v3float
+ %v1 = OpVariable %_ptr_Private_v3float Private %9
+ %void = OpTypeVoid
+ %12 = OpTypeFunction %void
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Private_float = OpTypePointer Private %float
+ %float_1 = OpConstant %float 1
+ %main = OpFunction %void None %12
+ %13 = OpLabel
+ %14 = OpAccessChain %_ptr_Uniform_uint %1 %uint_0
+ %17 = OpLoad %uint %14 None
+ %18 = OpAccessChain %_ptr_Private_float %v1 %17
+ OpStore %18 %float_1 None
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.wgsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.wgsl
new file mode 100644
index 0000000..bb0cc8b
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/private_var.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<uniform> i : u32;
+
+var<private> v1 : vec3<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ v1[i] = 1.0;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl
new file mode 100644
index 0000000..8d5931f
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl
@@ -0,0 +1,7 @@
+@group(0) @binding(0) var<uniform> i : u32;
+@group(0) @binding(1) var<storage, read_write> v1 : vec3<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ v1[i] = 1.0;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.dxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..4cd1c3f
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.dxc.hlsl
@@ -0,0 +1,10 @@
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+RWByteAddressBuffer v1 : register(u1);
+
+[numthreads(1, 1, 1)]
+void main() {
+ v1.Store((4u * i[0].x), asuint(1.0f));
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.fxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..4cd1c3f
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.fxc.hlsl
@@ -0,0 +1,10 @@
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+RWByteAddressBuffer v1 : register(u1);
+
+[numthreads(1, 1, 1)]
+void main() {
+ v1.Store((4u * i[0].x), asuint(1.0f));
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.glsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.glsl
new file mode 100644
index 0000000..a448515
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.glsl
@@ -0,0 +1,19 @@
+#version 310 es
+
+layout(binding = 0, std140) uniform i_block_ubo {
+ uint inner;
+} i;
+
+layout(binding = 1, std430) buffer v1_block_ssbo {
+ vec3 inner;
+} v1;
+
+void tint_symbol() {
+ v1.inner[i.inner] = 1.0f;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..e75cb75
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,11 @@
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+RWByteAddressBuffer v1 : register(u1);
+[numthreads(1, 1, 1)]
+void main() {
+ uint v = (uint(i[0u].x) * 4u);
+ v1.Store((0u + v), asuint(1.0f));
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..e75cb75
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,11 @@
+
+cbuffer cbuffer_i : register(b0) {
+ uint4 i[1];
+};
+RWByteAddressBuffer v1 : register(u1);
+[numthreads(1, 1, 1)]
+void main() {
+ uint v = (uint(i[0u].x) * 4u);
+ v1.Store((0u + v), asuint(1.0f));
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.glsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.glsl
new file mode 100644
index 0000000..3496e29
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.glsl
@@ -0,0 +1,14 @@
+#version 310 es
+
+layout(binding = 0, std140)
+uniform tint_symbol_2_1_ubo {
+ uint tint_symbol_1;
+} v;
+layout(binding = 1, std430)
+buffer tint_symbol_4_1_ssbo {
+ vec3 tint_symbol_3;
+} v_1;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ v_1.tint_symbol_3[v.tint_symbol_1] = 1.0f;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.msl
new file mode 100644
index 0000000..771c8d1
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.ir.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ const constant uint* i;
+ device packed_float3* v1;
+};
+
+kernel void tint_symbol(const constant uint* i [[buffer(0)]], device packed_float3* v1 [[buffer(1)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=i, .v1=v1};
+ (*tint_module_vars.v1)[(*tint_module_vars.i)] = 1.0f;
+}
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl
new file mode 100644
index 0000000..a6f8f48
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device packed_float3* tint_symbol_1 [[buffer(1)]], const constant uint* tint_symbol_2 [[buffer(0)]]) {
+ (*(tint_symbol_1))[*(tint_symbol_2)] = 1.0f;
+ return;
+}
+
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.spvasm b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.spvasm
new file mode 100644
index 0000000..0f29c05
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.spvasm
@@ -0,0 +1,49 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 23
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpMemberName %tint_symbol_3 0 "tint_symbol_2"
+ OpName %tint_symbol_3 "tint_symbol_3"
+ OpName %main "main"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 NonWritable
+ OpMemberDecorate %tint_symbol_3 0 Offset 0
+ OpDecorate %tint_symbol_3 Block
+ OpDecorate %5 DescriptorSet 0
+ OpDecorate %5 Binding 1
+ OpDecorate %5 Coherent
+ %uint = OpTypeInt 32 0
+%tint_symbol_1 = OpTypeStruct %uint
+%_ptr_Uniform_tint_symbol_1 = OpTypePointer Uniform %tint_symbol_1
+ %1 = OpVariable %_ptr_Uniform_tint_symbol_1 Uniform
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%tint_symbol_3 = OpTypeStruct %v3float
+%_ptr_StorageBuffer_tint_symbol_3 = OpTypePointer StorageBuffer %tint_symbol_3
+ %5 = OpVariable %_ptr_StorageBuffer_tint_symbol_3 StorageBuffer
+ %void = OpTypeVoid
+ %12 = OpTypeFunction %void
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %float_1 = OpConstant %float 1
+ %main = OpFunction %void None %12
+ %13 = OpLabel
+ %14 = OpAccessChain %_ptr_Uniform_uint %1 %uint_0
+ %17 = OpLoad %uint %14 None
+ %18 = OpAccessChain %_ptr_StorageBuffer_v3float %5 %uint_0
+ %20 = OpAccessChain %_ptr_StorageBuffer_float %18 %17
+ OpStore %20 %float_1 None
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.wgsl b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.wgsl
new file mode 100644
index 0000000..0362982
--- /dev/null
+++ b/test/tint/bug/fxc/vector_assignment_dynamic_index/storage_var.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+@group(0) @binding(0) var<uniform> i : u32;
+
+@group(0) @binding(1) var<storage, read_write> v1 : vec3<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ v1[i] = 1.0;
+}
