diff --git a/test/bug/tint/998.wgsl b/test/bug/tint/998.wgsl new file mode 100644 index 0000000000..2b3655f2c6 --- /dev/null +++ b/test/bug/tint/998.wgsl @@ -0,0 +1,19 @@ +[[block]] struct Constants { + zero: u32; +}; +[[group(1), binding(0)]] var constants: Constants; + +[[block]] struct Result { + value: u32; +}; +[[group(1), binding(1)]] var result: Result; + +struct S { + data: array; +}; +var s: S; + +[[stage(compute), workgroup_size(1)]] +fn main() { + s.data[constants.zero] = 0u; +} \ No newline at end of file diff --git a/test/bug/tint/998.wgsl.expected.hlsl b/test/bug/tint/998.wgsl.expected.hlsl new file mode 100644 index 0000000000..1acfe64bee --- /dev/null +++ b/test/bug/tint/998.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +SKIP: FAILED + +cbuffer cbuffer_constants : register(b0, space1) { + uint4 constants[1]; +}; + +RWByteAddressBuffer result : register(u1, space1); + +struct S { + uint data[3]; +}; + +static S s = (S)0; + +[numthreads(1, 1, 1)] +void main() { + s.data[constants[0].x] = 0u; + return; +} +C:\src\tint\test\Shader@0x0000015D0E1BAC50(15,3-24): error X3500: array reference cannot be used as an l-value; not natively addressable + diff --git a/test/bug/tint/998.wgsl.expected.msl b/test/bug/tint/998.wgsl.expected.msl new file mode 100644 index 0000000000..1c6fc8b053 --- /dev/null +++ b/test/bug/tint/998.wgsl.expected.msl @@ -0,0 +1,22 @@ +#include + +using namespace metal; +struct Constants { + /* 0x0000 */ uint zero; +}; +struct Result { + /* 0x0000 */ uint value; +}; +struct tint_array_wrapper { + uint arr[3]; +}; +struct S { + tint_array_wrapper data; +}; + +kernel void tint_symbol(constant Constants& constants [[buffer(0)]]) { + thread S tint_symbol_1 = {}; + tint_symbol_1.data.arr[constants.zero] = 0u; + return; +} + diff --git a/test/bug/tint/998.wgsl.expected.spvasm b/test/bug/tint/998.wgsl.expected.spvasm new file mode 100644 index 0000000000..36525c0b69 --- /dev/null +++ b/test/bug/tint/998.wgsl.expected.spvasm @@ -0,0 +1,57 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %Constants "Constants" + OpMemberName %Constants 0 "zero" + OpName %constants "constants" + OpName %Result "Result" + OpMemberName %Result 0 "value" + OpName %result "result" + OpName %S "S" + OpMemberName %S 0 "data" + OpName %s "s" + OpName %main "main" + OpDecorate %Constants Block + OpMemberDecorate %Constants 0 Offset 0 + OpDecorate %constants NonWritable + OpDecorate %constants DescriptorSet 1 + OpDecorate %constants Binding 0 + OpDecorate %Result Block + OpMemberDecorate %Result 0 Offset 0 + OpDecorate %result NonReadable + OpDecorate %result DescriptorSet 1 + OpDecorate %result Binding 1 + OpMemberDecorate %S 0 Offset 0 + OpDecorate %_arr_uint_uint_3 ArrayStride 4 + %uint = OpTypeInt 32 0 + %Constants = OpTypeStruct %uint +%_ptr_Uniform_Constants = OpTypePointer Uniform %Constants + %constants = OpVariable %_ptr_Uniform_Constants Uniform + %Result = OpTypeStruct %uint +%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result + %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer + %uint_3 = OpConstant %uint 3 +%_arr_uint_uint_3 = OpTypeArray %uint %uint_3 + %S = OpTypeStruct %_arr_uint_uint_3 +%_ptr_Private_S = OpTypePointer Private %S + %13 = OpConstantNull %S + %s = OpVariable %_ptr_Private_S Private %13 + %void = OpTypeVoid + %14 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint +%_ptr_Private_uint = OpTypePointer Private %uint + %main = OpFunction %void None %14 + %17 = OpLabel + %20 = OpAccessChain %_ptr_Uniform_uint %constants %uint_0 + %21 = OpLoad %uint %20 + %23 = OpAccessChain %_ptr_Private_uint %s %uint_0 %21 + OpStore %23 %uint_0 + OpReturn + OpFunctionEnd diff --git a/test/bug/tint/998.wgsl.expected.wgsl b/test/bug/tint/998.wgsl.expected.wgsl new file mode 100644 index 0000000000..acc05d2f42 --- /dev/null +++ b/test/bug/tint/998.wgsl.expected.wgsl @@ -0,0 +1,24 @@ +[[block]] +struct Constants { + zero : u32; +}; + +[[group(1), binding(0)]] var constants : Constants; + +[[block]] +struct Result { + value : u32; +}; + +[[group(1), binding(1)]] var result : Result; + +struct S { + data : array; +}; + +var s : S; + +[[stage(compute), workgroup_size(1)]] +fn main() { + s.data[constants.zero] = 0u; +}