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>
diff --git a/test/bug/fxc/dyn_array_idx/read/function.wgsl b/test/bug/fxc/dyn_array_idx/read/function.wgsl
new file mode 100644
index 0000000..b1f7ad0
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.hlsl
new file mode 100644
index 0000000..b9a4617
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl
new file mode 100644
index 0000000..da74b6f
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.spvasm
new file mode 100644
index 0000000..95c6f27
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.wgsl
new file mode 100644
index 0000000..1a66054
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/private.wgsl b/test/bug/fxc/dyn_array_idx/read/private.wgsl
new file mode 100644
index 0000000..731e57b
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.hlsl
new file mode 100644
index 0000000..ef50e8a
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl
new file mode 100644
index 0000000..24958dc
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.spvasm
new file mode 100644
index 0000000..54eb08e
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.wgsl
new file mode 100644
index 0000000..68074bd
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/storage.wgsl b/test/bug/fxc/dyn_array_idx/read/storage.wgsl
new file mode 100644
index 0000000..5237a5b
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.hlsl
new file mode 100644
index 0000000..dada6d4
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl
new file mode 100644
index 0000000..7bf987a
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.spvasm
new file mode 100644
index 0000000..b5c4ce4
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.wgsl
new file mode 100644
index 0000000..7610025
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl
new file mode 100644
index 0000000..6edb1df
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.hlsl
new file mode 100644
index 0000000..01e61f7
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
new file mode 100644
index 0000000..3466444
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm
new file mode 100644
index 0000000..1026040
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl
new file mode 100644
index 0000000..6612630
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl b/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl
new file mode 100644
index 0000000..b598cd3
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.hlsl
new file mode 100644
index 0000000..57a4140
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
new file mode 100644
index 0000000..f71e0a0
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.spvasm
new file mode 100644
index 0000000..09a3136
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.wgsl
new file mode 100644
index 0000000..51c70f7
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/function.wgsl b/test/bug/fxc/dyn_array_idx/write/function.wgsl
new file mode 100644
index 0000000..1fad5bc
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl
new file mode 100644
index 0000000..04d50a1
--- /dev/null
+++ b/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
+
diff --git a/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.msl
new file mode 100644
index 0000000..2130671
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.spvasm
new file mode 100644
index 0000000..46228e1
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.wgsl
new file mode 100644
index 0000000..1871726
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl
new file mode 100644
index 0000000..ab909db
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl
new file mode 100644
index 0000000..59669b3
--- /dev/null
+++ b/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
+
diff --git a/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.msl
new file mode 100644
index 0000000..74f9982
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm
new file mode 100644
index 0000000..446b655
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.wgsl
new file mode 100644
index 0000000..c106f66
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/private.wgsl b/test/bug/fxc/dyn_array_idx/write/private.wgsl
new file mode 100644
index 0000000..4f33168
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl
new file mode 100644
index 0000000..1ef5a38
--- /dev/null
+++ b/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
+
diff --git a/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.msl
new file mode 100644
index 0000000..4e7678b
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.spvasm
new file mode 100644
index 0000000..341b1b8
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.wgsl
new file mode 100644
index 0000000..42f5858
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl
new file mode 100644
index 0000000..6c20074
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl
new file mode 100644
index 0000000..ed009a4
--- /dev/null
+++ b/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
+
diff --git a/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.msl
new file mode 100644
index 0000000..f72d1a2
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm
new file mode 100644
index 0000000..224618f
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.wgsl
new file mode 100644
index 0000000..176c727
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/storage.wgsl b/test/bug/fxc/dyn_array_idx/write/storage.wgsl
new file mode 100644
index 0000000..7bedb67
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.hlsl
new file mode 100644
index 0000000..005f321
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl
new file mode 100644
index 0000000..9e0044f
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.spvasm
new file mode 100644
index 0000000..85915c5
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.wgsl
new file mode 100644
index 0000000..b06072b
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl b/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl
new file mode 100644
index 0000000..26efbe6
--- /dev/null
+++ b/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];
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.hlsl b/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.hlsl
new file mode 100644
index 0000000..0dd523f
--- /dev/null
+++ b/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;
+}
diff --git a/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
new file mode 100644
index 0000000..76981a0
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.spvasm
new file mode 100644
index 0000000..f395a39
--- /dev/null
+++ b/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
diff --git a/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.wgsl
new file mode 100644
index 0000000..4b11871
--- /dev/null
+++ b/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];
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl
similarity index 99%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl
index 8dc7515..a684c56 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl
+++ b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl
@@ -1,21 +1,21 @@
-var<private> v2f : vec2<f32>;
-var<private> v3i : vec3<i32>;
-var<private> v4u : vec4<u32>;
-var<private> v2b : vec2<bool>;
-
-fn foo() {
- for (var i : i32 = 0; i < 2; i = i + 1) {
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
- }
-}
-
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
-
- for (var i : i32 = 0; i < 2; i = i + 1) {
- foo();
- }
-}
+var<private> v2f : vec2<f32>;
+var<private> v3i : vec3<i32>;
+var<private> v4u : vec4<u32>;
+var<private> v2b : vec2<bool>;
+
+fn foo() {
+ for (var i : i32 = 0; i < 2; i = i + 1) {
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ }
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+
+ for (var i : i32 = 0; i < 2; i = i + 1) {
+ foo();
+ }
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.spvasm
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl
similarity index 99%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl
index 418155d..c85e301 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl
+++ b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl
@@ -1,20 +1,20 @@
-var<private> v2f : vec2<f32>;
-var<private> v3i : vec3<i32>;
-var<private> v4u : vec4<u32>;
-var<private> v2b : vec2<bool>;
-
-fn foo() {
- var i : i32 = 0;
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
-}
-
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
-
- for (var i : i32 = 0; i < 2; i = i + 1) {
- foo();
- }
-}
+var<private> v2f : vec2<f32>;
+var<private> v3i : vec3<i32>;
+var<private> v4u : vec4<u32>;
+var<private> v2b : vec2<bool>;
+
+fn foo() {
+ var i : i32 = 0;
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+
+ for (var i : i32 = 0; i < 2; i = i + 1) {
+ foo();
+ }
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.spvasm
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl
similarity index 99%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl
index ef2ec78..5d42838 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl
+++ b/test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl
@@ -1,30 +1,30 @@
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
- var v2f : vec2<f32>;
- var v3f : vec3<f32>;
- var v4f : vec4<f32>;
- var v2i : vec2<i32>;
- var v3i : vec3<i32>;
- var v4i : vec4<i32>;
- var v2u : vec2<u32>;
- var v3u : vec3<u32>;
- var v4u : vec4<u32>;
- var v2b : vec2<bool>;
- var v3b : vec3<bool>;
- var v4b : vec4<bool>;
-
- for (var i : i32 = 0; i < 2; i = i + 1) {
- v2f[i] = 1.0;
- v3f[i] = 1.0;
- v4f[i] = 1.0;
- v2i[i] = 1;
- v3i[i] = 1;
- v4i[i] = 1;
- v2u[i] = 1u;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v2b[i] = true;
- v3b[i] = true;
- v4b[i] = true;
- }
-}
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ var v2f : vec2<f32>;
+ var v3f : vec3<f32>;
+ var v4f : vec4<f32>;
+ var v2i : vec2<i32>;
+ var v3i : vec3<i32>;
+ var v4i : vec4<i32>;
+ var v2u : vec2<u32>;
+ var v3u : vec3<u32>;
+ var v4u : vec4<u32>;
+ var v2b : vec2<bool>;
+ var v3b : vec3<bool>;
+ var v4b : vec4<bool>;
+
+ for (var i : i32 = 0; i < 2; i = i + 1) {
+ v2f[i] = 1.0;
+ v3f[i] = 1.0;
+ v4f[i] = 1.0;
+ v2i[i] = 1;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v2u[i] = 1u;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v3b[i] = true;
+ v4b[i] = true;
+ }
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm b/test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.spvasm
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl
similarity index 99%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl
index d464d4c..7b91e48 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl
+++ b/test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl
@@ -1,26 +1,26 @@
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
- var v2f : vec2<f32>;
- var v2f_2 : vec2<f32>;
-
- var v3i : vec3<i32>;
- var v3i_2 : vec3<i32>;
-
- var v4u : vec4<u32>;
- var v4u_2 : vec4<u32>;
-
- var v2b : vec2<bool>;
- var v2b_2 : vec2<bool>;
-
- for (var i : i32 = 0; i < 2; i = i + 1) {
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
-
- v2f_2[i] = 1.0;
- v3i_2[i] = 1;
- v4u_2[i] = 1u;
- v2b_2[i] = true;
- }
-}
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ var v2f : vec2<f32>;
+ var v2f_2 : vec2<f32>;
+
+ var v3i : vec3<i32>;
+ var v3i_2 : vec3<i32>;
+
+ var v4u : vec4<u32>;
+ var v4u_2 : vec4<u32>;
+
+ var v2b : vec2<bool>;
+ var v2b_2 : vec2<bool>;
+
+ for (var i : i32 = 0; i < 2; i = i + 1) {
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+
+ v2f_2[i] = 1.0;
+ v3i_2[i] = 1;
+ v4u_2[i] = 1u;
+ v2b_2[i] = true;
+ }
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm b/test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.spvasm
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl
similarity index 99%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl
index d07e08d..d255bcc 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl
+++ b/test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl
@@ -1,32 +1,32 @@
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
- var v2f : vec2<f32>;
- var v3f : vec3<f32>;
- var v4f : vec4<f32>;
- var v2i : vec2<i32>;
- var v3i : vec3<i32>;
- var v4i : vec4<i32>;
- var v2u : vec2<u32>;
- var v3u : vec3<u32>;
- var v4u : vec4<u32>;
- var v2b : vec2<bool>;
- var v3b : vec3<bool>;
- var v4b : vec4<bool>;
-
- for (var i : i32 = 0; i < 2; i = i + 1) {
- v2f[i] = 1.0;
- v2i[i] = 1;
- v2u[i] = 1u;
- v2b[i] = true;
- }
-
- var i : i32 = 0;
- v3f[i] = 1.0;
- v4f[i] = 1.0;
- v3i[i] = 1;
- v4i[i] = 1;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v3b[i] = true;
- v4b[i] = true;
-}
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ var v2f : vec2<f32>;
+ var v3f : vec3<f32>;
+ var v4f : vec4<f32>;
+ var v2i : vec2<i32>;
+ var v3i : vec3<i32>;
+ var v4i : vec4<i32>;
+ var v2u : vec2<u32>;
+ var v3u : vec3<u32>;
+ var v4u : vec4<u32>;
+ var v2b : vec2<bool>;
+ var v3b : vec3<bool>;
+ var v4b : vec4<bool>;
+
+ for (var i : i32 = 0; i < 2; i = i + 1) {
+ v2f[i] = 1.0;
+ v2i[i] = 1;
+ v2u[i] = 1u;
+ v2b[i] = true;
+ }
+
+ var i : i32 = 0;
+ v3f[i] = 1.0;
+ v4f[i] = 1.0;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v3b[i] = true;
+ v4b[i] = true;
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm b/test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.spvasm
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl b/test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl b/test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl
similarity index 99%
rename from test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl
index 8103edf..276355c 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl
+++ b/test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl
@@ -1,29 +1,29 @@
-[[stage(compute), workgroup_size(1, 1, 1)]]
-fn main() {
- var v2f : vec2<f32>;
- var v3f : vec3<f32>;
- var v4f : vec4<f32>;
- var v2i : vec2<i32>;
- var v3i : vec3<i32>;
- var v4i : vec4<i32>;
- var v2u : vec2<u32>;
- var v3u : vec3<u32>;
- var v4u : vec4<u32>;
- var v2b : vec2<bool>;
- var v3b : vec3<bool>;
- var v4b : vec4<bool>;
-
- var i : i32 = 0;
- v2f[i] = 1.0;
- v3f[i] = 1.0;
- v4f[i] = 1.0;
- v2i[i] = 1;
- v3i[i] = 1;
- v4i[i] = 1;
- v2u[i] = 1u;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v2b[i] = true;
- v3b[i] = true;
- v4b[i] = true;
-}
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ var v2f : vec2<f32>;
+ var v3f : vec3<f32>;
+ var v4f : vec4<f32>;
+ var v2i : vec2<i32>;
+ var v3i : vec3<i32>;
+ var v4i : vec4<i32>;
+ var v2u : vec2<u32>;
+ var v3u : vec3<u32>;
+ var v4u : vec4<u32>;
+ var v2b : vec2<bool>;
+ var v3b : vec3<bool>;
+ var v4b : vec4<bool>;
+
+ var i : i32 = 0;
+ v2f[i] = 1.0;
+ v3f[i] = 1.0;
+ v4f[i] = 1.0;
+ v2i[i] = 1;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v2u[i] = 1u;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v3b[i] = true;
+ v4b[i] = true;
+}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl b/test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl
rename to test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.hlsl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.msl b/test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.msl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.msl
rename to test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.msl
diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm b/test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm
rename to test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.spvasm
diff --git a/test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl b/test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl
similarity index 100%
rename from test/fxc_bugs/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl
rename to test/bug/fxc/vector_assignment_in_loop/no_loop.wgsl.expected.wgsl
diff --git a/test/bug/tint/998.wgsl.expected.hlsl b/test/bug/tint/998.wgsl.expected.hlsl
index 88f4d37..9f7241a 100644
--- a/test/bug/tint/998.wgsl.expected.hlsl
+++ b/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 @@
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
+
