[tint][msl] Replace volatile loop conditional with outer volatile conditional.
Instead of attempting to preserve UB loops with volatile loads in the conditional, instead wrap the loop in a conditional using a volatile load.
This means the UB loop might still get removed, but control flow that comes after will not be affected.
This has the following benefits:
• Doesn't (appear to) trigger control flow bugs on older hardware.
• Doesn't require a volatile load per iteration (O(n) -> O(1)).
• Looks like it might actually work for all of Skia's tests.
Bug: tint:2125
Change-Id: Ibb17f77d1291eff6a0d5b3244be711bd8ecaff62
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/170801
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: David Neto <dneto@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 58af4e0..718770e 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -2122,7 +2122,7 @@
};
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
- EmitUnconditionalLoopHeader();
+ Line() << IsolateUB() << " while(true) {";
{
ScopedIndent si(this);
if (!EmitStatements(stmt->body->statements)) {
@@ -2192,7 +2192,7 @@
};
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
- EmitUnconditionalLoopHeader();
+ Line() << IsolateUB() << " while(true) {";
IncrementIndent();
TINT_DEFER({
DecrementIndent();
@@ -2215,7 +2215,7 @@
// For-loop can be generated.
{
auto out = Line();
- out << "for";
+ out << IsolateUB() << " for";
{
ScopedParen sp(out);
@@ -2225,8 +2225,7 @@
out << "; ";
}
- EmitLoopCondition(out, cond_buf.str());
- out << "; ";
+ out << cond_buf.str() << "; ";
if (!cont_buf.lines.empty()) {
out << tint::TrimSuffix(cont_buf.lines[0].content, ";");
@@ -2266,7 +2265,7 @@
// as a regular while in MSL. Instead we need to generate a `while(true)` loop.
bool emit_as_loop = cond_pre.lines.size() > 0;
if (emit_as_loop) {
- EmitUnconditionalLoopHeader();
+ Line() << IsolateUB() << " while(true) {";
IncrementIndent();
TINT_DEFER({
DecrementIndent();
@@ -2280,15 +2279,7 @@
}
} else {
// While can be generated.
- {
- auto out = Line();
- out << "while";
- {
- ScopedParen sp(out);
- EmitLoopCondition(out, cond_buf.str());
- }
- out << " {";
- }
+ Line() << IsolateUB() << " while(" << cond_buf.str() << ") {";
if (!EmitStatementsWithIndent(stmt->body->statements)) {
return false;
}
@@ -3027,25 +3018,16 @@
return true;
}
-std::string_view ASTPrinter::LoopPreservingVar() {
- if (loop_preserving_var_.empty()) {
- loop_preserving_var_ = UniqueIdentifier("tint_preserve_loop");
- Line(&helpers_) << "constant static volatile bool " << loop_preserving_var_ << " = true;";
+std::string_view ASTPrinter::IsolateUB() {
+ if (isolate_ub_macro_name_.empty()) {
+ isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
+ auto volatile_true = UniqueIdentifier("tint_volatile_true");
+ Line(&helpers_) << "#define " << isolate_ub_macro_name_ << " \\";
+ Line(&helpers_) << " if (volatile bool " << volatile_true << " = true; " << volatile_true
+ << ")";
Line(&helpers_);
}
- return loop_preserving_var_;
-}
-
-void ASTPrinter::EmitLoopCondition(StringStream& out, const std::string& cond) {
- if (cond.empty()) {
- out << LoopPreservingVar();
- } else {
- out << "(" << cond << ") == " << LoopPreservingVar();
- }
-}
-
-void ASTPrinter::EmitUnconditionalLoopHeader() {
- Line() << "while (" << LoopPreservingVar() << ") {";
+ return isolate_ub_macro_name_;
}
template <typename F>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.h b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
index a13f907c..70269c9 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
@@ -378,18 +378,10 @@
const ast::CallExpression* expr,
const sem::BuiltinFn* builtin);
- /// Lazily emits the global-scope declaration of the helper static volatile boolean variable
- /// used to fool the MSL compiler into thinking loops might exit.
- /// @return the name of the variable.
- std::string_view LoopPreservingVar();
-
- /// Emits a loop condition, but santized so that the MSL compiler can't infer that the loop
- /// never exits.
- void EmitLoopCondition(StringStream& out, const std::string& cond);
-
- /// Emits the header of an unconditional loop, but use the loop-preserving condition to fool the
- /// MSL compiler into thinking that the loop might exit.
- void EmitUnconditionalLoopHeader();
+ /// Lazilly generates the TINT_ISOLATE_UB macro, used to prevent UB statements from affecting
+ /// later logic.
+ /// @return the unique name of the TINT_ISOLATE_UB macro.
+ std::string_view IsolateUB();
/// Handles generating a builtin name
/// @param builtin the semantic info for the builtin
@@ -446,8 +438,9 @@
std::function<bool()> emit_continuing_;
- // Name of the variable used to ensure the MSL compiler thinks a loop will terminate.
- std::string loop_preserving_var_;
+ /// The name of the macro used to prevent UB affecting later control flow.
+ /// Do not use this directly, instead call IsolateUB().
+ std::string isolate_ub_macro_name_;
/// Name of atomicCompareExchangeWeak() helper for the given pointer storage
/// class and struct return type
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index b0a48d8..74d8fdd 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -226,14 +226,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_3 {
tint_array<float2x2, 4> m;
};
void comp_main_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
}
diff --git a/src/tint/lang/msl/writer/ast_printer/continue_test.cc b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
index 9b4dd02..727f05c 100644
--- a/src/tint/lang/msl/writer/ast_printer/continue_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
@@ -44,10 +44,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void test_function() {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (false) {
break;
}
diff --git a/src/tint/lang/msl/writer/ast_printer/loop_test.cc b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
index 73e7cdc..7aed227 100644
--- a/src/tint/lang/msl/writer/ast_printer/loop_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
@@ -51,10 +51,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
fragment void F() {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
break;
}
return;
@@ -79,13 +80,14 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void a_statement() {
}
fragment void F() {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
break;
{
a_statement();
@@ -113,13 +115,14 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void a_statement() {
}
fragment void F() {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
break;
{
a_statement();
@@ -153,8 +156,8 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(while (tint_preserve_loop) {
- while (tint_preserve_loop) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+ TINT_ISOLATE_UB while(true) {
break;
{
a_statement();
@@ -191,7 +194,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(while (tint_preserve_loop) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
float lhs = 2.5f;
float other = 0.0f;
break;
@@ -214,7 +217,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(for(; tint_preserve_loop; ) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(; ; ) {
return;
}
)");
@@ -232,7 +235,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(for(int i = 0; tint_preserve_loop; ) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(int i = 0; ; ) {
return;
}
)");
@@ -263,7 +266,7 @@
f(1);
f(2);
}
- for(; tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; ; ) {
return;
}
}
@@ -282,7 +285,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(for(; (true) == tint_preserve_loop; ) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(; true; ) {
return;
}
)");
@@ -301,9 +304,8 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(
- gen.Result(),
- R"(for(; tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ EXPECT_EQ(gen.Result(),
+ R"(TINT_ISOLATE_UB for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
return;
}
)");
@@ -329,7 +331,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(while (tint_preserve_loop) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
return;
{
f(1);
@@ -355,7 +357,7 @@
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
EXPECT_EQ(
gen.Result(),
- R"(for(int i = 0; (true) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ R"(TINT_ISOLATE_UB for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
a_statement();
}
)");
@@ -387,7 +389,7 @@
f(1);
f(2);
}
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!(true)) { break; }
return;
{
@@ -410,7 +412,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(while((true) == tint_preserve_loop) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
return;
}
)");
@@ -427,7 +429,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(while((true) == tint_preserve_loop) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
continue;
}
)");
@@ -447,7 +449,7 @@
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"(while(((t && false)) == tint_preserve_loop) {
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while((t && false)) {
return;
}
)");
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.msl b/test/tint/array/assign_to_function_var.wgsl.expected.msl
index 794e774..2c9731b 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
@@ -56,7 +57,7 @@
}
void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_8, const constant S* const tint_symbol_9, device S* const tint_symbol_10) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_8))[i] = int4(0);
}
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.msl b/test/tint/array/assign_to_private_var.wgsl.expected.msl
index bf83fb8..5e2e692 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
@@ -56,7 +57,7 @@
}
void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_8, const constant S* const tint_symbol_9, device S* const tint_symbol_10) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_8))[i] = int4(0);
}
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
index accc11a..e0a8b7d 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
@@ -58,7 +59,7 @@
}
void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_10, device S* const tint_symbol_11, const constant S* const tint_symbol_12, device S* const tint_symbol_13, device S_nested* const tint_symbol_14) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_10))[i] = int4(0);
}
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
index 3b0ae7e..bc6e5d4 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
@@ -54,12 +55,12 @@
}
void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_10, threadgroup tint_array<int4, 4>* const tint_symbol_11, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_12, const constant S* const tint_symbol_13, device S* const tint_symbol_14) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_10))[i] = int4(0);
(*(tint_symbol_11))[i] = int4(0);
}
- for(uint idx_1 = local_invocation_index; ((idx_1 < 24u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
+ TINT_ISOLATE_UB for(uint idx_1 = local_invocation_index; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
uint const i_1 = (idx_1 / 6u);
uint const i_2 = ((idx_1 % 6u) / 2u);
uint const i_3 = (idx_1 % 2u);
diff --git a/test/tint/array/strides.spvasm.expected.msl b/test/tint/array/strides.spvasm.expected.msl
index f67348f..3499687 100644
--- a/test/tint/array/strides.spvasm.expected.msl
+++ b/test/tint/array/strides.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct strided_arr {
/* 0x0000 */ float el;
@@ -35,13 +36,13 @@
}
void assign_and_preserve_padding_3(device tint_array<strided_arr, 2>* const dest, tint_array<strided_arr, 2> value) {
- for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 2u); i = (i + 1u)) {
assign_and_preserve_padding_4(&((*(dest))[i]), value[i]);
}
}
void assign_and_preserve_padding_2(device tint_array<tint_array<strided_arr, 2>, 3>* const dest, tint_array<tint_array<strided_arr, 2>, 3> value) {
- for(uint i = 0u; ((i < 3u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 3u); i = (i + 1u)) {
assign_and_preserve_padding_3(&((*(dest))[i]), value[i]);
}
}
@@ -51,7 +52,7 @@
}
void assign_and_preserve_padding(device tint_array<strided_arr_1, 4>* const dest, tint_array<strided_arr_1, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
index fc97070..d4470df 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -104,7 +105,7 @@
}
void assign_and_preserve_padding_3(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
- for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 2u); i = (i + 1u)) {
(*(dest))[i].elements = packed_float3(value[i]);
}
}
diff --git a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
index 204d507..ee7f114 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -159,7 +160,7 @@
}
void assign_and_preserve_padding_6(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
- for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 2u); i = (i + 1u)) {
(*(dest))[i].elements = packed_float3(value[i]);
}
}
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
index 705ddf0..7d6a510 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -106,7 +107,7 @@
}
void assign_and_preserve_padding_3(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
- for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 2u); i = (i + 1u)) {
(*(dest))[i].elements = packed_float3(value[i]);
}
}
diff --git a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
index 291a4c5..f28bf39 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -163,7 +164,7 @@
}
void assign_and_preserve_padding_6(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
- for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 2u); i = (i + 1u)) {
(*(dest))[i].elements = packed_float3(value[i]);
}
}
@@ -175,7 +176,7 @@
}
void assign_and_preserve_padding_8(device tint_array<Inner, 4>* const dest, tint_array<Inner, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_7(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
index be62ea6..d9ef1ab 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float2x2, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol, const constant tint_array<float2x2, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
index 14f6a5b..e09bee8 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
@@ -37,7 +38,7 @@
}
void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* const dest, tint_array<half2x3, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
index 26033d9..80b97c0 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_7 {
tint_array<half2x3, 4> w;
@@ -36,7 +37,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<half2x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* const tint_symbol_1, device half* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = half2x3(half3(0.0h), half3(0.0h));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
index 89464a7..6ca9dae 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -37,7 +38,7 @@
}
void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* const dest, tint_array<float2x3, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
index fce938a..01ed488 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float2x3, 4> w;
@@ -36,7 +37,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<float2x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x3(float3(0.0f), float3(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
index e9ffa1c..4dd72bb 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<half2x4, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<half2x4, 4>* const tint_symbol, const constant tint_array<half2x4, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = half2x4(half4(0.0h), half4(0.0h));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
index 1c75e49..8219d90 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float2x4, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float2x4, 4>* const tint_symbol, const constant tint_array<float2x4, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x4(float4(0.0f), float4(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
index 6699025..084c472 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -38,7 +39,7 @@
}
void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* const dest, tint_array<float3x3, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
index 13d0e9f..aacf68d 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float3x3, 4> w;
@@ -36,7 +37,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<float3x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
index e1bf4de..7d93adf 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float3x4, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float3x4, 4>* const tint_symbol, const constant tint_array<float3x4, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
index 9d0cb69..9c7eb06 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<half4x2, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<half4x2, 4>* const tint_symbol, const constant tint_array<half4x2, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
index 34a353f..346f524 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float4x2, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float4x2, 4>* const tint_symbol, const constant tint_array<float4x2, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
index b9ae76b..b83bc68 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
@@ -39,7 +40,7 @@
}
void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* const dest, tint_array<half4x3, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
index e4983aa..a887469 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<half4x3, 4> w;
@@ -36,7 +37,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<half4x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
index a13be34..7891af0 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -39,7 +40,7 @@
}
void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* const dest, tint_array<float4x3, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
index d6d36ab..efad62c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float4x3, 4> w;
@@ -36,7 +37,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<float4x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
index 3b9b765..1a5b6c4 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<half4x4, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<half4x4, 4>* const tint_symbol, const constant tint_array<half4x4, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
}
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
index 248ef96..d0a15df 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,14 +14,15 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_5 {
tint_array<float4x4, 4> w;
};
void f_inner(uint local_invocation_index, threadgroup tint_array<float4x4, 4>* const tint_symbol, const constant tint_array<float4x4, 4>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
index 57455e1..b9955f2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -31,7 +32,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
index dce4eac..5c82a33 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -29,7 +30,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
index d4f859f..34e488f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
index 5575d59..3f279dc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
index 8866890..56713e0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
@@ -66,7 +67,7 @@
}
void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
index a282ff6..588aaab 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int before;
@@ -59,7 +60,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
index 273e906..ff079b0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -66,7 +67,7 @@
}
void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
index e1b2d71..d33bd36 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int before;
@@ -59,7 +60,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
index 934101a..ad499d1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
index 70dfb67..d159c8c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
index 10d60c6..25e5698 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
index 93a2642..e5b4037 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
index 8f4151b..abdbcb8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -31,7 +32,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
index a83fab0..e42dc9b 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -29,7 +30,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
index a15a0bd..8044d64 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
index 786d13e..e48327d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
index da1c72f..82c29d0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
@@ -67,7 +68,7 @@
}
void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
index 53c0ed3..6933c8e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int before;
@@ -59,7 +60,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
index 18cf65e..b5946b4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -66,7 +67,7 @@
}
void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
index 7c3c470..f97c99e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int before;
@@ -58,7 +59,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
index 29eae07..76fa2d8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
index 0ab8876..be68271 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
index 4e9f71d..cd77cae 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -31,7 +32,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
index ae7c285..82d11af 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -29,7 +30,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
index 500386d..f02b27e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -31,7 +32,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
index 767775c..5e57962 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -29,7 +30,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
index cd0013d..4713bde 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
index ce195f0..0f591ce 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
index 4fc1cf1..752e4c5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
@@ -68,7 +69,7 @@
}
void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
index c53df85..6428fa2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int before;
@@ -59,7 +60,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
index f09b05a..ec467fa 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -68,7 +69,7 @@
}
void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
index 857b0bb..45ddab1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int before;
@@ -59,7 +60,7 @@
}
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
index 121388d..69895c5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
index 8a4308c..ad29c63 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
index a97b159..5153c38 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -32,7 +33,7 @@
}
void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
- for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 4u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
index 92d0210..a97f92f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
/* 0x0000 */ int before;
@@ -30,7 +31,7 @@
};
void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
S const tint_symbol = S{};
(*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/bug/chromium/1403752.wgsl.expected.msl b/test/tint/bug/chromium/1403752.wgsl.expected.msl
index 22558cd..7ae6187 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void d() {
int j = 0;
- for(; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; false; ) {
}
}
diff --git a/test/tint/bug/chromium/1449538.wgsl.expected.msl b/test/tint/bug/chromium/1449538.wgsl.expected.msl
index ed8df03..c9600a2 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.msl
@@ -2,24 +2,25 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
- for(int i0520 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i0520 = 0; false; ) {
}
- for(int i62 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i62 = 0; false; ) {
}
- for(int i0520 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i0520 = 0; false; ) {
}
- for(int i62 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i62 = 0; false; ) {
}
- for(int i62 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i62 = 0; false; ) {
}
- for(int i60 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i60 = 0; false; ) {
}
- for(int i62 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i62 = 0; false; ) {
}
- for(int i60 = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i60 = 0; false; ) {
}
}
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
index c934bb2..b14a8c8 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct UBO {
/* 0x0000 */ int dynamic_idx;
@@ -29,7 +30,7 @@
};
void f_inner(uint local_invocation_index, threadgroup S* const tint_symbol, device Result* const tint_symbol_1, const constant UBO* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 64u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol)).data[i] = 0;
}
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
index 1c3c26d..3ac2c9d 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct UBO {
/* 0x0000 */ int dynamic_idx;
@@ -29,7 +30,7 @@
};
void f_inner(uint local_invocation_index, threadgroup S* const tint_symbol, const constant UBO* const tint_symbol_1, device Result* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 64u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol)).data[i] = 0;
}
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
index e64a323..8a54b5e 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float2 vUV [[user(locn0)]];
@@ -15,7 +16,7 @@
float4 tint_symbol_inner(float2 vUV, texture2d<float, access::sample> tint_symbol_4, sampler tint_symbol_5) {
float3 const random = tint_symbol_4.sample(tint_symbol_5, vUV).rgb;
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((i < 1)) {
} else {
break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
index 89f7f51..da64fd9 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
float2 v2f;
@@ -12,7 +13,7 @@
};
void foo(thread tint_private_vars_struct* const tint_private_vars) {
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
(*(tint_private_vars)).v2f[i] = 1.0f;
(*(tint_private_vars)).v3i[i] = 1;
(*(tint_private_vars)).v4u[i] = 1u;
@@ -22,7 +23,7 @@
kernel void tint_symbol() {
thread tint_private_vars_struct tint_private_vars = {};
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
foo(&(tint_private_vars));
}
return;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
index 046f8c4..af0bc02 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
float2 v2f;
@@ -21,7 +22,7 @@
kernel void tint_symbol() {
thread tint_private_vars_struct tint_private_vars = {};
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
foo(&(tint_private_vars));
}
return;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
index c62bf40..0c01f6d 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
float2 v2f = 0.0f;
@@ -17,7 +18,7 @@
bool2 v2b = false;
bool3 v3b = false;
bool4 v4b = false;
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
v2f[i] = 1.0f;
v3f[i] = 1.0f;
v4f[i] = 1.0f;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
index 847e8f6..e455d47 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
float2 v2f = 0.0f;
@@ -13,7 +14,7 @@
uint4 v4u_2 = 0u;
bool2 v2b = false;
bool2 v2b_2 = false;
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
v2f[i] = 1.0f;
v3i[i] = 1;
v4u[i] = 1u;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
index 91ac2be..1904207 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
float2 v2f = 0.0f;
@@ -17,7 +18,7 @@
bool2 v2b = false;
bool3 v3b = false;
bool4 v4b = false;
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
v2f[i] = 1.0f;
v2i[i] = 1;
v2u[i] = 1u;
diff --git a/test/tint/bug/tint/1064.wgsl.expected.msl b/test/tint/bug/tint/1064.wgsl.expected.msl
index 7b81e57..3b9997e 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
fragment void tint_symbol() {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/bug/tint/1081.wgsl.expected.msl b/test/tint/bug/tint/1081.wgsl.expected.msl
index 8b120a8..d1ab933 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
bool tint_discarded;
@@ -25,7 +26,7 @@
int tint_symbol_inner(int3 x, thread tint_private_vars_struct* const tint_private_vars) {
int y = x[0];
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
int const r = f(y, tint_private_vars);
if ((r == 0)) {
break;
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index 2abb7d5..9f26018 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct LightData_tint_packed_vec3 {
/* 0x0000 */ float4 position;
@@ -86,8 +87,8 @@
int const TILE_SIZE = 16;
int const TILE_COUNT_X = 2;
int const TILE_COUNT_Y = 2;
- for(int y = 0; ((y < TILE_COUNT_Y)) == tint_preserve_loop; y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
- for(int x = 0; ((x < TILE_COUNT_X)) == tint_preserve_loop; x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int y = 0; (y < TILE_COUNT_Y); y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int x = 0; (x < TILE_COUNT_X); x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_SIZE))));
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / (*(tint_symbol_3)).fullScreenSize.xy) - float2(1.0f));
float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / (*(tint_symbol_3)).fullScreenSize.xy) - float2(1.0f));
@@ -98,7 +99,7 @@
frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord[1]) / viewNear), 0.0f);
frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord[1] / viewNear), 0.0f);
float dp = 0.0f;
- for(uint i = 0u; ((i < 6u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 6u); i = (i + 1u)) {
float4 p = 0.0f;
if ((frustumPlanes[i][0] > 0.0f)) {
p[0] = boxMax[0];
diff --git a/test/tint/bug/tint/1321.wgsl.expected.msl b/test/tint/bug/tint/1321.wgsl.expected.msl
index 99f3002..2d4a466 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int foo() {
return 1;
@@ -25,7 +26,7 @@
{
int const tint_symbol_1 = foo();
int const a_save = tint_symbol_1;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
{
float const x = arr[a_save];
break;
diff --git a/test/tint/bug/tint/1474-a.wgsl.expected.msl b/test/tint/bug/tint/1474-a.wgsl.expected.msl
index 02c5b90..dda1010 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
- while((true) == tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (true) {
break;
} else {
diff --git a/test/tint/bug/tint/1538.wgsl.expected.msl b/test/tint/bug/tint/1538.wgsl.expected.msl
index ea91f6c..bff79a4 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
template<typename T, size_t N>
struct tint_array {
@@ -21,7 +22,7 @@
}
int f() {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
g();
break;
}
@@ -30,7 +31,7 @@
}
kernel void tint_symbol(device tint_array<uint, 1>* tint_symbol_1 [[buffer(0)]]) {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (((*(tint_symbol_1))[0] == 0u)) {
break;
}
diff --git a/test/tint/bug/tint/1557.wgsl.expected.msl b/test/tint/bug/tint/1557.wgsl.expected.msl
index 9952e88..d47a9f9 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
return 0;
@@ -10,7 +11,7 @@
void g() {
int j = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((j >= 1)) {
break;
}
diff --git a/test/tint/bug/tint/1604.wgsl.expected.msl b/test/tint/bug/tint/1604.wgsl.expected.msl
index 2dfebe4..10273a4 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.msl
@@ -2,12 +2,13 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) {
switch(*(tint_symbol_1)) {
case 0: {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
return;
}
break;
diff --git a/test/tint/bug/tint/1605.wgsl.expected.msl b/test/tint/bug/tint/1605.wgsl.expected.msl
index 78953c3..bface6e 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
bool func_3(const constant int* const tint_symbol_1) {
- for(int i = 0; ((i < *(tint_symbol_1))) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
- for(int j = -1; ((j == 1)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < *(tint_symbol_1)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int j = -1; (j == 1); j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
return false;
}
}
diff --git a/test/tint/bug/tint/1764.wgsl.expected.msl b/test/tint/bug/tint/1764.wgsl.expected.msl
index f4dcf7f..ded1dcc 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.msl
@@ -14,10 +14,11 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<int, 246>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 246u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 246u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_1))[i] = 0;
}
diff --git a/test/tint/bug/tint/2010.spvasm.expected.msl b/test/tint/bug/tint/2010.spvasm.expected.msl
index 061aa4e..c528ebb 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint3 x_3;
@@ -48,7 +49,7 @@
uint x_88 = 0u;
uint const x_52 = (*(tint_private_vars)).x_3[0];
x_54 = 0u;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
uint x_55 = 0u;
x_58 = (*(tint_symbol_3)).field0.field0;
if ((x_54 < x_58)) {
@@ -80,7 +81,7 @@
}
x_85 = x_76.xyxy;
x_88 = 1u;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
float4 x_111 = 0.0f;
float4 x_86 = 0.0f;
uint x_89 = 0u;
@@ -130,7 +131,7 @@
atomic_store_explicit(tint_symbol_13, 0u, memory_order_relaxed);
atomic_store_explicit(tint_symbol_14, 0u, memory_order_relaxed);
}
- for(uint idx = local_invocation_index; ((idx < 4096u)) == tint_preserve_loop; idx = (idx + 32u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 32u)) {
uint const i = idx;
S const tint_symbol_1 = S{};
(*(tint_symbol_15))[i] = tint_symbol_1;
diff --git a/test/tint/bug/tint/2059.wgsl.expected.msl b/test/tint/bug/tint/2059.wgsl.expected.msl
index 097729c..3488383 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
@@ -64,7 +65,7 @@
}
void assign_and_preserve_padding_3(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1>* const dest, tint_array<float3x3, 1> value) {
- for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 1u); i = (i + 1u)) {
assign_and_preserve_padding(&((*(dest))[i]), value[i]);
}
}
@@ -78,7 +79,7 @@
}
void assign_and_preserve_padding_6(device tint_array<S_tint_packed_vec3, 1>* const dest, tint_array<S, 1> value) {
- for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 1u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
@@ -88,14 +89,14 @@
}
void assign_and_preserve_padding_7(device tint_array<S2_tint_packed_vec3, 1>* const dest, tint_array<S2, 1> value) {
- for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 1u); i = (i + 1u)) {
assign_and_preserve_padding_2(&((*(dest))[i]), value[i]);
}
}
kernel void tint_symbol(device tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_8 [[buffer(0)]], device S_tint_packed_vec3* tint_symbol_9 [[buffer(1)]], device S2_tint_packed_vec3* tint_symbol_10 [[buffer(2)]], device S3_tint_packed_vec3* tint_symbol_11 [[buffer(3)]], device S4_tint_packed_vec3* tint_symbol_12 [[buffer(4)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1>* tint_symbol_13 [[buffer(5)]], device tint_array<S_tint_packed_vec3, 1>* tint_symbol_14 [[buffer(6)]], device tint_array<S2_tint_packed_vec3, 1>* tint_symbol_15 [[buffer(7)]]) {
float3x3 m = float3x3(0.0f);
- for(uint c = 0u; ((c < 3u)) == tint_preserve_loop; c = (c + 1u)) {
+ TINT_ISOLATE_UB for(uint c = 0u; (c < 3u); c = (c + 1u)) {
m[c] = float3(float(((c * 3u) + 1u)), float(((c * 3u) + 2u)), float(((c * 3u) + 3u)));
}
{
diff --git a/test/tint/bug/tint/221.wgsl.expected.msl b/test/tint/bug/tint/221.wgsl.expected.msl
index c14e039..79c1851 100644
--- a/test/tint/bug/tint/221.wgsl.expected.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Buf {
/* 0x0000 */ uint count;
@@ -27,7 +28,7 @@
kernel void tint_symbol(device Buf* tint_symbol_1 [[buffer(0)]]) {
uint i = 0u;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((i >= (*(tint_symbol_1)).count)) {
break;
}
diff --git a/test/tint/bug/tint/534.wgsl.expected.msl b/test/tint/bug/tint/534.wgsl.expected.msl
index fd3dbbc..1594381 100644
--- a/test/tint/bug/tint/534.wgsl.expected.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
uint4 tint_ftou(float4 v) {
return select(uint4(4294967295u), select(uint4(v), uint4(0u), (v < float4(0.0f))), (v < float4(4294967040.0f)));
@@ -47,7 +48,7 @@
bool success = true;
uint4 srcColorBits = 0u;
uint4 dstColorBits = tint_ftou(dstColor);
- for(uint i = 0u; ((i < (*(tint_symbol_3)).channelCount)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
uint const tint_symbol_1 = i;
srcColorBits[tint_symbol_1] = ConvertToFp16FloatValue(srcColor[i]);
success = (success && (srcColorBits[i] == dstColorBits[i]));
diff --git a/test/tint/bug/tint/744.wgsl.expected.msl b/test/tint/bug/tint/744.wgsl.expected.msl
index 85ae3c4..9300f34 100644
--- a/test/tint/bug/tint/744.wgsl.expected.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Uniforms {
/* 0x0000 */ uint2 aShape;
@@ -31,7 +32,7 @@
uint const dimInner = (*(tint_symbol_1)).aShape[1];
uint const dimOutter = (*(tint_symbol_1)).outShape[1];
uint result = 0u;
- for(uint i = 0u; ((i < dimInner)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
uint const a = (i + (resultCell[0] * dimInner));
uint const b = (resultCell[1] + (i * dimOutter));
result = (result + ((*(tint_symbol_2)).numbers[a] * (*(tint_symbol_3)).numbers[b]));
diff --git a/test/tint/bug/tint/757.wgsl.expected.msl b/test/tint/bug/tint/757.wgsl.expected.msl
index 11910c8..b70aae4 100644
--- a/test/tint/bug/tint/757.wgsl.expected.msl
+++ b/test/tint/bug/tint/757.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Constants {
int level;
@@ -28,7 +29,7 @@
uint flatIndex = (((4u * GlobalInvocationID[2]) + (2u * GlobalInvocationID[1])) + GlobalInvocationID[0]);
flatIndex = (flatIndex * 1u);
float4 texel = tint_symbol_1.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
- for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 1u); i = (i + 1u)) {
(*(tint_symbol_2)).values[(flatIndex + i)] = texel[0];
}
}
diff --git a/test/tint/bug/tint/914.wgsl.expected.msl b/test/tint/bug/tint/914.wgsl.expected.msl
index 98d18dd..2c3568d 100644
--- a/test/tint/bug/tint/914.wgsl.expected.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Uniforms {
/* 0x0000 */ uint dimAOuter;
@@ -54,7 +55,7 @@
}
void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_11, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_12, const constant Uniforms* const tint_symbol_13, const device Matrix* const tint_symbol_14, const device Matrix* const tint_symbol_15, device Matrix* const tint_symbol_16) {
- for(uint idx = local_invocation_index; ((idx < 4096u)) == tint_preserve_loop; idx = (idx + 256u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
uint const i = (idx / 64u);
uint const i_1 = (idx % 64u);
(*(tint_symbol_11))[i][i_1] = 0.0f;
@@ -69,16 +70,16 @@
tint_array<float, 16> acc = {};
float ACached = 0.0f;
tint_array<float, 4> BCached = {};
- for(uint index = 0u; ((index < 16u)) == tint_preserve_loop; index = (index + 1u)) {
+ TINT_ISOLATE_UB for(uint index = 0u; (index < 16u); index = (index + 1u)) {
acc[index] = 0.0f;
}
uint const ColPerThreadA = 4u;
uint const tileColA = (local_id[0] * ColPerThreadA);
uint const RowPerThreadB = 4u;
uint const tileRowB = (local_id[1] * RowPerThreadB);
- for(uint t = 0u; ((t < numTiles)) == tint_preserve_loop; t = (t + 1u)) {
- for(uint innerRow = 0u; ((innerRow < 4u)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
- for(uint innerCol = 0u; ((innerCol < ColPerThreadA)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
+ TINT_ISOLATE_UB for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
+ TINT_ISOLATE_UB for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
+ TINT_ISOLATE_UB for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
uint const inputRow = (tileRow + innerRow);
uint const inputCol = (tileColA + innerCol);
uint const tint_symbol_1 = inputRow;
@@ -86,8 +87,8 @@
(*(tint_symbol_11))[tint_symbol_1][tint_symbol_2] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_13, tint_symbol_14);
}
}
- for(uint innerRow = 0u; ((innerRow < RowPerThreadB)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
- for(uint innerCol = 0u; ((innerCol < 4u)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
+ TINT_ISOLATE_UB for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+ TINT_ISOLATE_UB for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint const inputRow = (tileRowB + innerRow);
uint const inputCol = (tileCol + innerCol);
uint const tint_symbol_3 = innerCol;
@@ -96,13 +97,13 @@
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- for(uint k = 0u; ((k < 64u)) == tint_preserve_loop; k = (k + 1u)) {
- for(uint inner = 0u; ((inner < 4u)) == tint_preserve_loop; inner = (inner + 1u)) {
+ TINT_ISOLATE_UB for(uint k = 0u; (k < 64u); k = (k + 1u)) {
+ TINT_ISOLATE_UB for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
BCached[inner] = (*(tint_symbol_12))[k][(tileCol + inner)];
}
- for(uint innerRow = 0u; ((innerRow < 4u)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
+ TINT_ISOLATE_UB for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
ACached = (*(tint_symbol_11))[(tileRow + innerRow)][k];
- for(uint innerCol = 0u; ((innerCol < 4u)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
+ TINT_ISOLATE_UB for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * 4u) + innerCol);
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
}
@@ -110,8 +111,8 @@
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
- for(uint innerRow = 0u; ((innerRow < 4u)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
- for(uint innerCol = 0u; ((innerCol < 4u)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
+ TINT_ISOLATE_UB for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
+ TINT_ISOLATE_UB for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * 4u) + innerCol);
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_13, tint_symbol_16);
}
diff --git a/test/tint/bug/tint/942.wgsl.expected.msl b/test/tint/bug/tint/942.wgsl.expected.msl
index ea207c3..7a7d361 100644
--- a/test/tint/bug/tint/942.wgsl.expected.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Params {
/* 0x0000 */ uint filterDim;
@@ -30,7 +31,7 @@
}
void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocation_index, threadgroup tint_array<tint_array<float3, 256>, 4>* const tint_symbol_1, const constant Params* const tint_symbol_2, texture2d<float, access::sample> tint_symbol_3, const constant Flip* const tint_symbol_4, sampler tint_symbol_5, texture2d<float, access::write> tint_symbol_6) {
- for(uint idx = local_invocation_index; ((idx < 1024u)) == tint_preserve_loop; idx = (idx + 64u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 1024u); idx = (idx + 64u)) {
uint const i_1 = (idx / 256u);
uint const i_2 = (idx % 256u);
(*(tint_symbol_1))[i_1][i_2] = float3(0.0f);
@@ -39,8 +40,8 @@
uint const filterOffset = tint_div(((*(tint_symbol_2)).filterDim - 1u), 2u);
uint2 const dims = uint2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
uint2 const baseIndex = (((WorkGroupID.xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))) - uint2(filterOffset, 0u));
- for(uint r = 0u; ((r < 4u)) == tint_preserve_loop; r = (r + 1u)) {
- for(uint c = 0u; ((c < 4u)) == tint_preserve_loop; c = (c + 1u)) {
+ TINT_ISOLATE_UB for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+ TINT_ISOLATE_UB for(uint c = 0u; (c < 4u); c = (c + 1u)) {
uint2 loadIndex = (baseIndex + uint2(c, r));
if (((*(tint_symbol_4)).value != 0u)) {
loadIndex = loadIndex.yx;
@@ -49,8 +50,8 @@
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- for(uint r = 0u; ((r < 4u)) == tint_preserve_loop; r = (r + 1u)) {
- for(uint c = 0u; ((c < 4u)) == tint_preserve_loop; c = (c + 1u)) {
+ TINT_ISOLATE_UB for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+ TINT_ISOLATE_UB for(uint c = 0u; (c < 4u); c = (c + 1u)) {
uint2 writeIndex = (baseIndex + uint2(c, r));
if (((*(tint_symbol_4)).value != 0u)) {
writeIndex = writeIndex.yx;
@@ -58,7 +59,7 @@
uint const center = ((4u * LocalInvocationID[0]) + c);
if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
float3 acc = float3(0.0f);
- for(uint f = 0u; ((f < (*(tint_symbol_2)).filterDim)) == tint_preserve_loop; f = (f + 1u)) {
+ TINT_ISOLATE_UB for(uint f = 0u; (f < (*(tint_symbol_2)).filterDim); f = (f + 1u)) {
uint i = ((center + f) - filterOffset);
acc = (acc + ((1.0f / float((*(tint_symbol_2)).filterDim)) * (*(tint_symbol_1))[r][i]));
}
diff --git a/test/tint/bug/tint/948.wgsl.expected.msl b/test/tint/bug/tint/948.wgsl.expected.msl
index 510b303..0047f1c 100644
--- a/test/tint/bug/tint/948.wgsl.expected.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
float2 tUV;
@@ -100,7 +101,7 @@
float2 const x_111 = (*(tint_symbol_8)).stageSize;
stageUnits = (float2(1.0f) / x_111);
i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
int const x_122 = i;
if ((x_122 < 2)) {
} else {
@@ -136,7 +137,7 @@
float const x_184 = animationData[2];
(*(tint_private_vars)).mt = fmod((x_181 * x_184), 1.0f);
f = 0.0f;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
float const x_193 = f;
if ((x_193 < 8.0f)) {
} else {
diff --git a/test/tint/bug/tint/949.wgsl.expected.msl b/test/tint/bug/tint/949.wgsl.expected.msl
index b936ff4..d8b8c4f 100644
--- a/test/tint/bug/tint/949.wgsl.expected.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
float u_Float;
@@ -332,7 +333,7 @@
lastSampledHeight = 1.0f;
currSampledHeight = 1.0f;
i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
int const x_388 = i;
if ((x_388 < 15)) {
} else {
diff --git a/test/tint/bug/tint/990.wgsl.expected.msl b/test/tint/bug/tint/990.wgsl.expected.msl
index 7d1836d1..ec5a232 100644
--- a/test/tint/bug/tint/990.wgsl.expected.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
int i = 0;
- for(; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; false; ) {
}
}
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
index d612d9e5..9c324b7 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint local_invocation_index_1;
@@ -31,7 +32,7 @@
void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index_2;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((idx < 6u))) {
break;
}
@@ -55,7 +56,7 @@
}
void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
- for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 6u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
+ TINT_ISOLATE_UB for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
uint const i = (idx_1 / 2u);
uint const i_1 = (idx_1 % 2u);
uint const i_2 = (idx_1 % 1u);
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
index 7cf8ff3..3ae828f 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
@@ -14,10 +14,11 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
- for(uint idx = local_invocation_index; ((idx < 6u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
uint const i = (idx / 2u);
uint const i_1 = (idx % 2u);
uint const i_2 = (idx % 1u);
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
index 7177f3f..312a5cf 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint local_invocation_index_1;
@@ -23,7 +24,7 @@
void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index_2;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((idx < 4u))) {
break;
}
@@ -45,7 +46,7 @@
}
void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_2) {
- for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 4u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
+ TINT_ISOLATE_UB for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
uint const i = idx_1;
atomic_store_explicit(&((*(tint_symbol_2))[i]), 0u, memory_order_relaxed);
}
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
index a5fb3f5..d5d5e27 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
@@ -14,10 +14,11 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
atomic_store_explicit(&((*(tint_symbol))[i]), 0u, memory_order_relaxed);
}
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
index d612d9e5..9c324b7 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint local_invocation_index_1;
@@ -31,7 +32,7 @@
void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index_2;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((idx < 6u))) {
break;
}
@@ -55,7 +56,7 @@
}
void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
- for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 6u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
+ TINT_ISOLATE_UB for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
uint const i = (idx_1 / 2u);
uint const i_1 = (idx_1 % 2u);
uint const i_2 = (idx_1 % 1u);
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
index 7cf8ff3..3ae828f 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
@@ -14,10 +14,11 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
- for(uint idx = local_invocation_index; ((idx < 6u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
uint const i = (idx / 2u);
uint const i_1 = (idx % 2u);
uint const i_2 = (idx % 1u);
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
index 29ecbd8..431c5c2 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint local_invocation_index_1;
@@ -35,7 +36,7 @@
void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index_2;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((idx < 10u))) {
break;
}
@@ -59,7 +60,7 @@
}
void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<S_atomic, 10>* const tint_symbol_2) {
- for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 10u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
+ TINT_ISOLATE_UB for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
uint const i = idx_1;
(*(tint_symbol_2))[i].x = 0;
atomic_store_explicit(&((*(tint_symbol_2))[i].a), 0u, memory_order_relaxed);
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
index f46810f..04d1917 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int x;
@@ -23,7 +24,7 @@
};
void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S, 10>* const tint_symbol) {
- for(uint idx = local_invocation_index; ((idx < 10u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i].x = 0;
atomic_store_explicit(&((*(tint_symbol))[i].a), 0u, memory_order_relaxed);
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
index e543c39..1e6af34 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint local_invocation_index_1;
@@ -37,7 +38,7 @@
(*(tint_symbol)).x = 0;
(*(tint_symbol)).y = 0u;
idx = local_invocation_index_2;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((idx < 10u))) {
break;
}
@@ -63,7 +64,7 @@
(*(tint_symbol_2)).x = 0;
(*(tint_symbol_2)).y = 0u;
}
- for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 10u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
+ TINT_ISOLATE_UB for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
uint const i = idx_1;
atomic_store_explicit(&((*(tint_symbol_2)).a[i]), 0u, memory_order_relaxed);
}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
index d4e7b34..ad1e165 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int x;
@@ -27,7 +28,7 @@
(*(tint_symbol)).x = 0;
(*(tint_symbol)).y = 0u;
}
- for(uint idx = local_invocation_index; ((idx < 10u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
uint const i = idx;
atomic_store_explicit(&((*(tint_symbol)).a[i]), 0u, memory_order_relaxed);
}
diff --git a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
index 1998f18..2242f0f 100644
--- a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
+++ b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
@@ -2,12 +2,13 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void foo(texture2d<int, access::read_write> tint_symbol) {
{
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((i < 3))) {
break;
}
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
index 3fa0ee1..b840f6f 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int tint_workgroupUniformLoad(threadgroup int* const p) {
threadgroup_barrier(mem_flags::mem_threadgroup);
@@ -14,7 +15,7 @@
void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
{
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
int const tint_symbol = i;
int const tint_symbol_1 = tint_workgroupUniformLoad(tint_symbol_4);
if (!((tint_symbol < tint_symbol_1))) {
diff --git a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
index 4c582c9..071ba03 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
@@ -14,7 +14,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
@@ -22,7 +23,7 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
- for(; (((x > v[0]) && (dfdx(1.0f) > 0.0f))) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; ((x > v[0]) && (dfdx(1.0f) > 0.0f)); ) {
}
}
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
index 63ce7b4..f410fb6 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
@@ -14,7 +14,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
@@ -22,7 +23,7 @@
void tint_symbol_inner(float x, texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
float4 v = float4(0.0f);
- for(; ((x > v[0])) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; (x > v[0]); ) {
v = tint_symbol_3.sample(tint_symbol_4, float2(0.0f));
}
}
diff --git a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
index 7fb929f..3dae810 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
@@ -14,14 +14,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
{
if ((x > 0.0f)) { break; }
}
diff --git a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
index 7dc46dd..602c64b 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
@@ -14,14 +14,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
{
if ((x > 0.0f)) { break; }
}
diff --git a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
index 38bded8..38499a3 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
@@ -14,14 +14,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
{
if ((x > 0.0f)) { break; }
}
diff --git a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
index 6a4adfd..2c394d5 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
@@ -14,7 +14,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
@@ -22,7 +23,7 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
- while((((x > 0.0f) && (dfdx(1.0f) > 0.0f))) == tint_preserve_loop) {
+ TINT_ISOLATE_UB while(((x > 0.0f) && (dfdx(1.0f) > 0.0f))) {
}
}
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
index 9f81385..649759e 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
@@ -14,7 +14,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_2 {
float x [[user(locn0)]];
@@ -22,7 +23,7 @@
void tint_symbol_inner(float x, texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
float4 v = float4(0.0f);
- while(((x > v[0])) == tint_preserve_loop) {
+ TINT_ISOLATE_UB while((x > v[0])) {
v = tint_symbol_3.sample(tint_symbol_4, float2(0.0f));
}
}
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
index 666903f..d82e019 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct strided_arr {
/* 0x0000 */ float2 el;
@@ -41,7 +42,7 @@
}
void assign_and_preserve_padding(device tint_array<strided_arr, 2>* const dest, tint_array<strided_arr, 2> value) {
- for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 2u); i = (i + 1u)) {
assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
}
}
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.msl b/test/tint/loops/continue_in_switch.wgsl.expected.msl
index db2c859..e2991a3 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void f() {
- for(int i = 0; ((i < 4)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
switch(i) {
case 0: {
continue;
diff --git a/test/tint/loops/loop.wgsl.expected.msl b/test/tint/loops/loop.wgsl.expected.msl
index 096ffe4..b26d512 100644
--- a/test/tint/loops/loop.wgsl.expected.msl
+++ b/test/tint/loops/loop.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return i;
diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.msl
index 865b88d..1344663 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((i > 4)) {
return i;
}
diff --git a/test/tint/loops/loop_with_continuing.wgsl.expected.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.msl
index cdb7a82..024d5b9 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((i > 4)) {
return i;
}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.msl b/test/tint/loops/multiple_continues.wgsl.expected.msl
index 53d5235..803ee58 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
switch(i) {
case 0: {
continue;
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.msl b/test/tint/loops/multiple_switch.wgsl.expected.msl
index 2bb537d..5becfcc 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
int i = 0;
- for(int i_1 = 0; ((i_1 < 2)) == tint_preserve_loop; i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
switch(i_1) {
case 0: {
continue;
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
index dc0686f..01dd258 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
- for(int j = 0; ((j < 2)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
switch(i) {
case 0: {
continue;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
index 0dab2a6..e4d3cb7 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
@@ -2,13 +2,14 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
switch(i) {
case 0: {
- for(int j = 0; ((j < 2)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
switch(j) {
case 0: {
continue;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
index 23a482e..5aa6cb9 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
@@ -2,14 +2,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
int k = 0;
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
switch(i) {
case 0: {
- for(int j = 0; ((j < 2)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
switch(j) {
case 0: {
continue;
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
index 5885440..1adeeb0 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
int j = 0;
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
switch(i) {
case 0: {
switch(j) {
diff --git a/test/tint/loops/nested_loops.wgsl.expected.msl b/test/tint/loops/nested_loops.wgsl.expected.msl
index 2fd8895..3c0a637 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.msl
@@ -2,17 +2,18 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
int j = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return 1;
}
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
if ((j > 4)) {
return 2;
diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
index e0b172b..000b361 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
@@ -2,16 +2,17 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
int j = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((i > 4)) {
return 1;
}
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((j > 4)) {
return 2;
}
diff --git a/test/tint/loops/single_continue.wgsl.expected.msl b/test/tint/loops/single_continue.wgsl.expected.msl
index 2e85db7..f321a4b 100644
--- a/test/tint/loops/single_continue.wgsl.expected.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
kernel void tint_symbol() {
- for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
switch(i) {
case 0: {
continue;
diff --git a/test/tint/loops/while.wgsl.expected.msl b/test/tint/loops/while.wgsl.expected.msl
index e2655d3..f7428f9 100644
--- a/test/tint/loops/while.wgsl.expected.msl
+++ b/test/tint/loops/while.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
- while(((i < 4)) == tint_preserve_loop) {
+ TINT_ISOLATE_UB while((i < 4)) {
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
return i;
diff --git a/test/tint/loops/while_with_continue.wgsl.expected.msl b/test/tint/loops/while_with_continue.wgsl.expected.msl
index 68d388a..65fe6f5 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
int f() {
int i = 0;
- while(((i < 4)) == tint_preserve_loop) {
+ TINT_ISOLATE_UB while((i < 4)) {
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
continue;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
index 7541aea..457eb5c 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct str {
tint_array<int, 4> arr;
@@ -25,7 +26,7 @@
}
void tint_symbol_inner(uint local_invocation_index, threadgroup str* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_1)).arr[i] = 0;
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
index 28efaee..91fb413 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct str {
int i;
@@ -25,7 +26,7 @@
}
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<str, 4>* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i_1 = idx;
str const tint_symbol_1 = str{};
(*(tint_symbol_2))[i_1] = tint_symbol_1;
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
index fd889f8..140b544 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct str {
tint_array<int, 4> arr;
@@ -26,7 +27,7 @@
}
void tint_symbol_inner(uint local_invocation_index, threadgroup str* const tint_symbol_2) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_2)).arr[i] = 0;
}
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
index 1ff952b..054a738 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct str {
int i;
@@ -26,7 +27,7 @@
}
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<str, 4>* const tint_symbol_3) {
- for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i_1 = idx;
str const tint_symbol_1 = str{};
(*(tint_symbol_3))[i_1] = tint_symbol_1;
diff --git a/test/tint/samples/compute_boids.wgsl.expected.msl b/test/tint/samples/compute_boids.wgsl.expected.msl
index b7fe17a..cec623c 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_symbol_1 {
float2 a_particlePos [[attribute(0)]];
@@ -87,7 +88,7 @@
int cVelCount = 0;
float2 pos = 0.0f;
float2 vel = 0.0f;
- for(uint i = 0u; ((i < 5u)) == tint_preserve_loop; i = (i + 1u)) {
+ TINT_ISOLATE_UB for(uint i = 0u; (i < 5u); i = (i + 1u)) {
if ((i == index)) {
continue;
}
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
index 64296fb..6972c31 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Uniforms {
/* 0x0000 */ uint i;
@@ -31,7 +32,7 @@
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
InnerS v = {};
OuterS s1 = {};
- for(int i = 0; ((i < 4)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
s1.a1[(*(tint_symbol_1)).i] = v;
}
return;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
index 74b50fc..1bbb83c 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Uniforms {
/* 0x0000 */ uint i;
@@ -31,7 +32,7 @@
kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
InnerS v = {};
OuterS s1 = {};
- for(int i = 0; ((i < 4)) == tint_preserve_loop; s1.a1[(*(tint_symbol_1)).i] = v) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 4); s1.a1[(*(tint_symbol_1)).i] = v) {
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
return;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
index 76e6c8d..db29e84 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct Uniforms {
/* 0x0000 */ uint i;
@@ -32,7 +33,7 @@
InnerS v = {};
OuterS s1 = {};
int i = 0;
- for(s1.a1[(*(tint_symbol_1)).i] = v; ((i < 4)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(s1.a1[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
}
return;
}
diff --git a/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl b/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
index 99a18b6..4a7047b 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint i;
@@ -47,7 +48,7 @@
int const tint_symbol_2 = idx1(tint_private_vars);
int const tint_symbol_save = tint_symbol_2;
a[tint_symbol_save] = (a[tint_symbol_save] * 2.0f);
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
int const tint_symbol_3 = idx2(tint_private_vars);
if (!((a[tint_symbol_3] < 10.0f))) {
break;
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.msl b/test/tint/statements/decrement/complex.wgsl.expected.msl
index 7efbaf0..49634f7 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint v;
@@ -62,7 +63,7 @@
int const tint_symbol_2_save_1 = tint_symbol_7;
int const tint_symbol_3 = idx3(tint_private_vars);
(*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3] = as_type<int>((as_type<uint>((*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3]) - as_type<uint>(1)));
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!(((*(tint_private_vars)).v < 10u))) {
break;
}
diff --git a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
index 3c3e4da..195cf52 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol(device uint* const tint_symbol_1) {
- for(; ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; *(tint_symbol_1) = (*(tint_symbol_1) - 1u)) {
+ TINT_ISOLATE_UB for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) - 1u)) {
}
}
diff --git a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
index 89002b5..bee2a0c 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol(device uint* const tint_symbol_1) {
- for(*(tint_symbol_1) = (*(tint_symbol_1) - 1u); ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(*(tint_symbol_1) = (*(tint_symbol_1) - 1u); (*(tint_symbol_1) < 10u); ) {
}
}
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
index fd1c50f..9316b19 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
bool tint_discarded;
@@ -28,7 +29,7 @@
int result = tint_ftoi(tint_symbol_4.sample(tint_symbol_5, coord)[0]);
{
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!((i < 10))) {
break;
}
diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
index 75f3dd5..6eac933 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
bool tint_discarded;
@@ -19,7 +20,7 @@
}
if ((*(tint_symbol_2) < 0.0f)) {
int i = 0;
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if ((*(tint_symbol_2) > float(i))) {
if (!(tint_private_vars.tint_discarded)) {
*(tint_symbol_2) = float(i);
diff --git a/test/tint/statements/for/basic.wgsl.expected.msl b/test/tint/statements/for/basic.wgsl.expected.msl
index 453a4d3..ecdaaf5 100644
--- a/test/tint/statements/for/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.msl
@@ -2,13 +2,14 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void some_loop_body() {
}
void f() {
- for(int i = 0; ((i < 5)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; (i < 5); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
some_loop_body();
}
}
diff --git a/test/tint/statements/for/complex.wgsl.expected.msl b/test/tint/statements/for/complex.wgsl.expected.msl
index cda792d..38b2dcb 100644
--- a/test/tint/statements/for/complex.wgsl.expected.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.msl
@@ -2,14 +2,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void some_loop_body() {
}
void f() {
int j = 0;
- for(int i = 0; (((i < 5) && (j < 10))) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; ((i < 5) && (j < 10)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
some_loop_body();
j = as_type<int>((as_type<uint>(i) * as_type<uint>(30)));
}
diff --git a/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl b/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
index c2e0634..225ca95 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
int i = 0;
- for(; ((i < 1)) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; (i < 1); ) {
}
}
diff --git a/test/tint/statements/for/condition/basic.wgsl.expected.msl b/test/tint/statements/for/condition/basic.wgsl.expected.msl
index 7da63a4..9e4d48b 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
int i = 0;
- for(; ((i < 4)) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; (i < 4); ) {
}
}
diff --git a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
index ca21e5a..1ce1740 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
@@ -2,7 +2,8 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int i;
@@ -10,7 +11,7 @@
void f() {
int i = 0;
- for(; ((i < 1)) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; (i < 1); ) {
}
}
diff --git a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
index 2da7062..91eff5e 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
int i = 0;
- for(; (false) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
}
}
diff --git a/test/tint/statements/for/continuing/basic.wgsl.expected.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.msl
index 2da7062..91eff5e 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.msl
@@ -2,11 +2,12 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
int i = 0;
- for(; (false) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
}
}
diff --git a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
index 95ed7eb..6b5fe78 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
@@ -2,14 +2,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int i;
};
void f() {
- for(int i = 0; (false) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ TINT_ISOLATE_UB for(int i = 0; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
}
}
diff --git a/test/tint/statements/for/empty.wgsl.expected.msl b/test/tint/statements/for/empty.wgsl.expected.msl
index d05f6d5..92bb29d 100644
--- a/test/tint/statements/for/empty.wgsl.expected.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
- for(; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(; false; ) {
}
}
diff --git a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
index 6e66415..e9d4bc4 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
- for(int i = 1; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i = 1; false; ) {
}
}
diff --git a/test/tint/statements/for/initializer/basic.wgsl.expected.msl b/test/tint/statements/for/initializer/basic.wgsl.expected.msl
index c50d83b..f1ec896 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
- for(int i = 0; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i = 0; false; ) {
}
}
diff --git a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
index 60dbd8a..1af8e13 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
@@ -2,14 +2,15 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct S {
int i;
};
void f() {
- for(int i = 1; (false) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int i = 1; false; ) {
}
}
diff --git a/test/tint/statements/for/scoping.wgsl.expected.msl b/test/tint/statements/for/scoping.wgsl.expected.msl
index 0206ba1..74c7f3b 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void f() {
- for(int must_not_collide = 0; tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(int must_not_collide = 0; ; ) {
break;
}
int must_not_collide = 0;
diff --git a/test/tint/statements/increment/complex.wgsl.expected.msl b/test/tint/statements/increment/complex.wgsl.expected.msl
index 509a125..9f7b0fa 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.msl
@@ -14,7 +14,8 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
struct tint_private_vars_struct {
uint v;
@@ -62,7 +63,7 @@
int const tint_symbol_2_save_1 = tint_symbol_7;
int const tint_symbol_3 = idx3(tint_private_vars);
(*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3] = as_type<int>((as_type<uint>((*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3]) + as_type<uint>(1)));
- while (tint_preserve_loop) {
+ TINT_ISOLATE_UB while(true) {
if (!(((*(tint_private_vars)).v < 10u))) {
break;
}
diff --git a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
index 8092545..e452d03 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol(device uint* const tint_symbol_1) {
- for(; ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; *(tint_symbol_1) = (*(tint_symbol_1) + 1u)) {
+ TINT_ISOLATE_UB for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) + 1u)) {
}
}
diff --git a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
index 2cc9f36..f46177a 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
@@ -2,10 +2,11 @@
using namespace metal;
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol(device uint* const tint_symbol_1) {
- for(*(tint_symbol_1) = (*(tint_symbol_1) + 1u); ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; ) {
+ TINT_ISOLATE_UB for(*(tint_symbol_1) = (*(tint_symbol_1) + 1u); (*(tint_symbol_1) < 10u); ) {
}
}
diff --git a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
index a578747..37f33b9 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
@@ -14,10 +14,11 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<tint_array<int, 3>, 2>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 6u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
uint const i = (idx / 3u);
uint const i_1 = (idx % 3u);
(*(tint_symbol_1))[i][i_1] = 0;
diff --git a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
index 513d0a3..a1ca997 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
@@ -14,10 +14,11 @@
T elements[N];
};
-constant static volatile bool tint_preserve_loop = true;
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<int, 3>* const tint_symbol_1) {
- for(uint idx = local_invocation_index; ((idx < 3u)) == tint_preserve_loop; idx = (idx + 1u)) {
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol_1))[i] = 0;
}
