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>
This commit is contained in:
parent
03c8393213
commit
1b03f0a07a
|
@ -2071,7 +2071,7 @@ Expect<std::unique_ptr<ForHeader>> ParserImpl::expect_for_header() {
|
|||
|
||||
// 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 @@ Maybe<ast::Statement*> ParserImpl::for_stmt() {
|
|||
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
|
||||
|
|
|
@ -554,7 +554,7 @@ class ParserImpl {
|
|||
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();
|
||||
|
|
|
@ -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 {
|
||||
|
|
|
@ -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 @@ bool Resolver::ForLoopStatement(ast::ForLoopStatement* stmt) {
|
|||
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 @@ bool Resolver::BlockScope(const ast::BlockStatement* block, F&& callback) {
|
|||
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,
|
||||
|
|
|
@ -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 @@ Output Spirv::Run(const Program* in, const DataMap& data) {
|
|||
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>();
|
||||
|
|
|
@ -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::EmitLoop(ast::LoopStatement* stmt) {
|
|||
}
|
||||
|
||||
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 @@ bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
|
|||
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 @@ bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
|
|||
}
|
||||
}
|
||||
|
||||
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,8 +2630,11 @@ bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
|
|||
|
||||
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);
|
||||
|
@ -2635,8 +2648,6 @@ bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
|
|||
if (!emit_continuing()) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
line() << "}";
|
||||
} else {
|
||||
// For-loop can be generated.
|
||||
{
|
||||
|
@ -2666,12 +2677,6 @@ bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
|
|||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
line() << "}";
|
||||
}
|
||||
|
||||
if (multi_stmt_init) {
|
||||
decrement_indent();
|
||||
line() << "}";
|
||||
}
|
||||
|
||||
|
|
|
@ -161,9 +161,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoop) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( for(; ; ) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
for(; ; ) {
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
@ -180,9 +182,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleInit) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( for(int i = 0; ; ) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
for(int i = 0; ; ) {
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
@ -227,9 +231,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleCond) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( for(; true; ) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
for(; true; ) {
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
@ -248,7 +254,8 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtCond) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
while (true) {
|
||||
bool tint_tmp = true;
|
||||
if (tint_tmp) {
|
||||
tint_tmp = false;
|
||||
|
@ -256,6 +263,7 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtCond) {
|
|||
if (!((tint_tmp))) { break; }
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
@ -273,9 +281,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleCont) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( for(; ; i = (i + 1)) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
for(; ; i = (i + 1)) {
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
@ -295,7 +305,8 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtCont) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
while (true) {
|
||||
return;
|
||||
bool tint_tmp = true;
|
||||
if (tint_tmp) {
|
||||
|
@ -303,6 +314,7 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithMultiStmtCont) {
|
|||
}
|
||||
i = (tint_tmp);
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
@ -320,9 +332,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_ForLoopWithSimpleInitCondCont) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( for(int i = 0; true; i = (i + 1)) {
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
for(int i = 0; true; i = (i + 1)) {
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
|
|
|
@ -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 @@ bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
|
|||
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 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
|
|||
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);
|
||||
}
|
||||
|
|
|
@ -170,6 +170,10 @@ class GeneratorImpl : public TextGenerator {
|
|||
/// @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
|
||||
|
|
|
@ -141,6 +141,227 @@ TEST_F(MslGeneratorImplTest, Emit_LoopWithVarUsedInContinuing) {
|
|||
)");
|
||||
}
|
||||
|
||||
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
|
||||
|
|
|
@ -43,7 +43,7 @@ void main(tint_symbol_2 tint_symbol_1) {
|
|||
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) {
|
||||
|
|
|
@ -27,18 +27,9 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture
|
|||
bool success = true;
|
||||
uint4 srcColorBits = 0u;
|
||||
uint4 dstColorBits = uint4(dstColor);
|
||||
{
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
if (!((i < uniforms.channelCount))) {
|
||||
break;
|
||||
}
|
||||
for(uint i = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
|
||||
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
|
||||
success = (success && (srcColorBits[i] == dstColorBits[i]));
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
|
|
|
@ -36,19 +36,9 @@ fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
|||
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;
|
||||
}
|
||||
for(var i : u32 = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
|
||||
srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
|
||||
success = (success && (srcColorBits[i] == dstColorBits[i]));
|
||||
|
||||
continuing {
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x);
|
||||
if (success) {
|
||||
|
|
|
@ -19,8 +19,7 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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))));
|
||||
|
|
|
@ -15,19 +15,10 @@ kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Un
|
|||
uint const dimInner = uniforms.aShape.y;
|
||||
uint const dimOutter = uniforms.outShape.y;
|
||||
uint result = 0u;
|
||||
{
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
if (!((i < dimInner))) {
|
||||
break;
|
||||
}
|
||||
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]));
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
uint const index = (resultCell.y + (resultCell.x * dimOutter));
|
||||
resultMatrix.numbers[index] = result;
|
||||
|
|
|
@ -24,20 +24,10 @@ fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
|
|||
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;
|
||||
}
|
||||
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]));
|
||||
|
||||
continuing {
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
let index : u32 = (resultCell.y + (resultCell.x * dimOutter));
|
||||
resultMatrix.numbers[index] = result;
|
||||
|
|
|
@ -16,8 +16,7 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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));
|
||||
}
|
||||
}
|
||||
|
|
|
@ -12,17 +12,8 @@ kernel void tint_symbol(texture2d_array<float, access::sample> tint_symbol_2 [[t
|
|||
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;
|
||||
}
|
||||
for(uint i = 0u; (i < 1u); i = (i + 1u)) {
|
||||
result.values[(flatIndex + i)] = texel.r;
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -19,17 +19,7 @@ fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
|
|||
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;
|
||||
}
|
||||
for(var i : u32 = 0u; (i < 1u); i = (i + 1u)) {
|
||||
result.values[(flatIndex + i)] = texel.r;
|
||||
|
||||
continuing {
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -69,17 +69,25 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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_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;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
const uint tileRow = (local_id.y * RowPerThread);
|
||||
const uint tileCol = (local_id.x * ColPerThread);
|
||||
|
@ -91,8 +99,7 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
}
|
||||
}
|
||||
{
|
||||
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 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
}
|
||||
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 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
}
|
||||
}
|
||||
{
|
||||
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]);
|
||||
}
|
||||
|
|
|
@ -68,155 +68,47 @@ kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3
|
|||
tint_array_wrapper_2 acc = {};
|
||||
float ACached = 0.0f;
|
||||
tint_array_wrapper_3 BCached = {};
|
||||
{
|
||||
uint index = 0u;
|
||||
while (true) {
|
||||
if (!((index < (RowPerThread * ColPerThread)))) {
|
||||
break;
|
||||
}
|
||||
for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
acc.arr[index] = 0.0f;
|
||||
{
|
||||
index = (index + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
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;
|
||||
}
|
||||
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));
|
||||
{
|
||||
innerCol = (innerCol + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
innerRow = (innerRow + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
uint innerRow = 0u;
|
||||
while (true) {
|
||||
if (!((innerRow < RowPerThreadB))) {
|
||||
break;
|
||||
}
|
||||
{
|
||||
uint innerCol = 0u;
|
||||
while (true) {
|
||||
if (!((innerCol < ColPerThread))) {
|
||||
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));
|
||||
{
|
||||
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;
|
||||
}
|
||||
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)];
|
||||
{
|
||||
inner = (inner + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
uint innerRow = 0u;
|
||||
while (true) {
|
||||
if (!((innerRow < RowPerThread))) {
|
||||
break;
|
||||
}
|
||||
for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
|
||||
{
|
||||
uint innerCol = 0u;
|
||||
while (true) {
|
||||
if (!((innerCol < ColPerThread))) {
|
||||
break;
|
||||
}
|
||||
for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
uint innerRow = 0u;
|
||||
while (true) {
|
||||
if (!((innerRow < RowPerThread))) {
|
||||
break;
|
||||
}
|
||||
{
|
||||
uint innerCol = 0u;
|
||||
while (true) {
|
||||
if (!((innerCol < ColPerThread))) {
|
||||
break;
|
||||
}
|
||||
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]);
|
||||
{
|
||||
innerCol = (innerCol + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
innerRow = (innerRow + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
|
|
|
@ -65,167 +65,47 @@ fn main([[builtin(local_invocation_id)]] local_id : vec3<u32>, [[builtin(global_
|
|||
var acc : array<f32, 16>;
|
||||
var ACached : f32;
|
||||
var BCached : array<f32, 4>;
|
||||
{
|
||||
var index : u32 = 0u;
|
||||
loop {
|
||||
if (!((index < (RowPerThread * ColPerThread)))) {
|
||||
break;
|
||||
}
|
||||
for(var index : u32 = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
|
||||
acc[index] = 0.0;
|
||||
|
||||
continuing {
|
||||
index = (index + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
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;
|
||||
}
|
||||
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));
|
||||
|
||||
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;
|
||||
}
|
||||
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));
|
||||
|
||||
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;
|
||||
}
|
||||
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)];
|
||||
|
||||
continuing {
|
||||
inner = (inner + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
var innerRow : u32 = 0u;
|
||||
loop {
|
||||
if (!((innerRow < RowPerThread))) {
|
||||
break;
|
||||
}
|
||||
for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
|
||||
ACached = mm_Asub[(tileRow + innerRow)][k];
|
||||
{
|
||||
var innerCol : u32 = 0u;
|
||||
loop {
|
||||
if (!((innerCol < ColPerThread))) {
|
||||
break;
|
||||
}
|
||||
for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
var innerRow : u32 = 0u;
|
||||
loop {
|
||||
if (!((innerRow < RowPerThread))) {
|
||||
break;
|
||||
}
|
||||
{
|
||||
var innerCol : u32 = 0u;
|
||||
loop {
|
||||
if (!((innerCol < ColPerThread))) {
|
||||
break;
|
||||
}
|
||||
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]);
|
||||
|
||||
continuing {
|
||||
innerCol = (innerCol + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
continuing {
|
||||
innerRow = (innerRow + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -22,12 +22,16 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
const uint scalar_offset = (0u) / 4;
|
||||
const uint filterOffset = ((params[scalar_offset / 4][scalar_offset % 4] - 1u) / 2u);
|
||||
|
@ -37,11 +41,9 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
}
|
||||
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 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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]));
|
||||
|
|
|
@ -25,46 +25,18 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_4 [[texture
|
|||
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;
|
||||
}
|
||||
{
|
||||
uint c = 0u;
|
||||
while (true) {
|
||||
if (!((c < 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;
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
{
|
||||
uint r = 0u;
|
||||
while (true) {
|
||||
if (!((r < 4u))) {
|
||||
break;
|
||||
}
|
||||
{
|
||||
uint c = 0u;
|
||||
while (true) {
|
||||
if (!((c < 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;
|
||||
|
@ -72,29 +44,12 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_4 [[texture
|
|||
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;
|
||||
}
|
||||
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]));
|
||||
{
|
||||
f = (f + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
|
||||
}
|
||||
{
|
||||
c = (c + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
r = (r + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
|
|
|
@ -26,48 +26,18 @@ fn main([[builtin(workgroup_id)]] WorkGroupID : vec3<u32>, [[builtin(local_invoc
|
|||
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;
|
||||
}
|
||||
{
|
||||
var c : u32 = 0u;
|
||||
loop {
|
||||
if (!((c < 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;
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
workgroupBarrier();
|
||||
{
|
||||
var r : u32 = 0u;
|
||||
loop {
|
||||
if (!((r < 4u))) {
|
||||
break;
|
||||
}
|
||||
{
|
||||
var c : u32 = 0u;
|
||||
loop {
|
||||
if (!((c < 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;
|
||||
|
@ -75,32 +45,12 @@ fn main([[builtin(workgroup_id)]] WorkGroupID : vec3<u32>, [[builtin(local_invoc
|
|||
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;
|
||||
}
|
||||
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 {
|
||||
f = (f + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
|
||||
}
|
||||
|
||||
continuing {
|
||||
c = (c + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
continuing {
|
||||
r = (r + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -191,21 +191,28 @@ void mm_matMul_i1_i1_i1_(inout int dimAOuter, inout int dimInner, inout int dimB
|
|||
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;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
const uint x_187 = gl_LocalInvocationID.x;
|
||||
tileColA = (asint(x_187) * 64);
|
||||
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);
|
||||
|
@ -219,9 +226,13 @@ void mm_matMul_i1_i1_i1_(inout int dimAOuter, inout int dimInner, inout int dimB
|
|||
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);
|
||||
|
@ -235,20 +246,27 @@ void mm_matMul_i1_i1_i1_(inout int dimAOuter, inout int dimInner, inout int dimB
|
|||
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;
|
||||
|
@ -259,9 +277,14 @@ void mm_matMul_i1_i1_i1_(inout int dimAOuter, inout int dimInner, inout int dimB
|
|||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
GroupMemoryBarrierWithGroupSync();
|
||||
}
|
||||
}
|
||||
innerRow_4 = 0;
|
||||
{
|
||||
for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
|
||||
innerCol_4 = 0;
|
||||
while (true) {
|
||||
|
@ -299,6 +322,7 @@ void mm_matMul_i1_i1_i1_(inout int dimAOuter, inout int dimInner, inout int dimB
|
|||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
@ -336,16 +360,22 @@ void main(tint_symbol_1 tint_symbol) {
|
|||
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_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();
|
||||
gl_LocalInvocationID = gl_LocalInvocationID_param;
|
||||
gl_GlobalInvocationID = gl_GlobalInvocationID_param;
|
||||
|
|
|
@ -67,6 +67,7 @@ void main_1() {
|
|||
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: {
|
||||
|
@ -103,6 +104,7 @@ void main_1() {
|
|||
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)) {
|
||||
|
@ -117,6 +119,7 @@ void main_1() {
|
|||
animationData = x_217;
|
||||
}
|
||||
}
|
||||
}
|
||||
param = (frameID_1 + 0.5f);
|
||||
const float4x4 x_225 = getFrameData_f1_(param);
|
||||
frameData = x_225;
|
||||
|
@ -152,6 +155,7 @@ void main_1() {
|
|||
color = float4(x_298.x, x_298.y, x_298.z, alpha);
|
||||
}
|
||||
}
|
||||
}
|
||||
const uint scalar_offset_10 = (112u) / 4;
|
||||
const float3 x_310 = asfloat(x_20[scalar_offset_10 / 4].xyz);
|
||||
const float4 x_311 = color;
|
||||
|
|
|
@ -249,6 +249,7 @@ void main_1() {
|
|||
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;
|
||||
|
@ -265,6 +266,7 @@ void main_1() {
|
|||
lastSampledHeight = currSampledHeight;
|
||||
}
|
||||
}
|
||||
}
|
||||
parallaxOcclusion_0 = vCurrOffset;
|
||||
uvOffset = parallaxOcclusion_0;
|
||||
const float4 x_452 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + uvOffset));
|
||||
|
|
|
@ -32,8 +32,7 @@ static bool2 v2b = bool2(false, false);
|
|||
|
||||
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 @@ void foo() {
|
|||
[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();
|
||||
}
|
||||
}
|
||||
|
|
|
@ -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;
|
||||
}
|
||||
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;
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
@ -24,17 +15,8 @@ kernel void tint_symbol() {
|
|||
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;
|
||||
}
|
||||
for(int i = 0; (i < 2); i = (i + 1)) {
|
||||
foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -7,37 +7,17 @@ var<private> v4u : vec4<u32>;
|
|||
var<private> v2b : vec2<bool>;
|
||||
|
||||
fn foo() {
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0;
|
||||
v3i[i] = 1;
|
||||
v4u[i] = 1u;
|
||||
v2b[i] = true;
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
[[stage(compute), workgroup_size(1, 1, 1)]]
|
||||
fn main() {
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 2); i = (i + 1)) {
|
||||
foo();
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -14,8 +14,7 @@ void foo() {
|
|||
[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();
|
||||
}
|
||||
}
|
||||
|
|
|
@ -14,17 +14,8 @@ kernel void tint_symbol() {
|
|||
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;
|
||||
}
|
||||
for(int i = 0; (i < 2); i = (i + 1)) {
|
||||
foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -16,17 +16,7 @@ fn foo() {
|
|||
|
||||
[[stage(compute), workgroup_size(1, 1, 1)]]
|
||||
fn main() {
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 2); i = (i + 1)) {
|
||||
foo();
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -97,8 +97,7 @@ void main() {
|
|||
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);
|
||||
|
|
|
@ -14,12 +14,7 @@ kernel void tint_symbol() {
|
|||
bool2 v2b = false;
|
||||
bool3 v3b = false;
|
||||
bool4 v4b = false;
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(int i = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0f;
|
||||
v3f[i] = 1.0f;
|
||||
v4f[i] = 1.0f;
|
||||
|
@ -32,10 +27,6 @@ kernel void tint_symbol() {
|
|||
v2b[i] = true;
|
||||
v3b[i] = true;
|
||||
v4b[i] = true;
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -12,12 +12,7 @@ fn main() {
|
|||
var v2b : vec2<bool>;
|
||||
var v3b : vec3<bool>;
|
||||
var v4b : vec4<bool>;
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0;
|
||||
v3f[i] = 1.0;
|
||||
v4f[i] = 1.0;
|
||||
|
@ -30,10 +25,5 @@ fn main() {
|
|||
v2b[i] = true;
|
||||
v3b[i] = true;
|
||||
v4b[i] = true;
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -36,8 +36,7 @@ void main() {
|
|||
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);
|
||||
|
|
|
@ -10,12 +10,7 @@ kernel void tint_symbol() {
|
|||
uint4 v4u_2 = 0u;
|
||||
bool2 v2b = false;
|
||||
bool2 v2b_2 = false;
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(int i = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0f;
|
||||
v3i[i] = 1;
|
||||
v4u[i] = 1u;
|
||||
|
@ -24,10 +19,6 @@ kernel void tint_symbol() {
|
|||
v3i_2[i] = 1;
|
||||
v4u_2[i] = 1u;
|
||||
v2b_2[i] = true;
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -8,12 +8,7 @@ fn main() {
|
|||
var v4u_2 : vec4<u32>;
|
||||
var v2b : vec2<bool>;
|
||||
var v2b_2 : vec2<bool>;
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0;
|
||||
v3i[i] = 1;
|
||||
v4u[i] = 1u;
|
||||
|
@ -22,10 +17,5 @@ fn main() {
|
|||
v3i_2[i] = 1;
|
||||
v4u_2[i] = 1u;
|
||||
v2b_2[i] = true;
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -37,8 +37,7 @@ void main() {
|
|||
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);
|
||||
|
|
|
@ -14,20 +14,11 @@ kernel void tint_symbol() {
|
|||
bool2 v2b = false;
|
||||
bool3 v3b = false;
|
||||
bool4 v4b = false;
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(int i = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0f;
|
||||
v2i[i] = 1;
|
||||
v2u[i] = 1u;
|
||||
v2b[i] = true;
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
int i = 0;
|
||||
v3f[i] = 1.0f;
|
||||
|
|
|
@ -12,21 +12,11 @@ fn main() {
|
|||
var v2b : vec2<bool>;
|
||||
var v3b : vec3<bool>;
|
||||
var v4b : vec4<bool>;
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 2))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 2); i = (i + 1)) {
|
||||
v2f[i] = 1.0;
|
||||
v2i[i] = 1;
|
||||
v2u[i] = 1u;
|
||||
v2b[i] = true;
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
var i : i32 = 0;
|
||||
v3f[i] = 1.0;
|
||||
|
|
|
@ -53,8 +53,7 @@ void comp_main(tint_symbol_5 tint_symbol_4) {
|
|||
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;
|
||||
}
|
||||
|
|
|
@ -61,16 +61,8 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
|
|||
int cVelCount = 0;
|
||||
float2 pos = 0.0f;
|
||||
float2 vel = 0.0f;
|
||||
{
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
if (!((i < 5u))) {
|
||||
break;
|
||||
}
|
||||
for(uint i = 0u; (i < 5u); i = (i + 1u)) {
|
||||
if ((i == index)) {
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
continue;
|
||||
}
|
||||
pos = particlesA.particles.arr[i].pos.xy;
|
||||
|
@ -86,10 +78,6 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
|
|||
cVel = (cVel + vel);
|
||||
cVelCount = (cVelCount + 1);
|
||||
}
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
if ((cMassCount > 0)) {
|
||||
cMass = ((cMass / float2(float(cMassCount), float(cMassCount))) - vPos);
|
||||
|
|
|
@ -52,12 +52,7 @@ fn comp_main([[builtin(global_invocation_id)]] gl_GlobalInvocationID : vec3<u32>
|
|||
var cVelCount : i32 = 0;
|
||||
var pos : vec2<f32>;
|
||||
var vel : vec2<f32>;
|
||||
{
|
||||
var i : u32 = 0u;
|
||||
loop {
|
||||
if (!((i < 5u))) {
|
||||
break;
|
||||
}
|
||||
for(var i : u32 = 0u; (i < 5u); i = (i + 1u)) {
|
||||
if ((i == index)) {
|
||||
continue;
|
||||
}
|
||||
|
@ -74,11 +69,6 @@ fn comp_main([[builtin(global_invocation_id)]] gl_GlobalInvocationID : vec3<u32>
|
|||
cVel = (cVel + vel);
|
||||
cVelCount = (cVelCount + 1);
|
||||
}
|
||||
|
||||
continuing {
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
if ((cMassCount > 0)) {
|
||||
cMass = ((cMass / vec2<f32>(f32(cMassCount), f32(cMassCount))) - vPos);
|
||||
|
|
|
@ -8,8 +8,7 @@ void some_loop_body() {
|
|||
|
||||
void f() {
|
||||
{
|
||||
int i = 0;
|
||||
for(; !(!((i < 5))); i = (i + 1)) {
|
||||
for(int i = 0; (i < 5); i = (i + 1)) {
|
||||
some_loop_body();
|
||||
}
|
||||
}
|
||||
|
|
|
@ -5,17 +5,8 @@ void some_loop_body() {
|
|||
}
|
||||
|
||||
void f() {
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
if (!((i < 5))) {
|
||||
break;
|
||||
}
|
||||
for(int i = 0; (i < 5); i = (i + 1)) {
|
||||
some_loop_body();
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -2,17 +2,7 @@ fn some_loop_body() {
|
|||
}
|
||||
|
||||
fn f() {
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!((i < 5))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; (i < 5); i = (i + 1)) {
|
||||
some_loop_body();
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -15,7 +15,7 @@ void f() {
|
|||
if (tint_tmp) {
|
||||
tint_tmp = (j < 10);
|
||||
}
|
||||
if (!(!(!((tint_tmp))))) { break; }
|
||||
if (!((tint_tmp))) { break; }
|
||||
some_loop_body();
|
||||
j = (i * 30);
|
||||
i = (i + 1);
|
||||
|
|
|
@ -6,18 +6,9 @@ void some_loop_body() {
|
|||
|
||||
void f() {
|
||||
int j = 0;
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
if (!(((i < 5) && (j < 10)))) {
|
||||
break;
|
||||
}
|
||||
for(int i = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
|
||||
some_loop_body();
|
||||
j = (i * 30);
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -3,18 +3,8 @@ fn some_loop_body() {
|
|||
|
||||
fn f() {
|
||||
var j : i32;
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
if (!(((i < 5) && (j < 10)))) {
|
||||
break;
|
||||
}
|
||||
for(var i : i32 = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
|
||||
some_loop_body();
|
||||
j = (i * 30);
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -5,6 +5,8 @@ void unused_entry_point() {
|
|||
|
||||
void f() {
|
||||
int i = 0;
|
||||
for(; !(!((i < 4))); ) {
|
||||
{
|
||||
for(; (i < 4); ) {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -3,10 +3,7 @@
|
|||
using namespace metal;
|
||||
void f() {
|
||||
int i = 0;
|
||||
while (true) {
|
||||
if (!((i < 4))) {
|
||||
break;
|
||||
}
|
||||
for(; (i < 4); ) {
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -1,8 +1,5 @@
|
|||
fn f() {
|
||||
var i : i32;
|
||||
loop {
|
||||
if (!((i < 4))) {
|
||||
break;
|
||||
}
|
||||
for(; (i < 4); ;) {
|
||||
}
|
||||
}
|
||||
|
|
|
@ -5,9 +5,8 @@ void unused_entry_point() {
|
|||
|
||||
void f() {
|
||||
int i = 0;
|
||||
while (true) {
|
||||
{
|
||||
i = (i + 1);
|
||||
for(; ; i = (i + 1)) {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -3,10 +3,7 @@
|
|||
using namespace metal;
|
||||
void f() {
|
||||
int i = 0;
|
||||
while (true) {
|
||||
{
|
||||
i = (i + 1);
|
||||
}
|
||||
for(; ; i = (i + 1)) {
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -1,9 +1,5 @@
|
|||
fn f() {
|
||||
var i : i32;
|
||||
loop {
|
||||
|
||||
continuing {
|
||||
i = (i + 1);
|
||||
}
|
||||
for(; ; i = (i + 1)) {
|
||||
}
|
||||
}
|
||||
|
|
|
@ -4,6 +4,8 @@ void unused_entry_point() {
|
|||
}
|
||||
|
||||
void f() {
|
||||
while (true) {
|
||||
{
|
||||
for(; ; ) {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -2,7 +2,7 @@
|
|||
|
||||
using namespace metal;
|
||||
void f() {
|
||||
while (true) {
|
||||
for(; ; ) {
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -1,4 +1,4 @@
|
|||
fn f() {
|
||||
loop {
|
||||
for(; ; ;) {
|
||||
}
|
||||
}
|
||||
|
|
|
@ -5,8 +5,7 @@ void unused_entry_point() {
|
|||
|
||||
void f() {
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
for(int i = 0; ; ) {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -2,10 +2,7 @@
|
|||
|
||||
using namespace metal;
|
||||
void f() {
|
||||
{
|
||||
int i = 0;
|
||||
while (true) {
|
||||
}
|
||||
for(int i = 0; ; ) {
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -1,7 +1,4 @@
|
|||
fn f() {
|
||||
{
|
||||
var i : i32 = 0;
|
||||
loop {
|
||||
}
|
||||
for(var i : i32 = 0; ; ;) {
|
||||
}
|
||||
}
|
||||
|
|
|
@ -5,8 +5,7 @@ void unused_entry_point() {
|
|||
|
||||
void f() {
|
||||
{
|
||||
int must_not_collide = 0;
|
||||
while (true) {
|
||||
for(int must_not_collide = 0; ; ) {
|
||||
}
|
||||
}
|
||||
int must_not_collide = 0;
|
||||
|
|
|
@ -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;
|
||||
}
|
||||
|
|
|
@ -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;
|
||||
}
|
||||
|
|
Loading…
Reference in New Issue