encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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 2021-06-21 08:49:27 +00:00 committed by Tint LUCI CQ
parent c15baf695d
commit 663271dca4
16 changed files with 2061 additions and 2168 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

111
src/writer/msl/generator_impl.cc
View File

@ -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_);
make_indent();
auto emit_continuing = [this, stmt]() {
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)) {
if (!EmitBlock(stmt->continuing())) {
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;
}
return true;
};
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()) {
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;
}
}
out_ << ";" << std::endl;
continue;
}
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,7 +1867,6 @@ bool GeneratorImpl::EmitVariable(const sem::Variable* var,
out_ << " " << name;
}
if (!skip_constructor) {
if (decl->constructor() != nullptr) {
out_ << " = ";
if (!EmitExpression(decl->constructor())) {
@ -1939,7 +1880,6 @@ bool GeneratorImpl::EmitVariable(const sem::Variable* var,
return false;
}
}
}
out_ << ";" << std::endl;
return true;
@ -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

14
src/writer/msl/generator_impl.h
View File

@ -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

11
src/writer/msl/generator_impl_continue_test.cc
View File

@ -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

48
src/writer/msl/generator_impl_loop_test.cc
View File

@ -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;
EXPECT_EQ(gen.result(), R"( while (true) {
discard_fragment();
{
return;
}
}
)");
@ -84,25 +79,15 @@ 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();
}
}
{
lhs = rhs;
}
}
)");
@ -146,19 +131,12 @@ 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) {
EXPECT_EQ(gen.result(), R"( while (true) {
float lhs = 2.400000095f;
float other = 0.0f;
{
lhs = rhs;
}
tint_msl_is_first_1 = false;
lhs = 2.400000095f;
other = 0.0f;
}
}
)");
}

46
test/bug/tint/221.wgsl.expected.msl
View File

@ -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;
~~~~~~~~~~~~~~~~~~~^

58
test/bug/tint/744.wgsl.expected.msl
View File

@ -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;
~~~~~~~~~~~^

3262
test/bug/tint/749.spvasm.expected.msl
View File

File diff suppressed because it is too large Load Diff

11
test/bug/tint/757.wgsl.expected.msl
View File

@ -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;
while (true) {
if (!((i < 1u))) {
break;
}
result.values[(flatIndex + i)] = texel.r;
{
i = (i + 1u);
}
}
}

290
test/bug/tint/914.wgsl.expected.msl
View File

@ -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;
}

44
test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl
View File

@ -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

44
test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl
View File

@ -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

44
test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl
View File

@ -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

44
test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl
View File

@ -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

44
test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl
View File

@ -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

44
test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl
View File

@ -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

14
test/samples/compute_boids.wgsl.expected.msl
View File

@ -63,18 +63,14 @@ 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) {
i = (i + 1u);
}
tint_msl_is_first_1 = false;
while (true) {
if (!((i < 5u))) {
break;
}
if ((i == index)) {
{
i = (i + 1u);
}
continue;
}
pos = particlesA.particles.arr[i].pos.xy;
@ -90,6 +86,8 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
cVel = (cVel + vel);
cVelCount = (cVelCount + 1);
}
{
i = (i + 1u);
}
}
}