[msl] Use PreventInfiniteLoops
This replaces the TINT_ISOLATE_UB macro.
Bug: 380103825
Change-Id: Iddfd3abfc443ed164b41fdc10b78351d0e4598e3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/217257
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: 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 54b06cc..4d42a30 100644
--- a/src/tint/lang/msl/writer/loop_test.cc
+++ b/src/tint/lang/msl/writer/loop_test.cc
@@ -52,13 +52,13 @@
EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void a() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
break;
}
}
@@ -102,14 +102,18 @@
EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void a() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if (true) { break; }
}
continue;
@@ -162,15 +166,19 @@
EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void a() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool v = true;
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if (v) { break; }
}
continue;
@@ -229,15 +237,19 @@
EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void a() {
{
+ uint2 tint_loop_idx = 0u;
bool v = true;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if (v) { break; }
}
continue;
diff --git a/src/tint/lang/msl/writer/printer/BUILD.bazel b/src/tint/lang/msl/writer/printer/BUILD.bazel
index dd1018d..0013a3f 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/printer/BUILD.bazel
@@ -50,7 +50,6 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
- "//src/tint/lang/core/ir/analysis",
"//src/tint/lang/core/type",
"//src/tint/lang/msl",
"//src/tint/lang/msl/intrinsic",
diff --git a/src/tint/lang/msl/writer/printer/BUILD.cmake b/src/tint/lang/msl/writer/printer/BUILD.cmake
index 8969473..abbc249 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/printer/BUILD.cmake
@@ -51,7 +51,6 @@
tint_lang_core_constant
tint_lang_core_intrinsic
tint_lang_core_ir
- tint_lang_core_ir_analysis
tint_lang_core_type
tint_lang_msl
tint_lang_msl_intrinsic
diff --git a/src/tint/lang/msl/writer/printer/BUILD.gn b/src/tint/lang/msl/writer/printer/BUILD.gn
index 5c9f83a..7ce6ec9 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.gn
+++ b/src/tint/lang/msl/writer/printer/BUILD.gn
@@ -51,7 +51,6 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
- "${tint_src_dir}/lang/core/ir/analysis",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/msl",
"${tint_src_dir}/lang/msl/intrinsic",
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 5bae351..f493452 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -36,7 +36,6 @@
#include "src/tint/lang/core/constant/splat.h"
#include "src/tint/lang/core/fluent_types.h"
#include "src/tint/lang/core/ir/access.h"
-#include "src/tint/lang/core/ir/analysis/loop_analysis.h"
#include "src/tint/lang/core/ir/bitcast.h"
#include "src/tint/lang/core/ir/break_if.h"
#include "src/tint/lang/core/ir/constant.h"
@@ -146,9 +145,7 @@
// Emit functions.
for (auto* func : ir_.DependencyOrderedFunctions()) {
- loop_analysis_ = std::make_unique<core::ir::analysis::LoopAnalysis>(*func);
EmitFunction(func);
- loop_analysis_.reset();
}
StringStream ss;
@@ -169,8 +166,6 @@
/// MSL writer options
Options options_;
- std::unique_ptr<core::ir::analysis::LoopAnalysis> loop_analysis_;
-
/// A hashmap of value to name
Hashmap<const core::ir::Value*, std::string, 32> names_;
@@ -193,10 +188,6 @@
/// Non-empty only if the template has been generated.
std::string array_template_name_;
- /// The name of the macro used to prevent UB affecting later control flow.
- /// Do not use this directly, instead call IsolateUB().
- std::string isolate_ub_macro_name_;
-
/// Block to emit for a continuing
std::function<void()> emit_continuing_;
@@ -230,26 +221,6 @@
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()) {
- TINT_SCOPED_ASSIGNMENT(current_buffer_, &preamble_buffer_);
- isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
- Line();
- Line() << "#define " << isolate_ub_macro_name_ << "(VOLATILE_NAME) \\";
- Line() << " {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}";
- }
- StringStream ss;
- ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_false") << ")";
- return ss.str();
- }
-
/// Find all structures that are used in host-shareable address spaces and mark them as such so
/// that we know to pad the properly when we emit them.
void FindHostShareableStructs() {
@@ -682,19 +653,11 @@
Line() << "{";
{
ScopedIndent init(current_buffer_);
-
- // Analyze the loop to determine if we need to guard against undefined behavior caused
- // by infinite loops.
- auto* info = loop_analysis_->GetInfo(*l);
-
EmitBlock(l->Initializer());
Line() << "while(true) {";
{
ScopedIndent si(current_buffer_);
- if (!options_.disable_robustness && !info->IsFinite()) {
- Line() << IsolateUB();
- }
EmitBlock(l->Body());
}
Line() << "}";
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index 2d878ba..7a8533f 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -38,6 +38,7 @@
#include "src/tint/lang/core/ir/transform/demote_to_helper.h"
#include "src/tint/lang/core/ir/transform/multiplanar_external_texture.h"
#include "src/tint/lang/core/ir/transform/preserve_padding.h"
+#include "src/tint/lang/core/ir/transform/prevent_infinite_loops.h"
#include "src/tint/lang/core/ir/transform/remove_continue_in_switch.h"
#include "src/tint/lang/core/ir/transform/remove_terminator_args.h"
#include "src/tint/lang/core/ir/transform/rename_conflicts.h"
@@ -74,6 +75,10 @@
array_length_from_uniform_options);
RUN_TRANSFORM(core::ir::transform::BindingRemapper, module, remapper_data);
+ if (!options.disable_robustness) {
+ RUN_TRANSFORM(core::ir::transform::PreventInfiniteLoops, module);
+ }
+
{
core::ir::transform::BinaryPolyfillConfig binary_polyfills{};
binary_polyfills.int_div_mod = !options.disable_polyfill_integer_div_mod;
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 fc71736..978fdfb 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
@@ -24,9 +24,6 @@
device S* src_storage;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
};
@@ -63,7 +60,6 @@
uint v = 0u;
v = tint_local_index;
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 27c0213..13c9933 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
@@ -26,9 +26,6 @@
thread tint_array<tint_array<tint_array<int, 2>, 3>, 4>* dst_nested;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
};
@@ -63,7 +60,6 @@
uint v = 0u;
v = tint_local_index;
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 184604a..a5d2af7 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
@@ -30,9 +30,6 @@
device S_nested* dst_nested;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<int4, 4> tint_symbol_1;
};
@@ -67,7 +64,6 @@
uint v = 0u;
v = tint_local_index;
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 033a240..cade83f 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
@@ -26,9 +26,6 @@
threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* dst_nested;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_4 {
tint_array<int4, 4> tint_symbol_1;
tint_array<int4, 4> tint_symbol_2;
@@ -65,7 +62,6 @@
uint v = 0u;
v = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
@@ -82,7 +78,6 @@
uint v_2 = 0u;
v_2 = tint_local_index;
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 e0de60a..0c8c288 100644
--- a/test/tint/array/strides.spvasm.expected.ir.msl
+++ b/test/tint/array/strides.spvasm.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x0004 */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(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;
/* 0x0030 */ tint_array<int8_t, 80> tint_pad_1;
@@ -43,7 +40,6 @@
uint v = 0u;
v = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 2u)) {
break;
@@ -62,7 +58,6 @@
uint v_2 = 0u;
v_2 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
uint const v_3 = v_2;
if ((v_3 >= 3u)) {
break;
@@ -85,7 +80,6 @@
uint v_4 = 0u;
v_4 = 0u;
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 6543623..2c3f4bf 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct Inner_packed_vec3 {
/* 0x0000 */ float scalar_f32;
/* 0x0004 */ int scalar_i32;
@@ -66,7 +63,6 @@
uint v = 0u;
v = 0u;
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 43c0d52..c3d0fdd 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 packed_1;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad_1;
@@ -89,7 +86,6 @@
uint v = 0u;
v = 0u;
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 5fee146..4b1417c 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct Inner {
/* 0x0000 */ int scalar_i32;
/* 0x0004 */ float scalar_f32;
@@ -69,7 +66,6 @@
uint v = 0u;
v = 0u;
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 f6d86f4..30736b2 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
@@ -20,9 +20,6 @@
/* 0x000a */ tint_array<int8_t, 2> tint_pad;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 packed;
/* 0x000c */ tint_array<int8_t, 4> tint_pad_1;
@@ -100,7 +97,6 @@
uint v = 0u;
v = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
@@ -119,7 +115,6 @@
uint v_2 = 0u;
v_2 = 0u;
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 090228b..fae0314 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<float2x2, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<float2x2, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 c314155..b13326d 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* s;
@@ -36,7 +33,6 @@
uint v = 0u;
v = 0u;
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 fbddbd9..aac584d 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -24,9 +24,6 @@
threadgroup tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +59,6 @@
uint v_11 = 0u;
v_11 = tint_local_index;
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 9b5b263..6f28765 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* s;
@@ -36,7 +33,6 @@
uint v = 0u;
v = 0u;
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 5b4f038..7b8e6df 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
threadgroup tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +58,6 @@
uint v_11 = 0u;
v_11 = tint_local_index;
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 76f4709..6f8eb4d 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<half2x4, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<half2x4, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 c1edb07..9399b39 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<float2x4, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<float2x4, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 319cc4e..d4315e7 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* s;
@@ -37,7 +34,6 @@
uint v = 0u;
v = 0u;
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 7da7388..7501a2b 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
threadgroup tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +66,6 @@
uint v_15 = 0u;
v_15 = tint_local_index;
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 1bc2dfd..919fdfc 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<float3x4, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<float3x4, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 8bc1610..c98007e 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<half4x2, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<half4x2, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 aa1435a..85d79ac 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<float4x2, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<float4x2, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 20d20c2..9d02440 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* s;
@@ -38,7 +35,6 @@
uint v = 0u;
v = 0u;
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 f5f27df..388bef9 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
threadgroup tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +74,6 @@
uint v_19 = 0u;
v_19 = tint_local_index;
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 2eec94d..1ffcad0 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* s;
@@ -38,7 +35,6 @@
uint v = 0u;
v = 0u;
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 3fcde2f..5fd2133 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
threadgroup tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +74,6 @@
uint v_19 = 0u;
v_19 = tint_local_index;
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 b4adcfa..c0ba037 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<half4x4, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<half4x4, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 f00181f..109efff 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<float4x4, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<float4x4, 4> tint_symbol;
};
@@ -30,7 +27,6 @@
uint v = 0u;
v = tint_local_index;
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 0d61ebe..f7c76c5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -40,7 +37,6 @@
uint v = 0u;
v = 0u;
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 9e890cd..f203e72 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -26,9 +26,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -38,7 +35,6 @@
uint v = 0u;
v = tint_local_index;
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 d47aeb9..b182233 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 6e93929..034f6b8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 1e66f4e..2bda88b 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,9 +33,6 @@
int after;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<S_packed_vec3, 4>* s;
@@ -65,7 +62,6 @@
uint v_4 = 0u;
v_4 = 0u;
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 73f9ec1..35929fb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -38,9 +38,6 @@
threadgroup tint_array<S_packed_vec3, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +76,6 @@
uint v_7 = 0u;
v_7 = tint_local_index;
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 9af3b40..791ec19 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -33,9 +33,6 @@
int after;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<S_packed_vec3, 4>* s;
@@ -65,7 +62,6 @@
uint v_4 = 0u;
v_4 = 0u;
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 604a9fc..bc4c362 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -38,9 +38,6 @@
threadgroup tint_array<S_packed_vec3, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +76,6 @@
uint v_7 = 0u;
v_7 = tint_local_index;
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 6e0f92b..d605359 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 22c10fe..25f6f88 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 0dcf5fe..5738334 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 e3e8b78..32d5c0e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 3f065ae..371f30d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -40,7 +37,6 @@
uint v = 0u;
v = 0u;
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 3c163f0..06719be 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -26,9 +26,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -38,7 +35,6 @@
uint v = 0u;
v = tint_local_index;
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 261edbb..10dd8fe 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 7af8f12..ad7fb5e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 8087f74..73015ef 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,9 +33,6 @@
int after;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<S_packed_vec3, 4>* s;
@@ -67,7 +64,6 @@
uint v_5 = 0u;
v_5 = 0u;
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 31a73ff..5fff091 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -38,9 +38,6 @@
threadgroup tint_array<S_packed_vec3, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +78,6 @@
uint v_8 = 0u;
v_8 = tint_local_index;
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 d0c04d2..a492ec9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -32,9 +32,6 @@
int after;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<S_packed_vec3, 4>* s;
@@ -66,7 +63,6 @@
uint v_5 = 0u;
v_5 = 0u;
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 7f9fc85..d89f5c4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -37,9 +37,6 @@
threadgroup tint_array<S_packed_vec3, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +77,6 @@
uint v_8 = 0u;
v_8 = tint_local_index;
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 8c5a7ae..0b63816 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 176d1cf..df7d328 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 2bc64ef..c18e93c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -40,7 +37,6 @@
uint v = 0u;
v = 0u;
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 c3c152c..89f96d1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -26,9 +26,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -38,7 +35,6 @@
uint v = 0u;
v = tint_local_index;
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 54a2314..123b5cc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -40,7 +37,6 @@
uint v = 0u;
v = 0u;
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 cacf69bb..b0ffe26 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -26,9 +26,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -38,7 +35,6 @@
uint v = 0u;
v = tint_local_index;
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 1bd34e4..60d1bc8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 9d4b9ea..55c6b63c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 eb14cf6..4021b06 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,9 +33,6 @@
int after;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<S_packed_vec3, 4>* s;
@@ -69,7 +66,6 @@
uint v_6 = 0u;
v_6 = 0u;
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 502bdc2..7708382 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -38,9 +38,6 @@
threadgroup tint_array<S_packed_vec3, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +80,6 @@
uint v_9 = 0u;
v_9 = tint_local_index;
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 931f4ef..003d48f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -33,9 +33,6 @@
int after;
};
-#define TINT_ISOLATE_UB(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;
device tint_array<S_packed_vec3, 4>* s;
@@ -69,7 +66,6 @@
uint v_6 = 0u;
v_6 = 0u;
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 4a34690..e556e02 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -38,9 +38,6 @@
threadgroup tint_array<S_packed_vec3, 4>* w;
};
-#define TINT_ISOLATE_UB(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 +80,6 @@
uint v_9 = 0u;
v_9 = tint_local_index;
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 4126b7d..ee49f0d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 d52d4f1..c4331e9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 ae9d7f6..a66d52d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
/* 0x0084 */ tint_array<int8_t, 60> tint_pad_2;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant tint_array<S, 4>* u;
device tint_array<S, 4>* s;
@@ -41,7 +38,6 @@
uint v = 0u;
v = 0u;
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 32d3e05..12d3d65 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,9 +27,6 @@
threadgroup tint_array<S, 4>* w;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 4> tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 e8c2e82..411ced4 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void d() {
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
index b6f1da3..18b4edd 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
@@ -1,110 +1,163 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
int i0520 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i62 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_1);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i0520 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_2)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_2 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_2;
+ uint const tint_carry_2 = uint((tint_low_inc_2 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_2);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i62 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_3)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_3 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_3;
+ uint const tint_carry_3 = uint((tint_low_inc_3 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_3);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i62 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_4)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_4 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_4;
+ uint const tint_carry_4 = uint((tint_low_inc_4 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_4);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i60 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_5)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_5 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_5;
+ uint const tint_carry_5 = uint((tint_low_inc_5 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_5);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i62 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_6)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_6 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_6;
+ uint const tint_carry_6 = uint((tint_low_inc_6 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_6);
}
continue;
}
}
{
+ uint2 tint_loop_idx = 0u;
int i60 = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_7)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc_7 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_7;
+ uint const tint_carry_7 = uint((tint_low_inc_7 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_7);
}
continue;
}
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 82d2ab1..3546f28 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
@@ -31,9 +31,6 @@
threadgroup S* s;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
S tint_symbol;
};
@@ -43,7 +40,6 @@
uint v = 0u;
v = tint_local_index;
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 34f69d7..a517895 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
@@ -31,9 +31,6 @@
threadgroup S* s;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
S tint_symbol;
};
@@ -43,7 +40,6 @@
uint v = 0u;
v = tint_local_index;
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/gradient_in_varying_loop/1112.wgsl.expected.ir.msl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.ir.msl
index 0f31349d..30fe27d 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.ir.msl
@@ -7,9 +7,6 @@
texture2d<float, access::sample> depthTexture;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_outputs {
float4 tint_symbol_1 [[color(0)]];
};
@@ -22,8 +19,11 @@
float3 const random = tint_module_vars.randomTexture.sample(tint_module_vars.Sampler, vUV).xyz;
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 1)) {
} else {
break;
@@ -50,12 +50,20 @@
if (v_2) {
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
float const sampleDepth = 0.0f;
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/1064.wgsl.expected.ir.msl b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
index d7f982f..821663d 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
@@ -1,18 +1,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
fragment void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if (false) { break; }
}
continue;
diff --git a/test/tint/bug/tint/1081.wgsl.expected.ir.msl b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
index 1f866e8..6393e13 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
@@ -5,9 +5,6 @@
thread bool* continue_execution;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_outputs {
int tint_symbol_1 [[color(2)]];
};
@@ -26,13 +23,20 @@
int tint_symbol_inner(int3 x, tint_module_vars_struct tint_module_vars) {
int y = x.x;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
int const r = f(y, tint_module_vars);
if ((r == 0)) {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/1321.wgsl.expected.ir.msl b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
index 4694675..8e30a1f 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
@@ -13,9 +13,6 @@
T elements[N];
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int foo() {
return 1;
}
@@ -23,9 +20,12 @@
fragment void tint_symbol() {
tint_array<float, 4> arr = tint_array<float, 4>{};
{
+ uint2 tint_loop_idx = 0u;
thread float* const a = (&arr[min(uint(foo()), 3u)]);
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
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 4d10141..7f05e0e 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
@@ -1,13 +1,13 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (true) {
} else {
break;
@@ -18,6 +18,10 @@
return;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/1538.wgsl.expected.ir.msl b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
index b03a29d..4c800ae 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
@@ -1,9 +1,6 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -26,8 +23,11 @@
int f() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
g();
break;
}
@@ -39,14 +39,21 @@
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};
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.buf)[0u] == 0u)) {
break;
}
int s = f();
(*tint_module_vars.buf)[0u] = 0u;
{
+ uint const tint_low_inc_1 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_1);
}
continue;
}
diff --git a/test/tint/bug/tint/1557.wgsl.expected.ir.msl b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
index adc3222..1131e16 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
@@ -1,9 +1,6 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant int* u;
};
@@ -15,14 +12,21 @@
void g() {
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((j >= 1)) {
break;
}
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
int k = f();
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/1604.wgsl.expected.ir.msl b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
index afba81b..0537a76 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
@@ -5,17 +5,17 @@
const constant int* x;
};
-#define TINT_ISOLATE_UB(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};
switch((*tint_module_vars.x)) {
case 0:
{
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
return;
}
}
diff --git a/test/tint/bug/tint/1605.wgsl.expected.ir.msl b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
index 15bb159..8e463a4 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
@@ -5,22 +5,25 @@
const constant int* b;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
bool func_3(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < (*tint_module_vars.b))) {
} else {
break;
}
{
+ uint2 tint_loop_idx_1 = 0u;
int j = -1;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j == 1)) {
} else {
break;
@@ -29,6 +32,10 @@
}
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
diff --git a/test/tint/bug/tint/1764.wgsl.expected.ir.msl b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
index d5c3440..a925757 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<int, 246>* W;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<int, 246> tint_symbol_1;
};
@@ -29,7 +26,6 @@
uint v = 0u;
v = tint_local_index;
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 2d78c67..c024053 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.ir.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.ir.msl
@@ -47,9 +47,6 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_6 {
tint_array<S, 4096> tint_symbol_1;
atomic_uint tint_symbol_2;
@@ -66,8 +63,11 @@
uint const x_52 = (*tint_module_vars.x_3).x;
x_54 = 0u;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
uint x_55 = 0u;
x_58 = (*tint_module_vars.x_6).field0.field0;
if ((x_54 < x_58)) {
@@ -80,6 +80,10 @@
(*tint_module_vars.x_28)[min(x_62, 4095u)] = S{.field0=((x_67.xy + x_67.zw) * 0.5f), .field1=x_62};
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
x_55 = (x_54 + 32u);
x_54 = x_55;
}
@@ -101,8 +105,11 @@
x_85 = x_76.xyxy;
x_88 = 1u;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
float4 x_111 = 0.0f;
float4 x_86 = 0.0f;
uint x_89 = 0u;
@@ -130,6 +137,10 @@
x_86 = x_111;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry_1);
x_89 = (x_88 + 32u);
x_85 = x_86;
x_88 = x_89;
@@ -159,7 +170,6 @@
uint v_2 = 0u;
v_2 = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_2)
uint const v_3 = v_2;
if ((v_3 >= 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 11378a4..a31577c 100644
--- a/test/tint/bug/tint/2039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
uint out = 0u;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool tint_continue = false;
switch(2) {
case 1:
@@ -23,12 +23,20 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if (true) { break; }
}
continue;
}
out = (out + 1u);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if (true) { break; }
}
continue;
diff --git a/test/tint/bug/tint/2059.wgsl.expected.ir.msl b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
index 60a731d..c2b1451 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
};
-#define TINT_ISOLATE_UB(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 +72,6 @@
uint v = 0u;
v = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 1u)) {
break;
@@ -98,7 +94,6 @@
uint v_2 = 0u;
v_2 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
uint const v_3 = v_2;
if ((v_3 >= 1u)) {
break;
@@ -121,7 +116,6 @@
uint v_4 = 0u;
v_4 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_2)
uint const v_5 = v_4;
if ((v_5 >= 1u)) {
break;
diff --git a/test/tint/bug/tint/2201.wgsl.expected.ir.msl b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
index 2cc5558..f9f6e8e 100644
--- a/test/tint/bug/tint/2201.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
@@ -5,13 +5,13 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
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 b7f4c947..c1348e4 100644
--- a/test/tint/bug/tint/2202.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2202.wgsl.expected.ir.msl
@@ -5,16 +5,19 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
return;
}
}
diff --git a/test/tint/bug/tint/221.wgsl.expected.ir.msl b/test/tint/bug/tint/221.wgsl.expected.ir.msl
index 0492a2d..5f8e08f 100644
--- a/test/tint/bug/tint/221.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.ir.msl
@@ -22,9 +22,6 @@
device Buf* b;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
uint tint_mod_u32(uint lhs, uint rhs) {
return (lhs - ((lhs / select(rhs, 1u, (rhs == 0u))) * select(rhs, 1u, (rhs == 0u))));
}
@@ -33,14 +30,21 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.b=b};
uint i = 0u;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i >= (*tint_module_vars.b).count)) {
break;
}
device uint* const p = (&(*tint_module_vars.b).data[min(i, 49u)]);
if ((tint_mod_u32(i, 2u) == 0u)) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
(*p) = ((*p) * 2u);
i = (i + 1u);
}
@@ -48,6 +52,10 @@
}
(*p) = 0u;
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
(*p) = ((*p) * 2u);
i = (i + 1u);
}
diff --git a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
index 57661c0..25b8bde 100644
--- a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
@@ -5,20 +5,24 @@
texture2d<float, access::sample> tint_symbol;
};
-#define TINT_ISOLATE_UB(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};
{
+ uint2 tint_loop_idx = 0u;
uint level = tint_module_vars.tint_symbol.get_num_mip_levels();
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((level > 0u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
index fb5f664..428f1f4 100644
--- a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
@@ -5,19 +5,26 @@
device int* tint_symbol;
};
-#define TINT_ISOLATE_UB(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);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
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 85977d4..5830e23 100644
--- a/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
@@ -29,9 +29,6 @@
tint_array<uint3, 4> c;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_module_vars_struct {
const constant S_packed_vec3* ubuffer;
device S_packed_vec3* sbuffer;
@@ -56,7 +53,6 @@
uint v = 0u;
v = 0u;
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/379127084.wgsl.expected.ir.msl b/test/tint/bug/tint/379127084.wgsl.expected.ir.msl
index 021e2ef..5a08097 100644
--- a/test/tint/bug/tint/379127084.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/379127084.wgsl.expected.ir.msl
@@ -56,9 +56,6 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_outputs {
float4 FSOut_sk_FragColor [[color(0)]];
};
@@ -77,8 +74,11 @@
float _60_n = 1.0f;
int _61_o = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((_61_o < (*tint_module_vars._storage1).fsUniformData[min((*tint_module_vars.shadingSsboIndex), ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 128u) - 1u))].numOctaves_1)) {
float4 _62_f = 0.0f;
float2 const _skTemp2 = floor(_57_k);
@@ -107,8 +107,11 @@
float4 _71_g = 0.0f;
int _72_h = 0;
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
float const _73_i = ((float(_72_h) + 0.5f) * 0.25f);
float const v_2 = float(_67_p.x);
float2 const v_3 = float2(v_2, float(_73_i));
@@ -141,6 +144,10 @@
float const _skTemp13 = mix(_81_q, _82_r, _69_e.y);
_71_g[min(uint(_72_h), 3u)] = _skTemp13;
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
_72_h = as_type<int>((as_type<uint>(_72_h) + as_type<uint>(1)));
if ((_72_h >= 4)) { break; }
}
@@ -160,6 +167,10 @@
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
_61_o = as_type<int>((as_type<uint>(_61_o) + as_type<uint>(1)));
}
continue;
diff --git a/test/tint/bug/tint/379684039-2.wgsl.expected.ir.msl b/test/tint/bug/tint/379684039-2.wgsl.expected.ir.msl
index d0850f4..74a79a7 100644
--- a/test/tint/bug/tint/379684039-2.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/379684039-2.wgsl.expected.ir.msl
@@ -28,18 +28,22 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void tint_symbol(tint_module_vars_struct tint_module_vars) {
int2 tint_symbol_1 = int2(0);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((tint_symbol_1.y >= (*tint_module_vars._storage).fsUniformData[min((*tint_module_vars.idx), ((((*tint_module_vars.tint_storage_buffer_sizes)[0u].x - 0u) / 128u) - 1u))].size.y)) {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/379684039.wgsl.expected.ir.msl b/test/tint/bug/tint/379684039.wgsl.expected.ir.msl
index f4a0b35..a245b5c 100644
--- a/test/tint/bug/tint/379684039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/379684039.wgsl.expected.ir.msl
@@ -5,14 +5,14 @@
const device int2* _storage;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void tint_symbol(tint_module_vars_struct tint_module_vars) {
int2 tint_symbol_1 = int2(0);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((tint_symbol_1.y >= (*tint_module_vars._storage).y)) {
break;
}
@@ -20,6 +20,10 @@
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/bug/tint/534.wgsl.expected.ir.msl b/test/tint/bug/tint/534.wgsl.expected.ir.msl
index ff82816..865ddbf 100644
--- a/test/tint/bug/tint/534.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.ir.msl
@@ -32,9 +32,6 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
uint ConvertToFp16FloatValue(float fp32) {
return 1u;
}
@@ -60,9 +57,12 @@
uint4 srcColorBits = 0u;
uint4 dstColorBits = tint_v4f32_to_v4u32(dstColor);
{
+ uint2 tint_loop_idx = 0u;
uint i = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < (*tint_module_vars.uniforms).channelCount)) {
} else {
break;
@@ -77,6 +77,10 @@
}
success = v_5;
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = (i + 1u);
}
continue;
diff --git a/test/tint/bug/tint/914.wgsl.expected.ir.msl b/test/tint/bug/tint/914.wgsl.expected.ir.msl
index 28c473f7..a0fcba9 100644
--- a/test/tint/bug/tint/914.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.ir.msl
@@ -33,9 +33,6 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(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;
tint_array<tint_array<float, 64>, 64> tint_symbol_2;
@@ -91,7 +88,6 @@
uint v_3 = 0u;
v_3 = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_4 = v_3;
if ((v_4 >= 4096u)) {
break;
diff --git a/test/tint/bug/tint/942.wgsl.expected.ir.msl b/test/tint/bug/tint/942.wgsl.expected.ir.msl
index 30a9f19..20cd744 100644
--- a/test/tint/bug/tint/942.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.ir.msl
@@ -35,9 +35,6 @@
threadgroup tint_array<tint_array<tint_packed_vec3_f32_array_element, 256>, 4>* tile;
};
-#define TINT_ISOLATE_UB(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 +48,6 @@
uint v = 0u;
v = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 1024u)) {
break;
@@ -138,9 +134,12 @@
if (v_8) {
float3 acc = float3(0.0f);
{
+ uint2 tint_loop_idx = 0u;
uint f = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((f < (*tint_module_vars.params).filterDim)) {
} else {
break;
@@ -150,6 +149,10 @@
float const v_10 = (1.0f / float((*tint_module_vars.params).filterDim));
acc = (v_9 + (v_10 * float3((*tint_module_vars.tile)[min(r, 3u)][min(i, 255u)].packed)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
f = (f + 1u);
}
continue;
diff --git a/test/tint/bug/tint/948.wgsl.expected.ir.msl b/test/tint/bug/tint/948.wgsl.expected.ir.msl
index 8ac8d4d..cd13982 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl
@@ -48,9 +48,6 @@
thread float2* vUV;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct main_out {
float4 glFragColor_1;
};
@@ -121,8 +118,11 @@
stageUnits = (float2(1.0f) / x_111);
i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
int const x_122 = i;
if ((x_122 < 2)) {
} else {
@@ -163,8 +163,11 @@
(*tint_module_vars.mt) = fmod((x_181 * x_184), 1.0f);
f = 0.0f;
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
float const x_193 = f;
if ((x_193 < 8.0f)) {
} else {
@@ -183,6 +186,10 @@
float4 const x_217 = float4(0.0f);
animationData = x_217;
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
float const x_218 = f;
f = (x_218 + 1.0f);
}
@@ -236,6 +243,10 @@
color = float4(x_298.x, x_298.y, x_298.z, x_299);
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
int const x_304 = i;
i = as_type<int>((as_type<uint>(x_304) + as_type<uint>(1)));
}
diff --git a/test/tint/bug/tint/949.wgsl.expected.ir.msl b/test/tint/bug/tint/949.wgsl.expected.ir.msl
index 66c7372..7faf681 100644
--- a/test/tint/bug/tint/949.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.ir.msl
@@ -61,9 +61,6 @@
texture2d<float, access::sample> bumpSamplerTexture;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct main_out {
float4 glFragColor_1;
};
@@ -338,8 +335,11 @@
currSampledHeight = 1.0f;
i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
int const x_388 = i;
if ((x_388 < 15)) {
} else {
@@ -383,6 +383,10 @@
lastSampledHeight = x_440;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
int const x_441 = i;
i = as_type<int>((as_type<uint>(x_441) + as_type<uint>(1)));
}
diff --git a/test/tint/bug/tint/990.wgsl.expected.ir.msl b/test/tint/bug/tint/990.wgsl.expected.ir.msl
index b225169..b9d4695 100644
--- a/test/tint/bug/tint/990.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.ir.msl
@@ -1,20 +1,24 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
thread int* const p = (&i);
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 f5103c1..b644201 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
-#define TINT_ISOLATE_UB(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;
};
@@ -37,8 +34,11 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (!((idx < 6u))) {
break;
}
@@ -49,6 +49,10 @@
uint const v_1 = tint_mod_u32(x_33, 2u);
atomic_store_explicit((&(*tint_module_vars.wg)[min(v, 2u)][min(v_1, 1u)][min(tint_mod_u32(x_35, 1u), 0u)]), 0u, memory_order_relaxed);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
idx = (idx + 1u);
}
continue;
@@ -68,7 +72,6 @@
uint v_2 = 0u;
v_2 = local_invocation_index_1_param;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
uint const v_3 = v_2;
if ((v_3 >= 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 f1273b1..76de2ee 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
-#define TINT_ISOLATE_UB(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 +26,6 @@
uint v = 0u;
v = tint_local_index;
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 7466113..e37bc58 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<atomic_uint, 4>* wg;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<atomic_uint, 4> tint_symbol;
};
@@ -29,14 +26,21 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (!((idx < 4u))) {
break;
}
uint const x_26 = idx;
atomic_store_explicit((&(*tint_module_vars.wg)[min(x_26, 3u)]), 0u, memory_order_relaxed);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
idx = (idx + 1u);
}
continue;
@@ -56,7 +60,6 @@
uint v = 0u;
v = local_invocation_index_1_param;
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 b35d5d0..3716ffc 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<atomic_uint, 4>* wg;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<atomic_uint, 4> tint_symbol;
};
@@ -29,7 +26,6 @@
uint v = 0u;
v = tint_local_index;
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 f5103c1..b644201 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
@@ -18,9 +18,6 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
-#define TINT_ISOLATE_UB(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;
};
@@ -37,8 +34,11 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (!((idx < 6u))) {
break;
}
@@ -49,6 +49,10 @@
uint const v_1 = tint_mod_u32(x_33, 2u);
atomic_store_explicit((&(*tint_module_vars.wg)[min(v, 2u)][min(v_1, 1u)][min(tint_mod_u32(x_35, 1u), 0u)]), 0u, memory_order_relaxed);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
idx = (idx + 1u);
}
continue;
@@ -68,7 +72,6 @@
uint v_2 = 0u;
v_2 = local_invocation_index_1_param;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
uint const v_3 = v_2;
if ((v_3 >= 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 f1273b1..76de2ee 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* wg;
};
-#define TINT_ISOLATE_UB(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 +26,6 @@
uint v = 0u;
v = tint_local_index;
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 4d00e1a..6ef4e77 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
@@ -24,9 +24,6 @@
threadgroup tint_array<S_atomic, 10>* wg;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S_atomic, 10> tint_symbol;
};
@@ -35,8 +32,11 @@
uint idx = 0u;
idx = local_invocation_index_2;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (!((idx < 10u))) {
break;
}
@@ -45,6 +45,10 @@
atomic_store_explicit((&(*tint_module_vars.wg)[min(x_28, 9u)].a), 0u, memory_order_relaxed);
(*tint_module_vars.wg)[min(x_28, 9u)].y = 0u;
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
idx = (idx + 1u);
}
continue;
@@ -64,7 +68,6 @@
uint v = 0u;
v = local_invocation_index_1_param;
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 a20354c..7b57891 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
threadgroup tint_array<S, 10>* wg;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
tint_array<S, 10> tint_symbol;
};
@@ -35,7 +32,6 @@
uint v = 0u;
v = tint_local_index;
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 14b0564..dd110ac 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
@@ -24,9 +24,6 @@
threadgroup S_atomic* wg;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
S_atomic tint_symbol;
};
@@ -37,14 +34,21 @@
(*tint_module_vars.wg).y = 0u;
idx = local_invocation_index_2;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (!((idx < 10u))) {
break;
}
uint const x_35 = idx;
atomic_store_explicit((&(*tint_module_vars.wg).a[min(x_35, 9u)]), 0u, memory_order_relaxed);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
idx = (idx + 1u);
}
continue;
@@ -68,7 +72,6 @@
uint v = 0u;
v = local_invocation_index_1_param;
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 b9386d8..3488c05 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
threadgroup S* wg;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_1 {
S tint_symbol;
};
@@ -39,7 +36,6 @@
uint v = 0u;
v = tint_local_index;
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 c0a5876..3a773b6 100644
--- a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
+++ b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
@@ -5,19 +5,23 @@
texture2d<int, access::read_write> tex;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void foo(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 3)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
tint_module_vars.tex.write(int4(0), uint2(int2(0)));
const_cast<texture2d<int, access::read_write>thread &>(tint_module_vars.tex).fence();
}
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 df06f9b..4ec26bd 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
@@ -6,14 +6,14 @@
threadgroup int* b;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void foo(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
int const v = i;
threadgroup_barrier(mem_flags::mem_threadgroup);
int const v_1 = (*tint_module_vars.a);
@@ -23,6 +23,10 @@
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
threadgroup_barrier(mem_flags::mem_threadgroup);
int const v_2 = (*tint_module_vars.b);
threadgroup_barrier(mem_flags::mem_threadgroup);
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 7161fb2..52f75f9 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
@@ -13,9 +13,6 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -23,8 +20,11 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool v_1 = false;
if ((x > v.x)) {
v_1 = (dfdx(1.0f) > 0.0f);
@@ -36,6 +36,10 @@
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 7108bed..a18913c 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
sampler s;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -28,14 +25,21 @@
void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
float4 v = float4(0.0f);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((x > v.x)) {
} else {
break;
}
v = tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 de7b62f..9690ca1 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
@@ -13,19 +13,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
dfdx(1.0f);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if ((x > 0.0f)) { break; }
}
continue;
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 18c03d7..a9a2a62 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
@@ -13,19 +13,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
dfdx(1.0f);
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
if ((x > 0.0f)) { break; }
}
continue;
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 efa9d91..f9a4f82 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
@@ -13,18 +13,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
void tint_symbol_inner(float x) {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
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 4c7af50..0fd40ae 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
@@ -13,9 +13,6 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -23,8 +20,11 @@
void tint_symbol_inner(float x) {
float4 v = float4(0.0f);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool v_1 = false;
if ((x > 0.0f)) {
v_1 = (dfdx(1.0f) > 0.0f);
@@ -36,6 +36,10 @@
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 41b1284..c907128 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
@@ -18,9 +18,6 @@
sampler s;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_inputs {
float x [[user(locn0)]];
};
@@ -28,14 +25,21 @@
void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
float4 v = float4(0.0f);
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((x > v.x)) {
} else {
break;
}
v = tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 6a9c791..e23ac2f 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
@@ -18,9 +18,6 @@
/* 0x0008 */ tint_array<int8_t, 8> tint_pad;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct SSBO {
/* 0x0000 */ tint_array<strided_arr, 2> m;
};
@@ -47,7 +44,6 @@
uint v_1 = 0u;
v_1 = 0u;
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_with_breakif.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
index d9fbf58..2b08ee3 100644
--- a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool tint_continue = false;
switch(i) {
case 0:
@@ -23,12 +23,20 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i >= 4)) { break; }
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i >= 4)) { break; }
}
diff --git a/test/tint/loops/continue_in_switch_with_breakif_robustness.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch_with_breakif_robustness.wgsl.expected.ir.msl
index d9fbf58..2b08ee3 100644
--- a/test/tint/loops/continue_in_switch_with_breakif_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch_with_breakif_robustness.wgsl.expected.ir.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool tint_continue = false;
switch(i) {
case 0:
@@ -23,12 +23,20 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i >= 4)) { break; }
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i >= 4)) { break; }
}
diff --git a/test/tint/loops/loop.wgsl.expected.ir.msl b/test/tint/loops/loop.wgsl.expected.ir.msl
index d5e5fa2..8a55eb3 100644
--- a/test/tint/loops/loop.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return i;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/loops/loop_robustness.wgsl.expected.ir.msl b/test/tint/loops/loop_robustness.wgsl.expected.ir.msl
index d5e5fa2..8a55eb3 100644
--- a/test/tint/loops/loop_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_robustness.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return i;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 b32bbacb..f4578ae 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
@@ -1,18 +1,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i > 4)) {
return i;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i == 4)) { break; }
}
diff --git a/test/tint/loops/loop_with_break_if_robustness.wgsl.expected.ir.msl b/test/tint/loops/loop_with_break_if_robustness.wgsl.expected.ir.msl
index b32bbacb..f4578ae 100644
--- a/test/tint/loops/loop_with_break_if_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_break_if_robustness.wgsl.expected.ir.msl
@@ -1,18 +1,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i > 4)) {
return i;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i == 4)) { break; }
}
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 e21dbe4..9e61e13 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
@@ -1,18 +1,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i > 4)) {
return i;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
diff --git a/test/tint/loops/loop_with_continuing_robustness.wgsl.expected.ir.msl b/test/tint/loops/loop_with_continuing_robustness.wgsl.expected.ir.msl
index e21dbe4..9e61e13 100644
--- a/test/tint/loops/loop_with_continuing_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing_robustness.wgsl.expected.ir.msl
@@ -1,18 +1,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i > 4)) {
return i;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
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 3090c93..e2e9bcd 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
@@ -1,22 +1,25 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
}
{
+ uint2 tint_loop_idx_1 = 0u;
int j = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j < 2)) {
} else {
break;
@@ -35,17 +38,29 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
diff --git a/test/tint/loops/nested_loop_loop_switch_robustness.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_loop_switch_robustness.wgsl.expected.ir.msl
index 3090c93..e2e9bcd 100644
--- a/test/tint/loops/nested_loop_loop_switch_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_loop_switch_robustness.wgsl.expected.ir.msl
@@ -1,22 +1,25 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
}
{
+ uint2 tint_loop_idx_1 = 0u;
int j = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j < 2)) {
} else {
break;
@@ -35,17 +38,29 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
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 caf7752..4d89c60 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
@@ -18,9 +18,12 @@
case 0:
{
{
+ uint2 tint_loop_idx_1 = 0u;
int j = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j < 2)) {
} else {
break;
@@ -39,11 +42,19 @@
}
if (tint_continue_1) {
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
@@ -59,11 +70,19 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_robustness.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch_robustness.wgsl.expected.ir.msl
index caf7752..4d89c60 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_robustness.wgsl.expected.ir.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
@@ -18,9 +18,12 @@
case 0:
{
{
+ uint2 tint_loop_idx_1 = 0u;
int j = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j < 2)) {
} else {
break;
@@ -39,11 +42,19 @@
}
if (tint_continue_1) {
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
@@ -59,11 +70,19 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
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 0bb8ce8..f91ee20 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
@@ -1,15 +1,15 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
int k = 0;
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
@@ -19,9 +19,12 @@
case 0:
{
{
+ uint2 tint_loop_idx_1 = 0u;
int j = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j < 2)) {
} else {
break;
@@ -60,11 +63,19 @@
}
if (tint_continue_1) {
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
@@ -80,11 +91,19 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch_robustness.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch_robustness.wgsl.expected.ir.msl
index 0bb8ce8..f91ee20 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch_robustness.wgsl.expected.ir.msl
@@ -1,15 +1,15 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
int k = 0;
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
@@ -19,9 +19,12 @@
case 0:
{
{
+ uint2 tint_loop_idx_1 = 0u;
int j = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j < 2)) {
} else {
break;
@@ -60,11 +63,19 @@
}
if (tint_continue_1) {
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
}
continue;
@@ -80,11 +91,19 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
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 208823c..59cb654 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
@@ -1,15 +1,15 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
@@ -43,11 +43,19 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
diff --git a/test/tint/loops/nested_loop_switch_switch_robustness.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_switch_robustness.wgsl.expected.ir.msl
index 208823c..59cb654 100644
--- a/test/tint/loops/nested_loop_switch_switch_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_switch_robustness.wgsl.expected.ir.msl
@@ -1,15 +1,15 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
kernel void tint_symbol() {
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 2)) {
} else {
break;
@@ -43,11 +43,19 @@
}
if (tint_continue) {
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
}
continue;
diff --git a/test/tint/loops/nested_loops.wgsl.expected.ir.msl b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
index daa4a15..5041162 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
@@ -1,27 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return 1;
}
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
if ((j > 4)) {
return 2;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
}
continue;
}
diff --git a/test/tint/loops/nested_loops_robustness.wgsl.expected.ir.msl b/test/tint/loops/nested_loops_robustness.wgsl.expected.ir.msl
index daa4a15..5041162 100644
--- a/test/tint/loops/nested_loops_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_robustness.wgsl.expected.ir.msl
@@ -1,27 +1,34 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i > 4)) {
return 1;
}
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
if ((j > 4)) {
return 2;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
}
continue;
}
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 8d2609f..2b4726c 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
@@ -1,25 +1,32 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i > 4)) {
return 1;
}
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j > 4)) {
return 2;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
}
continue;
diff --git a/test/tint/loops/nested_loops_with_continuing_robustness.wgsl.expected.ir.msl b/test/tint/loops/nested_loops_with_continuing_robustness.wgsl.expected.ir.msl
index 8d2609f..2b4726c 100644
--- a/test/tint/loops/nested_loops_with_continuing_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_with_continuing_robustness.wgsl.expected.ir.msl
@@ -1,25 +1,32 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i > 4)) {
return 1;
}
{
+ uint2 tint_loop_idx_1 = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false_1)
+ if (all((tint_loop_idx_1 == uint2(4294967295u)))) {
+ break;
+ }
if ((j > 4)) {
return 2;
}
{
+ uint const tint_low_inc_1 = (tint_loop_idx_1.x + 1u);
+ tint_loop_idx_1.x = tint_low_inc_1;
+ uint const tint_carry_1 = uint((tint_low_inc_1 == 0u));
+ tint_loop_idx_1.y = (tint_loop_idx_1.y + tint_carry_1);
j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
}
continue;
diff --git a/test/tint/loops/while.wgsl.expected.ir.msl b/test/tint/loops/while.wgsl.expected.ir.msl
index 62623af..7b7e141 100644
--- a/test/tint/loops/while.wgsl.expected.ir.msl
+++ b/test/tint/loops/while.wgsl.expected.ir.msl
@@ -1,20 +1,24 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/loops/while_robustness.wgsl.expected.ir.msl b/test/tint/loops/while_robustness.wgsl.expected.ir.msl
index 62623af..7b7e141 100644
--- a/test/tint/loops/while_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_robustness.wgsl.expected.ir.msl
@@ -1,20 +1,24 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 62623af..7b7e141 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
@@ -1,20 +1,24 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/loops/while_with_continue_robustness.wgsl.expected.ir.msl b/test/tint/loops/while_with_continue_robustness.wgsl.expected.ir.msl
index 62623af..7b7e141 100644
--- a/test/tint/loops/while_with_continue_robustness.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue_robustness.wgsl.expected.ir.msl
@@ -1,20 +1,24 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
int f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 fbcdcaa..0301bb0 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
threadgroup str* S;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
str tint_symbol_1;
};
@@ -37,7 +34,6 @@
uint v = 0u;
v = tint_local_index;
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 f2fb702..30eca59 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
threadgroup tint_array<str, 4>* S;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<str, 4> tint_symbol_1;
};
@@ -37,7 +34,6 @@
uint v = 0u;
v = tint_local_index;
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 e3a3a94..2b87be7 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
threadgroup str* S;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
str tint_symbol_1;
};
@@ -37,7 +34,6 @@
uint v = 0u;
v = tint_local_index;
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 7c64eae..cac795e 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -21,9 +21,6 @@
threadgroup tint_array<str, 4>* S;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<str, 4> tint_symbol_1;
};
@@ -37,7 +34,6 @@
uint v = 0u;
v = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_1 = v;
if ((v_1 >= 4u)) {
break;
diff --git a/test/tint/shadowing/loop.wgsl.expected.ir.msl b/test/tint/shadowing/loop.wgsl.expected.ir.msl
index 73fe235..a54e1dd 100644
--- a/test/tint/shadowing/loop.wgsl.expected.ir.msl
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.msl
@@ -17,17 +17,21 @@
device tint_array<int, 10>* output;
};
-#define TINT_ISOLATE_UB(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;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
int x = (*tint_module_vars.output)[min(uint(i), 9u)];
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
int x_1 = (*tint_module_vars.output)[min(uint(x), 9u)];
i = as_type<int>((as_type<uint>(i) + as_type<uint>(x_1)));
if ((i > 10)) { 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 72037b9..748b54b 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
@@ -29,23 +29,27 @@
tint_array<InnerS, 8> a1;
};
-#define TINT_ISOLATE_UB(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};
InnerS v = {};
OuterS s1 = {};
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
s1.a1[min((*tint_module_vars.uniforms).i, 7u)] = v;
}
continue;
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 b6b66ea..a60c76c 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
@@ -29,23 +29,27 @@
tint_array<InnerS, 8> a1;
};
-#define TINT_ISOLATE_UB(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};
InnerS v = {};
OuterS s1 = {};
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
s1.a1[min((*tint_module_vars.uniforms).i, 7u)] = v;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
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 7a589be..54e2d44 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
@@ -24,9 +24,6 @@
T elements[N];
};
-#define TINT_ISOLATE_UB(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);
return 1;
@@ -45,15 +42,22 @@
void foo(tint_module_vars_struct tint_module_vars) {
tint_array<float, 4> a = tint_array<float, 4>{};
{
+ uint2 tint_loop_idx = 0u;
thread float* const v_1 = (&a[min(uint(idx1(tint_module_vars)), 3u)]);
(*v_1) = ((*v_1) * 2.0f);
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((a[min(uint(idx2(tint_module_vars)), 3u)] < 10.0f)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
thread float* const v_2 = (&a[min(uint(idx3(tint_module_vars)), 3u)]);
(*v_2) = ((*v_2) + 1.0f);
}
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
index 12e3e83..f2fcd00 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(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);
return 1;
@@ -58,6 +55,7 @@
void tint_symbol_1(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
int const v_1 = idx1(tint_module_vars);
int const v_2 = idx2(tint_module_vars);
uint const v_3 = (((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 64u) - 1u);
@@ -67,12 +65,18 @@
int const v_7 = as_type<int>((as_type<uint>((*v_5)[min(uint(v_6), 3u)]) - as_type<uint>(1)));
(*v_5)[min(uint(v_6), 3u)] = v_7;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.v) < 10u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
int const v_8 = idx4(tint_module_vars);
int const v_9 = idx5(tint_module_vars);
uint const v_10 = (((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 64u) - 1u);
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 df04b49..f2d8efa 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
@@ -5,18 +5,22 @@
device uint* i;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
(*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
}
continue;
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 6bb40ad..ce23f4d 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
@@ -5,19 +5,23 @@
device uint* i;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
(*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 76a56a7..5650e4b 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
@@ -8,9 +8,6 @@
thread bool* continue_execution;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct foo_outputs {
int tint_symbol [[color(0)]];
};
@@ -30,15 +27,22 @@
}
int result = tint_f32_to_i32(tint_module_vars.t.sample(tint_module_vars.s, coord).x);
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 10)) {
} else {
break;
}
result = as_type<int>((as_type<uint>(result) + as_type<uint>(i)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
int v = 0;
if ((*tint_module_vars.continue_execution)) {
v = atomic_fetch_add_explicit(tint_module_vars.a, 1, memory_order_relaxed);
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 752536e..c4a833a 100644
--- a/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
@@ -5,13 +5,13 @@
thread bool* continue_execution;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
(*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 1c9efef..01e3d1f 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
@@ -7,9 +7,6 @@
thread bool* continue_execution;
};
-#define TINT_ISOLATE_UB(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;
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.non_uniform_global=non_uniform_global, .output=output, .continue_execution=(&continue_execution)};
@@ -23,8 +20,11 @@
if (((*tint_module_vars.output) < 0.0f)) {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
float const v_1 = (*tint_module_vars.output);
if ((v_1 > float(i))) {
float const v_2 = float(i);
@@ -37,6 +37,10 @@
return;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
if ((i == 5)) { break; }
}
diff --git a/test/tint/statements/for/complex.wgsl.expected.ir.msl b/test/tint/statements/for/complex.wgsl.expected.ir.msl
index 4dd60ee..86ccef9 100644
--- a/test/tint/statements/for/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.ir.msl
@@ -1,18 +1,18 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void some_loop_body() {
}
void f() {
int j = 0;
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
bool v = false;
if ((i < 5)) {
v = (j < 10);
@@ -26,6 +26,10 @@
some_loop_body();
j = as_type<int>((as_type<uint>(i) * as_type<uint>(30)));
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
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 02a71f0..9d31c4e 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 1)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 fae57d4..d820734 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 4)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 02a71f0..9d31c4e 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if ((i < 1)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 76a659f..c8c7d5f 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
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 76a659f..c8c7d5f 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
int i = 0;
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
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 ae63ee6..4d25568 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
}
continue;
diff --git a/test/tint/statements/for/empty.wgsl.expected.ir.msl b/test/tint/statements/for/empty.wgsl.expected.ir.msl
index 7b3f2ce..fef5577 100644
--- a/test/tint/statements/for/empty.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.ir.msl
@@ -1,18 +1,22 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 a1a7b9a..a7ff235 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
int i = 1;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 c3e6673..d27c489 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
int i = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 a1a7b9a..a7ff235 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
@@ -1,19 +1,23 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
int i = 1;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (false) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
diff --git a/test/tint/statements/for/scoping.wgsl.expected.ir.msl b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
index 5b08034..ef61178 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void f() {
{
+ uint2 tint_loop_idx = 0u;
int must_not_collide = 0;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
break;
}
}
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.msl b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
index ef40ff4..9888f9e 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
@@ -23,9 +23,6 @@
const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
};
-#define TINT_ISOLATE_UB(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);
return 1;
@@ -58,6 +55,7 @@
void tint_symbol_1(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
int const v_1 = idx1(tint_module_vars);
int const v_2 = idx2(tint_module_vars);
uint const v_3 = (((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 64u) - 1u);
@@ -67,12 +65,18 @@
int const v_7 = as_type<int>((as_type<uint>((*v_5)[min(uint(v_6), 3u)]) + as_type<uint>(1)));
(*v_5)[min(uint(v_6), 3u)] = v_7;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.v) < 10u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
int const v_8 = idx4(tint_module_vars);
int const v_9 = idx5(tint_module_vars);
uint const v_10 = (((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 64u) - 1u);
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 1aa5603..a5d709e 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
@@ -5,18 +5,22 @@
device uint* i;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
(*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
}
continue;
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 550762b..2c57299 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
@@ -5,19 +5,23 @@
device uint* i;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
void tint_symbol(tint_module_vars_struct tint_module_vars) {
{
+ uint2 tint_loop_idx = 0u;
(*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
+ if (all((tint_loop_idx == uint2(4294967295u)))) {
+ break;
+ }
if (((*tint_module_vars.i) < 10u)) {
} else {
break;
}
{
+ uint const tint_low_inc = (tint_loop_idx.x + 1u);
+ tint_loop_idx.x = tint_low_inc;
+ uint const tint_carry = uint((tint_low_inc == 0u));
+ tint_loop_idx.y = (tint_loop_idx.y + tint_carry);
}
continue;
}
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 4d50db2..02aad41 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<tint_array<int, 3>, 2>* zero;
};
-#define TINT_ISOLATE_UB(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 +26,6 @@
uint v_1 = 0u;
v_1 = tint_local_index;
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 1fc2b21..ed037d3 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<int, 3>* zero;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<int, 3> tint_symbol_1;
};
@@ -29,7 +26,6 @@
uint v_1 = 0u;
v_1 = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_2 = v_1;
if ((v_2 >= 3u)) {
break;
diff --git a/test/tint/var/initialization/workgroup/array/u32_large.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/u32_large.wgsl.expected.ir.msl
index d49d33e..346fecc 100644
--- a/test/tint/var/initialization/workgroup/array/u32_large.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/u32_large.wgsl.expected.ir.msl
@@ -17,9 +17,6 @@
threadgroup tint_array<int, 23>* zero;
};
-#define TINT_ISOLATE_UB(VOLATILE_NAME) \
- {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
-
struct tint_symbol_2 {
tint_array<int, 23> tint_symbol_1;
};
@@ -29,7 +26,6 @@
uint v_1 = 0u;
v_1 = tint_local_index;
while(true) {
- TINT_ISOLATE_UB(tint_volatile_false)
uint const v_2 = v_1;
if ((v_2 >= 23u)) {
break;
