reader/wgsl: Generate ForLoopStatements
Instead of LoopStatements.
Update the writers to handle these.
Fixed: tint:952
Change-Id: Ibef66e133224810efc28c224d910b5e21f71f8d6
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57203
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index d80e8c7..a3b9b23 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -2071,7 +2071,7 @@
// for_statement
// : FOR PAREN_LEFT for_header PAREN_RIGHT BRACE_LEFT statements BRACE_RIGHT
-Maybe<ast::Statement*> ParserImpl::for_stmt() {
+Maybe<ast::ForLoopStatement*> ParserImpl::for_stmt() {
Source source;
if (!match(Token::Type::kFor, &source))
return Failure::kNoMatch;
@@ -2086,45 +2086,9 @@
if (stmts.errored)
return Failure::kErrored;
- // The for statement is a syntactic sugar on top of the loop statement.
- // We create corresponding nodes in ast with the exact same behaviour
- // as we would expect from the loop statement.
- if (header->condition != nullptr) {
- // !condition
- auto* not_condition = create<ast::UnaryOpExpression>(
- header->condition->source(), ast::UnaryOp::kNot, header->condition);
- // { break; }
- auto* break_stmt = create<ast::BreakStatement>(not_condition->source());
- auto* break_body =
- create<ast::BlockStatement>(not_condition->source(), ast::StatementList{
- break_stmt,
- });
- // if (!condition) { break; }
- auto* break_if_not_condition =
- create<ast::IfStatement>(not_condition->source(), not_condition,
- break_body, ast::ElseStatementList{});
- stmts.value.insert(stmts.value.begin(), break_if_not_condition);
- }
-
- ast::BlockStatement* continuing_body = nullptr;
- if (header->continuing != nullptr) {
- continuing_body = create<ast::BlockStatement>(header->continuing->source(),
- ast::StatementList{
- header->continuing,
- });
- }
-
- auto* body = create<ast::BlockStatement>(source, stmts.value);
- auto* loop = create<ast::LoopStatement>(source, body, continuing_body);
-
- if (header->initializer != nullptr) {
- return create<ast::BlockStatement>(source, ast::StatementList{
- header->initializer,
- loop,
- });
- }
-
- return loop;
+ return create<ast::ForLoopStatement>(
+ source, header->initializer, header->condition, header->continuing,
+ create<ast::BlockStatement>(stmts.value));
}
// func_call_stmt
diff --git a/src/reader/wgsl/parser_impl.h b/src/reader/wgsl/parser_impl.h
index 232d952..78192aa 100644
--- a/src/reader/wgsl/parser_impl.h
+++ b/src/reader/wgsl/parser_impl.h
@@ -554,7 +554,7 @@
Expect<std::unique_ptr<ForHeader>> expect_for_header();
/// Parses a `for_stmt` grammar element
/// @returns the parsed for loop or nullptr
- Maybe<ast::Statement*> for_stmt();
+ Maybe<ast::ForLoopStatement*> for_stmt();
/// Parses a `continuing_stmt` grammar element
/// @returns the parsed statements
Maybe<ast::BlockStatement*> continuing_stmt();
diff --git a/src/reader/wgsl/parser_impl_for_stmt_test.cc b/src/reader/wgsl/parser_impl_for_stmt_test.cc
index 5b72897..1aac60e 100644
--- a/src/reader/wgsl/parser_impl_for_stmt_test.cc
+++ b/src/reader/wgsl/parser_impl_for_stmt_test.cc
@@ -14,141 +14,154 @@
#include "src/reader/wgsl/parser_impl_test_helper.h"
+#include "src/ast/discard_statement.h"
+
namespace tint {
namespace reader {
namespace wgsl {
namespace {
-class ForStmtTest : public ParserImplTest {
- public:
- void TestForLoop(std::string loop_str, std::string for_str) {
- auto p_loop = parser(loop_str);
- auto e_loop = p_loop->expect_statements();
- EXPECT_FALSE(e_loop.errored);
- EXPECT_FALSE(p_loop->has_error()) << p_loop->error();
-
- auto p_for = parser(for_str);
- auto e_for = p_for->expect_statements();
- EXPECT_FALSE(e_for.errored);
- EXPECT_FALSE(p_for->has_error()) << p_for->error();
-
- std::string loop = ast::BlockStatement({}, {}, e_loop.value).str(Sem());
- std::string for_ = ast::BlockStatement({}, {}, e_for.value).str(Sem());
- EXPECT_EQ(loop, for_);
- }
-};
+using ForStmtTest = ParserImplTest;
// Test an empty for loop.
TEST_F(ForStmtTest, Empty) {
- std::string for_str = "for (;;) { }";
- std::string loop_str = "loop { }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (;;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_EQ(fl->initializer(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop with non-empty body.
TEST_F(ForStmtTest, Body) {
- std::string for_str = "for (;;) { discard; }";
- std::string loop_str = "loop { discard; }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (;;) { discard; }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_EQ(fl->initializer(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ ASSERT_EQ(fl->body()->size(), 1u);
+ EXPECT_TRUE(fl->body()->statements()[0]->Is<ast::DiscardStatement>());
}
// Test a for loop declaring a variable in the initializer statement.
TEST_F(ForStmtTest, InitializerStatementDecl) {
- std::string for_str = "for (var i: i32 ;;) { }";
- std::string loop_str = "{ var i: i32; loop { } }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (var i: i32 ;;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ ASSERT_TRUE(Is<ast::VariableDeclStatement>(fl->initializer()));
+ auto* var = fl->initializer()->As<ast::VariableDeclStatement>()->variable();
+ EXPECT_FALSE(var->is_const());
+ EXPECT_EQ(var->constructor(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop declaring and initializing a variable in the initializer
// statement.
TEST_F(ForStmtTest, InitializerStatementDeclEqual) {
- std::string for_str = "for (var i: i32 = 0 ;;) { }";
- std::string loop_str = "{ var i: i32 = 0; loop { } }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (var i: i32 = 0 ;;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ ASSERT_TRUE(Is<ast::VariableDeclStatement>(fl->initializer()));
+ auto* var = fl->initializer()->As<ast::VariableDeclStatement>()->variable();
+ EXPECT_FALSE(var->is_const());
+ EXPECT_NE(var->constructor(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop declaring a const variable in the initializer statement.
TEST_F(ForStmtTest, InitializerStatementConstDecl) {
- std::string for_str = "for (let i: i32 = 0 ;;) { }";
- std::string loop_str = "{ let i: i32 = 0; loop { } }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (let i: i32 = 0 ;;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ ASSERT_TRUE(Is<ast::VariableDeclStatement>(fl->initializer()));
+ auto* var = fl->initializer()->As<ast::VariableDeclStatement>()->variable();
+ EXPECT_TRUE(var->is_const());
+ EXPECT_NE(var->constructor(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop assigning a variable in the initializer statement.
TEST_F(ForStmtTest, InitializerStatementAssignment) {
- std::string for_str = "var i: i32; for (i = 0 ;;) { }";
- std::string loop_str = "var i: i32; { i = 0; loop { } }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (i = 0 ;;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_TRUE(Is<ast::AssignmentStatement>(fl->initializer()));
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop calling a function in the initializer statement.
TEST_F(ForStmtTest, InitializerStatementFuncCall) {
- std::string for_str = "for (a(b,c) ;;) { }";
- std::string loop_str = "{ a(b,c); loop { } }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (a(b,c) ;;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_TRUE(Is<ast::CallStatement>(fl->initializer()));
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop with a break condition
TEST_F(ForStmtTest, BreakCondition) {
- std::string for_str = "for (; 0 == 1;) { }";
- std::string loop_str = "loop { if (!(0 == 1)) { break; } }";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (; 0 == 1;) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_EQ(fl->initializer(), nullptr);
+ EXPECT_TRUE(Is<ast::BinaryExpression>(fl->condition()));
+ EXPECT_EQ(fl->continuing(), nullptr);
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop assigning a variable in the continuing statement.
TEST_F(ForStmtTest, ContinuingAssignment) {
- std::string for_str = "var x: i32; for (;; x = 2) { }";
- std::string loop_str = "var x: i32; loop { continuing { x = 2; }}";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (;; x = 2) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_EQ(fl->initializer(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_TRUE(Is<ast::AssignmentStatement>(fl->continuing()));
+ EXPECT_TRUE(fl->body()->empty());
}
// Test a for loop calling a function in the continuing statement.
TEST_F(ForStmtTest, ContinuingFuncCall) {
- std::string for_str = "for (;; a(b,c)) { }";
- std::string loop_str = "loop { continuing { a(b,c); }}";
-
- TestForLoop(loop_str, for_str);
-}
-
-// Test a for loop with all statements non-empty.
-TEST_F(ForStmtTest, All) {
- std::string for_str =
- R"(for(var i : i32 = 0; i < 4; i = i + 1) {
- if (a == 0) {
- continue;
- }
- a = a + 2;
- })";
-
- std::string loop_str =
- R"({ // Introduce new scope for loop variable i
- var i : i32 = 0;
- loop {
- if (!(i < 4)) {
- break;
- }
-
- if (a == 0) {
- continue;
- }
- a = a + 2;
-
- continuing {
- i = i + 1;
- }
- }
- };)";
-
- TestForLoop(loop_str, for_str);
+ auto p = parser("for (;; a(b,c)) { }");
+ auto fl = p->for_stmt();
+ EXPECT_FALSE(p->has_error()) << p->error();
+ EXPECT_FALSE(fl.errored);
+ ASSERT_TRUE(fl.matched);
+ EXPECT_EQ(fl->initializer(), nullptr);
+ EXPECT_EQ(fl->condition(), nullptr);
+ EXPECT_TRUE(Is<ast::CallStatement>(fl->continuing()));
+ EXPECT_TRUE(fl->body()->empty());
}
class ForStmtErrorTest : public ParserImplTest {
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index 65c0ab0..881d380 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -62,6 +62,7 @@
#include "src/sem/storage_texture_type.h"
#include "src/sem/struct.h"
#include "src/sem/variable.h"
+#include "src/utils/defer.h"
#include "src/utils/get_or_create.h"
#include "src/utils/math.h"
#include "src/utils/scoped_assignment.h"
@@ -1915,6 +1916,10 @@
stmt->body(), current_statement_);
builder_->Sem().Add(stmt->body(), sem_block_body);
TINT_SCOPED_ASSIGNMENT(current_statement_, sem_block_body);
+ TINT_SCOPED_ASSIGNMENT(current_block_, sem_block_body);
+
+ variable_stack_.push_scope();
+ TINT_DEFER(variable_stack_.pop_scope());
if (auto* initializer = stmt->initializer()) {
Mark(initializer);
@@ -4011,9 +4016,8 @@
TINT_SCOPED_ASSIGNMENT(current_block_,
const_cast<sem::BlockStatement*>(sem_block));
variable_stack_.push_scope();
- bool result = callback();
- variable_stack_.pop_scope();
- return result;
+ TINT_DEFER(variable_stack_.pop_scope());
+ return callback();
}
std::string Resolver::VectorPretty(uint32_t size,
diff --git a/src/transform/spirv.cc b/src/transform/spirv.cc
index 1d5a460..0d84459 100644
--- a/src/transform/spirv.cc
+++ b/src/transform/spirv.cc
@@ -29,6 +29,7 @@
#include "src/sem/variable.h"
#include "src/transform/external_texture_transform.h"
#include "src/transform/fold_constants.h"
+#include "src/transform/for_loop_to_loop.h"
#include "src/transform/inline_pointer_lets.h"
#include "src/transform/manager.h"
#include "src/transform/simplify.h"
@@ -50,6 +51,7 @@
manager.Add<Simplify>(); // Required for arrayLength()
manager.Add<FoldConstants>();
manager.Add<ExternalTextureTransform>();
+ manager.Add<ForLoopToLoop>(); // Must come after ZeroInitWorkgroupMemory
auto transformedInput = manager.Run(in, data);
auto* cfg = data.Get<Config>();
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index c93caa7..7ab9ea5 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -43,6 +43,7 @@
#include "src/sem/variable.h"
#include "src/transform/calculate_array_length.h"
#include "src/transform/hlsl.h"
+#include "src/utils/defer.h"
#include "src/utils/get_or_create.h"
#include "src/utils/scoped_assignment.h"
#include "src/writer/append_vector.h"
@@ -2574,6 +2575,15 @@
}
bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
+ // Nest a for loop with a new block. In HLSL the initializer scope is not
+ // nested by the for-loop, so we may get variable redefinitions.
+ line() << "{";
+ increment_indent();
+ TINT_DEFER({
+ decrement_indent();
+ line() << "}";
+ });
+
TextBuffer init_buf;
if (auto* init = stmt->initializer()) {
TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
@@ -2581,16 +2591,6 @@
return false;
}
}
- bool multi_stmt_init = init_buf.lines.size() > 1;
- // For-loop has multi-statement initializer.
- // This cannot be emitted with a regular for loop, so instead nest the loop in
- // a new block scope prefixed with these initializer statements.
- if (multi_stmt_init) {
- line() << "{";
- increment_indent();
- current_buffer_->Append(init_buf);
- init_buf.lines.clear(); // Don't emit the initializer again in the 'for'
- }
TextBuffer cond_pre;
std::stringstream cond_buf;
@@ -2609,10 +2609,20 @@
}
}
- if (cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1) {
- // For-loop has multi-statement conditional and / or continuing.
- // This cannot be emitted with a regular for loop, so instead generate a
- // `while(true)` loop.
+ // If the for-loop has a multi-statement conditional and / or continuing, then
+ // we cannot emit this as a regular for-loop in HLSL. Instead we need to
+ // generate a `while(true)` loop.
+ bool emit_as_loop = cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1;
+
+ // If the for-loop has multi-statement initializer, or is going to be emitted
+ // as a `while(true)` loop, then declare the initializer statement(s) before
+ // the loop.
+ if (init_buf.lines.size() > 1 || (stmt->initializer() && emit_as_loop)) {
+ current_buffer_->Append(init_buf);
+ init_buf.lines.clear(); // Don't emit the initializer again in the 'for'
+ }
+
+ if (emit_as_loop) {
auto emit_continuing = [&]() {
current_buffer_->Append(cont_buf);
return true;
@@ -2620,23 +2630,24 @@
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
line() << "while (true) {";
- {
- ScopedIndent si(this);
+ increment_indent();
+ TINT_DEFER({
+ decrement_indent();
+ line() << "}";
+ });
- if (stmt->condition()) {
- current_buffer_->Append(cond_pre);
- line() << "if (!(" << cond_buf.str() << ")) { break; }";
- }
-
- if (!EmitStatements(stmt->body()->statements())) {
- return false;
- }
-
- if (!emit_continuing()) {
- return false;
- }
+ if (stmt->condition()) {
+ current_buffer_->Append(cond_pre);
+ line() << "if (!(" << cond_buf.str() << ")) { break; }";
}
- line() << "}";
+
+ if (!EmitStatements(stmt->body()->statements())) {
+ return false;
+ }
+
+ if (!emit_continuing()) {
+ return false;
+ }
} else {
// For-loop can be generated.
{
@@ -2666,12 +2677,6 @@
return false;
}
}
-
- line() << "}";
- }
-
- if (multi_stmt_init) {
- decrement_indent();
line() << "}";
}
diff --git a/src/writer/hlsl/generator_impl_loop_test.cc b/src/writer/hlsl/generator_impl_loop_test.cc
index 0811a6e..b9d006f 100644
--- a/src/writer/hlsl/generator_impl_loop_test.cc
+++ b/src/writer/hlsl/generator_impl_loop_test.cc
@@ -161,8 +161,10 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( for(; ; ) {
- return;
+ EXPECT_EQ(gen.result(), R"( {
+ for(; ; ) {
+ return;
+ }
}
)");
}
@@ -180,8 +182,10 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( for(int i = 0; ; ) {
- return;
+ EXPECT_EQ(gen.result(), R"( {
+ for(int i = 0; ; ) {
+ return;
+ }
}
)");
}
@@ -227,8 +231,10 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( for(; true; ) {
- return;
+ EXPECT_EQ(gen.result(), R"( {
+ for(; true; ) {
+ return;
+ }
}
)");
}
@@ -248,13 +254,15 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( while (true) {
- bool tint_tmp = true;
- if (tint_tmp) {
- tint_tmp = false;
+ EXPECT_EQ(gen.result(), R"( {
+ while (true) {
+ bool tint_tmp = true;
+ if (tint_tmp) {
+ tint_tmp = false;
+ }
+ if (!((tint_tmp))) { break; }
+ return;
}
- if (!((tint_tmp))) { break; }
- return;
}
)");
}
@@ -273,8 +281,10 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( for(; ; i = (i + 1)) {
- return;
+ EXPECT_EQ(gen.result(), R"( {
+ for(; ; i = (i + 1)) {
+ return;
+ }
}
)");
}
@@ -295,13 +305,15 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( while (true) {
- return;
- bool tint_tmp = true;
- if (tint_tmp) {
- tint_tmp = false;
+ EXPECT_EQ(gen.result(), R"( {
+ while (true) {
+ return;
+ bool tint_tmp = true;
+ if (tint_tmp) {
+ tint_tmp = false;
+ }
+ i = (tint_tmp);
}
- i = (tint_tmp);
}
)");
}
@@ -320,8 +332,10 @@
gen.increment_indent();
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
- EXPECT_EQ(gen.result(), R"( for(int i = 0; true; i = (i + 1)) {
- return;
+ EXPECT_EQ(gen.result(), R"( {
+ for(int i = 0; true; i = (i + 1)) {
+ return;
+ }
}
)");
}
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 42d5a36..809ca10 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -54,6 +54,7 @@
#include "src/sem/vector_type.h"
#include "src/sem/void_type.h"
#include "src/transform/msl.h"
+#include "src/utils/defer.h"
#include "src/utils/scoped_assignment.h"
#include "src/writer/float_to_string.h"
@@ -1493,6 +1494,116 @@
return true;
}
+bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
+ TextBuffer init_buf;
+ if (auto* init = stmt->initializer()) {
+ TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
+ if (!EmitStatement(init)) {
+ return false;
+ }
+ }
+
+ TextBuffer cond_pre;
+ std::stringstream cond_buf;
+ if (auto* cond = stmt->condition()) {
+ TINT_SCOPED_ASSIGNMENT(current_buffer_, &cond_pre);
+ if (!EmitExpression(cond_buf, cond)) {
+ return false;
+ }
+ }
+
+ TextBuffer cont_buf;
+ if (auto* cont = stmt->continuing()) {
+ TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf);
+ if (!EmitStatement(cont)) {
+ return false;
+ }
+ }
+
+ // If the for-loop has a multi-statement conditional and / or continuing, then
+ // we cannot emit this as a regular for-loop in MSL. Instead we need to
+ // generate a `while(true)` loop.
+ bool emit_as_loop = cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1;
+
+ // If the for-loop has multi-statement initializer, or is going to be emitted
+ // as a `while(true)` loop, then declare the initializer statement(s) before
+ // the loop in a new block.
+ bool nest_in_block =
+ init_buf.lines.size() > 1 || (stmt->initializer() && emit_as_loop);
+ if (nest_in_block) {
+ line() << "{";
+ increment_indent();
+ current_buffer_->Append(init_buf);
+ init_buf.lines.clear(); // Don't emit the initializer again in the 'for'
+ }
+ TINT_DEFER({
+ if (nest_in_block) {
+ decrement_indent();
+ line() << "}";
+ }
+ });
+
+ if (emit_as_loop) {
+ auto emit_continuing = [&]() {
+ current_buffer_->Append(cont_buf);
+ return true;
+ };
+
+ TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
+ line() << "while (true) {";
+ increment_indent();
+ TINT_DEFER({
+ decrement_indent();
+ line() << "}";
+ });
+
+ if (stmt->condition()) {
+ current_buffer_->Append(cond_pre);
+ line() << "if (!(" << cond_buf.str() << ")) { break; }";
+ }
+
+ if (!EmitStatements(stmt->body()->statements())) {
+ return false;
+ }
+
+ if (!emit_continuing()) {
+ return false;
+ }
+ } else {
+ // For-loop can be generated.
+ {
+ auto out = line();
+ out << "for";
+ {
+ ScopedParen sp(out);
+
+ if (!init_buf.lines.empty()) {
+ out << init_buf.lines[0].content << " ";
+ } else {
+ out << "; ";
+ }
+
+ out << cond_buf.str() << "; ";
+
+ if (!cont_buf.lines.empty()) {
+ out << TrimSuffix(cont_buf.lines[0].content, ";");
+ }
+ }
+ out << " {";
+ }
+ {
+ auto emit_continuing = [] { return true; };
+ TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
+ if (!EmitStatementsWithIndent(stmt->body()->statements())) {
+ return false;
+ }
+ }
+ line() << "}";
+ }
+
+ return true;
+}
+
bool GeneratorImpl::EmitDiscard(ast::DiscardStatement*) {
// TODO(dsinclair): Verify this is correct when the discard semantics are
// defined for WGSL (https://github.com/gpuweb/gpuweb/issues/361)
@@ -1635,6 +1746,9 @@
if (auto* l = stmt->As<ast::LoopStatement>()) {
return EmitLoop(l);
}
+ if (auto* l = stmt->As<ast::ForLoopStatement>()) {
+ return EmitForLoop(l);
+ }
if (auto* r = stmt->As<ast::ReturnStatement>()) {
return EmitReturn(r);
}
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index 0f5bf1a..1495178 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -170,6 +170,10 @@
/// @param stmt the statement to emit
/// @returns true if the statement was emitted
bool EmitLoop(ast::LoopStatement* stmt);
+ /// Handles a for loop statement
+ /// @param stmt the statement to emit
+ /// @returns true if the statement was emitted
+ bool EmitForLoop(ast::ForLoopStatement* stmt);
/// Handles a member accessor expression
/// @param out the output of the expression stream
/// @param expr the member accessor expression
diff --git a/src/writer/msl/generator_impl_loop_test.cc b/src/writer/msl/generator_impl_loop_test.cc
index 0d18299..a6ae411 100644
--- a/src/writer/msl/generator_impl_loop_test.cc
+++ b/src/writer/msl/generator_impl_loop_test.cc
@@ -141,6 +141,227 @@
)");
}
+TEST_F(MslGeneratorImplTest, Emit_ForLoop) {
+ // for(; ; ) {
+ // return;
+ // }
+
+ auto* f = For(nullptr, nullptr, nullptr, Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( for(; ; ) {
+ return;
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleInit) {
+ // for(var i : i32; ; ) {
+ // return;
+ // }
+
+ auto* f = For(Decl(Var("i", ty.i32())), nullptr, nullptr, Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( for(int i = 0; ; ) {
+ return;
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtInit) {
+ // var<workgroup> a : atomic<i32>;
+ // for(var b = atomicCompareExchangeWeak(&a, 1, 2); ; ) {
+ // return;
+ // }
+ Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+ auto* multi_stmt = Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2);
+ auto* f = For(Decl(Var("b", nullptr, multi_stmt)), nullptr, nullptr,
+ Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( {
+ int prev_value = 1;
+ bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+ int2 b = int2(prev_value, matched);
+ for(; ; ) {
+ return;
+ }
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleCond) {
+ // for(; true; ) {
+ // return;
+ // }
+
+ auto* f = For(nullptr, true, nullptr, Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( for(; true; ) {
+ return;
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtCond) {
+ // var<workgroup> a : atomic<i32>;
+ // for(; atomicCompareExchangeWeak(&a, 1, 2).x == 0; ) {
+ // return;
+ // }
+
+ Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+ auto* multi_stmt = create<ast::BinaryExpression>(
+ ast::BinaryOp::kEqual,
+ MemberAccessor(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2),
+ "x"),
+ Expr(0));
+ auto* f = For(nullptr, multi_stmt, nullptr, Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
+ int prev_value = 1;
+ bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+ if (!((int2(prev_value, matched).x == 0))) { break; }
+ return;
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleCont) {
+ // for(; ; i = i + 1) {
+ // return;
+ // }
+
+ auto* v = Decl(Var("i", ty.i32()));
+ auto* f = For(nullptr, nullptr, Assign("i", Add("i", 1)), Block(Return()));
+ WrapInFunction(v, f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( for(; ; i = (i + 1)) {
+ return;
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtCont) {
+ // var<workgroup> a : atomic<i32>;
+ // for(; ; ignore(atomicCompareExchangeWeak(&a, 1, 2))) {
+ // return;
+ // }
+
+ Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+ auto* multi_stmt =
+ Ignore(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2));
+ auto* f = For(nullptr, nullptr, multi_stmt, Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( while (true) {
+ return;
+ int prev_value = 1;
+ bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+ (void) int2(prev_value, matched);
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleInitCondCont) {
+ // for(var i : i32; true; i = i + 1) {
+ // return;
+ // }
+
+ auto* f = For(Decl(Var("i", ty.i32())), true, Assign("i", Add("i", 1)),
+ Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( for(int i = 0; true; i = (i + 1)) {
+ return;
+ }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtInitCondCont) {
+ // var<workgroup> a : atomic<i32>;
+ // for(var b = atomicCompareExchangeWeak(&a, 1, 2);
+ // atomicCompareExchangeWeak(&a, 1, 2).x == 0;
+ // ignore(atomicCompareExchangeWeak(&a, 1, 2))) {
+ // return;
+ // }
+ Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+ auto* multi_stmt_a = Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2);
+ auto* multi_stmt_b = create<ast::BinaryExpression>(
+ ast::BinaryOp::kEqual,
+ MemberAccessor(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2),
+ "x"),
+ Expr(0));
+ auto* multi_stmt_c =
+ Ignore(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2));
+ auto* f = For(Decl(Var("b", nullptr, multi_stmt_a)), multi_stmt_b,
+ multi_stmt_c, Block(Return()));
+ WrapInFunction(f);
+
+ GeneratorImpl& gen = Build();
+
+ gen.increment_indent();
+
+ ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+ EXPECT_EQ(gen.result(), R"( {
+ int prev_value = 1;
+ bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+ int2 b = int2(prev_value, matched);
+ while (true) {
+ int prev_value_1 = 1;
+ bool matched_1 = atomic_compare_exchange_weak_explicit(&(a), &prev_value_1, 2, memory_order_relaxed, memory_order_relaxed);
+ if (!((int2(prev_value_1, matched_1).x == 0))) { break; }
+ return;
+ int prev_value_2 = 1;
+ bool matched_2 = atomic_compare_exchange_weak_explicit(&(a), &prev_value_2, 2, memory_order_relaxed, memory_order_relaxed);
+ (void) int2(prev_value_2, matched_2);
+ }
+ }
+)");
+}
+
} // namespace
} // namespace msl
} // namespace writer
diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl
index 2c9ea16..e7a735c 100644
--- a/test/bug/tint/534.wgsl.expected.hlsl
+++ b/test/bug/tint/534.wgsl.expected.hlsl
@@ -43,7 +43,7 @@
uint i = 0u;
while (true) {
const uint scalar_offset_1 = (12u) / 4;
- if (!(!(!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))))) { break; }
+ if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) { break; }
Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
bool tint_tmp_1 = success;
if (tint_tmp_1) {
diff --git a/test/bug/tint/534.wgsl.expected.msl b/test/bug/tint/534.wgsl.expected.msl
index 429c94b..356f44d 100644
--- a/test/bug/tint/534.wgsl.expected.msl
+++ b/test/bug/tint/534.wgsl.expected.msl
@@ -27,18 +27,9 @@
bool success = true;
uint4 srcColorBits = 0u;
uint4 dstColorBits = uint4(dstColor);
- {
- uint i = 0u;
- while (true) {
- if (!((i < uniforms.channelCount))) {
- break;
- }
- srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
- success = (success && (srcColorBits[i] == dstColorBits[i]));
- {
- i = (i + 1u);
- }
- }
+ for(uint i = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
+ srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+ success = (success && (srcColorBits[i] == dstColorBits[i]));
}
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
if (success) {
diff --git a/test/bug/tint/534.wgsl.expected.wgsl b/test/bug/tint/534.wgsl.expected.wgsl
index 91a330b..1fd409c 100644
--- a/test/bug/tint/534.wgsl.expected.wgsl
+++ b/test/bug/tint/534.wgsl.expected.wgsl
@@ -36,19 +36,9 @@
var success : bool = true;
var srcColorBits : vec4<u32>;
var dstColorBits : vec4<u32> = vec4<u32>(dstColor);
- {
- var i : u32 = 0u;
- loop {
- if (!((i < uniforms.channelCount))) {
- break;
- }
- srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
- success = (success && (srcColorBits[i] == dstColorBits[i]));
-
- continuing {
- i = (i + 1u);
- }
- }
+ for(var i : u32 = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
+ srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+ success = (success && (srcColorBits[i] == dstColorBits[i]));
}
var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x);
if (success) {
diff --git a/test/bug/tint/744.wgsl.expected.hlsl b/test/bug/tint/744.wgsl.expected.hlsl
index 0314864..47f47a7 100644
--- a/test/bug/tint/744.wgsl.expected.hlsl
+++ b/test/bug/tint/744.wgsl.expected.hlsl
@@ -19,8 +19,7 @@
const uint dimOutter = uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4];
uint result = 0u;
{
- uint i = 0u;
- for(; !(!((i < dimInner))); i = (i + 1u)) {
+ for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
const uint a = (i + (resultCell.x * dimInner));
const uint b = (resultCell.y + (i * dimOutter));
result = (result + (firstMatrix.Load((4u * a)) * secondMatrix.Load((4u * b))));
diff --git a/test/bug/tint/744.wgsl.expected.msl b/test/bug/tint/744.wgsl.expected.msl
index 7855d18..16ebd85 100644
--- a/test/bug/tint/744.wgsl.expected.msl
+++ b/test/bug/tint/744.wgsl.expected.msl
@@ -15,19 +15,10 @@
uint const dimInner = uniforms.aShape.y;
uint const dimOutter = uniforms.outShape.y;
uint result = 0u;
- {
- uint i = 0u;
- while (true) {
- if (!((i < dimInner))) {
- break;
- }
- uint const a = (i + (resultCell.x * dimInner));
- uint const b = (resultCell.y + (i * dimOutter));
- result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
- {
- i = (i + 1u);
- }
- }
+ for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
+ uint const a = (i + (resultCell.x * dimInner));
+ uint const b = (resultCell.y + (i * dimOutter));
+ result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
}
uint const index = (resultCell.y + (resultCell.x * dimOutter));
resultMatrix.numbers[index] = result;
diff --git a/test/bug/tint/744.wgsl.expected.wgsl b/test/bug/tint/744.wgsl.expected.wgsl
index d96d510..ad6c01e 100644
--- a/test/bug/tint/744.wgsl.expected.wgsl
+++ b/test/bug/tint/744.wgsl.expected.wgsl
@@ -24,20 +24,10 @@
let dimInner : u32 = uniforms.aShape.y;
let dimOutter : u32 = uniforms.outShape.y;
var result : u32 = 0u;
- {
- var i : u32 = 0u;
- loop {
- if (!((i < dimInner))) {
- break;
- }
- let a : u32 = (i + (resultCell.x * dimInner));
- let b : u32 = (resultCell.y + (i * dimOutter));
- result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
-
- continuing {
- i = (i + 1u);
- }
- }
+ for(var i : u32 = 0u; (i < dimInner); i = (i + 1u)) {
+ let a : u32 = (i + (resultCell.x * dimInner));
+ let b : u32 = (resultCell.y + (i * dimOutter));
+ result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
}
let index : u32 = (resultCell.y + (resultCell.x * dimOutter));
resultMatrix.numbers[index] = result;
diff --git a/test/bug/tint/757.wgsl.expected.hlsl b/test/bug/tint/757.wgsl.expected.hlsl
index 72ed5bd..da3073c 100644
--- a/test/bug/tint/757.wgsl.expected.hlsl
+++ b/test/bug/tint/757.wgsl.expected.hlsl
@@ -16,8 +16,7 @@
flatIndex = (flatIndex * 1u);
float4 texel = myTexture.Load(int4(GlobalInvocationID.xy, 0, 0));
{
- uint i = 0u;
- for(; !(!((i < 1u))); i = (i + 1u)) {
+ for(uint i = 0u; (i < 1u); i = (i + 1u)) {
result.Store((4u * (flatIndex + i)), asuint(texel.r));
}
}
diff --git a/test/bug/tint/757.wgsl.expected.msl b/test/bug/tint/757.wgsl.expected.msl
index 03e2067..8699607 100644
--- a/test/bug/tint/757.wgsl.expected.msl
+++ b/test/bug/tint/757.wgsl.expected.msl
@@ -12,17 +12,8 @@
uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
flatIndex = (flatIndex * 1u);
float4 texel = tint_symbol_2.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
- {
- uint i = 0u;
- while (true) {
- if (!((i < 1u))) {
- break;
- }
- result.values[(flatIndex + i)] = texel.r;
- {
- i = (i + 1u);
- }
- }
+ for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+ result.values[(flatIndex + i)] = texel.r;
}
return;
}
diff --git a/test/bug/tint/757.wgsl.expected.wgsl b/test/bug/tint/757.wgsl.expected.wgsl
index c7c0d8c..9690af5 100644
--- a/test/bug/tint/757.wgsl.expected.wgsl
+++ b/test/bug/tint/757.wgsl.expected.wgsl
@@ -19,17 +19,7 @@
var flatIndex : u32 = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
flatIndex = (flatIndex * 1u);
var texel : vec4<f32> = textureLoad(myTexture, vec2<i32>(GlobalInvocationID.xy), 0, 0);
- {
- var i : u32 = 0u;
- loop {
- if (!((i < 1u))) {
- break;
- }
- result.values[(flatIndex + i)] = texel.r;
-
- continuing {
- i = (i + 1u);
- }
- }
+ for(var i : u32 = 0u; (i < 1u); i = (i + 1u)) {
+ result.values[(flatIndex + i)] = texel.r;
}
}
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 3de9408..5e0ebb4 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -69,14 +69,22 @@
const uint3 global_id = tint_symbol.global_id;
const uint local_invocation_index = tint_symbol.local_invocation_index;
if ((local_invocation_index == 0u)) {
- for(int i = 0; (i < 64); i = (i + 1)) {
- for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
- mm_Asub[i][i_1] = 0.0f;
+ {
+ for(int i = 0; (i < 64); i = (i + 1)) {
+ {
+ for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+ mm_Asub[i][i_1] = 0.0f;
+ }
+ }
}
}
- for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
- for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
- mm_Bsub[i_2][i_3] = 0.0f;
+ {
+ for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+ {
+ for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
+ mm_Bsub[i_2][i_3] = 0.0f;
+ }
+ }
}
}
}
@@ -91,8 +99,7 @@
float ACached = 0.0f;
float BCached[4] = (float[4])0;
{
- uint index = 0u;
- for(; !(!((index < (RowPerThread * ColPerThread)))); index = (index + 1u)) {
+ for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
acc[index] = 0.0f;
}
}
@@ -101,14 +108,11 @@
const uint RowPerThreadB = (TileInner / 16u);
const uint tileRowB = (local_id.y * RowPerThreadB);
{
- uint t = 0u;
- for(; !(!((t < numTiles))); t = (t + 1u)) {
+ for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThreadA))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRow + innerRow);
const uint inputCol = (tileColA + innerCol);
mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
@@ -117,11 +121,9 @@
}
}
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThreadB))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint inputRow = (tileRowB + innerRow);
const uint inputCol = (tileCol + innerCol);
mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
@@ -131,21 +133,17 @@
}
GroupMemoryBarrierWithGroupSync();
{
- uint k = 0u;
- for(; !(!((k < TileInner))); k = (k + 1u)) {
+ for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
{
- uint inner = 0u;
- for(; !(!((inner < ColPerThread))); inner = (inner + 1u)) {
+ for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
BCached[inner] = mm_Bsub[k][(tileCol + inner)];
}
}
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
ACached = mm_Asub[(tileRow + innerRow)][k];
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint index = ((innerRow * ColPerThread) + innerCol);
acc[index] = (acc[index] + (ACached * BCached[innerCol]));
}
@@ -158,11 +156,9 @@
}
}
{
- uint innerRow = 0u;
- for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
{
- uint innerCol = 0u;
- for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
const uint index = ((innerRow * ColPerThread) + innerCol);
mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
}
diff --git a/test/bug/tint/914.wgsl.expected.msl b/test/bug/tint/914.wgsl.expected.msl
index d9a4b82..eb27d80 100644
--- a/test/bug/tint/914.wgsl.expected.msl
+++ b/test/bug/tint/914.wgsl.expected.msl
@@ -68,155 +68,47 @@
tint_array_wrapper_2 acc = {};
float ACached = 0.0f;
tint_array_wrapper_3 BCached = {};
- {
- uint index = 0u;
- while (true) {
- if (!((index < (RowPerThread * ColPerThread)))) {
- break;
- }
- acc.arr[index] = 0.0f;
- {
- index = (index + 1u);
- }
- }
+ for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
+ acc.arr[index] = 0.0f;
}
uint const ColPerThreadA = (TileInner / 16u);
uint const tileColA = (local_id.x * ColPerThreadA);
uint const RowPerThreadB = (TileInner / 16u);
uint const tileRowB = (local_id.y * RowPerThreadB);
- {
- uint t = 0u;
- while (true) {
- if (!((t < numTiles))) {
- break;
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThreadA))) {
- break;
- }
- uint const inputRow = (tileRow + innerRow);
- uint const inputCol = (tileColA + innerCol);
- tint_symbol_4.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
- {
- innerCol = (innerCol + 1u);
- }
- }
- }
- {
- innerRow = (innerRow + 1u);
- }
- }
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThreadB))) {
- break;
- }
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- uint const inputRow = (tileRowB + innerRow);
- uint const inputCol = (tileCol + innerCol);
- tint_symbol_5.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
- {
- innerCol = (innerCol + 1u);
- }
- }
- }
- {
- innerRow = (innerRow + 1u);
- }
- }
- }
- threadgroup_barrier(mem_flags::mem_threadgroup);
- {
- uint k = 0u;
- while (true) {
- if (!((k < TileInner))) {
- break;
- }
- {
- uint inner = 0u;
- while (true) {
- if (!((inner < ColPerThread))) {
- break;
- }
- BCached.arr[inner] = tint_symbol_5.arr[k].arr[(tileCol + inner)];
- {
- inner = (inner + 1u);
- }
- }
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- uint const index = ((innerRow * ColPerThread) + innerCol);
- acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
- {
- innerCol = (innerCol + 1u);
- }
- }
- }
- {
- innerRow = (innerRow + 1u);
- }
- }
- }
- {
- k = (k + 1u);
- }
- }
- }
- threadgroup_barrier(mem_flags::mem_threadgroup);
- {
- t = (t + 1u);
+ for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+ uint const inputRow = (tileRow + innerRow);
+ uint const inputCol = (tileColA + innerCol);
+ tint_symbol_4.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
}
}
- }
- {
- uint innerRow = 0u;
- while (true) {
- if (!((innerRow < RowPerThread))) {
- break;
+ for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ uint const inputRow = (tileRowB + innerRow);
+ uint const inputCol = (tileCol + innerCol);
+ tint_symbol_5.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
}
- {
- uint innerCol = 0u;
- while (true) {
- if (!((innerCol < ColPerThread))) {
- break;
- }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
+ for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
+ BCached.arr[inner] = tint_symbol_5.arr[k].arr[(tileCol + inner)];
+ }
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
uint const index = ((innerRow * ColPerThread) + innerCol);
- mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
- {
- innerCol = (innerCol + 1u);
- }
+ acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
}
}
- {
- innerRow = (innerRow + 1u);
- }
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ }
+ for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ uint const index = ((innerRow * ColPerThread) + innerCol);
+ mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
}
}
return;
diff --git a/test/bug/tint/914.wgsl.expected.wgsl b/test/bug/tint/914.wgsl.expected.wgsl
index ed45fab..17ca3f3 100644
--- a/test/bug/tint/914.wgsl.expected.wgsl
+++ b/test/bug/tint/914.wgsl.expected.wgsl
@@ -65,167 +65,47 @@
var acc : array<f32, 16>;
var ACached : f32;
var BCached : array<f32, 4>;
- {
- var index : u32 = 0u;
- loop {
- if (!((index < (RowPerThread * ColPerThread)))) {
- break;
- }
- acc[index] = 0.0;
-
- continuing {
- index = (index + 1u);
- }
- }
+ for(var index : u32 = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
+ acc[index] = 0.0;
}
let ColPerThreadA : u32 = (TileInner / 16u);
let tileColA : u32 = (local_id.x * ColPerThreadA);
let RowPerThreadB : u32 = (TileInner / 16u);
let tileRowB : u32 = (local_id.y * RowPerThreadB);
- {
- var t : u32 = 0u;
- loop {
- if (!((t < numTiles))) {
- break;
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThreadA))) {
- break;
- }
- let inputRow : u32 = (tileRow + innerRow);
- let inputCol : u32 = (tileColA + innerCol);
- mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
-
- continuing {
- innerCol = (innerCol + 1u);
- }
- }
- }
-
- continuing {
- innerRow = (innerRow + 1u);
- }
- }
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThreadB))) {
- break;
- }
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- let inputRow : u32 = (tileRowB + innerRow);
- let inputCol : u32 = (tileCol + innerCol);
- mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
-
- continuing {
- innerCol = (innerCol + 1u);
- }
- }
- }
-
- continuing {
- innerRow = (innerRow + 1u);
- }
- }
- }
- workgroupBarrier();
- {
- var k : u32 = 0u;
- loop {
- if (!((k < TileInner))) {
- break;
- }
- {
- var inner : u32 = 0u;
- loop {
- if (!((inner < ColPerThread))) {
- break;
- }
- BCached[inner] = mm_Bsub[k][(tileCol + inner)];
-
- continuing {
- inner = (inner + 1u);
- }
- }
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThread))) {
- break;
- }
- ACached = mm_Asub[(tileRow + innerRow)][k];
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThread))) {
- break;
- }
- let index : u32 = ((innerRow * ColPerThread) + innerCol);
- acc[index] = (acc[index] + (ACached * BCached[innerCol]));
-
- continuing {
- innerCol = (innerCol + 1u);
- }
- }
- }
-
- continuing {
- innerRow = (innerRow + 1u);
- }
- }
- }
-
- continuing {
- k = (k + 1u);
- }
- }
- }
- workgroupBarrier();
-
- continuing {
- t = (t + 1u);
+ for(var t : u32 = 0u; (t < numTiles); t = (t + 1u)) {
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+ let inputRow : u32 = (tileRow + innerRow);
+ let inputCol : u32 = (tileColA + innerCol);
+ mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
}
}
- }
- {
- var innerRow : u32 = 0u;
- loop {
- if (!((innerRow < RowPerThread))) {
- break;
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ let inputRow : u32 = (tileRowB + innerRow);
+ let inputCol : u32 = (tileCol + innerCol);
+ mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
}
- {
- var innerCol : u32 = 0u;
- loop {
- if (!((innerCol < ColPerThread))) {
- break;
- }
+ }
+ workgroupBarrier();
+ for(var k : u32 = 0u; (k < TileInner); k = (k + 1u)) {
+ for(var inner : u32 = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
+ BCached[inner] = mm_Bsub[k][(tileCol + inner)];
+ }
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ ACached = mm_Asub[(tileRow + innerRow)][k];
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
let index : u32 = ((innerRow * ColPerThread) + innerCol);
- mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
-
- continuing {
- innerCol = (innerCol + 1u);
- }
+ acc[index] = (acc[index] + (ACached * BCached[innerCol]));
}
}
-
- continuing {
- innerRow = (innerRow + 1u);
- }
+ }
+ workgroupBarrier();
+ }
+ for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+ for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+ let index : u32 = ((innerRow * ColPerThread) + innerCol);
+ mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
}
}
}
diff --git a/test/bug/tint/942.wgsl.expected.hlsl b/test/bug/tint/942.wgsl.expected.hlsl
index ff2739f..695e8ae 100644
--- a/test/bug/tint/942.wgsl.expected.hlsl
+++ b/test/bug/tint/942.wgsl.expected.hlsl
@@ -22,9 +22,13 @@
const uint3 LocalInvocationID = tint_symbol.LocalInvocationID;
const uint local_invocation_index = tint_symbol.local_invocation_index;
if ((local_invocation_index == 0u)) {
- for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
- for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
- tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+ {
+ for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
+ {
+ for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
+ tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+ }
+ }
}
}
}
@@ -37,11 +41,9 @@
const uint scalar_offset_1 = (4u) / 4;
const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[scalar_offset_1 / 4][scalar_offset_1 % 4], 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
{
- uint r = 0u;
- for(; !(!((r < 4u))); r = (r + 1u)) {
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
- uint c = 0u;
- for(; !(!((c < 4u))); c = (c + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 loadIndex = (baseIndex + int2(int(c), int(r)));
const uint scalar_offset_2 = (0u) / 4;
if ((flip[scalar_offset_2 / 4][scalar_offset_2 % 4] != 0u)) {
@@ -54,11 +56,9 @@
}
GroupMemoryBarrierWithGroupSync();
{
- uint r = 0u;
- for(; !(!((r < 4u))); r = (r + 1u)) {
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
{
- uint c = 0u;
- for(; !(!((c < 4u))); c = (c + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
int2 writeIndex = (baseIndex + int2(int(c), int(r)));
const uint scalar_offset_3 = (0u) / 4;
if ((flip[scalar_offset_3 / 4][scalar_offset_3 % 4] != 0u)) {
@@ -79,7 +79,7 @@
uint f = 0u;
while (true) {
const uint scalar_offset_4 = (0u) / 4;
- if (!(!(!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))))) { break; }
+ if (!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))) { break; }
uint i = ((center + f) - filterOffset);
const uint scalar_offset_5 = (0u) / 4;
acc = (acc + ((1.0f / float(params[scalar_offset_5 / 4][scalar_offset_5 % 4])) * tile[r][i]));
diff --git a/test/bug/tint/942.wgsl.expected.msl b/test/bug/tint/942.wgsl.expected.msl
index 16e440d..cecfa13 100644
--- a/test/bug/tint/942.wgsl.expected.msl
+++ b/test/bug/tint/942.wgsl.expected.msl
@@ -25,75 +25,30 @@
uint const filterOffset = ((params.filterDim - 1u) / 2u);
int2 const dims = int2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(0));
int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
- {
- uint r = 0u;
- while (true) {
- if (!((r < 4u))) {
- break;
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+ int2 loadIndex = (baseIndex + int2(int(c), int(r)));
+ if ((flip.value != 0u)) {
+ loadIndex = loadIndex.yx;
}
- {
- uint c = 0u;
- while (true) {
- if (!((c < 4u))) {
- break;
- }
- int2 loadIndex = (baseIndex + int2(int(c), int(r)));
- if ((flip.value != 0u)) {
- loadIndex = loadIndex.yx;
- }
- tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
- {
- c = (c + 1u);
- }
- }
- }
- {
- r = (r + 1u);
- }
+ tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- {
- uint r = 0u;
- while (true) {
- if (!((r < 4u))) {
- break;
+ for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+ for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+ int2 writeIndex = (baseIndex + int2(int(c), int(r)));
+ if ((flip.value != 0u)) {
+ writeIndex = writeIndex.yx;
}
- {
- uint c = 0u;
- while (true) {
- if (!((c < 4u))) {
- break;
- }
- int2 writeIndex = (baseIndex + int2(int(c), int(r)));
- if ((flip.value != 0u)) {
- writeIndex = writeIndex.yx;
- }
- uint const center = ((4u * LocalInvocationID.x) + c);
- if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
- float3 acc = float3(0.0f, 0.0f, 0.0f);
- {
- uint f = 0u;
- while (true) {
- if (!((f < params.filterDim))) {
- break;
- }
- uint i = ((center + f) - filterOffset);
- acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
- {
- f = (f + 1u);
- }
- }
- }
- tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
- }
- {
- c = (c + 1u);
- }
+ uint const center = ((4u * LocalInvocationID.x) + c);
+ if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+ float3 acc = float3(0.0f, 0.0f, 0.0f);
+ for(uint f = 0u; (f < params.filterDim); f = (f + 1u)) {
+ uint i = ((center + f) - filterOffset);
+ acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
}
- }
- {
- r = (r + 1u);
+ tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
}
}
}
diff --git a/test/bug/tint/942.wgsl.expected.wgsl b/test/bug/tint/942.wgsl.expected.wgsl
index d261497..b936d4e 100644
--- a/test/bug/tint/942.wgsl.expected.wgsl
+++ b/test/bug/tint/942.wgsl.expected.wgsl
@@ -26,80 +26,30 @@
let filterOffset : u32 = ((params.filterDim - 1u) / 2u);
let dims : vec2<i32> = textureDimensions(inputTex, 0);
let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
- {
- var r : u32 = 0u;
- loop {
- if (!((r < 4u))) {
- break;
+ for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
+ for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
+ var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+ if ((flip.value != 0u)) {
+ loadIndex = loadIndex.yx;
}
- {
- var c : u32 = 0u;
- loop {
- if (!((c < 4u))) {
- break;
- }
- var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
- if ((flip.value != 0u)) {
- loadIndex = loadIndex.yx;
- }
- tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
-
- continuing {
- c = (c + 1u);
- }
- }
- }
-
- continuing {
- r = (r + 1u);
- }
+ tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
}
}
workgroupBarrier();
- {
- var r : u32 = 0u;
- loop {
- if (!((r < 4u))) {
- break;
+ for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
+ for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
+ var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+ if ((flip.value != 0u)) {
+ writeIndex = writeIndex.yx;
}
- {
- var c : u32 = 0u;
- loop {
- if (!((c < 4u))) {
- break;
- }
- var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
- if ((flip.value != 0u)) {
- writeIndex = writeIndex.yx;
- }
- let center : u32 = ((4u * LocalInvocationID.x) + c);
- if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
- var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
- {
- var f : u32 = 0u;
- loop {
- if (!((f < params.filterDim))) {
- break;
- }
- var i : u32 = ((center + f) - filterOffset);
- acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
-
- continuing {
- f = (f + 1u);
- }
- }
- }
- textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
- }
-
- continuing {
- c = (c + 1u);
- }
+ let center : u32 = ((4u * LocalInvocationID.x) + c);
+ if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+ var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
+ for(var f : u32 = 0u; (f < params.filterDim); f = (f + 1u)) {
+ var i : u32 = ((center + f) - filterOffset);
+ acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
}
- }
-
- continuing {
- r = (r + 1u);
+ textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
}
}
}
diff --git a/test/bug/tint/943.spvasm.expected.hlsl b/test/bug/tint/943.spvasm.expected.hlsl
index 47f46cd..a317e46 100644
--- a/test/bug/tint/943.spvasm.expected.hlsl
+++ b/test/bug/tint/943.spvasm.expected.hlsl
@@ -191,10 +191,14 @@
const int x_152 = dimInner;
numTiles = (((x_152 - 1) / 64) + 1);
innerRow = 0;
- for(; (innerRow < 1); innerRow = (innerRow + 1)) {
- innerCol = 0;
- for(; (innerCol < 1); innerCol = (innerCol + 1)) {
- acc[innerRow][innerCol] = 0.0f;
+ {
+ for(; (innerRow < 1); innerRow = (innerRow + 1)) {
+ innerCol = 0;
+ {
+ for(; (innerCol < 1); innerCol = (innerCol + 1)) {
+ acc[innerRow][innerCol] = 0.0f;
+ }
+ }
}
}
const uint x_187 = gl_LocalInvocationID.x;
@@ -202,100 +206,120 @@
const uint x_192 = gl_LocalInvocationID.y;
tileRowB = (asint(x_192) * 1);
t = 0;
- for(; (t < numTiles); t = (t + 1)) {
- innerRow_1 = 0;
- for(; (innerRow_1 < 1); innerRow_1 = (innerRow_1 + 1)) {
- innerCol_1 = 0;
- for(; (innerCol_1 < 64); innerCol_1 = (innerCol_1 + 1)) {
- inputRow = (tileRow + innerRow_1);
- inputCol = (tileColA + innerCol_1);
- const int x_233 = inputRow;
- const int x_234 = inputCol;
- const int x_238 = t;
- const int x_240 = inputCol;
- param_3 = (globalRow + innerRow_1);
- param_4 = ((x_238 * 64) + x_240);
- const float x_244 = mm_readA_i1_i1_(param_3, param_4);
- mm_Asub[x_233][x_234] = x_244;
- }
- }
- innerRow_2 = 0;
- for(; (innerRow_2 < 1); innerRow_2 = (innerRow_2 + 1)) {
- innerCol_2 = 0;
- for(; (innerCol_2 < 1); innerCol_2 = (innerCol_2 + 1)) {
- inputRow_1 = (tileRowB + innerRow_2);
- inputCol_1 = (tileCol + innerCol_2);
- const int x_278 = inputRow_1;
- const int x_279 = inputCol_1;
- const int x_284 = globalCol;
- const int x_285 = innerCol_2;
- param_5 = ((t * 64) + inputRow_1);
- param_6 = (x_284 + x_285);
- const float x_289 = mm_readB_i1_i1_(param_5, param_6);
- mm_Bsub[x_278][x_279] = x_289;
- }
- }
- GroupMemoryBarrierWithGroupSync();
- k = 0;
- for(; (k < 64); k = (k + 1)) {
- inner = 0;
- for(; (inner < 1); inner = (inner + 1)) {
- const int x_314 = inner;
- const float x_320 = mm_Bsub[k][(tileCol + inner)];
- BCached[x_314] = x_320;
- }
- innerRow_3 = 0;
- for(; (innerRow_3 < 1); innerRow_3 = (innerRow_3 + 1)) {
- const float x_338 = mm_Asub[(tileRow + innerRow_3)][k];
- ACached = x_338;
- innerCol_3 = 0;
- for(; (innerCol_3 < 1); innerCol_3 = (innerCol_3 + 1)) {
- const int x_347 = innerRow_3;
- const int x_348 = innerCol_3;
- const float x_349 = ACached;
- const float x_352 = BCached[innerCol_3];
- const float x_355 = acc[x_347][x_348];
- acc[x_347][x_348] = (x_355 + (x_349 * x_352));
+ {
+ for(; (t < numTiles); t = (t + 1)) {
+ innerRow_1 = 0;
+ {
+ for(; (innerRow_1 < 1); innerRow_1 = (innerRow_1 + 1)) {
+ innerCol_1 = 0;
+ {
+ for(; (innerCol_1 < 64); innerCol_1 = (innerCol_1 + 1)) {
+ inputRow = (tileRow + innerRow_1);
+ inputCol = (tileColA + innerCol_1);
+ const int x_233 = inputRow;
+ const int x_234 = inputCol;
+ const int x_238 = t;
+ const int x_240 = inputCol;
+ param_3 = (globalRow + innerRow_1);
+ param_4 = ((x_238 * 64) + x_240);
+ const float x_244 = mm_readA_i1_i1_(param_3, param_4);
+ mm_Asub[x_233][x_234] = x_244;
+ }
+ }
}
}
+ innerRow_2 = 0;
+ {
+ for(; (innerRow_2 < 1); innerRow_2 = (innerRow_2 + 1)) {
+ innerCol_2 = 0;
+ {
+ for(; (innerCol_2 < 1); innerCol_2 = (innerCol_2 + 1)) {
+ inputRow_1 = (tileRowB + innerRow_2);
+ inputCol_1 = (tileCol + innerCol_2);
+ const int x_278 = inputRow_1;
+ const int x_279 = inputCol_1;
+ const int x_284 = globalCol;
+ const int x_285 = innerCol_2;
+ param_5 = ((t * 64) + inputRow_1);
+ param_6 = (x_284 + x_285);
+ const float x_289 = mm_readB_i1_i1_(param_5, param_6);
+ mm_Bsub[x_278][x_279] = x_289;
+ }
+ }
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ k = 0;
+ {
+ for(; (k < 64); k = (k + 1)) {
+ inner = 0;
+ {
+ for(; (inner < 1); inner = (inner + 1)) {
+ const int x_314 = inner;
+ const float x_320 = mm_Bsub[k][(tileCol + inner)];
+ BCached[x_314] = x_320;
+ }
+ }
+ innerRow_3 = 0;
+ {
+ for(; (innerRow_3 < 1); innerRow_3 = (innerRow_3 + 1)) {
+ const float x_338 = mm_Asub[(tileRow + innerRow_3)][k];
+ ACached = x_338;
+ innerCol_3 = 0;
+ {
+ for(; (innerCol_3 < 1); innerCol_3 = (innerCol_3 + 1)) {
+ const int x_347 = innerRow_3;
+ const int x_348 = innerCol_3;
+ const float x_349 = ACached;
+ const float x_352 = BCached[innerCol_3];
+ const float x_355 = acc[x_347][x_348];
+ acc[x_347][x_348] = (x_355 + (x_349 * x_352));
+ }
+ }
+ }
+ }
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
}
- GroupMemoryBarrierWithGroupSync();
}
innerRow_4 = 0;
- for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
- innerCol_4 = 0;
- while (true) {
- bool x_393 = false;
- bool x_394_phi = false;
- if ((innerCol_4 < 1)) {
- } else {
- break;
- }
- const int x_382 = globalCol;
- const int x_383 = innerCol_4;
- const int x_385 = dimBOuter;
- const bool x_386 = ((x_382 + x_383) < x_385);
- x_394_phi = x_386;
- if (x_386) {
- const int x_389 = globalRow;
- const int x_390 = innerRow_4;
- const int x_392 = dimAOuter;
- x_393 = ((x_389 + x_390) < x_392);
- x_394_phi = x_393;
- }
- if (x_394_phi) {
- const int x_400 = globalCol;
- const int x_401 = innerCol_4;
- const int x_403 = innerRow_4;
- const int x_404 = innerCol_4;
- param_7 = (globalRow + innerRow_4);
- param_8 = (x_400 + x_401);
- const float x_409 = acc[x_403][x_404];
- param_9 = x_409;
- mm_write_i1_i1_f1_(param_7, param_8, param_9);
- }
- {
- innerCol_4 = (innerCol_4 + 1);
+ {
+ for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
+ innerCol_4 = 0;
+ while (true) {
+ bool x_393 = false;
+ bool x_394_phi = false;
+ if ((innerCol_4 < 1)) {
+ } else {
+ break;
+ }
+ const int x_382 = globalCol;
+ const int x_383 = innerCol_4;
+ const int x_385 = dimBOuter;
+ const bool x_386 = ((x_382 + x_383) < x_385);
+ x_394_phi = x_386;
+ if (x_386) {
+ const int x_389 = globalRow;
+ const int x_390 = innerRow_4;
+ const int x_392 = dimAOuter;
+ x_393 = ((x_389 + x_390) < x_392);
+ x_394_phi = x_393;
+ }
+ if (x_394_phi) {
+ const int x_400 = globalCol;
+ const int x_401 = innerCol_4;
+ const int x_403 = innerRow_4;
+ const int x_404 = innerCol_4;
+ param_7 = (globalRow + innerRow_4);
+ param_8 = (x_400 + x_401);
+ const float x_409 = acc[x_403][x_404];
+ param_9 = x_409;
+ mm_write_i1_i1_f1_(param_7, param_8, param_9);
+ }
+ {
+ innerCol_4 = (innerCol_4 + 1);
+ }
}
}
}
@@ -336,14 +360,20 @@
const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
const uint local_invocation_index = tint_symbol.local_invocation_index;
if ((local_invocation_index == 0u)) {
- for(int i = 0; (i < 64); i = (i + 1)) {
- for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
- mm_Asub[i][i_1] = 0.0f;
+ {
+ for(int i = 0; (i < 64); i = (i + 1)) {
+ {
+ for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+ mm_Asub[i][i_1] = 0.0f;
+ }
+ }
}
}
- for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
- const float tint_symbol_6[1] = (float[1])0;
- mm_Bsub[i_2] = tint_symbol_6;
+ {
+ for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+ const float tint_symbol_6[1] = (float[1])0;
+ mm_Bsub[i_2] = tint_symbol_6;
+ }
}
}
GroupMemoryBarrierWithGroupSync();
diff --git a/test/bug/tint/948.wgsl.expected.hlsl b/test/bug/tint/948.wgsl.expected.hlsl
index 2454c3b..d6f3a02 100644
--- a/test/bug/tint/948.wgsl.expected.hlsl
+++ b/test/bug/tint/948.wgsl.expected.hlsl
@@ -67,89 +67,93 @@
const float2 x_111 = asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
stageUnits = (float2(1.0f, 1.0f) / x_111);
i = 0;
- for(; (i < 2); i = (i + 1)) {
- switch(i) {
- case 1: {
- const float2 x_150 = tileID;
- const uint scalar_offset_4 = (88u) / 4;
- uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
- const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
- const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
- frameID_1 = x_156.x;
- break;
- }
- case 0: {
- const float2 x_136 = tileID;
- const uint scalar_offset_5 = (88u) / 4;
- uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
- const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
- const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
- frameID_1 = x_142.x;
- break;
- }
- default: {
- break;
- }
- }
- const float x_166 = frameID_1;
- const uint scalar_offset_6 = (108u) / 4;
- const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
- const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
- animationData = x_172;
- const float x_174 = animationData.y;
- if ((x_174 > 0.0f)) {
- const uint scalar_offset_7 = (0u) / 4;
- const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
- const float x_184 = animationData.z;
- mt = ((x_181 * x_184) % 1.0f);
- f = 0.0f;
- for(; (f < 8.0f); f = (f + 1.0f)) {
- const float x_197 = animationData.y;
- if ((x_197 > mt)) {
- const float x_203 = animationData.x;
- frameID_1 = x_203;
+ {
+ for(; (i < 2); i = (i + 1)) {
+ switch(i) {
+ case 1: {
+ const float2 x_150 = tileID;
+ const uint scalar_offset_4 = (88u) / 4;
+ uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
+ const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
+ const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
+ frameID_1 = x_156.x;
break;
}
- const float x_208 = frameID_1;
- const uint scalar_offset_8 = (108u) / 4;
- const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
- const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
- animationData = x_217;
+ case 0: {
+ const float2 x_136 = tileID;
+ const uint scalar_offset_5 = (88u) / 4;
+ uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
+ const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
+ const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
+ frameID_1 = x_142.x;
+ break;
+ }
+ default: {
+ break;
+ }
}
- }
- param = (frameID_1 + 0.5f);
- const float4x4 x_225 = getFrameData_f1_(param);
- frameData = x_225;
- const float4 x_228 = frameData[0];
- const uint scalar_offset_9 = (96u) / 4;
- uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
- const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
- frameSize = (float2(x_228.w, x_228.z) / x_231);
- const float4 x_235 = frameData[0];
- offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
- const float4 x_241 = frameData[2];
- const float4 x_244 = frameData[0];
- ratio = (float2(x_241.x, x_241.y) / float2(x_244.w, x_244.z));
- const float x_248 = frameData[2].z;
- if ((x_248 == 1.0f)) {
- const float2 x_252 = tileUV;
- tileUV = float2(x_252.y, x_252.x);
- }
- if ((i == 0)) {
- const float4 x_268 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
- color = x_268;
- } else {
- const float4 x_279 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
- nc = x_279;
- const float x_283 = color.w;
- const float x_285 = nc.w;
- alpha = min((x_283 + x_285), 1.0f);
- const float4 x_290 = color;
- const float4 x_292 = nc;
- const float x_295 = nc.w;
- mixed = lerp(float3(x_290.x, x_290.y, x_290.z), float3(x_292.x, x_292.y, x_292.z), float3(x_295, x_295, x_295));
- const float3 x_298 = mixed;
- color = float4(x_298.x, x_298.y, x_298.z, alpha);
+ const float x_166 = frameID_1;
+ const uint scalar_offset_6 = (108u) / 4;
+ const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+ const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
+ animationData = x_172;
+ const float x_174 = animationData.y;
+ if ((x_174 > 0.0f)) {
+ const uint scalar_offset_7 = (0u) / 4;
+ const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+ const float x_184 = animationData.z;
+ mt = ((x_181 * x_184) % 1.0f);
+ f = 0.0f;
+ {
+ for(; (f < 8.0f); f = (f + 1.0f)) {
+ const float x_197 = animationData.y;
+ if ((x_197 > mt)) {
+ const float x_203 = animationData.x;
+ frameID_1 = x_203;
+ break;
+ }
+ const float x_208 = frameID_1;
+ const uint scalar_offset_8 = (108u) / 4;
+ const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+ const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
+ animationData = x_217;
+ }
+ }
+ }
+ param = (frameID_1 + 0.5f);
+ const float4x4 x_225 = getFrameData_f1_(param);
+ frameData = x_225;
+ const float4 x_228 = frameData[0];
+ const uint scalar_offset_9 = (96u) / 4;
+ uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
+ const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
+ frameSize = (float2(x_228.w, x_228.z) / x_231);
+ const float4 x_235 = frameData[0];
+ offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
+ const float4 x_241 = frameData[2];
+ const float4 x_244 = frameData[0];
+ ratio = (float2(x_241.x, x_241.y) / float2(x_244.w, x_244.z));
+ const float x_248 = frameData[2].z;
+ if ((x_248 == 1.0f)) {
+ const float2 x_252 = tileUV;
+ tileUV = float2(x_252.y, x_252.x);
+ }
+ if ((i == 0)) {
+ const float4 x_268 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
+ color = x_268;
+ } else {
+ const float4 x_279 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
+ nc = x_279;
+ const float x_283 = color.w;
+ const float x_285 = nc.w;
+ alpha = min((x_283 + x_285), 1.0f);
+ const float4 x_290 = color;
+ const float4 x_292 = nc;
+ const float x_295 = nc.w;
+ mixed = lerp(float3(x_290.x, x_290.y, x_290.z), float3(x_292.x, x_292.y, x_292.z), float3(x_295, x_295, x_295));
+ const float3 x_298 = mixed;
+ color = float4(x_298.x, x_298.y, x_298.z, alpha);
+ }
}
}
const uint scalar_offset_10 = (112u) / 4;
diff --git a/test/bug/tint/949.wgsl.expected.hlsl b/test/bug/tint/949.wgsl.expected.hlsl
index 926e020..2bb720a 100644
--- a/test/bug/tint/949.wgsl.expected.hlsl
+++ b/test/bug/tint/949.wgsl.expected.hlsl
@@ -249,20 +249,22 @@
lastSampledHeight = 1.0f;
currSampledHeight = 1.0f;
i = 0;
- for(; (i < 15); i = (i + 1)) {
- const float4 x_397 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + vCurrOffset));
- currSampledHeight = x_397.w;
- if ((currSampledHeight > currRayHeight)) {
- delta1 = (currSampledHeight - currRayHeight);
- delta2 = ((currRayHeight + stepSize) - lastSampledHeight);
- ratio = (delta1 / (delta1 + delta2));
- vCurrOffset = ((vLastOffset * ratio) + (vCurrOffset * (1.0f - ratio)));
- break;
- } else {
- currRayHeight = (currRayHeight - stepSize);
- vLastOffset = vCurrOffset;
- vCurrOffset = (vCurrOffset + (vMaxOffset * stepSize));
- lastSampledHeight = currSampledHeight;
+ {
+ for(; (i < 15); i = (i + 1)) {
+ const float4 x_397 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + vCurrOffset));
+ currSampledHeight = x_397.w;
+ if ((currSampledHeight > currRayHeight)) {
+ delta1 = (currSampledHeight - currRayHeight);
+ delta2 = ((currRayHeight + stepSize) - lastSampledHeight);
+ ratio = (delta1 / (delta1 + delta2));
+ vCurrOffset = ((vLastOffset * ratio) + (vCurrOffset * (1.0f - ratio)));
+ break;
+ } else {
+ currRayHeight = (currRayHeight - stepSize);
+ vLastOffset = vCurrOffset;
+ vCurrOffset = (vCurrOffset + (vMaxOffset * stepSize));
+ lastSampledHeight = currSampledHeight;
+ }
}
}
parallaxOcclusion_0 = vCurrOffset;
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
index 6496cf5..a222773 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
@@ -32,8 +32,7 @@
void foo() {
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_int3(v3i, i, 1);
Set_uint4(v4u, i, 1u);
@@ -45,8 +44,7 @@
[numthreads(1, 1, 1)]
void main() {
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
index ab45d31..2ab266f 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
@@ -2,20 +2,11 @@
using namespace metal;
void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) {
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- (*(tint_symbol_1))[i] = 1.0f;
- (*(tint_symbol_2))[i] = 1;
- (*(tint_symbol_3))[i] = 1u;
- (*(tint_symbol_4))[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ (*(tint_symbol_1))[i] = 1.0f;
+ (*(tint_symbol_2))[i] = 1;
+ (*(tint_symbol_3))[i] = 1u;
+ (*(tint_symbol_4))[i] = true;
}
}
@@ -24,17 +15,8 @@
thread int3 tint_symbol_6 = 0;
thread uint4 tint_symbol_7 = 0u;
thread bool2 tint_symbol_8 = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
index bc302f1..def10cb 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
@@ -7,37 +7,17 @@
var<private> v2b : vec2<bool>;
fn foo() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
}
}
[[stage(compute), workgroup_size(1, 1, 1)]]
fn main() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- foo();
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
index 33a99b9..ea05754 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
@@ -14,8 +14,7 @@
[numthreads(1, 1, 1)]
void main() {
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
index 2685e75..44a4fa0 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
@@ -14,17 +14,8 @@
thread int3 tint_symbol_6 = 0;
thread uint4 tint_symbol_7 = 0u;
thread bool2 tint_symbol_8 = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
index ae1667c4..21db9dd 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
@@ -16,17 +16,7 @@
[[stage(compute), workgroup_size(1, 1, 1)]]
fn main() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- foo();
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ foo();
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
index c8fa518..85ffc25 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
@@ -97,8 +97,7 @@
bool3 v3b = bool3(false, false, false);
bool4 v4b = bool4(false, false, false, false);
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_float3(v3f, i, 1.0f);
Set_float4(v4f, i, 1.0f);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
index a460250..b622316 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
@@ -14,28 +14,19 @@
bool2 v2b = false;
bool3 v3b = false;
bool4 v4b = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0f;
- v3f[i] = 1.0f;
- v4f[i] = 1.0f;
- v2i[i] = 1;
- v3i[i] = 1;
- v4i[i] = 1;
- v2u[i] = 1u;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v2b[i] = true;
- v3b[i] = true;
- v4b[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0f;
+ v3f[i] = 1.0f;
+ v4f[i] = 1.0f;
+ v2i[i] = 1;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v2u[i] = 1u;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v3b[i] = true;
+ v4b[i] = true;
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
index 7d73628..9f8f2e0 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
@@ -12,28 +12,18 @@
var v2b : vec2<bool>;
var v3b : vec3<bool>;
var v4b : vec4<bool>;
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v3f[i] = 1.0;
- v4f[i] = 1.0;
- v2i[i] = 1;
- v3i[i] = 1;
- v4i[i] = 1;
- v2u[i] = 1u;
- v3u[i] = 1u;
- v4u[i] = 1u;
- v2b[i] = true;
- v3b[i] = true;
- v4b[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v3f[i] = 1.0;
+ v4f[i] = 1.0;
+ v2i[i] = 1;
+ v3i[i] = 1;
+ v4i[i] = 1;
+ v2u[i] = 1u;
+ v3u[i] = 1u;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v3b[i] = true;
+ v4b[i] = true;
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
index cbc4e6a..9c60097 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
@@ -36,8 +36,7 @@
bool2 v2b = bool2(false, false);
bool2 v2b_2 = bool2(false, false);
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_int3(v3i, i, 1);
Set_uint4(v4u, i, 1u);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
index 9f28e5c..4c124ca 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
@@ -10,24 +10,15 @@
uint4 v4u_2 = 0u;
bool2 v2b = false;
bool2 v2b_2 = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0f;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
- v2f_2[i] = 1.0f;
- v3i_2[i] = 1;
- v4u_2[i] = 1u;
- v2b_2[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0f;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v2f_2[i] = 1.0f;
+ v3i_2[i] = 1;
+ v4u_2[i] = 1u;
+ v2b_2[i] = true;
}
return;
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
index 0229323..74950ca 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
@@ -8,24 +8,14 @@
var v4u_2 : vec4<u32>;
var v2b : vec2<bool>;
var v2b_2 : vec2<bool>;
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v3i[i] = 1;
- v4u[i] = 1u;
- v2b[i] = true;
- v2f_2[i] = 1.0;
- v3i_2[i] = 1;
- v4u_2[i] = 1u;
- v2b_2[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v3i[i] = 1;
+ v4u[i] = 1u;
+ v2b[i] = true;
+ v2f_2[i] = 1.0;
+ v3i_2[i] = 1;
+ v4u_2[i] = 1u;
+ v2b_2[i] = true;
}
}
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
index aed923f..2e6de37 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
@@ -37,8 +37,7 @@
bool3 v3b = bool3(false, false, false);
bool4 v4b = bool4(false, false, false, false);
{
- int i = 0;
- for(; !(!((i < 2))); i = (i + 1)) {
+ for(int i = 0; (i < 2); i = (i + 1)) {
Set_float2(v2f, i, 1.0f);
Set_int2(v2i, i, 1);
Set_uint2(v2u, i, 1u);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
index fb53199..2113ef6 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
@@ -14,20 +14,11 @@
bool2 v2b = false;
bool3 v3b = false;
bool4 v4b = false;
- {
- int i = 0;
- while (true) {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0f;
- v2i[i] = 1;
- v2u[i] = 1u;
- v2b[i] = true;
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0f;
+ v2i[i] = 1;
+ v2u[i] = 1u;
+ v2b[i] = true;
}
int i = 0;
v3f[i] = 1.0f;
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
index 8c54146..1fada75 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
@@ -12,21 +12,11 @@
var v2b : vec2<bool>;
var v3b : vec3<bool>;
var v4b : vec4<bool>;
- {
- var i : i32 = 0;
- loop {
- if (!((i < 2))) {
- break;
- }
- v2f[i] = 1.0;
- v2i[i] = 1;
- v2u[i] = 1u;
- v2b[i] = true;
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+ v2f[i] = 1.0;
+ v2i[i] = 1;
+ v2u[i] = 1u;
+ v2b[i] = true;
}
var i : i32 = 0;
v3f[i] = 1.0;
diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl
index d98a316..4643faa 100644
--- a/test/samples/compute_boids.wgsl.expected.hlsl
+++ b/test/samples/compute_boids.wgsl.expected.hlsl
@@ -53,8 +53,7 @@
float2 pos = float2(0.0f, 0.0f);
float2 vel = float2(0.0f, 0.0f);
{
- uint i = 0u;
- for(; !(!((i < 5u))); i = (i + 1u)) {
+ for(uint i = 0u; (i < 5u); i = (i + 1u)) {
if ((i == index)) {
continue;
}
diff --git a/test/samples/compute_boids.wgsl.expected.msl b/test/samples/compute_boids.wgsl.expected.msl
index dae2f1f..cc9a856 100644
--- a/test/samples/compute_boids.wgsl.expected.msl
+++ b/test/samples/compute_boids.wgsl.expected.msl
@@ -61,34 +61,22 @@
int cVelCount = 0;
float2 pos = 0.0f;
float2 vel = 0.0f;
- {
- uint i = 0u;
- while (true) {
- if (!((i < 5u))) {
- break;
- }
- if ((i == index)) {
- {
- i = (i + 1u);
- }
- continue;
- }
- pos = particlesA.particles.arr[i].pos.xy;
- vel = particlesA.particles.arr[i].vel.xy;
- if ((distance(pos, vPos) < params.rule1Distance)) {
- cMass = (cMass + pos);
- cMassCount = (cMassCount + 1);
- }
- if ((distance(pos, vPos) < params.rule2Distance)) {
- colVel = (colVel - (pos - vPos));
- }
- if ((distance(pos, vPos) < params.rule3Distance)) {
- cVel = (cVel + vel);
- cVelCount = (cVelCount + 1);
- }
- {
- i = (i + 1u);
- }
+ for(uint i = 0u; (i < 5u); i = (i + 1u)) {
+ if ((i == index)) {
+ continue;
+ }
+ pos = particlesA.particles.arr[i].pos.xy;
+ vel = particlesA.particles.arr[i].vel.xy;
+ if ((distance(pos, vPos) < params.rule1Distance)) {
+ cMass = (cMass + pos);
+ cMassCount = (cMassCount + 1);
+ }
+ if ((distance(pos, vPos) < params.rule2Distance)) {
+ colVel = (colVel - (pos - vPos));
+ }
+ if ((distance(pos, vPos) < params.rule3Distance)) {
+ cVel = (cVel + vel);
+ cVelCount = (cVelCount + 1);
}
}
if ((cMassCount > 0)) {
diff --git a/test/samples/compute_boids.wgsl.expected.wgsl b/test/samples/compute_boids.wgsl.expected.wgsl
index 09cb30f..dd981f4 100644
--- a/test/samples/compute_boids.wgsl.expected.wgsl
+++ b/test/samples/compute_boids.wgsl.expected.wgsl
@@ -52,32 +52,22 @@
var cVelCount : i32 = 0;
var pos : vec2<f32>;
var vel : vec2<f32>;
- {
- var i : u32 = 0u;
- loop {
- if (!((i < 5u))) {
- break;
- }
- if ((i == index)) {
- continue;
- }
- pos = particlesA.particles[i].pos.xy;
- vel = particlesA.particles[i].vel.xy;
- if ((distance(pos, vPos) < params.rule1Distance)) {
- cMass = (cMass + pos);
- cMassCount = (cMassCount + 1);
- }
- if ((distance(pos, vPos) < params.rule2Distance)) {
- colVel = (colVel - (pos - vPos));
- }
- if ((distance(pos, vPos) < params.rule3Distance)) {
- cVel = (cVel + vel);
- cVelCount = (cVelCount + 1);
- }
-
- continuing {
- i = (i + 1u);
- }
+ for(var i : u32 = 0u; (i < 5u); i = (i + 1u)) {
+ if ((i == index)) {
+ continue;
+ }
+ pos = particlesA.particles[i].pos.xy;
+ vel = particlesA.particles[i].vel.xy;
+ if ((distance(pos, vPos) < params.rule1Distance)) {
+ cMass = (cMass + pos);
+ cMassCount = (cMassCount + 1);
+ }
+ if ((distance(pos, vPos) < params.rule2Distance)) {
+ colVel = (colVel - (pos - vPos));
+ }
+ if ((distance(pos, vPos) < params.rule3Distance)) {
+ cVel = (cVel + vel);
+ cVelCount = (cVelCount + 1);
}
}
if ((cMassCount > 0)) {
diff --git a/test/statements/for/basic.wgsl.expected.hlsl b/test/statements/for/basic.wgsl.expected.hlsl
index c1e6281..f215f4b 100644
--- a/test/statements/for/basic.wgsl.expected.hlsl
+++ b/test/statements/for/basic.wgsl.expected.hlsl
@@ -8,8 +8,7 @@
void f() {
{
- int i = 0;
- for(; !(!((i < 5))); i = (i + 1)) {
+ for(int i = 0; (i < 5); i = (i + 1)) {
some_loop_body();
}
}
diff --git a/test/statements/for/basic.wgsl.expected.msl b/test/statements/for/basic.wgsl.expected.msl
index d2e498f..e36af9f 100644
--- a/test/statements/for/basic.wgsl.expected.msl
+++ b/test/statements/for/basic.wgsl.expected.msl
@@ -5,17 +5,8 @@
}
void f() {
- {
- int i = 0;
- while (true) {
- if (!((i < 5))) {
- break;
- }
- some_loop_body();
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; (i < 5); i = (i + 1)) {
+ some_loop_body();
}
}
diff --git a/test/statements/for/basic.wgsl.expected.wgsl b/test/statements/for/basic.wgsl.expected.wgsl
index 6a66161..3f9a88a 100644
--- a/test/statements/for/basic.wgsl.expected.wgsl
+++ b/test/statements/for/basic.wgsl.expected.wgsl
@@ -2,17 +2,7 @@
}
fn f() {
- {
- var i : i32 = 0;
- loop {
- if (!((i < 5))) {
- break;
- }
- some_loop_body();
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; (i < 5); i = (i + 1)) {
+ some_loop_body();
}
}
diff --git a/test/statements/for/complex.wgsl.expected.hlsl b/test/statements/for/complex.wgsl.expected.hlsl
index a3197c1..b48727f 100644
--- a/test/statements/for/complex.wgsl.expected.hlsl
+++ b/test/statements/for/complex.wgsl.expected.hlsl
@@ -15,7 +15,7 @@
if (tint_tmp) {
tint_tmp = (j < 10);
}
- if (!(!(!((tint_tmp))))) { break; }
+ if (!((tint_tmp))) { break; }
some_loop_body();
j = (i * 30);
i = (i + 1);
diff --git a/test/statements/for/complex.wgsl.expected.msl b/test/statements/for/complex.wgsl.expected.msl
index 2030618..9ed6331 100644
--- a/test/statements/for/complex.wgsl.expected.msl
+++ b/test/statements/for/complex.wgsl.expected.msl
@@ -6,18 +6,9 @@
void f() {
int j = 0;
- {
- int i = 0;
- while (true) {
- if (!(((i < 5) && (j < 10)))) {
- break;
- }
- some_loop_body();
- j = (i * 30);
- {
- i = (i + 1);
- }
- }
+ for(int i = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
+ some_loop_body();
+ j = (i * 30);
}
}
diff --git a/test/statements/for/complex.wgsl.expected.wgsl b/test/statements/for/complex.wgsl.expected.wgsl
index 1f3ce92..32bc4e0 100644
--- a/test/statements/for/complex.wgsl.expected.wgsl
+++ b/test/statements/for/complex.wgsl.expected.wgsl
@@ -3,18 +3,8 @@
fn f() {
var j : i32;
- {
- var i : i32 = 0;
- loop {
- if (!(((i < 5) && (j < 10)))) {
- break;
- }
- some_loop_body();
- j = (i * 30);
-
- continuing {
- i = (i + 1);
- }
- }
+ for(var i : i32 = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
+ some_loop_body();
+ j = (i * 30);
}
}
diff --git a/test/statements/for/condition.wgsl.expected.hlsl b/test/statements/for/condition.wgsl.expected.hlsl
index 5f51493..e4456e0 100644
--- a/test/statements/for/condition.wgsl.expected.hlsl
+++ b/test/statements/for/condition.wgsl.expected.hlsl
@@ -5,6 +5,8 @@
void f() {
int i = 0;
- for(; !(!((i < 4))); ) {
+ {
+ for(; (i < 4); ) {
+ }
}
}
diff --git a/test/statements/for/condition.wgsl.expected.msl b/test/statements/for/condition.wgsl.expected.msl
index c99148a..839541b 100644
--- a/test/statements/for/condition.wgsl.expected.msl
+++ b/test/statements/for/condition.wgsl.expected.msl
@@ -3,10 +3,7 @@
using namespace metal;
void f() {
int i = 0;
- while (true) {
- if (!((i < 4))) {
- break;
- }
+ for(; (i < 4); ) {
}
}
diff --git a/test/statements/for/condition.wgsl.expected.wgsl b/test/statements/for/condition.wgsl.expected.wgsl
index 14f48d5..7fc3c6f 100644
--- a/test/statements/for/condition.wgsl.expected.wgsl
+++ b/test/statements/for/condition.wgsl.expected.wgsl
@@ -1,8 +1,5 @@
fn f() {
var i : i32;
- loop {
- if (!((i < 4))) {
- break;
- }
+ for(; (i < 4); ;) {
}
}
diff --git a/test/statements/for/continuing.wgsl.expected.hlsl b/test/statements/for/continuing.wgsl.expected.hlsl
index 28b8b3b..b3ec736 100644
--- a/test/statements/for/continuing.wgsl.expected.hlsl
+++ b/test/statements/for/continuing.wgsl.expected.hlsl
@@ -5,9 +5,8 @@
void f() {
int i = 0;
- while (true) {
- {
- i = (i + 1);
+ {
+ for(; ; i = (i + 1)) {
}
}
}
diff --git a/test/statements/for/continuing.wgsl.expected.msl b/test/statements/for/continuing.wgsl.expected.msl
index 67a069f..a288b5d 100644
--- a/test/statements/for/continuing.wgsl.expected.msl
+++ b/test/statements/for/continuing.wgsl.expected.msl
@@ -3,10 +3,7 @@
using namespace metal;
void f() {
int i = 0;
- while (true) {
- {
- i = (i + 1);
- }
+ for(; ; i = (i + 1)) {
}
}
diff --git a/test/statements/for/continuing.wgsl.expected.wgsl b/test/statements/for/continuing.wgsl.expected.wgsl
index f3faf03..c76b514 100644
--- a/test/statements/for/continuing.wgsl.expected.wgsl
+++ b/test/statements/for/continuing.wgsl.expected.wgsl
@@ -1,9 +1,5 @@
fn f() {
var i : i32;
- loop {
-
- continuing {
- i = (i + 1);
- }
+ for(; ; i = (i + 1)) {
}
}
diff --git a/test/statements/for/empty.wgsl.expected.hlsl b/test/statements/for/empty.wgsl.expected.hlsl
index b642595..cd85e1a 100644
--- a/test/statements/for/empty.wgsl.expected.hlsl
+++ b/test/statements/for/empty.wgsl.expected.hlsl
@@ -4,6 +4,8 @@
}
void f() {
- while (true) {
+ {
+ for(; ; ) {
+ }
}
}
diff --git a/test/statements/for/empty.wgsl.expected.msl b/test/statements/for/empty.wgsl.expected.msl
index 622427b..e716521 100644
--- a/test/statements/for/empty.wgsl.expected.msl
+++ b/test/statements/for/empty.wgsl.expected.msl
@@ -2,7 +2,7 @@
using namespace metal;
void f() {
- while (true) {
+ for(; ; ) {
}
}
diff --git a/test/statements/for/empty.wgsl.expected.wgsl b/test/statements/for/empty.wgsl.expected.wgsl
index 3674ab1..bb4e48b 100644
--- a/test/statements/for/empty.wgsl.expected.wgsl
+++ b/test/statements/for/empty.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
fn f() {
- loop {
+ for(; ; ;) {
}
}
diff --git a/test/statements/for/initializer.wgsl.expected.hlsl b/test/statements/for/initializer.wgsl.expected.hlsl
index 8311c7b..e614afd 100644
--- a/test/statements/for/initializer.wgsl.expected.hlsl
+++ b/test/statements/for/initializer.wgsl.expected.hlsl
@@ -5,8 +5,7 @@
void f() {
{
- int i = 0;
- while (true) {
+ for(int i = 0; ; ) {
}
}
}
diff --git a/test/statements/for/initializer.wgsl.expected.msl b/test/statements/for/initializer.wgsl.expected.msl
index 70c5111..90082d6 100644
--- a/test/statements/for/initializer.wgsl.expected.msl
+++ b/test/statements/for/initializer.wgsl.expected.msl
@@ -2,10 +2,7 @@
using namespace metal;
void f() {
- {
- int i = 0;
- while (true) {
- }
+ for(int i = 0; ; ) {
}
}
diff --git a/test/statements/for/initializer.wgsl.expected.wgsl b/test/statements/for/initializer.wgsl.expected.wgsl
index e021556..ef9b991 100644
--- a/test/statements/for/initializer.wgsl.expected.wgsl
+++ b/test/statements/for/initializer.wgsl.expected.wgsl
@@ -1,7 +1,4 @@
fn f() {
- {
- var i : i32 = 0;
- loop {
- }
+ for(var i : i32 = 0; ; ;) {
}
}
diff --git a/test/statements/for/scoping.wgsl.expected.hlsl b/test/statements/for/scoping.wgsl.expected.hlsl
index 459eb43..143b1e3 100644
--- a/test/statements/for/scoping.wgsl.expected.hlsl
+++ b/test/statements/for/scoping.wgsl.expected.hlsl
@@ -5,8 +5,7 @@
void f() {
{
- int must_not_collide = 0;
- while (true) {
+ for(int must_not_collide = 0; ; ) {
}
}
int must_not_collide = 0;
diff --git a/test/statements/for/scoping.wgsl.expected.msl b/test/statements/for/scoping.wgsl.expected.msl
index d5c3f3a..0572921 100644
--- a/test/statements/for/scoping.wgsl.expected.msl
+++ b/test/statements/for/scoping.wgsl.expected.msl
@@ -2,10 +2,7 @@
using namespace metal;
void f() {
- {
- int must_not_collide = 0;
- while (true) {
- }
+ for(int must_not_collide = 0; ; ) {
}
int must_not_collide = 0;
}
diff --git a/test/statements/for/scoping.wgsl.expected.wgsl b/test/statements/for/scoping.wgsl.expected.wgsl
index ec202ae..0cac8520 100644
--- a/test/statements/for/scoping.wgsl.expected.wgsl
+++ b/test/statements/for/scoping.wgsl.expected.wgsl
@@ -1,8 +1,5 @@
fn f() {
- {
- var must_not_collide : i32 = 0;
- loop {
- }
+ for(var must_not_collide : i32 = 0; ; ;) {
}
var must_not_collide : i32;
}
