[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;