Update test/tint/array to store results
This CL updates the `test/tint/array` tests to make sure the results
are stored back into a storage variable. Also make sure each test has
an entry point.
Change-Id: I8a87129c22ed56bdbcecf69464aa03e9ce387659
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/164580
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/test/tint/array/assign_to_function_var.wgsl b/test/tint/array/assign_to_function_var.wgsl
index c89220d..ba54b54 100644
--- a/test/tint/array/assign_to_function_var.wgsl
+++ b/test/tint/array/assign_to_function_var.wgsl
@@ -50,3 +50,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let val = ArrayType();
+ foo(val);
+}
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.dxc.hlsl b/test/tint/array/assign_to_function_var.wgsl.expected.dxc.hlsl
index 569734e..cdd7609 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.dxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -16,22 +11,22 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
typedef int4 src_uniform_load_ret[4];
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- const uint scalar_offset = ((offset + (i * 16u))) / 4;
- arr_1[i] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
+ arr_1[i_1] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -41,8 +36,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- arr_2[i_1] = asint(src_storage.Load4((offset + (i_1 * 16u))));
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ arr_2[i_2] = asint(src_storage.Load4((offset + (i_2 * 16u))));
}
}
return arr_2;
@@ -51,8 +46,8 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
int4 tint_symbol[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol = tint_symbol_4;
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol = tint_symbol_6;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@@ -68,3 +63,25 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 val[4] = (int4[4])0;
+ foo(val);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.fxc.hlsl b/test/tint/array/assign_to_function_var.wgsl.expected.fxc.hlsl
index 569734e..cdd7609 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.fxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -16,22 +11,22 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
typedef int4 src_uniform_load_ret[4];
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- const uint scalar_offset = ((offset + (i * 16u))) / 4;
- arr_1[i] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
+ arr_1[i_1] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -41,8 +36,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- arr_2[i_1] = asint(src_storage.Load4((offset + (i_1 * 16u))));
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ arr_2[i_2] = asint(src_storage.Load4((offset + (i_2 * 16u))));
}
}
return arr_2;
@@ -51,8 +46,8 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
int4 tint_symbol[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol = tint_symbol_4;
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol = tint_symbol_6;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@@ -68,3 +63,25 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 val[4] = (int4[4])0;
+ foo(val);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.glsl b/test/tint/array/assign_to_function_var.wgsl.expected.glsl
index f3c51f8..9154bc7 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.glsl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void unused_entry_point() {
- return;
-}
struct S {
ivec4 arr[4];
};
@@ -19,20 +15,20 @@
} src_storage;
ivec4[4] ret_arr() {
- ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
- return tint_symbol_1;
+ ivec4 tint_symbol_2[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ return tint_symbol_2;
}
S ret_struct_arr() {
- S tint_symbol_2 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
- return tint_symbol_2;
+ S tint_symbol_3 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
+ return tint_symbol_3;
}
void foo(ivec4 src_param[4]) {
ivec4 src_function[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
ivec4 dst[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
- ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
- dst = tint_symbol_3;
+ ivec4 tint_symbol_4[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
+ dst = tint_symbol_4;
dst = src_param;
dst = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
@@ -40,8 +36,8 @@
dst = src_function;
dst = src_private;
dst = src_workgroup;
- S tint_symbol = ret_struct_arr();
- dst = tint_symbol.arr;
+ S tint_symbol_1 = ret_struct_arr();
+ dst = tint_symbol_1.arr;
dst = src_uniform.inner.arr;
dst = src_storage.inner.arr;
int dst_nested[4][3][2] = int[4][3][2](int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)));
@@ -49,3 +45,20 @@
dst_nested = src_nested;
}
+void tint_symbol(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint i = idx;
+ src_workgroup[i] = ivec4(0);
+ }
+ }
+ barrier();
+ ivec4 val[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ foo(val);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
index 446d9ec..918de54 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
@@ -1,6 +1,6 @@
SKIP: FAILED
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: S = struct @align(16) {
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: S = struct @align(16) {
arr:array<vec4<i32>, 4> @offset(0)
}
@@ -53,6 +53,34 @@
ret
}
}
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void -> %b5 {
+ %b5 = block {
+ loop [i: %b6, b: %b7, c: %b8] { # loop_1
+ %b6 = block { # initializer
+ next_iteration %b7 %tint_local_index
+ }
+ %b7 = block (%idx:u32) { # body
+ %28:bool = gte %idx:u32, 4u
+ if %28 [t: %b9] { # if_1
+ %b9 = block { # true
+ exit_loop # loop_1
+ }
+ }
+ %29:ptr<workgroup, vec4<i32>, read_write> = access %src_workgroup, %idx:u32
+ store %29, vec4<i32>(0i)
+ continue %b8
+ }
+ %b8 = block { # continuing
+ %30:u32 = add %idx:u32, 1u
+ next_iteration %b7 %30
+ }
+ }
+ %31:void = msl.threadgroup_barrier 4u
+ %val:array<vec4<i32>, 4> = let array<vec4<i32>, 4>(vec4<i32>(0i))
+ %33:void = call %foo, %val
+ ret
+ }
+}
unhandled variable address space
********************************************************************
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.msl b/test/tint/array/assign_to_function_var.wgsl.expected.msl
index a63f260..7c0cd4c 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.msl
@@ -23,33 +23,50 @@
};
tint_array<int4, 4> ret_arr() {
- tint_array<int4, 4> const tint_symbol_1 = tint_array<int4, 4>{};
- return tint_symbol_1;
-}
-
-S ret_struct_arr() {
- S const tint_symbol_2 = S{};
+ tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{};
return tint_symbol_2;
}
-void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
+S ret_struct_arr() {
+ S const tint_symbol_3 = S{};
+ return tint_symbol_3;
+}
+
+void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7) {
tint_array<int4, 4> src_function = {};
tint_array<int4, 4> dst = {};
- tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
- dst = tint_symbol_3;
+ tint_array<int4, 4> const tint_symbol_4 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
+ dst = tint_symbol_4;
dst = src_param;
dst = ret_arr();
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
dst = src_let;
dst = src_function;
dst = (*(tint_private_vars)).src_private;
- dst = *(tint_symbol_4);
- S const tint_symbol = ret_struct_arr();
- dst = tint_symbol.arr;
- dst = (*(tint_symbol_5)).arr;
+ dst = *(tint_symbol_5);
+ S const tint_symbol_1 = ret_struct_arr();
+ dst = tint_symbol_1.arr;
dst = (*(tint_symbol_6)).arr;
+ dst = (*(tint_symbol_7)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> dst_nested = {};
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
dst_nested = src_nested;
}
+void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_8, const constant S* const tint_symbol_9, device S* const tint_symbol_10) {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint const i = idx;
+ (*(tint_symbol_8))[i] = int4(0);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int4, 4> const val = tint_array<int4, 4>{};
+ foo(val, tint_private_vars, tint_symbol_8, tint_symbol_9, tint_symbol_10);
+}
+
+kernel void tint_symbol(const constant S* tint_symbol_12 [[buffer(0)]], device S* tint_symbol_13 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ threadgroup tint_array<int4, 4> tint_symbol_11;
+ tint_symbol_inner(local_invocation_index, &(tint_private_vars), &(tint_symbol_11), tint_symbol_12, tint_symbol_13);
+ return;
+}
+
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.spvasm b/test/tint/array/assign_to_function_var.wgsl.expected.spvasm
index 46e5c7f..35158e1 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.spvasm
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.spvasm
@@ -1,12 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 65
+; Bound: 96
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
- OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpEntryPoint GLCompute %main "main" %local_invocation_index_1
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %src_private "src_private"
OpName %src_workgroup "src_workgroup"
OpName %src_uniform_block "src_uniform_block"
@@ -15,7 +16,6 @@
OpMemberName %S 0 "arr"
OpName %src_uniform "src_uniform"
OpName %src_storage "src_storage"
- OpName %unused_entry_point "unused_entry_point"
OpName %ret_arr "ret_arr"
OpName %ret_struct_arr "ret_struct_arr"
OpName %foo "foo"
@@ -24,6 +24,11 @@
OpName %dst "dst"
OpName %dst_nested "dst_nested"
OpName %src_nested "src_nested"
+ OpName %main_inner "main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %main "main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %src_uniform_block Block
OpMemberDecorate %src_uniform_block 0 Offset 0
@@ -36,14 +41,16 @@
OpDecorate %_arr_int_uint_2 ArrayStride 4
OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
%int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
- %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
- %8 = OpConstantNull %_arr_v4int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+ %10 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %10
%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_v4int_uint_4
@@ -52,20 +59,19 @@
%src_uniform = OpVariable %_ptr_Uniform_src_uniform_block Uniform
%_ptr_StorageBuffer_src_uniform_block = OpTypePointer StorageBuffer %src_uniform_block
%src_storage = OpVariable %_ptr_StorageBuffer_src_uniform_block StorageBuffer
+ %19 = OpTypeFunction %_arr_v4int_uint_4
+ %22 = OpTypeFunction %S
+ %25 = OpConstantNull %S
%void = OpTypeVoid
- %17 = OpTypeFunction %void
- %21 = OpTypeFunction %_arr_v4int_uint_4
- %24 = OpTypeFunction %S
- %27 = OpConstantNull %S
- %28 = OpTypeFunction %void %_arr_v4int_uint_4
+ %26 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%int_1 = OpConstant %int 1
- %36 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
+ %35 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2
- %38 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
+ %37 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3
- %40 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
- %41 = OpConstantComposite %_arr_v4int_uint_4 %36 %38 %40 %40
+ %39 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+ %40 = OpConstantComposite %_arr_v4int_uint_4 %35 %37 %39 %39
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
@@ -75,47 +81,90 @@
%_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3
%_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
- %62 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %17
- %20 = OpLabel
+ %61 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+ %64 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %70 = OpConstantNull %uint
+ %bool = OpTypeBool
+%_ptr_Workgroup_v4int = OpTypePointer Workgroup %v4int
+ %84 = OpConstantNull %v4int
+ %uint_1 = OpConstant %uint 1
+ %uint_264 = OpConstant %uint 264
+ %91 = OpTypeFunction %void
+ %ret_arr = OpFunction %_arr_v4int_uint_4 None %19
+ %21 = OpLabel
+ OpReturnValue %10
+ OpFunctionEnd
+%ret_struct_arr = OpFunction %S None %22
+ %24 = OpLabel
+ OpReturnValue %25
+ OpFunctionEnd
+ %foo = OpFunction %void None %26
+ %src_param = OpFunctionParameter %_arr_v4int_uint_4
+ %30 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %10
+ %dst = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %10
+ %dst_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61
+ OpStore %dst %40
+ OpStore %dst %src_param
+ %41 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+ OpStore %dst %41
+ OpStore %dst %10
+ %42 = OpLoad %_arr_v4int_uint_4 %src_function
+ OpStore %dst %42
+ %43 = OpLoad %_arr_v4int_uint_4 %src_private
+ OpStore %dst %43
+ %44 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+ OpStore %dst %44
+ %45 = OpFunctionCall %S %ret_struct_arr
+ %46 = OpCompositeExtract %_arr_v4int_uint_4 %45 0
+ OpStore %dst %46
+ %49 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
+ %50 = OpLoad %_arr_v4int_uint_4 %49
+ OpStore %dst %50
+ %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
+ %53 = OpLoad %_arr_v4int_uint_4 %52
+ OpStore %dst %53
+ %63 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+ OpStore %dst_nested %63
OpReturn
OpFunctionEnd
- %ret_arr = OpFunction %_arr_v4int_uint_4 None %21
- %23 = OpLabel
- OpReturnValue %8
+ %main_inner = OpFunction %void None %64
+%local_invocation_index = OpFunctionParameter %uint
+ %67 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %70
+ OpStore %idx %local_invocation_index
+ OpBranch %71
+ %71 = OpLabel
+ OpLoopMerge %72 %73 None
+ OpBranch %74
+ %74 = OpLabel
+ %76 = OpLoad %uint %idx
+ %77 = OpULessThan %bool %76 %uint_4
+ %75 = OpLogicalNot %bool %77
+ OpSelectionMerge %79 None
+ OpBranchConditional %75 %80 %79
+ %80 = OpLabel
+ OpBranch %72
+ %79 = OpLabel
+ %81 = OpLoad %uint %idx
+ %83 = OpAccessChain %_ptr_Workgroup_v4int %src_workgroup %81
+ OpStore %83 %84
+ OpBranch %73
+ %73 = OpLabel
+ %85 = OpLoad %uint %idx
+ %87 = OpIAdd %uint %85 %uint_1
+ OpStore %idx %87
+ OpBranch %71
+ %72 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %90 = OpFunctionCall %void %foo %10
+ OpReturn
OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %24
- %26 = OpLabel
- OpReturnValue %27
- OpFunctionEnd
- %foo = OpFunction %void None %28
- %src_param = OpFunctionParameter %_arr_v4int_uint_4
- %31 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
- %dst = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
- %dst_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %62
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %62
- OpStore %dst %41
- OpStore %dst %src_param
- %42 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
- OpStore %dst %42
- OpStore %dst %8
- %43 = OpLoad %_arr_v4int_uint_4 %src_function
- OpStore %dst %43
- %44 = OpLoad %_arr_v4int_uint_4 %src_private
- OpStore %dst %44
- %45 = OpLoad %_arr_v4int_uint_4 %src_workgroup
- OpStore %dst %45
- %46 = OpFunctionCall %S %ret_struct_arr
- %47 = OpCompositeExtract %_arr_v4int_uint_4 %46 0
- OpStore %dst %47
- %50 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
- %51 = OpLoad %_arr_v4int_uint_4 %50
- OpStore %dst %51
- %53 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
- %54 = OpLoad %_arr_v4int_uint_4 %53
- OpStore %dst %54
- %64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
- OpStore %dst_nested %64
+ %main = OpFunction %void None %91
+ %93 = OpLabel
+ %95 = OpLoad %uint %local_invocation_index_1
+ %94 = OpFunctionCall %void %main_inner %95
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.wgsl b/test/tint/array/assign_to_function_var.wgsl.expected.wgsl
index ae87ab4..99d4b32 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.wgsl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.wgsl
@@ -38,3 +38,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let val = ArrayType();
+ foo(val);
+}
diff --git a/test/tint/array/assign_to_private_var.wgsl b/test/tint/array/assign_to_private_var.wgsl
index 8180341..49d4f81 100644
--- a/test/tint/array/assign_to_private_var.wgsl
+++ b/test/tint/array/assign_to_private_var.wgsl
@@ -50,3 +50,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let a = ArrayType();
+ foo(a);
+}
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.dxc.hlsl b/test/tint/array/assign_to_private_var.wgsl.expected.dxc.hlsl
index 61083f2..cef276f 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.dxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -18,22 +13,22 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
typedef int4 src_uniform_load_ret[4];
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- const uint scalar_offset = ((offset + (i * 16u))) / 4;
- arr_1[i] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
+ arr_1[i_1] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -43,8 +38,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- arr_2[i_1] = asint(src_storage.Load4((offset + (i_1 * 16u))));
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ arr_2[i_2] = asint(src_storage.Load4((offset + (i_2 * 16u))));
}
}
return arr_2;
@@ -52,8 +47,8 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol = tint_symbol_4;
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol = tint_symbol_6;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@@ -68,3 +63,25 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 a[4] = (int4[4])0;
+ foo(a);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.fxc.hlsl b/test/tint/array/assign_to_private_var.wgsl.expected.fxc.hlsl
index 61083f2..cef276f 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.fxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -18,22 +13,22 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
typedef int4 src_uniform_load_ret[4];
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- const uint scalar_offset = ((offset + (i * 16u))) / 4;
- arr_1[i] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
+ arr_1[i_1] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -43,8 +38,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- arr_2[i_1] = asint(src_storage.Load4((offset + (i_1 * 16u))));
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ arr_2[i_2] = asint(src_storage.Load4((offset + (i_2 * 16u))));
}
}
return arr_2;
@@ -52,8 +47,8 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol = tint_symbol_4;
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol = tint_symbol_6;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@@ -68,3 +63,25 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 a[4] = (int4[4])0;
+ foo(a);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.glsl b/test/tint/array/assign_to_private_var.wgsl.expected.glsl
index bc6ac31..b7dacea 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.glsl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void unused_entry_point() {
- return;
-}
struct S {
ivec4 arr[4];
};
@@ -21,19 +17,19 @@
ivec4 dst[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
int dst_nested[4][3][2] = int[4][3][2](int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)));
ivec4[4] ret_arr() {
- ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
- return tint_symbol_1;
+ ivec4 tint_symbol_2[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ return tint_symbol_2;
}
S ret_struct_arr() {
- S tint_symbol_2 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
- return tint_symbol_2;
+ S tint_symbol_3 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
+ return tint_symbol_3;
}
void foo(ivec4 src_param[4]) {
ivec4 src_function[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
- ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
- dst = tint_symbol_3;
+ ivec4 tint_symbol_4[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
+ dst = tint_symbol_4;
dst = src_param;
dst = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
@@ -41,11 +37,28 @@
dst = src_function;
dst = src_private;
dst = src_workgroup;
- S tint_symbol = ret_struct_arr();
- dst = tint_symbol.arr;
+ S tint_symbol_1 = ret_struct_arr();
+ dst = tint_symbol_1.arr;
dst = src_uniform.inner.arr;
dst = src_storage.inner.arr;
int src_nested[4][3][2] = int[4][3][2](int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)));
dst_nested = src_nested;
}
+void tint_symbol(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint i = idx;
+ src_workgroup[i] = ivec4(0);
+ }
+ }
+ barrier();
+ ivec4 a[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ foo(a);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
index baa11e0..2c0797c 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
@@ -1,6 +1,6 @@
SKIP: FAILED
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: S = struct @align(16) {
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: S = struct @align(16) {
arr:array<vec4<i32>, 4> @offset(0)
}
@@ -53,6 +53,34 @@
ret
}
}
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void -> %b5 {
+ %b5 = block {
+ loop [i: %b6, b: %b7, c: %b8] { # loop_1
+ %b6 = block { # initializer
+ next_iteration %b7 %tint_local_index
+ }
+ %b7 = block (%idx:u32) { # body
+ %28:bool = gte %idx:u32, 4u
+ if %28 [t: %b9] { # if_1
+ %b9 = block { # true
+ exit_loop # loop_1
+ }
+ }
+ %29:ptr<workgroup, vec4<i32>, read_write> = access %src_workgroup, %idx:u32
+ store %29, vec4<i32>(0i)
+ continue %b8
+ }
+ %b8 = block { # continuing
+ %30:u32 = add %idx:u32, 1u
+ next_iteration %b7 %30
+ }
+ }
+ %31:void = msl.threadgroup_barrier 4u
+ %a:array<vec4<i32>, 4> = let array<vec4<i32>, 4>(vec4<i32>(0i))
+ %33:void = call %foo, %a
+ ret
+ }
+}
unhandled variable address space
********************************************************************
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.msl b/test/tint/array/assign_to_private_var.wgsl.expected.msl
index d96bf73..bfe306f 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.msl
@@ -25,31 +25,48 @@
};
tint_array<int4, 4> ret_arr() {
- tint_array<int4, 4> const tint_symbol_1 = tint_array<int4, 4>{};
- return tint_symbol_1;
-}
-
-S ret_struct_arr() {
- S const tint_symbol_2 = S{};
+ tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{};
return tint_symbol_2;
}
-void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
+S ret_struct_arr() {
+ S const tint_symbol_3 = S{};
+ return tint_symbol_3;
+}
+
+void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7) {
tint_array<int4, 4> src_function = {};
- tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
- (*(tint_private_vars)).dst = tint_symbol_3;
+ tint_array<int4, 4> const tint_symbol_4 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
+ (*(tint_private_vars)).dst = tint_symbol_4;
(*(tint_private_vars)).dst = src_param;
(*(tint_private_vars)).dst = ret_arr();
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
(*(tint_private_vars)).dst = src_let;
(*(tint_private_vars)).dst = src_function;
(*(tint_private_vars)).dst = (*(tint_private_vars)).src_private;
- (*(tint_private_vars)).dst = *(tint_symbol_4);
- S const tint_symbol = ret_struct_arr();
- (*(tint_private_vars)).dst = tint_symbol.arr;
- (*(tint_private_vars)).dst = (*(tint_symbol_5)).arr;
+ (*(tint_private_vars)).dst = *(tint_symbol_5);
+ S const tint_symbol_1 = ret_struct_arr();
+ (*(tint_private_vars)).dst = tint_symbol_1.arr;
(*(tint_private_vars)).dst = (*(tint_symbol_6)).arr;
+ (*(tint_private_vars)).dst = (*(tint_symbol_7)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
(*(tint_private_vars)).dst_nested = src_nested;
}
+void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_8, const constant S* const tint_symbol_9, device S* const tint_symbol_10) {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint const i = idx;
+ (*(tint_symbol_8))[i] = int4(0);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int4, 4> const a = tint_array<int4, 4>{};
+ foo(a, tint_private_vars, tint_symbol_8, tint_symbol_9, tint_symbol_10);
+}
+
+kernel void tint_symbol(const constant S* tint_symbol_12 [[buffer(0)]], device S* tint_symbol_13 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ threadgroup tint_array<int4, 4> tint_symbol_11;
+ tint_symbol_inner(local_invocation_index, &(tint_private_vars), &(tint_symbol_11), tint_symbol_12, tint_symbol_13);
+ return;
+}
+
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.spvasm b/test/tint/array/assign_to_private_var.wgsl.expected.spvasm
index caae540..71f2976 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.spvasm
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.spvasm
@@ -1,12 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 66
+; Bound: 97
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
- OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpEntryPoint GLCompute %main "main" %local_invocation_index_1
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %src_private "src_private"
OpName %src_workgroup "src_workgroup"
OpName %src_uniform_block "src_uniform_block"
@@ -17,13 +18,17 @@
OpName %src_storage "src_storage"
OpName %dst "dst"
OpName %dst_nested "dst_nested"
- OpName %unused_entry_point "unused_entry_point"
OpName %ret_arr "ret_arr"
OpName %ret_struct_arr "ret_struct_arr"
OpName %foo "foo"
OpName %src_param "src_param"
OpName %src_function "src_function"
OpName %src_nested "src_nested"
+ OpName %main_inner "main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %main "main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %src_uniform_block Block
OpMemberDecorate %src_uniform_block 0 Offset 0
@@ -36,14 +41,16 @@
OpDecorate %_arr_int_uint_2 ArrayStride 4
OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
%int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
- %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
- %8 = OpConstantNull %_arr_v4int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+ %10 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %10
%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_v4int_uint_4
@@ -52,71 +59,113 @@
%src_uniform = OpVariable %_ptr_Uniform_src_uniform_block Uniform
%_ptr_StorageBuffer_src_uniform_block = OpTypePointer StorageBuffer %src_uniform_block
%src_storage = OpVariable %_ptr_StorageBuffer_src_uniform_block StorageBuffer
- %dst = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+ %dst = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %10
%uint_2 = OpConstant %uint 2
%_arr_int_uint_2 = OpTypeArray %int %uint_2
%uint_3 = OpConstant %uint 3
%_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3
%_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
%_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Private %_arr__arr__arr_int_uint_2_uint_3_uint_4
- %25 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
- %dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %25
+ %27 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+ %dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %27
+ %28 = OpTypeFunction %_arr_v4int_uint_4
+ %31 = OpTypeFunction %S
+ %34 = OpConstantNull %S
%void = OpTypeVoid
- %26 = OpTypeFunction %void
- %30 = OpTypeFunction %_arr_v4int_uint_4
- %33 = OpTypeFunction %S
- %36 = OpConstantNull %S
- %37 = OpTypeFunction %void %_arr_v4int_uint_4
+ %35 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%int_1 = OpConstant %int 1
- %44 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
+ %43 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2
- %46 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
+ %45 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3
- %48 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
- %49 = OpConstantComposite %_arr_v4int_uint_4 %44 %46 %48 %48
+ %47 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+ %48 = OpConstantComposite %_arr_v4int_uint_4 %43 %45 %47 %47
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %26
- %29 = OpLabel
+ %65 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %71 = OpConstantNull %uint
+ %bool = OpTypeBool
+%_ptr_Workgroup_v4int = OpTypePointer Workgroup %v4int
+ %85 = OpConstantNull %v4int
+ %uint_1 = OpConstant %uint 1
+ %uint_264 = OpConstant %uint 264
+ %92 = OpTypeFunction %void
+ %ret_arr = OpFunction %_arr_v4int_uint_4 None %28
+ %30 = OpLabel
+ OpReturnValue %10
+ OpFunctionEnd
+%ret_struct_arr = OpFunction %S None %31
+ %33 = OpLabel
+ OpReturnValue %34
+ OpFunctionEnd
+ %foo = OpFunction %void None %35
+ %src_param = OpFunctionParameter %_arr_v4int_uint_4
+ %39 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %10
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %27
+ OpStore %dst %48
+ OpStore %dst %src_param
+ %49 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+ OpStore %dst %49
+ OpStore %dst %10
+ %50 = OpLoad %_arr_v4int_uint_4 %src_function
+ OpStore %dst %50
+ %51 = OpLoad %_arr_v4int_uint_4 %src_private
+ OpStore %dst %51
+ %52 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+ OpStore %dst %52
+ %53 = OpFunctionCall %S %ret_struct_arr
+ %54 = OpCompositeExtract %_arr_v4int_uint_4 %53 0
+ OpStore %dst %54
+ %57 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
+ %58 = OpLoad %_arr_v4int_uint_4 %57
+ OpStore %dst %58
+ %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
+ %61 = OpLoad %_arr_v4int_uint_4 %60
+ OpStore %dst %61
+ %64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+ OpStore %dst_nested %64
OpReturn
OpFunctionEnd
- %ret_arr = OpFunction %_arr_v4int_uint_4 None %30
- %32 = OpLabel
- OpReturnValue %8
+ %main_inner = OpFunction %void None %65
+%local_invocation_index = OpFunctionParameter %uint
+ %68 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %71
+ OpStore %idx %local_invocation_index
+ OpBranch %72
+ %72 = OpLabel
+ OpLoopMerge %73 %74 None
+ OpBranch %75
+ %75 = OpLabel
+ %77 = OpLoad %uint %idx
+ %78 = OpULessThan %bool %77 %uint_4
+ %76 = OpLogicalNot %bool %78
+ OpSelectionMerge %80 None
+ OpBranchConditional %76 %81 %80
+ %81 = OpLabel
+ OpBranch %73
+ %80 = OpLabel
+ %82 = OpLoad %uint %idx
+ %84 = OpAccessChain %_ptr_Workgroup_v4int %src_workgroup %82
+ OpStore %84 %85
+ OpBranch %74
+ %74 = OpLabel
+ %86 = OpLoad %uint %idx
+ %88 = OpIAdd %uint %86 %uint_1
+ OpStore %idx %88
+ OpBranch %72
+ %73 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %91 = OpFunctionCall %void %foo %10
+ OpReturn
OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %33
- %35 = OpLabel
- OpReturnValue %36
- OpFunctionEnd
- %foo = OpFunction %void None %37
- %src_param = OpFunctionParameter %_arr_v4int_uint_4
- %40 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %25
- OpStore %dst %49
- OpStore %dst %src_param
- %50 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
- OpStore %dst %50
- OpStore %dst %8
- %51 = OpLoad %_arr_v4int_uint_4 %src_function
- OpStore %dst %51
- %52 = OpLoad %_arr_v4int_uint_4 %src_private
- OpStore %dst %52
- %53 = OpLoad %_arr_v4int_uint_4 %src_workgroup
- OpStore %dst %53
- %54 = OpFunctionCall %S %ret_struct_arr
- %55 = OpCompositeExtract %_arr_v4int_uint_4 %54 0
- OpStore %dst %55
- %58 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
- %59 = OpLoad %_arr_v4int_uint_4 %58
- OpStore %dst %59
- %61 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
- %62 = OpLoad %_arr_v4int_uint_4 %61
- OpStore %dst %62
- %65 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
- OpStore %dst_nested %65
+ %main = OpFunction %void None %92
+ %94 = OpLabel
+ %96 = OpLoad %uint %local_invocation_index_1
+ %95 = OpFunctionCall %void %main_inner %96
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.wgsl b/test/tint/array/assign_to_private_var.wgsl.expected.wgsl
index 63721f9..ba1a117 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.wgsl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.wgsl
@@ -40,3 +40,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let a = ArrayType();
+ foo(a);
+}
diff --git a/test/tint/array/assign_to_storage_var.wgsl b/test/tint/array/assign_to_storage_var.wgsl
index 69fc1ab..1f8f1e6 100644
--- a/test/tint/array/assign_to_storage_var.wgsl
+++ b/test/tint/array/assign_to_storage_var.wgsl
@@ -54,3 +54,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested.arr = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let ary = ArrayType();
+ foo(ary);
+}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl b/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
index dcc3c11..105b407 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.dxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -18,20 +13,20 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
void tint_symbol_store(uint offset, int4 value[4]) {
int4 array_1[4] = value;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- tint_symbol.Store4((offset + (i * 16u)), asuint(array_1[i]));
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ tint_symbol.Store4((offset + (i_1 * 16u)), asuint(array_1[i_1]));
}
}
}
@@ -40,9 +35,9 @@
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
- arr_1[i_1] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ const uint scalar_offset = ((offset + (i_2 * 16u))) / 4;
+ arr_1[i_2] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -52,8 +47,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
- arr_2[i_2] = asint(src_storage.Load4((offset + (i_2 * 16u))));
+ for(uint i_3 = 0u; (i_3 < 4u); i_3 = (i_3 + 1u)) {
+ arr_2[i_3] = asint(src_storage.Load4((offset + (i_3 * 16u))));
}
}
return arr_2;
@@ -62,8 +57,8 @@
void dst_nested_store_2(uint offset, int value[2]) {
int array_4[2] = value;
{
- for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) {
- dst_nested.Store((offset + (i_3 * 4u)), asuint(array_4[i_3]));
+ for(uint i_4 = 0u; (i_4 < 2u); i_4 = (i_4 + 1u)) {
+ dst_nested.Store((offset + (i_4 * 4u)), asuint(array_4[i_4]));
}
}
}
@@ -71,8 +66,8 @@
void dst_nested_store_1(uint offset, int value[3][2]) {
int array_3[3][2] = value;
{
- for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) {
- dst_nested_store_2((offset + (i_4 * 8u)), array_3[i_4]);
+ for(uint i_5 = 0u; (i_5 < 3u); i_5 = (i_5 + 1u)) {
+ dst_nested_store_2((offset + (i_5 * 8u)), array_3[i_5]);
}
}
}
@@ -80,16 +75,16 @@
void dst_nested_store(uint offset, int value[4][3][2]) {
int array_2[4][3][2] = value;
{
- for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
- dst_nested_store_1((offset + (i_5 * 24u)), array_2[i_5]);
+ for(uint i_6 = 0u; (i_6 < 4u); i_6 = (i_6 + 1u)) {
+ dst_nested_store_1((offset + (i_6 * 24u)), array_2[i_6]);
}
}
}
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol_store(0u, tint_symbol_4);
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol_store(0u, tint_symbol_6);
tint_symbol_store(0u, src_param);
tint_symbol_store(0u, ret_arr());
const int4 src_let[4] = (int4[4])0;
@@ -104,3 +99,25 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested_store(0u, src_nested);
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 ary[4] = (int4[4])0;
+ foo(ary);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl b/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
index dcc3c11..105b407 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.fxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -18,20 +13,20 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
void tint_symbol_store(uint offset, int4 value[4]) {
int4 array_1[4] = value;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- tint_symbol.Store4((offset + (i * 16u)), asuint(array_1[i]));
+ for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+ tint_symbol.Store4((offset + (i_1 * 16u)), asuint(array_1[i_1]));
}
}
}
@@ -40,9 +35,9 @@
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
- arr_1[i_1] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ const uint scalar_offset = ((offset + (i_2 * 16u))) / 4;
+ arr_1[i_2] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -52,8 +47,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
- arr_2[i_2] = asint(src_storage.Load4((offset + (i_2 * 16u))));
+ for(uint i_3 = 0u; (i_3 < 4u); i_3 = (i_3 + 1u)) {
+ arr_2[i_3] = asint(src_storage.Load4((offset + (i_3 * 16u))));
}
}
return arr_2;
@@ -62,8 +57,8 @@
void dst_nested_store_2(uint offset, int value[2]) {
int array_4[2] = value;
{
- for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) {
- dst_nested.Store((offset + (i_3 * 4u)), asuint(array_4[i_3]));
+ for(uint i_4 = 0u; (i_4 < 2u); i_4 = (i_4 + 1u)) {
+ dst_nested.Store((offset + (i_4 * 4u)), asuint(array_4[i_4]));
}
}
}
@@ -71,8 +66,8 @@
void dst_nested_store_1(uint offset, int value[3][2]) {
int array_3[3][2] = value;
{
- for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) {
- dst_nested_store_2((offset + (i_4 * 8u)), array_3[i_4]);
+ for(uint i_5 = 0u; (i_5 < 3u); i_5 = (i_5 + 1u)) {
+ dst_nested_store_2((offset + (i_5 * 8u)), array_3[i_5]);
}
}
}
@@ -80,16 +75,16 @@
void dst_nested_store(uint offset, int value[4][3][2]) {
int array_2[4][3][2] = value;
{
- for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
- dst_nested_store_1((offset + (i_5 * 24u)), array_2[i_5]);
+ for(uint i_6 = 0u; (i_6 < 4u); i_6 = (i_6 + 1u)) {
+ dst_nested_store_1((offset + (i_6 * 24u)), array_2[i_6]);
}
}
}
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol_store(0u, tint_symbol_4);
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol_store(0u, tint_symbol_6);
tint_symbol_store(0u, src_param);
tint_symbol_store(0u, ret_arr());
const int4 src_let[4] = (int4[4])0;
@@ -104,3 +99,25 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested_store(0u, src_nested);
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 ary[4] = (int4[4])0;
+ foo(ary);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.glsl b/test/tint/array/assign_to_storage_var.wgsl.expected.glsl
index 41ab9c1..9ca42b3 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.glsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void unused_entry_point() {
- return;
-}
struct S {
ivec4 arr[4];
};
@@ -31,19 +27,19 @@
} dst_nested;
ivec4[4] ret_arr() {
- ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
- return tint_symbol_1;
+ ivec4 tint_symbol_2[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ return tint_symbol_2;
}
S ret_struct_arr() {
- S tint_symbol_2 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
- return tint_symbol_2;
+ S tint_symbol_3 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
+ return tint_symbol_3;
}
void foo(ivec4 src_param[4]) {
ivec4 src_function[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
- ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
- dst.inner.arr = tint_symbol_3;
+ ivec4 tint_symbol_4[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
+ dst.inner.arr = tint_symbol_4;
dst.inner.arr = src_param;
dst.inner.arr = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
@@ -51,11 +47,28 @@
dst.inner.arr = src_function;
dst.inner.arr = src_private;
dst.inner.arr = src_workgroup;
- S tint_symbol = ret_struct_arr();
- dst.inner.arr = tint_symbol.arr;
+ S tint_symbol_1 = ret_struct_arr();
+ dst.inner.arr = tint_symbol_1.arr;
dst.inner.arr = src_uniform.inner.arr;
dst.inner.arr = src_storage.inner.arr;
int src_nested[4][3][2] = int[4][3][2](int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)));
dst_nested.inner.arr = src_nested;
}
+void tint_symbol(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint i = idx;
+ src_workgroup[i] = ivec4(0);
+ }
+ }
+ barrier();
+ ivec4 ary[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ foo(ary);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
index c248454..42ffcfd 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
@@ -1,6 +1,6 @@
SKIP: FAILED
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: S = struct @align(16) {
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: S = struct @align(16) {
arr:array<vec4<i32>, 4> @offset(0)
}
@@ -68,6 +68,34 @@
ret
}
}
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void -> %b5 {
+ %b5 = block {
+ loop [i: %b6, b: %b7, c: %b8] { # loop_1
+ %b6 = block { # initializer
+ next_iteration %b7 %tint_local_index
+ }
+ %b7 = block (%idx:u32) { # body
+ %39:bool = gte %idx:u32, 4u
+ if %39 [t: %b9] { # if_1
+ %b9 = block { # true
+ exit_loop # loop_1
+ }
+ }
+ %40:ptr<workgroup, vec4<i32>, read_write> = access %src_workgroup, %idx:u32
+ store %40, vec4<i32>(0i)
+ continue %b8
+ }
+ %b8 = block { # continuing
+ %41:u32 = add %idx:u32, 1u
+ next_iteration %b7 %41
+ }
+ }
+ %42:void = msl.threadgroup_barrier 4u
+ %ary:array<vec4<i32>, 4> = let array<vec4<i32>, 4>(vec4<i32>(0i))
+ %44:void = call %foo, %ary
+ ret
+ }
+}
unhandled variable address space
********************************************************************
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
index 06a713d..b5d6233 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
@@ -27,31 +27,48 @@
};
tint_array<int4, 4> ret_arr() {
- tint_array<int4, 4> const tint_symbol_1 = tint_array<int4, 4>{};
- return tint_symbol_1;
-}
-
-S ret_struct_arr() {
- S const tint_symbol_2 = S{};
+ tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{};
return tint_symbol_2;
}
-void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, device S* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, device S_nested* const tint_symbol_8) {
+S ret_struct_arr() {
+ S const tint_symbol_3 = S{};
+ return tint_symbol_3;
+}
+
+void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, device S* const tint_symbol_5, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, device S_nested* const tint_symbol_9) {
tint_array<int4, 4> src_function = {};
- tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
- (*(tint_symbol_4)).arr = tint_symbol_3;
- (*(tint_symbol_4)).arr = src_param;
- (*(tint_symbol_4)).arr = ret_arr();
+ tint_array<int4, 4> const tint_symbol_4 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
+ (*(tint_symbol_5)).arr = tint_symbol_4;
+ (*(tint_symbol_5)).arr = src_param;
+ (*(tint_symbol_5)).arr = ret_arr();
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
- (*(tint_symbol_4)).arr = src_let;
- (*(tint_symbol_4)).arr = src_function;
- (*(tint_symbol_4)).arr = (*(tint_private_vars)).src_private;
- (*(tint_symbol_4)).arr = *(tint_symbol_5);
- S const tint_symbol = ret_struct_arr();
- (*(tint_symbol_4)).arr = tint_symbol.arr;
- (*(tint_symbol_4)).arr = (*(tint_symbol_6)).arr;
- (*(tint_symbol_4)).arr = (*(tint_symbol_7)).arr;
+ (*(tint_symbol_5)).arr = src_let;
+ (*(tint_symbol_5)).arr = src_function;
+ (*(tint_symbol_5)).arr = (*(tint_private_vars)).src_private;
+ (*(tint_symbol_5)).arr = *(tint_symbol_6);
+ S const tint_symbol_1 = ret_struct_arr();
+ (*(tint_symbol_5)).arr = tint_symbol_1.arr;
+ (*(tint_symbol_5)).arr = (*(tint_symbol_7)).arr;
+ (*(tint_symbol_5)).arr = (*(tint_symbol_8)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
- (*(tint_symbol_8)).arr = src_nested;
+ (*(tint_symbol_9)).arr = src_nested;
+}
+
+void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_10, device S* const tint_symbol_11, const constant S* const tint_symbol_12, device S* const tint_symbol_13, device S_nested* const tint_symbol_14) {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint const i = idx;
+ (*(tint_symbol_10))[i] = int4(0);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int4, 4> const ary = tint_array<int4, 4>{};
+ foo(ary, tint_private_vars, tint_symbol_11, tint_symbol_10, tint_symbol_12, tint_symbol_13, tint_symbol_14);
+}
+
+kernel void tint_symbol(device S* tint_symbol_16 [[buffer(1)]], const constant S* tint_symbol_17 [[buffer(0)]], device S* tint_symbol_18 [[buffer(2)]], device S_nested* tint_symbol_19 [[buffer(3)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ threadgroup tint_array<int4, 4> tint_symbol_15;
+ tint_symbol_inner(local_invocation_index, &(tint_private_vars), &(tint_symbol_15), tint_symbol_16, tint_symbol_17, tint_symbol_18, tint_symbol_19);
+ return;
}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm b/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
index c2fb753..46d2806 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
@@ -1,12 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 80
+; Bound: 111
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
- OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpEntryPoint GLCompute %main "main" %local_invocation_index_1
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %src_private "src_private"
OpName %src_workgroup "src_workgroup"
OpName %src_uniform_block "src_uniform_block"
@@ -21,13 +22,17 @@
OpName %S_nested "S_nested"
OpMemberName %S_nested 0 "arr"
OpName %dst_nested "dst_nested"
- OpName %unused_entry_point "unused_entry_point"
OpName %ret_arr "ret_arr"
OpName %ret_struct_arr "ret_struct_arr"
OpName %foo "foo"
OpName %src_param "src_param"
OpName %src_function "src_function"
OpName %src_nested "src_nested"
+ OpName %main_inner "main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %main "main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %src_uniform_block Block
OpMemberDecorate %src_uniform_block 0 Offset 0
@@ -47,14 +52,16 @@
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
OpDecorate %dst_nested DescriptorSet 0
OpDecorate %dst_nested Binding 3
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
%int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
- %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
- %8 = OpConstantNull %_arr_v4int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+ %10 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %10
%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_v4int_uint_4
@@ -73,75 +80,117 @@
%dst_nested_block = OpTypeStruct %S_nested
%_ptr_StorageBuffer_dst_nested_block = OpTypePointer StorageBuffer %dst_nested_block
%dst_nested = OpVariable %_ptr_StorageBuffer_dst_nested_block StorageBuffer
+ %29 = OpTypeFunction %_arr_v4int_uint_4
+ %32 = OpTypeFunction %S
+ %35 = OpConstantNull %S
%void = OpTypeVoid
- %27 = OpTypeFunction %void
- %31 = OpTypeFunction %_arr_v4int_uint_4
- %34 = OpTypeFunction %S
- %37 = OpConstantNull %S
- %38 = OpTypeFunction %void %_arr_v4int_uint_4
+ %36 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%int_1 = OpConstant %int 1
- %48 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
+ %47 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2
- %50 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
+ %49 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3
- %52 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
- %53 = OpConstantComposite %_arr_v4int_uint_4 %48 %50 %52 %52
+ %51 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+ %52 = OpConstantComposite %_arr_v4int_uint_4 %47 %49 %51 %51
%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
- %76 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+ %75 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
%_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer StorageBuffer %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %27
- %30 = OpLabel
+ %79 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %85 = OpConstantNull %uint
+ %bool = OpTypeBool
+%_ptr_Workgroup_v4int = OpTypePointer Workgroup %v4int
+ %99 = OpConstantNull %v4int
+ %uint_1 = OpConstant %uint 1
+ %uint_264 = OpConstant %uint 264
+ %106 = OpTypeFunction %void
+ %ret_arr = OpFunction %_arr_v4int_uint_4 None %29
+ %31 = OpLabel
+ OpReturnValue %10
+ OpFunctionEnd
+%ret_struct_arr = OpFunction %S None %32
+ %34 = OpLabel
+ OpReturnValue %35
+ OpFunctionEnd
+ %foo = OpFunction %void None %36
+ %src_param = OpFunctionParameter %_arr_v4int_uint_4
+ %40 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %10
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %75
+ %45 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ OpStore %45 %52
+ %53 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ OpStore %53 %src_param
+ %54 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %55 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+ OpStore %54 %55
+ %56 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ OpStore %56 %10
+ %57 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %58 = OpLoad %_arr_v4int_uint_4 %src_function
+ OpStore %57 %58
+ %59 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %60 = OpLoad %_arr_v4int_uint_4 %src_private
+ OpStore %59 %60
+ %61 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %62 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+ OpStore %61 %62
+ %63 = OpFunctionCall %S %ret_struct_arr
+ %64 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %65 = OpCompositeExtract %_arr_v4int_uint_4 %63 0
+ OpStore %64 %65
+ %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %68 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
+ %69 = OpLoad %_arr_v4int_uint_4 %68
+ OpStore %66 %69
+ %70 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
+ %71 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
+ %72 = OpLoad %_arr_v4int_uint_4 %71
+ OpStore %70 %72
+ %77 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0 %uint_0
+ %78 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+ OpStore %77 %78
OpReturn
OpFunctionEnd
- %ret_arr = OpFunction %_arr_v4int_uint_4 None %31
- %33 = OpLabel
- OpReturnValue %8
+ %main_inner = OpFunction %void None %79
+%local_invocation_index = OpFunctionParameter %uint
+ %82 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %85
+ OpStore %idx %local_invocation_index
+ OpBranch %86
+ %86 = OpLabel
+ OpLoopMerge %87 %88 None
+ OpBranch %89
+ %89 = OpLabel
+ %91 = OpLoad %uint %idx
+ %92 = OpULessThan %bool %91 %uint_4
+ %90 = OpLogicalNot %bool %92
+ OpSelectionMerge %94 None
+ OpBranchConditional %90 %95 %94
+ %95 = OpLabel
+ OpBranch %87
+ %94 = OpLabel
+ %96 = OpLoad %uint %idx
+ %98 = OpAccessChain %_ptr_Workgroup_v4int %src_workgroup %96
+ OpStore %98 %99
+ OpBranch %88
+ %88 = OpLabel
+ %100 = OpLoad %uint %idx
+ %102 = OpIAdd %uint %100 %uint_1
+ OpStore %idx %102
+ OpBranch %86
+ %87 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %105 = OpFunctionCall %void %foo %10
+ OpReturn
OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %34
- %36 = OpLabel
- OpReturnValue %37
- OpFunctionEnd
- %foo = OpFunction %void None %38
- %src_param = OpFunctionParameter %_arr_v4int_uint_4
- %41 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %76
- %46 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- OpStore %46 %53
- %54 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- OpStore %54 %src_param
- %55 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %56 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
- OpStore %55 %56
- %57 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- OpStore %57 %8
- %58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %59 = OpLoad %_arr_v4int_uint_4 %src_function
- OpStore %58 %59
- %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %61 = OpLoad %_arr_v4int_uint_4 %src_private
- OpStore %60 %61
- %62 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %63 = OpLoad %_arr_v4int_uint_4 %src_workgroup
- OpStore %62 %63
- %64 = OpFunctionCall %S %ret_struct_arr
- %65 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %66 = OpCompositeExtract %_arr_v4int_uint_4 %64 0
- OpStore %65 %66
- %67 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %69 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
- %70 = OpLoad %_arr_v4int_uint_4 %69
- OpStore %67 %70
- %71 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 %uint_0
- %72 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
- %73 = OpLoad %_arr_v4int_uint_4 %72
- OpStore %71 %73
- %78 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0 %uint_0
- %79 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
- OpStore %78 %79
+ %main = OpFunction %void None %106
+ %108 = OpLabel
+ %110 = OpLoad %uint %local_invocation_index_1
+ %109 = OpFunctionCall %void %main_inner %110
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.wgsl b/test/tint/array/assign_to_storage_var.wgsl.expected.wgsl
index 8647639..e793a56 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.wgsl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.wgsl
@@ -44,3 +44,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested.arr = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let ary = ArrayType();
+ foo(ary);
+}
diff --git a/test/tint/array/assign_to_subexpr.wgsl b/test/tint/array/assign_to_subexpr.wgsl
index ce4b50a..c020713 100644
--- a/test/tint/array/assign_to_subexpr.wgsl
+++ b/test/tint/array/assign_to_subexpr.wgsl
@@ -1,10 +1,12 @@
+@group(0) @binding(0) var<storage, read_write> s: i32;
+
alias ArrayType = array<i32, 4>;
struct S {
arr : array<i32, 4>,
};
-fn foo() {
+fn foo() -> i32 {
let src : ArrayType = ArrayType();
var dst : ArrayType;
@@ -24,4 +26,11 @@
*dst_ptr = src;
(*dst_struct_ptr).arr = src;
(*dst_array_ptr)[0] = src;
+
+ return (*dst_ptr)[0] + (*dst_struct_ptr).arr[0] + (*dst_array_ptr)[0][0];
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ s = foo();
}
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.dxc.hlsl b/test/tint/array/assign_to_subexpr.wgsl.expected.dxc.hlsl
index 38802ee..eaffe68 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.dxc.hlsl
@@ -1,13 +1,10 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
+RWByteAddressBuffer s : register(u0);
struct S {
int arr[4];
};
-void foo() {
+int foo() {
const int src[4] = (int[4])0;
int tint_symbol[4] = (int[4])0;
S dst_struct = (S)0;
@@ -17,4 +14,11 @@
tint_symbol = src;
dst_struct.arr = src;
dst_array[0] = src;
+ return ((tint_symbol[0] + dst_struct.arr[0]) + dst_array[0][0]);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ s.Store(0u, asuint(foo()));
+ return;
}
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.fxc.hlsl b/test/tint/array/assign_to_subexpr.wgsl.expected.fxc.hlsl
index 38802ee..eaffe68 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.fxc.hlsl
@@ -1,13 +1,10 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
+RWByteAddressBuffer s : register(u0);
struct S {
int arr[4];
};
-void foo() {
+int foo() {
const int src[4] = (int[4])0;
int tint_symbol[4] = (int[4])0;
S dst_struct = (S)0;
@@ -17,4 +14,11 @@
tint_symbol = src;
dst_struct.arr = src;
dst_array[0] = src;
+ return ((tint_symbol[0] + dst_struct.arr[0]) + dst_array[0][0]);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ s.Store(0u, asuint(foo()));
+ return;
}
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.glsl b/test/tint/array/assign_to_subexpr.wgsl.expected.glsl
index 9dbb72b..5532162 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.glsl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.glsl
@@ -1,14 +1,14 @@
#version 310 es
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void unused_entry_point() {
- return;
-}
+layout(binding = 0, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
struct S {
int arr[4];
};
-void foo() {
+int foo() {
int src[4] = int[4](0, 0, 0, 0);
int dst[4] = int[4](0, 0, 0, 0);
S dst_struct = S(int[4](0, 0, 0, 0));
@@ -18,5 +18,15 @@
dst = src;
dst_struct.arr = src;
dst_array[0] = src;
+ return ((dst[0] + dst_struct.arr[0]) + dst_array[0][0]);
}
+void tint_symbol() {
+ s.inner = foo();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl b/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl
index 28251dc..9f8ee81 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl
@@ -1,6 +1,51 @@
SKIP: FAILED
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:247 internal compiler error: Switch() matched no cases. Type: tint::core::ir::Access
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: S = struct @align(4) {
+ arr:array<i32, 4> @offset(0)
+}
+
+%b1 = block { # root
+ %s:ptr<storage, i32, read_write> = var @binding_point(0, 0)
+}
+
+%foo = func():i32 -> %b2 {
+ %b2 = block {
+ %src:array<i32, 4> = let array<i32, 4>(0i)
+ %dst:ptr<function, array<i32, 4>, read_write> = var
+ %dst_struct:ptr<function, S, read_write> = var
+ %dst_array:ptr<function, array<array<i32, 4>, 2>, read_write> = var
+ %dst_ptr:ptr<function, array<i32, 4>, read_write> = let %dst
+ %dst_struct_ptr:ptr<function, S, read_write> = let %dst_struct
+ %dst_array_ptr:ptr<function, array<array<i32, 4>, 2>, read_write> = let %dst_array
+ %10:ptr<function, array<i32, 4>, read_write> = access %dst_struct, 0u
+ store %10, %src
+ %11:ptr<function, array<i32, 4>, read_write> = access %dst_array, 1i
+ store %11, %src
+ store %dst_ptr, %src
+ %12:ptr<function, array<i32, 4>, read_write> = access %dst_struct_ptr, 0u
+ store %12, %src
+ %13:ptr<function, array<i32, 4>, read_write> = access %dst_array_ptr, 0i
+ store %13, %src
+ %14:ptr<function, i32, read_write> = access %dst_ptr, 0i
+ %15:i32 = load %14
+ %16:ptr<function, i32, read_write> = access %dst_struct_ptr, 0u, 0i
+ %17:i32 = load %16
+ %18:i32 = add %15, %17
+ %19:ptr<function, i32, read_write> = access %dst_array_ptr, 0i, 0i
+ %20:i32 = load %19
+ %21:i32 = add %18, %20
+ ret %21
+ }
+}
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func():void -> %b3 {
+ %b3 = block {
+ %23:i32 = call %foo
+ store %s, %23
+ ret
+ }
+}
+
+unhandled variable address space
********************************************************************
* The tint shader compiler has encountered an unexpected error. *
* *
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.msl b/test/tint/array/assign_to_subexpr.wgsl.expected.msl
index 8851fcd..4f791ac 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.msl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.msl
@@ -18,7 +18,7 @@
tint_array<int, 4> arr;
};
-void foo() {
+int foo() {
tint_array<int, 4> const src = tint_array<int, 4>{};
tint_array<int, 4> dst = {};
S dst_struct = {};
@@ -28,5 +28,11 @@
dst = src;
dst_struct.arr = src;
dst_array[0] = src;
+ return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(dst[0]) + as_type<uint>(dst_struct.arr[0])))) + as_type<uint>(dst_array[0][0])));
+}
+
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(0)]]) {
+ *(tint_symbol_1) = foo();
+ return;
}
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.spvasm b/test/tint/array/assign_to_subexpr.wgsl.expected.spvasm
index c6ba3d6..9f5fc07 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.spvasm
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.spvasm
@@ -1,25 +1,34 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 30
+; Bound: 46
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
- OpExecutionMode %unused_entry_point LocalSize 1 1 1
- OpName %unused_entry_point "unused_entry_point"
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %foo "foo"
OpName %dst "dst"
OpName %S "S"
OpMemberName %S 0 "arr"
OpName %dst_struct "dst_struct"
OpName %dst_array "dst_array"
+ OpName %main "main"
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
OpDecorate %_arr_int_uint_4 ArrayStride 4
OpMemberDecorate %S 0 Offset 0
OpDecorate %_arr__arr_int_uint_4_uint_2 ArrayStride 16
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
%int = OpTypeInt 32 1
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %5 = OpTypeFunction %int
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4
@@ -35,12 +44,12 @@
%uint_0 = OpConstant %uint 0
%int_1 = OpConstant %int 1
%28 = OpConstantNull %int
-%unused_entry_point = OpFunction %void None %1
- %4 = OpLabel
- OpReturn
- OpFunctionEnd
- %foo = OpFunction %void None %1
- %6 = OpLabel
+%_ptr_Function_int = OpTypePointer Function %int
+ %void = OpTypeVoid
+ %39 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %foo = OpFunction %int None %5
+ %7 = OpLabel
%dst = OpVariable %_ptr_Function__arr_int_uint_4 Function %11
%dst_struct = OpVariable %_ptr_Function_S Function %17
%dst_array = OpVariable %_ptr_Function__arr__arr_int_uint_4_uint_2 Function %22
@@ -53,5 +62,20 @@
OpStore %27 %11
%29 = OpAccessChain %_ptr_Function__arr_int_uint_4 %dst_array %28
OpStore %29 %11
+ %31 = OpAccessChain %_ptr_Function_int %dst %28
+ %32 = OpLoad %int %31
+ %33 = OpAccessChain %_ptr_Function_int %dst_struct %uint_0 %28
+ %34 = OpLoad %int %33
+ %35 = OpIAdd %int %32 %34
+ %36 = OpAccessChain %_ptr_Function_int %dst_array %28 %28
+ %37 = OpLoad %int %36
+ %38 = OpIAdd %int %35 %37
+ OpReturnValue %38
+ OpFunctionEnd
+ %main = OpFunction %void None %39
+ %42 = OpLabel
+ %44 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %45 = OpFunctionCall %int %foo
+ OpStore %44 %45
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.wgsl b/test/tint/array/assign_to_subexpr.wgsl.expected.wgsl
index 91dc949..645d0b3 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.wgsl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.wgsl
@@ -1,10 +1,12 @@
+@group(0) @binding(0) var<storage, read_write> s : i32;
+
alias ArrayType = array<i32, 4>;
struct S {
arr : array<i32, 4>,
}
-fn foo() {
+fn foo() -> i32 {
let src : ArrayType = ArrayType();
var dst : ArrayType;
var dst_struct : S;
@@ -17,4 +19,10 @@
*(dst_ptr) = src;
(*(dst_struct_ptr)).arr = src;
(*(dst_array_ptr))[0] = src;
+ return (((*(dst_ptr))[0] + (*(dst_struct_ptr)).arr[0]) + (*(dst_array_ptr))[0][0]);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ s = foo();
}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl b/test/tint/array/assign_to_workgroup_var.wgsl
index 74f3e67..356ca0e 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl
@@ -50,3 +50,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let val = ArrayType();
+ foo(val);
+}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.dxc.hlsl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.dxc.hlsl
index 08c9b20..9393b1c 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.dxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -18,22 +13,22 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
typedef int4 src_uniform_load_ret[4];
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- const uint scalar_offset = ((offset + (i * 16u))) / 4;
- arr_1[i] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_4 = 0u; (i_4 < 4u); i_4 = (i_4 + 1u)) {
+ const uint scalar_offset = ((offset + (i_4 * 16u))) / 4;
+ arr_1[i_4] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -43,8 +38,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- arr_2[i_1] = asint(src_storage.Load4((offset + (i_1 * 16u))));
+ for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
+ arr_2[i_5] = asint(src_storage.Load4((offset + (i_5 * 16u))));
}
}
return arr_2;
@@ -52,8 +47,8 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol = tint_symbol_4;
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol = tint_symbol_6;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@@ -68,3 +63,34 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ tint_symbol[i] = (0).xxxx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ {
+ for(uint idx_1 = local_invocation_index; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
+ const uint i_1 = (idx_1 / 6u);
+ const uint i_2 = ((idx_1 % 6u) / 2u);
+ const uint i_3 = (idx_1 % 2u);
+ dst_nested[i_1][i_2][i_3] = 0;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 val[4] = (int4[4])0;
+ foo(val);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.fxc.hlsl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.fxc.hlsl
index 08c9b20..9393b1c 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.fxc.hlsl
@@ -1,8 +1,3 @@
-[numthreads(1, 1, 1)]
-void unused_entry_point() {
- return;
-}
-
struct S {
int4 arr[4];
};
@@ -18,22 +13,22 @@
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
- const int4 tint_symbol_2[4] = (int4[4])0;
- return tint_symbol_2;
+ const int4 tint_symbol_4[4] = (int4[4])0;
+ return tint_symbol_4;
}
S ret_struct_arr() {
- const S tint_symbol_3 = (S)0;
- return tint_symbol_3;
+ const S tint_symbol_5 = (S)0;
+ return tint_symbol_5;
}
typedef int4 src_uniform_load_ret[4];
src_uniform_load_ret src_uniform_load(uint offset) {
int4 arr_1[4] = (int4[4])0;
{
- for(uint i = 0u; (i < 4u); i = (i + 1u)) {
- const uint scalar_offset = ((offset + (i * 16u))) / 4;
- arr_1[i] = asint(src_uniform[scalar_offset / 4]);
+ for(uint i_4 = 0u; (i_4 < 4u); i_4 = (i_4 + 1u)) {
+ const uint scalar_offset = ((offset + (i_4 * 16u))) / 4;
+ arr_1[i_4] = asint(src_uniform[scalar_offset / 4]);
}
}
return arr_1;
@@ -43,8 +38,8 @@
src_storage_load_ret src_storage_load(uint offset) {
int4 arr_2[4] = (int4[4])0;
{
- for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
- arr_2[i_1] = asint(src_storage.Load4((offset + (i_1 * 16u))));
+ for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
+ arr_2[i_5] = asint(src_storage.Load4((offset + (i_5 * 16u))));
}
}
return arr_2;
@@ -52,8 +47,8 @@
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
- const int4 tint_symbol_4[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
- tint_symbol = tint_symbol_4;
+ const int4 tint_symbol_6[4] = {(1).xxxx, (2).xxxx, (3).xxxx, (3).xxxx};
+ tint_symbol = tint_symbol_6;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@@ -68,3 +63,34 @@
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}
+
+struct tint_symbol_3 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void main_inner(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ tint_symbol[i] = (0).xxxx;
+ src_workgroup[i] = (0).xxxx;
+ }
+ }
+ {
+ for(uint idx_1 = local_invocation_index; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
+ const uint i_1 = (idx_1 / 6u);
+ const uint i_2 = ((idx_1 % 6u) / 2u);
+ const uint i_3 = (idx_1 % 2u);
+ dst_nested[i_1][i_2][i_3] = 0;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ const int4 val[4] = (int4[4])0;
+ foo(val);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_3 tint_symbol_2) {
+ main_inner(tint_symbol_2.local_invocation_index);
+ return;
+}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.glsl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.glsl
index ba09877..eb8f310 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.glsl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void unused_entry_point() {
- return;
-}
struct S {
ivec4 arr[4];
};
@@ -21,19 +17,19 @@
shared ivec4 dst[4];
shared int dst_nested[4][3][2];
ivec4[4] ret_arr() {
- ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
- return tint_symbol_1;
+ ivec4 tint_symbol_2[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ return tint_symbol_2;
}
S ret_struct_arr() {
- S tint_symbol_2 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
- return tint_symbol_2;
+ S tint_symbol_3 = S(ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0)));
+ return tint_symbol_3;
}
void foo(ivec4 src_param[4]) {
ivec4 src_function[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
- ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
- dst = tint_symbol_3;
+ ivec4 tint_symbol_4[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
+ dst = tint_symbol_4;
dst = src_param;
dst = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
@@ -41,11 +37,37 @@
dst = src_function;
dst = src_private;
dst = src_workgroup;
- S tint_symbol = ret_struct_arr();
- dst = tint_symbol.arr;
+ S tint_symbol_1 = ret_struct_arr();
+ dst = tint_symbol_1.arr;
dst = src_uniform.inner.arr;
dst = src_storage.inner.arr;
int src_nested[4][3][2] = int[4][3][2](int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)), int[3][2](int[2](0, 0), int[2](0, 0), int[2](0, 0)));
dst_nested = src_nested;
}
+void tint_symbol(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint i = idx;
+ dst[i] = ivec4(0);
+ src_workgroup[i] = ivec4(0);
+ }
+ }
+ {
+ for(uint idx_1 = local_invocation_index; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
+ uint i_1 = (idx_1 / 6u);
+ uint i_2 = ((idx_1 % 6u) / 2u);
+ uint i_3 = (idx_1 % 2u);
+ dst_nested[i_1][i_2][i_3] = 0;
+ }
+ }
+ barrier();
+ ivec4 val[4] = ivec4[4](ivec4(0), ivec4(0), ivec4(0), ivec4(0));
+ foo(val);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
index 1f6127b..6daea82 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
@@ -1,6 +1,6 @@
SKIP: FAILED
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: S = struct @align(16) {
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: S = struct @align(16) {
arr:array<vec4<i32>, 4> @offset(0)
}
@@ -53,6 +53,60 @@
ret
}
}
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void -> %b5 {
+ %b5 = block {
+ loop [i: %b6, b: %b7, c: %b8] { # loop_1
+ %b6 = block { # initializer
+ next_iteration %b7 %tint_local_index
+ }
+ %b7 = block (%idx:u32) { # body
+ %28:bool = gte %idx:u32, 4u
+ if %28 [t: %b9] { # if_1
+ %b9 = block { # true
+ exit_loop # loop_1
+ }
+ }
+ %29:ptr<workgroup, vec4<i32>, read_write> = access %src_workgroup, %idx:u32
+ store %29, vec4<i32>(0i)
+ %30:ptr<workgroup, vec4<i32>, read_write> = access %dst, %idx:u32
+ store %30, vec4<i32>(0i)
+ continue %b8
+ }
+ %b8 = block { # continuing
+ %31:u32 = add %idx:u32, 1u
+ next_iteration %b7 %31
+ }
+ }
+ loop [i: %b10, b: %b11, c: %b12] { # loop_2
+ %b10 = block { # initializer
+ next_iteration %b11 %tint_local_index
+ }
+ %b11 = block (%idx_1:u32) { # body
+ %33:bool = gte %idx_1:u32, 24u
+ if %33 [t: %b13] { # if_2
+ %b13 = block { # true
+ exit_loop # loop_2
+ }
+ }
+ %34:u32 = mod %idx_1:u32, 2u
+ %35:u32 = div %idx_1:u32, 2u
+ %36:u32 = mod %35, 3u
+ %37:u32 = div %idx_1:u32, 6u
+ %38:ptr<workgroup, i32, read_write> = access %dst_nested, %37, %36, %34
+ store %38, 0i
+ continue %b12
+ }
+ %b12 = block { # continuing
+ %39:u32 = add %idx_1:u32, 1u
+ next_iteration %b11 %39
+ }
+ }
+ %40:void = msl.threadgroup_barrier 4u
+ %val:array<vec4<i32>, 4> = let array<vec4<i32>, 4>(vec4<i32>(0i))
+ %42:void = call %foo, %val
+ ret
+ }
+}
unhandled variable address space
********************************************************************
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
index a6a6a5a..3ad0e3b 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
@@ -23,31 +23,57 @@
};
tint_array<int4, 4> ret_arr() {
- tint_array<int4, 4> const tint_symbol_1 = tint_array<int4, 4>{};
- return tint_symbol_1;
-}
-
-S ret_struct_arr() {
- S const tint_symbol_2 = S{};
+ tint_array<int4, 4> const tint_symbol_2 = tint_array<int4, 4>{};
return tint_symbol_2;
}
-void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_8) {
+S ret_struct_arr() {
+ S const tint_symbol_3 = S{};
+ return tint_symbol_3;
+}
+
+void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_5, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_9) {
tint_array<int4, 4> src_function = {};
- tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
- *(tint_symbol_4) = tint_symbol_3;
- *(tint_symbol_4) = src_param;
- *(tint_symbol_4) = ret_arr();
+ tint_array<int4, 4> const tint_symbol_4 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
+ *(tint_symbol_5) = tint_symbol_4;
+ *(tint_symbol_5) = src_param;
+ *(tint_symbol_5) = ret_arr();
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
- *(tint_symbol_4) = src_let;
- *(tint_symbol_4) = src_function;
- *(tint_symbol_4) = (*(tint_private_vars)).src_private;
- *(tint_symbol_4) = *(tint_symbol_5);
- S const tint_symbol = ret_struct_arr();
- *(tint_symbol_4) = tint_symbol.arr;
- *(tint_symbol_4) = (*(tint_symbol_6)).arr;
- *(tint_symbol_4) = (*(tint_symbol_7)).arr;
+ *(tint_symbol_5) = src_let;
+ *(tint_symbol_5) = src_function;
+ *(tint_symbol_5) = (*(tint_private_vars)).src_private;
+ *(tint_symbol_5) = *(tint_symbol_6);
+ S const tint_symbol_1 = ret_struct_arr();
+ *(tint_symbol_5) = tint_symbol_1.arr;
+ *(tint_symbol_5) = (*(tint_symbol_7)).arr;
+ *(tint_symbol_5) = (*(tint_symbol_8)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
- *(tint_symbol_8) = src_nested;
+ *(tint_symbol_9) = src_nested;
+}
+
+void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_10, threadgroup tint_array<int4, 4>* const tint_symbol_11, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_12, const constant S* const tint_symbol_13, device S* const tint_symbol_14) {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint const i = idx;
+ (*(tint_symbol_10))[i] = int4(0);
+ (*(tint_symbol_11))[i] = int4(0);
+ }
+ for(uint idx_1 = local_invocation_index; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
+ uint const i_1 = (idx_1 / 6u);
+ uint const i_2 = ((idx_1 % 6u) / 2u);
+ uint const i_3 = (idx_1 % 2u);
+ (*(tint_symbol_12))[i_1][i_2][i_3] = 0;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ tint_array<int4, 4> const val = tint_array<int4, 4>{};
+ foo(val, tint_private_vars, tint_symbol_10, tint_symbol_11, tint_symbol_13, tint_symbol_14, tint_symbol_12);
+}
+
+kernel void tint_symbol(const constant S* tint_symbol_18 [[buffer(0)]], device S* tint_symbol_19 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ threadgroup tint_array<int4, 4> tint_symbol_15;
+ threadgroup tint_array<int4, 4> tint_symbol_16;
+ threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4> tint_symbol_17;
+ tint_symbol_inner(local_invocation_index, &(tint_private_vars), &(tint_symbol_15), &(tint_symbol_16), &(tint_symbol_17), tint_symbol_18, tint_symbol_19);
+ return;
}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.spvasm b/test/tint/array/assign_to_workgroup_var.wgsl.expected.spvasm
index 85496ae..bcdc29a 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.spvasm
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.spvasm
@@ -1,12 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 66
+; Bound: 122
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
- OpExecutionMode %unused_entry_point LocalSize 1 1 1
+ OpEntryPoint GLCompute %main "main" %local_invocation_index_1
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %src_private "src_private"
OpName %src_workgroup "src_workgroup"
OpName %src_uniform_block "src_uniform_block"
@@ -17,13 +18,18 @@
OpName %src_storage "src_storage"
OpName %dst "dst"
OpName %dst_nested "dst_nested"
- OpName %unused_entry_point "unused_entry_point"
OpName %ret_arr "ret_arr"
OpName %ret_struct_arr "ret_struct_arr"
OpName %foo "foo"
OpName %src_param "src_param"
OpName %src_function "src_function"
OpName %src_nested "src_nested"
+ OpName %main_inner "main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %idx_1 "idx_1"
+ OpName %main "main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %src_uniform_block Block
OpMemberDecorate %src_uniform_block 0 Offset 0
@@ -36,14 +42,16 @@
OpDecorate %_arr_int_uint_2 ArrayStride 4
OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
%int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
- %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
- %8 = OpConstantNull %_arr_v4int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+ %10 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %10
%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_v4int_uint_4
@@ -60,63 +68,142 @@
%_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
%_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Workgroup %_arr__arr__arr_int_uint_2_uint_3_uint_4
%dst_nested = OpVariable %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 Workgroup
+ %27 = OpTypeFunction %_arr_v4int_uint_4
+ %30 = OpTypeFunction %S
+ %33 = OpConstantNull %S
%void = OpTypeVoid
- %25 = OpTypeFunction %void
- %29 = OpTypeFunction %_arr_v4int_uint_4
- %32 = OpTypeFunction %S
- %35 = OpConstantNull %S
- %36 = OpTypeFunction %void %_arr_v4int_uint_4
+ %34 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%int_1 = OpConstant %int 1
- %43 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
+ %42 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2
- %45 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
+ %44 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3
- %47 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
- %48 = OpConstantComposite %_arr_v4int_uint_4 %43 %45 %47 %47
+ %46 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+ %47 = OpConstantComposite %_arr_v4int_uint_4 %42 %44 %46 %46
%uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
- %64 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %25
- %28 = OpLabel
+ %63 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+ %65 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %71 = OpConstantNull %uint
+ %bool = OpTypeBool
+%_ptr_Workgroup_v4int = OpTypePointer Workgroup %v4int
+ %85 = OpConstantNull %v4int
+ %uint_1 = OpConstant %uint 1
+ %uint_24 = OpConstant %uint 24
+ %uint_6 = OpConstant %uint 6
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %111 = OpConstantNull %int
+ %uint_264 = OpConstant %uint 264
+ %117 = OpTypeFunction %void
+ %ret_arr = OpFunction %_arr_v4int_uint_4 None %27
+ %29 = OpLabel
+ OpReturnValue %10
+ OpFunctionEnd
+%ret_struct_arr = OpFunction %S None %30
+ %32 = OpLabel
+ OpReturnValue %33
+ OpFunctionEnd
+ %foo = OpFunction %void None %34
+ %src_param = OpFunctionParameter %_arr_v4int_uint_4
+ %38 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %10
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %63
+ OpStore %dst %47
+ OpStore %dst %src_param
+ %48 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+ OpStore %dst %48
+ OpStore %dst %10
+ %49 = OpLoad %_arr_v4int_uint_4 %src_function
+ OpStore %dst %49
+ %50 = OpLoad %_arr_v4int_uint_4 %src_private
+ OpStore %dst %50
+ %51 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+ OpStore %dst %51
+ %52 = OpFunctionCall %S %ret_struct_arr
+ %53 = OpCompositeExtract %_arr_v4int_uint_4 %52 0
+ OpStore %dst %53
+ %56 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
+ %57 = OpLoad %_arr_v4int_uint_4 %56
+ OpStore %dst %57
+ %59 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
+ %60 = OpLoad %_arr_v4int_uint_4 %59
+ OpStore %dst %60
+ %64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+ OpStore %dst_nested %64
OpReturn
OpFunctionEnd
- %ret_arr = OpFunction %_arr_v4int_uint_4 None %29
- %31 = OpLabel
- OpReturnValue %8
+ %main_inner = OpFunction %void None %65
+%local_invocation_index = OpFunctionParameter %uint
+ %68 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %71
+ %idx_1 = OpVariable %_ptr_Function_uint Function %71
+ OpStore %idx %local_invocation_index
+ OpBranch %72
+ %72 = OpLabel
+ OpLoopMerge %73 %74 None
+ OpBranch %75
+ %75 = OpLabel
+ %77 = OpLoad %uint %idx
+ %78 = OpULessThan %bool %77 %uint_4
+ %76 = OpLogicalNot %bool %78
+ OpSelectionMerge %80 None
+ OpBranchConditional %76 %81 %80
+ %81 = OpLabel
+ OpBranch %73
+ %80 = OpLabel
+ %82 = OpLoad %uint %idx
+ %84 = OpAccessChain %_ptr_Workgroup_v4int %dst %82
+ OpStore %84 %85
+ %86 = OpAccessChain %_ptr_Workgroup_v4int %src_workgroup %82
+ OpStore %86 %85
+ OpBranch %74
+ %74 = OpLabel
+ %87 = OpLoad %uint %idx
+ %89 = OpIAdd %uint %87 %uint_1
+ OpStore %idx %89
+ OpBranch %72
+ %73 = OpLabel
+ OpStore %idx_1 %local_invocation_index
+ OpBranch %91
+ %91 = OpLabel
+ OpLoopMerge %92 %93 None
+ OpBranch %94
+ %94 = OpLabel
+ %96 = OpLoad %uint %idx_1
+ %98 = OpULessThan %bool %96 %uint_24
+ %95 = OpLogicalNot %bool %98
+ OpSelectionMerge %99 None
+ OpBranchConditional %95 %100 %99
+ %100 = OpLabel
+ OpBranch %92
+ %99 = OpLabel
+ %101 = OpLoad %uint %idx_1
+ %103 = OpUDiv %uint %101 %uint_6
+ %104 = OpLoad %uint %idx_1
+ %105 = OpUMod %uint %104 %uint_6
+ %106 = OpUDiv %uint %105 %uint_2
+ %107 = OpLoad %uint %idx_1
+ %108 = OpUMod %uint %107 %uint_2
+ %110 = OpAccessChain %_ptr_Workgroup_int %dst_nested %103 %106 %108
+ OpStore %110 %111
+ OpBranch %93
+ %93 = OpLabel
+ %112 = OpLoad %uint %idx_1
+ %113 = OpIAdd %uint %112 %uint_1
+ OpStore %idx_1 %113
+ OpBranch %91
+ %92 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %116 = OpFunctionCall %void %foo %10
+ OpReturn
OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %32
- %34 = OpLabel
- OpReturnValue %35
- OpFunctionEnd
- %foo = OpFunction %void None %36
- %src_param = OpFunctionParameter %_arr_v4int_uint_4
- %39 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %64
- OpStore %dst %48
- OpStore %dst %src_param
- %49 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
- OpStore %dst %49
- OpStore %dst %8
- %50 = OpLoad %_arr_v4int_uint_4 %src_function
- OpStore %dst %50
- %51 = OpLoad %_arr_v4int_uint_4 %src_private
- OpStore %dst %51
- %52 = OpLoad %_arr_v4int_uint_4 %src_workgroup
- OpStore %dst %52
- %53 = OpFunctionCall %S %ret_struct_arr
- %54 = OpCompositeExtract %_arr_v4int_uint_4 %53 0
- OpStore %dst %54
- %57 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 %uint_0
- %58 = OpLoad %_arr_v4int_uint_4 %57
- OpStore %dst %58
- %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 %uint_0
- %61 = OpLoad %_arr_v4int_uint_4 %60
- OpStore %dst %61
- %65 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
- OpStore %dst_nested %65
+ %main = OpFunction %void None %117
+ %119 = OpLabel
+ %121 = OpLoad %uint %local_invocation_index_1
+ %120 = OpFunctionCall %void %main_inner %121
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.wgsl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.wgsl
index 301b212..f430134 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.wgsl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.wgsl
@@ -40,3 +40,9 @@
var src_nested : array<array<array<i32, 2>, 3>, 4>;
dst_nested = src_nested;
}
+
+@compute @workgroup_size(1)
+fn main() {
+ let val = ArrayType();
+ foo(val);
+}
diff --git a/test/tint/array/function_parameter.wgsl b/test/tint/array/function_parameter.wgsl
index be06c51..d39a43d 100644
--- a/test/tint/array/function_parameter.wgsl
+++ b/test/tint/array/function_parameter.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s: f32;
+
fn f1(a : array<f32, 4>) -> f32 {
return a[3];
}
@@ -18,4 +20,6 @@
let v1 : f32 = f1(a1);
let v2 : f32 = f2(a2);
let v3 : f32 = f3(a3);
+
+ s = v1 + v2 + v3;
}
diff --git a/test/tint/array/function_parameter.wgsl.expected.dxc.hlsl b/test/tint/array/function_parameter.wgsl.expected.dxc.hlsl
index b2e7f41..3d959e8 100644
--- a/test/tint/array/function_parameter.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/function_parameter.wgsl.expected.dxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
float f1(float a[4]) {
return a[3];
}
@@ -18,5 +20,6 @@
const float v1 = f1(a1);
const float v2 = f2(a2);
const float v3 = f3(a3);
+ s.Store(0u, asuint(((v1 + v2) + v3)));
return;
}
diff --git a/test/tint/array/function_parameter.wgsl.expected.fxc.hlsl b/test/tint/array/function_parameter.wgsl.expected.fxc.hlsl
index b2e7f41..3d959e8 100644
--- a/test/tint/array/function_parameter.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/function_parameter.wgsl.expected.fxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
float f1(float a[4]) {
return a[3];
}
@@ -18,5 +20,6 @@
const float v1 = f1(a1);
const float v2 = f2(a2);
const float v3 = f3(a3);
+ s.Store(0u, asuint(((v1 + v2) + v3)));
return;
}
diff --git a/test/tint/array/function_parameter.wgsl.expected.glsl b/test/tint/array/function_parameter.wgsl.expected.glsl
index 6b9fe79..661101e 100644
--- a/test/tint/array/function_parameter.wgsl.expected.glsl
+++ b/test/tint/array/function_parameter.wgsl.expected.glsl
@@ -1,5 +1,9 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ float inner;
+} s;
+
float f1(float a[4]) {
return a[3];
}
@@ -19,6 +23,7 @@
float v1 = f1(a1);
float v2 = f2(a2);
float v3 = f3(a3);
+ s.inner = ((v1 + v2) + v3);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/array/function_parameter.wgsl.expected.ir.msl b/test/tint/array/function_parameter.wgsl.expected.ir.msl
index 238a512..c241422 100644
--- a/test/tint/array/function_parameter.wgsl.expected.ir.msl
+++ b/test/tint/array/function_parameter.wgsl.expected.ir.msl
@@ -1,32 +1,49 @@
-#include <metal_stdlib>
-using namespace metal;
-template<typename T, size_t N>
-struct tint_array {
- const constant T& operator[](size_t i) const constant { return elements[i]; }
- device T& operator[](size_t i) device { return elements[i]; }
- const device T& operator[](size_t i) const device { return elements[i]; }
- thread T& operator[](size_t i) thread { return elements[i]; }
- const thread T& operator[](size_t i) const thread { return elements[i]; }
- threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
- const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
- T elements[N];
-};
+SKIP: FAILED
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: %b1 = block { # root
+ %s:ptr<storage, f32, read_write> = var @binding_point(0, 0)
+}
-float f1(tint_array<float, 4> a) {
- return a[3];
+%f1 = func(%a:array<f32, 4>):f32 -> %b2 {
+ %b2 = block {
+ %4:f32 = access %a, 3i
+ ret %4
+ }
}
-float f2(tint_array<tint_array<float, 4>, 3> a) {
- return a[2][3];
+%f2 = func(%a_1:array<array<f32, 4>, 3>):f32 -> %b3 { # %a_1: 'a'
+ %b3 = block {
+ %7:f32 = access %a_1, 2i, 3i
+ ret %7
+ }
}
-float f3(tint_array<tint_array<tint_array<float, 4>, 3>, 2> a) {
- return a[1][2][3];
+%f3 = func(%a_2:array<array<array<f32, 4>, 3>, 2>):f32 -> %b4 { # %a_2: 'a'
+ %b4 = block {
+ %10:f32 = access %a_2, 1i, 2i, 3i
+ ret %10
+ }
}
-kernel void tint_symbol() {
- tint_array<float, 4> const a1 = tint_array<float, 4>{};
- tint_array<tint_array<float, 4>, 3> const a2 = tint_array<tint_array<float, 4>, 3>{};
- tint_array<tint_array<tint_array<float, 4>, 3>, 2> const a3 = tint_array<tint_array<tint_array<float, 4>, 3>, 2>{};
- float const v1 = f1(a1);
- float const v2 = f2(a2);
- float const v3 = f3(a3);
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func():void -> %b5 {
+ %b5 = block {
+ %a1:array<f32, 4> = let array<f32, 4>(0.0f)
+ %a2:array<array<f32, 4>, 3> = let array<array<f32, 4>, 3>(array<f32, 4>(0.0f))
+ %a3:array<array<array<f32, 4>, 3>, 2> = let array<array<array<f32, 4>, 3>, 2>(array<array<f32, 4>, 3>(array<f32, 4>(0.0f)))
+ %15:f32 = call %f1, %a1
+ %v1:f32 = let %15
+ %17:f32 = call %f2, %a2
+ %v2:f32 = let %17
+ %19:f32 = call %f3, %a3
+ %v3:f32 = let %19
+ %21:f32 = add %v1, %v2
+ %22:f32 = add %21, %v3
+ store %s, %22
+ ret
+ }
}
+
+unhandled variable address space
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/array/function_parameter.wgsl.expected.msl b/test/tint/array/function_parameter.wgsl.expected.msl
index bf68a61..f66c5e4 100644
--- a/test/tint/array/function_parameter.wgsl.expected.msl
+++ b/test/tint/array/function_parameter.wgsl.expected.msl
@@ -26,13 +26,14 @@
return a[1][2][3];
}
-kernel void tint_symbol() {
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(0)]]) {
tint_array<float, 4> const a1 = tint_array<float, 4>{};
tint_array<tint_array<float, 4>, 3> const a2 = tint_array<tint_array<float, 4>, 3>{};
tint_array<tint_array<tint_array<float, 4>, 3>, 2> const a3 = tint_array<tint_array<tint_array<float, 4>, 3>, 2>{};
float const v1 = f1(a1);
float const v2 = f2(a2);
float const v3 = f3(a3);
+ *(tint_symbol_1) = ((v1 + v2) + v3);
return;
}
diff --git a/test/tint/array/function_parameter.wgsl.expected.spvasm b/test/tint/array/function_parameter.wgsl.expected.spvasm
index 406d4b0..90962fa 100644
--- a/test/tint/array/function_parameter.wgsl.expected.spvasm
+++ b/test/tint/array/function_parameter.wgsl.expected.spvasm
@@ -1,12 +1,15 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 41
+; Bound: 49
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %f1 "f1"
OpName %a "a"
OpName %f2 "f2"
@@ -14,54 +17,67 @@
OpName %f3 "f3"
OpName %a_1 "a"
OpName %main "main"
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
OpDecorate %_arr_float_uint_4 ArrayStride 4
OpDecorate %_arr__arr_float_uint_4_uint_3 ArrayStride 16
OpDecorate %_arr__arr__arr_float_uint_4_uint_3_uint_2 ArrayStride 48
%float = OpTypeFloat 32
+ %s_block = OpTypeStruct %float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_float_uint_4 = OpTypeArray %float %uint_4
- %1 = OpTypeFunction %float %_arr_float_uint_4
+ %5 = OpTypeFunction %float %_arr_float_uint_4
%int = OpTypeInt 32 1
%int_3 = OpConstant %int 3
%uint_3 = OpConstant %uint 3
%_arr__arr_float_uint_4_uint_3 = OpTypeArray %_arr_float_uint_4 %uint_3
- %12 = OpTypeFunction %float %_arr__arr_float_uint_4_uint_3
+ %15 = OpTypeFunction %float %_arr__arr_float_uint_4_uint_3
%int_2 = OpConstant %int 2
%uint_2 = OpConstant %uint 2
%_arr__arr__arr_float_uint_4_uint_3_uint_2 = OpTypeArray %_arr__arr_float_uint_4_uint_3 %uint_2
- %21 = OpTypeFunction %float %_arr__arr__arr_float_uint_4_uint_3_uint_2
+ %24 = OpTypeFunction %float %_arr__arr__arr_float_uint_4_uint_3_uint_2
%int_1 = OpConstant %int 1
%void = OpTypeVoid
- %31 = OpTypeFunction %void
- %35 = OpConstantNull %_arr_float_uint_4
- %36 = OpConstantNull %_arr__arr_float_uint_4_uint_3
- %37 = OpConstantNull %_arr__arr__arr_float_uint_4_uint_3_uint_2
- %f1 = OpFunction %float None %1
+ %34 = OpTypeFunction %void
+ %38 = OpConstantNull %_arr_float_uint_4
+ %39 = OpConstantNull %_arr__arr_float_uint_4_uint_3
+ %40 = OpConstantNull %_arr__arr__arr_float_uint_4_uint_3_uint_2
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %f1 = OpFunction %float None %5
%a = OpFunctionParameter %_arr_float_uint_4
- %8 = OpLabel
- %11 = OpCompositeExtract %float %a 3
- OpReturnValue %11
+ %11 = OpLabel
+ %14 = OpCompositeExtract %float %a 3
+ OpReturnValue %14
OpFunctionEnd
- %f2 = OpFunction %float None %12
+ %f2 = OpFunction %float None %15
%a_0 = OpFunctionParameter %_arr__arr_float_uint_4_uint_3
- %17 = OpLabel
- %19 = OpCompositeExtract %_arr_float_uint_4 %a_0 2
- %20 = OpCompositeExtract %float %19 3
- OpReturnValue %20
+ %20 = OpLabel
+ %22 = OpCompositeExtract %_arr_float_uint_4 %a_0 2
+ %23 = OpCompositeExtract %float %22 3
+ OpReturnValue %23
OpFunctionEnd
- %f3 = OpFunction %float None %21
+ %f3 = OpFunction %float None %24
%a_1 = OpFunctionParameter %_arr__arr__arr_float_uint_4_uint_3_uint_2
- %26 = OpLabel
- %28 = OpCompositeExtract %_arr__arr_float_uint_4_uint_3 %a_1 1
- %29 = OpCompositeExtract %_arr_float_uint_4 %28 2
- %30 = OpCompositeExtract %float %29 3
- OpReturnValue %30
+ %29 = OpLabel
+ %31 = OpCompositeExtract %_arr__arr_float_uint_4_uint_3 %a_1 1
+ %32 = OpCompositeExtract %_arr_float_uint_4 %31 2
+ %33 = OpCompositeExtract %float %32 3
+ OpReturnValue %33
OpFunctionEnd
- %main = OpFunction %void None %31
- %34 = OpLabel
- %38 = OpFunctionCall %float %f1 %35
- %39 = OpFunctionCall %float %f2 %36
- %40 = OpFunctionCall %float %f3 %37
+ %main = OpFunction %void None %34
+ %37 = OpLabel
+ %41 = OpFunctionCall %float %f1 %38
+ %42 = OpFunctionCall %float %f2 %39
+ %43 = OpFunctionCall %float %f3 %40
+ %46 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0
+ %47 = OpFAdd %float %41 %42
+ %48 = OpFAdd %float %47 %43
+ OpStore %46 %48
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/function_parameter.wgsl.expected.wgsl b/test/tint/array/function_parameter.wgsl.expected.wgsl
index be06c51..9305ffe 100644
--- a/test/tint/array/function_parameter.wgsl.expected.wgsl
+++ b/test/tint/array/function_parameter.wgsl.expected.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s : f32;
+
fn f1(a : array<f32, 4>) -> f32 {
return a[3];
}
@@ -18,4 +20,5 @@
let v1 : f32 = f1(a1);
let v2 : f32 = f2(a2);
let v3 : f32 = f3(a3);
+ s = ((v1 + v2) + v3);
}
diff --git a/test/tint/array/function_return_type.wgsl b/test/tint/array/function_return_type.wgsl
index 4b5dddc..2b5f9e7 100644
--- a/test/tint/array/function_return_type.wgsl
+++ b/test/tint/array/function_return_type.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s: f32;
+
fn f1() -> array<f32, 4> {
return array<f32, 4>();
}
@@ -15,4 +17,6 @@
let a1 : array<f32, 4> = f1();
let a2 : array<array<f32, 4>, 3> = f2();
let a3 : array<array<array<f32, 4>, 3>, 2> = f3();
+
+ s = a1[0] + a2[0][0] + a3[0][0][0];
}
diff --git a/test/tint/array/function_return_type.wgsl.expected.dxc.hlsl b/test/tint/array/function_return_type.wgsl.expected.dxc.hlsl
index 954b060..db96ae5 100644
--- a/test/tint/array/function_return_type.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/function_return_type.wgsl.expected.dxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
typedef float f1_ret[4];
f1_ret f1() {
const float tint_symbol_5[4] = (float[4])0;
@@ -26,5 +28,6 @@
const float a1[4] = f1();
const float a2[3][4] = f2();
const float a3[2][3][4] = f3();
+ s.Store(0u, asuint(((a1[0] + a2[0][0]) + a3[0][0][0])));
return;
}
diff --git a/test/tint/array/function_return_type.wgsl.expected.fxc.hlsl b/test/tint/array/function_return_type.wgsl.expected.fxc.hlsl
index 954b060..db96ae5 100644
--- a/test/tint/array/function_return_type.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/function_return_type.wgsl.expected.fxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
typedef float f1_ret[4];
f1_ret f1() {
const float tint_symbol_5[4] = (float[4])0;
@@ -26,5 +28,6 @@
const float a1[4] = f1();
const float a2[3][4] = f2();
const float a3[2][3][4] = f3();
+ s.Store(0u, asuint(((a1[0] + a2[0][0]) + a3[0][0][0])));
return;
}
diff --git a/test/tint/array/function_return_type.wgsl.expected.glsl b/test/tint/array/function_return_type.wgsl.expected.glsl
index 497ca18..f716ca7 100644
--- a/test/tint/array/function_return_type.wgsl.expected.glsl
+++ b/test/tint/array/function_return_type.wgsl.expected.glsl
@@ -1,5 +1,9 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ float inner;
+} s;
+
float[4] f1() {
float tint_symbol_6[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
return tint_symbol_6;
@@ -24,6 +28,7 @@
float a1[4] = f1();
float a2[3][4] = f2();
float a3[2][3][4] = f3();
+ s.inner = ((a1[0] + a2[0][0]) + a3[0][0][0]);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/array/function_return_type.wgsl.expected.ir.msl b/test/tint/array/function_return_type.wgsl.expected.ir.msl
index 440724d..ffb2175 100644
--- a/test/tint/array/function_return_type.wgsl.expected.ir.msl
+++ b/test/tint/array/function_return_type.wgsl.expected.ir.msl
@@ -1,32 +1,56 @@
-#include <metal_stdlib>
-using namespace metal;
-template<typename T, size_t N>
-struct tint_array {
- const constant T& operator[](size_t i) const constant { return elements[i]; }
- device T& operator[](size_t i) device { return elements[i]; }
- const device T& operator[](size_t i) const device { return elements[i]; }
- thread T& operator[](size_t i) thread { return elements[i]; }
- const thread T& operator[](size_t i) const thread { return elements[i]; }
- threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
- const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
- T elements[N];
-};
+SKIP: FAILED
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: %b1 = block { # root
+ %s:ptr<storage, f32, read_write> = var @binding_point(0, 0)
+}
-tint_array<float, 4> f1() {
- return tint_array<float, 4>{};
+%f1 = func():array<f32, 4> -> %b2 {
+ %b2 = block {
+ ret array<f32, 4>(0.0f)
+ }
}
-tint_array<tint_array<float, 4>, 3> f2() {
- tint_array<float, 4> const v = f1();
- tint_array<float, 4> const v_1 = f1();
- return tint_array<tint_array<float, 4>, 3>{v, v_1, f1()};
+%f2 = func():array<array<f32, 4>, 3> -> %b3 {
+ %b3 = block {
+ %4:array<f32, 4> = call %f1
+ %5:array<f32, 4> = let %4
+ %6:array<f32, 4> = call %f1
+ %7:array<f32, 4> = let %6
+ %8:array<f32, 4> = call %f1
+ %9:array<array<f32, 4>, 3> = construct %5, %7, %8
+ ret %9
+ }
}
-tint_array<tint_array<tint_array<float, 4>, 3>, 2> f3() {
- tint_array<tint_array<float, 4>, 3> const v_2 = f2();
- return tint_array<tint_array<tint_array<float, 4>, 3>, 2>{v_2, f2()};
+%f3 = func():array<array<array<f32, 4>, 3>, 2> -> %b4 {
+ %b4 = block {
+ %11:array<array<f32, 4>, 3> = call %f2
+ %12:array<array<f32, 4>, 3> = let %11
+ %13:array<array<f32, 4>, 3> = call %f2
+ %14:array<array<array<f32, 4>, 3>, 2> = construct %12, %13
+ ret %14
+ }
}
-kernel void tint_symbol() {
- tint_array<float, 4> const a1 = f1();
- tint_array<tint_array<float, 4>, 3> const a2 = f2();
- tint_array<tint_array<tint_array<float, 4>, 3>, 2> const a3 = f3();
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func():void -> %b5 {
+ %b5 = block {
+ %16:array<f32, 4> = call %f1
+ %a1:array<f32, 4> = let %16
+ %18:array<array<f32, 4>, 3> = call %f2
+ %a2:array<array<f32, 4>, 3> = let %18
+ %20:array<array<array<f32, 4>, 3>, 2> = call %f3
+ %a3:array<array<array<f32, 4>, 3>, 2> = let %20
+ %22:f32 = access %a1, 0i
+ %23:f32 = access %a2, 0i, 0i
+ %24:f32 = add %22, %23
+ %25:f32 = access %a3, 0i, 0i, 0i
+ %26:f32 = add %24, %25
+ store %s, %26
+ ret
+ }
}
+
+unhandled variable address space
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/array/function_return_type.wgsl.expected.msl b/test/tint/array/function_return_type.wgsl.expected.msl
index 703fc8867..793db60 100644
--- a/test/tint/array/function_return_type.wgsl.expected.msl
+++ b/test/tint/array/function_return_type.wgsl.expected.msl
@@ -34,10 +34,11 @@
return tint_symbol_8;
}
-kernel void tint_symbol() {
+kernel void tint_symbol(device float* tint_symbol_9 [[buffer(0)]]) {
tint_array<float, 4> const a1 = f1();
tint_array<tint_array<float, 4>, 3> const a2 = f2();
tint_array<tint_array<tint_array<float, 4>, 3>, 2> const a3 = f3();
+ *(tint_symbol_9) = ((a1[0] + a2[0][0]) + a3[0][0][0]);
return;
}
diff --git a/test/tint/array/function_return_type.wgsl.expected.spvasm b/test/tint/array/function_return_type.wgsl.expected.spvasm
index a1a6115..690ce05 100644
--- a/test/tint/array/function_return_type.wgsl.expected.spvasm
+++ b/test/tint/array/function_return_type.wgsl.expected.spvasm
@@ -1,56 +1,80 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 49
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %f1 "f1"
OpName %f2 "f2"
OpName %f3 "f3"
OpName %main "main"
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
OpDecorate %_arr_float_uint_4 ArrayStride 4
OpDecorate %_arr__arr_float_uint_4_uint_3 ArrayStride 16
OpDecorate %_arr__arr__arr_float_uint_4_uint_3_uint_2 ArrayStride 48
%float = OpTypeFloat 32
+ %s_block = OpTypeStruct %float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_float_uint_4 = OpTypeArray %float %uint_4
- %1 = OpTypeFunction %_arr_float_uint_4
- %8 = OpConstantNull %_arr_float_uint_4
+ %5 = OpTypeFunction %_arr_float_uint_4
+ %11 = OpConstantNull %_arr_float_uint_4
%uint_3 = OpConstant %uint 3
%_arr__arr_float_uint_4_uint_3 = OpTypeArray %_arr_float_uint_4 %uint_3
- %9 = OpTypeFunction %_arr__arr_float_uint_4_uint_3
+ %12 = OpTypeFunction %_arr__arr_float_uint_4_uint_3
%uint_2 = OpConstant %uint 2
%_arr__arr__arr_float_uint_4_uint_3_uint_2 = OpTypeArray %_arr__arr_float_uint_4_uint_3 %uint_2
- %18 = OpTypeFunction %_arr__arr__arr_float_uint_4_uint_3_uint_2
+ %21 = OpTypeFunction %_arr__arr__arr_float_uint_4_uint_3_uint_2
%void = OpTypeVoid
- %26 = OpTypeFunction %void
- %f1 = OpFunction %_arr_float_uint_4 None %1
- %7 = OpLabel
- OpReturnValue %8
+ %29 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %int = OpTypeInt 32 1
+ %40 = OpConstantNull %int
+ %f1 = OpFunction %_arr_float_uint_4 None %5
+ %10 = OpLabel
+ OpReturnValue %11
OpFunctionEnd
- %f2 = OpFunction %_arr__arr_float_uint_4_uint_3 None %9
- %13 = OpLabel
- %14 = OpFunctionCall %_arr_float_uint_4 %f1
- %15 = OpFunctionCall %_arr_float_uint_4 %f1
- %16 = OpFunctionCall %_arr_float_uint_4 %f1
- %17 = OpCompositeConstruct %_arr__arr_float_uint_4_uint_3 %14 %15 %16
- OpReturnValue %17
+ %f2 = OpFunction %_arr__arr_float_uint_4_uint_3 None %12
+ %16 = OpLabel
+ %17 = OpFunctionCall %_arr_float_uint_4 %f1
+ %18 = OpFunctionCall %_arr_float_uint_4 %f1
+ %19 = OpFunctionCall %_arr_float_uint_4 %f1
+ %20 = OpCompositeConstruct %_arr__arr_float_uint_4_uint_3 %17 %18 %19
+ OpReturnValue %20
OpFunctionEnd
- %f3 = OpFunction %_arr__arr__arr_float_uint_4_uint_3_uint_2 None %18
- %22 = OpLabel
- %23 = OpFunctionCall %_arr__arr_float_uint_4_uint_3 %f2
- %24 = OpFunctionCall %_arr__arr_float_uint_4_uint_3 %f2
- %25 = OpCompositeConstruct %_arr__arr__arr_float_uint_4_uint_3_uint_2 %23 %24
- OpReturnValue %25
+ %f3 = OpFunction %_arr__arr__arr_float_uint_4_uint_3_uint_2 None %21
+ %25 = OpLabel
+ %26 = OpFunctionCall %_arr__arr_float_uint_4_uint_3 %f2
+ %27 = OpFunctionCall %_arr__arr_float_uint_4_uint_3 %f2
+ %28 = OpCompositeConstruct %_arr__arr__arr_float_uint_4_uint_3_uint_2 %26 %27
+ OpReturnValue %28
OpFunctionEnd
- %main = OpFunction %void None %26
- %29 = OpLabel
- %30 = OpFunctionCall %_arr_float_uint_4 %f1
- %31 = OpFunctionCall %_arr__arr_float_uint_4_uint_3 %f2
- %32 = OpFunctionCall %_arr__arr__arr_float_uint_4_uint_3_uint_2 %f3
+ %main = OpFunction %void None %29
+ %32 = OpLabel
+ %33 = OpFunctionCall %_arr_float_uint_4 %f1
+ %34 = OpFunctionCall %_arr__arr_float_uint_4_uint_3 %f2
+ %35 = OpFunctionCall %_arr__arr__arr_float_uint_4_uint_3_uint_2 %f3
+ %38 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0
+ %41 = OpCompositeExtract %float %33 0
+ %42 = OpCompositeExtract %_arr_float_uint_4 %34 0
+ %43 = OpCompositeExtract %float %42 0
+ %44 = OpFAdd %float %41 %43
+ %45 = OpCompositeExtract %_arr__arr_float_uint_4_uint_3 %35 0
+ %46 = OpCompositeExtract %_arr_float_uint_4 %45 0
+ %47 = OpCompositeExtract %float %46 0
+ %48 = OpFAdd %float %44 %47
+ OpStore %38 %48
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/function_return_type.wgsl.expected.wgsl b/test/tint/array/function_return_type.wgsl.expected.wgsl
index 4b5dddc..03cd1b7 100644
--- a/test/tint/array/function_return_type.wgsl.expected.wgsl
+++ b/test/tint/array/function_return_type.wgsl.expected.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s : f32;
+
fn f1() -> array<f32, 4> {
return array<f32, 4>();
}
@@ -15,4 +17,5 @@
let a1 : array<f32, 4> = f1();
let a2 : array<array<f32, 4>, 3> = f2();
let a3 : array<array<array<f32, 4>, 3>, 2> = f3();
+ s = ((a1[0] + a2[0][0]) + a3[0][0][0]);
}
diff --git a/test/tint/array/size.wgsl b/test/tint/array/size.wgsl
index 4a7c5ee..ac028bd 100644
--- a/test/tint/array/size.wgsl
+++ b/test/tint/array/size.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s: f32;
+
const slen = 4;
const ulen = 4u;
@@ -14,4 +16,7 @@
signed_constant = signed_literal;
unsigned_constant = signed_literal;
shr_const_expr = signed_literal;
+
+ s = signed_literal[0] + unsigned_literal[0] + signed_constant[0] + unsigned_constant[0] +
+ shr_const_expr[0];
}
diff --git a/test/tint/array/size.wgsl.expected.dxc.hlsl b/test/tint/array/size.wgsl.expected.dxc.hlsl
index 34f5448..3ccb283 100644
--- a/test/tint/array/size.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/size.wgsl.expected.dxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
void main() {
float signed_literal[4] = (float[4])0;
float unsigned_literal[4] = (float[4])0;
@@ -8,5 +10,6 @@
signed_constant = signed_literal;
unsigned_constant = signed_literal;
shr_const_expr = signed_literal;
+ s.Store(0u, asuint(((((signed_literal[0] + unsigned_literal[0]) + signed_constant[0]) + unsigned_constant[0]) + shr_const_expr[0])));
return;
}
diff --git a/test/tint/array/size.wgsl.expected.fxc.hlsl b/test/tint/array/size.wgsl.expected.fxc.hlsl
index 34f5448..3ccb283 100644
--- a/test/tint/array/size.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/size.wgsl.expected.fxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
void main() {
float signed_literal[4] = (float[4])0;
float unsigned_literal[4] = (float[4])0;
@@ -8,5 +10,6 @@
signed_constant = signed_literal;
unsigned_constant = signed_literal;
shr_const_expr = signed_literal;
+ s.Store(0u, asuint(((((signed_literal[0] + unsigned_literal[0]) + signed_constant[0]) + unsigned_constant[0]) + shr_const_expr[0])));
return;
}
diff --git a/test/tint/array/size.wgsl.expected.glsl b/test/tint/array/size.wgsl.expected.glsl
index 00d6381..2628e2a 100644
--- a/test/tint/array/size.wgsl.expected.glsl
+++ b/test/tint/array/size.wgsl.expected.glsl
@@ -1,6 +1,10 @@
#version 310 es
precision highp float;
+layout(binding = 0, std430) buffer s_block_ssbo {
+ float inner;
+} s;
+
void tint_symbol() {
float signed_literal[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
float unsigned_literal[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
@@ -11,6 +15,7 @@
signed_constant = signed_literal;
unsigned_constant = signed_literal;
shr_const_expr = signed_literal;
+ s.inner = ((((signed_literal[0] + unsigned_literal[0]) + signed_constant[0]) + unsigned_constant[0]) + shr_const_expr[0]);
}
void main() {
diff --git a/test/tint/array/size.wgsl.expected.ir.msl b/test/tint/array/size.wgsl.expected.ir.msl
index 0d242e4..c60b0f6 100644
--- a/test/tint/array/size.wgsl.expected.ir.msl
+++ b/test/tint/array/size.wgsl.expected.ir.msl
@@ -1,26 +1,47 @@
-#include <metal_stdlib>
-using namespace metal;
-template<typename T, size_t N>
-struct tint_array {
- const constant T& operator[](size_t i) const constant { return elements[i]; }
- device T& operator[](size_t i) device { return elements[i]; }
- const device T& operator[](size_t i) const device { return elements[i]; }
- thread T& operator[](size_t i) thread { return elements[i]; }
- const thread T& operator[](size_t i) const thread { return elements[i]; }
- threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
- const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
- T elements[N];
-};
+SKIP: FAILED
-
-fragment void tint_symbol() {
- tint_array<float, 4> signed_literal = {};
- tint_array<float, 4> unsigned_literal = {};
- tint_array<float, 4> signed_constant = {};
- tint_array<float, 4> unsigned_constant = {};
- tint_array<float, 4> shr_const_expr = {};
- unsigned_literal = signed_literal;
- signed_constant = signed_literal;
- unsigned_constant = signed_literal;
- shr_const_expr = signed_literal;
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: %b1 = block { # root
+ %s:ptr<storage, f32, read_write> = var @binding_point(0, 0)
}
+
+%tint_symbol = @fragment func():void -> %b2 {
+ %b2 = block {
+ %signed_literal:ptr<function, array<f32, 4>, read_write> = var
+ %unsigned_literal:ptr<function, array<f32, 4>, read_write> = var
+ %signed_constant:ptr<function, array<f32, 4>, read_write> = var
+ %unsigned_constant:ptr<function, array<f32, 4>, read_write> = var
+ %shr_const_expr:ptr<function, array<f32, 4>, read_write> = var
+ %8:array<f32, 4> = load %signed_literal
+ store %unsigned_literal, %8
+ %9:array<f32, 4> = load %signed_literal
+ store %signed_constant, %9
+ %10:array<f32, 4> = load %signed_literal
+ store %unsigned_constant, %10
+ %11:array<f32, 4> = load %signed_literal
+ store %shr_const_expr, %11
+ %12:ptr<function, f32, read_write> = access %signed_literal, 0i
+ %13:f32 = load %12
+ %14:ptr<function, f32, read_write> = access %unsigned_literal, 0i
+ %15:f32 = load %14
+ %16:f32 = add %13, %15
+ %17:ptr<function, f32, read_write> = access %signed_constant, 0i
+ %18:f32 = load %17
+ %19:f32 = add %16, %18
+ %20:ptr<function, f32, read_write> = access %unsigned_constant, 0i
+ %21:f32 = load %20
+ %22:f32 = add %19, %21
+ %23:ptr<function, f32, read_write> = access %shr_const_expr, 0i
+ %24:f32 = load %23
+ %25:f32 = add %22, %24
+ store %s, %25
+ ret
+ }
+}
+
+unhandled variable address space
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/array/size.wgsl.expected.msl b/test/tint/array/size.wgsl.expected.msl
index f3277f2..0b2bd02 100644
--- a/test/tint/array/size.wgsl.expected.msl
+++ b/test/tint/array/size.wgsl.expected.msl
@@ -14,7 +14,7 @@
T elements[N];
};
-fragment void tint_symbol() {
+fragment void tint_symbol(device float* tint_symbol_1 [[buffer(0)]]) {
tint_array<float, 4> signed_literal = {};
tint_array<float, 4> unsigned_literal = {};
tint_array<float, 4> signed_constant = {};
@@ -24,6 +24,7 @@
signed_constant = signed_literal;
unsigned_constant = signed_literal;
shr_const_expr = signed_literal;
+ *(tint_symbol_1) = ((((signed_literal[0] + unsigned_literal[0]) + signed_constant[0]) + unsigned_constant[0]) + shr_const_expr[0]);
return;
}
diff --git a/test/tint/array/size.wgsl.expected.spvasm b/test/tint/array/size.wgsl.expected.spvasm
index 2bc675d..aa25b8f 100644
--- a/test/tint/array/size.wgsl.expected.spvasm
+++ b/test/tint/array/size.wgsl.expected.spvasm
@@ -1,41 +1,72 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 20
+; Bound: 43
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main"
OpExecutionMode %main OriginUpperLeft
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %main "main"
OpName %signed_literal "signed_literal"
OpName %unsigned_literal "unsigned_literal"
OpName %signed_constant "signed_constant"
OpName %unsigned_constant "unsigned_constant"
OpName %shr_const_expr "shr_const_expr"
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
OpDecorate %_arr_float_uint_4 ArrayStride 4
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
%float = OpTypeFloat 32
+ %s_block = OpTypeStruct %float
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_float_uint_4 = OpTypeArray %float %uint_4
%_ptr_Function__arr_float_uint_4 = OpTypePointer Function %_arr_float_uint_4
- %11 = OpConstantNull %_arr_float_uint_4
- %main = OpFunction %void None %1
- %4 = OpLabel
-%signed_literal = OpVariable %_ptr_Function__arr_float_uint_4 Function %11
-%unsigned_literal = OpVariable %_ptr_Function__arr_float_uint_4 Function %11
-%signed_constant = OpVariable %_ptr_Function__arr_float_uint_4 Function %11
-%unsigned_constant = OpVariable %_ptr_Function__arr_float_uint_4 Function %11
-%shr_const_expr = OpVariable %_ptr_Function__arr_float_uint_4 Function %11
- %16 = OpLoad %_arr_float_uint_4 %signed_literal
- OpStore %unsigned_literal %16
- %17 = OpLoad %_arr_float_uint_4 %signed_literal
- OpStore %signed_constant %17
- %18 = OpLoad %_arr_float_uint_4 %signed_literal
- OpStore %unsigned_constant %18
+ %14 = OpConstantNull %_arr_float_uint_4
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %int = OpTypeInt 32 1
+ %27 = OpConstantNull %int
+%_ptr_Function_float = OpTypePointer Function %float
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+%signed_literal = OpVariable %_ptr_Function__arr_float_uint_4 Function %14
+%unsigned_literal = OpVariable %_ptr_Function__arr_float_uint_4 Function %14
+%signed_constant = OpVariable %_ptr_Function__arr_float_uint_4 Function %14
+%unsigned_constant = OpVariable %_ptr_Function__arr_float_uint_4 Function %14
+%shr_const_expr = OpVariable %_ptr_Function__arr_float_uint_4 Function %14
%19 = OpLoad %_arr_float_uint_4 %signed_literal
- OpStore %shr_const_expr %19
+ OpStore %unsigned_literal %19
+ %20 = OpLoad %_arr_float_uint_4 %signed_literal
+ OpStore %signed_constant %20
+ %21 = OpLoad %_arr_float_uint_4 %signed_literal
+ OpStore %unsigned_constant %21
+ %22 = OpLoad %_arr_float_uint_4 %signed_literal
+ OpStore %shr_const_expr %22
+ %25 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0
+ %29 = OpAccessChain %_ptr_Function_float %signed_literal %27
+ %30 = OpLoad %float %29
+ %31 = OpAccessChain %_ptr_Function_float %unsigned_literal %27
+ %32 = OpLoad %float %31
+ %33 = OpFAdd %float %30 %32
+ %34 = OpAccessChain %_ptr_Function_float %signed_constant %27
+ %35 = OpLoad %float %34
+ %36 = OpFAdd %float %33 %35
+ %37 = OpAccessChain %_ptr_Function_float %unsigned_constant %27
+ %38 = OpLoad %float %37
+ %39 = OpFAdd %float %36 %38
+ %40 = OpAccessChain %_ptr_Function_float %shr_const_expr %27
+ %41 = OpLoad %float %40
+ %42 = OpFAdd %float %39 %41
+ OpStore %25 %42
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/size.wgsl.expected.wgsl b/test/tint/array/size.wgsl.expected.wgsl
index f9eee82..011829b 100644
--- a/test/tint/array/size.wgsl.expected.wgsl
+++ b/test/tint/array/size.wgsl.expected.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s : f32;
+
const slen = 4;
const ulen = 4u;
@@ -13,4 +15,5 @@
signed_constant = signed_literal;
unsigned_constant = signed_literal;
shr_const_expr = signed_literal;
+ s = ((((signed_literal[0] + unsigned_literal[0]) + signed_constant[0]) + unsigned_constant[0]) + shr_const_expr[0]);
}
diff --git a/test/tint/array/type_initializer.wgsl b/test/tint/array/type_initializer.wgsl
index e975089..2397240 100644
--- a/test/tint/array/type_initializer.wgsl
+++ b/test/tint/array/type_initializer.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s: i32;
+
@compute @workgroup_size(1)
fn main() {
let x : i32 = 42;
@@ -45,4 +47,10 @@
array<i32, 4>(1, x, x + 1, nonempty[3]),
nested_nonempty[1][2],
)[1];
+
+
+ s = empty[0] + nonempty[0] + nonempty_with_expr[0] + nested_empty[0][0][0] +
+ nested_nonempty[0][0][0] + nested_nonempty_with_expr[0][0][0] +
+ subexpr_empty + subexpr_nonempty + subexpr_nonempty_with_expr + subexpr_nested_empty[0] +
+ subexpr_nested_nonempty[0] + subexpr_nested_nonempty_with_expr[0];
}
diff --git a/test/tint/array/type_initializer.wgsl.expected.dxc.hlsl b/test/tint/array/type_initializer.wgsl.expected.dxc.hlsl
index eb59a88..b43a887 100644
--- a/test/tint/array/type_initializer.wgsl.expected.dxc.hlsl
+++ b/test/tint/array/type_initializer.wgsl.expected.dxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
const int x = 42;
@@ -19,5 +21,6 @@
const int tint_symbol_4[4] = {1, x, (x + 1), nonempty[3]};
const int tint_symbol_5[2][4] = {tint_symbol_4, nested_nonempty[1][2]};
const int subexpr_nested_nonempty_with_expr[4] = tint_symbol_5[1];
+ s.Store(0u, asuint((((((((((((empty[0] + nonempty[0]) + nonempty_with_expr[0]) + nested_empty[0][0][0]) + nested_nonempty[0][0][0]) + nested_nonempty_with_expr[0][0][0]) + subexpr_empty) + subexpr_nonempty) + subexpr_nonempty_with_expr) + subexpr_nested_empty[0]) + subexpr_nested_nonempty[0]) + subexpr_nested_nonempty_with_expr[0])));
return;
}
diff --git a/test/tint/array/type_initializer.wgsl.expected.fxc.hlsl b/test/tint/array/type_initializer.wgsl.expected.fxc.hlsl
index eb59a88..b43a887 100644
--- a/test/tint/array/type_initializer.wgsl.expected.fxc.hlsl
+++ b/test/tint/array/type_initializer.wgsl.expected.fxc.hlsl
@@ -1,3 +1,5 @@
+RWByteAddressBuffer s : register(u0);
+
[numthreads(1, 1, 1)]
void main() {
const int x = 42;
@@ -19,5 +21,6 @@
const int tint_symbol_4[4] = {1, x, (x + 1), nonempty[3]};
const int tint_symbol_5[2][4] = {tint_symbol_4, nested_nonempty[1][2]};
const int subexpr_nested_nonempty_with_expr[4] = tint_symbol_5[1];
+ s.Store(0u, asuint((((((((((((empty[0] + nonempty[0]) + nonempty_with_expr[0]) + nested_empty[0][0][0]) + nested_nonempty[0][0][0]) + nested_nonempty_with_expr[0][0][0]) + subexpr_empty) + subexpr_nonempty) + subexpr_nonempty_with_expr) + subexpr_nested_empty[0]) + subexpr_nested_nonempty[0]) + subexpr_nested_nonempty_with_expr[0])));
return;
}
diff --git a/test/tint/array/type_initializer.wgsl.expected.glsl b/test/tint/array/type_initializer.wgsl.expected.glsl
index 65a79d4..d02dfc0 100644
--- a/test/tint/array/type_initializer.wgsl.expected.glsl
+++ b/test/tint/array/type_initializer.wgsl.expected.glsl
@@ -1,5 +1,9 @@
#version 310 es
+layout(binding = 0, std430) buffer s_block_ssbo {
+ int inner;
+} s;
+
void tint_symbol() {
int x = 42;
int empty[4] = int[4](0, 0, 0, 0);
@@ -20,6 +24,7 @@
int tint_symbol_5[4] = int[4](1, x, (x + 1), nonempty[3]);
int tint_symbol_6[2][4] = int[2][4](tint_symbol_5, nested_nonempty[1][2]);
int subexpr_nested_nonempty_with_expr[4] = tint_symbol_6[1];
+ s.inner = (((((((((((empty[0] + nonempty[0]) + nonempty_with_expr[0]) + nested_empty[0][0][0]) + nested_nonempty[0][0][0]) + nested_nonempty_with_expr[0][0][0]) + subexpr_empty) + subexpr_nonempty) + subexpr_nonempty_with_expr) + subexpr_nested_empty[0]) + subexpr_nested_nonempty[0]) + subexpr_nested_nonempty_with_expr[0]);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/test/tint/array/type_initializer.wgsl.expected.ir.msl b/test/tint/array/type_initializer.wgsl.expected.ir.msl
index 8b89e3d..e00cec5 100644
--- a/test/tint/array/type_initializer.wgsl.expected.ir.msl
+++ b/test/tint/array/type_initializer.wgsl.expected.ir.msl
@@ -1,31 +1,76 @@
-#include <metal_stdlib>
-using namespace metal;
-template<typename T, size_t N>
-struct tint_array {
- const constant T& operator[](size_t i) const constant { return elements[i]; }
- device T& operator[](size_t i) device { return elements[i]; }
- const device T& operator[](size_t i) const device { return elements[i]; }
- thread T& operator[](size_t i) thread { return elements[i]; }
- const thread T& operator[](size_t i) const thread { return elements[i]; }
- threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
- const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
- T elements[N];
-};
+SKIP: FAILED
-
-kernel void tint_symbol() {
- int const x = 42;
- tint_array<int, 4> const empty = tint_array<int, 4>{};
- tint_array<int, 4> const nonempty = tint_array<int, 4>{1, 2, 3, 4};
- tint_array<int, 4> const nonempty_with_expr = tint_array<int, 4>{1, x, (x + 1), nonempty[3]};
- tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_empty = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{};
- tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_array<tint_array<int, 4>, 3>{tint_array<int, 4>{1, 2, 3, 4}, tint_array<int, 4>{5, 6, 7, 8}, tint_array<int, 4>{9, 10, 11, 12}}, tint_array<tint_array<int, 4>, 3>{tint_array<int, 4>{13, 14, 15, 16}, tint_array<int, 4>{17, 18, 19, 20}, tint_array<int, 4>{21, 22, 23, 24}}};
- tint_array<int, 4> const v = tint_array<int, 4>{1, 2, x, (x + 1)};
- tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty_with_expr = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_array<tint_array<int, 4>, 3>{v, tint_array<int, 4>{5, 6, nonempty[2], (nonempty[3] + 1)}, nonempty}, nested_nonempty[1]};
- int const subexpr_empty = 0;
- int const subexpr_nonempty = 3;
- int const subexpr_nonempty_with_expr = tint_array<int, 4>{1, x, (x + 1), nonempty[3]}[2];
- tint_array<int, 4> const subexpr_nested_empty = tint_array<int, 4>{};
- tint_array<int, 4> const subexpr_nested_nonempty = tint_array<int, 4>{5, 6, 7, 8};
- tint_array<int, 4> const subexpr_nested_nonempty_with_expr = tint_array<tint_array<int, 4>, 2>{tint_array<int, 4>{1, x, (x + 1), nonempty[3]}, nested_nonempty[1][2]}[1];
+<dawn>/src/tint/lang/msl/writer/printer/printer.cc:493 internal compiler error: %b1 = block { # root
+ %s:ptr<storage, i32, read_write> = var @binding_point(0, 0)
}
+
+%tint_symbol = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+ %b2 = block {
+ %x:i32 = let 42i
+ %empty:array<i32, 4> = let array<i32, 4>(0i)
+ %nonempty:array<i32, 4> = let array<i32, 4>(1i, 2i, 3i, 4i)
+ %6:i32 = add %x, 1i
+ %7:i32 = access %nonempty, 3i
+ %8:array<i32, 4> = construct 1i, %x, %6, %7
+ %nonempty_with_expr:array<i32, 4> = let %8
+ %nested_empty:array<array<array<i32, 4>, 3>, 2> = let array<array<array<i32, 4>, 3>, 2>(array<array<i32, 4>, 3>(array<i32, 4>(0i)))
+ %nested_nonempty:array<array<array<i32, 4>, 3>, 2> = let array<array<array<i32, 4>, 3>, 2>(array<array<i32, 4>, 3>(array<i32, 4>(1i, 2i, 3i, 4i), array<i32, 4>(5i, 6i, 7i, 8i), array<i32, 4>(9i, 10i, 11i, 12i)), array<array<i32, 4>, 3>(array<i32, 4>(13i, 14i, 15i, 16i), array<i32, 4>(17i, 18i, 19i, 20i), array<i32, 4>(21i, 22i, 23i, 24i)))
+ %12:i32 = add %x, 1i
+ %13:array<i32, 4> = construct 1i, 2i, %x, %12
+ %14:array<i32, 4> = let %13
+ %15:i32 = access %nonempty, 2i
+ %16:i32 = access %nonempty, 3i
+ %17:i32 = add %16, 1i
+ %18:array<i32, 4> = construct 5i, 6i, %15, %17
+ %19:array<array<i32, 4>, 3> = construct %14, %18, %nonempty
+ %20:array<array<i32, 4>, 3> = access %nested_nonempty, 1i
+ %21:array<array<array<i32, 4>, 3>, 2> = construct %19, %20
+ %nested_nonempty_with_expr:array<array<array<i32, 4>, 3>, 2> = let %21
+ %subexpr_empty:i32 = let 0i
+ %subexpr_nonempty:i32 = let 3i
+ %25:i32 = add %x, 1i
+ %26:i32 = access %nonempty, 3i
+ %27:array<i32, 4> = construct 1i, %x, %25, %26
+ %28:i32 = access %27, 2i
+ %subexpr_nonempty_with_expr:i32 = let %28
+ %subexpr_nested_empty:array<i32, 4> = let array<i32, 4>(0i)
+ %subexpr_nested_nonempty:array<i32, 4> = let array<i32, 4>(5i, 6i, 7i, 8i)
+ %32:i32 = add %x, 1i
+ %33:i32 = access %nonempty, 3i
+ %34:array<i32, 4> = construct 1i, %x, %32, %33
+ %35:array<i32, 4> = access %nested_nonempty, 1i, 2i
+ %36:array<array<i32, 4>, 2> = construct %34, %35
+ %37:array<i32, 4> = access %36, 1i
+ %subexpr_nested_nonempty_with_expr:array<i32, 4> = let %37
+ %39:i32 = access %empty, 0i
+ %40:i32 = access %nonempty, 0i
+ %41:i32 = add %39, %40
+ %42:i32 = access %nonempty_with_expr, 0i
+ %43:i32 = add %41, %42
+ %44:i32 = access %nested_empty, 0i, 0i, 0i
+ %45:i32 = add %43, %44
+ %46:i32 = access %nested_nonempty, 0i, 0i, 0i
+ %47:i32 = add %45, %46
+ %48:i32 = access %nested_nonempty_with_expr, 0i, 0i, 0i
+ %49:i32 = add %47, %48
+ %50:i32 = add %49, %subexpr_empty
+ %51:i32 = add %50, %subexpr_nonempty
+ %52:i32 = add %51, %subexpr_nonempty_with_expr
+ %53:i32 = access %subexpr_nested_empty, 0i
+ %54:i32 = add %52, %53
+ %55:i32 = access %subexpr_nested_nonempty, 0i
+ %56:i32 = add %54, %55
+ %57:i32 = access %subexpr_nested_nonempty_with_expr, 0i
+ %58:i32 = add %56, %57
+ store %s, %58
+ ret
+ }
+}
+
+unhandled variable address space
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/array/type_initializer.wgsl.expected.msl b/test/tint/array/type_initializer.wgsl.expected.msl
index 266d4b1..6ead166 100644
--- a/test/tint/array/type_initializer.wgsl.expected.msl
+++ b/test/tint/array/type_initializer.wgsl.expected.msl
@@ -14,7 +14,7 @@
T elements[N];
};
-kernel void tint_symbol() {
+kernel void tint_symbol(device int* tint_symbol_7 [[buffer(0)]]) {
int const x = 42;
tint_array<int, 4> const empty = tint_array<int, 4>{};
tint_array<int, 4> const nonempty = tint_array<int, 4>{1, 2, 3, 4};
@@ -34,6 +34,7 @@
tint_array<int, 4> const tint_symbol_5 = tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty[3]};
tint_array<tint_array<int, 4>, 2> const tint_symbol_6 = tint_array<tint_array<int, 4>, 2>{tint_symbol_5, nested_nonempty[1][2]};
tint_array<int, 4> const subexpr_nested_nonempty_with_expr = tint_symbol_6[1];
+ *(tint_symbol_7) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(empty[0]) + as_type<uint>(nonempty[0])))) + as_type<uint>(nonempty_with_expr[0])))) + as_type<uint>(nested_empty[0][0][0])))) + as_type<uint>(nested_nonempty[0][0][0])))) + as_type<uint>(nested_nonempty_with_expr[0][0][0])))) + as_type<uint>(subexpr_empty)))) + as_type<uint>(subexpr_nonempty)))) + as_type<uint>(subexpr_nonempty_with_expr)))) + as_type<uint>(subexpr_nested_empty[0])))) + as_type<uint>(subexpr_nested_nonempty[0])))) + as_type<uint>(subexpr_nested_nonempty_with_expr[0])));
return;
}
diff --git a/test/tint/array/type_initializer.wgsl.expected.spvasm b/test/tint/array/type_initializer.wgsl.expected.spvasm
index 5d4084d..d63b08f 100644
--- a/test/tint/array/type_initializer.wgsl.expected.spvasm
+++ b/test/tint/array/type_initializer.wgsl.expected.spvasm
@@ -1,89 +1,129 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 74
+; Bound: 106
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
+ OpName %s_block "s_block"
+ OpMemberName %s_block 0 "inner"
+ OpName %s "s"
OpName %main "main"
+ OpDecorate %s_block Block
+ OpMemberDecorate %s_block 0 Offset 0
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
OpDecorate %_arr_int_uint_4 ArrayStride 4
OpDecorate %_arr__arr_int_uint_4_uint_3 ArrayStride 16
OpDecorate %_arr__arr__arr_int_uint_4_uint_3_uint_2 ArrayStride 48
OpDecorate %_arr__arr_int_uint_4_uint_2 ArrayStride 16
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
%int = OpTypeInt 32 1
+ %s_block = OpTypeStruct %int
+%_ptr_StorageBuffer_s_block = OpTypePointer StorageBuffer %s_block
+ %s = OpVariable %_ptr_StorageBuffer_s_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
%int_42 = OpConstant %int 42
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4
- %10 = OpConstantNull %_arr_int_uint_4
+ %13 = OpConstantNull %_arr_int_uint_4
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%int_4 = OpConstant %int 4
- %15 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_4
+ %18 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_4
%uint_3 = OpConstant %uint 3
%_arr__arr_int_uint_4_uint_3 = OpTypeArray %_arr_int_uint_4 %uint_3
%uint_2 = OpConstant %uint 2
%_arr__arr__arr_int_uint_4_uint_3_uint_2 = OpTypeArray %_arr__arr_int_uint_4_uint_3 %uint_2
- %23 = OpConstantNull %_arr__arr__arr_int_uint_4_uint_3_uint_2
+ %26 = OpConstantNull %_arr__arr__arr_int_uint_4_uint_3_uint_2
%int_5 = OpConstant %int 5
%int_6 = OpConstant %int 6
%int_7 = OpConstant %int 7
%int_8 = OpConstant %int 8
- %28 = OpConstantComposite %_arr_int_uint_4 %int_5 %int_6 %int_7 %int_8
+ %31 = OpConstantComposite %_arr_int_uint_4 %int_5 %int_6 %int_7 %int_8
%int_9 = OpConstant %int 9
%int_10 = OpConstant %int 10
%int_11 = OpConstant %int 11
%int_12 = OpConstant %int 12
- %33 = OpConstantComposite %_arr_int_uint_4 %int_9 %int_10 %int_11 %int_12
- %34 = OpConstantComposite %_arr__arr_int_uint_4_uint_3 %15 %28 %33
+ %36 = OpConstantComposite %_arr_int_uint_4 %int_9 %int_10 %int_11 %int_12
+ %37 = OpConstantComposite %_arr__arr_int_uint_4_uint_3 %18 %31 %36
%int_13 = OpConstant %int 13
%int_14 = OpConstant %int 14
%int_15 = OpConstant %int 15
%int_16 = OpConstant %int 16
- %39 = OpConstantComposite %_arr_int_uint_4 %int_13 %int_14 %int_15 %int_16
+ %42 = OpConstantComposite %_arr_int_uint_4 %int_13 %int_14 %int_15 %int_16
%int_17 = OpConstant %int 17
%int_18 = OpConstant %int 18
%int_19 = OpConstant %int 19
%int_20 = OpConstant %int 20
- %44 = OpConstantComposite %_arr_int_uint_4 %int_17 %int_18 %int_19 %int_20
+ %47 = OpConstantComposite %_arr_int_uint_4 %int_17 %int_18 %int_19 %int_20
%int_21 = OpConstant %int 21
%int_22 = OpConstant %int 22
%int_23 = OpConstant %int 23
%int_24 = OpConstant %int 24
- %49 = OpConstantComposite %_arr_int_uint_4 %int_21 %int_22 %int_23 %int_24
- %50 = OpConstantComposite %_arr__arr_int_uint_4_uint_3 %39 %44 %49
- %51 = OpConstantComposite %_arr__arr__arr_int_uint_4_uint_3_uint_2 %34 %50
- %61 = OpConstantNull %int
+ %52 = OpConstantComposite %_arr_int_uint_4 %int_21 %int_22 %int_23 %int_24
+ %53 = OpConstantComposite %_arr__arr_int_uint_4_uint_3 %42 %47 %52
+ %54 = OpConstantComposite %_arr__arr__arr_int_uint_4_uint_3_uint_2 %37 %53
+ %64 = OpConstantNull %int
%_arr__arr_int_uint_4_uint_2 = OpTypeArray %_arr_int_uint_4 %uint_2
- %main = OpFunction %void None %1
- %4 = OpLabel
- %16 = OpIAdd %int %int_42 %int_1
- %17 = OpCompositeExtract %int %15 3
- %18 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_42 %16 %17
- %52 = OpIAdd %int %int_42 %int_1
- %53 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_2 %int_42 %52
- %54 = OpCompositeExtract %int %15 2
- %55 = OpCompositeExtract %int %15 3
- %56 = OpIAdd %int %55 %int_1
- %57 = OpCompositeConstruct %_arr_int_uint_4 %int_5 %int_6 %54 %56
- %58 = OpCompositeConstruct %_arr__arr_int_uint_4_uint_3 %53 %57 %15
- %59 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %51 1
- %60 = OpCompositeConstruct %_arr__arr__arr_int_uint_4_uint_3_uint_2 %58 %59
- %62 = OpIAdd %int %int_42 %int_1
- %63 = OpCompositeExtract %int %15 3
- %64 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_42 %62 %63
- %65 = OpCompositeExtract %int %64 2
- %67 = OpIAdd %int %int_42 %int_1
- %68 = OpCompositeExtract %int %15 3
- %69 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_42 %67 %68
- %70 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %51 1
- %71 = OpCompositeExtract %_arr_int_uint_4 %70 2
- %72 = OpCompositeConstruct %_arr__arr_int_uint_4_uint_2 %69 %71
- %73 = OpCompositeExtract %_arr_int_uint_4 %72 1
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %19 = OpIAdd %int %int_42 %int_1
+ %20 = OpCompositeExtract %int %18 3
+ %21 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_42 %19 %20
+ %55 = OpIAdd %int %int_42 %int_1
+ %56 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_2 %int_42 %55
+ %57 = OpCompositeExtract %int %18 2
+ %58 = OpCompositeExtract %int %18 3
+ %59 = OpIAdd %int %58 %int_1
+ %60 = OpCompositeConstruct %_arr_int_uint_4 %int_5 %int_6 %57 %59
+ %61 = OpCompositeConstruct %_arr__arr_int_uint_4_uint_3 %56 %60 %18
+ %62 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %54 1
+ %63 = OpCompositeConstruct %_arr__arr__arr_int_uint_4_uint_3_uint_2 %61 %62
+ %65 = OpIAdd %int %int_42 %int_1
+ %66 = OpCompositeExtract %int %18 3
+ %67 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_42 %65 %66
+ %68 = OpCompositeExtract %int %67 2
+ %70 = OpIAdd %int %int_42 %int_1
+ %71 = OpCompositeExtract %int %18 3
+ %72 = OpCompositeConstruct %_arr_int_uint_4 %int_1 %int_42 %70 %71
+ %73 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %54 1
+ %74 = OpCompositeExtract %_arr_int_uint_4 %73 2
+ %75 = OpCompositeConstruct %_arr__arr_int_uint_4_uint_2 %72 %74
+ %76 = OpCompositeExtract %_arr_int_uint_4 %75 1
+ %79 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+ %80 = OpCompositeExtract %int %13 0
+ %81 = OpCompositeExtract %int %18 0
+ %82 = OpIAdd %int %80 %81
+ %83 = OpCompositeExtract %int %21 0
+ %84 = OpIAdd %int %82 %83
+ %85 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %26 0
+ %86 = OpCompositeExtract %_arr_int_uint_4 %85 0
+ %87 = OpCompositeExtract %int %86 0
+ %88 = OpIAdd %int %84 %87
+ %89 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %54 0
+ %90 = OpCompositeExtract %_arr_int_uint_4 %89 0
+ %91 = OpCompositeExtract %int %90 0
+ %92 = OpIAdd %int %88 %91
+ %93 = OpCompositeExtract %_arr__arr_int_uint_4_uint_3 %63 0
+ %94 = OpCompositeExtract %_arr_int_uint_4 %93 0
+ %95 = OpCompositeExtract %int %94 0
+ %96 = OpIAdd %int %92 %95
+ %97 = OpIAdd %int %96 %64
+ %98 = OpIAdd %int %97 %int_3
+ %99 = OpIAdd %int %98 %68
+ %100 = OpCompositeExtract %int %13 0
+ %101 = OpIAdd %int %99 %100
+ %102 = OpCompositeExtract %int %31 0
+ %103 = OpIAdd %int %101 %102
+ %104 = OpCompositeExtract %int %76 0
+ %105 = OpIAdd %int %103 %104
+ OpStore %79 %105
OpReturn
OpFunctionEnd
diff --git a/test/tint/array/type_initializer.wgsl.expected.wgsl b/test/tint/array/type_initializer.wgsl.expected.wgsl
index a1a942a..fc876d5 100644
--- a/test/tint/array/type_initializer.wgsl.expected.wgsl
+++ b/test/tint/array/type_initializer.wgsl.expected.wgsl
@@ -1,3 +1,5 @@
+@group(0) @binding(0) var<storage, read_write> s : i32;
+
@compute @workgroup_size(1)
fn main() {
let x : i32 = 42;
@@ -13,4 +15,5 @@
let subexpr_nested_empty : array<i32, 4> = array<array<i32, 4>, 2>()[1];
let subexpr_nested_nonempty : array<i32, 4> = array<array<i32, 4>, 2>(array<i32, 4>(1, 2, 3, 4), array<i32, 4>(5, 6, 7, 8))[1];
let subexpr_nested_nonempty_with_expr : array<i32, 4> = array<array<i32, 4>, 2>(array<i32, 4>(1, x, (x + 1), nonempty[3]), nested_nonempty[1][2])[1];
+ s = (((((((((((empty[0] + nonempty[0]) + nonempty_with_expr[0]) + nested_empty[0][0][0]) + nested_nonempty[0][0][0]) + nested_nonempty_with_expr[0][0][0]) + subexpr_empty) + subexpr_nonempty) + subexpr_nonempty_with_expr) + subexpr_nested_empty[0]) + subexpr_nested_nonempty[0]) + subexpr_nested_nonempty_with_expr[0]);
}
