[msl] Emit TINT_ISOLATE_UB macro
Add some basic loop unit tests as well.
Bug: 41493494
Bug: 42251016
Change-Id: I018b7a90736100b2bbbe23bd6f9824afded3d096
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/204679
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/writer/BUILD.bazel b/src/tint/lang/msl/writer/BUILD.bazel
index a9e6eb5..adf42fa 100644
--- a/src/tint/lang/msl/writer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/BUILD.bazel
@@ -95,6 +95,7 @@
"helper_test.h",
"if_test.cc",
"let_test.cc",
+ "loop_test.cc",
"return_test.cc",
"type_test.cc",
"var_test.cc",
diff --git a/src/tint/lang/msl/writer/BUILD.cmake b/src/tint/lang/msl/writer/BUILD.cmake
index 214c2d8..40cde20 100644
--- a/src/tint/lang/msl/writer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/BUILD.cmake
@@ -109,6 +109,7 @@
lang/msl/writer/helper_test.h
lang/msl/writer/if_test.cc
lang/msl/writer/let_test.cc
+ lang/msl/writer/loop_test.cc
lang/msl/writer/return_test.cc
lang/msl/writer/type_test.cc
lang/msl/writer/var_test.cc
diff --git a/src/tint/lang/msl/writer/BUILD.gn b/src/tint/lang/msl/writer/BUILD.gn
index bf7afcd..d5b2bfa 100644
--- a/src/tint/lang/msl/writer/BUILD.gn
+++ b/src/tint/lang/msl/writer/BUILD.gn
@@ -99,6 +99,7 @@
"helper_test.h",
"if_test.cc",
"let_test.cc",
+ "loop_test.cc",
"return_test.cc",
"type_test.cc",
"var_test.cc",
diff --git a/src/tint/lang/msl/writer/loop_test.cc b/src/tint/lang/msl/writer/loop_test.cc
new file mode 100644
index 0000000..330dd2c
--- /dev/null
+++ b/src/tint/lang/msl/writer/loop_test.cc
@@ -0,0 +1,165 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/msl/writer/helper_test.h"
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+namespace tint::msl::writer {
+namespace {
+
+TEST_F(MslWriterTest, Loop) {
+ auto* func = b.Function("a", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* l = b.Loop();
+ b.Append(l->Body(), [&] { b.ExitLoop(l); });
+ b.Append(l->Continuing(), [&] { b.NextIteration(l); });
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+void a() {
+ {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ break;
+ }
+ }
+}
+)");
+}
+
+TEST_F(MslWriterTest, LoopContinueAndBreakIf) {
+ auto* func = b.Function("a", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* l = b.Loop();
+ b.Append(l->Body(), [&] { b.Continue(l); });
+ b.Append(l->Continuing(), [&] { b.BreakIf(l, true); });
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+void a() {
+ {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ {
+ if (true) { break; }
+ }
+ continue;
+ }
+ }
+}
+)");
+}
+
+TEST_F(MslWriterTest, LoopBodyVarInContinue) {
+ auto* func = b.Function("a", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* l = b.Loop();
+ b.Append(l->Body(), [&] {
+ auto* v = b.Var("v", true);
+ b.Continue(l);
+
+ b.Append(l->Continuing(), [&] { b.BreakIf(l, v); });
+ });
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+void a() {
+ {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ bool v = true;
+ {
+ if (v) { break; }
+ }
+ continue;
+ }
+ }
+}
+)");
+}
+
+TEST_F(MslWriterTest, LoopInitializer) {
+ auto* func = b.Function("a", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* l = b.Loop();
+ b.Append(l->Initializer(), [&] {
+ auto* v = b.Var("v", true);
+ b.NextIteration(l);
+
+ b.Append(l->Body(), [&] { b.Continue(l); });
+ b.Append(l->Continuing(), [&] { b.BreakIf(l, v); });
+ });
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+void a() {
+ {
+ bool v = true;
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ {
+ if (v) { break; }
+ }
+ continue;
+ }
+ }
+}
+)");
+}
+
+} // namespace
+} // namespace tint::msl::writer
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 5ab53b9..0578753 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -182,6 +182,10 @@
/// Non-empty only if the template has been generated.
std::string array_template_name_;
+ /// 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_;
+
/// Block to emit for a continuing
std::function<void()> emit_continuing_;
@@ -215,6 +219,21 @@
return array_template_name_;
}
+ /// @returns a call to the TINT_ISOLATE_UB macro, creating that macro on first call
+ std::string IsolateUB() {
+ if (isolate_ub_macro_name_.empty()) {
+ TINT_SCOPED_ASSIGNMENT(current_buffer_, &preamble_buffer_);
+ isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
+ Line();
+ Line() << "#define " << isolate_ub_macro_name_ << "(VOLATILE_NAME) \\";
+ Line() << " volatile bool VOLATILE_NAME = true; \\";
+ Line() << " if (VOLATILE_NAME)";
+ }
+ StringStream ss;
+ ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_true") << ")";
+ return ss.str();
+ }
+
/// Find all structures that are used in host-shareable address spaces and mark them as such so
/// that we know to pad the properly when we emit them.
void FindHostShareableStructs() {
@@ -652,7 +671,7 @@
ScopedIndent init(current_buffer_);
EmitBlock(l->Initializer());
- Line() << "while(true) {";
+ Line() << IsolateUB() << " while(true) {";
{
ScopedIndent si(current_buffer_);
EmitBlock(l->Body());
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
index 978fdfb..14a5111 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
@@ -24,6 +24,10 @@
device S* src_storage;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
};
@@ -59,7 +63,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
index 13c9933..09ba79a 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
@@ -26,6 +26,10 @@
thread tint_array<tint_array<tint_array<int, 2>, 3>, 4>* dst_nested;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
};
@@ -59,7 +63,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
index a5d2af7..eeaddb1 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
@@ -30,6 +30,10 @@
device S_nested* dst_nested;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
};
@@ -63,7 +67,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
index cade83f..9b28283 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
@@ -26,6 +26,10 @@
threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* dst_nested;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_4 {
tint_array<int4, 4> tint_symbol_1;
tint_array<int4, 4> tint_symbol_2;
@@ -61,7 +65,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
@@ -77,7 +81,7 @@
{
uint v_2 = 0u;
v_2 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_3 = v_2;
if ((v_3 >= 24u)) {
break;
diff --git a/test/tint/array/strides.spvasm.expected.ir.msl b/test/tint/array/strides.spvasm.expected.ir.msl
index d84e88c..431a06a 100644
--- a/test/tint/array/strides.spvasm.expected.ir.msl
+++ b/test/tint/array/strides.spvasm.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct strided_arr_1 {
/* 0x0000 */ tint_array<tint_array<strided_arr, 2>, 3> el;
/* 0x0030 */ tint_array<int8_t, 80> tint_pad_1;
@@ -39,7 +43,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 2u)) {
break;
@@ -57,7 +61,7 @@
{
uint v_2 = 0u;
v_2 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_3 = v_2;
if ((v_3 >= 3u)) {
break;
@@ -79,7 +83,7 @@
{
uint v_4 = 0u;
v_4 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
uint const v_5 = v_4;
if ((v_5 >= 4u)) {
break;
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
index 3d7a05d..351609c 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct Inner_packed_vec3 {
/* 0x0000 */ float scalar_f32;
/* 0x0004 */ int scalar_i32;
@@ -61,7 +65,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 2u)) {
break;
diff --git a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
index e5473dd..378d780 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 packed_1;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad_1;
@@ -84,7 +88,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 2u)) {
break;
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
index 4b1417c..df814f3 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct Inner {
/* 0x0000 */ int scalar_i32;
/* 0x0004 */ float scalar_f32;
@@ -65,7 +69,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 2u)) {
break;
diff --git a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
index 30736b2..ea90738 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
@@ -20,6 +20,10 @@
/* 0x000a */ tint_array<int8_t, 2> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 packed;
/* 0x000c */ tint_array<int8_t, 4> tint_pad_1;
@@ -96,7 +100,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
@@ -114,7 +118,7 @@
{
uint v_2 = 0u;
v_2 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_3 = v_2;
if ((v_3 >= 2u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index ba5453b..ba3dfa8 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<float2x2, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float2x2, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
index f94c189..dec692e 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u;
device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* s;
@@ -32,7 +36,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index 0ea7608..749dea2 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -24,6 +24,10 @@
threadgroup tint_array<half2x3, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<half2x3, 4> tint_symbol;
};
@@ -47,7 +51,7 @@
{
uint v_11 = 0u;
v_11 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_12 = v_11;
if ((v_12 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
index 9576b4b..0a23ec0 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u;
device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* s;
@@ -32,7 +36,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index 6def1c5..2f7aae2 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -23,6 +23,10 @@
threadgroup tint_array<float2x3, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float2x3, 4> tint_symbol;
};
@@ -46,7 +50,7 @@
{
uint v_11 = 0u;
v_11 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_12 = v_11;
if ((v_12 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index 0c46632..3a70542 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<half2x4, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<half2x4, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index fc5bc06..e3cde16 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<float2x4, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float2x4, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
index 592c479..33d13f1 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u;
device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* s;
@@ -33,7 +37,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index 51a5e15..fb01a13 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -23,6 +23,10 @@
threadgroup tint_array<float3x3, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float3x3, 4> tint_symbol;
};
@@ -50,7 +54,7 @@
{
uint v_15 = 0u;
v_15 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_16 = v_15;
if ((v_16 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index ab712f0..aee2d2e 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<float3x4, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float3x4, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index 32af314..b4321d6 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<half4x2, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<half4x2, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index eaf51f2..ab5e665 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<float4x2, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float4x2, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
index c850e03..b00d8c5 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u;
device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* s;
@@ -34,7 +38,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index f0ab510..1dcf6b3 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -23,6 +23,10 @@
threadgroup tint_array<half4x3, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<half4x3, 4> tint_symbol;
};
@@ -54,7 +58,7 @@
{
uint v_19 = 0u;
v_19 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_20 = v_19;
if ((v_20 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
index 001068f..a6ad8aa 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u;
device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* s;
@@ -34,7 +38,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index 4f98077..6daa1c2 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -23,6 +23,10 @@
threadgroup tint_array<float4x3, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float4x3, 4> tint_symbol;
};
@@ -54,7 +58,7 @@
{
uint v_19 = 0u;
v_19 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_20 = v_19;
if ((v_20 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 6df1c65..4f0e022 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<half4x4, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<half4x4, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 3e6728a..35ccec5 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<float4x4, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<float4x4, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
index ef86daa..89d8544 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -36,7 +40,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
index cdb0a34..201ecd7 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -26,6 +26,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -34,7 +38,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
index e4e4e10..f51c619 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index aae5830..81a2c8d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
index 269e006..2fa25c3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,6 +33,10 @@
int after;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
device tint_array<S_packed_vec3, 4>* s;
@@ -61,7 +65,7 @@
{
uint v_4 = 0u;
v_4 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_5 = v_4;
if ((v_5 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index 061fc36..359ceec 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -38,6 +38,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -61,7 +65,7 @@
{
uint v_7 = 0u;
v_7 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_8 = v_7;
if ((v_8 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
index 62fde8f..f080673 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -33,6 +33,10 @@
int after;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
device tint_array<S_packed_vec3, 4>* s;
@@ -61,7 +65,7 @@
{
uint v_4 = 0u;
v_4 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_5 = v_4;
if ((v_5 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index c704370..c04ef3c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -38,6 +38,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -61,7 +65,7 @@
{
uint v_7 = 0u;
v_7 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_8 = v_7;
if ((v_8 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
index 4ae220e..913cd10 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index 2af3623..ad9574a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
index 64214a6..5d712dc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index 21c7386..7d3170c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
index c1ce431..e70ceb1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -36,7 +40,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
index 3e4e0ec..35f763e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -26,6 +26,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -34,7 +38,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
index 07922c7..e425614 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
index ec49517..b845d9e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
index 2a77a01..07623e3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,6 +33,10 @@
int after;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
device tint_array<S_packed_vec3, 4>* s;
@@ -63,7 +67,7 @@
{
uint v_5 = 0u;
v_5 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_6 = v_5;
if ((v_6 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
index 2de80e0..02ba688 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -38,6 +38,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -62,7 +66,7 @@
{
uint v_8 = 0u;
v_8 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_9 = v_8;
if ((v_9 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
index 886f65e..1d18969d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -32,6 +32,10 @@
int after;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
device tint_array<S_packed_vec3, 4>* s;
@@ -62,7 +66,7 @@
{
uint v_5 = 0u;
v_5 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_6 = v_5;
if ((v_6 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index d682233..864fd2a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -37,6 +37,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -61,7 +65,7 @@
{
uint v_8 = 0u;
v_8 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_9 = v_8;
if ((v_9 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
index b9e1126..6b56138 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
index 8713553..0b5d8a2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
index 6565537..8ff7e82 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -36,7 +40,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index 0b30ecf..6c637af 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -26,6 +26,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -34,7 +38,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
index 72c17d4..139cb9e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -36,7 +40,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index 83dee0c..6f8d3af 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -26,6 +26,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -34,7 +38,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
index 5953e94..0a827db 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index 4a0d100..9a43ba5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
index 9dec492..c71d8eb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,6 +33,10 @@
int after;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
device tint_array<S_packed_vec3, 4>* s;
@@ -65,7 +69,7 @@
{
uint v_6 = 0u;
v_6 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_7 = v_6;
if ((v_7 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index 6dffa29..2ef73c3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -38,6 +38,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -63,7 +67,7 @@
{
uint v_9 = 0u;
v_9 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_10 = v_9;
if ((v_10 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
index d57a1fc..d9bf522 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -33,6 +33,10 @@
int after;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
device tint_array<S_packed_vec3, 4>* s;
@@ -65,7 +69,7 @@
{
uint v_6 = 0u;
v_6 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_7 = v_6;
if ((v_7 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index 8694393..9318dc9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -38,6 +38,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -63,7 +67,7 @@
{
uint v_9 = 0u;
v_9 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_10 = v_9;
if ((v_10 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
index a867e72..198b3a0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 47f3faf..ea18e24 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
index 10fb498..473a54d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
/* 0x0084 */ tint_array<int8_t, 60> tint_pad_2;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -37,7 +41,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 714b11d..08984e5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,6 +27,10 @@
threadgroup tint_array<S, 4>* w;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
index 0362f07..fe4107f 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void d() {
int j = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
index d0ebe5e..3e5d3f6 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
int i0520 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
@@ -16,7 +20,7 @@
}
{
int i62 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if (false) {
} else {
break;
@@ -28,7 +32,7 @@
}
{
int i0520 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
if (false) {
} else {
break;
@@ -40,7 +44,7 @@
}
{
int i62 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
if (false) {
} else {
break;
@@ -52,7 +56,7 @@
}
{
int i62 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
if (false) {
} else {
break;
@@ -64,7 +68,7 @@
}
{
int i60 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
if (false) {
} else {
break;
@@ -76,7 +80,7 @@
}
{
int i62 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_6) while(true) {
if (false) {
} else {
break;
@@ -88,7 +92,7 @@
}
{
int i60 = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_7) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
index b0f5e65..5fc61ca 100644
--- a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
@@ -13,6 +13,10 @@
T elements[N];
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
device tint_array<int, 4>* tint_symbol;
};
@@ -20,7 +24,7 @@
void foo(device tint_array<int, 4>* const arg) {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
index 955f46a..3a5398f 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
@@ -31,6 +31,10 @@
threadgroup S* s;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
S tint_symbol;
};
@@ -39,7 +43,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 64u)) {
break;
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
index 6808f32..d725486 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
@@ -31,6 +31,10 @@
threadgroup S* s;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
S tint_symbol;
};
@@ -39,7 +43,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 64u)) {
break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
index ea965fb..31c8097 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
@@ -8,10 +8,14 @@
thread bool2* v2b;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void foo(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
@@ -36,7 +40,7 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v2f=(&v2f), .v3i=(&v3i), .v4u=(&v4u), .v2b=(&v2b)};
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
index 0cc9f48..d2cc280 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
@@ -8,6 +8,10 @@
thread bool2* v2b;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void foo(tint_module_vars_struct tint_module_vars) {
int i = 0;
(*tint_module_vars.v2f)[i] = 1.0f;
@@ -24,7 +28,7 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v2f=(&v2f), .v3i=(&v3i), .v4u=(&v4u), .v2b=(&v2b)};
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
index a53d1ad..3efa51b 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
@@ -1,6 +1,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
float2 v2f = 0.0f;
float3 v3f = 0.0f;
@@ -16,7 +20,7 @@
bool4 v4b = false;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
index 5c6fc4c3..3c97ae1 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
@@ -1,6 +1,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
float2 v2f = 0.0f;
float2 v2f_2 = 0.0f;
@@ -12,7 +16,7 @@
bool2 v2b_2 = false;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
index 259100d..d779d9c 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
@@ -1,6 +1,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
float2 v2f = 0.0f;
float3 v3f = 0.0f;
@@ -16,7 +20,7 @@
bool4 v4b = false;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/bug/tint/1064.wgsl.expected.ir.msl b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
index cb79529..e728d67 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
@@ -1,9 +1,13 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
fragment void tint_symbol() {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/bug/tint/1081.wgsl.expected.ir.msl b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
index f44ab5b..5b8f06d 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
@@ -5,6 +5,10 @@
thread bool* continue_execution;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_outputs {
int tint_symbol_1 [[color(2)]];
};
@@ -23,7 +27,7 @@
int tint_symbol_inner(int3 x, tint_module_vars_struct tint_module_vars) {
int y = x[0u];
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
int const r = f(y, tint_module_vars);
if ((r == 0)) {
break;
diff --git a/test/tint/bug/tint/1121.wgsl.expected.ir.msl b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
index 00753a9..53964d3 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
@@ -56,6 +56,10 @@
const constant Uniforms* uniforms;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
uint index = GlobalInvocationID[0u];
if ((index >= (*tint_module_vars.config).numLights)) {
@@ -87,14 +91,14 @@
int const TILE_COUNT_Y = 2;
{
int y = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((y < TILE_COUNT_Y)) {
} else {
break;
}
{
int x = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((x < TILE_COUNT_X)) {
} else {
break;
@@ -114,7 +118,7 @@
float dp = 0.0f;
{
uint i = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
if ((i < 6u)) {
} else {
break;
diff --git a/test/tint/bug/tint/1321.wgsl.expected.ir.msl b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
index c842da8..537c53a 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
@@ -13,6 +13,10 @@
T elements[N];
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int foo() {
return 1;
}
@@ -21,7 +25,7 @@
tint_array<float, 4> arr = tint_array<float, 4>{};
{
thread float* const a = (&arr[foo()]);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
float const x = (*a);
break;
}
diff --git a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
index 84d42b6..9444e09 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
@@ -1,9 +1,13 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (true) {
} else {
break;
diff --git a/test/tint/bug/tint/1538.wgsl.expected.ir.msl b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
index a178dd0..2c5da95 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
@@ -1,6 +1,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -23,7 +27,7 @@
int f() {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
g();
break;
}
@@ -35,7 +39,7 @@
kernel void tint_symbol(device tint_array<uint, 1>* buf [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.buf=buf};
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if (((*tint_module_vars.buf)[0] == 0u)) {
break;
}
diff --git a/test/tint/bug/tint/1557.wgsl.expected.ir.msl b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
index 2c8d850..81900a2 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
@@ -1,6 +1,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_module_vars_struct {
const constant int* u;
};
@@ -12,7 +16,7 @@
void g() {
int j = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((j >= 1)) {
break;
}
diff --git a/test/tint/bug/tint/1604.wgsl.expected.ir.msl b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
index da22029..80c92e3 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
@@ -5,13 +5,17 @@
const constant int* x;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol(const constant int* x [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x=x};
switch((*tint_module_vars.x)) {
case 0:
{
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
return;
}
}
diff --git a/test/tint/bug/tint/1605.wgsl.expected.ir.msl b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
index 25ac1bd..f618569 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
@@ -5,17 +5,21 @@
const constant int* b;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
bool func_3(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < (*tint_module_vars.b))) {
} else {
break;
}
{
int j = -1;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((j == 1)) {
} else {
break;
diff --git a/test/tint/bug/tint/1764.wgsl.expected.ir.msl b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
index 8e5721e..49a49c8 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
@@ -17,6 +17,10 @@
threadgroup tint_array<int, 246>* W;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<int, 246> tint_symbol_1;
};
@@ -25,7 +29,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 246u)) {
break;
diff --git a/test/tint/bug/tint/2010.spvasm.expected.ir.msl b/test/tint/bug/tint/2010.spvasm.expected.ir.msl
index 57eed71..3e1fb27 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.ir.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.ir.msl
@@ -46,6 +46,10 @@
device S_4* x_12;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_6 {
tint_array<S, 4096> tint_symbol_1;
atomic_uint tint_symbol_2;
@@ -62,7 +66,7 @@
uint const x_52 = (*tint_module_vars.x_3)[0u];
x_54 = 0u;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint x_55 = 0u;
x_58 = (*tint_module_vars.x_6).field0.field0;
if ((x_54 < x_58)) {
@@ -96,7 +100,7 @@
x_85 = x_76.xyxy;
x_88 = 1u;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
float4 x_111 = 0.0f;
float4 x_86 = 0.0f;
uint x_89 = 0u;
@@ -153,7 +157,7 @@
{
uint v_3 = 0u;
v_3 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
uint const v_4 = v_3;
if ((v_4 >= 4096u)) {
break;
diff --git a/test/tint/bug/tint/2039.wgsl.expected.ir.msl b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
index fd6830b..8401406a 100644
--- a/test/tint/bug/tint/2039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
uint out = 0u;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
bool tint_continue = false;
switch(2) {
case 1:
diff --git a/test/tint/bug/tint/2059.wgsl.expected.ir.msl b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
index 12a3ecb..0d038a4 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct S2_packed_vec3 {
/* 0x0000 */ tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1> m;
};
@@ -71,7 +75,7 @@
{
uint v = 0u;
v = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 1u)) {
break;
@@ -93,7 +97,7 @@
{
uint v_2 = 0u;
v_2 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_3 = v_2;
if ((v_3 >= 1u)) {
break;
@@ -115,7 +119,7 @@
{
uint v_4 = 0u;
v_4 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
uint const v_5 = v_4;
if ((v_5 >= 1u)) {
break;
@@ -142,7 +146,7 @@
float3x3 m = float3x3(0.0f);
{
uint c = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
if ((c < 3u)) {
} else {
break;
diff --git a/test/tint/bug/tint/2201.wgsl.expected.ir.msl b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
index 9e9e983..aae30ea 100644
--- a/test/tint/bug/tint/2201.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
@@ -5,9 +5,13 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (true) {
break;
} else {
diff --git a/test/tint/bug/tint/2202.wgsl.expected.ir.msl b/test/tint/bug/tint/2202.wgsl.expected.ir.msl
index 8217706..f4bd798 100644
--- a/test/tint/bug/tint/2202.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2202.wgsl.expected.ir.msl
@@ -5,11 +5,15 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
return;
}
}
diff --git a/test/tint/bug/tint/221.wgsl.expected.ir.msl b/test/tint/bug/tint/221.wgsl.expected.ir.msl
index f3e91c9..5fb7153 100644
--- a/test/tint/bug/tint/221.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
device Buf* b;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
uint tint_mod_u32(uint lhs, uint rhs) {
uint const v = select(rhs, 1u, (rhs == 0u));
return (lhs - ((lhs / v) * v));
@@ -31,7 +35,7 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.b=b};
uint i = 0u;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i >= (*tint_module_vars.b).count)) {
break;
}
diff --git a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
index e8798fc..be80066 100644
--- a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
@@ -5,11 +5,15 @@
texture2d<float, access::sample> tint_symbol;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void e(texture2d<float, access::sample> tint_symbol [[texture(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_symbol=tint_symbol};
{
uint level = tint_module_vars.tint_symbol.get_num_mip_levels();
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((level > 0u)) {
} else {
break;
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
index 806cf8e..cab996a 100644
--- a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
@@ -5,14 +5,18 @@
device int* tint_symbol;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol_1(device int* tint_symbol [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_symbol=tint_symbol};
int i = (*tint_module_vars.tint_symbol);
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
{
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((i > 5)) {
i = (i * 2);
break;
diff --git a/test/tint/bug/tint/534.wgsl.expected.ir.msl b/test/tint/bug/tint/534.wgsl.expected.ir.msl
index c008335..bf985e2 100644
--- a/test/tint/bug/tint/534.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.ir.msl
@@ -31,6 +31,10 @@
const constant Uniforms* uniforms;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
uint ConvertToFp16FloatValue(float fp32) {
return 1u;
}
@@ -54,7 +58,7 @@
uint4 dstColorBits = tint_v4f32_to_v4u32(dstColor);
{
uint i = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < (*tint_module_vars.uniforms).channelCount)) {
} else {
break;
diff --git a/test/tint/bug/tint/744.wgsl.expected.ir.msl b/test/tint/bug/tint/744.wgsl.expected.ir.msl
index b1d6f8d..f86490f 100644
--- a/test/tint/bug/tint/744.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.ir.msl
@@ -30,6 +30,10 @@
const constant Uniforms* uniforms;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol_inner(uint3 global_id, tint_module_vars_struct tint_module_vars) {
uint2 const resultCell = uint2(global_id[1u], global_id[0u]);
uint const dimInner = (*tint_module_vars.uniforms).aShape[1u];
@@ -37,7 +41,7 @@
uint result = 0u;
{
uint i = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < dimInner)) {
} else {
break;
diff --git a/test/tint/bug/tint/757.wgsl.expected.ir.msl b/test/tint/bug/tint/757.wgsl.expected.ir.msl
index 12d3138..c4b51b5 100644
--- a/test/tint/bug/tint/757.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/757.wgsl.expected.ir.msl
@@ -27,13 +27,17 @@
device Result* result;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
uint flatIndex = (((4u * GlobalInvocationID[2u]) + (2u * GlobalInvocationID[1u])) + GlobalInvocationID[0u]);
flatIndex = (flatIndex * 1u);
float4 texel = tint_module_vars.myTexture.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
{
uint i = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 1u)) {
} else {
break;
diff --git a/test/tint/bug/tint/914.wgsl.expected.ir.msl b/test/tint/bug/tint/914.wgsl.expected.ir.msl
index 0745782..ebedfcc 100644
--- a/test/tint/bug/tint/914.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.ir.msl
@@ -32,6 +32,10 @@
threadgroup tint_array<tint_array<float, 64>, 64>* mm_Bsub;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_3 {
tint_array<tint_array<float, 64>, 64> tint_symbol_1;
tint_array<tint_array<float, 64>, 64> tint_symbol_2;
@@ -86,7 +90,7 @@
{
uint v_3 = 0u;
v_3 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_4 = v_3;
if ((v_4 >= 4096u)) {
break;
@@ -110,7 +114,7 @@
tint_array<float, 4> BCached = {};
{
uint index = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((index < 16u)) {
} else {
break;
@@ -128,21 +132,21 @@
uint const tileRowB = (local_id[1u] * RowPerThreadB);
{
uint t = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
if ((t < numTiles)) {
} else {
break;
}
{
uint innerRow = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
if ((innerRow < 4u)) {
} else {
break;
}
{
uint innerCol = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
if ((innerCol < ColPerThreadA)) {
} else {
break;
@@ -164,14 +168,14 @@
}
{
uint innerRow = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
if ((innerRow < RowPerThreadB)) {
} else {
break;
}
{
uint innerCol = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_6) while(true) {
if ((innerCol < 4u)) {
} else {
break;
@@ -195,14 +199,14 @@
threadgroup_barrier(mem_flags::mem_threadgroup);
{
uint k = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_7) while(true) {
if ((k < 64u)) {
} else {
break;
}
{
uint inner = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_8) while(true) {
if ((inner < 4u)) {
} else {
break;
@@ -216,7 +220,7 @@
}
{
uint innerRow = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_9) while(true) {
if ((innerRow < 4u)) {
} else {
break;
@@ -224,7 +228,7 @@
ACached = (*tint_module_vars.mm_Asub)[(tileRow + innerRow)][k];
{
uint innerCol = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_10) while(true) {
if ((innerCol < 4u)) {
} else {
break;
@@ -258,14 +262,14 @@
}
{
uint innerRow = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_11) while(true) {
if ((innerRow < 4u)) {
} else {
break;
}
{
uint innerCol = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_12) while(true) {
if ((innerCol < 4u)) {
} else {
break;
diff --git a/test/tint/bug/tint/942.wgsl.expected.ir.msl b/test/tint/bug/tint/942.wgsl.expected.ir.msl
index b694881..fb3264f 100644
--- a/test/tint/bug/tint/942.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.ir.msl
@@ -31,6 +31,10 @@
threadgroup tint_array<tint_array<float3, 256>, 4>* tile;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<tint_array<float3, 256>, 4> tint_symbol_1;
};
@@ -43,7 +47,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 1024u)) {
break;
@@ -64,14 +68,14 @@
uint2 const baseIndex = (v_4 - uint2(filterOffset, 0u));
{
uint r = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((r < 4u)) {
} else {
break;
}
{
uint c = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
if ((c < 4u)) {
} else {
break;
@@ -99,14 +103,14 @@
threadgroup_barrier(mem_flags::mem_threadgroup);
{
uint r = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
if ((r < 4u)) {
} else {
break;
}
{
uint c = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
if ((c < 4u)) {
} else {
break;
@@ -132,7 +136,7 @@
float3 acc = float3(0.0f);
{
uint f = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
if ((f < (*tint_module_vars.params).filterDim)) {
} else {
break;
diff --git a/test/tint/bug/tint/948.wgsl.expected.ir.msl b/test/tint/bug/tint/948.wgsl.expected.ir.msl
index 7763667..af5b423 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl
@@ -48,6 +48,10 @@
thread float2* vUV;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct main_out {
float4 glFragColor_1;
};
@@ -118,7 +122,7 @@
stageUnits = (float2(1.0f) / x_111);
i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
int const x_122 = i;
if ((x_122 < 2)) {
} else {
@@ -159,7 +163,7 @@
(*tint_module_vars.mt) = fmod((x_181 * x_184), 1.0f);
f = 0.0f;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
float const x_193 = f;
if ((x_193 < 8.0f)) {
} else {
diff --git a/test/tint/bug/tint/990.wgsl.expected.ir.msl b/test/tint/bug/tint/990.wgsl.expected.ir.msl
index c72f776..b2ba140 100644
--- a/test/tint/bug/tint/990.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.ir.msl
@@ -1,11 +1,15 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
int i = 0;
{
thread int* const p = (&i);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
index 3f11ef4..57f860a 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
};
@@ -35,7 +39,7 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (!((idx < 6u))) {
break;
}
@@ -64,7 +68,7 @@
{
uint v_3 = 0u;
v_3 = local_invocation_index_1_param;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_4 = v_3;
if ((v_4 >= 6u)) {
break;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
index 325a251..ce7ebc2 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
@@ -17,6 +17,10 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
};
@@ -25,7 +29,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 6u)) {
break;
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
index d66b4cd..8a9a376 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<atomic_uint, 4>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<atomic_uint, 4> tint_symbol;
};
@@ -26,7 +30,7 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (!((idx < 4u))) {
break;
}
@@ -51,7 +55,7 @@
{
uint v = 0u;
v = local_invocation_index_1_param;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
index 04f2299..62f1481 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
@@ -17,6 +17,10 @@
threadgroup tint_array<atomic_uint, 4>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<atomic_uint, 4> tint_symbol;
};
@@ -25,7 +29,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
index 3f11ef4..57f860a 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
@@ -18,6 +18,10 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
};
@@ -35,7 +39,7 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (!((idx < 6u))) {
break;
}
@@ -64,7 +68,7 @@
{
uint v_3 = 0u;
v_3 = local_invocation_index_1_param;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_4 = v_3;
if ((v_4 >= 6u)) {
break;
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
index 325a251..ce7ebc2 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
@@ -17,6 +17,10 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
};
@@ -25,7 +29,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 6u)) {
break;
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
index ee08efb..cefd49b 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
@@ -24,6 +24,10 @@
threadgroup tint_array<S_atomic, 10>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S_atomic, 10> tint_symbol;
};
@@ -32,7 +36,7 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (!((idx < 10u))) {
break;
}
@@ -59,7 +63,7 @@
{
uint v = 0u;
v = local_invocation_index_1_param;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_1 = v;
if ((v_1 >= 10u)) {
break;
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
index a4f77ac..a3173ad 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
@@ -23,6 +23,10 @@
threadgroup tint_array<S, 10>* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
tint_array<S, 10> tint_symbol;
};
@@ -31,7 +35,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 10u)) {
break;
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
index 8e5fa30..2f3f91d 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
@@ -24,6 +24,10 @@
threadgroup S_atomic* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
S_atomic tint_symbol;
};
@@ -34,7 +38,7 @@
(*tint_module_vars.wg).y = 0u;
idx = local_invocation_index_2;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (!((idx < 10u))) {
break;
}
@@ -63,7 +67,7 @@
{
uint v = 0u;
v = local_invocation_index_1_param;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
uint const v_1 = v;
if ((v_1 >= 10u)) {
break;
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
index c3f6974..0e4dbc6 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
@@ -23,6 +23,10 @@
threadgroup S* wg;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_1 {
S tint_symbol;
};
@@ -35,7 +39,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 10u)) {
break;
diff --git a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
index 92032aa..031e5f1 100644
--- a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
+++ b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
@@ -5,10 +5,14 @@
texture2d<int, access::read_write> tex;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void foo(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 3)) {
} else {
break;
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
index 3a15b6b..4f47486 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
@@ -6,10 +6,14 @@
threadgroup int* b;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void foo(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
int const v = i;
threadgroup_barrier(mem_flags::mem_threadgroup);
int const v_1 = (*tint_module_vars.a);
diff --git a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
index e736db7..ce5135c 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
@@ -13,6 +13,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -20,7 +24,7 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
bool v_1 = false;
if ((x > v[0u])) {
v_1 = (dfdx(1.0f) > 0.0f);
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
index 34ccd45..3d98bff 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
sampler s;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -25,7 +29,7 @@
void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
float4 v = float4(0.0f);
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((x > v[0u])) {
} else {
break;
diff --git a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
index 507f55c..eefa0a6 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
@@ -13,13 +13,17 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
dfdx(1.0f);
{
if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
index 8d45d19..5aca6c4 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
@@ -13,13 +13,17 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
dfdx(1.0f);
{
if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
index 8501abe..d05523a 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
@@ -13,13 +13,17 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
{
dfdx(1.0f);
if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
index 246451e..0edc31c 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
@@ -13,6 +13,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -20,7 +24,7 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
bool v_1 = false;
if ((x > 0.0f)) {
v_1 = (dfdx(1.0f) > 0.0f);
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
index beef7de..8e828b2 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
@@ -18,6 +18,10 @@
sampler s;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -25,7 +29,7 @@
void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
float4 v = float4(0.0f);
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((x > v[0u])) {
} else {
break;
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
index e23ac2f..b966a93 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
@@ -18,6 +18,10 @@
/* 0x0008 */ tint_array<int8_t, 8> tint_pad;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct SSBO {
/* 0x0000 */ tint_array<strided_arr, 2> m;
};
@@ -43,7 +47,7 @@
{
uint v_1 = 0u;
v_1 = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_2 = v_1;
if ((v_2 >= 2u)) {
break;
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
index 9114b25..b030c35 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void f() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
index d73f2a3..e5a8abd 100644
--- a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
bool tint_continue = false;
switch(i) {
case 0:
diff --git a/test/tint/loops/loop.wgsl.expected.ir.msl b/test/tint/loops/loop.wgsl.expected.ir.msl
index 777ab74..4b11fde 100644
--- a/test/tint/loops/loop.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
i = (i + 1);
if ((i > 4)) {
return i;
diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
index 164cacc..5be8b31 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i > 4)) {
return i;
}
diff --git a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
index 66d9c1c..ca6575b 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i > 4)) {
return i;
}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
index c5e78ae..d75eccc 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
index d8a8569..24b152c 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
@@ -1,11 +1,15 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
int i = 0;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
index 548ba27..6309932 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
@@ -1,17 +1,21 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
}
{
int j = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((j < 2)) {
} else {
break;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
index ae0b853..7e16bb1 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
@@ -14,7 +18,7 @@
{
{
int j = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((j < 2)) {
} else {
break;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
index 6c06c1e..73f3fab 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
@@ -1,11 +1,15 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
int k = 0;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
@@ -15,7 +19,7 @@
{
{
int j = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((j < 2)) {
} else {
break;
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
index 703e006..c48d971 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
@@ -1,11 +1,15 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
int j = 0;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/loops/nested_loops.wgsl.expected.ir.msl b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
index 939ab1f..fe853a5 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
@@ -1,17 +1,21 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
int j = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
i = (i + 1);
if ((i > 4)) {
return 1;
}
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
j = (j + 1);
if ((j > 4)) {
return 2;
diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
index 53fe43a..71435a3 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
@@ -1,16 +1,20 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
int j = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i > 4)) {
return 1;
}
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
if ((j > 4)) {
return 2;
}
diff --git a/test/tint/loops/single_continue.wgsl.expected.ir.msl b/test/tint/loops/single_continue.wgsl.expected.ir.msl
index 7fecaf4..6f505df 100644
--- a/test/tint/loops/single_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 2)) {
} else {
break;
diff --git a/test/tint/loops/while.wgsl.expected.ir.msl b/test/tint/loops/while.wgsl.expected.ir.msl
index fdf1dd6..dbebdfb 100644
--- a/test/tint/loops/while.wgsl.expected.ir.msl
+++ b/test/tint/loops/while.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
index fdf1dd6..dbebdfb 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
index 0301bb0..88d037f 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
threadgroup str* S;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
str tint_symbol_1;
};
@@ -33,7 +37,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
index 021d045..6410c42 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
threadgroup tint_array<str, 4>* S;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<str, 4> tint_symbol_1;
};
@@ -33,7 +37,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
index 2b87be7..801cdb0 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
threadgroup str* S;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
str tint_symbol_1;
};
@@ -33,7 +37,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
index 8b0bbca..458af39 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -21,6 +21,10 @@
threadgroup tint_array<str, 4>* S;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<str, 4> tint_symbol_1;
};
@@ -33,7 +37,7 @@
{
uint v = 0u;
v = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/samples/compute_boids.wgsl.expected.ir.msl b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
index 1b4125b..408eb69 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.ir.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
@@ -38,6 +38,10 @@
device Particles* particlesB;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct vert_main_outputs {
float4 tint_symbol [[position]];
};
@@ -81,7 +85,7 @@
float2 vel = 0.0f;
{
uint i = 0u;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 5u)) {
} else {
break;
diff --git a/test/tint/shadowing/loop.wgsl.expected.ir.msl b/test/tint/shadowing/loop.wgsl.expected.ir.msl
index 6a50233..f34d31d 100644
--- a/test/tint/shadowing/loop.wgsl.expected.ir.msl
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.msl
@@ -17,11 +17,15 @@
device tint_array<int, 10>* output;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void foo(device tint_array<int, 10>* output [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.output=output};
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
int x = (*tint_module_vars.output)[i];
{
int x = (*tint_module_vars.output)[x];
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
index 1c0c0c2..632ed2a 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
@@ -29,13 +29,17 @@
tint_array<InnerS, 8> a1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol(const constant Uniforms* uniforms [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
InnerS v = {};
OuterS s1 = {};
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
index 69d664d..28e960a 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
@@ -29,13 +29,17 @@
tint_array<InnerS, 8> a1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol(const constant Uniforms* uniforms [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
InnerS v = {};
OuterS s1 = {};
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
index f7c2674..5f983f3 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
@@ -29,6 +29,10 @@
tint_array<InnerS, 8> a1;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
kernel void tint_symbol(const constant Uniforms* uniforms [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
InnerS v = {};
@@ -36,7 +40,7 @@
int i = 0;
{
s1.a1[(*tint_module_vars.uniforms).i] = v;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
index 4ec904b..c84338f 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
@@ -24,6 +24,10 @@
T elements[N];
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int idx1(tint_module_vars_struct tint_module_vars) {
(*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
return 1;
@@ -44,7 +48,7 @@
{
thread float* const v_1 = (&a[idx1(tint_module_vars)]);
(*v_1) = ((*v_1) * 2.0f);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((a[idx2(tint_module_vars)] < 10.0f)) {
} else {
break;
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
index 6054410..0c8623d 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
thread uint* v;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int idx1(tint_module_vars_struct tint_module_vars) {
(*tint_module_vars.v) = ((*tint_module_vars.v) - 1u);
return 1;
@@ -58,7 +62,7 @@
device int4* const v_2 = (&(*tint_module_vars.tint_symbol)[v_1].a[idx2(tint_module_vars)]);
int const v_3 = idx3(tint_module_vars);
(*v_2)[v_3] = ((*v_2)[v_3] - 1);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (((*tint_module_vars.v) < 10u)) {
} else {
break;
diff --git a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
index d61534f..a809d8b 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
@@ -5,9 +5,13 @@
device uint* i;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
diff --git a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
index 7bcb668..aecb1e3 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
@@ -5,10 +5,14 @@
device uint* i;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
(*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
index bb7625e..a742574 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
@@ -8,6 +8,10 @@
thread bool* continue_execution;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct foo_outputs {
int tint_symbol [[color(0)]];
};
@@ -28,7 +32,7 @@
int result = tint_f32_to_i32(tint_module_vars.t.sample(tint_module_vars.s, coord)[0u]);
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 10)) {
} else {
break;
diff --git a/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl b/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
index 26ef48b..44045f7 100644
--- a/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
@@ -5,9 +5,13 @@
thread bool* continue_execution;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f(tint_module_vars_struct tint_module_vars) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
(*tint_module_vars.continue_execution) = false;
return;
}
diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
index b57789e..3de5ec2 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
@@ -7,6 +7,10 @@
thread bool* continue_execution;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
fragment void tint_symbol(device int* non_uniform_global [[buffer(0)]], device float* output [[buffer(1)]]) {
thread bool continue_execution = true;
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.non_uniform_global=non_uniform_global, .output=output, .continue_execution=(&continue_execution)};
@@ -20,7 +24,7 @@
if (((*tint_module_vars.output) < 0.0f)) {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
float const v_1 = (*tint_module_vars.output);
if ((v_1 > float(i))) {
float const v_2 = float(i);
diff --git a/test/tint/statements/for/basic.wgsl.expected.ir.msl b/test/tint/statements/for/basic.wgsl.expected.ir.msl
index 03e6f6f..569ca8b 100644
--- a/test/tint/statements/for/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.ir.msl
@@ -1,13 +1,17 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void some_loop_body() {
}
void f() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 5)) {
} else {
break;
diff --git a/test/tint/statements/for/complex.wgsl.expected.ir.msl b/test/tint/statements/for/complex.wgsl.expected.ir.msl
index 70d2ffc..15f2f24 100644
--- a/test/tint/statements/for/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.ir.msl
@@ -1,6 +1,10 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void some_loop_body() {
}
@@ -8,7 +12,7 @@
int j = 0;
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
bool v = false;
if ((i < 5)) {
v = (j < 10);
diff --git a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
index 0acd4a1..3204d4a 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 1)) {
} else {
break;
diff --git a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
index c42baba..52e6a2f 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 4)) {
} else {
break;
diff --git a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
index 0acd4a1..3204d4a 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if ((i < 1)) {
} else {
break;
diff --git a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
index ec98763..43b21e1 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
index ec98763..43b21e1 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
int i = 0;
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
index bce67aa..3152937 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/empty.wgsl.expected.ir.msl b/test/tint/statements/for/empty.wgsl.expected.ir.msl
index 47641f0..986f6cc 100644
--- a/test/tint/statements/for/empty.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.ir.msl
@@ -1,9 +1,13 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
index db40de0..6be977a 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
int i = 1;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
index 61e7a43..b86deae 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
int i = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
index db40de0..6be977a 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
int i = 1;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (false) {
} else {
break;
diff --git a/test/tint/statements/for/scoping.wgsl.expected.ir.msl b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
index 2947505..efaba43 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
@@ -1,10 +1,14 @@
#include <metal_stdlib>
using namespace metal;
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void f() {
{
int must_not_collide = 0;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
break;
}
}
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.msl b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
index 40b5eb8..132300a 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
@@ -22,6 +22,10 @@
thread uint* v;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
int idx1(tint_module_vars_struct tint_module_vars) {
(*tint_module_vars.v) = ((*tint_module_vars.v) + 1u);
return 1;
@@ -58,7 +62,7 @@
device int4* const v_2 = (&(*tint_module_vars.tint_symbol)[v_1].a[idx2(tint_module_vars)]);
int const v_3 = idx3(tint_module_vars);
(*v_2)[v_3] = ((*v_2)[v_3] + 1);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (((*tint_module_vars.v) < 10u)) {
} else {
break;
diff --git a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
index 4620c07..8c17b12 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
@@ -5,9 +5,13 @@
device uint* i;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
diff --git a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
index afca7bb..0a1e473 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
@@ -5,10 +5,14 @@
device uint* i;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
(*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
diff --git a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
index 02aad41..fdd9f42 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
@@ -17,6 +17,10 @@
threadgroup tint_array<tint_array<int, 3>, 2>* zero;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<tint_array<int, 3>, 2> tint_symbol_1;
};
@@ -25,7 +29,7 @@
{
uint v_1 = 0u;
v_1 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_2 = v_1;
if ((v_2 >= 6u)) {
break;
diff --git a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
index ed037d3..a400836 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
@@ -17,6 +17,10 @@
threadgroup tint_array<int, 3>* zero;
};
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
struct tint_symbol_2 {
tint_array<int, 3> tint_symbol_1;
};
@@ -25,7 +29,7 @@
{
uint v_1 = 0u;
v_1 = tint_local_index;
- while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
uint const v_2 = v_1;
if ((v_2 >= 3u)) {
break;
