blob: 5a1e8c0e8d2194d331a09bf5927698469718ede9 [file] [log] [blame]
// Copyright 2020 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/fluent_types.h"
#include "src/tint/lang/msl/writer/ast_printer/helper_test.h"
#include "src/tint/lang/msl/writer/writer.h"
#include "src/tint/lang/wgsl/ast/stage_attribute.h"
using namespace tint::core::number_suffixes; // NOLINT
using namespace tint::core::fluent_types; // NOLINT
namespace tint::msl::writer {
namespace {
using MslASTPrinterTest = TestHelper;
TEST_F(MslASTPrinterTest, InvalidProgram) {
Diagnostics().AddError(diag::System::Writer, Source{}) << "make the program invalid";
ASSERT_FALSE(IsValid());
auto program = resolver::Resolve(*this);
ASSERT_FALSE(program.IsValid());
auto result = Generate(program, Options{});
EXPECT_NE(result, Success);
EXPECT_EQ(result.Failure().reason.Str(), "error: make the program invalid");
}
TEST_F(MslASTPrinterTest, UnsupportedExtension) {
Enable(Source{{12, 34}}, wgsl::Extension::kChromiumExperimentalPushConstant);
ASTPrinter& gen = Build();
ASSERT_FALSE(gen.Generate());
EXPECT_EQ(
gen.Diagnostics().Str(),
R"(12:34 error: MSL backend does not support extension 'chromium_experimental_push_constant')");
}
TEST_F(MslASTPrinterTest, RequiresDirective) {
Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures);
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
)");
}
TEST_F(MslASTPrinterTest, Generate) {
Func("my_func", tint::Empty, ty.void_(), tint::Empty,
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
kernel void my_func() {
return;
}
)");
}
TEST_F(MslASTPrinterTest, HasInvariantAttribute_True) {
auto* out = Structure("Out", Vector{
Member("pos", ty.vec4<f32>(),
Vector{
Builtin(core::BuiltinValue::kPosition),
Invariant(),
}),
});
Func("vert_main", tint::Empty, ty.Of(out), Vector{Return(Call(ty.Of(out)))},
Vector{
Stage(ast::PipelineStage::kVertex),
});
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_TRUE(gen.HasInvariant());
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
#if __METAL_VERSION__ >= 210
#define TINT_INVARIANT [[invariant]]
#else
#define TINT_INVARIANT
#endif
struct Out {
float4 pos [[position]] TINT_INVARIANT;
};
vertex Out vert_main() {
return Out{};
}
)");
}
TEST_F(MslASTPrinterTest, HasInvariantAttribute_False) {
auto* out = Structure("Out", Vector{
Member("pos", ty.vec4<f32>(),
Vector{
Builtin(core::BuiltinValue::kPosition),
}),
});
Func("vert_main", tint::Empty, ty.Of(out), Vector{Return(Call(ty.Of(out)))},
Vector{
Stage(ast::PipelineStage::kVertex),
});
ASTPrinter& gen = Build();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_FALSE(gen.HasInvariant());
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
struct Out {
float4 pos [[position]];
};
vertex Out vert_main() {
return Out{};
}
)");
}
TEST_F(MslASTPrinterTest, WorkgroupMatrix) {
GlobalVar("m", ty.mat2x2<f32>(), core::AddressSpace::kWorkgroup);
Func("comp_main", tint::Empty, ty.void_(), Vector{Decl(Let("x", Expr("m")))},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
ASTPrinter& gen = SanitizeAndBuild();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
struct tint_symbol_4 {
float2x2 m;
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup float2x2* const tint_symbol) {
if ((local_idx < 1u)) {
*(tint_symbol) = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
void comp_main_inner(uint local_invocation_index, threadgroup float2x2* const tint_symbol_1) {
tint_zero_workgroup_memory(local_invocation_index, tint_symbol_1);
float2x2 const x = *(tint_symbol_1);
}
kernel void comp_main(threadgroup tint_symbol_4* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup float2x2* const tint_symbol_2 = &((*(tint_symbol_3)).m);
comp_main_inner(local_invocation_index, tint_symbol_2);
return;
}
)");
auto allocations = gen.DynamicWorkgroupAllocations();
ASSERT_TRUE(allocations.count("comp_main"));
ASSERT_EQ(allocations.at("comp_main").size(), 1u);
EXPECT_EQ(allocations.at("comp_main")[0], 2u * 2u * sizeof(float));
}
TEST_F(MslASTPrinterTest, WorkgroupMatrixInArray) {
GlobalVar("m", ty.array(ty.mat2x2<f32>(), 4_i), core::AddressSpace::kWorkgroup);
Func("comp_main", tint::Empty, ty.void_(), Vector{Decl(Let("x", Expr("m")))},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
ASTPrinter& gen = SanitizeAndBuild();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
template<typename T, size_t N>
struct tint_array {
const constant T& operator[](size_t i) const constant { return elements[i]; }
device T& operator[](size_t i) device { return elements[i]; }
const device T& operator[](size_t i) const device { return elements[i]; }
thread T& operator[](size_t i) thread { return elements[i]; }
const thread T& operator[](size_t i) const thread { return elements[i]; }
threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
T elements[N];
};
#define TINT_ISOLATE_UB(VOLATILE_NAME) \
volatile bool VOLATILE_NAME = true; \
if (VOLATILE_NAME)
struct tint_symbol_4 {
tint_array<float2x2, 4> m;
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
void comp_main_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol_1) {
tint_zero_workgroup_memory(local_invocation_index, tint_symbol_1);
tint_array<float2x2, 4> const x = *(tint_symbol_1);
}
kernel void comp_main(threadgroup tint_symbol_4* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup tint_array<float2x2, 4>* const tint_symbol_2 = &((*(tint_symbol_3)).m);
comp_main_inner(local_invocation_index, tint_symbol_2);
return;
}
)");
auto allocations = gen.DynamicWorkgroupAllocations();
ASSERT_TRUE(allocations.count("comp_main"));
ASSERT_EQ(allocations.at("comp_main").size(), 1u);
EXPECT_EQ(allocations.at("comp_main")[0], 4u * 2u * 2u * sizeof(float));
}
TEST_F(MslASTPrinterTest, WorkgroupMatrixInStruct) {
Structure("S1", Vector{
Member("m1", ty.mat2x2<f32>()),
Member("m2", ty.mat4x4<f32>()),
});
Structure("S2", Vector{
Member("s", ty("S1")),
});
GlobalVar("s", ty("S2"), core::AddressSpace::kWorkgroup);
Func("comp_main", tint::Empty, ty.void_(), Vector{Decl(Let("x", Expr("s")))},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
ASTPrinter& gen = SanitizeAndBuild();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
struct S1 {
float2x2 m1;
float4x4 m2;
};
struct S2 {
S1 s;
};
struct tint_symbol_5 {
S2 s;
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup S2* const tint_symbol_1) {
if ((local_idx < 1u)) {
S2 const tint_symbol = S2{};
*(tint_symbol_1) = tint_symbol;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
void comp_main_inner(uint local_invocation_index, threadgroup S2* const tint_symbol_2) {
tint_zero_workgroup_memory(local_invocation_index, tint_symbol_2);
S2 const x = *(tint_symbol_2);
}
kernel void comp_main(threadgroup tint_symbol_5* tint_symbol_4 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup S2* const tint_symbol_3 = &((*(tint_symbol_4)).s);
comp_main_inner(local_invocation_index, tint_symbol_3);
return;
}
)");
auto allocations = gen.DynamicWorkgroupAllocations();
ASSERT_TRUE(allocations.count("comp_main"));
ASSERT_EQ(allocations.at("comp_main").size(), 1u);
EXPECT_EQ(allocations.at("comp_main")[0], (2 * 2 * sizeof(float)) + (4u * 4u * sizeof(float)));
}
TEST_F(MslASTPrinterTest, WorkgroupMatrix_Multiples) {
GlobalVar("m1", ty.mat2x2<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m2", ty.mat2x3<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m3", ty.mat2x4<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m4", ty.mat3x2<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m5", ty.mat3x3<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m6", ty.mat3x4<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m7", ty.mat4x2<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m8", ty.mat4x3<f32>(), core::AddressSpace::kWorkgroup);
GlobalVar("m9", ty.mat4x4<f32>(), core::AddressSpace::kWorkgroup);
Func("main1", tint::Empty, ty.void_(),
Vector{
Decl(Let("a1", Expr("m1"))),
Decl(Let("a2", Expr("m2"))),
Decl(Let("a3", Expr("m3"))),
},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
Func("main2", tint::Empty, ty.void_(),
Vector{
Decl(Let("a1", Expr("m4"))),
Decl(Let("a2", Expr("m5"))),
Decl(Let("a3", Expr("m6"))),
},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
Func("main3", tint::Empty, ty.void_(),
Vector{
Decl(Let("a1", Expr("m7"))),
Decl(Let("a2", Expr("m8"))),
Decl(Let("a3", Expr("m9"))),
},
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
Func("main4_no_usages", tint::Empty, ty.void_(), tint::Empty,
Vector{
Stage(ast::PipelineStage::kCompute),
WorkgroupSize(1_i),
});
ASTPrinter& gen = SanitizeAndBuild();
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
using namespace metal;
struct tint_symbol_16 {
float2x2 m1;
float2x3 m2;
float2x4 m3;
};
struct tint_symbol_24 {
float3x2 m4;
float3x3 m5;
float3x4 m6;
};
struct tint_symbol_32 {
float4x2 m7;
float4x3 m8;
float4x4 m9;
};
void tint_zero_workgroup_memory(uint local_idx, threadgroup float2x2* const tint_symbol, threadgroup float2x3* const tint_symbol_1, threadgroup float2x4* const tint_symbol_2) {
if ((local_idx < 1u)) {
*(tint_symbol) = float2x2(float2(0.0f), float2(0.0f));
*(tint_symbol_1) = float2x3(float3(0.0f), float3(0.0f));
*(tint_symbol_2) = float2x4(float4(0.0f), float4(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
void tint_zero_workgroup_memory_1(uint local_idx_1, threadgroup float3x2* const tint_symbol_3, threadgroup float3x3* const tint_symbol_4, threadgroup float3x4* const tint_symbol_5) {
if ((local_idx_1 < 1u)) {
*(tint_symbol_3) = float3x2(float2(0.0f), float2(0.0f), float2(0.0f));
*(tint_symbol_4) = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
*(tint_symbol_5) = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
void tint_zero_workgroup_memory_2(uint local_idx_2, threadgroup float4x2* const tint_symbol_6, threadgroup float4x3* const tint_symbol_7, threadgroup float4x4* const tint_symbol_8) {
if ((local_idx_2 < 1u)) {
*(tint_symbol_6) = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
*(tint_symbol_7) = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
*(tint_symbol_8) = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
void main1_inner(uint local_invocation_index, threadgroup float2x2* const tint_symbol_9, threadgroup float2x3* const tint_symbol_10, threadgroup float2x4* const tint_symbol_11) {
tint_zero_workgroup_memory(local_invocation_index, tint_symbol_9, tint_symbol_10, tint_symbol_11);
float2x2 const a1 = *(tint_symbol_9);
float2x3 const a2 = *(tint_symbol_10);
float2x4 const a3 = *(tint_symbol_11);
}
kernel void main1(threadgroup tint_symbol_16* tint_symbol_13 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
threadgroup float2x2* const tint_symbol_12 = &((*(tint_symbol_13)).m1);
threadgroup float2x3* const tint_symbol_14 = &((*(tint_symbol_13)).m2);
threadgroup float2x4* const tint_symbol_15 = &((*(tint_symbol_13)).m3);
main1_inner(local_invocation_index, tint_symbol_12, tint_symbol_14, tint_symbol_15);
return;
}
void main2_inner(uint local_invocation_index_1, threadgroup float3x2* const tint_symbol_17, threadgroup float3x3* const tint_symbol_18, threadgroup float3x4* const tint_symbol_19) {
tint_zero_workgroup_memory_1(local_invocation_index_1, tint_symbol_17, tint_symbol_18, tint_symbol_19);
float3x2 const a1 = *(tint_symbol_17);
float3x3 const a2 = *(tint_symbol_18);
float3x4 const a3 = *(tint_symbol_19);
}
kernel void main2(threadgroup tint_symbol_24* tint_symbol_21 [[threadgroup(0)]], uint local_invocation_index_1 [[thread_index_in_threadgroup]]) {
threadgroup float3x2* const tint_symbol_20 = &((*(tint_symbol_21)).m4);
threadgroup float3x3* const tint_symbol_22 = &((*(tint_symbol_21)).m5);
threadgroup float3x4* const tint_symbol_23 = &((*(tint_symbol_21)).m6);
main2_inner(local_invocation_index_1, tint_symbol_20, tint_symbol_22, tint_symbol_23);
return;
}
void main3_inner(uint local_invocation_index_2, threadgroup float4x2* const tint_symbol_25, threadgroup float4x3* const tint_symbol_26, threadgroup float4x4* const tint_symbol_27) {
tint_zero_workgroup_memory_2(local_invocation_index_2, tint_symbol_25, tint_symbol_26, tint_symbol_27);
float4x2 const a1 = *(tint_symbol_25);
float4x3 const a2 = *(tint_symbol_26);
float4x4 const a3 = *(tint_symbol_27);
}
kernel void main3(threadgroup tint_symbol_32* tint_symbol_29 [[threadgroup(0)]], uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
threadgroup float4x2* const tint_symbol_28 = &((*(tint_symbol_29)).m7);
threadgroup float4x3* const tint_symbol_30 = &((*(tint_symbol_29)).m8);
threadgroup float4x4* const tint_symbol_31 = &((*(tint_symbol_29)).m9);
main3_inner(local_invocation_index_2, tint_symbol_28, tint_symbol_30, tint_symbol_31);
return;
}
kernel void main4_no_usages() {
return;
}
)");
auto allocations = gen.DynamicWorkgroupAllocations();
ASSERT_TRUE(allocations.count("main1"));
ASSERT_TRUE(allocations.count("main2"));
ASSERT_TRUE(allocations.count("main3"));
ASSERT_EQ(allocations.at("main1").size(), 1u);
EXPECT_EQ(allocations.at("main1")[0], 20u * sizeof(float));
ASSERT_EQ(allocations.at("main2").size(), 1u);
EXPECT_EQ(allocations.at("main2")[0], 32u * sizeof(float));
ASSERT_EQ(allocations.at("main3").size(), 1u);
EXPECT_EQ(allocations.at("main3")[0], 40u * sizeof(float));
EXPECT_EQ(allocations.at("main4_no_usages").size(), 0u);
}
} // namespace
} // namespace tint::msl::writer