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_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;
+}
+
