encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

PromoteSideEffectsToDecl: add to 4 backends and regen tests 

Added to hlsl, msl, glsl, and spirv backends.

Bug: tint:1300
Change-Id: I06062bd8e4b32acbc4a8670b060a7a22bc1ae034
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/80600
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
Antonio Maiorano 2022-03-15 15:35:13 +00:00 committed by Tint LUCI CQ
parent c25ddf4b1c
commit 93baaae60b
80 changed files with 915 additions and 742 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/tint/transform/glsl.cc
View File

@ -29,6 +29,7 @@
#include "src/tint/transform/manager.h"
#include "src/tint/transform/pad_array_elements.h"
#include "src/tint/transform/promote_initializers_to_const_var.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
#include "src/tint/transform/renamer.h"
#include "src/tint/transform/simplify_pointers.h"
@ -83,6 +84,7 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const {
manager.Add<ZeroInitWorkgroupMemory>();
}
manager.Add<CanonicalizeEntryPointIO>();
manager.Add<PromoteSideEffectsToDecl>();
manager.Add<SimplifyPointers>();
manager.Add<RemovePhonies>();

2
src/tint/writer/hlsl/generator_impl.cc
View File

@ -59,6 +59,7 @@
#include "src/tint/transform/num_workgroups_from_uniform.h"
#include "src/tint/transform/pad_array_elements.h"
#include "src/tint/transform/promote_initializers_to_const_var.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
#include "src/tint/transform/simplify_pointers.h"
#include "src/tint/transform/unshadow.h"
@ -186,6 +187,7 @@ SanitizedResult Sanitize(
// assumes that num_workgroups builtins only appear as struct members and are
// only accessed directly via member accessors.
manager.Add<transform::NumWorkgroupsFromUniform>();
manager.Add<transform::PromoteSideEffectsToDecl>();
manager.Add<transform::SimplifyPointers>();
manager.Add<transform::RemovePhonies>();
// ArrayLengthFromUniform must come after InlinePointerLets and Simplify, as

2
src/tint/writer/msl/generator_impl.cc
View File

@ -65,6 +65,7 @@
#include "src/tint/transform/module_scope_var_to_entry_point_param.h"
#include "src/tint/transform/pad_array_elements.h"
#include "src/tint/transform/promote_initializers_to_const_var.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
#include "src/tint/transform/simplify_pointers.h"
#include "src/tint/transform/unshadow.h"
@ -173,6 +174,7 @@ SanitizedResult Sanitize(
manager.Add<transform::ZeroInitWorkgroupMemory>();
}
manager.Add<transform::CanonicalizeEntryPointIO>();
manager.Add<transform::PromoteSideEffectsToDecl>();
manager.Add<transform::PromoteInitializersToConstVar>();
manager.Add<transform::VectorizeScalarMatrixConstructors>();

2
src/tint/writer/spirv/builder.cc
View File

@ -48,6 +48,7 @@
#include "src/tint/transform/fold_constants.h"
#include "src/tint/transform/for_loop_to_loop.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_unreachable_statements.h"
#include "src/tint/transform/simplify_pointers.h"
#include "src/tint/transform/unshadow.h"
@ -276,6 +277,7 @@ SanitizedResult Sanitize(const Program* in,
manager.Add<transform::ZeroInitWorkgroupMemory>();
}
manager.Add<transform::RemoveUnreachableStatements>();
manager.Add<transform::PromoteSideEffectsToDecl>();
manager.Add<transform::SimplifyPointers>(); // Required for arrayLength()
manager.Add<transform::FoldConstants>();
manager.Add<transform::VectorizeScalarMatrixConstructors>();

15
test/tint/array/assign_to_function_var.wgsl.expected.glsl
View File

@ -18,20 +18,20 @@ layout(binding = 1, std430) buffer S_2 {
ivec4 arr[4];
} src_storage;
ivec4[4] ret_arr() {
ivec4 tint_symbol[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol;
ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol_1;
}
S ret_struct_arr() {
S tint_symbol_1 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_1;
S tint_symbol_2 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_2;
}
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_2[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst = tint_symbol_2;
ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst = tint_symbol_3;
dst = src_param;
dst = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
@ -39,7 +39,8 @@ void foo(ivec4 src_param[4]) {
dst = src_function;
dst = src_private;
dst = src_workgroup;
dst = ret_struct_arr().arr;
S tint_symbol = ret_struct_arr();
dst = tint_symbol.arr;
dst = src_uniform.arr;
dst = src_storage.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)));

31
test/tint/array/assign_to_function_var.wgsl.expected.hlsl
View File

@ -16,17 +16,17 @@ RWByteAddressBuffer src_storage : register(u1, space0);
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
const int4 tint_symbol_5[4] = (int4[4])0;
return tint_symbol_5;
}
S ret_struct_arr() {
const S tint_symbol_6 = (S)0;
const int4 tint_symbol_6[4] = (int4[4])0;
return tint_symbol_6;
}
typedef int4 tint_symbol_1_ret[4];
tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
S ret_struct_arr() {
const S tint_symbol_7 = (S)0;
return tint_symbol_7;
}
typedef int4 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
int4 arr_1[4] = (int4[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
@ -37,8 +37,8 @@ tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
return arr_1;
}
typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
typedef int4 tint_symbol_4_ret[4];
tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
int4 arr_2[4] = (int4[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
@ -51,8 +51,8 @@ tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
int4 tint_symbol[4] = (int4[4])0;
const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_7;
const int4 tint_symbol_8[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_8;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@ -60,9 +60,10 @@ void foo(int4 src_param[4]) {
tint_symbol = src_function;
tint_symbol = src_private;
tint_symbol = src_workgroup;
tint_symbol = ret_struct_arr().arr;
tint_symbol = tint_symbol_1(src_uniform, 0u);
tint_symbol = tint_symbol_3(src_storage, 0u);
const S tint_symbol_1 = ret_struct_arr();
tint_symbol = tint_symbol_1.arr;
tint_symbol = tint_symbol_2(src_uniform, 0u);
tint_symbol = tint_symbol_4(src_storage, 0u);
int dst_nested[4][3][2] = (int[4][3][2])0;
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;

21
test/tint/array/assign_to_function_var.wgsl.expected.msl
View File

@ -10,13 +10,13 @@ struct S {
};
tint_array_wrapper ret_arr() {
tint_array_wrapper const tint_symbol = {.arr={}};
return tint_symbol;
tint_array_wrapper const tint_symbol_1 = {.arr={}};
return tint_symbol_1;
}
S ret_struct_arr() {
S const tint_symbol_1 = {};
return tint_symbol_1;
S const tint_symbol_2 = {};
return tint_symbol_2;
}
struct tint_array_wrapper_3 {
@ -31,21 +31,22 @@ struct tint_array_wrapper_1 {
tint_array_wrapper_2 arr[4];
};
void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, threadgroup tint_array_wrapper* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7) {
tint_array_wrapper src_function = {};
tint_array_wrapper dst = {};
tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
dst = tint_symbol_2;
tint_array_wrapper const tint_symbol_3 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
dst = tint_symbol_3;
dst = src_param;
dst = ret_arr();
tint_array_wrapper const src_let = {.arr={}};
dst = src_let;
dst = src_function;
dst = *(tint_symbol_3);
dst = *(tint_symbol_4);
dst = ret_struct_arr().arr;
dst = (*(tint_symbol_5)).arr;
dst = *(tint_symbol_5);
S const tint_symbol = ret_struct_arr();
dst = tint_symbol.arr;
dst = (*(tint_symbol_6)).arr;
dst = (*(tint_symbol_7)).arr;
tint_array_wrapper_1 dst_nested = {};
tint_array_wrapper_1 src_nested = {};
dst_nested = src_nested;

15
test/tint/array/assign_to_private_var.wgsl.expected.glsl
View File

@ -20,19 +20,19 @@ layout(binding = 1, std430) buffer S_2 {
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[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol;
ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol_1;
}
S ret_struct_arr() {
S tint_symbol_1 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_1;
S tint_symbol_2 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_2;
}
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_2[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst = tint_symbol_2;
ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst = tint_symbol_3;
dst = src_param;
dst = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
@ -40,7 +40,8 @@ void foo(ivec4 src_param[4]) {
dst = src_function;
dst = src_private;
dst = src_workgroup;
dst = ret_struct_arr().arr;
S tint_symbol = ret_struct_arr();
dst = tint_symbol.arr;
dst = src_uniform.arr;
dst = src_storage.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)));

31
test/tint/array/assign_to_private_var.wgsl.expected.hlsl
View File

@ -18,17 +18,17 @@ static int dst_nested[4][3][2] = (int[4][3][2])0;
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
const int4 tint_symbol_5[4] = (int4[4])0;
return tint_symbol_5;
}
S ret_struct_arr() {
const S tint_symbol_6 = (S)0;
const int4 tint_symbol_6[4] = (int4[4])0;
return tint_symbol_6;
}
typedef int4 tint_symbol_1_ret[4];
tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
S ret_struct_arr() {
const S tint_symbol_7 = (S)0;
return tint_symbol_7;
}
typedef int4 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
int4 arr_1[4] = (int4[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
@ -39,8 +39,8 @@ tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
return arr_1;
}
typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
typedef int4 tint_symbol_4_ret[4];
tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
int4 arr_2[4] = (int4[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
@ -52,8 +52,8 @@ tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_7;
const int4 tint_symbol_8[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_8;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@ -61,9 +61,10 @@ void foo(int4 src_param[4]) {
tint_symbol = src_function;
tint_symbol = src_private;
tint_symbol = src_workgroup;
tint_symbol = ret_struct_arr().arr;
tint_symbol = tint_symbol_1(src_uniform, 0u);
tint_symbol = tint_symbol_3(src_storage, 0u);
const S tint_symbol_1 = ret_struct_arr();
tint_symbol = tint_symbol_1.arr;
tint_symbol = tint_symbol_2(src_uniform, 0u);
tint_symbol = tint_symbol_4(src_storage, 0u);
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}

45
test/tint/array/assign_to_private_var.wgsl.expected.msl
View File

@ -22,30 +22,31 @@ struct tint_array_wrapper_1 {
};
tint_array_wrapper ret_arr() {
tint_array_wrapper const tint_symbol = {.arr={}};
return tint_symbol;
}
S ret_struct_arr() {
S const tint_symbol_1 = {};
tint_array_wrapper const tint_symbol_1 = {.arr={}};
return tint_symbol_1;
}
void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, thread tint_array_wrapper_1* const tint_symbol_8) {
tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
*(tint_symbol_3) = tint_symbol_2;
*(tint_symbol_3) = src_param;
*(tint_symbol_3) = ret_arr();
tint_array_wrapper const src_let = {.arr={}};
*(tint_symbol_3) = src_let;
*(tint_symbol_3) = src_function;
*(tint_symbol_3) = *(tint_symbol_4);
*(tint_symbol_3) = *(tint_symbol_5);
*(tint_symbol_3) = ret_struct_arr().arr;
*(tint_symbol_3) = (*(tint_symbol_6)).arr;
*(tint_symbol_3) = (*(tint_symbol_7)).arr;
tint_array_wrapper_1 src_nested = {};
*(tint_symbol_8) = src_nested;
S ret_struct_arr() {
S const tint_symbol_2 = {};
return tint_symbol_2;
}
void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_4, thread tint_array_wrapper* const tint_symbol_5, threadgroup tint_array_wrapper* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, thread tint_array_wrapper_1* const tint_symbol_9) {
tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_3 = {.arr={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_wrapper const src_let = {.arr={}};
*(tint_symbol_4) = src_let;
*(tint_symbol_4) = src_function;
*(tint_symbol_4) = *(tint_symbol_5);
*(tint_symbol_4) = *(tint_symbol_6);
S const tint_symbol = ret_struct_arr();
*(tint_symbol_4) = tint_symbol.arr;
*(tint_symbol_4) = (*(tint_symbol_7)).arr;
*(tint_symbol_4) = (*(tint_symbol_8)).arr;
tint_array_wrapper_1 src_nested = {};
*(tint_symbol_9) = src_nested;
}

18
test/tint/array/assign_to_storage_var.wgsl.expected.glsl
View File

@ -28,27 +28,29 @@ layout(binding = 3, std430) buffer S_nested_1 {
int arr[4][3][2];
} dst_nested;
ivec4[4] ret_arr() {
ivec4 tint_symbol[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol;
ivec4 tint_symbol_2[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol_2;
}
S ret_struct_arr() {
S tint_symbol_1 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_1;
S tint_symbol_3 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 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_2[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst.arr = tint_symbol_2;
ivec4 tint_symbol_4[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst.arr = tint_symbol_4;
dst.arr = src_param;
dst.arr = ret_arr();
ivec4 tint_symbol[4] = ret_arr();
dst.arr = tint_symbol;
ivec4 src_let[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
dst.arr = src_let;
dst.arr = src_function;
dst.arr = src_private;
dst.arr = src_workgroup;
dst.arr = ret_struct_arr().arr;
S tint_symbol_1 = ret_struct_arr();
dst.arr = tint_symbol_1.arr;
dst.arr = src_uniform.arr;
dst.arr = src_storage.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)));

54
test/tint/array/assign_to_storage_var.wgsl.expected.hlsl
View File

@ -18,16 +18,16 @@ RWByteAddressBuffer dst_nested : register(u3, space0);
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
const int4 tint_symbol_11[4] = (int4[4])0;
return tint_symbol_11;
const int4 tint_symbol_13[4] = (int4[4])0;
return tint_symbol_13;
}
S ret_struct_arr() {
const S tint_symbol_12 = (S)0;
return tint_symbol_12;
const S tint_symbol_14 = (S)0;
return tint_symbol_14;
}
void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
int4 array[4] = value;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
@ -36,8 +36,8 @@ void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
}
}
typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
typedef int4 tint_symbol_5_ret[4];
tint_symbol_5_ret tint_symbol_5(uint4 buffer[4], uint offset) {
int4 arr_1[4] = (int4[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
@ -48,8 +48,8 @@ tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
return arr_1;
}
typedef int4 tint_symbol_5_ret[4];
tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) {
typedef int4 tint_symbol_7_ret[4];
tint_symbol_7_ret tint_symbol_7(RWByteAddressBuffer buffer, uint offset) {
int4 arr_2[4] = (int4[4])0;
{
[loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
@ -59,7 +59,7 @@ tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) {
return arr_2;
}
void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[2]) {
void tint_symbol_11(RWByteAddressBuffer buffer, uint offset, int value[2]) {
int array_3[2] = value;
{
[loop] for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) {
@ -68,38 +68,40 @@ void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[2]) {
}
}
void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, int value[3][2]) {
void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, int value[3][2]) {
int array_2[3][2] = value;
{
[loop] for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) {
tint_symbol_9(buffer, (offset + (i_4 * 8u)), array_2[i_4]);
tint_symbol_11(buffer, (offset + (i_4 * 8u)), array_2[i_4]);
}
}
}
void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) {
void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) {
int array_1[4][3][2] = value;
{
[loop] for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
tint_symbol_8(buffer, (offset + (i_5 * 24u)), array_1[i_5]);
tint_symbol_10(buffer, (offset + (i_5 * 24u)), array_1[i_5]);
}
}
}
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
const int4 tint_symbol_13[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol_1(tint_symbol, 0u, tint_symbol_13);
tint_symbol_1(tint_symbol, 0u, src_param);
tint_symbol_1(tint_symbol, 0u, ret_arr());
const int4 tint_symbol_15[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol_3(tint_symbol, 0u, tint_symbol_15);
tint_symbol_3(tint_symbol, 0u, src_param);
const int4 tint_symbol_1[4] = ret_arr();
tint_symbol_3(tint_symbol, 0u, tint_symbol_1);
const int4 src_let[4] = (int4[4])0;
tint_symbol_1(tint_symbol, 0u, src_let);
tint_symbol_1(tint_symbol, 0u, src_function);
tint_symbol_1(tint_symbol, 0u, src_private);
tint_symbol_1(tint_symbol, 0u, src_workgroup);
tint_symbol_1(tint_symbol, 0u, ret_struct_arr().arr);
tint_symbol_1(tint_symbol, 0u, tint_symbol_3(src_uniform, 0u));
tint_symbol_1(tint_symbol, 0u, tint_symbol_5(src_storage, 0u));
tint_symbol_3(tint_symbol, 0u, src_let);
tint_symbol_3(tint_symbol, 0u, src_function);
tint_symbol_3(tint_symbol, 0u, src_private);
tint_symbol_3(tint_symbol, 0u, src_workgroup);
const S tint_symbol_2 = ret_struct_arr();
tint_symbol_3(tint_symbol, 0u, tint_symbol_2.arr);
tint_symbol_3(tint_symbol, 0u, tint_symbol_5(src_uniform, 0u));
tint_symbol_3(tint_symbol, 0u, tint_symbol_7(src_storage, 0u));
int src_nested[4][3][2] = (int[4][3][2])0;
tint_symbol_7(dst_nested, 0u, src_nested);
tint_symbol_9(dst_nested, 0u, src_nested);
}

36
test/tint/array/assign_to_storage_var.wgsl.expected.msl
View File

@ -26,30 +26,32 @@ struct S_nested {
};
tint_array_wrapper ret_arr() {
tint_array_wrapper const tint_symbol = {.arr={}};
return tint_symbol;
tint_array_wrapper const tint_symbol_2 = {.arr={}};
return tint_symbol_2;
}
S ret_struct_arr() {
S const tint_symbol_1 = {};
return tint_symbol_1;
S const tint_symbol_3 = {};
return tint_symbol_3;
}
void foo(tint_array_wrapper src_param, device S* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, device S_nested* const tint_symbol_8) {
void foo(tint_array_wrapper src_param, device S* const tint_symbol_5, thread tint_array_wrapper* const tint_symbol_6, threadgroup tint_array_wrapper* const tint_symbol_7, const constant S* const tint_symbol_8, device S* const tint_symbol_9, device S_nested* const tint_symbol_10) {
tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
(*(tint_symbol_3)).arr = tint_symbol_2;
(*(tint_symbol_3)).arr = src_param;
(*(tint_symbol_3)).arr = ret_arr();
tint_array_wrapper const tint_symbol_4 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
(*(tint_symbol_5)).arr = tint_symbol_4;
(*(tint_symbol_5)).arr = src_param;
tint_array_wrapper const tint_symbol = ret_arr();
(*(tint_symbol_5)).arr = tint_symbol;
tint_array_wrapper const src_let = {.arr={}};
(*(tint_symbol_3)).arr = src_let;
(*(tint_symbol_3)).arr = src_function;
(*(tint_symbol_3)).arr = *(tint_symbol_4);
(*(tint_symbol_3)).arr = *(tint_symbol_5);
(*(tint_symbol_3)).arr = ret_struct_arr().arr;
(*(tint_symbol_3)).arr = (*(tint_symbol_6)).arr;
(*(tint_symbol_3)).arr = (*(tint_symbol_7)).arr;
(*(tint_symbol_5)).arr = src_let;
(*(tint_symbol_5)).arr = src_function;
(*(tint_symbol_5)).arr = *(tint_symbol_6);
(*(tint_symbol_5)).arr = *(tint_symbol_7);
S const tint_symbol_1 = ret_struct_arr();
(*(tint_symbol_5)).arr = tint_symbol_1.arr;
(*(tint_symbol_5)).arr = (*(tint_symbol_8)).arr;
(*(tint_symbol_5)).arr = (*(tint_symbol_9)).arr;
tint_array_wrapper_1 src_nested = {};
(*(tint_symbol_8)).arr = src_nested;
(*(tint_symbol_10)).arr = src_nested;
}

14
test/tint/array/assign_to_storage_var.wgsl.expected.spvasm
View File

@ -106,9 +106,9 @@
OpStore %44 %51
%52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %52 %src_param
%53 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%54 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
OpStore %53 %54
%53 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
%54 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %54 %53
%55 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %55 %8
%56 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
@ -120,10 +120,10 @@
%60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%61 = OpLoad %_arr_v4int_uint_4 %src_workgroup
OpStore %60 %61
%62 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%63 = OpFunctionCall %S %ret_struct_arr
%64 = OpCompositeExtract %_arr_v4int_uint_4 %63 0
OpStore %62 %64
%62 = OpFunctionCall %S %ret_struct_arr
%63 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%64 = OpCompositeExtract %_arr_v4int_uint_4 %62 0
OpStore %63 %64
%65 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%67 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
%68 = OpLoad %_arr_v4int_uint_4 %67

15
test/tint/array/assign_to_workgroup_var.wgsl.expected.glsl
View File

@ -20,19 +20,19 @@ layout(binding = 1, std430) buffer S_2 {
shared ivec4 dst[4];
shared int dst_nested[4][3][2];
ivec4[4] ret_arr() {
ivec4 tint_symbol[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol;
ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
return tint_symbol_1;
}
S ret_struct_arr() {
S tint_symbol_1 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_1;
S tint_symbol_2 = S(ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0)));
return tint_symbol_2;
}
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_2[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst = tint_symbol_2;
ivec4 tint_symbol_3[4] = ivec4[4](ivec4(1), ivec4(2), ivec4(3), ivec4(3));
dst = tint_symbol_3;
dst = src_param;
dst = ret_arr();
ivec4 src_let[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
@ -40,7 +40,8 @@ void foo(ivec4 src_param[4]) {
dst = src_function;
dst = src_private;
dst = src_workgroup;
dst = ret_struct_arr().arr;
S tint_symbol = ret_struct_arr();
dst = tint_symbol.arr;
dst = src_uniform.arr;
dst = src_storage.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)));

31
test/tint/array/assign_to_workgroup_var.wgsl.expected.hlsl
View File

@ -18,17 +18,17 @@ groupshared int dst_nested[4][3][2];
typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() {
const int4 tint_symbol_5[4] = (int4[4])0;
return tint_symbol_5;
}
S ret_struct_arr() {
const S tint_symbol_6 = (S)0;
const int4 tint_symbol_6[4] = (int4[4])0;
return tint_symbol_6;
}
typedef int4 tint_symbol_1_ret[4];
tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
S ret_struct_arr() {
const S tint_symbol_7 = (S)0;
return tint_symbol_7;
}
typedef int4 tint_symbol_2_ret[4];
tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
int4 arr_1[4] = (int4[4])0;
{
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
@ -39,8 +39,8 @@ tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
return arr_1;
}
typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
typedef int4 tint_symbol_4_ret[4];
tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
int4 arr_2[4] = (int4[4])0;
{
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
@ -52,8 +52,8 @@ tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
void foo(int4 src_param[4]) {
int4 src_function[4] = (int4[4])0;
const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_7;
const int4 tint_symbol_8[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_8;
tint_symbol = src_param;
tint_symbol = ret_arr();
const int4 src_let[4] = (int4[4])0;
@ -61,9 +61,10 @@ void foo(int4 src_param[4]) {
tint_symbol = src_function;
tint_symbol = src_private;
tint_symbol = src_workgroup;
tint_symbol = ret_struct_arr().arr;
tint_symbol = tint_symbol_1(src_uniform, 0u);
tint_symbol = tint_symbol_3(src_storage, 0u);
const S tint_symbol_1 = ret_struct_arr();
tint_symbol = tint_symbol_1.arr;
tint_symbol = tint_symbol_2(src_uniform, 0u);
tint_symbol = tint_symbol_4(src_storage, 0u);
int src_nested[4][3][2] = (int[4][3][2])0;
dst_nested = src_nested;
}

45
test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
View File

@ -22,30 +22,31 @@ struct tint_array_wrapper_1 {
};
tint_array_wrapper ret_arr() {
tint_array_wrapper const tint_symbol = {.arr={}};
return tint_symbol;
}
S ret_struct_arr() {
S const tint_symbol_1 = {};
tint_array_wrapper const tint_symbol_1 = {.arr={}};
return tint_symbol_1;
}
void foo(tint_array_wrapper src_param, threadgroup tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, threadgroup tint_array_wrapper_1* const tint_symbol_8) {
tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
*(tint_symbol_3) = tint_symbol_2;
*(tint_symbol_3) = src_param;
*(tint_symbol_3) = ret_arr();
tint_array_wrapper const src_let = {.arr={}};
*(tint_symbol_3) = src_let;
*(tint_symbol_3) = src_function;
*(tint_symbol_3) = *(tint_symbol_4);
*(tint_symbol_3) = *(tint_symbol_5);
*(tint_symbol_3) = ret_struct_arr().arr;
*(tint_symbol_3) = (*(tint_symbol_6)).arr;
*(tint_symbol_3) = (*(tint_symbol_7)).arr;
tint_array_wrapper_1 src_nested = {};
*(tint_symbol_8) = src_nested;
S ret_struct_arr() {
S const tint_symbol_2 = {};
return tint_symbol_2;
}
void foo(tint_array_wrapper src_param, threadgroup tint_array_wrapper* const tint_symbol_4, thread tint_array_wrapper* const tint_symbol_5, threadgroup tint_array_wrapper* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, threadgroup tint_array_wrapper_1* const tint_symbol_9) {
tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_3 = {.arr={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_wrapper const src_let = {.arr={}};
*(tint_symbol_4) = src_let;
*(tint_symbol_4) = src_function;
*(tint_symbol_4) = *(tint_symbol_5);
*(tint_symbol_4) = *(tint_symbol_6);
S const tint_symbol = ret_struct_arr();
*(tint_symbol_4) = tint_symbol.arr;
*(tint_symbol_4) = (*(tint_symbol_7)).arr;
*(tint_symbol_4) = (*(tint_symbol_8)).arr;
tint_array_wrapper_1 src_nested = {};
*(tint_symbol_9) = src_nested;
}

17
test/tint/array/function_return_type.wgsl.expected.glsl
View File

@ -1,18 +1,23 @@
#version 310 es
float[4] f1() {
float tint_symbol_1[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
return tint_symbol_1;
float tint_symbol_6[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
return tint_symbol_6;
}
float[3][4] f2() {
float tint_symbol_2[3][4] = float[3][4](f1(), f1(), f1());
return tint_symbol_2;
float tint_symbol_1[4] = f1();
float tint_symbol_2[4] = f1();
float tint_symbol_3[4] = f1();
float tint_symbol_7[3][4] = float[3][4](tint_symbol_1, tint_symbol_2, tint_symbol_3);
return tint_symbol_7;
}
float[2][3][4] f3() {
float tint_symbol_3[2][3][4] = float[2][3][4](f2(), f2());
return tint_symbol_3;
float tint_symbol_4[3][4] = f2();
float tint_symbol_5[3][4] = f2();
float tint_symbol_8[2][3][4] = float[2][3][4](tint_symbol_4, tint_symbol_5);
return tint_symbol_8;
}
void tint_symbol() {

17
test/tint/array/function_return_type.wgsl.expected.hlsl
View File

@ -1,19 +1,24 @@
typedef float f1_ret[4];
f1_ret f1() {
const float tint_symbol[4] = (float[4])0;
return tint_symbol;
const float tint_symbol_5[4] = (float[4])0;
return tint_symbol_5;
}
typedef float f2_ret[3][4];
f2_ret f2() {
const float tint_symbol_1[3][4] = {f1(), f1(), f1()};
return tint_symbol_1;
const float tint_symbol[4] = f1();
const float tint_symbol_1[4] = f1();
const float tint_symbol_2[4] = f1();
const float tint_symbol_6[3][4] = {tint_symbol, tint_symbol_1, tint_symbol_2};
return tint_symbol_6;
}
typedef float f3_ret[2][3][4];
f3_ret f3() {
const float tint_symbol_2[2][3][4] = {f2(), f2()};
return tint_symbol_2;
const float tint_symbol_3[3][4] = f2();
const float tint_symbol_4[3][4] = f2();
const float tint_symbol_7[2][3][4] = {tint_symbol_3, tint_symbol_4};
return tint_symbol_7;
}
[numthreads(1, 1, 1)]

17
test/tint/array/function_return_type.wgsl.expected.msl
View File

@ -6,8 +6,8 @@ struct tint_array_wrapper {
};
tint_array_wrapper f1() {
tint_array_wrapper const tint_symbol_1 = {.arr={}};
return tint_symbol_1;
tint_array_wrapper const tint_symbol_6 = {.arr={}};
return tint_symbol_6;
}
struct tint_array_wrapper_1 {
@ -15,8 +15,11 @@ struct tint_array_wrapper_1 {
};
tint_array_wrapper_1 f2() {
tint_array_wrapper_1 const tint_symbol_2 = {.arr={f1(), f1(), f1()}};
return tint_symbol_2;
tint_array_wrapper const tint_symbol_1 = f1();
tint_array_wrapper const tint_symbol_2 = f1();
tint_array_wrapper const tint_symbol_3 = f1();
tint_array_wrapper_1 const tint_symbol_7 = {.arr={tint_symbol_1, tint_symbol_2, tint_symbol_3}};
return tint_symbol_7;
}
struct tint_array_wrapper_2 {
@ -24,8 +27,10 @@ struct tint_array_wrapper_2 {
};
tint_array_wrapper_2 f3() {
tint_array_wrapper_2 const tint_symbol_3 = {.arr={f2(), f2()}};
return tint_symbol_3;
tint_array_wrapper_1 const tint_symbol_4 = f2();
tint_array_wrapper_1 const tint_symbol_5 = f2();
tint_array_wrapper_2 const tint_symbol_8 = {.arr={tint_symbol_4, tint_symbol_5}};
return tint_symbol_8;
}
kernel void tint_symbol() {

25
test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl
View File

@ -83,7 +83,10 @@ float linearDepth(float depthSample) {
uvec3 getTile(vec4 fragCoord) {
float sliceScale = (float(tileCount.z) / log2((camera.zFar / camera.zNear)));
float sliceBias = -(((float(tileCount.z) * log2(camera.zNear)) / log2((camera.zFar / camera.zNear))));
uint zTile = uint(max(((log2(linearDepth(fragCoord.z)) * sliceScale) + sliceBias), 0.0f));
float tint_symbol_3 = linearDepth(fragCoord.z);
float tint_symbol_4 = log2(tint_symbol_3);
float tint_symbol_5 = max(((tint_symbol_4 * sliceScale) + sliceBias), 0.0f);
uint zTile = uint(tint_symbol_5);
return uvec3(uint((fragCoord.x / (camera.outputSize.x / float(tileCount.x)))), uint((fragCoord.y / (camera.outputSize.y / float(tileCount.y)))), zTile);
}
@ -157,7 +160,9 @@ float pointLightVisibility(uint lightIndex, vec3 worldPos, vec3 pointToLight) {
if ((shadowIndex == -1)) {
return 1.0f;
}
shadowIndex = (shadowIndex + getCubeFace((pointToLight * -1.0f)));
int tint_symbol_6 = shadowIndex;
int tint_symbol_7 = getCubeFace((pointToLight * -1.0f));
shadowIndex = (tint_symbol_6 + tint_symbol_7);
vec4 viewport = shadow.properties[shadowIndex].viewport;
vec4 lightPos = (shadow.properties[shadowIndex].viewProj * vec4(worldPos, 1.0f));
vec3 shadowPos = vec3((((lightPos.xy / lightPos.w) * vec2(0.5f, -0.5f)) + vec2(0.5f, 0.5f)), (lightPos.z / lightPos.w));
@ -312,7 +317,9 @@ vec3 lightRadiance(PuctualLight light, SurfaceInfo surface) {
vec3 numerator = ((NDF * G) * F);
float denominator = max(((4.0f * max(dot(surface.normal, surface.v), 0.0f)) * NdotL), 0.001f);
vec3 specular = (numerator / vec3(denominator));
vec3 radiance = ((light.color * light.intensity) * lightAttenuation(light));
vec3 tint_symbol_8 = (light.color * light.intensity);
float tint_symbol_9 = lightAttenuation(light);
vec3 radiance = (tint_symbol_8 * tint_symbol_9);
return (((((kD * surface.albedo) / vec3(PI)) + specular) * radiance) * NdotL);
}
@ -333,7 +340,9 @@ FragmentOutput fragmentMain(VertexOutput tint_symbol) {
light.color = globalLights.dirColor;
light.intensity = globalLights.dirIntensity;
float lightVis = dirLightVisibility(tint_symbol.worldPos);
Lo = (Lo + (lightRadiance(light, surface) * lightVis));
vec3 tint_symbol_10 = Lo;
vec3 tint_symbol_11 = lightRadiance(light, surface);
Lo = (tint_symbol_10 + (tint_symbol_11 * lightVis));
}
uint clusterIndex = getClusterIndex(tint_symbol.position);
uint lightOffset = clusterLights.lights[clusterIndex].offset;
@ -348,7 +357,9 @@ FragmentOutput fragmentMain(VertexOutput tint_symbol) {
light.color = globalLights.lights[i].color;
light.intensity = globalLights.lights[i].intensity;
float lightVis = pointLightVisibility(i, tint_symbol.worldPos, light.pointToLight);
Lo = (Lo + (lightRadiance(light, surface) * lightVis));
vec3 tint_symbol_12 = Lo;
vec3 tint_symbol_13 = lightRadiance(light, surface);
Lo = (tint_symbol_12 + (tint_symbol_13 * lightVis));
}
}
vec2 ssaoCoord = (tint_symbol.position.xy / vec2(textureSize(ssaoTexture_1, 0).xy));
@ -362,8 +373,8 @@ FragmentOutput fragmentMain(VertexOutput tint_symbol) {
}
void main() {
VertexOutput tint_symbol_3 = VertexOutput(gl_FragCoord, worldPos_1, view_1, texcoord_1, texcoord2_1, color_1, instanceColor_1, normal_1, tangent_1, bitangent_1);
FragmentOutput inner_result = fragmentMain(tint_symbol_3);
VertexOutput tint_symbol_14 = VertexOutput(gl_FragCoord, worldPos_1, view_1, texcoord_1, texcoord2_1, color_1, instanceColor_1, normal_1, tangent_1, bitangent_1);
FragmentOutput inner_result = fragmentMain(tint_symbol_14);
color_2 = inner_result.color;
emissive_1 = inner_result.emissive;
return;

8
test/tint/bug/tint/1076.wgsl.expected.msl
View File

@ -20,13 +20,13 @@ FragIn tint_symbol_inner(FragIn in, float b) {
if ((in.mask == 0u)) {
return in;
}
FragIn const tint_symbol_5 = {.a=b, .mask=1u};
return tint_symbol_5;
FragIn const tint_symbol_4 = {.a=b, .mask=1u};
return tint_symbol_4;
}
fragment tint_symbol_3 tint_symbol(uint mask [[sample_mask]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
FragIn const tint_symbol_4 = {.a=tint_symbol_1.a, .mask=mask};
FragIn const inner_result = tint_symbol_inner(tint_symbol_4, tint_symbol_1.b);
FragIn const tint_symbol_5 = {.a=tint_symbol_1.a, .mask=mask};
FragIn const inner_result = tint_symbol_inner(tint_symbol_5, tint_symbol_1.b);
tint_symbol_3 wrapper_result = {};
wrapper_result.a = inner_result.a;
wrapper_result.mask = inner_result.mask;

3
test/tint/bug/tint/1113.wgsl.expected.glsl
View File

@ -199,7 +199,8 @@ void main_create_lut(uvec3 GlobalInvocationID) {
uint numTriangles = atomicOr(counters.values[voxelIndex], 0u);
int offset = -1;
if ((numTriangles > 0u)) {
offset = int(atomicAdd(dbg.offsetCounter, numTriangles));
uint tint_symbol = atomicAdd(dbg.offsetCounter, numTriangles);
offset = int(tint_symbol);
}
atomicExchange(LUT.values[voxelIndex], offset);
}

3
test/tint/bug/tint/1113.wgsl.expected.hlsl
View File

@ -126,7 +126,8 @@ void main_create_lut_inner(uint3 GlobalInvocationID) {
uint numTriangles = atomicLoad_1(counters, (4u * voxelIndex));
int offset = -1;
if ((numTriangles > 0u)) {
offset = int(atomicAdd_1(dbg, 0u, numTriangles));
const uint tint_symbol_6 = atomicAdd_1(dbg, 0u, numTriangles);
offset = int(tint_symbol_6);
}
atomicStore_1(LUT, (4u * voxelIndex), offset);
}

113
test/tint/bug/tint/1113.wgsl.expected.msl
View File

@ -58,15 +58,15 @@ struct AI32s {
/* 0x0000 */ atomic_int values[1];
};
float3 toVoxelPos(float3 position, const constant Uniforms* const tint_symbol) {
float3 bbMin = float3((*(tint_symbol)).bbMin[0], (*(tint_symbol)).bbMin[1], (*(tint_symbol)).bbMin[2]);
float3 bbMax = float3((*(tint_symbol)).bbMax[0], (*(tint_symbol)).bbMax[1], (*(tint_symbol)).bbMax[2]);
float3 toVoxelPos(float3 position, const constant Uniforms* const tint_symbol_1) {
float3 bbMin = float3((*(tint_symbol_1)).bbMin[0], (*(tint_symbol_1)).bbMin[1], (*(tint_symbol_1)).bbMin[2]);
float3 bbMax = float3((*(tint_symbol_1)).bbMax[0], (*(tint_symbol_1)).bbMax[1], (*(tint_symbol_1)).bbMax[2]);
float3 bbSize = (bbMax - bbMin);
float cubeSize = fmax(fmax(bbSize[0], bbSize[1]), bbSize[2]);
float gridSize = float((*(tint_symbol)).gridSize);
float gx = ((gridSize * (position[0] - (*(tint_symbol)).bbMin[0])) / cubeSize);
float gy = ((gridSize * (position[1] - (*(tint_symbol)).bbMin[1])) / cubeSize);
float gz = ((gridSize * (position[2] - (*(tint_symbol)).bbMin[2])) / cubeSize);
float gridSize = float((*(tint_symbol_1)).gridSize);
float gx = ((gridSize * (position[0] - (*(tint_symbol_1)).bbMin[0])) / cubeSize);
float gy = ((gridSize * (position[1] - (*(tint_symbol_1)).bbMin[1])) / cubeSize);
float gz = ((gridSize * (position[2] - (*(tint_symbol_1)).bbMin[2])) / cubeSize);
return float3(gx, gy, gz);
}
@ -82,89 +82,90 @@ uint3 toIndex3D(uint gridSize, uint index) {
return uint3(x_1, y_1, z_1);
}
float3 loadPosition(uint vertexIndex, device F32s* const tint_symbol_1) {
float3 position = float3((*(tint_symbol_1)).values[((3u * vertexIndex) + 0u)], (*(tint_symbol_1)).values[((3u * vertexIndex) + 1u)], (*(tint_symbol_1)).values[((3u * vertexIndex) + 2u)]);
float3 loadPosition(uint vertexIndex, device F32s* const tint_symbol_2) {
float3 position = float3((*(tint_symbol_2)).values[((3u * vertexIndex) + 0u)], (*(tint_symbol_2)).values[((3u * vertexIndex) + 1u)], (*(tint_symbol_2)).values[((3u * vertexIndex) + 2u)]);
return position;
}
void doIgnore(const constant Uniforms* const tint_symbol_2, device Dbg* const tint_symbol_3, device AU32s* const tint_symbol_4, device U32s* const tint_symbol_5, device F32s* const tint_symbol_6, device AI32s* const tint_symbol_7) {
uint g42 = (*(tint_symbol_2)).numTriangles;
uint kj6 = (*(tint_symbol_3)).value1;
uint b53 = atomic_load_explicit(&((*(tint_symbol_4)).values[0]), memory_order_relaxed);
uint rwg = (*(tint_symbol_5)).values[0];
float rb5 = (*(tint_symbol_6)).values[0];
int g55 = atomic_load_explicit(&((*(tint_symbol_7)).values[0]), memory_order_relaxed);
void doIgnore(const constant Uniforms* const tint_symbol_3, device Dbg* const tint_symbol_4, device AU32s* const tint_symbol_5, device U32s* const tint_symbol_6, device F32s* const tint_symbol_7, device AI32s* const tint_symbol_8) {
uint g42 = (*(tint_symbol_3)).numTriangles;
uint kj6 = (*(tint_symbol_4)).value1;
uint b53 = atomic_load_explicit(&((*(tint_symbol_5)).values[0]), memory_order_relaxed);
uint rwg = (*(tint_symbol_6)).values[0];
float rb5 = (*(tint_symbol_7)).values[0];
int g55 = atomic_load_explicit(&((*(tint_symbol_8)).values[0]), memory_order_relaxed);
}
void main_count_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_8, device Dbg* const tint_symbol_9, device AU32s* const tint_symbol_10, device U32s* const tint_symbol_11, device F32s* const tint_symbol_12, device AI32s* const tint_symbol_13) {
void main_count_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_9, device Dbg* const tint_symbol_10, device AU32s* const tint_symbol_11, device U32s* const tint_symbol_12, device F32s* const tint_symbol_13, device AI32s* const tint_symbol_14) {
uint triangleIndex = GlobalInvocationID[0];
if ((triangleIndex >= (*(tint_symbol_8)).numTriangles)) {
if ((triangleIndex >= (*(tint_symbol_9)).numTriangles)) {
return;
}
doIgnore(tint_symbol_8, tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12, tint_symbol_13);
uint i0 = (*(tint_symbol_11)).values[((3u * triangleIndex) + 0u)];
uint i1 = (*(tint_symbol_11)).values[((3u * triangleIndex) + 1u)];
uint i2 = (*(tint_symbol_11)).values[((3u * triangleIndex) + 2u)];
float3 p0 = loadPosition(i0, tint_symbol_12);
float3 p1 = loadPosition(i1, tint_symbol_12);
float3 p2 = loadPosition(i2, tint_symbol_12);
doIgnore(tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12, tint_symbol_13, tint_symbol_14);
uint i0 = (*(tint_symbol_12)).values[((3u * triangleIndex) + 0u)];
uint i1 = (*(tint_symbol_12)).values[((3u * triangleIndex) + 1u)];
uint i2 = (*(tint_symbol_12)).values[((3u * triangleIndex) + 2u)];
float3 p0 = loadPosition(i0, tint_symbol_13);
float3 p1 = loadPosition(i1, tint_symbol_13);
float3 p2 = loadPosition(i2, tint_symbol_13);
float3 center = (((p0 + p1) + p2) / 3.0f);
float3 voxelPos = toVoxelPos(center, tint_symbol_8);
uint voxelIndex = toIndex1D((*(tint_symbol_8)).gridSize, voxelPos);
uint acefg = atomic_fetch_add_explicit(&((*(tint_symbol_10)).values[voxelIndex]), 1u, memory_order_relaxed);
float3 voxelPos = toVoxelPos(center, tint_symbol_9);
uint voxelIndex = toIndex1D((*(tint_symbol_9)).gridSize, voxelPos);
uint acefg = atomic_fetch_add_explicit(&((*(tint_symbol_11)).values[voxelIndex]), 1u, memory_order_relaxed);
if ((triangleIndex == 0u)) {
(*(tint_symbol_9)).value0 = (*(tint_symbol_8)).gridSize;
(*(tint_symbol_9)).value_f32_0 = center[0];
(*(tint_symbol_9)).value_f32_1 = center[1];
(*(tint_symbol_9)).value_f32_2 = center[2];
(*(tint_symbol_10)).value0 = (*(tint_symbol_9)).gridSize;
(*(tint_symbol_10)).value_f32_0 = center[0];
(*(tint_symbol_10)).value_f32_1 = center[1];
(*(tint_symbol_10)).value_f32_2 = center[2];
}
}
kernel void main_count(const constant Uniforms* tint_symbol_14 [[buffer(0)]], device Dbg* tint_symbol_15 [[buffer(1)]], device AU32s* tint_symbol_16 [[buffer(2)]], device U32s* tint_symbol_17 [[buffer(3)]], device F32s* tint_symbol_18 [[buffer(4)]], device AI32s* tint_symbol_19 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
main_count_inner(GlobalInvocationID, tint_symbol_14, tint_symbol_15, tint_symbol_16, tint_symbol_17, tint_symbol_18, tint_symbol_19);
kernel void main_count(const constant Uniforms* tint_symbol_15 [[buffer(0)]], device Dbg* tint_symbol_16 [[buffer(1)]], device AU32s* tint_symbol_17 [[buffer(2)]], device U32s* tint_symbol_18 [[buffer(3)]], device F32s* tint_symbol_19 [[buffer(4)]], device AI32s* tint_symbol_20 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
main_count_inner(GlobalInvocationID, tint_symbol_15, tint_symbol_16, tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
return;
}
void main_create_lut_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_20, device Dbg* const tint_symbol_21, device AU32s* const tint_symbol_22, device U32s* const tint_symbol_23, device F32s* const tint_symbol_24, device AI32s* const tint_symbol_25) {
void main_create_lut_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_21, device Dbg* const tint_symbol_22, device AU32s* const tint_symbol_23, device U32s* const tint_symbol_24, device F32s* const tint_symbol_25, device AI32s* const tint_symbol_26) {
uint voxelIndex = GlobalInvocationID[0];
doIgnore(tint_symbol_20, tint_symbol_21, tint_symbol_22, tint_symbol_23, tint_symbol_24, tint_symbol_25);
uint maxVoxels = (((*(tint_symbol_20)).gridSize * (*(tint_symbol_20)).gridSize) * (*(tint_symbol_20)).gridSize);
doIgnore(tint_symbol_21, tint_symbol_22, tint_symbol_23, tint_symbol_24, tint_symbol_25, tint_symbol_26);
uint maxVoxels = (((*(tint_symbol_21)).gridSize * (*(tint_symbol_21)).gridSize) * (*(tint_symbol_21)).gridSize);
if ((voxelIndex >= maxVoxels)) {
return;
}
uint numTriangles = atomic_load_explicit(&((*(tint_symbol_22)).values[voxelIndex]), memory_order_relaxed);
uint numTriangles = atomic_load_explicit(&((*(tint_symbol_23)).values[voxelIndex]), memory_order_relaxed);
int offset = -1;
if ((numTriangles > 0u)) {
offset = int(atomic_fetch_add_explicit(&((*(tint_symbol_21)).offsetCounter), numTriangles, memory_order_relaxed));
uint const tint_symbol = atomic_fetch_add_explicit(&((*(tint_symbol_22)).offsetCounter), numTriangles, memory_order_relaxed);
offset = int(tint_symbol);
}
atomic_store_explicit(&((*(tint_symbol_25)).values[voxelIndex]), offset, memory_order_relaxed);
atomic_store_explicit(&((*(tint_symbol_26)).values[voxelIndex]), offset, memory_order_relaxed);
}
kernel void main_create_lut(const constant Uniforms* tint_symbol_26 [[buffer(0)]], device Dbg* tint_symbol_27 [[buffer(1)]], device AU32s* tint_symbol_28 [[buffer(2)]], device U32s* tint_symbol_29 [[buffer(3)]], device F32s* tint_symbol_30 [[buffer(4)]], device AI32s* tint_symbol_31 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
main_create_lut_inner(GlobalInvocationID, tint_symbol_26, tint_symbol_27, tint_symbol_28, tint_symbol_29, tint_symbol_30, tint_symbol_31);
kernel void main_create_lut(const constant Uniforms* tint_symbol_27 [[buffer(0)]], device Dbg* tint_symbol_28 [[buffer(1)]], device AU32s* tint_symbol_29 [[buffer(2)]], device U32s* tint_symbol_30 [[buffer(3)]], device F32s* tint_symbol_31 [[buffer(4)]], device AI32s* tint_symbol_32 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
main_create_lut_inner(GlobalInvocationID, tint_symbol_27, tint_symbol_28, tint_symbol_29, tint_symbol_30, tint_symbol_31, tint_symbol_32);
return;
}
void main_sort_triangles_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_32, device Dbg* const tint_symbol_33, device AU32s* const tint_symbol_34, device U32s* const tint_symbol_35, device F32s* const tint_symbol_36, device AI32s* const tint_symbol_37) {
void main_sort_triangles_inner(uint3 GlobalInvocationID, const constant Uniforms* const tint_symbol_33, device Dbg* const tint_symbol_34, device AU32s* const tint_symbol_35, device U32s* const tint_symbol_36, device F32s* const tint_symbol_37, device AI32s* const tint_symbol_38) {
uint triangleIndex = GlobalInvocationID[0];
doIgnore(tint_symbol_32, tint_symbol_33, tint_symbol_34, tint_symbol_35, tint_symbol_36, tint_symbol_37);
if ((triangleIndex >= (*(tint_symbol_32)).numTriangles)) {
doIgnore(tint_symbol_33, tint_symbol_34, tint_symbol_35, tint_symbol_36, tint_symbol_37, tint_symbol_38);
if ((triangleIndex >= (*(tint_symbol_33)).numTriangles)) {
return;
}
uint i0 = (*(tint_symbol_35)).values[((3u * triangleIndex) + 0u)];
uint i1 = (*(tint_symbol_35)).values[((3u * triangleIndex) + 1u)];
uint i2 = (*(tint_symbol_35)).values[((3u * triangleIndex) + 2u)];
float3 p0 = loadPosition(i0, tint_symbol_36);
float3 p1 = loadPosition(i1, tint_symbol_36);
float3 p2 = loadPosition(i2, tint_symbol_36);
uint i0 = (*(tint_symbol_36)).values[((3u * triangleIndex) + 0u)];
uint i1 = (*(tint_symbol_36)).values[((3u * triangleIndex) + 1u)];
uint i2 = (*(tint_symbol_36)).values[((3u * triangleIndex) + 2u)];
float3 p0 = loadPosition(i0, tint_symbol_37);
float3 p1 = loadPosition(i1, tint_symbol_37);
float3 p2 = loadPosition(i2, tint_symbol_37);
float3 center = (((p0 + p1) + p2) / 3.0f);
float3 voxelPos = toVoxelPos(center, tint_symbol_32);
uint voxelIndex = toIndex1D((*(tint_symbol_32)).gridSize, voxelPos);
int triangleOffset = atomic_fetch_add_explicit(&((*(tint_symbol_37)).values[voxelIndex]), 1, memory_order_relaxed);
float3 voxelPos = toVoxelPos(center, tint_symbol_33);
uint voxelIndex = toIndex1D((*(tint_symbol_33)).gridSize, voxelPos);
int triangleOffset = atomic_fetch_add_explicit(&((*(tint_symbol_38)).values[voxelIndex]), 1, memory_order_relaxed);
}
kernel void main_sort_triangles(const constant Uniforms* tint_symbol_38 [[buffer(0)]], device Dbg* tint_symbol_39 [[buffer(1)]], device AU32s* tint_symbol_40 [[buffer(2)]], device U32s* tint_symbol_41 [[buffer(3)]], device F32s* tint_symbol_42 [[buffer(4)]], device AI32s* tint_symbol_43 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
main_sort_triangles_inner(GlobalInvocationID, tint_symbol_38, tint_symbol_39, tint_symbol_40, tint_symbol_41, tint_symbol_42, tint_symbol_43);
kernel void main_sort_triangles(const constant Uniforms* tint_symbol_39 [[buffer(0)]], device Dbg* tint_symbol_40 [[buffer(1)]], device AU32s* tint_symbol_41 [[buffer(2)]], device U32s* tint_symbol_42 [[buffer(3)]], device F32s* tint_symbol_43 [[buffer(4)]], device AI32s* tint_symbol_44 [[buffer(5)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
main_sort_triangles_inner(GlobalInvocationID, tint_symbol_39, tint_symbol_40, tint_symbol_41, tint_symbol_42, tint_symbol_43, tint_symbol_44);
return;
}

10
test/tint/bug/tint/1113.wgsl.expected.spvasm
View File

@ -542,11 +542,11 @@
OpSelectionMerge %321 None
OpBranchConditional %320 %322 %321
%322 = OpLabel
%326 = OpAccessChain %_ptr_StorageBuffer_uint_0 %dbg %uint_0
%327 = OpLoad %uint %numTriangles
%324 = OpAtomicIAdd %uint %326 %uint_1 %uint_0 %327
%323 = OpBitcast %int %324
OpStore %offset %323
%325 = OpAccessChain %_ptr_StorageBuffer_uint_0 %dbg %uint_0
%326 = OpLoad %uint %numTriangles
%323 = OpAtomicIAdd %uint %325 %uint_1 %uint_0 %326
%327 = OpBitcast %int %323
OpStore %offset %327
OpBranch %321
%321 = OpLabel
%330 = OpLoad %uint %voxelIndex_0

3
test/tint/bug/tint/1321.wgsl.expected.glsl
View File

@ -7,7 +7,8 @@ int foo() {
void tint_symbol() {
float arr[4] = float[4](0.0f, 0.0f, 0.0f, 0.0f);
int a_save = foo();
int tint_symbol_1 = foo();
int a_save = tint_symbol_1;
{
for(; ; ) {
float x = arr[a_save];

3
test/tint/bug/tint/1321.wgsl.expected.msl
View File

@ -11,7 +11,8 @@ struct tint_array_wrapper {
fragment void tint_symbol() {
tint_array_wrapper arr = {.arr={}};
int const a_save = foo();
int const tint_symbol_1 = foo();
int const a_save = tint_symbol_1;
for(; ; ) {
float const x = arr.arr[a_save];
break;

3
test/tint/bug/tint/534.wgsl.expected.glsl
View File

@ -37,7 +37,8 @@ void tint_symbol_1(uvec3 GlobalInvocationID) {
uvec4 dstColorBits = uvec4(dstColor);
{
for(uint i = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
uint tint_symbol_2 = ConvertToFp16FloatValue(srcColor[i]);
srcColorBits[i] = tint_symbol_2;
bool tint_tmp = success;
if (tint_tmp) {
tint_tmp = (srcColorBits[i] == dstColorBits[i]);

3
test/tint/bug/tint/534.wgsl.expected.hlsl
View File

@ -33,7 +33,8 @@ void main_inner(uint3 GlobalInvocationID) {
uint4 dstColorBits = uint4(dstColor);
{
[loop] for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
const uint tint_symbol_3 = ConvertToFp16FloatValue(srcColor[i]);
set_uint4(srcColorBits, i, tint_symbol_3);
bool tint_tmp_1 = success;
if (tint_tmp_1) {
tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);

23
test/tint/bug/tint/534.wgsl.expected.msl
View File

@ -16,32 +16,33 @@ uint ConvertToFp16FloatValue(float fp32) {
return 1u;
}
void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_1, const constant Uniforms* const tint_symbol_2, texture2d<float, access::sample> tint_symbol_3, device OutputBuf* const tint_symbol_4) {
int2 size = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_2, const constant Uniforms* const tint_symbol_3, texture2d<float, access::sample> tint_symbol_4, device OutputBuf* const tint_symbol_5) {
int2 size = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
int2 dstTexCoord = int2(uint3(GlobalInvocationID).xy);
int2 srcTexCoord = dstTexCoord;
if (((*(tint_symbol_2)).dstTextureFlipY == 1u)) {
if (((*(tint_symbol_3)).dstTextureFlipY == 1u)) {
srcTexCoord[1] = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(size[1]) - as_type<uint>(dstTexCoord[1])))) - as_type<uint>(1)));
}
float4 srcColor = tint_symbol_1.read(uint2(srcTexCoord), 0);
float4 dstColor = tint_symbol_3.read(uint2(dstTexCoord), 0);
float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
float4 dstColor = tint_symbol_4.read(uint2(dstTexCoord), 0);
bool success = true;
uint4 srcColorBits = 0u;
uint4 dstColorBits = uint4(dstColor);
for(uint i = 0u; (i < (*(tint_symbol_2)).channelCount); i = (i + 1u)) {
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
uint const tint_symbol_1 = ConvertToFp16FloatValue(srcColor[i]);
srcColorBits[i] = tint_symbol_1;
success = (success && (srcColorBits[i] == dstColorBits[i]));
}
uint outputIndex = ((GlobalInvocationID[1] * uint(size[0])) + GlobalInvocationID[0]);
if (success) {
(*(tint_symbol_4)).result[outputIndex] = uint(1);
(*(tint_symbol_5)).result[outputIndex] = uint(1);
} else {
(*(tint_symbol_4)).result[outputIndex] = uint(0);
(*(tint_symbol_5)).result[outputIndex] = uint(0);
}
}
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]], const constant Uniforms* tint_symbol_6 [[buffer(0)]], texture2d<float, access::sample> tint_symbol_7 [[texture(1)]], device OutputBuf* tint_symbol_8 [[buffer(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
tint_symbol_inner(GlobalInvocationID, tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_6 [[texture(0)]], const constant Uniforms* tint_symbol_7 [[buffer(0)]], texture2d<float, access::sample> tint_symbol_8 [[texture(1)]], device OutputBuf* tint_symbol_9 [[buffer(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
tint_symbol_inner(GlobalInvocationID, tint_symbol_6, tint_symbol_7, tint_symbol_8, tint_symbol_9);
return;
}

14
test/tint/bug/tint/534.wgsl.expected.spvasm
View File

@ -169,13 +169,13 @@
%94 = OpLabel
OpBranch %84
%93 = OpLabel
%95 = OpLoad %uint %i
%96 = OpAccessChain %_ptr_Function_uint %srcColorBits %95
%98 = OpLoad %uint %i
%100 = OpAccessChain %_ptr_Function_float %srcColor %98
%101 = OpLoad %float %100
%97 = OpFunctionCall %uint %ConvertToFp16FloatValue %101
OpStore %96 %97
%96 = OpLoad %uint %i
%98 = OpAccessChain %_ptr_Function_float %srcColor %96
%99 = OpLoad %float %98
%95 = OpFunctionCall %uint %ConvertToFp16FloatValue %99
%100 = OpLoad %uint %i
%101 = OpAccessChain %_ptr_Function_uint %srcColorBits %100
OpStore %101 %95
%102 = OpLoad %bool %success
OpSelectionMerge %103 None
OpBranchConditional %102 %104 %103

248
test/tint/bug/tint/749.spvasm.expected.msl
View File

@ -61,18 +61,18 @@ void swap_i1_i1_(thread int* const i, thread int* const j, thread QuicksortObjec
int const x_34_save = x_33;
int const x_35 = (*(tint_symbol_81)).numbers.arr[x_34_save];
QuicksortObject const x_943 = *(tint_symbol_81);
tint_array_wrapper const tint_symbol_3 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_4 = {.numbers=tint_symbol_3};
*(tint_symbol_81) = tint_symbol_4;
tint_array_wrapper const tint_symbol_2 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_3 = {.numbers=tint_symbol_2};
*(tint_symbol_81) = tint_symbol_3;
*(tint_symbol_81) = x_943;
float2 const x_527 = float2(x_526[0], x_526[0]);
int const x_36_save = x_32;
float3 const x_528 = float3(x_524[0], x_524[2], x_524[0]);
(*(tint_symbol_81)).numbers.arr[x_36_save] = x_35;
QuicksortObject const x_944 = *(tint_symbol_81);
tint_array_wrapper const tint_symbol_5 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_6 = {.numbers=tint_symbol_5};
*(tint_symbol_81) = tint_symbol_6;
tint_array_wrapper const tint_symbol_4 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_5 = {.numbers=tint_symbol_4};
*(tint_symbol_81) = tint_symbol_5;
*(tint_symbol_81) = x_944;
float3 const x_529 = float3(x_526[1], x_526[2], x_526[0]);
int const x_945 = *(i);
@ -95,9 +95,9 @@ void swap_i1_i1_(thread int* const i, thread int* const j, thread QuicksortObjec
(*(tint_symbol_81)).numbers.arr[x_36_save] = 0;
(*(tint_symbol_81)).numbers.arr[x_36_save] = x_949;
QuicksortObject const x_950 = *(tint_symbol_81);
tint_array_wrapper const tint_symbol_7 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_8 = {.numbers=tint_symbol_7};
*(tint_symbol_81) = tint_symbol_8;
tint_array_wrapper const tint_symbol_6 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_7 = {.numbers=tint_symbol_6};
*(tint_symbol_81) = tint_symbol_7;
*(tint_symbol_81) = x_950;
float3 const x_532 = float3(x_528[0], x_528[1], x_528[0]);
int const x_951 = (*(tint_symbol_81)).numbers.arr[x_34_save];
@ -153,9 +153,9 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
float3 const x_536 = float3(x_534[0], x_534[2], x_535[0]);
j_1 = 10;
QuicksortObject const x_960 = *(tint_symbol_82);
tint_array_wrapper const tint_symbol_9 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_10 = {.numbers=tint_symbol_9};
*(tint_symbol_82) = tint_symbol_10;
tint_array_wrapper const tint_symbol_8 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_9 = {.numbers=tint_symbol_8};
*(tint_symbol_82) = tint_symbol_9;
*(tint_symbol_82) = x_960;
while (true) {
int const x_961 = pivot;
@ -170,9 +170,9 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
pivot = x_963;
x_537 = float2(float3(1.0f, 2.0f, 3.0f)[1], float3(1.0f, 2.0f, 3.0f)[2]);
QuicksortObject const x_964 = *(tint_symbol_82);
tint_array_wrapper const tint_symbol_11 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_12 = {.numbers=tint_symbol_11};
*(tint_symbol_82) = tint_symbol_12;
tint_array_wrapper const tint_symbol_10 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_11 = {.numbers=tint_symbol_10};
*(tint_symbol_82) = tint_symbol_11;
*(tint_symbol_82) = x_964;
int const x_56 = *(h);
int const x_965 = *(h);
@ -206,9 +206,9 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
param_1 = x_971;
int const x_62 = (*(tint_symbol_82)).numbers.arr[x_61_save];
QuicksortObject const x_972 = *(tint_symbol_82);
tint_array_wrapper const tint_symbol_13 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_14 = {.numbers=tint_symbol_13};
*(tint_symbol_82) = tint_symbol_14;
tint_array_wrapper const tint_symbol_12 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_13 = {.numbers=tint_symbol_12};
*(tint_symbol_82) = tint_symbol_13;
*(tint_symbol_82) = x_972;
int const x_63 = pivot;
float2 const x_540 = float2(float3(1.0f, 2.0f, 3.0f)[1], x_534[2]);
@ -267,9 +267,9 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
param_1 = x_985;
}
QuicksortObject const x_986 = *(tint_symbol_82);
tint_array_wrapper const tint_symbol_15 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_16 = {.numbers=tint_symbol_15};
*(tint_symbol_82) = tint_symbol_16;
tint_array_wrapper const tint_symbol_14 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_15 = {.numbers=tint_symbol_14};
*(tint_symbol_82) = tint_symbol_15;
*(tint_symbol_82) = x_986;
{
int const x_987 = *(h);
@ -302,9 +302,9 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
(*(tint_symbol_82)).numbers.arr[x_42_save] = x_993;
float2 const x_549 = float2(x_534[0], x_534[1]);
QuicksortObject const x_994 = *(tint_symbol_82);
tint_array_wrapper const tint_symbol_17 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_18 = {.numbers=tint_symbol_17};
*(tint_symbol_82) = tint_symbol_18;
tint_array_wrapper const tint_symbol_16 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_17 = {.numbers=tint_symbol_16};
*(tint_symbol_82) = tint_symbol_17;
*(tint_symbol_82) = x_994;
int const x_995 = *(h);
*(h) = 0;
@ -372,8 +372,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
param_5 = x_1007;
h_1 = 9;
tint_array_wrapper const x_1008 = stack;
tint_array_wrapper const tint_symbol_19 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_19;
tint_array_wrapper const tint_symbol_18 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_18;
stack = x_1008;
float2 const x_556 = float2(float3(1.0f, 2.0f, 3.0f)[1], float3(1.0f, 2.0f, 3.0f)[1]);
int const x_1009 = param_5;
@ -406,15 +406,15 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
param_4 = x_1015;
int const x_95 = l_1;
QuicksortObject const x_1016 = *(tint_symbol_83);
tint_array_wrapper const tint_symbol_20 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_21 = {.numbers=tint_symbol_20};
*(tint_symbol_83) = tint_symbol_21;
tint_array_wrapper const tint_symbol_19 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_20 = {.numbers=tint_symbol_19};
*(tint_symbol_83) = tint_symbol_20;
*(tint_symbol_83) = x_1016;
float3 const x_560 = float3(x_559[1], x_559[0], x_557[0]);
int const x_96_save = x_94;
tint_array_wrapper const x_1017 = stack;
tint_array_wrapper const tint_symbol_22 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_22;
tint_array_wrapper const tint_symbol_21 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_21;
stack = x_1017;
float3 const x_561 = float3(x_556[1], x_556[1], x_556[1]);
int const x_1018 = l_1;
@ -464,13 +464,13 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
h_1 = 0;
h_1 = x_1028;
tint_array_wrapper const x_1029 = stack;
tint_array_wrapper const tint_symbol_23 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_23;
tint_array_wrapper const tint_symbol_22 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_22;
stack = x_1029;
int const x_106 = top;
tint_array_wrapper const x_1030 = stack;
tint_array_wrapper const tint_symbol_24 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_24;
tint_array_wrapper const tint_symbol_23 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_23;
stack = x_1030;
float2 const x_567 = float2(x_558[0], x_564[2]);
int const x_1031 = param_4;
@ -481,9 +481,9 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
break;
}
QuicksortObject const x_1032 = *(tint_symbol_83);
tint_array_wrapper const tint_symbol_25 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_26 = {.numbers=tint_symbol_25};
*(tint_symbol_83) = tint_symbol_26;
tint_array_wrapper const tint_symbol_24 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_25 = {.numbers=tint_symbol_24};
*(tint_symbol_83) = tint_symbol_25;
*(tint_symbol_83) = x_1032;
float3 const x_568 = float3(x_559[1], x_559[0], x_563[1]);
int const x_1033 = param_4;
@ -508,8 +508,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
stack.arr[x_96_save] = x_1037;
int const x_111 = stack.arr[x_110_save];
tint_array_wrapper const x_1038 = stack;
tint_array_wrapper const tint_symbol_27 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_27;
tint_array_wrapper const tint_symbol_26 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_26;
stack = x_1038;
float3 const x_571 = float3(x_559[1], x_559[0], x_564[1]);
int const x_1039 = l_1;
@ -517,8 +517,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
l_1 = x_1039;
h_1 = x_111;
tint_array_wrapper const x_1040 = stack;
tint_array_wrapper const tint_symbol_28 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_28;
tint_array_wrapper const tint_symbol_27 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_27;
stack = x_1040;
float2 const x_572 = float2(x_562[1], x_561[1]);
int const x_1041 = p;
@ -610,8 +610,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
stack.arr[x_100_save] = 0;
stack.arr[x_100_save] = x_1061;
tint_array_wrapper const x_1062 = stack;
tint_array_wrapper const tint_symbol_29 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_29;
tint_array_wrapper const tint_symbol_28 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_28;
stack = x_1062;
float2 const x_584 = float2(x_569[2], x_569[1]);
float3 const x_585 = float3(x_580[1], x_577[0], x_577[0]);
@ -650,8 +650,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
h_1 = x_1070;
top = x_133;
tint_array_wrapper const x_1071 = stack;
tint_array_wrapper const tint_symbol_30 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_30;
tint_array_wrapper const tint_symbol_29 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_29;
stack = x_1071;
int const x_134 = p;
float2 const x_590 = float2(x_576[0], x_573[1]);
@ -676,9 +676,9 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
stack.arr[x_96_save] = x_1076;
float2 const x_592 = float2(float3(1.0f, 2.0f, 3.0f)[0], float3(1.0f, 2.0f, 3.0f)[1]);
QuicksortObject const x_1077 = *(tint_symbol_83);
tint_array_wrapper const tint_symbol_31 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_32 = {.numbers=tint_symbol_31};
*(tint_symbol_83) = tint_symbol_32;
tint_array_wrapper const tint_symbol_30 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_31 = {.numbers=tint_symbol_30};
*(tint_symbol_83) = tint_symbol_31;
*(tint_symbol_83) = x_1077;
int const x_137 = p;
int const x_1078 = stack.arr[x_114_save];
@ -743,8 +743,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
float2 const x_601 = float2(x_563[0], x_563[1]);
stack.arr[x_147_save] = as_type<int>((1u + as_type<uint>(x_145)));
tint_array_wrapper const x_1093 = stack;
tint_array_wrapper const tint_symbol_33 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_33;
tint_array_wrapper const tint_symbol_32 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_32;
stack = x_1093;
int const x_148 = top;
int const x_1094 = stack.arr[x_114_save];
@ -752,8 +752,8 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
stack.arr[x_114_save] = x_1094;
float2 const x_602 = float2(x_565[1], x_599[1]);
tint_array_wrapper const x_1095 = stack;
tint_array_wrapper const tint_symbol_34 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_34;
tint_array_wrapper const tint_symbol_33 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
stack = tint_symbol_33;
stack = x_1095;
int const x_149 = as_type<int>((as_type<uint>(x_148) + as_type<uint>(as_type<int>(1u))));
int const x_1096 = stack.arr[x_147_save];
@ -788,9 +788,9 @@ void quicksort_(thread QuicksortObject* const tint_symbol_83) {
l_1 = x_1103;
float2 const x_604 = float2(x_563[2], x_564[0]);
QuicksortObject const x_1104 = *(tint_symbol_83);
tint_array_wrapper const tint_symbol_35 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_36 = {.numbers=tint_symbol_35};
*(tint_symbol_83) = tint_symbol_36;
tint_array_wrapper const tint_symbol_34 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_35 = {.numbers=tint_symbol_34};
*(tint_symbol_83) = tint_symbol_35;
*(tint_symbol_83) = x_1104;
}
}
@ -809,15 +809,15 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv = x_717;
i_2 = 0;
QuicksortObject const x_721 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_37 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_38 = {.numbers=tint_symbol_37};
*(tint_symbol_84) = tint_symbol_38;
tint_array_wrapper const tint_symbol_36 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_37 = {.numbers=tint_symbol_36};
*(tint_symbol_84) = tint_symbol_37;
*(tint_symbol_84) = x_721;
if (true) {
QuicksortObject const x_722 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_39 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_40 = {.numbers=tint_symbol_39};
*(tint_symbol_84) = tint_symbol_40;
tint_array_wrapper const tint_symbol_38 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_39 = {.numbers=tint_symbol_38};
*(tint_symbol_84) = tint_symbol_39;
*(tint_symbol_84) = x_722;
float2 const x_431 = float2(float3(1.0f, 2.0f, 3.0f)[0], float3(1.0f, 2.0f, 3.0f)[0]);
int const x_158 = i_2;
@ -829,15 +829,15 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
color = x_725;
float2 const x_432 = float2(x_431[1], x_431[1]);
QuicksortObject const x_726 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_41 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_42 = {.numbers=tint_symbol_41};
*(tint_symbol_84) = tint_symbol_42;
tint_array_wrapper const tint_symbol_40 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_41 = {.numbers=tint_symbol_40};
*(tint_symbol_84) = tint_symbol_41;
*(tint_symbol_84) = x_726;
}
QuicksortObject const x_756 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_43 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_44 = {.numbers=tint_symbol_43};
*(tint_symbol_84) = tint_symbol_44;
tint_array_wrapper const tint_symbol_42 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_43 = {.numbers=tint_symbol_42};
*(tint_symbol_84) = tint_symbol_43;
*(tint_symbol_84) = x_756;
float2 const x_446 = float2(float2()[0], float2()[0]);
int const x_757 = i_2;
@ -845,9 +845,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
i_2 = x_757;
quicksort_(tint_symbol_84);
QuicksortObject const x_758 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_45 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_46 = {.numbers=tint_symbol_45};
*(tint_symbol_84) = tint_symbol_46;
tint_array_wrapper const tint_symbol_44 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_45 = {.numbers=tint_symbol_44};
*(tint_symbol_84) = tint_symbol_45;
*(tint_symbol_84) = x_758;
float4 const x_184 = *(tint_symbol_85);
float2 const x_759 = uv;
@ -860,18 +860,18 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
float2 const x_185 = float2(x_184[0], x_184[1]);
float3 const x_448 = float3(x_185[1], x_446[1], x_446[1]);
QuicksortObject const x_761 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_47 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_48 = {.numbers=tint_symbol_47};
*(tint_symbol_84) = tint_symbol_48;
tint_array_wrapper const tint_symbol_46 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_47 = {.numbers=tint_symbol_46};
*(tint_symbol_84) = tint_symbol_47;
*(tint_symbol_84) = x_761;
float2 const x_762 = uv;
uv = float2(0.0f, 0.0f);
uv = x_762;
float2 const x_191 = (*(tint_symbol_86)).resolution;
QuicksortObject const x_763 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_49 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_50 = {.numbers=tint_symbol_49};
*(tint_symbol_84) = tint_symbol_50;
tint_array_wrapper const tint_symbol_48 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_49 = {.numbers=tint_symbol_48};
*(tint_symbol_84) = tint_symbol_49;
*(tint_symbol_84) = x_763;
float3 const x_449 = float3(x_184[1], float3(1.0f, 2.0f, 3.0f)[2], x_184[3]);
float3 const x_764 = color;
@ -879,9 +879,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
color = x_764;
float2 const x_192 = (x_185 / x_191);
QuicksortObject const x_765 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_51 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_52 = {.numbers=tint_symbol_51};
*(tint_symbol_84) = tint_symbol_52;
tint_array_wrapper const tint_symbol_50 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_51 = {.numbers=tint_symbol_50};
*(tint_symbol_84) = tint_symbol_51;
*(tint_symbol_84) = x_765;
float2 const x_450 = float2(x_447[0], x_185[1]);
float3 const x_766 = color;
@ -897,18 +897,18 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
color = x_768;
float3 const x_451 = float3(x_185[0], x_185[1], x_446[1]);
QuicksortObject const x_769 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_53 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_54 = {.numbers=tint_symbol_53};
*(tint_symbol_84) = tint_symbol_54;
tint_array_wrapper const tint_symbol_52 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_53 = {.numbers=tint_symbol_52};
*(tint_symbol_84) = tint_symbol_53;
*(tint_symbol_84) = x_769;
int const x_770 = (*(tint_symbol_84)).numbers.arr[0u];
(*(tint_symbol_84)).numbers.arr[0u] = 0;
(*(tint_symbol_84)).numbers.arr[0u] = x_770;
int const x_201 = (*(tint_symbol_84)).numbers.arr[0u];
QuicksortObject const x_771 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_55 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_56 = {.numbers=tint_symbol_55};
*(tint_symbol_84) = tint_symbol_56;
tint_array_wrapper const tint_symbol_54 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_55 = {.numbers=tint_symbol_54};
*(tint_symbol_84) = tint_symbol_55;
*(tint_symbol_84) = x_771;
int const x_772 = (*(tint_symbol_84)).numbers.arr[0u];
(*(tint_symbol_84)).numbers.arr[0u] = 0;
@ -922,9 +922,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
i_2 = 0;
i_2 = x_774;
QuicksortObject const x_775 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_57 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_58 = {.numbers=tint_symbol_57};
*(tint_symbol_84) = tint_symbol_58;
tint_array_wrapper const tint_symbol_56 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_57 = {.numbers=tint_symbol_56};
*(tint_symbol_84) = tint_symbol_57;
*(tint_symbol_84) = x_775;
float3 const x_453 = float3(x_451[0], x_450[0], x_450[1]);
color[0] = (x_206 + float(x_201));
@ -941,9 +941,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv[0] = 0.0f;
uv[0] = x_778;
QuicksortObject const x_779 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_59 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_60 = {.numbers=tint_symbol_59};
*(tint_symbol_84) = tint_symbol_60;
tint_array_wrapper const tint_symbol_58 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_59 = {.numbers=tint_symbol_58};
*(tint_symbol_84) = tint_symbol_59;
*(tint_symbol_84) = x_779;
if ((x_210 > 0.25f)) {
int const x_780 = i_2;
@ -958,18 +958,18 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv[0] = x_782;
int const x_216 = (*(tint_symbol_84)).numbers.arr[1];
QuicksortObject const x_783 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_61 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_62 = {.numbers=tint_symbol_61};
*(tint_symbol_84) = tint_symbol_62;
tint_array_wrapper const tint_symbol_60 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_61 = {.numbers=tint_symbol_60};
*(tint_symbol_84) = tint_symbol_61;
*(tint_symbol_84) = x_783;
float2 const x_457 = float2(x_454[0], x_454[0]);
float2 const x_784 = uv;
uv = float2(0.0f, 0.0f);
uv = x_784;
QuicksortObject const x_785 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_63 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_64 = {.numbers=tint_symbol_63};
*(tint_symbol_84) = tint_symbol_64;
tint_array_wrapper const tint_symbol_62 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_63 = {.numbers=tint_symbol_62};
*(tint_symbol_84) = tint_symbol_63;
*(tint_symbol_84) = x_785;
float2 const x_458 = float2(float3(1.0f, 2.0f, 3.0f)[2], float2()[1]);
int const x_786 = i_2;
@ -1087,9 +1087,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
color[0] = 0.0f;
color[0] = x_816;
QuicksortObject const x_817 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_65 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_66 = {.numbers=tint_symbol_65};
*(tint_symbol_84) = tint_symbol_66;
tint_array_wrapper const tint_symbol_64 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_65 = {.numbers=tint_symbol_64};
*(tint_symbol_84) = tint_symbol_65;
*(tint_symbol_84) = x_817;
float3 const x_468 = float3(x_467[0], x_467[0], x_467[0]);
float const x_818 = uv[0];
@ -1198,9 +1198,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv[0] = x_844;
float3 const x_482 = float3(x_455[0], x_475[1], x_455[1]);
QuicksortObject const x_845 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_67 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_68 = {.numbers=tint_symbol_67};
*(tint_symbol_84) = tint_symbol_68;
tint_array_wrapper const tint_symbol_66 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_67 = {.numbers=tint_symbol_66};
*(tint_symbol_84) = tint_symbol_67;
*(tint_symbol_84) = x_845;
float const x_846 = uv[1];
uv[1] = 0.0f;
@ -1271,9 +1271,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
(*(tint_symbol_84)).numbers.arr[6u] = x_863;
float2 const x_490 = float2(x_480[2], x_480[2]);
QuicksortObject const x_864 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_69 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_70 = {.numbers=tint_symbol_69};
*(tint_symbol_84) = tint_symbol_70;
tint_array_wrapper const tint_symbol_68 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_69 = {.numbers=tint_symbol_68};
*(tint_symbol_84) = tint_symbol_69;
*(tint_symbol_84) = x_864;
color[1] = (float(x_280) + x_283);
float const x_865 = color[0];
@ -1290,9 +1290,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
color[0] = x_867;
float const x_287 = uv[1];
QuicksortObject const x_868 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_71 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_72 = {.numbers=tint_symbol_71};
*(tint_symbol_84) = tint_symbol_72;
tint_array_wrapper const tint_symbol_70 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_71 = {.numbers=tint_symbol_70};
*(tint_symbol_84) = tint_symbol_71;
*(tint_symbol_84) = x_868;
float2 const x_493 = float2(x_475[0], x_475[1]);
float const x_869 = uv[0];
@ -1452,9 +1452,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv[0] = 0.0f;
uv[0] = x_910;
QuicksortObject const x_911 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_73 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_74 = {.numbers=tint_symbol_73};
*(tint_symbol_84) = tint_symbol_74;
tint_array_wrapper const tint_symbol_72 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_73 = {.numbers=tint_symbol_72};
*(tint_symbol_84) = tint_symbol_73;
*(tint_symbol_84) = x_911;
float3 const x_513 = float3(x_505[2], x_505[0], x_448[0]);
int const x_912 = (*(tint_symbol_84)).numbers.arr[8];
@ -1506,14 +1506,14 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv[0] = 0.0f;
uv[0] = x_923;
QuicksortObject const x_924 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_75 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_76 = {.numbers=tint_symbol_75};
*(tint_symbol_84) = tint_symbol_76;
tint_array_wrapper const tint_symbol_74 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_75 = {.numbers=tint_symbol_74};
*(tint_symbol_84) = tint_symbol_75;
*(tint_symbol_84) = x_924;
QuicksortObject const x_925 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_77 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_78 = {.numbers=tint_symbol_77};
*(tint_symbol_84) = tint_symbol_78;
tint_array_wrapper const tint_symbol_76 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_77 = {.numbers=tint_symbol_76};
*(tint_symbol_84) = tint_symbol_77;
*(tint_symbol_84) = x_925;
float const x_926 = color[1];
color[1] = 0.0f;
@ -1532,9 +1532,9 @@ void main_1(thread QuicksortObject* const tint_symbol_84, thread float4* const t
uv[0] = x_929;
*(tint_symbol_87) = x_330;
QuicksortObject const x_930 = *(tint_symbol_84);
tint_array_wrapper const tint_symbol_79 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_80 = {.numbers=tint_symbol_79};
*(tint_symbol_84) = tint_symbol_80;
tint_array_wrapper const tint_symbol_78 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
QuicksortObject const tint_symbol_79 = {.numbers=tint_symbol_78};
*(tint_symbol_84) = tint_symbol_79;
*(tint_symbol_84) = x_930;
float3 const x_522 = float3(x_330[3], x_330[1], x_493[0]);
float const x_931 = color[0];
@ -1554,8 +1554,8 @@ struct tint_symbol_1 {
main_out tint_symbol_inner(float4 gl_FragCoord_param, thread float4* const tint_symbol_88, thread QuicksortObject* const tint_symbol_89, const constant buf0* const tint_symbol_90, thread float4* const tint_symbol_91) {
*(tint_symbol_88) = gl_FragCoord_param;
main_1(tint_symbol_89, tint_symbol_88, tint_symbol_90, tint_symbol_91);
main_out const tint_symbol_2 = {.x_GLF_color_1=*(tint_symbol_91)};
return tint_symbol_2;
main_out const tint_symbol_80 = {.x_GLF_color_1=*(tint_symbol_91)};
return tint_symbol_80;
}
fragment tint_symbol_1 tint_symbol(const constant buf0* tint_symbol_94 [[buffer(0)]], float4 gl_FragCoord_param [[position]]) {

40
test/tint/bug/tint/913.wgsl.expected.glsl
View File

@ -57,33 +57,33 @@ void tint_symbol_1(uvec3 GlobalInvocationID) {
vec4 srcColor = texelFetch(src_1, ivec2(srcTexCoord), 0);
vec4 dstColor = texelFetch(dst_1, ivec2(dstTexCoord), 0);
if ((uniforms.channelCount == 2u)) {
bool tint_tmp_5 = success;
if (tint_tmp_5) {
tint_tmp_5 = aboutEqual(dstColor.r, srcColor.r);
bool tint_symbol_3 = success;
if (tint_symbol_3) {
tint_symbol_3 = aboutEqual(dstColor.r, srcColor.r);
}
bool tint_tmp_4 = (tint_tmp_5);
if (tint_tmp_4) {
tint_tmp_4 = aboutEqual(dstColor.g, srcColor.g);
bool tint_symbol_2 = tint_symbol_3;
if (tint_symbol_2) {
tint_symbol_2 = aboutEqual(dstColor.g, srcColor.g);
}
success = (tint_tmp_4);
success = tint_symbol_2;
} else {
bool tint_tmp_9 = success;
if (tint_tmp_9) {
tint_tmp_9 = aboutEqual(dstColor.r, srcColor.r);
bool tint_symbol_7 = success;
if (tint_symbol_7) {
tint_symbol_7 = aboutEqual(dstColor.r, srcColor.r);
}
bool tint_tmp_8 = (tint_tmp_9);
if (tint_tmp_8) {
tint_tmp_8 = aboutEqual(dstColor.g, srcColor.g);
bool tint_symbol_6 = tint_symbol_7;
if (tint_symbol_6) {
tint_symbol_6 = aboutEqual(dstColor.g, srcColor.g);
}
bool tint_tmp_7 = (tint_tmp_8);
if (tint_tmp_7) {
tint_tmp_7 = aboutEqual(dstColor.b, srcColor.b);
bool tint_symbol_5 = tint_symbol_6;
if (tint_symbol_5) {
tint_symbol_5 = aboutEqual(dstColor.b, srcColor.b);
}
bool tint_tmp_6 = (tint_tmp_7);
if (tint_tmp_6) {
tint_tmp_6 = aboutEqual(dstColor.a, srcColor.a);
bool tint_symbol_4 = tint_symbol_5;
if (tint_symbol_4) {
tint_symbol_4 = aboutEqual(dstColor.a, srcColor.a);
}
success = (tint_tmp_6);
success = tint_symbol_4;
}
}
uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);

40
test/tint/bug/tint/913.wgsl.expected.hlsl
View File

@ -49,33 +49,33 @@ void main_inner(uint3 GlobalInvocationID) {
const float4 srcColor = src.Load(int3(int2(srcTexCoord), 0));
const float4 dstColor = tint_symbol.Load(int3(int2(dstTexCoord), 0));
if ((uniforms[0].y == 2u)) {
bool tint_tmp_7 = success;
if (tint_tmp_7) {
tint_tmp_7 = aboutEqual(dstColor.r, srcColor.r);
bool tint_symbol_4 = success;
if (tint_symbol_4) {
tint_symbol_4 = aboutEqual(dstColor.r, srcColor.r);
}
bool tint_tmp_6 = (tint_tmp_7);
if (tint_tmp_6) {
tint_tmp_6 = aboutEqual(dstColor.g, srcColor.g);
bool tint_symbol_3 = tint_symbol_4;
if (tint_symbol_3) {
tint_symbol_3 = aboutEqual(dstColor.g, srcColor.g);
}
success = (tint_tmp_6);
success = tint_symbol_3;
} else {
bool tint_tmp_11 = success;
if (tint_tmp_11) {
tint_tmp_11 = aboutEqual(dstColor.r, srcColor.r);
bool tint_symbol_8 = success;
if (tint_symbol_8) {
tint_symbol_8 = aboutEqual(dstColor.r, srcColor.r);
}
bool tint_tmp_10 = (tint_tmp_11);
if (tint_tmp_10) {
tint_tmp_10 = aboutEqual(dstColor.g, srcColor.g);
bool tint_symbol_7 = tint_symbol_8;
if (tint_symbol_7) {
tint_symbol_7 = aboutEqual(dstColor.g, srcColor.g);
}
bool tint_tmp_9 = (tint_tmp_10);
if (tint_tmp_9) {
tint_tmp_9 = aboutEqual(dstColor.b, srcColor.b);
bool tint_symbol_6 = tint_symbol_7;
if (tint_symbol_6) {
tint_symbol_6 = aboutEqual(dstColor.b, srcColor.b);
}
bool tint_tmp_8 = (tint_tmp_9);
if (tint_tmp_8) {
tint_tmp_8 = aboutEqual(dstColor.a, srcColor.a);
bool tint_symbol_5 = tint_symbol_6;
if (tint_symbol_5) {
tint_symbol_5 = aboutEqual(dstColor.a, srcColor.a);
}
success = (tint_tmp_8);
success = tint_symbol_5;
}
}
const uint outputIndex = ((GlobalInvocationID.y * uint(dstSize.x)) + GlobalInvocationID.x);

56
test/tint/bug/tint/913.wgsl.expected.msl
View File

@ -17,37 +17,61 @@ bool aboutEqual(float value, float expect) {
return (fabs((value - expect)) < 0.001f);
}
void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2, const constant Uniforms* const tint_symbol_3, device OutputBuf* const tint_symbol_4) {
int2 const srcSize = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
int2 const dstSize = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
void tint_symbol_inner(uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_7, texture2d<float, access::sample> tint_symbol_8, const constant Uniforms* const tint_symbol_9, device OutputBuf* const tint_symbol_10) {
int2 const srcSize = int2(tint_symbol_7.get_width(), tint_symbol_7.get_height());
int2 const dstSize = int2(tint_symbol_8.get_width(), tint_symbol_8.get_height());
uint2 const dstTexCoord = uint2(uint3(GlobalInvocationID).xy);
float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
bool success = true;
if (((((dstTexCoord[0] < (*(tint_symbol_3)).dstCopyOrigin[0]) || (dstTexCoord[1] < (*(tint_symbol_3)).dstCopyOrigin[1])) || (dstTexCoord[0] >= ((*(tint_symbol_3)).dstCopyOrigin[0] + (*(tint_symbol_3)).copySize[0]))) || (dstTexCoord[1] >= ((*(tint_symbol_3)).dstCopyOrigin[1] + (*(tint_symbol_3)).copySize[1])))) {
success = (success && all((tint_symbol_2.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
if (((((dstTexCoord[0] < (*(tint_symbol_9)).dstCopyOrigin[0]) || (dstTexCoord[1] < (*(tint_symbol_9)).dstCopyOrigin[1])) || (dstTexCoord[0] >= ((*(tint_symbol_9)).dstCopyOrigin[0] + (*(tint_symbol_9)).copySize[0]))) || (dstTexCoord[1] >= ((*(tint_symbol_9)).dstCopyOrigin[1] + (*(tint_symbol_9)).copySize[1])))) {
success = (success && all((tint_symbol_8.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
} else {
uint2 srcTexCoord = ((dstTexCoord - (*(tint_symbol_3)).dstCopyOrigin) + (*(tint_symbol_3)).srcCopyOrigin);
if (((*(tint_symbol_3)).dstTextureFlipY == 1u)) {
uint2 srcTexCoord = ((dstTexCoord - (*(tint_symbol_9)).dstCopyOrigin) + (*(tint_symbol_9)).srcCopyOrigin);
if (((*(tint_symbol_9)).dstTextureFlipY == 1u)) {
srcTexCoord[1] = ((uint(srcSize[1]) - srcTexCoord[1]) - 1u);
}
float4 const srcColor = tint_symbol_1.read(uint2(int2(srcTexCoord)), 0);
float4 const dstColor = tint_symbol_2.read(uint2(int2(dstTexCoord)), 0);
if (((*(tint_symbol_3)).channelCount == 2u)) {
success = ((success && aboutEqual(dstColor[0], srcColor[0])) && aboutEqual(dstColor[1], srcColor[1]));
float4 const srcColor = tint_symbol_7.read(uint2(int2(srcTexCoord)), 0);
float4 const dstColor = tint_symbol_8.read(uint2(int2(dstTexCoord)), 0);
if (((*(tint_symbol_9)).channelCount == 2u)) {
bool tint_symbol_2 = success;
if (tint_symbol_2) {
tint_symbol_2 = aboutEqual(dstColor[0], srcColor[0]);
}
bool tint_symbol_1 = tint_symbol_2;
if (tint_symbol_1) {
tint_symbol_1 = aboutEqual(dstColor[1], srcColor[1]);
}
success = tint_symbol_1;
} else {
success = ((((success && aboutEqual(dstColor[0], srcColor[0])) && aboutEqual(dstColor[1], srcColor[1])) && aboutEqual(dstColor[2], srcColor[2])) && aboutEqual(dstColor[3], srcColor[3]));
bool tint_symbol_6 = success;
if (tint_symbol_6) {
tint_symbol_6 = aboutEqual(dstColor[0], srcColor[0]);
}
bool tint_symbol_5 = tint_symbol_6;
if (tint_symbol_5) {
tint_symbol_5 = aboutEqual(dstColor[1], srcColor[1]);
}
bool tint_symbol_4 = tint_symbol_5;
if (tint_symbol_4) {
tint_symbol_4 = aboutEqual(dstColor[2], srcColor[2]);
}
bool tint_symbol_3 = tint_symbol_4;
if (tint_symbol_3) {
tint_symbol_3 = aboutEqual(dstColor[3], srcColor[3]);
}
success = tint_symbol_3;
}
}
uint const outputIndex = ((GlobalInvocationID[1] * uint(dstSize[0])) + GlobalInvocationID[0]);
if (success) {
(*(tint_symbol_4)).result[outputIndex] = 1u;
(*(tint_symbol_10)).result[outputIndex] = 1u;
} else {
(*(tint_symbol_4)).result[outputIndex] = 0u;
(*(tint_symbol_10)).result[outputIndex] = 0u;
}
}
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]], texture2d<float, access::sample> tint_symbol_6 [[texture(1)]], const constant Uniforms* tint_symbol_7 [[buffer(0)]], device OutputBuf* tint_symbol_8 [[buffer(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
tint_symbol_inner(GlobalInvocationID, tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_11 [[texture(0)]], texture2d<float, access::sample> tint_symbol_12 [[texture(1)]], const constant Uniforms* tint_symbol_13 [[buffer(0)]], device OutputBuf* tint_symbol_14 [[buffer(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
tint_symbol_inner(GlobalInvocationID, tint_symbol_11, tint_symbol_12, tint_symbol_13, tint_symbol_14);
return;
}

180
test/tint/bug/tint/913.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 193
; Bound: 205
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
@ -29,6 +29,12 @@
OpName %GlobalInvocationID "GlobalInvocationID"
OpName %success "success"
OpName %srcTexCoord "srcTexCoord"
OpName %tint_symbol_1 "tint_symbol_1"
OpName %tint_symbol "tint_symbol"
OpName %tint_symbol_5 "tint_symbol_5"
OpName %tint_symbol_4 "tint_symbol_4"
OpName %tint_symbol_3 "tint_symbol_3"
OpName %tint_symbol_2 "tint_symbol_2"
OpName %main "main"
OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
OpDecorate %src DescriptorSet 0
@ -93,7 +99,7 @@
%110 = OpConstantNull %v2uint
%_ptr_Function_uint = OpTypePointer Function %uint
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%188 = OpTypeFunction %void
%200 = OpTypeFunction %void
%aboutEqual = OpFunction %bool None %18
%value = OpFunctionParameter %float
%expect = OpFunctionParameter %float
@ -108,6 +114,12 @@
%33 = OpLabel
%success = OpVariable %_ptr_Function_bool Function %50
%srcTexCoord = OpVariable %_ptr_Function_v2uint Function %110
%tint_symbol_1 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_5 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_4 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_3 = OpVariable %_ptr_Function_bool Function %50
%tint_symbol_2 = OpVariable %_ptr_Function_bool Function %50
%37 = OpLoad %7 %src
%34 = OpImageQuerySizeLod %v2int %37 %int_0
%40 = OpLoad %7 %dst
@ -209,92 +221,110 @@
OpBranchConditional %133 %135 %136
%135 = OpLabel
%137 = OpLoad %bool %success
OpSelectionMerge %138 None
OpBranchConditional %137 %139 %138
%139 = OpLabel
%141 = OpCompositeExtract %float %128 0
%142 = OpCompositeExtract %float %124 0
%140 = OpFunctionCall %bool %aboutEqual %141 %142
OpBranch %138
%138 = OpLabel
%143 = OpPhi %bool %137 %135 %140 %139
OpSelectionMerge %144 None
OpBranchConditional %143 %145 %144
%145 = OpLabel
%147 = OpCompositeExtract %float %128 1
%148 = OpCompositeExtract %float %124 1
%146 = OpFunctionCall %bool %aboutEqual %147 %148
OpBranch %144
%144 = OpLabel
%149 = OpPhi %bool %143 %138 %146 %145
OpStore %success %149
OpStore %tint_symbol_1 %137
%139 = OpLoad %bool %tint_symbol_1
OpSelectionMerge %140 None
OpBranchConditional %139 %141 %140
%141 = OpLabel
%143 = OpCompositeExtract %float %128 0
%144 = OpCompositeExtract %float %124 0
%142 = OpFunctionCall %bool %aboutEqual %143 %144
OpStore %tint_symbol_1 %142
OpBranch %140
%140 = OpLabel
%145 = OpLoad %bool %tint_symbol_1
OpStore %tint_symbol %145
%147 = OpLoad %bool %tint_symbol
OpSelectionMerge %148 None
OpBranchConditional %147 %149 %148
%149 = OpLabel
%151 = OpCompositeExtract %float %128 1
%152 = OpCompositeExtract %float %124 1
%150 = OpFunctionCall %bool %aboutEqual %151 %152
OpStore %tint_symbol %150
OpBranch %148
%148 = OpLabel
%153 = OpLoad %bool %tint_symbol
OpStore %success %153
OpBranch %134
%136 = OpLabel
%150 = OpLoad %bool %success
OpSelectionMerge %151 None
OpBranchConditional %150 %152 %151
%152 = OpLabel
%154 = OpCompositeExtract %float %128 0
%155 = OpCompositeExtract %float %124 0
%153 = OpFunctionCall %bool %aboutEqual %154 %155
OpBranch %151
%151 = OpLabel
%156 = OpPhi %bool %150 %136 %153 %152
%154 = OpLoad %bool %success
OpStore %tint_symbol_5 %154
%156 = OpLoad %bool %tint_symbol_5
OpSelectionMerge %157 None
OpBranchConditional %156 %158 %157
%158 = OpLabel
%160 = OpCompositeExtract %float %128 1
%161 = OpCompositeExtract %float %124 1
%160 = OpCompositeExtract %float %128 0
%161 = OpCompositeExtract %float %124 0
%159 = OpFunctionCall %bool %aboutEqual %160 %161
OpStore %tint_symbol_5 %159
OpBranch %157
%157 = OpLabel
%162 = OpPhi %bool %156 %151 %159 %158
OpSelectionMerge %163 None
OpBranchConditional %162 %164 %163
%164 = OpLabel
%166 = OpCompositeExtract %float %128 2
%167 = OpCompositeExtract %float %124 2
%165 = OpFunctionCall %bool %aboutEqual %166 %167
OpBranch %163
%163 = OpLabel
%168 = OpPhi %bool %162 %157 %165 %164
OpSelectionMerge %169 None
OpBranchConditional %168 %170 %169
%170 = OpLabel
%172 = OpCompositeExtract %float %128 3
%173 = OpCompositeExtract %float %124 3
%171 = OpFunctionCall %bool %aboutEqual %172 %173
OpBranch %169
%169 = OpLabel
%174 = OpPhi %bool %168 %163 %171 %170
OpStore %success %174
%162 = OpLoad %bool %tint_symbol_5
OpStore %tint_symbol_4 %162
%164 = OpLoad %bool %tint_symbol_4
OpSelectionMerge %165 None
OpBranchConditional %164 %166 %165
%166 = OpLabel
%168 = OpCompositeExtract %float %128 1
%169 = OpCompositeExtract %float %124 1
%167 = OpFunctionCall %bool %aboutEqual %168 %169
OpStore %tint_symbol_4 %167
OpBranch %165
%165 = OpLabel
%170 = OpLoad %bool %tint_symbol_4
OpStore %tint_symbol_3 %170
%172 = OpLoad %bool %tint_symbol_3
OpSelectionMerge %173 None
OpBranchConditional %172 %174 %173
%174 = OpLabel
%176 = OpCompositeExtract %float %128 2
%177 = OpCompositeExtract %float %124 2
%175 = OpFunctionCall %bool %aboutEqual %176 %177
OpStore %tint_symbol_3 %175
OpBranch %173
%173 = OpLabel
%178 = OpLoad %bool %tint_symbol_3
OpStore %tint_symbol_2 %178
%180 = OpLoad %bool %tint_symbol_2
OpSelectionMerge %181 None
OpBranchConditional %180 %182 %181
%182 = OpLabel
%184 = OpCompositeExtract %float %128 3
%185 = OpCompositeExtract %float %124 3
%183 = OpFunctionCall %bool %aboutEqual %184 %185
OpStore %tint_symbol_2 %183
OpBranch %181
%181 = OpLabel
%186 = OpLoad %bool %tint_symbol_2
OpStore %success %186
OpBranch %134
%134 = OpLabel
OpBranch %87
%87 = OpLabel
%175 = OpCompositeExtract %uint %GlobalInvocationID 1
%177 = OpCompositeExtract %int %39 0
%176 = OpBitcast %uint %177
%178 = OpIMul %uint %175 %176
%179 = OpCompositeExtract %uint %GlobalInvocationID 0
%180 = OpIAdd %uint %178 %179
%181 = OpLoad %bool %success
OpSelectionMerge %182 None
OpBranchConditional %181 %183 %184
%183 = OpLabel
%186 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %180
OpStore %186 %uint_1
OpBranch %182
%184 = OpLabel
%187 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %180
OpStore %187 %uint_0
OpBranch %182
%182 = OpLabel
%187 = OpCompositeExtract %uint %GlobalInvocationID 1
%189 = OpCompositeExtract %int %39 0
%188 = OpBitcast %uint %189
%190 = OpIMul %uint %187 %188
%191 = OpCompositeExtract %uint %GlobalInvocationID 0
%192 = OpIAdd %uint %190 %191
%193 = OpLoad %bool %success
OpSelectionMerge %194 None
OpBranchConditional %193 %195 %196
%195 = OpLabel
%198 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %192
OpStore %198 %uint_1
OpBranch %194
%196 = OpLabel
%199 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %192
OpStore %199 %uint_0
OpBranch %194
%194 = OpLabel
OpReturn
OpFunctionEnd
%main = OpFunction %void None %188
%190 = OpLabel
%192 = OpLoad %v3uint %GlobalInvocationID_1
%191 = OpFunctionCall %void %main_inner %192
%main = OpFunction %void None %200
%202 = OpLabel
%204 = OpLoad %v3uint %GlobalInvocationID_1
%203 = OpFunctionCall %void %main_inner %204
OpReturn
OpFunctionEnd

6
test/tint/bug/tint/914.wgsl.expected.glsl
View File

@ -96,7 +96,8 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
uint inputRow = (tileRow + innerRow);
uint inputCol = (tileColA + innerCol);
mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
float tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
mm_Asub[inputRow][inputCol] = tint_symbol_1;
}
}
}
@ -107,7 +108,8 @@ void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
uint inputRow = (tileRowB + innerRow);
uint inputCol = (tileCol + innerCol);
mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
float tint_symbol_2 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
mm_Bsub[innerCol][inputCol] = tint_symbol_2;
}
}
}

6
test/tint/bug/tint/914.wgsl.expected.hlsl
View File

@ -89,7 +89,8 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
[loop] for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRow + innerRow);
const uint inputCol = (tileColA + innerCol);
mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
const float tint_symbol_2 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
mm_Asub[inputRow][inputCol] = tint_symbol_2;
}
}
}
@ -100,7 +101,8 @@ void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
[loop] for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRowB + innerRow);
const uint inputCol = (tileCol + innerCol);
mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
const float tint_symbol_3 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
mm_Bsub[innerCol][inputCol] = tint_symbol_3;
}
}
}

48
test/tint/bug/tint/914.wgsl.expected.msl
View File

@ -11,26 +11,26 @@ struct Matrix {
/* 0x0000 */ float numbers[1];
};
float mm_readA(uint row, uint col, const constant Uniforms* const tint_symbol_1, const device Matrix* const tint_symbol_2) {
if (((row < (*(tint_symbol_1)).dimAOuter) && (col < (*(tint_symbol_1)).dimInner))) {
float const result = (*(tint_symbol_2)).numbers[((row * (*(tint_symbol_1)).dimInner) + col)];
float mm_readA(uint row, uint col, const constant Uniforms* const tint_symbol_3, const device Matrix* const tint_symbol_4) {
if (((row < (*(tint_symbol_3)).dimAOuter) && (col < (*(tint_symbol_3)).dimInner))) {
float const result = (*(tint_symbol_4)).numbers[((row * (*(tint_symbol_3)).dimInner) + col)];
return result;
}
return 0.0f;
}
float mm_readB(uint row, uint col, const constant Uniforms* const tint_symbol_3, const device Matrix* const tint_symbol_4) {
if (((row < (*(tint_symbol_3)).dimInner) && (col < (*(tint_symbol_3)).dimBOuter))) {
float const result = (*(tint_symbol_4)).numbers[((row * (*(tint_symbol_3)).dimBOuter) + col)];
float mm_readB(uint row, uint col, const constant Uniforms* const tint_symbol_5, const device Matrix* const tint_symbol_6) {
if (((row < (*(tint_symbol_5)).dimInner) && (col < (*(tint_symbol_5)).dimBOuter))) {
float const result = (*(tint_symbol_6)).numbers[((row * (*(tint_symbol_5)).dimBOuter) + col)];
return result;
}
return 0.0f;
}
void mm_write(uint row, uint col, float value, const constant Uniforms* const tint_symbol_5, device Matrix* const tint_symbol_6) {
if (((row < (*(tint_symbol_5)).dimAOuter) && (col < (*(tint_symbol_5)).dimBOuter))) {
uint const index = (col + (row * (*(tint_symbol_5)).dimBOuter));
(*(tint_symbol_6)).numbers[index] = value;
void mm_write(uint row, uint col, float value, const constant Uniforms* const tint_symbol_7, device Matrix* const tint_symbol_8) {
if (((row < (*(tint_symbol_7)).dimAOuter) && (col < (*(tint_symbol_7)).dimBOuter))) {
uint const index = (col + (row * (*(tint_symbol_7)).dimBOuter));
(*(tint_symbol_8)).numbers[index] = value;
}
}
@ -60,19 +60,19 @@ struct tint_array_wrapper_3 {
float arr[4];
};
void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array_wrapper* const tint_symbol_7, threadgroup tint_array_wrapper* const tint_symbol_8, const constant Uniforms* const tint_symbol_9, const device Matrix* const tint_symbol_10, const device Matrix* const tint_symbol_11, device Matrix* const tint_symbol_12) {
void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array_wrapper* const tint_symbol_9, threadgroup tint_array_wrapper* const tint_symbol_10, const constant Uniforms* const tint_symbol_11, const device Matrix* const tint_symbol_12, const device Matrix* const tint_symbol_13, device Matrix* const tint_symbol_14) {
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
uint const i = (idx / 64u);
uint const i_1 = (idx % 64u);
(*(tint_symbol_7)).arr[i].arr[i_1] = float();
(*(tint_symbol_8)).arr[i].arr[i_1] = float();
(*(tint_symbol_9)).arr[i].arr[i_1] = float();
(*(tint_symbol_10)).arr[i].arr[i_1] = float();
}
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const tileRow = (local_id[1] * RowPerThread);
uint const tileCol = (local_id[0] * ColPerThread);
uint const globalRow = (global_id[1] * RowPerThread);
uint const globalCol = (global_id[0] * ColPerThread);
uint const numTiles = ((((*(tint_symbol_9)).dimInner - 1u) / TileInner) + 1u);
uint const numTiles = ((((*(tint_symbol_11)).dimInner - 1u) / TileInner) + 1u);
tint_array_wrapper_2 acc = {};
float ACached = 0.0f;
tint_array_wrapper_3 BCached = {};
@ -88,23 +88,25 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in
for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
uint const inputRow = (tileRow + innerRow);
uint const inputCol = (tileColA + innerCol);
(*(tint_symbol_7)).arr[inputRow].arr[inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol), tint_symbol_9, tint_symbol_10);
float const tint_symbol_1 = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol), tint_symbol_11, tint_symbol_12);
(*(tint_symbol_9)).arr[inputRow].arr[inputCol] = tint_symbol_1;
}
}
for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
uint const inputRow = (tileRowB + innerRow);
uint const inputCol = (tileCol + innerCol);
(*(tint_symbol_8)).arr[innerCol].arr[inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol), tint_symbol_9, tint_symbol_11);
float const tint_symbol_2 = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol), tint_symbol_11, tint_symbol_13);
(*(tint_symbol_10)).arr[innerCol].arr[inputCol] = tint_symbol_2;
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
BCached.arr[inner] = (*(tint_symbol_8)).arr[k].arr[(tileCol + inner)];
BCached.arr[inner] = (*(tint_symbol_10)).arr[k].arr[(tileCol + inner)];
}
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
ACached = (*(tint_symbol_7)).arr[(tileRow + innerRow)].arr[k];
ACached = (*(tint_symbol_9)).arr[(tileRow + innerRow)].arr[k];
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * ColPerThread) + innerCol);
acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
@ -116,15 +118,15 @@ void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_in
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * ColPerThread) + innerCol);
mm_write((globalRow + innerRow), (globalCol + innerCol), acc.arr[index], tint_symbol_9, tint_symbol_12);
mm_write((globalRow + innerRow), (globalCol + innerCol), acc.arr[index], tint_symbol_11, tint_symbol_14);
}
}
}
kernel void tint_symbol(const constant Uniforms* tint_symbol_15 [[buffer(0)]], const device Matrix* tint_symbol_16 [[buffer(2)]], const device Matrix* tint_symbol_17 [[buffer(3)]], device Matrix* tint_symbol_18 [[buffer(1)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup tint_array_wrapper tint_symbol_13;
threadgroup tint_array_wrapper tint_symbol_14;
tint_symbol_inner(local_id, global_id, local_invocation_index, &(tint_symbol_13), &(tint_symbol_14), tint_symbol_15, tint_symbol_16, tint_symbol_17, tint_symbol_18);
kernel void tint_symbol(const constant Uniforms* tint_symbol_17 [[buffer(0)]], const device Matrix* tint_symbol_18 [[buffer(2)]], const device Matrix* tint_symbol_19 [[buffer(3)]], device Matrix* tint_symbol_20 [[buffer(1)]], uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup tint_array_wrapper tint_symbol_15;
threadgroup tint_array_wrapper tint_symbol_16;
tint_symbol_inner(local_id, global_id, local_invocation_index, &(tint_symbol_15), &(tint_symbol_16), tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
return;
}

32
test/tint/bug/tint/914.wgsl.expected.spvasm
View File

@ -364,14 +364,14 @@
%206 = OpIAdd %uint %131 %205
%207 = OpLoad %uint %innerCol
%208 = OpIAdd %uint %171 %207
%209 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %206 %208
%211 = OpLoad %uint %innerRow
%212 = OpIAdd %uint %135 %211
%213 = OpLoad %uint %t
%214 = OpIMul %uint %213 %TileAOuter
%215 = OpIAdd %uint %214 %208
%210 = OpFunctionCall %float %mm_readA %212 %215
OpStore %209 %210
%210 = OpLoad %uint %innerRow
%211 = OpIAdd %uint %135 %210
%212 = OpLoad %uint %t
%213 = OpIMul %uint %212 %TileAOuter
%214 = OpIAdd %uint %213 %208
%209 = OpFunctionCall %float %mm_readA %211 %214
%215 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %206 %208
OpStore %215 %209
OpBranch %198
%198 = OpLabel
%216 = OpLoad %uint %innerCol
@ -418,15 +418,15 @@
%241 = OpIAdd %uint %174 %240
%242 = OpLoad %uint %innerCol_0
%243 = OpIAdd %uint %133 %242
%244 = OpLoad %uint %innerCol_0
%245 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %244 %243
%247 = OpLoad %uint %t
%248 = OpIMul %uint %247 %TileAOuter
%249 = OpIAdd %uint %248 %241
%245 = OpLoad %uint %t
%246 = OpIMul %uint %245 %TileAOuter
%247 = OpIAdd %uint %246 %241
%248 = OpLoad %uint %innerCol_0
%249 = OpIAdd %uint %137 %248
%244 = OpFunctionCall %float %mm_readB %247 %249
%250 = OpLoad %uint %innerCol_0
%251 = OpIAdd %uint %137 %250
%246 = OpFunctionCall %float %mm_readB %249 %251
OpStore %245 %246
%251 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %250 %243
OpStore %251 %244
OpBranch %233
%233 = OpLabel
%252 = OpLoad %uint %innerCol_0

3
test/tint/bug/tint/980.wgsl.expected.glsl
View File

@ -16,7 +16,8 @@ layout(binding = 0, std430) buffer S_1 {
uint i;
} io;
void tint_symbol(uint idx) {
io.v = Bad(io.i, io.v);
vec3 tint_symbol_1 = Bad(io.i, io.v);
io.v = tint_symbol_1;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

3
test/tint/bug/tint/980.wgsl.expected.hlsl
View File

@ -15,7 +15,8 @@ struct tint_symbol_1 {
};
void main_inner(uint idx) {
io.Store3(0u, asuint(Bad(io.Load(12u), asfloat(io.Load3(0u)))));
const float3 tint_symbol_2 = Bad(io.Load(12u), asfloat(io.Load3(0u)));
io.Store3(0u, asuint(tint_symbol_2));
}
[numthreads(1, 1, 1)]

9
test/tint/bug/tint/980.wgsl.expected.msl
View File

@ -23,12 +23,13 @@ struct S {
/* 0x000c */ uint i;
};
void tint_symbol_inner(uint idx, device S* const tint_symbol_1) {
(*(tint_symbol_1)).v = Bad((*(tint_symbol_1)).i, (*(tint_symbol_1)).v);
void tint_symbol_inner(uint idx, device S* const tint_symbol_2) {
float3 const tint_symbol_1 = Bad((*(tint_symbol_2)).i, (*(tint_symbol_2)).v);
(*(tint_symbol_2)).v = tint_symbol_1;
}
kernel void tint_symbol(device S* tint_symbol_2 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
tint_symbol_inner(idx, tint_symbol_2);
kernel void tint_symbol(device S* tint_symbol_3 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
tint_symbol_inner(idx, tint_symbol_3);
return;
}

18
test/tint/bug/tint/980.wgsl.expected.spvasm
View File

@ -42,10 +42,10 @@
%_ptr_Function_float = OpTypePointer Function %float
%void = OpTypeVoid
%27 = OpTypeFunction %void %uint
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%uint_1 = OpConstant %uint 1
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
%42 = OpTypeFunction %void
%Bad = OpFunction %v3float None %9
%index = OpFunctionParameter %uint
@ -65,13 +65,13 @@
%main_inner = OpFunction %void None %27
%idx = OpFunctionParameter %uint
%31 = OpLabel
%34 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0
%38 = OpAccessChain %_ptr_StorageBuffer_uint %io %uint_1
%39 = OpLoad %uint %38
%40 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0
%41 = OpLoad %v3float %40
%35 = OpFunctionCall %v3float %Bad %39 %41
OpStore %34 %35
%35 = OpAccessChain %_ptr_StorageBuffer_uint %io %uint_1
%36 = OpLoad %uint %35
%39 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0
%40 = OpLoad %v3float %39
%32 = OpFunctionCall %v3float %Bad %36 %40
%41 = OpAccessChain %_ptr_StorageBuffer_v3float %io %uint_0
OpStore %41 %32
OpReturn
OpFunctionEnd
%main = OpFunction %void None %42

4
test/tint/bug/tint/993.wgsl.expected.glsl
View File

@ -27,7 +27,9 @@ int runTest() {
}
void tint_symbol() {
result.value = uint(runTest());
int tint_symbol_1 = runTest();
uint tint_symbol_2 = uint(tint_symbol_1);
result.value = tint_symbol_2;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

4
test/tint/bug/tint/993.wgsl.expected.hlsl
View File

@ -18,6 +18,8 @@ int runTest() {
[numthreads(1, 1, 1)]
void main() {
result.Store(0u, asuint(uint(runTest())));
const int tint_symbol = runTest();
const uint tint_symbol_1 = uint(tint_symbol);
result.Store(0u, asuint(tint_symbol_1));
return;
}

10
test/tint/bug/tint/993.wgsl.expected.msl
View File

@ -17,12 +17,14 @@ struct TestData {
/* 0x0000 */ tint_array_wrapper data;
};
int runTest(device TestData* const tint_symbol_1, const constant Constants* const tint_symbol_2) {
return atomic_load_explicit(&((*(tint_symbol_1)).data.arr[(0u + uint((*(tint_symbol_2)).zero))]), memory_order_relaxed);
int runTest(device TestData* const tint_symbol_3, const constant Constants* const tint_symbol_4) {
return atomic_load_explicit(&((*(tint_symbol_3)).data.arr[(0u + uint((*(tint_symbol_4)).zero))]), memory_order_relaxed);
}
kernel void tint_symbol(device Result* tint_symbol_3 [[buffer(1)]], device TestData* tint_symbol_4 [[buffer(2)]], const constant Constants* tint_symbol_5 [[buffer(0)]]) {
(*(tint_symbol_3)).value = uint(runTest(tint_symbol_4, tint_symbol_5));
kernel void tint_symbol(device TestData* tint_symbol_5 [[buffer(2)]], const constant Constants* tint_symbol_6 [[buffer(0)]], device Result* tint_symbol_7 [[buffer(1)]]) {
int const tint_symbol_1 = runTest(tint_symbol_5, tint_symbol_6);
uint const tint_symbol_2 = uint(tint_symbol_1);
(*(tint_symbol_7)).value = tint_symbol_2;
return;
}

8
test/tint/bug/tint/993.wgsl.expected.spvasm
View File

@ -65,9 +65,9 @@
OpFunctionEnd
%main = OpFunction %void None %28
%31 = OpLabel
%33 = OpAccessChain %_ptr_StorageBuffer_uint %result %uint_0
%35 = OpFunctionCall %int %runTest
%34 = OpBitcast %uint %35
OpStore %33 %34
%32 = OpFunctionCall %int %runTest
%33 = OpBitcast %uint %32
%35 = OpAccessChain %_ptr_StorageBuffer_uint %result %uint_0
OpStore %35 %33
OpReturn
OpFunctionEnd

9
test/tint/expressions/splat/call/bool.wgsl.expected.glsl
View File

@ -9,8 +9,11 @@ bool get_bool() {
}
void f() {
bvec2 v2 = bvec2(get_bool());
bvec3 v3 = bvec3(get_bool());
bvec4 v4 = bvec4(get_bool());
bool tint_symbol = get_bool();
bvec2 v2 = bvec2(tint_symbol);
bool tint_symbol_1 = get_bool();
bvec3 v3 = bvec3(tint_symbol_1);
bool tint_symbol_2 = get_bool();
bvec4 v4 = bvec4(tint_symbol_2);
}

9
test/tint/expressions/splat/call/bool.wgsl.expected.hlsl
View File

@ -8,7 +8,10 @@ bool get_bool() {
}
void f() {
bool2 v2 = bool2((get_bool()).xx);
bool3 v3 = bool3((get_bool()).xxx);
bool4 v4 = bool4((get_bool()).xxxx);
const bool tint_symbol = get_bool();
bool2 v2 = bool2((tint_symbol).xx);
const bool tint_symbol_1 = get_bool();
bool3 v3 = bool3((tint_symbol_1).xxx);
const bool tint_symbol_2 = get_bool();
bool4 v4 = bool4((tint_symbol_2).xxxx);
}

9
test/tint/expressions/splat/call/bool.wgsl.expected.msl
View File

@ -6,8 +6,11 @@ bool get_bool() {
}
void f() {
bool2 v2 = bool2(get_bool());
bool3 v3 = bool3(get_bool());
bool4 v4 = bool4(get_bool());
bool const tint_symbol = get_bool();
bool2 v2 = bool2(tint_symbol);
bool const tint_symbol_1 = get_bool();
bool3 v3 = bool3(tint_symbol_1);
bool const tint_symbol_2 = get_bool();
bool4 v4 = bool4(tint_symbol_2);
}

12
test/tint/expressions/splat/call/bool.wgsl.expected.spvasm
View File

@ -40,14 +40,14 @@
%v2 = OpVariable %_ptr_Function_v2bool Function %17
%v3 = OpVariable %_ptr_Function_v3bool Function %23
%v4 = OpVariable %_ptr_Function_v4bool Function %29
%13 = OpFunctionCall %bool %get_bool
%14 = OpCompositeConstruct %v2bool %13 %13
%12 = OpFunctionCall %bool %get_bool
%14 = OpCompositeConstruct %v2bool %12 %12
OpStore %v2 %14
%19 = OpFunctionCall %bool %get_bool
%20 = OpCompositeConstruct %v3bool %19 %19 %19
%18 = OpFunctionCall %bool %get_bool
%20 = OpCompositeConstruct %v3bool %18 %18 %18
OpStore %v3 %20
%25 = OpFunctionCall %bool %get_bool
%26 = OpCompositeConstruct %v4bool %25 %25 %25 %25
%24 = OpFunctionCall %bool %get_bool
%26 = OpCompositeConstruct %v4bool %24 %24 %24 %24
OpStore %v4 %26
OpReturn
OpFunctionEnd

9
test/tint/expressions/splat/call/f32.wgsl.expected.glsl
View File

@ -9,8 +9,11 @@ float get_f32() {
}
void f() {
vec2 v2 = vec2(get_f32());
vec3 v3 = vec3(get_f32());
vec4 v4 = vec4(get_f32());
float tint_symbol = get_f32();
vec2 v2 = vec2(tint_symbol);
float tint_symbol_1 = get_f32();
vec3 v3 = vec3(tint_symbol_1);
float tint_symbol_2 = get_f32();
vec4 v4 = vec4(tint_symbol_2);
}

9
test/tint/expressions/splat/call/f32.wgsl.expected.hlsl
View File

@ -8,7 +8,10 @@ float get_f32() {
}
void f() {
float2 v2 = float2((get_f32()).xx);
float3 v3 = float3((get_f32()).xxx);
float4 v4 = float4((get_f32()).xxxx);
const float tint_symbol = get_f32();
float2 v2 = float2((tint_symbol).xx);
const float tint_symbol_1 = get_f32();
float3 v3 = float3((tint_symbol_1).xxx);
const float tint_symbol_2 = get_f32();
float4 v4 = float4((tint_symbol_2).xxxx);
}

9
test/tint/expressions/splat/call/f32.wgsl.expected.msl
View File

@ -6,8 +6,11 @@ float get_f32() {
}
void f() {
float2 v2 = float2(get_f32());
float3 v3 = float3(get_f32());
float4 v4 = float4(get_f32());
float const tint_symbol = get_f32();
float2 v2 = float2(tint_symbol);
float const tint_symbol_1 = get_f32();
float3 v3 = float3(tint_symbol_1);
float const tint_symbol_2 = get_f32();
float4 v4 = float4(tint_symbol_2);
}

12
test/tint/expressions/splat/call/f32.wgsl.expected.spvasm
View File

@ -40,14 +40,14 @@
%v2 = OpVariable %_ptr_Function_v2float Function %17
%v3 = OpVariable %_ptr_Function_v3float Function %23
%v4 = OpVariable %_ptr_Function_v4float Function %29
%13 = OpFunctionCall %float %get_f32
%14 = OpCompositeConstruct %v2float %13 %13
%12 = OpFunctionCall %float %get_f32
%14 = OpCompositeConstruct %v2float %12 %12
OpStore %v2 %14
%19 = OpFunctionCall %float %get_f32
%20 = OpCompositeConstruct %v3float %19 %19 %19
%18 = OpFunctionCall %float %get_f32
%20 = OpCompositeConstruct %v3float %18 %18 %18
OpStore %v3 %20
%25 = OpFunctionCall %float %get_f32
%26 = OpCompositeConstruct %v4float %25 %25 %25 %25
%24 = OpFunctionCall %float %get_f32
%26 = OpCompositeConstruct %v4float %24 %24 %24 %24
OpStore %v4 %26
OpReturn
OpFunctionEnd

9
test/tint/expressions/splat/call/i32.wgsl.expected.glsl
View File

@ -9,8 +9,11 @@ int get_i32() {
}
void f() {
ivec2 v2 = ivec2(get_i32());
ivec3 v3 = ivec3(get_i32());
ivec4 v4 = ivec4(get_i32());
int tint_symbol = get_i32();
ivec2 v2 = ivec2(tint_symbol);
int tint_symbol_1 = get_i32();
ivec3 v3 = ivec3(tint_symbol_1);
int tint_symbol_2 = get_i32();
ivec4 v4 = ivec4(tint_symbol_2);
}

9
test/tint/expressions/splat/call/i32.wgsl.expected.hlsl
View File

@ -8,7 +8,10 @@ int get_i32() {
}
void f() {
int2 v2 = int2((get_i32()).xx);
int3 v3 = int3((get_i32()).xxx);
int4 v4 = int4((get_i32()).xxxx);
const int tint_symbol = get_i32();
int2 v2 = int2((tint_symbol).xx);
const int tint_symbol_1 = get_i32();
int3 v3 = int3((tint_symbol_1).xxx);
const int tint_symbol_2 = get_i32();
int4 v4 = int4((tint_symbol_2).xxxx);
}

9
test/tint/expressions/splat/call/i32.wgsl.expected.msl
View File

@ -6,8 +6,11 @@ int get_i32() {
}
void f() {
int2 v2 = int2(get_i32());
int3 v3 = int3(get_i32());
int4 v4 = int4(get_i32());
int const tint_symbol = get_i32();
int2 v2 = int2(tint_symbol);
int const tint_symbol_1 = get_i32();
int3 v3 = int3(tint_symbol_1);
int const tint_symbol_2 = get_i32();
int4 v4 = int4(tint_symbol_2);
}

12
test/tint/expressions/splat/call/i32.wgsl.expected.spvasm
View File

@ -40,14 +40,14 @@
%v2 = OpVariable %_ptr_Function_v2int Function %17
%v3 = OpVariable %_ptr_Function_v3int Function %23
%v4 = OpVariable %_ptr_Function_v4int Function %29
%13 = OpFunctionCall %int %get_i32
%14 = OpCompositeConstruct %v2int %13 %13
%12 = OpFunctionCall %int %get_i32
%14 = OpCompositeConstruct %v2int %12 %12
OpStore %v2 %14
%19 = OpFunctionCall %int %get_i32
%20 = OpCompositeConstruct %v3int %19 %19 %19
%18 = OpFunctionCall %int %get_i32
%20 = OpCompositeConstruct %v3int %18 %18 %18
OpStore %v3 %20
%25 = OpFunctionCall %int %get_i32
%26 = OpCompositeConstruct %v4int %25 %25 %25 %25
%24 = OpFunctionCall %int %get_i32
%26 = OpCompositeConstruct %v4int %24 %24 %24 %24
OpStore %v4 %26
OpReturn
OpFunctionEnd

9
test/tint/expressions/splat/call/u32.wgsl.expected.glsl
View File

@ -9,8 +9,11 @@ uint get_u32() {
}
void f() {
uvec2 v2 = uvec2(get_u32());
uvec3 v3 = uvec3(get_u32());
uvec4 v4 = uvec4(get_u32());
uint tint_symbol = get_u32();
uvec2 v2 = uvec2(tint_symbol);
uint tint_symbol_1 = get_u32();
uvec3 v3 = uvec3(tint_symbol_1);
uint tint_symbol_2 = get_u32();
uvec4 v4 = uvec4(tint_symbol_2);
}

9
test/tint/expressions/splat/call/u32.wgsl.expected.hlsl
View File

@ -8,7 +8,10 @@ uint get_u32() {
}
void f() {
uint2 v2 = uint2((get_u32()).xx);
uint3 v3 = uint3((get_u32()).xxx);
uint4 v4 = uint4((get_u32()).xxxx);
const uint tint_symbol = get_u32();
uint2 v2 = uint2((tint_symbol).xx);
const uint tint_symbol_1 = get_u32();
uint3 v3 = uint3((tint_symbol_1).xxx);
const uint tint_symbol_2 = get_u32();
uint4 v4 = uint4((tint_symbol_2).xxxx);
}

9
test/tint/expressions/splat/call/u32.wgsl.expected.msl
View File

@ -6,8 +6,11 @@ uint get_u32() {
}
void f() {
uint2 v2 = uint2(get_u32());
uint3 v3 = uint3(get_u32());
uint4 v4 = uint4(get_u32());
uint const tint_symbol = get_u32();
uint2 v2 = uint2(tint_symbol);
uint const tint_symbol_1 = get_u32();
uint3 v3 = uint3(tint_symbol_1);
uint const tint_symbol_2 = get_u32();
uint4 v4 = uint4(tint_symbol_2);
}

12
test/tint/expressions/splat/call/u32.wgsl.expected.spvasm
View File

@ -40,14 +40,14 @@
%v2 = OpVariable %_ptr_Function_v2uint Function %17
%v3 = OpVariable %_ptr_Function_v3uint Function %23
%v4 = OpVariable %_ptr_Function_v4uint Function %29
%13 = OpFunctionCall %uint %get_u32
%14 = OpCompositeConstruct %v2uint %13 %13
%12 = OpFunctionCall %uint %get_u32
%14 = OpCompositeConstruct %v2uint %12 %12
OpStore %v2 %14
%19 = OpFunctionCall %uint %get_u32
%20 = OpCompositeConstruct %v3uint %19 %19 %19
%18 = OpFunctionCall %uint %get_u32
%20 = OpCompositeConstruct %v3uint %18 %18 %18
OpStore %v3 %20
%25 = OpFunctionCall %uint %get_u32
%26 = OpCompositeConstruct %v4uint %25 %25 %25 %25
%24 = OpFunctionCall %uint %get_u32
%26 = OpCompositeConstruct %v4uint %24 %24 %24 %24
OpStore %v4 %26
OpReturn
OpFunctionEnd

11
test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.glsl
View File

@ -16,15 +16,16 @@ mat2 arr_to_mat2x2_stride_16(strided_arr arr[2]) {
}
strided_arr[2] mat2x2_stride_16_to_arr(mat2 m) {
strided_arr tint_symbol = strided_arr(m[0u]);
strided_arr tint_symbol_1 = strided_arr(m[1u]);
strided_arr tint_symbol_2[2] = strided_arr[2](tint_symbol, tint_symbol_1);
return tint_symbol_2;
strided_arr tint_symbol_1 = strided_arr(m[0u]);
strided_arr tint_symbol_2 = strided_arr(m[1u]);
strided_arr tint_symbol_3[2] = strided_arr[2](tint_symbol_1, tint_symbol_2);
return tint_symbol_3;
}
void f_1() {
mat2 x_15 = arr_to_mat2x2_stride_16(ssbo.m);
ssbo.m = mat2x2_stride_16_to_arr(x_15);
strided_arr tint_symbol[2] = mat2x2_stride_16_to_arr(x_15);
ssbo.m = tint_symbol;
return;
}

33
test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl
View File

@ -10,44 +10,45 @@ float2x2 arr_to_mat2x2_stride_16(strided_arr arr[2]) {
typedef strided_arr mat2x2_stride_16_to_arr_ret[2];
mat2x2_stride_16_to_arr_ret mat2x2_stride_16_to_arr(float2x2 m) {
const strided_arr tint_symbol_6 = {m[0u]};
const strided_arr tint_symbol_7 = {m[1u]};
const strided_arr tint_symbol_8[2] = {tint_symbol_6, tint_symbol_7};
return tint_symbol_8;
}
strided_arr tint_symbol_1(RWByteAddressBuffer buffer, uint offset) {
const strided_arr tint_symbol_9 = {asfloat(buffer.Load2((offset + 0u)))};
const strided_arr tint_symbol_7 = {m[0u]};
const strided_arr tint_symbol_8 = {m[1u]};
const strided_arr tint_symbol_9[2] = {tint_symbol_7, tint_symbol_8};
return tint_symbol_9;
}
typedef strided_arr tint_symbol_ret[2];
tint_symbol_ret tint_symbol(RWByteAddressBuffer buffer, uint offset) {
strided_arr tint_symbol_2(RWByteAddressBuffer buffer, uint offset) {
const strided_arr tint_symbol_10 = {asfloat(buffer.Load2((offset + 0u)))};
return tint_symbol_10;
}
typedef strided_arr tint_symbol_1_ret[2];
tint_symbol_1_ret tint_symbol_1(RWByteAddressBuffer buffer, uint offset) {
strided_arr arr_1[2] = (strided_arr[2])0;
{
[loop] for(uint i = 0u; (i < 2u); i = (i + 1u)) {
arr_1[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
arr_1[i] = tint_symbol_2(buffer, (offset + (i * 16u)));
}
}
return arr_1;
}
void tint_symbol_4(RWByteAddressBuffer buffer, uint offset, strided_arr value) {
void tint_symbol_5(RWByteAddressBuffer buffer, uint offset, strided_arr value) {
buffer.Store2((offset + 0u), asuint(value.el));
}
void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, strided_arr value[2]) {
void tint_symbol_4(RWByteAddressBuffer buffer, uint offset, strided_arr value[2]) {
strided_arr array[2] = value;
{
[loop] for(uint i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
tint_symbol_4(buffer, (offset + (i_1 * 16u)), array[i_1]);
tint_symbol_5(buffer, (offset + (i_1 * 16u)), array[i_1]);
}
}
}
void f_1() {
const float2x2 x_15 = arr_to_mat2x2_stride_16(tint_symbol(ssbo, 0u));
tint_symbol_3(ssbo, 0u, mat2x2_stride_16_to_arr(x_15));
const float2x2 x_15 = arr_to_mat2x2_stride_16(tint_symbol_1(ssbo, 0u));
const strided_arr tint_symbol[2] = mat2x2_stride_16_to_arr(x_15);
tint_symbol_4(ssbo, 0u, tint_symbol);
return;
}

19
test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
View File

@ -19,20 +19,21 @@ float2x2 arr_to_mat2x2_stride_16(tint_array_wrapper arr) {
}
tint_array_wrapper mat2x2_stride_16_to_arr(float2x2 m) {
strided_arr const tint_symbol = {.el=m[0u]};
strided_arr const tint_symbol_1 = {.el=m[1u]};
tint_array_wrapper const tint_symbol_2 = {.arr={tint_symbol, tint_symbol_1}};
return tint_symbol_2;
strided_arr const tint_symbol_1 = {.el=m[0u]};
strided_arr const tint_symbol_2 = {.el=m[1u]};
tint_array_wrapper const tint_symbol_3 = {.arr={tint_symbol_1, tint_symbol_2}};
return tint_symbol_3;
}
void f_1(device SSBO* const tint_symbol_3) {
float2x2 const x_15 = arr_to_mat2x2_stride_16((*(tint_symbol_3)).m);
(*(tint_symbol_3)).m = mat2x2_stride_16_to_arr(x_15);
void f_1(device SSBO* const tint_symbol_4) {
float2x2 const x_15 = arr_to_mat2x2_stride_16((*(tint_symbol_4)).m);
tint_array_wrapper const tint_symbol = mat2x2_stride_16_to_arr(x_15);
(*(tint_symbol_4)).m = tint_symbol;
return;
}
kernel void f(device SSBO* tint_symbol_4 [[buffer(0)]]) {
f_1(tint_symbol_4);
kernel void f(device SSBO* tint_symbol_5 [[buffer(0)]]) {
f_1(tint_symbol_5);
return;
}

6
test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm
View File

@ -66,9 +66,9 @@
%37 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0
%38 = OpLoad %_arr_strided_arr_uint_2 %37
%35 = OpFunctionCall %mat2v2float %arr_to_mat2x2_stride_16 %38
%39 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0
%40 = OpFunctionCall %_arr_strided_arr_uint_2 %mat2x2_stride_16_to_arr %35
OpStore %39 %40
%39 = OpFunctionCall %_arr_strided_arr_uint_2 %mat2x2_stride_16_to_arr %35
%40 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0
OpStore %40 %39
OpReturn
OpFunctionEnd
%f = OpFunction %void None %31

4
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.glsl
View File

@ -31,7 +31,9 @@ layout(binding = 4) uniform Uniforms_1 {
void tint_symbol() {
InnerS v = InnerS(0);
OuterS s = OuterS(S1[8](S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0))), S1(InnerS[8](InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0), InnerS(0)))));
s.a1[getNextIndex()].a2[uniforms.j] = v;
InnerS tint_symbol_1 = v;
uint tint_symbol_2 = getNextIndex();
s.a1[tint_symbol_2].a2[uniforms.j] = tint_symbol_1;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

3
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.hlsl
View File

@ -25,7 +25,8 @@ void main() {
OuterS s = (OuterS)0;
{
S1 tint_symbol_1[8] = s.a1;
const uint tint_symbol_2_save = getNextIndex();
const uint tint_symbol_4 = getNextIndex();
const uint tint_symbol_2_save = tint_symbol_4;
InnerS tint_symbol_3[8] = tint_symbol_1[tint_symbol_2_save].a2;
tint_symbol_3[uniforms[0].y] = v;
tint_symbol_1[tint_symbol_2_save].a2 = tint_symbol_3;

14
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl
View File

@ -26,16 +26,18 @@ struct OuterS {
tint_array_wrapper_1 a1;
};
uint getNextIndex(thread uint* const tint_symbol_1) {
*(tint_symbol_1) = (*(tint_symbol_1) + 1u);
return *(tint_symbol_1);
uint getNextIndex(thread uint* const tint_symbol_3) {
*(tint_symbol_3) = (*(tint_symbol_3) + 1u);
return *(tint_symbol_3);
}
kernel void tint_symbol(const constant Uniforms* tint_symbol_3 [[buffer(0)]]) {
thread uint tint_symbol_2 = 0u;
kernel void tint_symbol(const constant Uniforms* tint_symbol_5 [[buffer(0)]]) {
thread uint tint_symbol_4 = 0u;
InnerS v = {};
OuterS s = {};
s.a1.arr[getNextIndex(&(tint_symbol_2))].a2.arr[(*(tint_symbol_3)).j] = v;
InnerS const tint_symbol_1 = v;
uint const tint_symbol_2 = getNextIndex(&(tint_symbol_4));
s.a1.arr[tint_symbol_2].a2.arr[(*(tint_symbol_5)).j] = tint_symbol_1;
return;
}

10
test/tint/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm
View File

@ -69,11 +69,11 @@
%18 = OpLabel
%v = OpVariable %_ptr_Function_InnerS Function %23
%s = OpVariable %_ptr_Function_OuterS Function %31
%32 = OpLoad %InnerS %v
%33 = OpFunctionCall %uint %getNextIndex
%35 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
%36 = OpLoad %uint %35
%37 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %33 %uint_0 %36
%38 = OpLoad %InnerS %v
OpStore %37 %38
%36 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
%37 = OpLoad %uint %36
%38 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %33 %uint_0 %37
OpStore %38 %32
OpReturn
OpFunctionEnd

3
test/tint/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.glsl
View File

@ -20,7 +20,8 @@ void tint_symbol() {
OuterS s1 = OuterS(uint[8](0u, 0u, 0u, 0u, 0u, 0u, 0u, 0u));
vec3 v = vec3(0.0f, 0.0f, 0.0f);
v[s1.a1[uniforms.i]] = 1.0f;
v[f(s1.a1[uniforms.i])] = 1.0f;
uint tint_symbol_1 = f(s1.a1[uniforms.i]);
v[tint_symbol_1] = 1.0f;
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

3
test/tint/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.hlsl
View File

@ -19,6 +19,7 @@ void main() {
OuterS s1 = (OuterS)0;
float3 v = float3(0.0f, 0.0f, 0.0f);
set_float3(v, s1.a1[uniforms[0].x], 1.0f);
set_float3(v, f(s1.a1[uniforms[0].x]), 1.0f);
const uint tint_symbol = f(s1.a1[uniforms[0].x]);
set_float3(v, tint_symbol, 1.0f);
return;
}

7
test/tint/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.msl
View File

@ -17,11 +17,12 @@ uint f(uint i) {
return (i + 1u);
}
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
OuterS s1 = {};
float3 v = 0.0f;
v[s1.a1.arr[(*(tint_symbol_1)).i]] = 1.0f;
v[f(s1.a1.arr[(*(tint_symbol_1)).i])] = 1.0f;
v[s1.a1.arr[(*(tint_symbol_2)).i]] = 1.0f;
uint const tint_symbol_1 = f(s1.a1.arr[(*(tint_symbol_2)).i]);
v[tint_symbol_1] = 1.0f;
return;
}

8
test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.glsl
View File

@ -4,11 +4,11 @@ int f(int a, int b, int c) {
return ((a * b) + c);
}
void phony_sink(int p0, int p1, int p2) {
}
void tint_symbol() {
phony_sink(f(1, 2, 3), f(4, 5, 6), f(7, f(8, 9, 10), 11));
int tint_symbol_1 = f(1, 2, 3);
int tint_symbol_2 = f(4, 5, 6);
int tint_symbol_3 = f(8, 9, 10);
int tint_symbol_4 = f(7, tint_symbol_3, 11);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

8
test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.hlsl
View File

@ -2,11 +2,11 @@ int f(int a, int b, int c) {
return ((a * b) + c);
}
void phony_sink(int p0, int p1, int p2) {
}
[numthreads(1, 1, 1)]
void main() {
phony_sink(f(1, 2, 3), f(4, 5, 6), f(7, f(8, 9, 10), 11));
const int tint_symbol = f(1, 2, 3);
const int tint_symbol_1 = f(4, 5, 6);
const int tint_symbol_2 = f(8, 9, 10);
const int tint_symbol_3 = f(7, tint_symbol_2, 11);
return;
}

8
test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.msl
View File

@ -5,11 +5,11 @@ int f(int a, int b, int c) {
return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * as_type<uint>(b)))) + as_type<uint>(c)));
}
void phony_sink(int p0, int p1, int p2) {
}
kernel void tint_symbol() {
phony_sink(f(1, 2, 3), f(4, 5, 6), f(7, f(8, 9, 10), 11));
int const tint_symbol_1 = f(1, 2, 3);
int const tint_symbol_2 = f(4, 5, 6);
int const tint_symbol_3 = f(8, 9, 10);
int const tint_symbol_4 = f(7, tint_symbol_3, 11);
return;
}

8
test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.spvasm
View File

@ -22,10 +22,10 @@
%int_4 = OpConstant %int 4
%int_5 = OpConstant %int 5
%int_6 = OpConstant %int 6
%int_7 = OpConstant %int 7
%int_8 = OpConstant %int 8
%int_9 = OpConstant %int 9
%int_10 = OpConstant %int 10
%int_7 = OpConstant %int 7
%int_11 = OpConstant %int 11
%f = OpFunction %int None %1
%a = OpFunctionParameter %int
@ -40,9 +40,9 @@
%13 = OpLabel
%14 = OpFunctionCall %int %f %int_1 %int_2 %int_3
%18 = OpFunctionCall %int %f %int_4 %int_5 %int_6
%24 = OpFunctionCall %int %f %int_8 %int_9 %int_10
%22 = OpFunctionCall %int %f %int_7 %24 %int_11
%29 = OpIMul %int %18 %22
%22 = OpFunctionCall %int %f %int_8 %int_9 %int_10
%26 = OpFunctionCall %int %f %int_7 %22 %int_11
%29 = OpIMul %int %18 %26
%30 = OpIAdd %int %14 %29
OpReturn
OpFunctionEnd