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 0000000000..b1f7ad024e --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var 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 0000000000..b9a46178c5 --- /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 0000000000..da74b6f472 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/read/function.wgsl.expected.msl @@ -0,0 +1,22 @@ +#include + +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 0000000000..95c6f27817 --- /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 0000000000..1a660541e3 --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var 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 0000000000..731e57b530 --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var result: Result; + +var 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 0000000000..ef50e8a3ce --- /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 0000000000..24958dc128 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/read/private.wgsl.expected.msl @@ -0,0 +1,22 @@ +#include + +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 0000000000..54eb08ef2a --- /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 0000000000..68074bd832 --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var result : Result; + +var 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 0000000000..5237a5b760 --- /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 ubo: UBO; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(2)]] var result: Result; + +[[block]] +struct SSBO { + data: array; +}; +[[group(0), binding(1)]] var 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 0000000000..dada6d4c9d --- /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 0000000000..7bf987a437 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/read/storage.wgsl.expected.msl @@ -0,0 +1,21 @@ +#include + +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 0000000000..b5c4ce429a --- /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 0000000000..76100254de --- /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 ubo : UBO; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(2)]] var result : Result; + +[[block]] +struct SSBO { + data : array; +}; + +[[group(0), binding(1)]] var 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 0000000000..6edb1df484 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl @@ -0,0 +1,16 @@ +[[block]] +struct UBO { + data: [[stride(16)]] array; + dynamic_idx: i32; +}; +[[group(0), binding(0)]] var ubo: UBO; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(2)]] var 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 0000000000..01e61f78b5 --- /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 0000000000..34664447b2 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl @@ -0,0 +1,23 @@ +#include + +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 0000000000..1026040d3b --- /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 0000000000..661263003d --- /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; + dynamic_idx : i32; +}; + +[[group(0), binding(0)]] var ubo : UBO; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(2)]] var 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 0000000000..b598cd3851 --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var result: Result; + +var 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 0000000000..57a414019c --- /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 0000000000..f71e0a016f --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl @@ -0,0 +1,27 @@ +#include + +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 0000000000..09a3136c59 --- /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 0000000000..51c70f7365 --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var result : Result; + +var 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 0000000000..1fad5bcc1c --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var 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 0000000000..04d50a17ac --- /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 0000000000..21306716d0 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.msl @@ -0,0 +1,23 @@ +#include + +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 0000000000..46228e19d0 --- /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 0000000000..1871726cb7 --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var 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 0000000000..ab909db4bf --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var result: Result; + +fn x(p : ptr) { + (*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 0000000000..59669b372a --- /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 0000000000..74f99826d0 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.msl @@ -0,0 +1,27 @@ +#include + +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 0000000000..446b655cd8 --- /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 0000000000..c106f661d1 --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var result : Result; + +fn x(p : ptr) { + (*(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 0000000000..4f33168a27 --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var result: Result; + +var 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 0000000000..1ef5a388bc --- /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 0000000000..4e7678bc4a --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.msl @@ -0,0 +1,23 @@ +#include + +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 0000000000..341b1b8fcc --- /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 0000000000..42f58588ac --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var result : Result; + +var 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 0000000000..6c200746ea --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var result: Result; + +var s : S; + +fn x(p : ptr) { + (*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 0000000000..ed009a4d77 --- /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 0000000000..f72d1a2d9d --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.msl @@ -0,0 +1,27 @@ +#include + +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 0000000000..224618fad3 --- /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 0000000000..176c7274b9 --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var result : Result; + +var s : S; + +fn x(p : ptr) { + (*(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 0000000000..7bedb67a4f --- /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 ubo: UBO; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(2)]] var result: Result; + +[[block]] +struct SSBO { + data: array; +}; +[[group(0), binding(1)]] var 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 0000000000..005f32199f --- /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 0000000000..9e0044ffa5 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/storage.wgsl.expected.msl @@ -0,0 +1,22 @@ +#include + +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 0000000000..85915c5102 --- /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 0000000000..b06072b6e2 --- /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 ubo : UBO; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(2)]] var result : Result; + +[[block]] +struct SSBO { + data : array; +}; + +[[group(0), binding(1)]] var 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 0000000000..26efbe6492 --- /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 ubo: UBO; +struct S { + data: array; +}; +[[block]] +struct Result { + out: i32; +}; +[[group(0), binding(1)]] var result: Result; + +var 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 0000000000..0dd523fa14 --- /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 0000000000..76981a03c8 --- /dev/null +++ b/test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl @@ -0,0 +1,28 @@ +#include + +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 0000000000..f395a39a47 --- /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 0000000000..4b118715bb --- /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 ubo : UBO; + +struct S { + data : array; +}; + +[[block]] +struct Result { + out : i32; +}; + +[[group(0), binding(1)]] var result : Result; + +var 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 94% 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 8dc751527a..a684c5675c 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 v2f : vec2; -var v3i : vec3; -var v4u : vec4; -var v2b : vec2; - -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 v2f : vec2; +var v3i : vec3; +var v4u : vec4; +var v2b : vec2; + +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 94% 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 418155d2c5..c85e301ac0 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 v2f : vec2; -var v3i : vec3; -var v4u : vec4; -var v2b : vec2; - -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 v2f : vec2; +var v3i : vec3; +var v4u : vec4; +var v2b : vec2; + +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 95% 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 ef2ec78d6a..5d428383a6 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; - var v3f : vec3; - var v4f : vec4; - var v2i : vec2; - var v3i : vec3; - var v4i : vec4; - var v2u : vec2; - var v3u : vec3; - var v4u : vec4; - var v2b : vec2; - var v3b : vec3; - var v4b : vec4; - - 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; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + + 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 94% 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 d464d4c8d1..7b91e4887f 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; - var v2f_2 : vec2; - - var v3i : vec3; - var v3i_2 : vec3; - - var v4u : vec4; - var v4u_2 : vec4; - - var v2b : vec2; - var v2b_2 : vec2; - - 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; + var v2f_2 : vec2; + + var v3i : vec3; + var v3i_2 : vec3; + + var v4u : vec4; + var v4u_2 : vec4; + + var v2b : vec2; + var v2b_2 : vec2; + + 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 94% 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 d07e08d46b..d255bcc6bb 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; - var v3f : vec3; - var v4f : vec4; - var v2i : vec2; - var v3i : vec3; - var v4i : vec4; - var v2u : vec2; - var v3u : vec3; - var v4u : vec4; - var v2b : vec2; - var v3b : vec3; - var v4b : vec4; - - 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; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + + 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 94% 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 8103edfa4f..276355c706 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; - var v3f : vec3; - var v4f : vec4; - var v2i : vec2; - var v3i : vec3; - var v4i : vec4; - var v2u : vec2; - var v3u : vec3; - var v4u : vec4; - var v2b : vec2; - var v3b : vec3; - var v4b : vec4; - - 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; + var v3f : vec3; + var v4f : vec4; + var v2i : vec2; + var v3i : vec3; + var v4i : vec4; + var v2u : vec2; + var v3u : vec3; + var v4u : vec4; + var v2b : vec2; + var v3b : vec3; + var v4b : vec4; + + 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 88f4d37b64..9f7241a887 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 @@ void main() { s.data[constants[0].x] = 0u; return; } +C:\src\tint\test\Shader@0x000002B2827B9810(15,3-24): error X3500: array reference cannot be used as an l-value; not natively addressable +