writer/msl: Fix continuing block emission
Inline the `continuing` block in the places where `continue` is called. Simplifies the emission, and fixes emission of `let` statements in the loop. This fix matches the same approach in writer/hlsl. See: https://dawn-review.googlesource.com/c/tint/+/51784 Fixed: tint:833 Fixed: tint:914 Change-Id: If4d8cde62dfaf8efa24272854ca7ff5edc0a8234 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55341 Commit-Queue: Ben Clayton <bclayton@chromium.org> Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Ben Clayton
Tint LUCI CQ
parent
c15baf695d
commit
663271dca4
|
|
@ -50,6 +50,7 @@
|
|||
#include "src/sem/variable.h"
|
||||
#include "src/sem/vector_type.h"
|
||||
#include "src/sem/void_type.h"
|
||||
#include "src/utils/scoped_assignment.h"
|
||||
#include "src/writer/float_to_string.h"
|
||||
|
||||
namespace tint {
|
||||
|
|
@ -875,6 +876,9 @@ bool GeneratorImpl::EmitConstructor(ast::ConstructorExpression* expr) {
|
|||
}
|
||||
|
||||
bool GeneratorImpl::EmitContinue(ast::ContinueStatement*) {
|
||||
if (!emit_continuing_()) {
|
||||
return false;
|
||||
}
|
||||
make_indent();
|
||||
out_ << "continue;" << std::endl;
|
||||
return true;
|
||||
|
|
@ -1279,91 +1283,30 @@ bool GeneratorImpl::EmitIdentifier(ast::IdentifierExpression* expr) {
|
|||
}
|
||||
|
||||
bool GeneratorImpl::EmitLoop(ast::LoopStatement* stmt) {
|
||||
loop_emission_counter_++;
|
||||
|
||||
std::string guard =
|
||||
"tint_msl_is_first_" + std::to_string(loop_emission_counter_);
|
||||
|
||||
if (stmt->has_continuing()) {
|
||||
make_indent();
|
||||
|
||||
// Continuing variables get their own scope.
|
||||
out_ << "{" << std::endl;
|
||||
increment_indent();
|
||||
|
||||
make_indent();
|
||||
out_ << "bool " << guard << " = true;" << std::endl;
|
||||
|
||||
// A continuing block may use variables declared in the method body. As a
|
||||
// first pass, if we have a continuing, we pull all declarations outside
|
||||
// the for loop into the continuing scope. Then, the variable declarations
|
||||
// will be turned into assignments.
|
||||
for (auto* s : *(stmt->body())) {
|
||||
if (auto* decl = s->As<ast::VariableDeclStatement>()) {
|
||||
if (!EmitVariable(program_->Sem().Get(decl->variable()), true)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
make_indent();
|
||||
out_ << "for(;;) {" << std::endl;
|
||||
increment_indent();
|
||||
|
||||
if (stmt->has_continuing()) {
|
||||
make_indent();
|
||||
out_ << "if (!" << guard << ") ";
|
||||
|
||||
if (!EmitBlockAndNewline(stmt->continuing())) {
|
||||
return false;
|
||||
}
|
||||
|
||||
make_indent();
|
||||
out_ << guard << " = false;" << std::endl;
|
||||
out_ << std::endl;
|
||||
}
|
||||
|
||||
for (auto* s : *(stmt->body())) {
|
||||
// If we have a continuing block we've already emitted the variable
|
||||
// declaration before the loop, so treat it as an assignment.
|
||||
auto* decl = s->As<ast::VariableDeclStatement>();
|
||||
if (decl != nullptr && stmt->has_continuing()) {
|
||||
auto emit_continuing = [this, stmt]() {
|
||||
if (stmt->has_continuing()) {
|
||||
make_indent();
|
||||
|
||||
auto* var = decl->variable();
|
||||
out_ << program_->Symbols().NameFor(var->symbol()) << " = ";
|
||||
if (var->constructor() != nullptr) {
|
||||
if (!EmitExpression(var->constructor())) {
|
||||
return false;
|
||||
}
|
||||
} else {
|
||||
auto* type = program_->Sem().Get(var)->Type()->UnwrapRef();
|
||||
if (!EmitZeroValue(type)) {
|
||||
return false;
|
||||
}
|
||||
if (!EmitBlock(stmt->continuing())) {
|
||||
return false;
|
||||
}
|
||||
out_ << ";" << std::endl;
|
||||
continue;
|
||||
out_ << std::endl;
|
||||
}
|
||||
return true;
|
||||
};
|
||||
|
||||
if (!EmitStatement(s)) {
|
||||
return false;
|
||||
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
|
||||
bool ok = EmitBlockBraces("while (true)", [&] {
|
||||
for (auto* s : stmt->body()->statements()) {
|
||||
if (!EmitStatement(s)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
decrement_indent();
|
||||
make_indent();
|
||||
out_ << "}" << std::endl;
|
||||
|
||||
// Close the scope for any continuing variables.
|
||||
if (stmt->has_continuing()) {
|
||||
decrement_indent();
|
||||
make_indent();
|
||||
out_ << "}" << std::endl;
|
||||
}
|
||||
|
||||
return true;
|
||||
return emit_continuing();
|
||||
});
|
||||
out_ << std::endl;
|
||||
return ok;
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitDiscard(ast::DiscardStatement*) {
|
||||
|
|
@ -1530,7 +1473,7 @@ bool GeneratorImpl::EmitStatement(ast::Statement* stmt) {
|
|||
}
|
||||
if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
|
||||
auto* var = program_->Sem().Get(v->variable());
|
||||
return EmitVariable(var, false);
|
||||
return EmitVariable(var);
|
||||
}
|
||||
|
||||
diagnostics_.add_error("unknown statement type: " + program_->str(stmt));
|
||||
|
|
@ -1882,8 +1825,7 @@ bool GeneratorImpl::EmitUnaryOp(ast::UnaryOpExpression* expr) {
|
|||
return true;
|
||||
}
|
||||
|
||||
bool GeneratorImpl::EmitVariable(const sem::Variable* var,
|
||||
bool skip_constructor) {
|
||||
bool GeneratorImpl::EmitVariable(const sem::Variable* var) {
|
||||
make_indent();
|
||||
|
||||
auto* decl = var->Declaration();
|
||||
|
|
@ -1925,19 +1867,17 @@ bool GeneratorImpl::EmitVariable(const sem::Variable* var,
|
|||
out_ << " " << name;
|
||||
}
|
||||
|
||||
if (!skip_constructor) {
|
||||
if (decl->constructor() != nullptr) {
|
||||
out_ << " = ";
|
||||
if (!EmitExpression(decl->constructor())) {
|
||||
return false;
|
||||
}
|
||||
} else if (var->StorageClass() == ast::StorageClass::kPrivate ||
|
||||
var->StorageClass() == ast::StorageClass::kFunction ||
|
||||
var->StorageClass() == ast::StorageClass::kNone) {
|
||||
out_ << " = ";
|
||||
if (!EmitZeroValue(type)) {
|
||||
return false;
|
||||
}
|
||||
if (decl->constructor() != nullptr) {
|
||||
out_ << " = ";
|
||||
if (!EmitExpression(decl->constructor())) {
|
||||
return false;
|
||||
}
|
||||
} else if (var->StorageClass() == ast::StorageClass::kPrivate ||
|
||||
var->StorageClass() == ast::StorageClass::kFunction ||
|
||||
var->StorageClass() == ast::StorageClass::kNone) {
|
||||
out_ << " = ";
|
||||
if (!EmitZeroValue(type)) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
out_ << ";" << std::endl;
|
||||
|
|
@ -2046,6 +1986,21 @@ GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(
|
|||
return {};
|
||||
}
|
||||
|
||||
template <typename F>
|
||||
bool GeneratorImpl::EmitBlockBraces(const std::string& prefix, F&& cb) {
|
||||
out_ << prefix << (prefix.empty() ? "{" : " {") << std::endl;
|
||||
increment_indent();
|
||||
|
||||
if (!cb()) {
|
||||
return false;
|
||||
}
|
||||
|
||||
decrement_indent();
|
||||
make_indent();
|
||||
out_ << "}";
|
||||
return true;
|
||||
}
|
||||
|
||||
} // namespace msl
|
||||
} // namespace writer
|
||||
} // namespace tint
|
||||
|
|
|
|||
|
|
@ -209,9 +209,8 @@ class GeneratorImpl : public TextGenerator {
|
|||
bool EmitUnaryOp(ast::UnaryOpExpression* expr);
|
||||
/// Handles generating a variable
|
||||
/// @param var the variable to generate
|
||||
/// @param skip_constructor set true if the constructor should be skipped
|
||||
/// @returns true if the variable was emitted
|
||||
bool EmitVariable(const sem::Variable* var, bool skip_constructor);
|
||||
bool EmitVariable(const sem::Variable* var);
|
||||
/// Handles generating a program scope constant variable
|
||||
/// @param var the variable to emit
|
||||
/// @returns true if the variable was emitted
|
||||
|
|
@ -260,8 +259,17 @@ class GeneratorImpl : public TextGenerator {
|
|||
/// type.
|
||||
SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
|
||||
|
||||
/// Emits `prefix`, followed by an opening brace `{`, then calls `cb` to emit
|
||||
/// the block body, then finally emits the closing brace `}`.
|
||||
/// @param prefix the string to emit before the opening brace
|
||||
/// @param cb a function or function-like object with the signature `bool()`
|
||||
/// that emits the block body.
|
||||
/// @returns the return value of `cb`.
|
||||
template <typename F>
|
||||
bool EmitBlockBraces(const std::string& prefix, F&& cb);
|
||||
|
||||
const Program* program_ = nullptr;
|
||||
uint32_t loop_emission_counter_ = 0;
|
||||
std::function<bool()> emit_continuing_;
|
||||
};
|
||||
|
||||
} // namespace msl
|
||||
|
|
|
|||
|
|
@ -22,15 +22,18 @@ namespace {
|
|||
using MslGeneratorImplTest = TestHelper;
|
||||
|
||||
TEST_F(MslGeneratorImplTest, Emit_Continue) {
|
||||
auto* c = create<ast::ContinueStatement>();
|
||||
WrapInFunction(Loop(Block(c)));
|
||||
auto* loop = Loop(Block(create<ast::ContinueStatement>()));
|
||||
WrapInFunction(loop);
|
||||
|
||||
GeneratorImpl& gen = Build();
|
||||
|
||||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(c)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), " continue;\n");
|
||||
ASSERT_TRUE(gen.EmitStatement(loop)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
continue;
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
|
|
|||
|
|
@ -33,7 +33,7 @@ TEST_F(MslGeneratorImplTest, Emit_Loop) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(l)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( for(;;) {
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
discard_fragment();
|
||||
}
|
||||
)");
|
||||
|
|
@ -50,15 +50,10 @@ TEST_F(MslGeneratorImplTest, Emit_LoopWithContinuing) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(l)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
bool tint_msl_is_first_1 = true;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
return;
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
discard_fragment();
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
discard_fragment();
|
||||
{
|
||||
return;
|
||||
}
|
||||
}
|
||||
)");
|
||||
|
|
@ -84,26 +79,16 @@ TEST_F(MslGeneratorImplTest, Emit_LoopNestedWithContinuing) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(outer)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
bool tint_msl_is_first_1 = true;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
lhs = rhs;
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
while (true) {
|
||||
discard_fragment();
|
||||
{
|
||||
bool tint_msl_is_first_2 = true;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_2) {
|
||||
return;
|
||||
}
|
||||
tint_msl_is_first_2 = false;
|
||||
|
||||
discard_fragment();
|
||||
}
|
||||
return;
|
||||
}
|
||||
}
|
||||
{
|
||||
lhs = rhs;
|
||||
}
|
||||
}
|
||||
)");
|
||||
}
|
||||
|
|
@ -146,18 +131,11 @@ TEST_F(MslGeneratorImplTest, Emit_LoopWithVarUsedInContinuing) {
|
|||
gen.increment_indent();
|
||||
|
||||
ASSERT_TRUE(gen.EmitStatement(outer)) << gen.error();
|
||||
EXPECT_EQ(gen.result(), R"( {
|
||||
bool tint_msl_is_first_1 = true;
|
||||
float lhs;
|
||||
float other;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
lhs = rhs;
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
lhs = 2.400000095f;
|
||||
other = 0.0f;
|
||||
EXPECT_EQ(gen.result(), R"( while (true) {
|
||||
float lhs = 2.400000095f;
|
||||
float other = 0.0f;
|
||||
{
|
||||
lhs = rhs;
|
||||
}
|
||||
}
|
||||
)");
|
||||
|
|
|
|||
|
|
@ -1,18 +1,34 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct tint_array_wrapper {
|
||||
/* 0x0000 */ uint arr[50];
|
||||
};
|
||||
struct Buf {
|
||||
/* 0x0000 */ uint count;
|
||||
/* 0x0004 */ tint_array_wrapper data;
|
||||
};
|
||||
|
||||
kernel void tint_symbol(device Buf& b [[buffer(0)]]) {
|
||||
uint i = 0u;
|
||||
while (true) {
|
||||
if ((i >= b.count)) {
|
||||
break;
|
||||
}
|
||||
uint const p_save = i;
|
||||
if (((i % 2u) == 0u)) {
|
||||
{
|
||||
b.data.arr[p_save] = (b.data.arr[p_save] * 2u);
|
||||
i = (i + 1u);
|
||||
}
|
||||
continue;
|
||||
}
|
||||
b.data.arr[p_save] = 0u;
|
||||
{
|
||||
b.data.arr[p_save] = (b.data.arr[p_save] * 2u);
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:16:24: error: default initialization of an object of const type 'device uint *const' (aka 'device unsigned int *const')
|
||||
device uint* const p;
|
||||
^
|
||||
= nullptr
|
||||
program_source:27:9: error: cannot assign to variable 'p' with const-qualified type 'device uint *const' (aka 'device unsigned int *const')
|
||||
p = &(b.data.array[i]);
|
||||
~ ^
|
||||
program_source:16:24: note: variable 'p' declared const here
|
||||
device uint* const p;
|
||||
~~~~~~~~~~~~~~~~~~~^
|
||||
|
|
|
|||
|
|
@ -1,28 +1,36 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ packed_uint2 aShape;
|
||||
/* 0x0008 */ packed_uint2 bShape;
|
||||
/* 0x0010 */ packed_uint2 outShape;
|
||||
};
|
||||
struct Matrix {
|
||||
/* 0x0000 */ uint numbers[1];
|
||||
};
|
||||
|
||||
kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
|
||||
uint2 const resultCell = uint2(global_id.y, global_id.x);
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
uint const index = (resultCell.y + (resultCell.x * dimOutter));
|
||||
resultMatrix.numbers[index] = result;
|
||||
return;
|
||||
}
|
||||
|
||||
Validation Failure:
|
||||
|
||||
Compilation failed:
|
||||
|
||||
program_source:22:18: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const a;
|
||||
^
|
||||
= 0
|
||||
program_source:23:18: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const b;
|
||||
^
|
||||
= 0
|
||||
program_source:33:11: error: cannot assign to variable 'a' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
a = (i + (resultCell.x * dimInner));
|
||||
~ ^
|
||||
program_source:22:18: note: variable 'a' declared const here
|
||||
uint const a;
|
||||
~~~~~~~~~~~^
|
||||
program_source:34:11: error: cannot assign to variable 'b' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
b = (resultCell.y + (i * dimOutter));
|
||||
~ ^
|
||||
program_source:23:18: note: variable 'b' declared const here
|
||||
uint const b;
|
||||
~~~~~~~~~~~^
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
|
|
@ -14,18 +14,13 @@ kernel void tint_symbol(texture2d_array<float, access::sample> tint_symbol_2 [[t
|
|||
float4 texel = tint_symbol_2.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
|
||||
{
|
||||
uint i = 0u;
|
||||
{
|
||||
bool tint_msl_is_first_1 = true;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
i = (i + 1u);
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
if (!((i < 1u))) {
|
||||
break;
|
||||
}
|
||||
result.values[(flatIndex + i)] = texel.r;
|
||||
while (true) {
|
||||
if (!((i < 1u))) {
|
||||
break;
|
||||
}
|
||||
result.values[(flatIndex + i)] = texel.r;
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,78 +1,224 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
using namespace metal;
|
||||
struct Uniforms {
|
||||
/* 0x0000 */ uint dimAOuter;
|
||||
/* 0x0004 */ uint dimInner;
|
||||
/* 0x0008 */ uint dimBOuter;
|
||||
};
|
||||
struct Matrix {
|
||||
/* 0x0000 */ float numbers[1];
|
||||
};
|
||||
struct tint_array_wrapper_1 {
|
||||
float arr[64];
|
||||
};
|
||||
struct tint_array_wrapper {
|
||||
tint_array_wrapper_1 arr[64];
|
||||
};
|
||||
struct tint_array_wrapper_2 {
|
||||
float arr[16];
|
||||
};
|
||||
struct tint_array_wrapper_3 {
|
||||
float arr[4];
|
||||
};
|
||||
|
||||
constant uint RowPerThread = 4u;
|
||||
constant uint ColPerThread = 4u;
|
||||
constant uint TileAOuter = 64u;
|
||||
constant uint TileBOuter = 64u;
|
||||
constant uint TileInner = 64u;
|
||||
float mm_readA(constant Uniforms& uniforms, const device Matrix& firstMatrix, uint row, uint col) {
|
||||
if (((row < uniforms.dimAOuter) && (col < uniforms.dimInner))) {
|
||||
float const result = firstMatrix.numbers[((row * uniforms.dimInner) + col)];
|
||||
return result;
|
||||
}
|
||||
return 0.0f;
|
||||
}
|
||||
|
||||
Validation Failure:
|
||||
float mm_readB(constant Uniforms& uniforms, const device Matrix& secondMatrix, uint row, uint col) {
|
||||
if (((row < uniforms.dimInner) && (col < uniforms.dimBOuter))) {
|
||||
float const result = secondMatrix.numbers[((row * uniforms.dimBOuter) + col)];
|
||||
return result;
|
||||
}
|
||||
return 0.0f;
|
||||
}
|
||||
|
||||
Compilation failed:
|
||||
void mm_write(constant Uniforms& uniforms, device Matrix& resultMatrix, uint row, uint col, float value) {
|
||||
if (((row < uniforms.dimAOuter) && (col < uniforms.dimBOuter))) {
|
||||
uint const index = (col + (row * uniforms.dimBOuter));
|
||||
resultMatrix.numbers[index] = value;
|
||||
}
|
||||
}
|
||||
|
||||
program_source:56:31: warning: equality comparison with extraneous parentheses
|
||||
kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
|
||||
threadgroup tint_array_wrapper tint_symbol_4;
|
||||
threadgroup tint_array_wrapper tint_symbol_5;
|
||||
if ((local_invocation_index == 0u)) {
|
||||
~~~~~~~~~~~~~~~~~~~~~~~^~~~~
|
||||
program_source:56:31: note: remove extraneous parentheses around the comparison to silence this warning
|
||||
if ((local_invocation_index == 0u)) {
|
||||
~ ^ ~
|
||||
program_source:56:31: note: use '=' to turn this equality comparison into an assignment
|
||||
if ((local_invocation_index == 0u)) {
|
||||
^~
|
||||
=
|
||||
program_source:122:30: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const inputRow;
|
||||
^
|
||||
= 0
|
||||
program_source:123:30: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const inputCol;
|
||||
^
|
||||
= 0
|
||||
program_source:133:30: error: cannot assign to variable 'inputRow' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
inputRow = (tileRow + innerRow);
|
||||
~~~~~~~~ ^
|
||||
program_source:122:30: note: variable 'inputRow' declared const here
|
||||
uint const inputRow;
|
||||
~~~~~~~~~~~^~~~~~~~
|
||||
program_source:134:30: error: cannot assign to variable 'inputCol' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
inputCol = (tileColA + innerCol);
|
||||
~~~~~~~~ ^
|
||||
program_source:123:30: note: variable 'inputCol' declared const here
|
||||
uint const inputCol;
|
||||
~~~~~~~~~~~^~~~~~~~
|
||||
program_source:159:30: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const inputRow;
|
||||
^
|
||||
= 0
|
||||
program_source:160:30: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const inputCol;
|
||||
^
|
||||
= 0
|
||||
program_source:170:30: error: cannot assign to variable 'inputRow' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
inputRow = (tileRowB + innerRow);
|
||||
~~~~~~~~ ^
|
||||
program_source:159:30: note: variable 'inputRow' declared const here
|
||||
uint const inputRow;
|
||||
~~~~~~~~~~~^~~~~~~~
|
||||
program_source:171:30: error: cannot assign to variable 'inputCol' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
inputCol = (tileCol + innerCol);
|
||||
~~~~~~~~ ^
|
||||
program_source:160:30: note: variable 'inputCol' declared const here
|
||||
uint const inputCol;
|
||||
~~~~~~~~~~~^~~~~~~~
|
||||
program_source:228:36: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const index;
|
||||
^
|
||||
= 0
|
||||
program_source:238:33: error: cannot assign to variable 'index' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
index = ((innerRow * ColPerThread) + innerCol);
|
||||
~~~~~ ^
|
||||
program_source:228:36: note: variable 'index' declared const here
|
||||
uint const index;
|
||||
~~~~~~~~~~~^~~~~
|
||||
program_source:270:24: error: default initialization of an object of const type 'const uint' (aka 'const unsigned int')
|
||||
uint const index;
|
||||
^
|
||||
= 0
|
||||
program_source:280:21: error: cannot assign to variable 'index' with const-qualified type 'const uint' (aka 'const unsigned int')
|
||||
index = ((innerRow * ColPerThread) + innerCol);
|
||||
~~~~~ ^
|
||||
program_source:270:24: note: variable 'index' declared const here
|
||||
uint const index;
|
||||
~~~~~~~~~~~^~~~~
|
||||
tint_array_wrapper const tint_symbol_2 = {.arr={}};
|
||||
tint_symbol_4 = tint_symbol_2;
|
||||
tint_array_wrapper const tint_symbol_3 = {.arr={}};
|
||||
tint_symbol_5 = tint_symbol_3;
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
uint const tileRow = (local_id.y * RowPerThread);
|
||||
uint const tileCol = (local_id.x * ColPerThread);
|
||||
uint const globalRow = (global_id.y * RowPerThread);
|
||||
uint const globalCol = (global_id.x * ColPerThread);
|
||||
uint const numTiles = (((uniforms.dimInner - 1u) / TileInner) + 1u);
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
uint innerRow = 0u;
|
||||
while (true) {
|
||||
if (!((innerRow < RowPerThread))) {
|
||||
break;
|
||||
}
|
||||
{
|
||||
uint innerCol = 0u;
|
||||
while (true) {
|
||||
if (!((innerCol < ColPerThread))) {
|
||||
break;
|
||||
}
|
||||
uint const index = ((innerRow * ColPerThread) + innerCol);
|
||||
mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
|
||||
{
|
||||
innerCol = (innerCol + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
{
|
||||
innerRow = (innerRow + 1u);
|
||||
}
|
||||
}
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -1,37 +1,33 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
|
||||
[[block]]
|
||||
using namespace metal;
|
||||
struct tint_symbol_2 {
|
||||
/* 0x0000 */ uint buffer_size[2];
|
||||
};
|
||||
struct SB_RO {
|
||||
arg_0 : array<i32>;
|
||||
/* 0x0000 */ int arg_0[1];
|
||||
};
|
||||
struct tint_symbol {
|
||||
float4 value [[position]];
|
||||
};
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> sb_ro : SB_RO;
|
||||
|
||||
fn arrayLength_1588cd() {
|
||||
var res : u32 = arrayLength(&(sb_ro.arg_0));
|
||||
void arrayLength_1588cd(constant tint_symbol_2& tint_symbol_3) {
|
||||
uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
[[builtin(position)]]
|
||||
value : vec4<f32>;
|
||||
};
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn vertex_main() -> tint_symbol {
|
||||
arrayLength_1588cd();
|
||||
let tint_symbol_1 : tint_symbol = tint_symbol(vec4<f32>());
|
||||
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_1588cd(tint_symbol_3);
|
||||
tint_symbol const tint_symbol_1 = {.value=float4()};
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn fragment_main() {
|
||||
arrayLength_1588cd();
|
||||
fragment void fragment_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_1588cd(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute)]]
|
||||
fn compute_main() {
|
||||
arrayLength_1588cd();
|
||||
kernel void compute_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_1588cd(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
Failed to generate: error: Unknown import method: arrayLength
|
||||
|
|
|
|||
|
|
@ -1,37 +1,33 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
|
||||
[[block]]
|
||||
using namespace metal;
|
||||
struct tint_symbol_2 {
|
||||
/* 0x0000 */ uint buffer_size[1];
|
||||
};
|
||||
struct SB_RW {
|
||||
arg_0 : array<i32>;
|
||||
/* 0x0000 */ int arg_0[1];
|
||||
};
|
||||
struct tint_symbol {
|
||||
float4 value [[position]];
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
|
||||
|
||||
fn arrayLength_61b1c7() {
|
||||
var res : u32 = arrayLength(&(sb_rw.arg_0));
|
||||
void arrayLength_61b1c7(constant tint_symbol_2& tint_symbol_3) {
|
||||
uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
[[builtin(position)]]
|
||||
value : vec4<f32>;
|
||||
};
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn vertex_main() -> tint_symbol {
|
||||
arrayLength_61b1c7();
|
||||
let tint_symbol_1 : tint_symbol = tint_symbol(vec4<f32>());
|
||||
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_61b1c7(tint_symbol_3);
|
||||
tint_symbol const tint_symbol_1 = {.value=float4()};
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn fragment_main() {
|
||||
arrayLength_61b1c7();
|
||||
fragment void fragment_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_61b1c7(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute)]]
|
||||
fn compute_main() {
|
||||
arrayLength_61b1c7();
|
||||
kernel void compute_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_61b1c7(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
Failed to generate: error: Unknown import method: arrayLength
|
||||
|
|
|
|||
|
|
@ -1,37 +1,33 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
|
||||
[[block]]
|
||||
using namespace metal;
|
||||
struct tint_symbol_2 {
|
||||
/* 0x0000 */ uint buffer_size[2];
|
||||
};
|
||||
struct SB_RO {
|
||||
arg_0 : array<f32>;
|
||||
/* 0x0000 */ float arg_0[1];
|
||||
};
|
||||
struct tint_symbol {
|
||||
float4 value [[position]];
|
||||
};
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> sb_ro : SB_RO;
|
||||
|
||||
fn arrayLength_a0f5ca() {
|
||||
var res : u32 = arrayLength(&(sb_ro.arg_0));
|
||||
void arrayLength_a0f5ca(constant tint_symbol_2& tint_symbol_3) {
|
||||
uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
[[builtin(position)]]
|
||||
value : vec4<f32>;
|
||||
};
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn vertex_main() -> tint_symbol {
|
||||
arrayLength_a0f5ca();
|
||||
let tint_symbol_1 : tint_symbol = tint_symbol(vec4<f32>());
|
||||
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_a0f5ca(tint_symbol_3);
|
||||
tint_symbol const tint_symbol_1 = {.value=float4()};
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn fragment_main() {
|
||||
arrayLength_a0f5ca();
|
||||
fragment void fragment_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_a0f5ca(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute)]]
|
||||
fn compute_main() {
|
||||
arrayLength_a0f5ca();
|
||||
kernel void compute_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_a0f5ca(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
Failed to generate: error: Unknown import method: arrayLength
|
||||
|
|
|
|||
|
|
@ -1,37 +1,33 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
|
||||
[[block]]
|
||||
using namespace metal;
|
||||
struct tint_symbol_2 {
|
||||
/* 0x0000 */ uint buffer_size[1];
|
||||
};
|
||||
struct SB_RW {
|
||||
arg_0 : array<f32>;
|
||||
/* 0x0000 */ float arg_0[1];
|
||||
};
|
||||
struct tint_symbol {
|
||||
float4 value [[position]];
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
|
||||
|
||||
fn arrayLength_cdd123() {
|
||||
var res : u32 = arrayLength(&(sb_rw.arg_0));
|
||||
void arrayLength_cdd123(constant tint_symbol_2& tint_symbol_3) {
|
||||
uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
[[builtin(position)]]
|
||||
value : vec4<f32>;
|
||||
};
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn vertex_main() -> tint_symbol {
|
||||
arrayLength_cdd123();
|
||||
let tint_symbol_1 : tint_symbol = tint_symbol(vec4<f32>());
|
||||
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_cdd123(tint_symbol_3);
|
||||
tint_symbol const tint_symbol_1 = {.value=float4()};
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn fragment_main() {
|
||||
arrayLength_cdd123();
|
||||
fragment void fragment_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_cdd123(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute)]]
|
||||
fn compute_main() {
|
||||
arrayLength_cdd123();
|
||||
kernel void compute_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_cdd123(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
Failed to generate: error: Unknown import method: arrayLength
|
||||
|
|
|
|||
|
|
@ -1,37 +1,33 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
|
||||
[[block]]
|
||||
using namespace metal;
|
||||
struct tint_symbol_2 {
|
||||
/* 0x0000 */ uint buffer_size[2];
|
||||
};
|
||||
struct SB_RO {
|
||||
arg_0 : array<u32>;
|
||||
/* 0x0000 */ uint arg_0[1];
|
||||
};
|
||||
struct tint_symbol {
|
||||
float4 value [[position]];
|
||||
};
|
||||
|
||||
[[group(0), binding(1)]] var<storage, read> sb_ro : SB_RO;
|
||||
|
||||
fn arrayLength_cfca0a() {
|
||||
var res : u32 = arrayLength(&(sb_ro.arg_0));
|
||||
void arrayLength_cfca0a(constant tint_symbol_2& tint_symbol_3) {
|
||||
uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
[[builtin(position)]]
|
||||
value : vec4<f32>;
|
||||
};
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn vertex_main() -> tint_symbol {
|
||||
arrayLength_cfca0a();
|
||||
let tint_symbol_1 : tint_symbol = tint_symbol(vec4<f32>());
|
||||
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_cfca0a(tint_symbol_3);
|
||||
tint_symbol const tint_symbol_1 = {.value=float4()};
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn fragment_main() {
|
||||
arrayLength_cfca0a();
|
||||
fragment void fragment_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_cfca0a(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute)]]
|
||||
fn compute_main() {
|
||||
arrayLength_cfca0a();
|
||||
kernel void compute_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_cfca0a(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
Failed to generate: error: Unknown import method: arrayLength
|
||||
|
|
|
|||
|
|
@ -1,37 +1,33 @@
|
|||
SKIP: FAILED
|
||||
#include <metal_stdlib>
|
||||
|
||||
|
||||
[[block]]
|
||||
using namespace metal;
|
||||
struct tint_symbol_2 {
|
||||
/* 0x0000 */ uint buffer_size[1];
|
||||
};
|
||||
struct SB_RW {
|
||||
arg_0 : array<u32>;
|
||||
/* 0x0000 */ uint arg_0[1];
|
||||
};
|
||||
struct tint_symbol {
|
||||
float4 value [[position]];
|
||||
};
|
||||
|
||||
[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
|
||||
|
||||
fn arrayLength_eb510f() {
|
||||
var res : u32 = arrayLength(&(sb_rw.arg_0));
|
||||
void arrayLength_eb510f(constant tint_symbol_2& tint_symbol_3) {
|
||||
uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
|
||||
}
|
||||
|
||||
struct tint_symbol {
|
||||
[[builtin(position)]]
|
||||
value : vec4<f32>;
|
||||
};
|
||||
|
||||
[[stage(vertex)]]
|
||||
fn vertex_main() -> tint_symbol {
|
||||
arrayLength_eb510f();
|
||||
let tint_symbol_1 : tint_symbol = tint_symbol(vec4<f32>());
|
||||
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_eb510f(tint_symbol_3);
|
||||
tint_symbol const tint_symbol_1 = {.value=float4()};
|
||||
return tint_symbol_1;
|
||||
}
|
||||
|
||||
[[stage(fragment)]]
|
||||
fn fragment_main() {
|
||||
arrayLength_eb510f();
|
||||
fragment void fragment_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_eb510f(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
[[stage(compute)]]
|
||||
fn compute_main() {
|
||||
arrayLength_eb510f();
|
||||
kernel void compute_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
|
||||
arrayLength_eb510f(tint_symbol_3);
|
||||
return;
|
||||
}
|
||||
|
||||
Failed to generate: error: Unknown import method: arrayLength
|
||||
|
|
|
|||
|
|
@ -63,33 +63,31 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
|
|||
float2 vel = 0.0f;
|
||||
{
|
||||
uint i = 0u;
|
||||
{
|
||||
bool tint_msl_is_first_1 = true;
|
||||
for(;;) {
|
||||
if (!tint_msl_is_first_1) {
|
||||
while (true) {
|
||||
if (!((i < 5u))) {
|
||||
break;
|
||||
}
|
||||
if ((i == index)) {
|
||||
{
|
||||
i = (i + 1u);
|
||||
}
|
||||
tint_msl_is_first_1 = false;
|
||||
|
||||
if (!((i < 5u))) {
|
||||
break;
|
||||
}
|
||||
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);
|
||||
}
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue