[msl ir] Convince the Metal compiler loops are never infinite
Bug: crbug.com/371840056
Change-Id: I3f335762fd5fcf7af4d296c7e132f138fbfc8863
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/210894
Commit-Queue: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/msl/writer/loop_test.cc b/src/tint/lang/msl/writer/loop_test.cc
index 330dd2c..f9e0d17 100644
--- a/src/tint/lang/msl/writer/loop_test.cc
+++ b/src/tint/lang/msl/writer/loop_test.cc
@@ -47,12 +47,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void a() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
break;
}
}
@@ -74,12 +74,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void a() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
{
if (true) { break; }
}
@@ -108,12 +108,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void a() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
bool v = true;
{
if (v) { break; }
@@ -144,13 +144,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void a() {
{
bool v = true;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
{
if (v) { break; }
}
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 401d3f4..aa84321 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -219,6 +219,12 @@
return array_template_name_;
}
+ /// Lazily generates the TINT_ISOLATE_UB macro, and returns a call to
+ /// the macro, passing in a unique identifier. The call tricks the MSL
+ /// compiler into thinking it might execute a `break`, but otherwise
+ /// has no effect in the generated code.
+ /// Invoke this inside the body of a loop to prevent the MSL compiler
+ /// from inferring the loop never terminates.
/// @returns a call to the TINT_ISOLATE_UB macro, creating that macro on first call
std::string IsolateUB() {
if (isolate_ub_macro_name_.empty()) {
@@ -226,11 +232,10 @@
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)";
+ Line() << " {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}";
}
StringStream ss;
- ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_true") << ")";
+ ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_false") << ")";
return ss.str();
}
@@ -668,9 +673,10 @@
ScopedIndent init(current_buffer_);
EmitBlock(l->Initializer());
- Line() << IsolateUB() << " while(true) {";
+ Line() << "while(true) {";
{
ScopedIndent si(current_buffer_);
+ Line() << IsolateUB();
EmitBlock(l->Body());
}
Line() << "}";
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 14a5111..fc71736 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
@@ -25,8 +25,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
@@ -63,7 +62,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 09ba79a..27c0213 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
@@ -27,8 +27,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
@@ -63,7 +62,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 eeaddb1..184604a 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
@@ -31,8 +31,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
@@ -67,7 +66,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 9b28283..033a240 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
@@ -27,8 +27,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_4 {
tint_array<int4, 4> tint_symbol_1;
@@ -65,7 +64,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
@@ -81,7 +81,8 @@
{
uint v_2 = 0u;
v_2 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 431a06a..b355790 100644
--- a/test/tint/array/strides.spvasm.expected.ir.msl
+++ b/test/tint/array/strides.spvasm.expected.ir.msl
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct strided_arr_1 {
/* 0x0000 */ tint_array<tint_array<strided_arr, 2>, 3> el;
@@ -43,7 +42,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 2u)) {
break;
@@ -61,7 +61,8 @@
{
uint v_2 = 0u;
v_2 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
uint const v_3 = v_2;
if ((v_3 >= 3u)) {
break;
@@ -83,7 +84,8 @@
{
uint v_4 = 0u;
v_4 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
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 351609c..d798fea 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct Inner_packed_vec3 {
/* 0x0000 */ float scalar_f32;
@@ -65,7 +64,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 378d780..40572dd 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 packed_1;
@@ -88,7 +87,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 df814f3..5fee146 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct Inner {
/* 0x0000 */ int scalar_i32;
@@ -69,7 +68,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ea90738..f6d86f4 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
@@ -21,8 +21,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 packed;
@@ -100,7 +99,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
@@ -118,7 +118,8 @@
{
uint v_2 = 0u;
v_2 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 ba3dfa8..4903a1f 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<float2x2, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 dec692e..0ef6ce0 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u;
@@ -36,7 +35,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 9940bdc0..9ba26c5b 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
@@ -25,8 +25,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4> tint_symbol;
@@ -62,7 +61,8 @@
{
uint v_11 = 0u;
v_11 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 0a23ec0..010faf1 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u;
@@ -36,7 +35,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 121862f..6eede57 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
@@ -24,8 +24,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4> tint_symbol;
@@ -61,7 +60,8 @@
{
uint v_11 = 0u;
v_11 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3a70542..430020c 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<half2x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 e3cde16..61fc3d3 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<float2x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 33d13f1..8c5e9a6 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u;
@@ -37,7 +36,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 026c275..b3d3a51 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
@@ -24,8 +24,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4> tint_symbol;
@@ -69,7 +68,8 @@
{
uint v_15 = 0u;
v_15 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 aee2d2e..1a6b5dc 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<float3x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 b4321d6..c99d751 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<half4x2, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ab5e665..e185b2a 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<float4x2, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 b00d8c5..76face9 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u;
@@ -38,7 +37,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 52c9321..e2ab200 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
@@ -24,8 +24,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4> tint_symbol;
@@ -77,7 +76,8 @@
{
uint v_19 = 0u;
v_19 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 a6ad8aa..125ce31 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u;
@@ -38,7 +37,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 1d30101..13068e7 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
@@ -24,8 +24,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4> tint_symbol;
@@ -77,7 +76,8 @@
{
uint v_19 = 0u;
v_19 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 4f0e022..3505775 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<half4x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 35ccec5..440abbd 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<float4x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 89d8544..8b67dc7 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 201ecd7..02b3af7 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
@@ -27,8 +27,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 f51c619..7ebea3d 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 81a2c8d..1ebc4a5 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 2fa25c3..a4fd815 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
@@ -34,8 +34,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
@@ -65,7 +64,8 @@
{
uint v_4 = 0u;
v_4 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8ba6fb0..e3ef41a 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
@@ -39,8 +39,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_packed_vec3, 4> tint_symbol;
@@ -79,7 +78,8 @@
{
uint v_7 = 0u;
v_7 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 f080673..1ea3778 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
@@ -34,8 +34,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
@@ -65,7 +64,8 @@
{
uint v_4 = 0u;
v_4 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ccddfbf..e8fa2b3 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
@@ -39,8 +39,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_packed_vec3, 4> tint_symbol;
@@ -79,7 +78,8 @@
{
uint v_7 = 0u;
v_7 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 913cd10..851d6d2 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ad9574a..9b35ae4 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 5d712dc..0c3eb5b 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 7d3170c..2d3491e 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 e70ceb1..bf52720 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 35f763e..44ff961 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
@@ -27,8 +27,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 e425614..bba164c 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 b845d9e..d3191c8 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 07623e3..3c4752f 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
@@ -34,8 +34,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
@@ -67,7 +66,8 @@
{
uint v_5 = 0u;
v_5 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 1316725..bc96eb1 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
@@ -39,8 +39,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_packed_vec3, 4> tint_symbol;
@@ -81,7 +80,8 @@
{
uint v_8 = 0u;
v_8 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 1d18969d..083aa88 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
@@ -33,8 +33,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
@@ -66,7 +65,8 @@
{
uint v_5 = 0u;
v_5 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 c4cc786..70ffec1 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
@@ -38,8 +38,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_packed_vec3, 4> tint_symbol;
@@ -80,7 +79,8 @@
{
uint v_8 = 0u;
v_8 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6b56138..31faeaa 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 0b5d8a2..8153372 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8ff7e82..fb2085f 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6c637af..acad8ac 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
@@ -27,8 +27,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 139cb9e..2355cd3 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6f8d3af..9510bf9 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
@@ -27,8 +27,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 0a827db..ca17dbd 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 9a43ba5..04e94dc 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 c71d8eb..8dd5d90 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
@@ -34,8 +34,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
@@ -69,7 +68,8 @@
{
uint v_6 = 0u;
v_6 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 0c76569..98652ab 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
@@ -39,8 +39,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_packed_vec3, 4> tint_symbol;
@@ -83,7 +82,8 @@
{
uint v_9 = 0u;
v_9 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 d9bf522..79faedd 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
@@ -34,8 +34,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S_packed_vec3, 4>* u;
@@ -69,7 +68,8 @@
{
uint v_6 = 0u;
v_6 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 eb89dec..e1ca55e 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
@@ -39,8 +39,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_packed_vec3, 4> tint_symbol;
@@ -83,7 +82,8 @@
{
uint v_9 = 0u;
v_9 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 198b3a0..494cafd 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ea18e24..a5d3929 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 473a54d..8bc1c3a 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
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 08984e5..9e2e17a 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
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 fe4107f..e8c2e82 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void d() {
int j = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3e5d3f6..b6f1da3 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
int i0520 = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if (false) {
} else {
break;
@@ -20,7 +20,8 @@
}
{
int i62 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
if (false) {
} else {
break;
@@ -32,7 +33,8 @@
}
{
int i0520 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
if (false) {
} else {
break;
@@ -44,7 +46,8 @@
}
{
int i62 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_3)
if (false) {
} else {
break;
@@ -56,7 +59,8 @@
}
{
int i62 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_4)
if (false) {
} else {
break;
@@ -68,7 +72,8 @@
}
{
int i60 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_5)
if (false) {
} else {
break;
@@ -80,7 +85,8 @@
}
{
int i62 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_6) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_6)
if (false) {
} else {
break;
@@ -92,7 +98,8 @@
}
{
int i60 = 0;
- TINT_ISOLATE_UB(tint_volatile_true_7) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_7)
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 4617798..f1b46a3 100644
--- a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
@@ -14,8 +14,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
device tint_array<int, 4>* tint_symbol;
@@ -24,7 +23,8 @@
void foo(device tint_array<int, 4>* const arg) {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3a5398f..b8b2bc7 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
@@ -32,8 +32,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
S tint_symbol;
@@ -43,7 +42,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 d725486..c5c0ef0 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
@@ -32,8 +32,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
S tint_symbol;
@@ -43,7 +42,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 69b47a5..9ad691d 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
@@ -9,13 +9,13 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void foo(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((i < 2)) {
} else {
break;
@@ -40,7 +40,8 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v2f=(&v2f), .v3i=(&v3i), .v4u=(&v4u), .v2b=(&v2b)};
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 4eb9cf4..0e0353d 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
@@ -9,8 +9,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void foo(tint_module_vars_struct tint_module_vars) {
int i = 0;
@@ -28,7 +27,8 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v2f=(&v2f), .v3i=(&v3i), .v4u=(&v4u), .v2b=(&v2b)};
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 34a3cd7..0c1bd3b 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
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
float2 v2f = 0.0f;
@@ -20,7 +19,8 @@
bool4 v4b = false;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3484c4e..44bc789 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
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
float2 v2f = 0.0f;
@@ -16,7 +15,8 @@
bool2 v2b_2 = false;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 70f3124..e33ec1c 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
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
float2 v2f = 0.0f;
@@ -20,7 +19,8 @@
bool4 v4b = false;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 e728d67..d7f982f 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
@@ -2,12 +2,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
fragment void tint_symbol() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8bf645b..f562caa 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
@@ -6,8 +6,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_outputs {
int tint_symbol_1 [[color(2)]];
@@ -27,7 +26,8 @@
int tint_symbol_inner(int3 x, tint_module_vars_struct tint_module_vars) {
int y = x[0u];
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 b0ca2dd..9c99a38 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
@@ -57,8 +57,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
uint index = GlobalInvocationID[0u];
@@ -91,14 +90,16 @@
int const TILE_COUNT_Y = 2;
{
int y = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((y < TILE_COUNT_Y)) {
} else {
break;
}
{
int x = 0;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
if ((x < TILE_COUNT_X)) {
} else {
break;
@@ -118,7 +119,8 @@
float dp = 0.0f;
{
uint i = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
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 537c53a..58ac15d 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
@@ -14,8 +14,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int foo() {
return 1;
@@ -25,7 +24,8 @@
tint_array<float, 4> arr = tint_array<float, 4>{};
{
thread float* const a = (&arr[foo()]);
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 9444e09..4d10141 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
@@ -2,12 +2,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 2c5da95..102080b 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
template<typename T, size_t N>
struct tint_array {
@@ -27,7 +26,8 @@
int f() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
g();
break;
}
@@ -39,7 +39,8 @@
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};
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 ad98ec0..adc3222 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant int* u;
@@ -16,7 +15,8 @@
void g() {
int j = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 80c92e3..afba81b 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
@@ -6,8 +6,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol(const constant int* x [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x=x};
@@ -15,7 +14,8 @@
case 0:
{
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
return;
}
}
diff --git a/test/tint/bug/tint/1605.wgsl.expected.ir.msl b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
index 11318e0f..15bb159 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
@@ -6,20 +6,21 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
bool func_3(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((i < (*tint_module_vars.b))) {
} else {
break;
}
{
int j = -1;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 49a49c8..f23a45c 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
@@ -18,8 +18,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<int, 246> tint_symbol_1;
@@ -29,7 +28,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3e1fb27..51fd847 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.ir.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.ir.msl
@@ -47,8 +47,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_6 {
tint_array<S, 4096> tint_symbol_1;
@@ -66,7 +65,8 @@
uint const x_52 = (*tint_module_vars.x_3)[0u];
x_54 = 0u;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint x_55 = 0u;
x_58 = (*tint_module_vars.x_6).field0.field0;
if ((x_54 < x_58)) {
@@ -100,7 +100,8 @@
x_85 = x_76.xyxy;
x_88 = 1u;
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
float4 x_111 = 0.0f;
float4 x_86 = 0.0f;
uint x_89 = 0u;
@@ -157,7 +158,8 @@
{
uint v_3 = 0u;
v_3 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
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 8401406a..11378a4 100644
--- a/test/tint/bug/tint/2039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
uint out = 0u;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 0d038a4..7fdf29c 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct S2_packed_vec3 {
/* 0x0000 */ tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1> m;
@@ -75,7 +74,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 1u)) {
break;
@@ -97,7 +97,8 @@
{
uint v_2 = 0u;
v_2 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
uint const v_3 = v_2;
if ((v_3 >= 1u)) {
break;
@@ -119,7 +120,8 @@
{
uint v_4 = 0u;
v_4 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
uint const v_5 = v_4;
if ((v_5 >= 1u)) {
break;
@@ -146,7 +148,8 @@
float3x3 m = float3x3(0.0f);
{
uint c = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_3)
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 aae30ea..2cc5558 100644
--- a/test/tint/bug/tint/2201.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
@@ -6,12 +6,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 f4bd798..b7f4c947 100644
--- a/test/tint/bug/tint/2202.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2202.wgsl.expected.ir.msl
@@ -6,14 +6,15 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
return;
}
}
diff --git a/test/tint/bug/tint/221.wgsl.expected.ir.msl b/test/tint/bug/tint/221.wgsl.expected.ir.msl
index 5fb7153..0025dd0 100644
--- a/test/tint/bug/tint/221.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.ir.msl
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
uint tint_mod_u32(uint lhs, uint rhs) {
uint const v = select(rhs, 1u, (rhs == 0u));
@@ -35,7 +34,8 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.b=b};
uint i = 0u;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 be80066..57661c0 100644
--- a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
@@ -6,14 +6,14 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
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();
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 568f3b8..fb5f664 100644
--- a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
@@ -6,17 +6,18 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
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);
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
{
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
if ((i > 5)) {
i = as_type<int>((as_type<uint>(i) * as_type<uint>(2)));
break;
diff --git a/test/tint/bug/tint/366037039.wgsl.expected.ir.msl b/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
index 3da7d69..85977d4 100644
--- a/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
@@ -30,8 +30,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_module_vars_struct {
const constant S_packed_vec3* ubuffer;
@@ -56,7 +55,8 @@
{
uint v = 0u;
v = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/bug/tint/534.wgsl.expected.ir.msl b/test/tint/bug/tint/534.wgsl.expected.ir.msl
index bf985e2..270c65a 100644
--- a/test/tint/bug/tint/534.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.ir.msl
@@ -32,8 +32,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
uint ConvertToFp16FloatValue(float fp32) {
return 1u;
@@ -58,7 +57,8 @@
uint4 dstColorBits = tint_v4f32_to_v4u32(dstColor);
{
uint i = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 f86490f..1a98eae 100644
--- a/test/tint/bug/tint/744.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.ir.msl
@@ -31,8 +31,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol_inner(uint3 global_id, tint_module_vars_struct tint_module_vars) {
uint2 const resultCell = uint2(global_id[1u], global_id[0u]);
@@ -41,7 +40,8 @@
uint result = 0u;
{
uint i = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 c4b51b5..0acd061 100644
--- a/test/tint/bug/tint/757.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/757.wgsl.expected.ir.msl
@@ -28,8 +28,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
uint flatIndex = (((4u * GlobalInvocationID[2u]) + (2u * GlobalInvocationID[1u])) + GlobalInvocationID[0u]);
@@ -37,7 +36,8 @@
float4 texel = tint_module_vars.myTexture.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
{
uint i = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ebedfcc..ac94aea 100644
--- a/test/tint/bug/tint/914.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.ir.msl
@@ -33,8 +33,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_3 {
tint_array<tint_array<float, 64>, 64> tint_symbol_1;
@@ -90,7 +89,8 @@
{
uint v_3 = 0u;
v_3 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_4 = v_3;
if ((v_4 >= 4096u)) {
break;
@@ -114,7 +114,8 @@
tint_array<float, 4> BCached = {};
{
uint index = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
if ((index < 16u)) {
} else {
break;
@@ -132,21 +133,24 @@
uint const tileRowB = (local_id[1u] * RowPerThreadB);
{
uint t = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
if ((t < numTiles)) {
} else {
break;
}
{
uint innerRow = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_3)
if ((innerRow < 4u)) {
} else {
break;
}
{
uint innerCol = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_4)
if ((innerCol < ColPerThreadA)) {
} else {
break;
@@ -168,14 +172,16 @@
}
{
uint innerRow = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_5)
if ((innerRow < RowPerThreadB)) {
} else {
break;
}
{
uint innerCol = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_6) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_6)
if ((innerCol < 4u)) {
} else {
break;
@@ -199,14 +205,16 @@
threadgroup_barrier(mem_flags::mem_threadgroup);
{
uint k = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_7) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_7)
if ((k < 64u)) {
} else {
break;
}
{
uint inner = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_8) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_8)
if ((inner < 4u)) {
} else {
break;
@@ -220,7 +228,8 @@
}
{
uint innerRow = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_9) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_9)
if ((innerRow < 4u)) {
} else {
break;
@@ -228,7 +237,8 @@
ACached = (*tint_module_vars.mm_Asub)[(tileRow + innerRow)][k];
{
uint innerCol = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_10) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_10)
if ((innerCol < 4u)) {
} else {
break;
@@ -262,14 +272,16 @@
}
{
uint innerRow = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_11) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_11)
if ((innerRow < 4u)) {
} else {
break;
}
{
uint innerCol = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_12) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_12)
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 d053cc5..998db89 100644
--- a/test/tint/bug/tint/942.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.ir.msl
@@ -36,8 +36,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<tint_array<tint_packed_vec3_f32_array_element, 256>, 4> tint_symbol_1;
@@ -51,7 +50,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 1024u)) {
break;
@@ -72,14 +72,16 @@
uint2 const baseIndex = (v_4 - uint2(filterOffset, 0u));
{
uint r = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
if ((r < 4u)) {
} else {
break;
}
{
uint c = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_2)
if ((c < 4u)) {
} else {
break;
@@ -107,14 +109,16 @@
threadgroup_barrier(mem_flags::mem_threadgroup);
{
uint r = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_3)
if ((r < 4u)) {
} else {
break;
}
{
uint c = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_4)
if ((c < 4u)) {
} else {
break;
@@ -140,7 +144,8 @@
float3 acc = float3(0.0f);
{
uint f = 0u;
- TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_5)
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 5a500d2..d738f09 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl
@@ -49,8 +49,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct main_out {
float4 glFragColor_1;
@@ -122,7 +121,8 @@
stageUnits = (float2(1.0f) / x_111);
i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
int const x_122 = i;
if ((x_122 < 2)) {
} else {
@@ -163,7 +163,8 @@
(*tint_module_vars.mt) = fmod((x_181 * x_184), 1.0f);
f = 0.0f;
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 b2ba140..b225169 100644
--- a/test/tint/bug/tint/990.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.ir.msl
@@ -2,14 +2,14 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
int i = 0;
{
thread int* const p = (&i);
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 57f860a..e14cb85 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -39,7 +38,8 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if (!((idx < 6u))) {
break;
}
@@ -68,7 +68,8 @@
{
uint v_3 = 0u;
v_3 = local_invocation_index_1_param;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 ce7ebc2..85326e1 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
@@ -18,8 +18,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -29,7 +28,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8a9a376..cc5f715 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<atomic_uint, 4> tint_symbol;
@@ -30,7 +29,8 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if (!((idx < 4u))) {
break;
}
@@ -55,7 +55,8 @@
{
uint v = 0u;
v = local_invocation_index_1_param;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 62f1481..eede400 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
@@ -18,8 +18,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<atomic_uint, 4> tint_symbol;
@@ -29,7 +28,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 57f860a..e14cb85 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -39,7 +38,8 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if (!((idx < 6u))) {
break;
}
@@ -68,7 +68,8 @@
{
uint v_3 = 0u;
v_3 = local_invocation_index_1_param;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 ce7ebc2..85326e1 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
@@ -18,8 +18,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -29,7 +28,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 cefd49b..ca989a3 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
@@ -25,8 +25,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S_atomic, 10> tint_symbol;
@@ -36,7 +35,8 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if (!((idx < 10u))) {
break;
}
@@ -63,7 +63,8 @@
{
uint v = 0u;
v = local_invocation_index_1_param;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 a3173ad..6140181 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
@@ -24,8 +24,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
tint_array<S, 10> tint_symbol;
@@ -35,7 +34,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 2f3f91d..1ef5578 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
@@ -25,8 +25,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
S_atomic tint_symbol;
@@ -38,7 +37,8 @@
(*tint_module_vars.wg).y = 0u;
idx = local_invocation_index_2;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if (!((idx < 10u))) {
break;
}
@@ -67,7 +67,8 @@
{
uint v = 0u;
v = local_invocation_index_1_param;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 0e4dbc6..e123238 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
@@ -24,8 +24,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_1 {
S tint_symbol;
@@ -39,7 +38,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 84b3346..c0a5876 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
@@ -6,13 +6,13 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void foo(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 efd369f..df06f9b 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
@@ -7,13 +7,13 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void foo(tint_module_vars_struct tint_module_vars) {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ce5135c..49b79da 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
@@ -14,8 +14,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -24,7 +23,8 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3d98bff..7d349b6 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -29,7 +28,8 @@
void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
float4 v = float4(0.0f);
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 eefa0a6..de7b62f 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
@@ -14,8 +14,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -23,7 +22,8 @@
void tint_symbol_inner(float x) {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 5aca6c4..18c03d7 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
@@ -14,8 +14,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -23,7 +22,8 @@
void tint_symbol_inner(float x) {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 d05523a..efa9d91 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
@@ -14,8 +14,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -23,7 +22,8 @@
void tint_symbol_inner(float x) {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
{
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 0edc31c..4c7af50 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
@@ -14,8 +14,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -24,7 +23,8 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8e828b2..93cc724 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_inputs {
float x [[user(locn0)]];
@@ -29,7 +28,8 @@
void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
float4 v = float4(0.0f);
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 b966a93..6a9c791 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
@@ -19,8 +19,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct SSBO {
/* 0x0000 */ tint_array<strided_arr, 2> m;
@@ -47,7 +46,8 @@
{
uint v_1 = 0u;
v_1 = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 ec8684d..8c5fb99 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void f() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 5755d1f..d9fbf58 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 5062fd6..d5e5fa2 100644
--- a/test/tint/loops/loop.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return i;
diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
index 6348b86..b32bbacb 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6d7a2b3..e21dbe4 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8ef35c5..a0c6a80 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 2a0046b..5edf8d5 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
@@ -2,14 +2,14 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
int i = 0;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 1eff924..3090c93 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
@@ -2,20 +2,21 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((i < 2)) {
} else {
break;
}
{
int j = 0;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 08e2cec..eb3d97c 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((i < 2)) {
} else {
break;
@@ -19,7 +19,8 @@
{
{
int j = 0;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 14e9c8b..6bc5dec 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
@@ -2,14 +2,14 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
int k = 0;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((i < 2)) {
} else {
break;
@@ -20,7 +20,8 @@
{
{
int j = 0;
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 6ed9ba5..208823c 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
@@ -2,14 +2,14 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
int j = 0;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 fe13807..daa4a15 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
@@ -2,20 +2,21 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
int j = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return 1;
}
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
if ((j > 4)) {
return 2;
diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
index 1e711ce..8d2609f 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
@@ -2,19 +2,20 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
int j = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
if ((i > 4)) {
return 1;
}
{
- TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false_1)
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 f830ab8..5321463 100644
--- a/test/tint/loops/single_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 a2e7eb1..62623af 100644
--- a/test/tint/loops/while.wgsl.expected.ir.msl
+++ b/test/tint/loops/while.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 a2e7eb1..62623af 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 88d037f..fbcdcaa 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
str tint_symbol_1;
@@ -37,7 +36,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6410c42..a50727f 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<str, 4> tint_symbol_1;
@@ -37,7 +36,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 801cdb0..e3a3a94 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
str tint_symbol_1;
@@ -37,7 +36,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 458af39..cb2fb04 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
@@ -22,8 +22,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<str, 4> tint_symbol_1;
@@ -37,7 +36,8 @@
{
uint v = 0u;
v = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 394da95..5144d1b 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.ir.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
@@ -39,8 +39,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct vert_main_outputs {
float4 tint_symbol [[position]];
@@ -85,7 +84,8 @@
float2 vel = 0.0f;
{
uint i = 0u;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 32f654a..bf61e09 100644
--- a/test/tint/shadowing/loop.wgsl.expected.ir.msl
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.msl
@@ -18,14 +18,14 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
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;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 c4065b5..a5ef9e6 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
@@ -30,8 +30,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol(const constant Uniforms* uniforms [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
@@ -39,7 +38,8 @@
OuterS s1 = {};
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 5c483ed..140855f 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
@@ -30,8 +30,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol(const constant Uniforms* uniforms [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
@@ -39,7 +38,8 @@
OuterS s1 = {};
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 1226caa..b2a5dc7 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
@@ -30,8 +30,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
kernel void tint_symbol(const constant Uniforms* uniforms [[buffer(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
@@ -40,7 +39,8 @@
int i = 0;
{
s1.a1[(*tint_module_vars.uniforms).i] = v;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 c84338f..f3c5980 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
@@ -25,8 +25,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int idx1(tint_module_vars_struct tint_module_vars) {
(*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
@@ -48,7 +47,8 @@
{
thread float* const v_1 = (&a[idx1(tint_module_vars)]);
(*v_1) = ((*v_1) * 2.0f);
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 d932ca4..e9505f5 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int idx1(tint_module_vars_struct tint_module_vars) {
(*tint_module_vars.v) = ((*tint_module_vars.v) - 1u);
@@ -62,7 +61,8 @@
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] = as_type<int>((as_type<uint>((*v_2)[v_3]) - as_type<uint>(1)));
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 a809d8b..df04b49 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
@@ -6,12 +6,12 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 aecb1e3..6bb40ad 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
@@ -6,13 +6,13 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
(*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 39ed593..7e4bc07 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
@@ -9,8 +9,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct foo_outputs {
int tint_symbol [[color(0)]];
@@ -32,7 +31,8 @@
int result = tint_f32_to_i32(tint_module_vars.t.sample(tint_module_vars.s, coord)[0u]);
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 44045f7..752536e 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
@@ -6,12 +6,12 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f(tint_module_vars_struct tint_module_vars) {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
(*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 7303f09..1c9efef 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
@@ -8,8 +8,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
fragment void tint_symbol(device int* non_uniform_global [[buffer(0)]], device float* output [[buffer(1)]]) {
thread bool continue_execution = true;
@@ -24,7 +23,8 @@
if (((*tint_module_vars.output) < 0.0f)) {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 fb42656..97465a8 100644
--- a/test/tint/statements/for/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.ir.msl
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void some_loop_body() {
}
@@ -11,7 +10,8 @@
void f() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 16f4e4e..4dd60ee 100644
--- a/test/tint/statements/for/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.ir.msl
@@ -2,8 +2,7 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void some_loop_body() {
}
@@ -12,7 +11,8 @@
int j = 0;
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3204d4a..02a71f0 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 52e6a2f..fae57d4 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 3204d4a..02a71f0 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 9fe44a7..76a659f 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 9fe44a7..76a659f 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
int i = 0;
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 c5a4b75..ae63ee6 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 986f6cc..7b3f2ce 100644
--- a/test/tint/statements/for/empty.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.ir.msl
@@ -2,12 +2,12 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6be977a..a1a7b9a 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
int i = 1;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 b86deae..c3e6673 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
int i = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 6be977a..a1a7b9a 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
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
int i = 1;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 efaba43..5b08034 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
@@ -2,13 +2,13 @@
using namespace metal;
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void f() {
{
int must_not_collide = 0;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
break;
}
}
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.msl b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
index 7f62d2b..b4cc61f 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
@@ -23,8 +23,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
int idx1(tint_module_vars_struct tint_module_vars) {
(*tint_module_vars.v) = ((*tint_module_vars.v) + 1u);
@@ -62,7 +61,8 @@
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] = as_type<int>((as_type<uint>((*v_2)[v_3]) + as_type<uint>(1)));
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 8c17b12..1aa5603 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
@@ -6,12 +6,12 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 0a1e473..550762b 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
@@ -6,13 +6,13 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
(*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 fdd9f42..4d50db2 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
@@ -18,8 +18,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<tint_array<int, 3>, 2> tint_symbol_1;
@@ -29,7 +28,8 @@
{
uint v_1 = 0u;
v_1 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
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 a400836..1fc2b21 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
@@ -18,8 +18,7 @@
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- volatile bool VOLATILE_NAME = true; \
- if (VOLATILE_NAME)
+ {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
struct tint_symbol_2 {
tint_array<int, 3> tint_symbol_1;
@@ -29,7 +28,8 @@
{
uint v_1 = 0u;
v_1 = tint_local_index;
- TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ while(true) {
+ TINT_ISOLATE_UB(tint_volatile_false)
uint const v_2 = v_1;
if ((v_2 >= 3u)) {
break;
